diff options
Diffstat (limited to 'libavfilter/deshake_opencl_kernel.h')
-rw-r--r-- | libavfilter/deshake_opencl_kernel.h | 225 |
1 files changed, 225 insertions, 0 deletions
diff --git a/libavfilter/deshake_opencl_kernel.h b/libavfilter/deshake_opencl_kernel.h new file mode 100644 index 0000000..dd45d6f --- /dev/null +++ b/libavfilter/deshake_opencl_kernel.h @@ -0,0 +1,225 @@ +/* + * 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_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, int x, int y, + int w, int h,int stride, unsigned char def) +{ + return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[x + 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.5f), (int)(y + 0.5f), 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; + x_f = (int)x; + y_f = (int)y; + x_c = x_f + 1; + y_c = y_f + 1; + + if (x_f < -1 || x_f > width || y_f < -1 || y_f > height) { + return def; + } else { + v4 = pixel(src, x_f, y_f, 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); + v1 = pixel(src, x_c, y_c, 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; + x_f = (int)x; + y_f = (int)y; + x_c = x_f + 1; + y_c = y_f + 1; + + if (x_f < - 1 || x_f > width || y_f < -1 || y_f > height) + return def; + else { + v4 = pixel(src, x_f, y_f, 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); + v1 = pixel(src, x_c, y_c, 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_luma(global unsigned char *src, + global unsigned char *dst, + float4 matrix, + int interpolate, + int fill, + int src_stride_lu, + int dst_stride_lu, + int height, + int width) +{ + int x = get_global_id(0); + int y = get_global_id(1); + int idx_dst = y * dst_stride_lu + x; + unsigned char def = 0; + float x_s = x * matrix.x + y * matrix.y + matrix.z; + float y_s = x * (-matrix.y) + y * matrix.x + matrix.w; + + if (x < width && y < height) { + switch (fill) { + case 0: //FILL_BLANK + def = 0; + break; + case 1: //FILL_ORIGINAL + def = src[y*src_stride_lu + x]; + break; + case 2: //FILL_CLAMP + y_s = clipf(y_s, 0, height - 1); + x_s = clipf(x_s, 0, width - 1); + def = src[(int)y_s * src_stride_lu + (int)x_s]; + break; + case 3: //FILL_MIRROR + y_s = mirror(y_s, height - 1); + x_s = mirror(x_s, width - 1); + def = src[(int)y_s * src_stride_lu + (int)x_s]; + break; + } + switch (interpolate) { + case 0: //INTERPOLATE_NEAREST + dst[idx_dst] = interpolate_nearest(x_s, y_s, src, width, height, src_stride_lu, def); + break; + case 1: //INTERPOLATE_BILINEAR + dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, width, height, src_stride_lu, def); + break; + case 2: //INTERPOLATE_BIQUADRATIC + dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, width, height, src_stride_lu, def); + break; + default: + return; + } + } +} + +kernel void avfilter_transform_chroma(global unsigned char *src, + global unsigned char *dst, + float4 matrix, + int interpolate, + int fill, + 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 x = get_global_id(0); + int y = get_global_id(1); + int pad_ch = get_global_size(1)>>1; + global unsigned char *dst_u = dst + height * dst_stride_lu; + global unsigned char *src_u = src + height * src_stride_lu; + global unsigned char *dst_v = dst_u + ch * dst_stride_ch; + global unsigned char *src_v = src_u + ch * src_stride_ch; + src = y < pad_ch ? src_u : src_v; + dst = y < pad_ch ? dst_u : dst_v; + y = select(y - pad_ch, y, y < pad_ch); + float x_s = x * matrix.x + y * matrix.y + matrix.z; + float y_s = x * (-matrix.y) + y * matrix.x + matrix.w; + int idx_dst = y * dst_stride_ch + x; + unsigned char def; + + if (x < cw && y < ch) { + switch (fill) { + case 0: //FILL_BLANK + def = 0; + break; + case 1: //FILL_ORIGINAL + def = src[y*src_stride_ch + x]; + break; + case 2: //FILL_CLAMP + y_s = clipf(y_s, 0, ch - 1); + x_s = clipf(x_s, 0, cw - 1); + def = src[(int)y_s * src_stride_ch + (int)x_s]; + break; + case 3: //FILL_MIRROR + y_s = mirror(y_s, ch - 1); + x_s = mirror(x_s, cw - 1); + def = src[(int)y_s * src_stride_ch + (int)x_s]; + break; + } + switch (interpolate) { + case 0: //INTERPOLATE_NEAREST + dst[idx_dst] = interpolate_nearest(x_s, y_s, src, cw, ch, src_stride_ch, def); + break; + case 1: //INTERPOLATE_BILINEAR + dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, cw, ch, src_stride_ch, def); + break; + case 2: //INTERPOLATE_BIQUADRATIC + dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, cw, ch, src_stride_ch, def); + break; + default: + return; + } + } +} +); + +#endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */ |