git.sesse.net Git - ffmpeg/blob - libavfilter/unsharp_opencl_kernel.h

   1 /*
   2  * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
   3  * Copyright (C) 2013 Lenny Wang
   4  *
   5  * This file is part of FFmpeg.
   6  *
   7  * FFmpeg is free software; you can redistribute it and/or
   8  * modify it under the terms of the GNU Lesser General Public
   9  * License as published by the Free Software Foundation; either
  10  * version 2.1 of the License, or (at your option) any later version.
  11  *
  12  * FFmpeg is distributed in the hope that it will be useful,
  13  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  14  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  15  * Lesser General Public License for more details.
  16  *
  17  * You should have received a copy of the GNU Lesser General Public
  18  * License along with FFmpeg; if not, write to the Free Software
  19  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  20  */
  21
  22 #ifndef AVFILTER_UNSHARP_OPENCL_KERNEL_H
  23 #define AVFILTER_UNSHARP_OPENCL_KERNEL_H
  24
  25 #include "libavutil/opencl.h"
  26
  27 const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
  28 inline unsigned char clip_uint8(int a)
  29 {
  30     if (a & (~0xFF))
  31         return (-a)>>31;
  32     else
  33         return a;
  34 }
  35
  36 kernel void unsharp_luma(
  37                     global unsigned char *src,
  38                     global unsigned char *dst,
  39                     global int *mask,
  40                     int amount,
  41                     int scalebits,
  42                     int halfscale,
  43                     int src_stride,
  44                     int dst_stride,
  45                     int width,
  46                     int height)
  47 {
  48     int2 threadIdx, blockIdx, globalIdx;
  49     threadIdx.x = get_local_id(0);
  50     threadIdx.y = get_local_id(1);
  51     blockIdx.x = get_group_id(0);
  52     blockIdx.y = get_group_id(1);
  53     globalIdx.x = get_global_id(0);
  54     globalIdx.y = get_global_id(1);
  55
  56     if (!amount) {
  57         if (globalIdx.x < width && globalIdx.y < height)
  58             dst[globalIdx.x + globalIdx.y*dst_stride] = src[globalIdx.x + globalIdx.y*src_stride];
  59         return;
  60     }
  61
  62     local uchar l[32][32];
  63     local int lc[LU_RADIUS_X*LU_RADIUS_Y];
  64     int indexIx, indexIy, i, j;
  65
  66     for(i = 0; i <= 1; i++) {
  67         indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
  68         indexIy = indexIy < 0 ? 0 : indexIy;
  69         indexIy = indexIy >= height ? height - 1: indexIy;
  70         for(j = 0; j <= 1; j++) {
  71             indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
  72             indexIx = indexIx < 0 ? 0 : indexIx;
  73             indexIx = indexIx >= width ? width - 1: indexIx;
  74             l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy*src_stride + indexIx];
  75         }
  76     }
  77
  78     int indexL = threadIdx.y*16 + threadIdx.x;
  79     if (indexL < LU_RADIUS_X*LU_RADIUS_Y)
  80         lc[indexL] = mask[indexL];
  81     barrier(CLK_LOCAL_MEM_FENCE);
  82
  83     int idx, idy, maskIndex;
  84     int sum = 0;
  85     int steps_x = LU_RADIUS_X/2;
  86     int steps_y = LU_RADIUS_Y/2;
  87
  88     \n#pragma unroll\n
  89     for (i = -steps_y; i <= steps_y; i++) {
  90         idy = 8 + i + threadIdx.y;
  91         \n#pragma unroll\n
  92         for (j = -steps_x; j <= steps_x; j++) {
  93             idx = 8 + j + threadIdx.x;
  94             maskIndex = (i + steps_y)*LU_RADIUS_X + j + steps_x;
  95             sum += (int)l[idy][idx] * lc[maskIndex];
  96         }
  97     }
  98     int temp = (int)l[threadIdx.y + 8][threadIdx.x + 8];
  99     int res = temp + (((temp - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
 100     if (globalIdx.x < width && globalIdx.y < height)
 101         dst[globalIdx.x + globalIdx.y*dst_stride] = clip_uint8(res);
 102 }
 103
 104 kernel void unsharp_chroma(
 105                     global unsigned char *src_y,
 106                     global unsigned char *dst_y,
 107                     global int *mask,
 108                     int amount,
 109                     int scalebits,
 110                     int halfscale,
 111                     int src_stride_lu,
 112                     int src_stride_ch,
 113                     int dst_stride_lu,
 114                     int dst_stride_ch,
 115                     int width,
 116                     int height,
 117                     int cw,
 118                     int ch)
 119 {
 120     global unsigned char *dst_u = dst_y + height * dst_stride_lu;
 121     global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
 122     global unsigned char *src_u = src_y + height * src_stride_lu;
 123     global unsigned char *src_v = src_u + ch * src_stride_ch;
 124     int2 threadIdx, blockIdx, globalIdx;
 125     threadIdx.x = get_local_id(0);
 126     threadIdx.y = get_local_id(1);
 127     blockIdx.x = get_group_id(0);
 128     blockIdx.y = get_group_id(1);
 129     globalIdx.x = get_global_id(0);
 130     globalIdx.y = get_global_id(1);
 131     int padch = get_global_size(1)/2;
 132     global unsigned char *src = globalIdx.y>=padch ? src_v : src_u;
 133     global unsigned char *dst = globalIdx.y>=padch ? dst_v : dst_u;
 134
 135     blockIdx.y = globalIdx.y>=padch ? blockIdx.y - get_num_groups(1)/2 : blockIdx.y;
 136     globalIdx.y = globalIdx.y>=padch ? globalIdx.y - padch : globalIdx.y;
 137
 138     if (!amount) {
 139         if (globalIdx.x < cw && globalIdx.y < ch)
 140             dst[globalIdx.x + globalIdx.y*dst_stride_ch] = src[globalIdx.x + globalIdx.y*src_stride_ch];
 141         return;
 142     }
 143
 144     local uchar l[32][32];
 145     local int lc[CH_RADIUS_X*CH_RADIUS_Y];
 146     int indexIx, indexIy, i, j;
 147     for(i = 0; i <= 1; i++) {
 148         indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
 149         indexIy = indexIy < 0 ? 0 : indexIy;
 150         indexIy = indexIy >= ch ? ch - 1: indexIy;
 151         for(j = 0; j <= 1; j++) {
 152             indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
 153             indexIx = indexIx < 0 ? 0 : indexIx;
 154             indexIx = indexIx >= cw ? cw - 1: indexIx;
 155             l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy * src_stride_ch + indexIx];
 156         }
 157     }
 158
 159     int indexL = threadIdx.y*16 + threadIdx.x;
 160     if (indexL < CH_RADIUS_X*CH_RADIUS_Y)
 161         lc[indexL] = mask[indexL];
 162     barrier(CLK_LOCAL_MEM_FENCE);
 163
 164     int idx, idy, maskIndex;
 165     int sum = 0;
 166     int steps_x = CH_RADIUS_X/2;
 167     int steps_y = CH_RADIUS_Y/2;
 168
 169     \n#pragma unroll\n
 170     for (i = -steps_y; i <= steps_y; i++) {
 171         idy = 8 + i + threadIdx.y;
 172         \n#pragma unroll\n
 173         for (j = -steps_x; j <= steps_x; j++) {
 174             idx = 8 + j + threadIdx.x;
 175             maskIndex = (i + steps_y)*CH_RADIUS_X + j + steps_x;
 176             sum += (int)l[idy][idx] * lc[maskIndex];
 177         }
 178     }
 179     int temp = (int)l[threadIdx.y + 8][threadIdx.x + 8];
 180     int res = temp + (((temp - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
 181     if (globalIdx.x < cw && globalIdx.y < ch)
 182         dst[globalIdx.x + globalIdx.y*dst_stride_ch] = clip_uint8(res);
 183 }
 184
 185 kernel void unsharp_default(global  unsigned char *src,
 186                     global  unsigned char *dst,
 187                     const global  unsigned int *mask_lu,
 188                     const global  unsigned int *mask_ch,
 189                     int amount_lu,
 190                     int amount_ch,
 191                     int step_x_lu,
 192                     int step_y_lu,
 193                     int step_x_ch,
 194                     int step_y_ch,
 195                     int scalebits_lu,
 196                     int scalebits_ch,
 197                     int halfscale_lu,
 198                     int halfscale_ch,
 199                     int src_stride_lu,
 200                     int src_stride_ch,
 201                     int dst_stride_lu,
 202                     int dst_stride_ch,
 203                     int height,
 204                     int width,
 205                     int ch,
 206                     int cw)
 207 {
 208     global unsigned char *dst_y = dst;
 209     global unsigned char *dst_u = dst_y + height * dst_stride_lu;
 210     global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
 211
 212     global unsigned char *src_y = src;
 213     global unsigned char *src_u = src_y + height * src_stride_lu;
 214     global unsigned char *src_v = src_u + ch * src_stride_ch;
 215
 216     global unsigned char *temp_dst;
 217     global unsigned char *temp_src;
 218     const global unsigned int *temp_mask;
 219     int global_id = get_global_id(0);
 220     int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
 221         temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
 222     if (global_id < width * height) {
 223         y = global_id / width;
 224         x = global_id % width;
 225         temp_dst = dst_y;
 226         temp_src = src_y;
 227         temp_src_stride = src_stride_lu;
 228         temp_dst_stride = dst_stride_lu;
 229         temp_height = height;
 230         temp_width = width;
 231         temp_steps_x = step_x_lu;
 232         temp_steps_y = step_y_lu;
 233         temp_mask = mask_lu;
 234         temp_amount = amount_lu;
 235         temp_scalebits = scalebits_lu;
 236         temp_halfscale = halfscale_lu;
 237     } else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
 238         y = (global_id - width * height) / cw;
 239         x = (global_id - width * height) % cw;
 240         temp_dst = dst_u;
 241         temp_src = src_u;
 242         temp_src_stride = src_stride_ch;
 243         temp_dst_stride = dst_stride_ch;
 244         temp_height = ch;
 245         temp_width = cw;
 246         temp_steps_x = step_x_ch;
 247         temp_steps_y = step_y_ch;
 248         temp_mask = mask_ch;
 249         temp_amount = amount_ch;
 250         temp_scalebits = scalebits_ch;
 251         temp_halfscale = halfscale_ch;
 252     } else {
 253         y = (global_id - width * height - ch * cw) / cw;
 254         x = (global_id - width * height - ch * cw) % cw;
 255         temp_dst = dst_v;
 256         temp_src = src_v;
 257         temp_src_stride = src_stride_ch;
 258         temp_dst_stride = dst_stride_ch;
 259         temp_height = ch;
 260         temp_width = cw;
 261         temp_steps_x = step_x_ch;
 262         temp_steps_y = step_y_ch;
 263         temp_mask = mask_ch;
 264         temp_amount = amount_ch;
 265         temp_scalebits = scalebits_ch;
 266         temp_halfscale = halfscale_ch;
 267     }
 268     if (temp_amount) {
 269         sum = 0;
 270         for (j = 0; j <= 2 * temp_steps_y; j++) {
 271             idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
 272             for (i = 0; i <= 2 * temp_steps_x; i++) {
 273                 idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
 274                 sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
 275             }
 276         }
 277         temp = (int)temp_src[x + y * temp_src_stride];
 278         res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);
 279         temp_dst[x + y * temp_dst_stride] = clip_uint8(res);
 280     } else {
 281         temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
 282     }
 283 }
 284 );
 285
 286 #endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */