Go to most recent revision | Details | Last modification | View Log | RSS feed
Rev | Author | Line No. | Line |
---|---|---|---|
4349 | Serge | 1 | /* |
2 | * Copyright (C) 2013 Wei Gao |
||
3 | * |
||
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_DESHAKE_OPENCL_KERNEL_H |
||
23 | #define AVFILTER_DESHAKE_OPENCL_KERNEL_H |
||
24 | |||
25 | #include "libavutil/opencl.h" |
||
26 | |||
27 | const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL( |
||
28 | |||
29 | inline unsigned char pixel(global const unsigned char *src, float x, float y, |
||
30 | int w, int h,int stride, unsigned char def) |
||
31 | { |
||
32 | return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[(int)x + (int)y * stride]; |
||
33 | } |
||
34 | unsigned char interpolate_nearest(float x, float y, global const unsigned char *src, |
||
35 | int width, int height, int stride, unsigned char def) |
||
36 | { |
||
37 | return pixel(src, (int)(x + 0.5), (int)(y + 0.5), width, height, stride, def); |
||
38 | } |
||
39 | |||
40 | unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src, |
||
41 | int width, int height, int stride, unsigned char def) |
||
42 | { |
||
43 | int x_c, x_f, y_c, y_f; |
||
44 | int v1, v2, v3, v4; |
||
45 | |||
46 | if (x < -1 || x > width || y < -1 || y > height) { |
||
47 | return def; |
||
48 | } else { |
||
49 | x_f = (int)x; |
||
50 | x_c = x_f + 1; |
||
51 | |||
52 | y_f = (int)y; |
||
53 | y_c = y_f + 1; |
||
54 | |||
55 | v1 = pixel(src, x_c, y_c, width, height, stride, def); |
||
56 | v2 = pixel(src, x_c, y_f, width, height, stride, def); |
||
57 | v3 = pixel(src, x_f, y_c, width, height, stride, def); |
||
58 | v4 = pixel(src, x_f, y_f, width, height, stride, def); |
||
59 | |||
60 | return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) + |
||
61 | v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y))); |
||
62 | } |
||
63 | } |
||
64 | |||
65 | unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src, |
||
66 | int width, int height, int stride, unsigned char def) |
||
67 | { |
||
68 | int x_c, x_f, y_c, y_f; |
||
69 | unsigned char v1, v2, v3, v4; |
||
70 | float f1, f2, f3, f4; |
||
71 | |||
72 | if (x < - 1 || x > width || y < -1 || y > height) |
||
73 | return def; |
||
74 | else { |
||
75 | x_f = (int)x; |
||
76 | x_c = x_f + 1; |
||
77 | y_f = (int)y; |
||
78 | y_c = y_f + 1; |
||
79 | |||
80 | v1 = pixel(src, x_c, y_c, width, height, stride, def); |
||
81 | v2 = pixel(src, x_c, y_f, width, height, stride, def); |
||
82 | v3 = pixel(src, x_f, y_c, width, height, stride, def); |
||
83 | v4 = pixel(src, x_f, y_f, width, height, stride, def); |
||
84 | |||
85 | f1 = 1 - sqrt((x_c - x) * (y_c - y)); |
||
86 | f2 = 1 - sqrt((x_c - x) * (y - y_f)); |
||
87 | f3 = 1 - sqrt((x - x_f) * (y_c - y)); |
||
88 | f4 = 1 - sqrt((x - x_f) * (y - y_f)); |
||
89 | return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4); |
||
90 | } |
||
91 | } |
||
92 | |||
93 | inline const float clipf(float a, float amin, float amax) |
||
94 | { |
||
95 | if (a < amin) return amin; |
||
96 | else if (a > amax) return amax; |
||
97 | else return a; |
||
98 | } |
||
99 | |||
100 | inline int mirror(int v, int m) |
||
101 | { |
||
102 | while ((unsigned)v > (unsigned)m) { |
||
103 | v = -v; |
||
104 | if (v < 0) |
||
105 | v += 2 * m; |
||
106 | } |
||
107 | return v; |
||
108 | } |
||
109 | |||
110 | kernel void avfilter_transform(global unsigned char *src, |
||
111 | global unsigned char *dst, |
||
112 | global float *matrix, |
||
113 | global float *matrix2, |
||
114 | int interpolate, |
||
115 | int fillmethod, |
||
116 | int src_stride_lu, |
||
117 | int dst_stride_lu, |
||
118 | int src_stride_ch, |
||
119 | int dst_stride_ch, |
||
120 | int height, |
||
121 | int width, |
||
122 | int ch, |
||
123 | int cw) |
||
124 | { |
||
125 | int global_id = get_global_id(0); |
||
126 | |||
127 | global unsigned char *dst_y = dst; |
||
128 | global unsigned char *dst_u = dst_y + height * dst_stride_lu; |
||
129 | global unsigned char *dst_v = dst_u + ch * dst_stride_ch; |
||
130 | |||
131 | global unsigned char *src_y = src; |
||
132 | global unsigned char *src_u = src_y + height * src_stride_lu; |
||
133 | global unsigned char *src_v = src_u + ch * src_stride_ch; |
||
134 | |||
135 | global unsigned char *tempdst; |
||
136 | global unsigned char *tempsrc; |
||
137 | |||
138 | int x; |
||
139 | int y; |
||
140 | float x_s; |
||
141 | float y_s; |
||
142 | int tempsrc_stride; |
||
143 | int tempdst_stride; |
||
144 | int temp_height; |
||
145 | int temp_width; |
||
146 | int curpos; |
||
147 | unsigned char def = 0; |
||
148 | if (global_id < width*height) { |
||
149 | y = global_id/width; |
||
150 | x = global_id%width; |
||
151 | x_s = x * matrix[0] + y * matrix[1] + matrix[2]; |
||
152 | y_s = x * matrix[3] + y * matrix[4] + matrix[5]; |
||
153 | tempdst = dst_y; |
||
154 | tempsrc = src_y; |
||
155 | tempsrc_stride = src_stride_lu; |
||
156 | tempdst_stride = dst_stride_lu; |
||
157 | temp_height = height; |
||
158 | temp_width = width; |
||
159 | } else if ((global_id >= width*height)&&(global_id < width*height + ch*cw)) { |
||
160 | y = (global_id - width*height)/cw; |
||
161 | x = (global_id - width*height)%cw; |
||
162 | x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2]; |
||
163 | y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5]; |
||
164 | tempdst = dst_u; |
||
165 | tempsrc = src_u; |
||
166 | tempsrc_stride = src_stride_ch; |
||
167 | tempdst_stride = dst_stride_ch; |
||
168 | temp_height = ch; |
||
169 | temp_width = cw; |
||
170 | } else { |
||
171 | y = (global_id - width*height - ch*cw)/cw; |
||
172 | x = (global_id - width*height - ch*cw)%cw; |
||
173 | x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2]; |
||
174 | y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5]; |
||
175 | tempdst = dst_v; |
||
176 | tempsrc = src_v; |
||
177 | tempsrc_stride = src_stride_ch; |
||
178 | tempdst_stride = dst_stride_ch; |
||
179 | temp_height = ch; |
||
180 | temp_width = cw; |
||
181 | } |
||
182 | curpos = y * tempdst_stride + x; |
||
183 | switch (fillmethod) { |
||
184 | case 0: //FILL_BLANK |
||
185 | def = 0; |
||
186 | break; |
||
187 | case 1: //FILL_ORIGINAL |
||
188 | def = tempsrc[y*tempsrc_stride+x]; |
||
189 | break; |
||
190 | case 2: //FILL_CLAMP |
||
191 | y_s = clipf(y_s, 0, temp_height - 1); |
||
192 | x_s = clipf(x_s, 0, temp_width - 1); |
||
193 | def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s]; |
||
194 | break; |
||
195 | case 3: //FILL_MIRROR |
||
196 | y_s = mirror(y_s,temp_height - 1); |
||
197 | x_s = mirror(x_s,temp_width - 1); |
||
198 | def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s]; |
||
199 | break; |
||
200 | } |
||
201 | switch (interpolate) { |
||
202 | case 0: //INTERPOLATE_NEAREST |
||
203 | tempdst[curpos] = interpolate_nearest(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); |
||
204 | break; |
||
205 | case 1: //INTERPOLATE_BILINEAR |
||
206 | tempdst[curpos] = interpolate_bilinear(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); |
||
207 | break; |
||
208 | case 2: //INTERPOLATE_BIQUADRATIC |
||
209 | tempdst[curpos] = interpolate_biquadratic(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def); |
||
210 | break; |
||
211 | default: |
||
212 | return; |
||
213 | } |
||
214 | } |
||
215 | ); |
||
216 | |||
217 | #endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */>>>>>>>>>> |