Skip to content
Permalink
Browse files

lavfi: add sobel, prewitt, roberts filters

Add opencl version of sobel, prewitt, roberts filters.
  • Loading branch information...
nafarya authored and fhvwy committed Jun 30, 2018
1 parent 54b425a commit 581bafa83b9cb4f438940e0b10a9f9832ebd57b6
Showing with 283 additions and 46 deletions.
  1. +3 −0 configure
  2. +7 −1 libavfilter/Makefile
  3. +3 −0 libavfilter/allfilters.c
  4. +82 −0 libavfilter/opencl/convolution.cl
  5. +188 −45 libavfilter/vf_convolution_opencl.c
@@ -3372,12 +3372,14 @@ perspective_filter_deps="gpl"
phase_filter_deps="gpl"
pp7_filter_deps="gpl"
pp_filter_deps="gpl postproc"
prewitt_opencl_filter_deps="opencl"
procamp_vaapi_filter_deps="vaapi VAProcPipelineParameterBuffer"
program_opencl_filter_deps="opencl"
pullup_filter_deps="gpl"
removelogo_filter_deps="avcodec avformat swscale"
repeatfields_filter_deps="gpl"
resample_filter_deps="avresample"
roberts_opencl_filter_deps="opencl"
rubberband_filter_deps="librubberband"
sab_filter_deps="gpl swscale"
scale2ref_filter_deps="swscale"
@@ -3396,6 +3398,7 @@ showspectrumpic_filter_deps="avcodec"
showspectrumpic_filter_select="fft"
signature_filter_deps="gpl avcodec avformat"
smartblur_filter_deps="gpl swscale"
sobel_opencl_filter_deps="opencl"
sofalizer_filter_deps="libmysofa avcodec"
sofalizer_filter_select="fft"
spectrumsynth_filter_deps="avcodec"
@@ -171,7 +171,7 @@ OBJS-$(CONFIG_COLORMATRIX_FILTER) += vf_colormatrix.o
OBJS-$(CONFIG_COLORSPACE_FILTER) += vf_colorspace.o colorspace.o colorspacedsp.o
OBJS-$(CONFIG_CONVOLUTION_FILTER) += vf_convolution.o
OBJS-$(CONFIG_CONVOLUTION_OPENCL_FILTER) += vf_convolution_opencl.o opencl.o \
opencl/convolution.o
opencl/convolution.o
OBJS-$(CONFIG_CONVOLVE_FILTER) += vf_convolve.o framesync.o
OBJS-$(CONFIG_COPY_FILTER) += vf_copy.o
OBJS-$(CONFIG_COREIMAGE_FILTER) += vf_coreimage.o
@@ -294,6 +294,8 @@ OBJS-$(CONFIG_PP_FILTER) += vf_pp.o
OBJS-$(CONFIG_PP7_FILTER) += vf_pp7.o
OBJS-$(CONFIG_PREMULTIPLY_FILTER) += vf_premultiply.o framesync.o
OBJS-$(CONFIG_PREWITT_FILTER) += vf_convolution.o
OBJS-$(CONFIG_PREWITT_OPENCL_FILTER) += vf_convolution_opencl.o opencl.o \
opencl/convolution.o
OBJS-$(CONFIG_PROCAMP_VAAPI_FILTER) += vf_procamp_vaapi.o vaapi_vpp.o
OBJS-$(CONFIG_PROGRAM_OPENCL_FILTER) += vf_program_opencl.o opencl.o framesync.o
OBJS-$(CONFIG_PSEUDOCOLOR_FILTER) += vf_pseudocolor.o
@@ -310,6 +312,8 @@ OBJS-$(CONFIG_REMOVELOGO_FILTER) += bbox.o lswsutils.o lavfutils.o v
OBJS-$(CONFIG_REPEATFIELDS_FILTER) += vf_repeatfields.o
OBJS-$(CONFIG_REVERSE_FILTER) += f_reverse.o
OBJS-$(CONFIG_ROBERTS_FILTER) += vf_convolution.o
OBJS-$(CONFIG_ROBERTS_OPENCL_FILTER) += vf_convolution_opencl.o opencl.o \
opencl/convolution.o
OBJS-$(CONFIG_ROTATE_FILTER) += vf_rotate.o
OBJS-$(CONFIG_SAB_FILTER) += vf_sab.o
OBJS-$(CONFIG_SCALE_FILTER) += vf_scale.o scale.o
@@ -338,6 +342,8 @@ OBJS-$(CONFIG_SIGNALSTATS_FILTER) += vf_signalstats.o
OBJS-$(CONFIG_SIGNATURE_FILTER) += vf_signature.o
OBJS-$(CONFIG_SMARTBLUR_FILTER) += vf_smartblur.o
OBJS-$(CONFIG_SOBEL_FILTER) += vf_convolution.o
OBJS-$(CONFIG_SOBEL_OPENCL_FILTER) += vf_convolution_opencl.o opencl.o \
opencl/convolution.o
OBJS-$(CONFIG_SPLIT_FILTER) += split.o
OBJS-$(CONFIG_SPP_FILTER) += vf_spp.o
OBJS-$(CONFIG_SR_FILTER) += vf_sr.o
@@ -282,6 +282,7 @@ extern AVFilter ff_vf_pp;
extern AVFilter ff_vf_pp7;
extern AVFilter ff_vf_premultiply;
extern AVFilter ff_vf_prewitt;
extern AVFilter ff_vf_prewitt_opencl;
extern AVFilter ff_vf_procamp_vaapi;
extern AVFilter ff_vf_program_opencl;
extern AVFilter ff_vf_pseudocolor;
@@ -298,6 +299,7 @@ extern AVFilter ff_vf_removelogo;
extern AVFilter ff_vf_repeatfields;
extern AVFilter ff_vf_reverse;
extern AVFilter ff_vf_roberts;
extern AVFilter ff_vf_roberts_opencl;
extern AVFilter ff_vf_rotate;
extern AVFilter ff_vf_sab;
extern AVFilter ff_vf_scale;
@@ -326,6 +328,7 @@ extern AVFilter ff_vf_signalstats;
extern AVFilter ff_vf_signature;
extern AVFilter ff_vf_smartblur;
extern AVFilter ff_vf_sobel;
extern AVFilter ff_vf_sobel_opencl;
extern AVFilter ff_vf_split;
extern AVFilter ff_vf_spp;
extern AVFilter ff_vf_sr;
@@ -43,3 +43,85 @@ __kernel void convolution_global(__write_only image2d_t dst,
float4 dstPix = convPix * div + bias;
write_imagef(dst, loc, dstPix);
}


