Add the first OpenCL shader for extensional VPP

Signed-off-by: Zhao Yakui <yakui.zhao@intel.com>

1 files changed, 39 insertions, 0 deletions

diff --git a/src/opencl/climage-mirror.cl b/src/opencl/climage-mirror.cl
new file mode 100644
index 0000000..f1e5cc7
--- /dev/null
+++ b/src/opencl/climage-mirror.cl
@@ -0,0 +1,39 @@
+struct vpp_image_dim {
+    unsigned int width;
+    unsigned int height;
+};
+
+struct vpp_rect_region {
+    unsigned int x;
+    unsigned int y;
+    unsigned int width;
+    unsigned int height;
+};
+
+__kernel void
+compiler_mirror_effect(__read_only image2d_t src_y,
+            __read_only image2d_t src_uv,
+            struct vpp_image_dim src_dim,
+            struct vpp_rect_region src_region,
+            __write_only image2d_t dst_y,
+            __write_only image2d_t dst_uv,
+            struct vpp_image_dim dst_dim,
+            struct vpp_rect_region dst_region)
+{
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST;
+    int2 loc = (int2)(get_global_id(0), get_global_id(1));
+    uint4 color_y, color_uv;
+
+    if (loc.y < src_dim.height/2) {
+        color_y = read_imageui(src_y, sampler, loc);
+        color_uv = read_imageui(src_uv, sampler, loc/2);
+    } else {
+        int2 newloc = (int2)(loc.x, src_dim.height - loc.y);
+        color_y = read_imageui(src_y, sampler, newloc);
+        color_uv = read_imageui(src_uv, sampler, newloc/2);
+    }
+
+    write_imageui(dst_y, loc, color_y);
+    write_imageui(dst_uv, loc/2, color_uv);
+}
+