2 * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
3 * Copyright (C) 2013 Lenny Wang
6 * This file is part of FFmpeg.
8 * FFmpeg is free software; you can redistribute it and/or
9 * modify it under the terms of the GNU Lesser General Public
10 * License as published by the Free Software Foundation; either
11 * version 2.1 of the License, or (at your option) any later version.
13 * FFmpeg is distributed in the hope that it will be useful,
14 * but WITHOUT ANY WARRANTY; without even the implied warranty of
15 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
16 * Lesser General Public License for more details.
18 * You should have received a copy of the GNU Lesser General Public
19 * License along with FFmpeg; if not, write to the Free Software
20 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
23 #ifndef AVFILTER_DESHAKE_OPENCL_KERNEL_H
24 #define AVFILTER_DESHAKE_OPENCL_KERNEL_H
26 #include "libavutil/opencl.h"
28 const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL(
29 inline unsigned char pixel(global const unsigned char *src, int x, int y,
30 int w, int h,int stride, unsigned char def)
32 return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[x + y * stride];
35 unsigned char interpolate_nearest(float x, float y, global const unsigned char *src,
36 int width, int height, int stride, unsigned char def)
38 return pixel(src, (int)(x + 0.5f), (int)(y + 0.5f), width, height, stride, def);
41 unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src,
42 int width, int height, int stride, unsigned char def)
44 int x_c, x_f, y_c, y_f;
51 if (x_f < -1 || x_f > width || y_f < -1 || y_f > height) {
54 v4 = pixel(src, x_f, y_f, width, height, stride, def);
55 v2 = pixel(src, x_c, y_f, width, height, stride, def);
56 v3 = pixel(src, x_f, y_c, width, height, stride, def);
57 v1 = pixel(src, x_c, y_c, width, height, stride, def);
58 return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) +
59 v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y)));
63 unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src,
64 int width, int height, int stride, unsigned char def)
66 int x_c, x_f, y_c, y_f;
67 unsigned char v1, v2, v3, v4;
74 if (x_f < - 1 || x_f > width || y_f < -1 || y_f > height)
77 v4 = pixel(src, x_f, y_f, width, height, stride, def);
78 v2 = pixel(src, x_c, y_f, width, height, stride, def);
79 v3 = pixel(src, x_f, y_c, width, height, stride, def);
80 v1 = pixel(src, x_c, y_c, width, height, stride, def);
82 f1 = 1 - sqrt((x_c - x) * (y_c - y));
83 f2 = 1 - sqrt((x_c - x) * (y - y_f));
84 f3 = 1 - sqrt((x - x_f) * (y_c - y));
85 f4 = 1 - sqrt((x - x_f) * (y - y_f));
86 return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4);
90 inline const float clipf(float a, float amin, float amax)
92 if (a < amin) return amin;
93 else if (a > amax) return amax;
97 inline int mirror(int v, int m)
99 while ((unsigned)v > (unsigned)m) {
107 kernel void avfilter_transform_luma(global unsigned char *src,
108 global unsigned char *dst,
117 int x = get_global_id(0);
118 int y = get_global_id(1);
119 int idx_dst = y * dst_stride_lu + x;
120 unsigned char def = 0;
121 float x_s = x * matrix.x + y * matrix.y + matrix.z;
122 float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
124 if (x < width && y < height) {
129 case 1: //FILL_ORIGINAL
130 def = src[y*src_stride_lu + x];
133 y_s = clipf(y_s, 0, height - 1);
134 x_s = clipf(x_s, 0, width - 1);
135 def = src[(int)y_s * src_stride_lu + (int)x_s];
137 case 3: //FILL_MIRROR
138 y_s = mirror(y_s, height - 1);
139 x_s = mirror(x_s, width - 1);
140 def = src[(int)y_s * src_stride_lu + (int)x_s];
143 switch (interpolate) {
144 case 0: //INTERPOLATE_NEAREST
145 dst[idx_dst] = interpolate_nearest(x_s, y_s, src, width, height, src_stride_lu, def);
147 case 1: //INTERPOLATE_BILINEAR
148 dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, width, height, src_stride_lu, def);
150 case 2: //INTERPOLATE_BIQUADRATIC
151 dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, width, height, src_stride_lu, def);
159 kernel void avfilter_transform_chroma(global unsigned char *src,
160 global unsigned char *dst,
174 int x = get_global_id(0);
175 int y = get_global_id(1);
176 int pad_ch = get_global_size(1)>>1;
177 global unsigned char *dst_u = dst + height * dst_stride_lu;
178 global unsigned char *src_u = src + height * src_stride_lu;
179 global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
180 global unsigned char *src_v = src_u + ch * src_stride_ch;
181 src = y < pad_ch ? src_u : src_v;
182 dst = y < pad_ch ? dst_u : dst_v;
183 y = select(y - pad_ch, y, y < pad_ch);
184 float x_s = x * matrix.x + y * matrix.y + matrix.z;
185 float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
186 int idx_dst = y * dst_stride_ch + x;
189 if (x < cw && y < ch) {
194 case 1: //FILL_ORIGINAL
195 def = src[y*src_stride_ch + x];
198 y_s = clipf(y_s, 0, ch - 1);
199 x_s = clipf(x_s, 0, cw - 1);
200 def = src[(int)y_s * src_stride_ch + (int)x_s];
202 case 3: //FILL_MIRROR
203 y_s = mirror(y_s, ch - 1);
204 x_s = mirror(x_s, cw - 1);
205 def = src[(int)y_s * src_stride_ch + (int)x_s];
208 switch (interpolate) {
209 case 0: //INTERPOLATE_NEAREST
210 dst[idx_dst] = interpolate_nearest(x_s, y_s, src, cw, ch, src_stride_ch, def);
212 case 1: //INTERPOLATE_BILINEAR
213 dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, cw, ch, src_stride_ch, def);
215 case 2: //INTERPOLATE_BIQUADRATIC
216 dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, cw, ch, src_stride_ch, def);
225 #endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */