[gegl] video-degradation: add opencl support

[Øyvind Kolås <ok src gnome org>]

commit 4c5ef4b52287093dc0549fa99a79d78aa3bac1e9
Author: Nanley Chery <nchery17 gmail com>
Date:   Mon Nov 17 02:02:24 2014 -0500

    video-degradation: add opencl support

 opencl/video-degradation.cl           |   31 +++++++++++++++++++
 opencl/video-degradation.cl.h         |   33 +++++++++++++++++++++
 operations/common/video-degradation.c |   52 ++++++++++++++++++++++++++++++++-
 3 files changed, 115 insertions(+), 1 deletions(-)
---
diff --git a/opencl/video-degradation.cl b/opencl/video-degradation.cl
new file mode 100644
index 0000000..445153c
--- /dev/null
+++ b/opencl/video-degradation.cl
@@ -0,0 +1,31 @@
+__kernel void gegl_video_degradation (__global const float4 *input,
+                                      __global float4 *output,
+                                      __global const int *pattern,
+                                      const int pat_w,
+                                      const int pat_h,
+                                      const int additive,
+                                      const int rotated)
+{
+  const size_t gidx   = get_global_id(0);
+  const size_t gidy   = get_global_id(1);
+  const size_t gid    = gidx - get_global_offset(0) +
+                       (gidy - get_global_offset(1)) *
+                        get_global_size(0);
+  const float4 indata = input[gid];
+
+  /* Get channel to keep in this input pixel */
+  const int sel_b = pattern[rotated ? pat_w * (gidx % pat_h) + gidy % pat_w:
+                                      pat_w * (gidy % pat_h) + gidx % pat_w];
+
+  /* Mask channels according to sel_b variable */
+  float4 value = select(0.f, indata, sel_b == (int4)(0, 1, 2, 3));
+
+  /* Add original pixel if enabled */
+  if (additive)
+      value = fmin(value + indata, 1.0f);
+
+  /* Maintain original alpha channel */
+  value.w = indata.w;
+
+  output[gid] = value;
+}
diff --git a/opencl/video-degradation.cl.h b/opencl/video-degradation.cl.h
new file mode 100644
index 0000000..d3fa625
--- /dev/null
+++ b/opencl/video-degradation.cl.h
@@ -0,0 +1,33 @@
+static const char* video_degradation_cl_source =
+"__kernel void gegl_video_degradation (__global const float4 *input,           \n"
+"                                      __global float4 *output,                \n"
+"                                      __global const int *pattern,            \n"
+"                                      const int pat_w,                        \n"
+"                                      const int pat_h,                        \n"
+"                                      const int additive,                     \n"
+"                                      const int rotated)                      \n"
+"{                                                                             \n"
+"  const size_t gidx   = get_global_id(0);                                     \n"
+"  const size_t gidy   = get_global_id(1);                                     \n"
+"  const size_t gid    = gidx - get_global_offset(0) +                         \n"
+"                       (gidy - get_global_offset(1)) *                        \n"
+"                        get_global_size(0);                                   \n"
+"  const float4 indata = input[gid];                                           \n"
+"                                                                              \n"
+"  /* Get channel to keep in this input pixel */                               \n"
+"  const int sel_b = pattern[rotated ? pat_w * (gidx % pat_h) + gidy % pat_w:  \n"
+"                                      pat_w * (gidy % pat_h) + gidx % pat_w]; \n"
+"                                                                              \n"
+"  /* Mask channels according to sel_b variable */                             \n"
+"  float4 value = select(0.f, indata, sel_b == (int4)(0, 1, 2, 3));            \n"
+"                                                                              \n"
+"  /* Add original pixel if enabled */                                         \n"
+"  if (additive)                                                               \n"
+"      value = fmin(value + indata, 1.0f);                                     \n"
+"                                                                              \n"
+"  /* Maintain original alpha channel */                                       \n"
+"  value.w = indata.w;                                                         \n"
+"                                                                              \n"
+"  output[gid] = value;                                                        \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/video-degradation.c b/operations/common/video-degradation.c
index bd40187..29197ff 100644
--- a/operations/common/video-degradation.c
+++ b/operations/common/video-degradation.c
@@ -187,6 +187,53 @@ prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "output", format);
 }
 
+static gboolean
+cl_process (GeglOperation       *operation,
+            cl_mem               in_buf,
+            cl_mem               out_buf,
+            const size_t         n_pixels,
+            const GeglRectangle *roi,
+            gint                 level)
+{
+  GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
+  GeglClRunData *cl_data = operation_class->cl_data;
+  GeglProperties *o = GEGL_PROPERTIES (operation);
+  const size_t gbl_size[2] = {roi->width, roi->height};
+  const size_t gbl_off[2]  = {roi->x, roi->y};
+  cl_int cl_err = 0;
+
+  cl_mem filter_pat = gegl_clCreateBuffer (gegl_cl_get_context (),
+                                           CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+                                           pattern_width[o->pattern] *
+                                           pattern_height[o->pattern] * sizeof(cl_int),
+                                           (void*)pattern[o->pattern],
+                                           &cl_err);
+  CL_CHECK;
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
+                                    sizeof(cl_mem), &in_buf,
+                                    sizeof(cl_mem), &out_buf,
+                                    sizeof(cl_mem), &filter_pat,
+                                    sizeof(cl_int), &pattern_width[o->pattern],
+                                    sizeof(cl_int), &pattern_height[o->pattern],
+                                    sizeof(cl_int), &o->additive,
+                                    sizeof(cl_int), &o->rotated,
+                                    NULL);
+  CL_CHECK;
+  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+                                        cl_data->kernel[0], 2,
+                                        gbl_off, gbl_size, NULL,
+                                        0, NULL, NULL);
+  CL_CHECK;
+  cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
+  CL_CHECK;
+  cl_err = gegl_clReleaseMemObject (filter_pat);
+  CL_CHECK;
+  return FALSE;
+  error:
+    if (filter_pat)
+      gegl_clReleaseMemObject (filter_pat);
+    return TRUE;
+}
 
 static gboolean
 process (GeglOperation       *operation,
@@ -250,6 +297,8 @@ process (GeglOperation       *operation,
   return TRUE;
 }
 
+#include "opencl/video-degradation.cl.h"
+
 static void
 gegl_op_class_init (GeglOpClass *klass)
 {
@@ -260,9 +309,9 @@ gegl_op_class_init (GeglOpClass *klass)
   filter_class    = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
   operation_class->prepare = prepare;
-  operation_class->opencl_support = FALSE;
 
   filter_class->process    = process;
+  filter_class->cl_process = cl_process;
 
   gegl_operation_class_set_keys (operation_class,
     "name",        "gegl:video-degradation",
@@ -271,6 +320,7 @@ gegl_op_class_init (GeglOpClass *klass)
     "license",     "GPL3+",
     "description", _("This function simulates the degradation of "
                      "being on an old low-dotpitch RGB video monitor."),
+    "cl-source"  , video_degradation_cl_source,
     NULL);
 }