FFmpeg
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups Pages
unsharp_opencl_kernel.h
Go to the documentation of this file.
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 */
Generated on Sun Jul 20 2014 23:06:01 for FFmpeg by   doxygen 1.8.2