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  * 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 
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 */
#define AV_OPENCL_KERNEL(...)
Definition: opencl.h:46
float w
float x
static void interpolate(float *out, float v1, float v2, int size)
Definition: twinvq.c:84
#define height
#define width
#define src
Definition: vp9dsp.c:530
float y
uint8_t pixel
Definition: tiny_ssim.c:42
float z
GLint GLenum GLboolean GLsizei stride
Definition: opengl_enc.c:105
if(ret< 0)
Definition: vf_mcdeint.c:282
uint8_t pi<< 24) CONV_FUNC(AV_SAMPLE_FMT_S64, int64_t, AV_SAMPLE_FMT_U8,(uint64_t)((*(constuint8_t *) pi-0x80U))<< 56) CONV_FUNC(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_U8,(*(constuint8_t *) pi-0x80)*(1.0f/(1<< 7))) CONV_FUNC(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_U8,(*(constuint8_t *) pi-0x80)*(1.0/(1<< 7))) CONV_FUNC(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S16,(*(constint16_t *) pi >>8)+0x80) CONV_FUNC(AV_SAMPLE_FMT_S64, int64_t, AV_SAMPLE_FMT_S16,(uint64_t)(*(constint16_t *) pi)<< 48) CONV_FUNC(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S16,*(constint16_t *) pi *(1.0f/(1<< 15))) CONV_FUNC(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S16,*(constint16_t *) pi *(1.0/(1<< 15))) CONV_FUNC(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S32,(*(constint32_t *) pi >>24)+0x80) CONV_FUNC(AV_SAMPLE_FMT_S64, int64_t, AV_SAMPLE_FMT_S32,(uint64_t)(*(constint32_t *) pi)<< 32) CONV_FUNC(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S32,*(constint32_t *) pi *(1.0f/(1U<< 31))) CONV_FUNC(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S32,*(constint32_t *) pi *(1.0/(1U<< 31))) CONV_FUNC(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S64,(*(constint64_t *) pi >>56)+0x80) CONV_FUNC(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S64,*(constint64_t *) pi *(1.0f/(INT64_C(1)<< 63))) CONV_FUNC(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S64,*(constint64_t *) pi *(1.0/(INT64_C(1)<< 63))) CONV_FUNC(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_FLT, av_clip_uint8(lrintf(*(constfloat *) pi *(1<< 7))+0x80)) CONV_FUNC(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_FLT, av_clip_int16(lrintf(*(constfloat *) pi *(1<< 15)))) CONV_FUNC(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_FLT, av_clipl_int32(llrintf(*(constfloat *) pi *(1U<< 31)))) CONV_FUNC(AV_SAMPLE_FMT_S64, int64_t, AV_SAMPLE_FMT_FLT, llrintf(*(constfloat *) pi *(INT64_C(1)<< 63))) CONV_FUNC(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_DBL, av_clip_uint8(lrint(*(constdouble *) pi *(1<< 7))+0x80)) CONV_FUNC(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_DBL, av_clip_int16(lrint(*(constdouble *) pi *(1<< 15)))) CONV_FUNC(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_DBL, av_clipl_int32(llrint(*(constdouble *) pi *(1U<< 31)))) CONV_FUNC(AV_SAMPLE_FMT_S64, int64_t, AV_SAMPLE_FMT_DBL, llrint(*(constdouble *) pi *(INT64_C(1)<< 63)))#defineFMT_PAIR_FUNC(out, in) staticconv_func_type *constfmt_pair_to_conv_functions[AV_SAMPLE_FMT_NB *AV_SAMPLE_FMT_NB]={FMT_PAIR_FUNC(AV_SAMPLE_FMT_U8, AV_SAMPLE_FMT_U8), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S16, AV_SAMPLE_FMT_U8), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S32, AV_SAMPLE_FMT_U8), FMT_PAIR_FUNC(AV_SAMPLE_FMT_FLT, AV_SAMPLE_FMT_U8), FMT_PAIR_FUNC(AV_SAMPLE_FMT_DBL, AV_SAMPLE_FMT_U8), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S64, AV_SAMPLE_FMT_U8), FMT_PAIR_FUNC(AV_SAMPLE_FMT_U8, AV_SAMPLE_FMT_S16), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S16, AV_SAMPLE_FMT_S16), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S32, AV_SAMPLE_FMT_S16), FMT_PAIR_FUNC(AV_SAMPLE_FMT_FLT, AV_SAMPLE_FMT_S16), FMT_PAIR_FUNC(AV_SAMPLE_FMT_DBL, AV_SAMPLE_FMT_S16), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S64, AV_SAMPLE_FMT_S16), FMT_PAIR_FUNC(AV_SAMPLE_FMT_U8, AV_SAMPLE_FMT_S32), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S16, AV_SAMPLE_FMT_S32), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S32, AV_SAMPLE_FMT_S32), FMT_PAIR_FUNC(AV_SAMPLE_FMT_FLT, AV_SAMPLE_FMT_S32), FMT_PAIR_FUNC(AV_SAMPLE_FMT_DBL, AV_SAMPLE_FMT_S32), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S64, AV_SAMPLE_FMT_S32), FMT_PAIR_FUNC(AV_SAMPLE_FMT_U8, AV_SAMPLE_FMT_FLT), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S16, AV_SAMPLE_FMT_FLT), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S32, AV_SAMPLE_FMT_FLT), FMT_PAIR_FUNC(AV_SAMPLE_FMT_FLT, AV_SAMPLE_FMT_FLT), FMT_PAIR_FUNC(AV_SAMPLE_FMT_DBL, AV_SAMPLE_FMT_FLT), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S64, AV_SAMPLE_FMT_FLT), FMT_PAIR_FUNC(AV_SAMPLE_FMT_U8, AV_SAMPLE_FMT_DBL), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S16, AV_SAMPLE_FMT_DBL), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S32, AV_SAMPLE_FMT_DBL), FMT_PAIR_FUNC(AV_SAMPLE_FMT_FLT, AV_SAMPLE_FMT_DBL), FMT_PAIR_FUNC(AV_SAMPLE_FMT_DBL, AV_SAMPLE_FMT_DBL), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S64, AV_SAMPLE_FMT_DBL), FMT_PAIR_FUNC(AV_SAMPLE_FMT_U8, AV_SAMPLE_FMT_S64), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S16, AV_SAMPLE_FMT_S64), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S32, AV_SAMPLE_FMT_S64), FMT_PAIR_FUNC(AV_SAMPLE_FMT_FLT, AV_SAMPLE_FMT_S64), FMT_PAIR_FUNC(AV_SAMPLE_FMT_DBL, AV_SAMPLE_FMT_S64), FMT_PAIR_FUNC(AV_SAMPLE_FMT_S64, AV_SAMPLE_FMT_S64),};staticvoidcpy1(uint8_t **dst, constuint8_t **src, intlen){memcpy(*dst,*src, len);}staticvoidcpy2(uint8_t **dst, constuint8_t **src, intlen){memcpy(*dst,*src, 2 *len);}staticvoidcpy4(uint8_t **dst, constuint8_t **src, intlen){memcpy(*dst,*src, 4 *len);}staticvoidcpy8(uint8_t **dst, constuint8_t **src, intlen){memcpy(*dst,*src, 8 *len);}AudioConvert *swri_audio_convert_alloc(enumAVSampleFormatout_fmt, enumAVSampleFormatin_fmt, intchannels, constint *ch_map, intflags){AudioConvert *ctx;conv_func_type *f=fmt_pair_to_conv_functions[av_get_packed_sample_fmt(out_fmt)+AV_SAMPLE_FMT_NB *av_get_packed_sample_fmt(in_fmt)];if(!f) returnNULL;ctx=av_mallocz(sizeof(*ctx));if(!ctx) returnNULL;if(channels==1){in_fmt=av_get_planar_sample_fmt(in_fmt);out_fmt=av_get_planar_sample_fmt(out_fmt);}ctx->channels=channels;ctx->conv_f=f;ctx->ch_map=ch_map;if(in_fmt==AV_SAMPLE_FMT_U8||in_fmt==AV_SAMPLE_FMT_U8P) memset(ctx->silence, 0x80, sizeof(ctx->silence));if(out_fmt==in_fmt &&!ch_map){switch(av_get_bytes_per_sample(in_fmt)){case1:ctx->simd_f=cpy1;break;case2:ctx->simd_f=cpy2;break;case4:ctx->simd_f=cpy4;break;case8:ctx->simd_f=cpy8;break;}}if(HAVE_YASM &&1) swri_audio_convert_init_x86(ctx, out_fmt, in_fmt, channels);if(ARCH_ARM) swri_audio_convert_init_arm(ctx, out_fmt, in_fmt, channels);if(ARCH_AARCH64) swri_audio_convert_init_aarch64(ctx, out_fmt, in_fmt, channels);returnctx;}voidswri_audio_convert_free(AudioConvert **ctx){av_freep(ctx);}intswri_audio_convert(AudioConvert *ctx, AudioData *out, AudioData *in, intlen){intch;intoff=0;constintos=(out->planar?1:out->ch_count)*out->bps;unsignedmisaligned=0;av_assert0(ctx->channels==out->ch_count);if(ctx->in_simd_align_mask){intplanes=in->planar?in->ch_count:1;unsignedm=0;for(ch=0;ch< planes;ch++) m|=(intptr_t) in->ch[ch];misaligned|=m &ctx->in_simd_align_mask;}if(ctx->out_simd_align_mask){intplanes=out->planar?out->ch_count:1;unsignedm=0;for(ch=0;ch< planes;ch++) m|=(intptr_t) out->ch[ch];misaligned|=m &ctx->out_simd_align_mask;}if(ctx->simd_f &&!ctx->ch_map &&!misaligned){off=len &~15;av_assert1(off >=0);av_assert1(off<=len);av_assert2(ctx->channels==SWR_CH_MAX||!in->ch[ctx->channels]);if(off >0){if(out->planar==in->planar){intplanes=out->planar?out->ch_count:1;for(ch=0;ch< planes;ch++){ctx->simd_f(out-> ch ch
Definition: audioconvert.c:56
OpenCL wrapper.
const char * ff_kernel_deshake_opencl