From 63e8fc4d8f0873fd23ffb945d63297c038f10c4e Mon Sep 17 00:00:00 2001 From: highgod0401 Date: Wed, 8 May 2013 08:53:09 +0800 Subject: lavfi: modify names of deshake_kernel.h and unsharp_kernel.h Reviewed-by: Stefano Sabatini Signed-off-by: Michael Niedermayer --- libavfilter/Makefile | 2 +- libavfilter/deshake_kernel.h | 217 ------------------------------------ libavfilter/deshake_opencl_kernel.h | 217 ++++++++++++++++++++++++++++++++++++ libavfilter/opencl_allkernels.c | 4 +- libavfilter/unsharp_kernel.h | 137 ----------------------- libavfilter/unsharp_opencl_kernel.h | 137 +++++++++++++++++++++++ 6 files changed, 357 insertions(+), 357 deletions(-) delete mode 100644 libavfilter/deshake_kernel.h create mode 100644 libavfilter/deshake_opencl_kernel.h delete mode 100644 libavfilter/unsharp_kernel.h create mode 100644 libavfilter/unsharp_opencl_kernel.h diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 913f50211e..73a24dd9b5 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -232,7 +232,7 @@ OBJS-$(CONFIG_AMOVIE_FILTER) += src_movie.o OBJS-$(CONFIG_MOVIE_FILTER) += src_movie.o SKIPHEADERS-$(CONFIG_LIBVIDSTAB) += vidstabutils.h -SKIPHEADERS-$(CONFIG_OPENCL) += opencl_internal.h deshake_kernel.h unsharp_kernel.h +SKIPHEADERS-$(CONFIG_OPENCL) += opencl_internal.h deshake_opencl_kernel.h unsharp_opencl_kernel.h TOOLS = graph2dot TESTPROGS = drawutils filtfmts formats diff --git a/libavfilter/deshake_kernel.h b/libavfilter/deshake_kernel.h deleted file mode 100644 index 1eb06fee39..0000000000 --- a/libavfilter/deshake_kernel.h +++ /dev/null @@ -1,217 +0,0 @@ -/* - * Copyright (C) 2013 Wei Gao - * - * - * 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_DESHAKE_KERNEL_H -#define AVFILTER_DESHAKE_KERNEL_H - -#include "libavutil/opencl.h" - -const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL( - -inline unsigned char pixel(global const unsigned char *src, float x, float y, - int w, int h,int stride, unsigned char def) -{ - return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[(int)x + (int)y * stride]; -} -unsigned char interpolate_nearest(float x, float y, global const unsigned char *src, - int width, int height, int stride, unsigned char def) -{ - return pixel(src, (int)(x + 0.5), (int)(y + 0.5), width, height, stride, def); -} - -unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src, - int width, int height, int stride, unsigned char def) -{ - int x_c, x_f, y_c, y_f; - int v1, v2, v3, v4; - - if (x < -1 || x > width || y < -1 || y > height) { - return def; - } else { - x_f = (int)x; - x_c = x_f + 1; - - y_f = (int)y; - y_c = y_f + 1; - - v1 = pixel(src, x_c, y_c, width, height, stride, def); - v2 = pixel(src, x_c, y_f, width, height, stride, def); - v3 = pixel(src, x_f, y_c, width, height, stride, def); - v4 = pixel(src, x_f, y_f, width, height, stride, def); - - return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) + - v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y))); - } -} - -unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src, - int width, int height, int stride, unsigned char def) -{ - int x_c, x_f, y_c, y_f; - unsigned char v1, v2, v3, v4; - float f1, f2, f3, f4; - - if (x < - 1 || x > width || y < -1 || y > height) - return def; - else { - x_f = (int)x; - x_c = x_f + 1; - y_f = (int)y; - y_c = y_f + 1; - - v1 = pixel(src, x_c, y_c, width, height, stride, def); - v2 = pixel(src, x_c, y_f, width, height, stride, def); - v3 = pixel(src, x_f, y_c, width, height, stride, def); - v4 = pixel(src, x_f, y_f, width, height, stride, def); - - f1 = 1 - sqrt((x_c - x) * (y_c - y)); - f2 = 1 - sqrt((x_c - x) * (y - y_f)); - f3 = 1 - sqrt((x - x_f) * (y_c - y)); - f4 = 1 - sqrt((x - x_f) * (y - y_f)); - return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4); - } -} - -inline const float clipf(float a, float amin, float amax) -{ - if (a < amin) return amin; - else if (a > amax) return amax; - else return a; -} - -inline int mirror(int v, int m) -{ - while ((unsigned)v > (unsigned)m) { - v = -v; - if (v < 0) - v += 2 * m; - } - return v; -} - -kernel void avfilter_transform(global unsigned char *src, - global unsigned char *dst, - global float *matrix, - global float *matrix2, - int interpolate, - int fillmethod, - int src_stride_lu, - int dst_stride_lu, - int src_stride_ch, - int dst_stride_ch, - int height, - int width, - int ch, - int cw) -{ - int global_id = get_global_id(0); - - 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 *tempdst; - global unsigned char *tempsrc; - - int x; - int y; - float x_s; - float y_s; - int tempsrc_stride; - int tempdst_stride; - int temp_height; - int temp_width; - int curpos; - unsigned char def = 0; - if (global_id < width*height) { - y = global_id/width; - x = global_id%width; - x_s = x * matrix[0] + y * matrix[1] + matrix[2]; - y_s = x * matrix[3] + y * matrix[4] + matrix[5]; - tempdst = dst_y; - tempsrc = src_y; - tempsrc_stride = src_stride_lu; - tempdst_stride = dst_stride_lu; - temp_height = height; - temp_width = width; - } else if ((global_id >= width*height)&&(global_id < width*height + ch*cw)) { - y = (global_id - width*height)/cw; - x = (global_id - width*height)%cw; - x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2]; - y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5]; - tempdst = dst_u; - tempsrc = src_u; - tempsrc_stride = src_stride_ch; - tempdst_stride = dst_stride_ch; - temp_height = ch; - temp_width = cw; - } else { - y = (global_id - width*height - ch*cw)/cw; - x = (global_id - width*height - ch*cw)%cw; - x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2]; - y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5]; - tempdst = dst_v; - tempsrc = src_v; - tempsrc_stride = src_stride_ch; - tempdst_stride = dst_stride_ch; - temp_height = ch; - temp_width = cw; - } - curpos = y * tempdst_stride + x; - switch (fillmethod) { - case 0: //FILL_BLANK - def = 0; - break; - case 1: //FILL_ORIGINAL - def = tempsrc[y*tempsrc_stride+x]; - break; - case 2: //FILL_CLAMP - y_s = clipf(y_s, 0, temp_height - 1); - x_s = clipf(x_s, 0, temp_width - 1); - def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s]; - break; - case 3: //FILL_MIRROR - y_s = mirror(y_s,temp_height - 1); - x_s = mirror(x_s,temp_width - 1); - def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s]; - break; - } - switch (interpolate) { - case 0: //INTERPOLATE_NEAREST - tempdst[curpos] = interpolate_nearest(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); - break; - case 1: //INTERPOLATE_BILINEAR - tempdst[curpos] = interpolate_bilinear(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); - break; - case 2: //INTERPOLATE_BIQUADRATIC - tempdst[curpos] = interpolate_biquadratic(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); - break; - default: - return; - } -} -); - -#endif /* AVFILTER_DESHAKE_KERNEL_H */ diff --git a/libavfilter/deshake_opencl_kernel.h b/libavfilter/deshake_opencl_kernel.h new file mode 100644 index 0000000000..ca0bf839b1 --- /dev/null +++ b/libavfilter/deshake_opencl_kernel.h @@ -0,0 +1,217 @@ +/* + * Copyright (C) 2013 Wei Gao + * + * + * 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_DESHAKE_OPENCL_KERNEL_H +#define AVFILTER_DESHAKE_OPENCL_KERNEL_H + +#include "libavutil/opencl.h" + +const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL( + +inline unsigned char pixel(global const unsigned char *src, float x, float y, + int w, int h,int stride, unsigned char def) +{ + return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[(int)x + (int)y * stride]; +} +unsigned char interpolate_nearest(float x, float y, global const unsigned char *src, + int width, int height, int stride, unsigned char def) +{ + return pixel(src, (int)(x + 0.5), (int)(y + 0.5), width, height, stride, def); +} + +unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src, + int width, int height, int stride, unsigned char def) +{ + int x_c, x_f, y_c, y_f; + int v1, v2, v3, v4; + + if (x < -1 || x > width || y < -1 || y > height) { + return def; + } else { + x_f = (int)x; + x_c = x_f + 1; + + y_f = (int)y; + y_c = y_f + 1; + + v1 = pixel(src, x_c, y_c, width, height, stride, def); + v2 = pixel(src, x_c, y_f, width, height, stride, def); + v3 = pixel(src, x_f, y_c, width, height, stride, def); + v4 = pixel(src, x_f, y_f, width, height, stride, def); + + return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) + + v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y))); + } +} + +unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src, + int width, int height, int stride, unsigned char def) +{ + int x_c, x_f, y_c, y_f; + unsigned char v1, v2, v3, v4; + float f1, f2, f3, f4; + + if (x < - 1 || x > width || y < -1 || y > height) + return def; + else { + x_f = (int)x; + x_c = x_f + 1; + y_f = (int)y; + y_c = y_f + 1; + + v1 = pixel(src, x_c, y_c, width, height, stride, def); + v2 = pixel(src, x_c, y_f, width, height, stride, def); + v3 = pixel(src, x_f, y_c, width, height, stride, def); + v4 = pixel(src, x_f, y_f, width, height, stride, def); + + f1 = 1 - sqrt((x_c - x) * (y_c - y)); + f2 = 1 - sqrt((x_c - x) * (y - y_f)); + f3 = 1 - sqrt((x - x_f) * (y_c - y)); + f4 = 1 - sqrt((x - x_f) * (y - y_f)); + return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4); + } +} + +inline const float clipf(float a, float amin, float amax) +{ + if (a < amin) return amin; + else if (a > amax) return amax; + else return a; +} + +inline int mirror(int v, int m) +{ + while ((unsigned)v > (unsigned)m) { + v = -v; + if (v < 0) + v += 2 * m; + } + return v; +} + +kernel void avfilter_transform(global unsigned char *src, + global unsigned char *dst, + global float *matrix, + global float *matrix2, + int interpolate, + int fillmethod, + int src_stride_lu, + int dst_stride_lu, + int src_stride_ch, + int dst_stride_ch, + int height, + int width, + int ch, + int cw) +{ + int global_id = get_global_id(0); + + 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 *tempdst; + global unsigned char *tempsrc; + + int x; + int y; + float x_s; + float y_s; + int tempsrc_stride; + int tempdst_stride; + int temp_height; + int temp_width; + int curpos; + unsigned char def = 0; + if (global_id < width*height) { + y = global_id/width; + x = global_id%width; + x_s = x * matrix[0] + y * matrix[1] + matrix[2]; + y_s = x * matrix[3] + y * matrix[4] + matrix[5]; + tempdst = dst_y; + tempsrc = src_y; + tempsrc_stride = src_stride_lu; + tempdst_stride = dst_stride_lu; + temp_height = height; + temp_width = width; + } else if ((global_id >= width*height)&&(global_id < width*height + ch*cw)) { + y = (global_id - width*height)/cw; + x = (global_id - width*height)%cw; + x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2]; + y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5]; + tempdst = dst_u; + tempsrc = src_u; + tempsrc_stride = src_stride_ch; + tempdst_stride = dst_stride_ch; + temp_height = ch; + temp_width = cw; + } else { + y = (global_id - width*height - ch*cw)/cw; + x = (global_id - width*height - ch*cw)%cw; + x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2]; + y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5]; + tempdst = dst_v; + tempsrc = src_v; + tempsrc_stride = src_stride_ch; + tempdst_stride = dst_stride_ch; + temp_height = ch; + temp_width = cw; + } + curpos = y * tempdst_stride + x; + switch (fillmethod) { + case 0: //FILL_BLANK + def = 0; + break; + case 1: //FILL_ORIGINAL + def = tempsrc[y*tempsrc_stride+x]; + break; + case 2: //FILL_CLAMP + y_s = clipf(y_s, 0, temp_height - 1); + x_s = clipf(x_s, 0, temp_width - 1); + def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s]; + break; + case 3: //FILL_MIRROR + y_s = mirror(y_s,temp_height - 1); + x_s = mirror(x_s,temp_width - 1); + def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s]; + break; + } + switch (interpolate) { + case 0: //INTERPOLATE_NEAREST + tempdst[curpos] = interpolate_nearest(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); + break; + case 1: //INTERPOLATE_BILINEAR + tempdst[curpos] = interpolate_bilinear(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); + break; + case 2: //INTERPOLATE_BIQUADRATIC + tempdst[curpos] = interpolate_biquadratic(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); + break; + default: + return; + } +} +); + +#endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */ diff --git a/libavfilter/opencl_allkernels.c b/libavfilter/opencl_allkernels.c index b4530c0b2e..6d80fa8598 100644 --- a/libavfilter/opencl_allkernels.c +++ b/libavfilter/opencl_allkernels.c @@ -21,8 +21,8 @@ #include "opencl_allkernels.h" #if CONFIG_OPENCL #include "libavutil/opencl.h" -#include "deshake_kernel.h" -#include "unsharp_kernel.h" +#include "deshake_opencl_kernel.h" +#include "unsharp_opencl_kernel.h" #endif #define OPENCL_REGISTER_KERNEL_CODE(X, x) \ diff --git a/libavfilter/unsharp_kernel.h b/libavfilter/unsharp_kernel.h deleted file mode 100644 index dd65bc0a9a..0000000000 --- a/libavfilter/unsharp_kernel.h +++ /dev/null @@ -1,137 +0,0 @@ -/* - * Copyright (C) 2013 Wei Gao - * - * 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_KERNEL_H -#define AVFILTER_UNSHARP_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(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_KERNEL_H */ diff --git a/libavfilter/unsharp_opencl_kernel.h b/libavfilter/unsharp_opencl_kernel.h new file mode 100644 index 0000000000..0cc8e906b4 --- /dev/null +++ b/libavfilter/unsharp_opencl_kernel.h @@ -0,0 +1,137 @@ +/* + * Copyright (C) 2013 Wei Gao + * + * 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(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 */ -- cgit v1.2.3