diff options
| author | Danil Iashchenko <danyaschenko@gmail.com> | 2018-06-30 03:46:24 +0300 |
|---|---|---|
| committer | Mark Thompson <sw@jkqxz.net> | 2018-07-03 00:11:28 +0100 |
| commit | 581bafa83b9cb4f438940e0b10a9f9832ebd57b6 (patch) | |
| tree | c1350d31be655a8912f698cac846cda43e6e03cd /libavfilter/opencl | |
| parent | 54b425a7fad3e35ebb4757108fef5e3ee0e3d957 (diff) | |
| download | ffmpeg-streaming-581bafa83b9cb4f438940e0b10a9f9832ebd57b6.zip ffmpeg-streaming-581bafa83b9cb4f438940e0b10a9f9832ebd57b6.tar.gz | |
lavfi: add sobel, prewitt, roberts filters
Add opencl version of sobel, prewitt, roberts filters.
Diffstat (limited to 'libavfilter/opencl')
| -rw-r--r-- | libavfilter/opencl/convolution.cl | 82 |
1 files changed, 82 insertions, 0 deletions
diff --git a/libavfilter/opencl/convolution.cl b/libavfilter/opencl/convolution.cl index 03ef4ef..815c779 100644 --- a/libavfilter/opencl/convolution.cl +++ b/libavfilter/opencl/convolution.cl @@ -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); +} |
