| 90793591 |
/*
* Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com> |
| d943330d |
* Copyright (C) 2013 Lenny Wang |
| 90793591 |
*
*
* This file is part of FFmpeg.
*
* FFmpeg is free software; you can redistribute it and/or
* modify it under the terms of the GNU Lesser General Public
* License as published by the Free Software Foundation; either
* version 2.1 of the License, or (at your option) any later version.
*
* FFmpeg is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
* Lesser General Public License for more details.
*
* You should have received a copy of the GNU Lesser General Public
* License along with FFmpeg; if not, write to the Free Software
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
|
| 63e8fc4d |
#ifndef AVFILTER_DESHAKE_OPENCL_KERNEL_H
#define AVFILTER_DESHAKE_OPENCL_KERNEL_H |
| 90793591 |
#include "libavutil/opencl.h"
const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL( |
| d943330d |
inline unsigned char pixel(global const unsigned char *src, int x, int y, |
| 90793591 |
int w, int h,int stride, unsigned char def)
{ |
| d943330d |
return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[x + y * stride]; |
| 90793591 |
} |
| d943330d |
|
| 90793591 |
unsigned char interpolate_nearest(float x, float y, global const unsigned char *src,
int width, int height, int stride, unsigned char def)
{ |
| d943330d |
return pixel(src, (int)(x + 0.5f), (int)(y + 0.5f), width, height, stride, def); |
| 90793591 |
}
unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src,
int width, int height, int stride, unsigned char def)
{
int x_c, x_f, y_c, y_f;
int v1, v2, v3, v4; |
| d943330d |
x_f = (int)x;
y_f = (int)y;
x_c = x_f + 1;
y_c = y_f + 1; |
| 90793591 |
|
| d943330d |
if (x_f < -1 || x_f > width || y_f < -1 || y_f > height) { |
| 90793591 |
return def;
} else { |
| d943330d |
v4 = pixel(src, x_f, y_f, width, height, stride, def); |
| 90793591 |
v2 = pixel(src, x_c, y_f, width, height, stride, def);
v3 = pixel(src, x_f, y_c, width, height, stride, def); |
| d943330d |
v1 = pixel(src, x_c, y_c, width, height, stride, def); |
| 90793591 |
return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) +
v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y)));
}
}
unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src,
int width, int height, int stride, unsigned char def)
{
int x_c, x_f, y_c, y_f;
unsigned char v1, v2, v3, v4;
float f1, f2, f3, f4; |
| d943330d |
x_f = (int)x;
y_f = (int)y;
x_c = x_f + 1;
y_c = y_f + 1; |
| 90793591 |
|
| d943330d |
if (x_f < - 1 || x_f > width || y_f < -1 || y_f > height) |
| 90793591 |
return def;
else { |
| d943330d |
v4 = pixel(src, x_f, y_f, width, height, stride, def); |
| 90793591 |
v2 = pixel(src, x_c, y_f, width, height, stride, def);
v3 = pixel(src, x_f, y_c, width, height, stride, def); |
| d943330d |
v1 = pixel(src, x_c, y_c, width, height, stride, def); |
| 90793591 |
f1 = 1 - sqrt((x_c - x) * (y_c - y));
f2 = 1 - sqrt((x_c - x) * (y - y_f));
f3 = 1 - sqrt((x - x_f) * (y_c - y));
f4 = 1 - sqrt((x - x_f) * (y - y_f));
return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4);
}
}
inline const float clipf(float a, float amin, float amax)
{
if (a < amin) return amin;
else if (a > amax) return amax;
else return a;
}
inline int mirror(int v, int m)
{
while ((unsigned)v > (unsigned)m) {
v = -v;
if (v < 0)
v += 2 * m;
}
return v;
}
|
| d943330d |
kernel void avfilter_transform_luma(global unsigned char *src,
global unsigned char *dst,
float4 matrix,
int interpolate,
int fill,
int src_stride_lu,
int dst_stride_lu,
int height,
int width) |
| 90793591 |
{ |
| d943330d |
int x = get_global_id(0);
int y = get_global_id(1);
int idx_dst = y * dst_stride_lu + x;
unsigned char def = 0;
float x_s = x * matrix.x + y * matrix.y + matrix.z;
float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
if (x < width && y < height) {
switch (fill) {
case 0: //FILL_BLANK
def = 0;
break;
case 1: //FILL_ORIGINAL
def = src[y*src_stride_lu + x];
break;
case 2: //FILL_CLAMP
y_s = clipf(y_s, 0, height - 1);
x_s = clipf(x_s, 0, width - 1);
def = src[(int)y_s * src_stride_lu + (int)x_s];
break;
case 3: //FILL_MIRROR
y_s = mirror(y_s, height - 1);
x_s = mirror(x_s, width - 1);
def = src[(int)y_s * src_stride_lu + (int)x_s];
break;
}
switch (interpolate) {
case 0: //INTERPOLATE_NEAREST
dst[idx_dst] = interpolate_nearest(x_s, y_s, src, width, height, src_stride_lu, def);
break;
case 1: //INTERPOLATE_BILINEAR
dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, width, height, src_stride_lu, def);
break;
case 2: //INTERPOLATE_BIQUADRATIC
dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, width, height, src_stride_lu, def);
break;
default:
return;
} |
| 90793591 |
} |
| d943330d |
}
kernel void avfilter_transform_chroma(global unsigned char *src,
global unsigned char *dst,
float4 matrix,
int interpolate,
int fill,
int src_stride_lu,
int dst_stride_lu,
int src_stride_ch,
int dst_stride_ch,
int height,
int width,
int ch,
int cw)
{
int x = get_global_id(0);
int y = get_global_id(1);
int pad_ch = get_global_size(1)>>1;
global unsigned char *dst_u = dst + height * dst_stride_lu;
global unsigned char *src_u = src + height * src_stride_lu;
global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
global unsigned char *src_v = src_u + ch * src_stride_ch;
src = y < pad_ch ? src_u : src_v;
dst = y < pad_ch ? dst_u : dst_v;
y = select(y - pad_ch, y, y < pad_ch);
float x_s = x * matrix.x + y * matrix.y + matrix.z;
float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
int idx_dst = y * dst_stride_ch + x;
unsigned char def;
if (x < cw && y < ch) {
switch (fill) {
case 0: //FILL_BLANK
def = 0;
break;
case 1: //FILL_ORIGINAL
def = src[y*src_stride_ch + x];
break;
case 2: //FILL_CLAMP
y_s = clipf(y_s, 0, ch - 1);
x_s = clipf(x_s, 0, cw - 1);
def = src[(int)y_s * src_stride_ch + (int)x_s];
break;
case 3: //FILL_MIRROR
y_s = mirror(y_s, ch - 1);
x_s = mirror(x_s, cw - 1);
def = src[(int)y_s * src_stride_ch + (int)x_s];
break;
}
switch (interpolate) {
case 0: //INTERPOLATE_NEAREST
dst[idx_dst] = interpolate_nearest(x_s, y_s, src, cw, ch, src_stride_ch, def);
break;
case 1: //INTERPOLATE_BILINEAR
dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, cw, ch, src_stride_ch, def);
break;
case 2: //INTERPOLATE_BIQUADRATIC
dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, cw, ch, src_stride_ch, def);
break;
default:
return;
} |
| 90793591 |
}
}
);
|
| 63e8fc4d |
#endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */ |