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 */