FFmpeg
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups Pages
deshake_opencl_kernel.h
Go to the documentation of this file.
1 /*
2  * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
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 
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 */