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 27const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL( 28inline unsigned char clip_uint8(int a) 29{ 30 if (a & (~0xFF)) 31 return (-a)>>31; 32 else 33 return a; 34} 35 36kernel 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 104kernel 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 185kernel 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 */ 287