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 28const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL( 29inline 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 35unsigned 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 41unsigned 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 63unsigned 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 90inline 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 97inline 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 107kernel 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 159kernel 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 */ 226