avfilter/vf_overlay: Restore shorthand option order
[ffmpeg.git] / 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_x,
40                     global int *mask_y,
41                     int amount,
42                     int scalebits,
43                     int halfscale,
44                     int src_stride,
45                     int dst_stride,
46                     int width,
47                     int height)
48 {
49     int2 threadIdx, blockIdx, globalIdx;
50     threadIdx.x = get_local_id(0);
51     threadIdx.y = get_local_id(1);
52     blockIdx.x = get_group_id(0);
53     blockIdx.y = get_group_id(1);
54     globalIdx.x = get_global_id(0);
55     globalIdx.y = get_global_id(1);
56
57     if (!amount) {
58         if (globalIdx.x < width && globalIdx.y < height)
59             dst[globalIdx.x + globalIdx.y*dst_stride] = src[globalIdx.x + globalIdx.y*src_stride];
60         return;
61     }
62
63     local unsigned int l[32][32];
64     local unsigned int lcx[LU_RADIUS_X];
65     local unsigned int lcy[LU_RADIUS_Y];
66     int indexIx, indexIy, i, j;
67
68     //load up tile: actual workspace + halo of 8 points in x and y \n
69     for(i = 0; i <= 1; i++) {
70         indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
71         indexIy = indexIy < 0 ? 0 : indexIy;
72         indexIy = indexIy >= height ? height - 1: indexIy;
73         for(j = 0; j <= 1; j++) {
74             indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
75             indexIx = indexIx < 0 ? 0 : indexIx;
76             indexIx = indexIx >= width ? width - 1: indexIx;
77             l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy*src_stride + indexIx];
78         }
79     }
80
81     int indexL = threadIdx.y*16 + threadIdx.x;
82     if (indexL < LU_RADIUS_X)
83         lcx[indexL] = mask_x[indexL];
84     if (indexL < LU_RADIUS_Y)
85         lcy[indexL] = mask_y[indexL];
86     barrier(CLK_LOCAL_MEM_FENCE);
87
88     //needed for unsharp mask application in the end \n
89     int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];
90
91     int idx, idy, maskIndex;
92     int temp[2] = {0};
93     int steps_x = (LU_RADIUS_X-1)/2;
94     int steps_y = (LU_RADIUS_Y-1)/2;
95
96     // compute the actual workspace + left&right halos \n
97       \n#pragma unroll\n
98     for (j = 0; j <=1; j++) {
99       //extra work to cover left and right halos \n
100       idx = 16*j + threadIdx.x;
101       \n#pragma unroll\n
102         for (i = -steps_y; i <= steps_y; i++) {
103           idy = 8 + i + threadIdx.y;
104           maskIndex = (i + steps_y);
105           temp[j] += (int)l[idy][idx] * lcy[maskIndex];
106         }
107     }
108     barrier(CLK_LOCAL_MEM_FENCE);
109     //save results from the vertical filter in local memory \n
110     idy = 8 + threadIdx.y;
111       \n#pragma unroll\n
112     for (j = 0; j <=1; j++) {
113       idx = 16*j + threadIdx.x;
114       l[idy][idx] = temp[j];
115     }
116     barrier(CLK_LOCAL_MEM_FENCE);
117
118     //compute results with the horizontal filter \n
119     int sum = 0;
120     idy = 8 + threadIdx.y;
121     \n#pragma unroll\n
122       for (j = -steps_x; j <= steps_x; j++) {
123         idx = 8 + j + threadIdx.x;
124         maskIndex = j + steps_x;
125         sum += (int)l[idy][idx] * lcx[maskIndex];
126       }
127
128     int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
129
130     if (globalIdx.x < width && globalIdx.y < height)
131         dst[globalIdx.x + globalIdx.y*dst_stride] = clip_uint8(res);
132 }
133
134 kernel void unsharp_chroma(
135                     global unsigned char *src_y,
136                     global unsigned char *dst_y,
137                     global int *mask_x,
138                     global int *mask_y,
139                     int amount,
140                     int scalebits,
141                     int halfscale,
142                     int src_stride_lu,
143                     int src_stride_ch,
144                     int dst_stride_lu,
145                     int dst_stride_ch,
146                     int width,
147                     int height,
148                     int cw,
149                     int ch)
150 {
151     global unsigned char *dst_u = dst_y + height * dst_stride_lu;
152     global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
153     global unsigned char *src_u = src_y + height * src_stride_lu;
154     global unsigned char *src_v = src_u + ch * src_stride_ch;
155     int2 threadIdx, blockIdx, globalIdx;
156     threadIdx.x = get_local_id(0);
157     threadIdx.y = get_local_id(1);
158     blockIdx.x = get_group_id(0);
159     blockIdx.y = get_group_id(1);
160     globalIdx.x = get_global_id(0);
161     globalIdx.y = get_global_id(1);
162     int padch = get_global_size(1)/2;
163     global unsigned char *src = globalIdx.y>=padch ? src_v : src_u;
164     global unsigned char *dst = globalIdx.y>=padch ? dst_v : dst_u;
165
166     blockIdx.y = globalIdx.y>=padch ? blockIdx.y - get_num_groups(1)/2 : blockIdx.y;
167     globalIdx.y = globalIdx.y>=padch ? globalIdx.y - padch : globalIdx.y;
168
169     if (!amount) {
170         if (globalIdx.x < cw && globalIdx.y < ch)
171             dst[globalIdx.x + globalIdx.y*dst_stride_ch] = src[globalIdx.x + globalIdx.y*src_stride_ch];
172         return;
173     }
174
175     local unsigned int l[32][32];
176     local unsigned int lcx[CH_RADIUS_X];
177     local unsigned int lcy[CH_RADIUS_Y];
178     int indexIx, indexIy, i, j;
179     for(i = 0; i <= 1; i++) {
180         indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
181         indexIy = indexIy < 0 ? 0 : indexIy;
182         indexIy = indexIy >= ch ? ch - 1: indexIy;
183         for(j = 0; j <= 1; j++) {
184             indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
185             indexIx = indexIx < 0 ? 0 : indexIx;
186             indexIx = indexIx >= cw ? cw - 1: indexIx;
187             l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy * src_stride_ch + indexIx];
188         }
189     }
190
191     int indexL = threadIdx.y*16 + threadIdx.x;
192     if (indexL < CH_RADIUS_X)
193         lcx[indexL] = mask_x[indexL];
194     if (indexL < CH_RADIUS_Y)
195         lcy[indexL] = mask_y[indexL];
196     barrier(CLK_LOCAL_MEM_FENCE);
197
198     int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];
199
200     int idx, idy, maskIndex;
201     int steps_x = CH_RADIUS_X/2;
202     int steps_y = CH_RADIUS_Y/2;
203     int temp[2] = {0,0};
204
205     \n#pragma unroll\n
206       for (j = 0; j <= 1; j++) {
207         idx = 16*j + threadIdx.x;
208         \n#pragma unroll\n
209           for (i = -steps_y; i <= steps_y; i++) {
210             idy = 8 + i + threadIdx.y;
211             maskIndex = i + steps_y;
212             temp[j] += (int)l[idy][idx] * lcy[maskIndex];
213           }
214       }
215
216     barrier(CLK_LOCAL_MEM_FENCE);
217     idy = 8 + threadIdx.y;
218     \n#pragma unroll\n
219     for (j = 0; j <= 1; j++) {
220       idx = 16*j + threadIdx.x;
221       l[idy][idx] = temp[j];
222     }
223     barrier(CLK_LOCAL_MEM_FENCE);
224
225     //compute results with the horizontal filter \n
226     int sum = 0;
227     idy = 8 + threadIdx.y;
228     \n#pragma unroll\n
229       for (j = -steps_x; j <= steps_x; j++) {
230         idx = 8 + j + threadIdx.x;
231         maskIndex = j + steps_x;
232         sum += (int)l[idy][idx] * lcx[maskIndex];
233       }
234
235     int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
236
237     if (globalIdx.x < cw && globalIdx.y < ch)
238         dst[globalIdx.x + globalIdx.y*dst_stride_ch] = clip_uint8(res);
239 }
240
241 kernel void unsharp_default(global  unsigned char *src,
242                     global  unsigned char *dst,
243                     const global  unsigned int *mask_lu,
244                     const global  unsigned int *mask_ch,
245                     int amount_lu,
246                     int amount_ch,
247                     int step_x_lu,
248                     int step_y_lu,
249                     int step_x_ch,
250                     int step_y_ch,
251                     int scalebits_lu,
252                     int scalebits_ch,
253                     int halfscale_lu,
254                     int halfscale_ch,
255                     int src_stride_lu,
256                     int src_stride_ch,
257                     int dst_stride_lu,
258                     int dst_stride_ch,
259                     int height,
260                     int width,
261                     int ch,
262                     int cw)
263 {
264     global unsigned char *dst_y = dst;
265     global unsigned char *dst_u = dst_y + height * dst_stride_lu;
266     global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
267
268     global unsigned char *src_y = src;
269     global unsigned char *src_u = src_y + height * src_stride_lu;
270     global unsigned char *src_v = src_u + ch * src_stride_ch;
271
272     global unsigned char *temp_dst;
273     global unsigned char *temp_src;
274     const global unsigned int *temp_mask;
275     int global_id = get_global_id(0);
276     int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
277         temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
278     if (global_id < width * height) {
279         y = global_id / width;
280         x = global_id % width;
281         temp_dst = dst_y;
282         temp_src = src_y;
283         temp_src_stride = src_stride_lu;
284         temp_dst_stride = dst_stride_lu;
285         temp_height = height;
286         temp_width = width;
287         temp_steps_x = step_x_lu;
288         temp_steps_y = step_y_lu;
289         temp_mask = mask_lu;
290         temp_amount = amount_lu;
291         temp_scalebits = scalebits_lu;
292         temp_halfscale = halfscale_lu;
293     } else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
294         y = (global_id - width * height) / cw;
295         x = (global_id - width * height) % cw;
296         temp_dst = dst_u;
297         temp_src = src_u;
298         temp_src_stride = src_stride_ch;
299         temp_dst_stride = dst_stride_ch;
300         temp_height = ch;
301         temp_width = cw;
302         temp_steps_x = step_x_ch;
303         temp_steps_y = step_y_ch;
304         temp_mask = mask_ch;
305         temp_amount = amount_ch;
306         temp_scalebits = scalebits_ch;
307         temp_halfscale = halfscale_ch;
308     } else {
309         y = (global_id - width * height - ch * cw) / cw;
310         x = (global_id - width * height - ch * cw) % cw;
311         temp_dst = dst_v;
312         temp_src = src_v;
313         temp_src_stride = src_stride_ch;
314         temp_dst_stride = dst_stride_ch;
315         temp_height = ch;
316         temp_width = cw;
317         temp_steps_x = step_x_ch;
318         temp_steps_y = step_y_ch;
319         temp_mask = mask_ch;
320         temp_amount = amount_ch;
321         temp_scalebits = scalebits_ch;
322         temp_halfscale = halfscale_ch;
323     }
324     if (temp_amount) {
325         sum = 0;
326         for (j = 0; j <= 2 * temp_steps_y; j++) {
327             idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
328             for (i = 0; i <= 2 * temp_steps_x; i++) {
329                 idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
330                 sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
331             }
332         }
333         temp = (int)temp_src[x + y * temp_src_stride];
334         res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);
335         temp_dst[x + y * temp_dst_stride] = clip_uint8(res);
336     } else {
337         temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
338     }
339 }
340 );
341
342 #endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */