libavfilter/deshake_opencl_kernel.h

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