Commit | Line | Data |
---|---|---|
2ba45a60 DM |
1 | /* |
2 | * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com> | |
3 | * Copyright (C) 2013 Lenny Wang | |
4 | * | |
5 | * | |
6 | * This file is part of FFmpeg. | |
7 | * | |
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. | |
12 | * | |
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. | |
17 | * | |
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 | |
21 | */ | |
22 | ||
23 | #ifndef AVFILTER_DESHAKE_OPENCL_KERNEL_H | |
24 | #define AVFILTER_DESHAKE_OPENCL_KERNEL_H | |
25 | ||
26 | #include "libavutil/opencl.h" | |
27 | ||
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) | |
31 | { | |
32 | return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[x + y * stride]; | |
33 | } | |
34 | ||
35 | unsigned char interpolate_nearest(float x, float y, global const unsigned char *src, | |
36 | int width, int height, int stride, unsigned char def) | |
37 | { | |
38 | return pixel(src, (int)(x + 0.5f), (int)(y + 0.5f), width, height, stride, def); | |
39 | } | |
40 | ||
41 | unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src, | |
42 | int width, int height, int stride, unsigned char def) | |
43 | { | |
44 | int x_c, x_f, y_c, y_f; | |
45 | int v1, v2, v3, v4; | |
46 | x_f = (int)x; | |
47 | y_f = (int)y; | |
48 | x_c = x_f + 1; | |
49 | y_c = y_f + 1; | |
50 | ||
51 | if (x_f < -1 || x_f > width || y_f < -1 || y_f > height) { | |
52 | return def; | |
53 | } else { | |
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))); | |
60 | } | |
61 | } | |
62 | ||
63 | unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src, | |
64 | int width, int height, int stride, unsigned char def) | |
65 | { | |
66 | int x_c, x_f, y_c, y_f; | |
67 | unsigned char v1, v2, v3, v4; | |
68 | float f1, f2, f3, f4; | |
69 | x_f = (int)x; | |
70 | y_f = (int)y; | |
71 | x_c = x_f + 1; | |
72 | y_c = y_f + 1; | |
73 | ||
74 | if (x_f < - 1 || x_f > width || y_f < -1 || y_f > height) | |
75 | return def; | |
76 | else { | |
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); | |
81 | ||
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); | |
87 | } | |
88 | } | |
89 | ||
90 | inline const float clipf(float a, float amin, float amax) | |
91 | { | |
92 | if (a < amin) return amin; | |
93 | else if (a > amax) return amax; | |
94 | else return a; | |
95 | } | |
96 | ||
97 | inline int mirror(int v, int m) | |
98 | { | |
99 | while ((unsigned)v > (unsigned)m) { | |
100 | v = -v; | |
101 | if (v < 0) | |
102 | v += 2 * m; | |
103 | } | |
104 | return v; | |
105 | } | |
106 | ||
107 | kernel void avfilter_transform_luma(global unsigned char *src, | |
108 | global unsigned char *dst, | |
109 | float4 matrix, | |
110 | int interpolate, | |
111 | int fill, | |
112 | int src_stride_lu, | |
113 | int dst_stride_lu, | |
114 | int height, | |
115 | int width) | |
116 | { | |
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; | |
123 | ||
124 | if (x < width && y < height) { | |
125 | switch (fill) { | |
126 | case 0: //FILL_BLANK | |
127 | def = 0; | |
128 | break; | |
129 | case 1: //FILL_ORIGINAL | |
130 | def = src[y*src_stride_lu + x]; | |
131 | break; | |
132 | case 2: //FILL_CLAMP | |
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]; | |
136 | break; | |
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]; | |
141 | break; | |
142 | } | |
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); | |
146 | break; | |
147 | case 1: //INTERPOLATE_BILINEAR | |
148 | dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, width, height, src_stride_lu, def); | |
149 | break; | |
150 | case 2: //INTERPOLATE_BIQUADRATIC | |
151 | dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, width, height, src_stride_lu, def); | |
152 | break; | |
153 | default: | |
154 | return; | |
155 | } | |
156 | } | |
157 | } | |
158 | ||
159 | kernel void avfilter_transform_chroma(global unsigned char *src, | |
160 | global unsigned char *dst, | |
161 | float4 matrix, | |
162 | int interpolate, | |
163 | int fill, | |
164 | int src_stride_lu, | |
165 | int dst_stride_lu, | |
166 | int src_stride_ch, | |
167 | int dst_stride_ch, | |
168 | int height, | |
169 | int width, | |
170 | int ch, | |
171 | int cw) | |
172 | { | |
173 | ||
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; | |
187 | unsigned char def; | |
188 | ||
189 | if (x < cw && y < ch) { | |
190 | switch (fill) { | |
191 | case 0: //FILL_BLANK | |
192 | def = 0; | |
193 | break; | |
194 | case 1: //FILL_ORIGINAL | |
195 | def = src[y*src_stride_ch + x]; | |
196 | break; | |
197 | case 2: //FILL_CLAMP | |
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]; | |
201 | break; | |
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]; | |
206 | break; | |
207 | } | |
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); | |
211 | break; | |
212 | case 1: //INTERPOLATE_BILINEAR | |
213 | dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, cw, ch, src_stride_ch, def); | |
214 | break; | |
215 | case 2: //INTERPOLATE_BIQUADRATIC | |
216 | dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, cw, ch, src_stride_ch, def); | |
217 | break; | |
218 | default: | |
219 | return; | |
220 | } | |
221 | } | |
222 | } | |
223 | ); | |
224 | ||
225 | #endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */ |