[deb_ffmpeg.git] / ffmpeg / libavfilter / unsharp_opencl_kernel.h

/*
 * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
 * Copyright (C) 2013 Lenny Wang
 *
 * This file is part of FFmpeg.
 *
 * FFmpeg is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Lesser General Public
 * License as published by the Free Software Foundation; either
 * version 2.1 of the License, or (at your option) any later version.
 *
 * FFmpeg is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Lesser General Public License for more details.
 *
 * You should have received a copy of the GNU Lesser General Public
 * License along with FFmpeg; if not, write to the Free Software
 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
 */

#ifndef AVFILTER_UNSHARP_OPENCL_KERNEL_H
#define AVFILTER_UNSHARP_OPENCL_KERNEL_H

#include "libavutil/opencl.h"

const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
inline unsigned char clip_uint8(int a)
{
    if (a & (~0xFF))
        return (-a)>>31;
    else
        return a;
}

kernel void unsharp_luma(
                    global unsigned char *src,
                    global unsigned char *dst,
                    global int *mask,
                    int amount,
                    int scalebits,
                    int halfscale,
                    int src_stride,
                    int dst_stride,
                    int width,
                    int height)
{
    int2 threadIdx, blockIdx, globalIdx;
    threadIdx.x = get_local_id(0);
    threadIdx.y = get_local_id(1);
    blockIdx.x = get_group_id(0);
    blockIdx.y = get_group_id(1);
    globalIdx.x = get_global_id(0);
    globalIdx.y = get_global_id(1);

    if (!amount) {
        if (globalIdx.x < width && globalIdx.y < height)
            dst[globalIdx.x + globalIdx.y*dst_stride] = src[globalIdx.x + globalIdx.y*src_stride];
        return;
    }

    local uchar l[32][32];
    local int lc[LU_RADIUS_X*LU_RADIUS_Y];
    int indexIx, indexIy, i, j;

    for(i = 0; i <= 1; i++) {
        indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
        indexIy = indexIy < 0 ? 0 : indexIy;
        indexIy = indexIy >= height ? height - 1: indexIy;
        for(j = 0; j <= 1; j++) {
            indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
            indexIx = indexIx < 0 ? 0 : indexIx;
            indexIx = indexIx >= width ? width - 1: indexIx;
            l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy*src_stride + indexIx];
        }
    }

    int indexL = threadIdx.y*16 + threadIdx.x;
    if (indexL < LU_RADIUS_X*LU_RADIUS_Y)
        lc[indexL] = mask[indexL];
    barrier(CLK_LOCAL_MEM_FENCE);

    int idx, idy, maskIndex;
    int sum = 0;
    int steps_x = LU_RADIUS_X/2;
    int steps_y = LU_RADIUS_Y/2;

    \n#pragma unroll\n
    for (i = -steps_y; i <= steps_y; i++) {
        idy = 8 + i + threadIdx.y;
        \n#pragma unroll\n
        for (j = -steps_x; j <= steps_x; j++) {
            idx = 8 + j + threadIdx.x;
            maskIndex = (i + steps_y)*LU_RADIUS_X + j + steps_x;
            sum += (int)l[idy][idx] * lc[maskIndex];
        }
    }
    int temp = (int)l[threadIdx.y + 8][threadIdx.x + 8];
    int res = temp + (((temp - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
    if (globalIdx.x < width && globalIdx.y < height)
        dst[globalIdx.x + globalIdx.y*dst_stride] = clip_uint8(res);
}

kernel void unsharp_chroma(
                    global unsigned char *src_y,
                    global unsigned char *dst_y,
                    global int *mask,
                    int amount,
                    int scalebits,
                    int halfscale,
                    int src_stride_lu,
                    int src_stride_ch,
                    int dst_stride_lu,
                    int dst_stride_ch,
                    int width,
                    int height,
                    int cw,
                    int ch)
{
    global unsigned char *dst_u = dst_y + height * dst_stride_lu;
    global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
    global unsigned char *src_u = src_y + height * src_stride_lu;
    global unsigned char *src_v = src_u + ch * src_stride_ch;
    int2 threadIdx, blockIdx, globalIdx;
    threadIdx.x = get_local_id(0);
    threadIdx.y = get_local_id(1);
    blockIdx.x = get_group_id(0);
    blockIdx.y = get_group_id(1);
    globalIdx.x = get_global_id(0);
    globalIdx.y = get_global_id(1);
    int padch = get_global_size(1)/2;
    global unsigned char *src = globalIdx.y>=padch ? src_v : src_u;
    global unsigned char *dst = globalIdx.y>=padch ? dst_v : dst_u;

    blockIdx.y = globalIdx.y>=padch ? blockIdx.y - get_num_groups(1)/2 : blockIdx.y;
    globalIdx.y = globalIdx.y>=padch ? globalIdx.y - padch : globalIdx.y;

    if (!amount) {
        if (globalIdx.x < cw && globalIdx.y < ch)
            dst[globalIdx.x + globalIdx.y*dst_stride_ch] = src[globalIdx.x + globalIdx.y*src_stride_ch];
        return;
    }

    local uchar l[32][32];
    local int lc[CH_RADIUS_X*CH_RADIUS_Y];
    int indexIx, indexIy, i, j;
    for(i = 0; i <= 1; i++) {
        indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
        indexIy = indexIy < 0 ? 0 : indexIy;
        indexIy = indexIy >= ch ? ch - 1: indexIy;
        for(j = 0; j <= 1; j++) {
            indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
            indexIx = indexIx < 0 ? 0 : indexIx;
            indexIx = indexIx >= cw ? cw - 1: indexIx;
            l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy * src_stride_ch + indexIx];
        }
    }

    int indexL = threadIdx.y*16 + threadIdx.x;
    if (indexL < CH_RADIUS_X*CH_RADIUS_Y)
        lc[indexL] = mask[indexL];
    barrier(CLK_LOCAL_MEM_FENCE);

    int idx, idy, maskIndex;
    int sum = 0;
    int steps_x = CH_RADIUS_X/2;
    int steps_y = CH_RADIUS_Y/2;

    \n#pragma unroll\n
    for (i = -steps_y; i <= steps_y; i++) {
        idy = 8 + i + threadIdx.y;
        \n#pragma unroll\n
        for (j = -steps_x; j <= steps_x; j++) {
            idx = 8 + j + threadIdx.x;
            maskIndex = (i + steps_y)*CH_RADIUS_X + j + steps_x;
            sum += (int)l[idy][idx] * lc[maskIndex];
        }
    }
    int temp = (int)l[threadIdx.y + 8][threadIdx.x + 8];
    int res = temp + (((temp - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
    if (globalIdx.x < cw && globalIdx.y < ch)
        dst[globalIdx.x + globalIdx.y*dst_stride_ch] = clip_uint8(res);
}

kernel void unsharp_default(global  unsigned char *src,
                    global  unsigned char *dst,
                    const global  unsigned int *mask_lu,
                    const global  unsigned int *mask_ch,
                    int amount_lu,
                    int amount_ch,
                    int step_x_lu,
                    int step_y_lu,
                    int step_x_ch,
                    int step_y_ch,
                    int scalebits_lu,
                    int scalebits_ch,
                    int halfscale_lu,
                    int halfscale_ch,
                    int src_stride_lu,
                    int src_stride_ch,
                    int dst_stride_lu,
                    int dst_stride_ch,
                    int height,
                    int width,
                    int ch,
                    int cw)
{
    global unsigned char *dst_y = dst;
    global unsigned char *dst_u = dst_y + height * dst_stride_lu;
    global unsigned char *dst_v = dst_u + ch * dst_stride_ch;

    global unsigned char *src_y = src;
    global unsigned char *src_u = src_y + height * src_stride_lu;
    global unsigned char *src_v = src_u + ch * src_stride_ch;

    global unsigned char *temp_dst;
    global unsigned char *temp_src;
    const global unsigned int *temp_mask;
    int global_id = get_global_id(0);
    int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
        temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
    if (global_id < width * height) {
        y = global_id / width;
        x = global_id % width;
        temp_dst = dst_y;
        temp_src = src_y;
        temp_src_stride = src_stride_lu;
        temp_dst_stride = dst_stride_lu;
        temp_height = height;
        temp_width = width;
        temp_steps_x = step_x_lu;
        temp_steps_y = step_y_lu;
        temp_mask = mask_lu;
        temp_amount = amount_lu;
        temp_scalebits = scalebits_lu;
        temp_halfscale = halfscale_lu;
    } else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
        y = (global_id - width * height) / cw;
        x = (global_id - width * height) % cw;
        temp_dst = dst_u;
        temp_src = src_u;
        temp_src_stride = src_stride_ch;
        temp_dst_stride = dst_stride_ch;
        temp_height = ch;
        temp_width = cw;
        temp_steps_x = step_x_ch;
        temp_steps_y = step_y_ch;
        temp_mask = mask_ch;
        temp_amount = amount_ch;
        temp_scalebits = scalebits_ch;
        temp_halfscale = halfscale_ch;
    } else {
        y = (global_id - width * height - ch * cw) / cw;
        x = (global_id - width * height - ch * cw) % cw;
        temp_dst = dst_v;
        temp_src = src_v;
        temp_src_stride = src_stride_ch;
        temp_dst_stride = dst_stride_ch;
        temp_height = ch;
        temp_width = cw;
        temp_steps_x = step_x_ch;
        temp_steps_y = step_y_ch;
        temp_mask = mask_ch;
        temp_amount = amount_ch;
        temp_scalebits = scalebits_ch;
        temp_halfscale = halfscale_ch;
    }
    if (temp_amount) {
        sum = 0;
        for (j = 0; j <= 2 * temp_steps_y; j++) {
            idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
            for (i = 0; i <= 2 * temp_steps_x; i++) {
                idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
                sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
            }
        }
        temp = (int)temp_src[x + y * temp_src_stride];
        res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);
        temp_dst[x + y * temp_dst_stride] = clip_uint8(res);
    } else {
        temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
    }
}
);

#endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */
Commit	Line	Data
2ba45a60 DM	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.ydst_stride] = src[globalIdx.x + globalIdx.ysrc_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[i16 + threadIdx.y][j16 + 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.ydst_stride_ch] = src[globalIdx.x + globalIdx.ysrc_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[i16 + threadIdx.y][j16 + 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 */