summaryrefslogtreecommitdiff
path: root/libavfilter/unsharp_opencl_kernel.h (plain)
blob: 307d0f1814ca833ed65fd3dc0c4d121ce5c44c0b
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_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
134kernel 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
241kernel 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 */
343
generated by cgit at 2024-05-18 04:42:44 (GMT)