__kernel void sobel_global(__write_only image2d_t dst,
__read_only image2d_t src,
float div,
float bias)
{
const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST);

int2 loc = (int2)(get_global_id(0), get_global_id(1));

float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * -1 +
read_imagef(src, sampler, loc + (int2)( 0,-1)) * -2 +
read_imagef(src, sampler, loc + (int2)( 1,-1)) * -1 +
read_imagef(src, sampler, loc + (int2)(-1, 1)) * 1 +
read_imagef(src, sampler, loc + (int2)( 0, 1)) * 2 +
read_imagef(src, sampler, loc + (int2)( 1, 1)) * 1;

float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * -1 +
read_imagef(src, sampler, loc + (int2)(-1, 0)) * -2 +
read_imagef(src, sampler, loc + (int2)(-1, 1)) * -1 +
read_imagef(src, sampler, loc + (int2)( 1,-1)) * 1 +
read_imagef(src, sampler, loc + (int2)( 1, 0)) * 2 +
read_imagef(src, sampler, loc + (int2)( 1, 1)) * 1;

float4 dstPix = hypot(sum1, sum2) * div + bias;
write_imagef(dst, loc, dstPix);
}

__kernel void prewitt_global(__write_only image2d_t dst,
__read_only image2d_t src,
float div,
float bias)
{
const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST);

int2 loc = (int2)(get_global_id(0), get_global_id(1));

float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +
read_imagef(src, sampler, loc + (int2)( 0,-1)) * 1 +
read_imagef(src, sampler, loc + (int2)( 1,-1)) * 1 +
read_imagef(src, sampler, loc + (int2)(-1, 1)) * -1 +
read_imagef(src, sampler, loc + (int2)( 0, 1)) * -1 +
read_imagef(src, sampler, loc + (int2)( 1, 1)) * -1;

float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +
read_imagef(src, sampler, loc + (int2)(-1, 0)) * 1 +
read_imagef(src, sampler, loc + (int2)(-1, 1)) * 1 +
read_imagef(src, sampler, loc + (int2)( 1,-1)) * -1 +
read_imagef(src, sampler, loc + (int2)( 1, 0)) * -1 +
read_imagef(src, sampler, loc + (int2)( 1, 1)) * -1;

float4 dstPix = hypot(sum1, sum2) * div + bias;
write_imagef(dst, loc, dstPix);
}

__kernel void roberts_global(__write_only image2d_t dst,
__read_only image2d_t src,
float div,
float bias)
{
const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST);

int2 loc = (int2)(get_global_id(0), get_global_id(1));

float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +
read_imagef(src, sampler, loc + (int2)( 0,-1)) * -1;


float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1, 0)) * -1 +
read_imagef(src, sampler, loc + (int2)( 0, 0)) * 1;


float4 dstPix = hypot(sum1, sum2) * div + bias;
write_imagef(dst, loc, dstPix);
}
Oops, something went wrong.

0 comments on commit 581bafa

Please sign in to comment.
You can’t perform that action at this time.