Merge commit '7e929dac100916fc45cb95e231025f3439c20156'

[ffmpeg] / libavfilter / opencl / convolution.cl

diff --git a/libavfilter/opencl/convolution.cl b/libavfilter/opencl/convolution.cl
index 03ef4eff1bc7236f9f9d7defea700cc982e01a31..815c7792725923bf7eda148da7aa0dd2bb26791f 100644 (file)
--- 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);
+}