[gegl] Simple OpenCL API for point filter

[Victor Matheus de Araujo Oliveira <vmaolive src gnome org>]

commit 9687b4fe0c28cb7bc5a460316465c6546f18c42c
Author: Victor Oliveira <victormatheus gmail com>
Date:   Sat Apr 14 13:01:23 2012 -0300

    Simple OpenCL API for point filter
    
    It can't get easier that that :)

 gegl/operation/gegl-operation-point-filter.c |   35 +++++++-
 gegl/operation/gegl-operation.c              |  107 ++++++++++++++++++++++++++
 gegl/operation/gegl-operation.h              |   11 +++-
 operations/common/brightness-contrast.c      |   54 +------------
 4 files changed, 152 insertions(+), 55 deletions(-)
---
diff --git a/gegl/operation/gegl-operation-point-filter.c b/gegl/operation/gegl-operation-point-filter.c
index 211563e..cc7f48d 100644
--- a/gegl/operation/gegl-operation-point-filter.c
+++ b/gegl/operation/gegl-operation-point-filter.c
@@ -85,6 +85,7 @@ gegl_operation_point_filter_cl_process (GeglOperation       *operation,
   const Babl *in_format  = gegl_operation_get_format (operation, "input");
   const Babl *out_format = gegl_operation_get_format (operation, "output");
 
+  GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
   GeglOperationPointFilterClass *point_filter_class = GEGL_OPERATION_POINT_FILTER_GET_CLASS (operation);
 
   gint j;
@@ -116,12 +117,35 @@ gegl_operation_point_filter_cl_process (GeglOperation       *operation,
         if (err) return FALSE;
         for (j=0; j < i->n; j++)
           {
-            cl_err = point_filter_class->cl_process(operation, i->tex[read][j], i->tex[0][j],
-                                                    i->size[0][j], &i->roi[0][j], level);
+            if (point_filter_class->cl_process)
+              {
+                cl_err = point_filter_class->cl_process(operation, i->tex[read][j], i->tex[0][j],
+                                                        i->size[0][j], &i->roi[0][j], level);
+              }
+            else if (operation_class->cl_data)
+              {
+                gint p = 0;
+                gegl_cl_run_data *cl_data = operation_class->cl_data;
+
+                cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read][j]);
+                cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[  0 ][j]);
+
+                gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
+
+                cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                                     cl_data->kernel[0], 1,
+                                                     NULL, &i->size[0][j], NULL,
+                                                     0, NULL, NULL);
+              }
+            else
+              {
+                g_warning ("OpenCL support enabled, but no way to execute");
+                return FALSE;
+              }
+
             if (cl_err != CL_SUCCESS)
               {
-                GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in %s [GeglOperationPointFilter] Kernel",
-                           GEGL_OPERATION_CLASS (operation)->name);
+                GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointFilter Kernel: %sS", gegl_cl_errstring(cl_err));
                 return FALSE;
               }
           }
@@ -141,11 +165,12 @@ gegl_operation_point_filter_process (GeglOperation       *operation,
   const Babl *out_format = gegl_operation_get_format (operation, "output");
   GeglOperationPointFilterClass *point_filter_class;
 
+  GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
   point_filter_class = GEGL_OPERATION_POINT_FILTER_GET_CLASS (operation);
 
   if ((result->width > 0) && (result->height > 0))
     {
-      if (gegl_cl_is_accelerated () && point_filter_class->cl_process)
+      if (gegl_cl_is_accelerated () && operation_class->opencl_support)
         {
           if (gegl_operation_point_filter_cl_process (operation, input, output, result, level))
             return TRUE;
diff --git a/gegl/operation/gegl-operation.c b/gegl/operation/gegl-operation.c
index dc24e61..d8f4873 100644
--- a/gegl/operation/gegl-operation.c
+++ b/gegl/operation/gegl-operation.c
@@ -92,6 +92,7 @@ gegl_operation_class_init (GeglOperationClass *klass)
   klass->get_bounding_box          = get_bounding_box;
   klass->get_invalidated_by_change = get_invalidated_by_change;
   klass->get_required_for_output   = get_required_for_output;
+  klass->cl_data                   = NULL;
 }
 
 static void
@@ -283,6 +284,18 @@ gegl_operation_prepare (GeglOperation *self)
 
   klass = GEGL_OPERATION_GET_CLASS (self);
 
+  /* build OpenCL kernel */
+  {
+    const gchar *cl_source = gegl_operation_class_get_key (klass, "cl-source");
+    const gchar *cl_kernel = gegl_operation_class_get_key (klass, "cl-kernel");
+    if (cl_source && cl_kernel)
+      {
+        const char *kernel_name[] = {cl_kernel, NULL};
+        gegl_cl_run_data *cl_data = gegl_cl_compile_and_build (cl_source, kernel_name);
+        klass->cl_data = cl_data;
+      }
+  }
+
   if (klass->prepare)
     klass->prepare (self);
 }
@@ -494,6 +507,100 @@ gegl_operation_invalidate (GeglOperation       *operation,
   gegl_node_invalidated (node, roi, TRUE);
 }
 
+gboolean
+gegl_operation_cl_set_kernel_args (GeglOperation *operation,
+                                   cl_kernel      kernel,
+                                   gint          *p,
+                                   cl_int        *err)
+{
+  GParamSpec **self;
+  GParamSpec **parent;
+  guint n_self;
+  guint n_parent;
+  gint prop_no;
+
+  self = g_object_class_list_properties (
+            G_OBJECT_CLASS (g_type_class_ref (G_OBJECT_CLASS_TYPE (GEGL_OPERATION_GET_CLASS(operation)))),
+            &n_self);
+
+  parent = g_object_class_list_properties (
+            G_OBJECT_CLASS (g_type_class_ref (GEGL_TYPE_OPERATION)),
+            &n_parent);
+
+  for (prop_no=0;prop_no<n_self;prop_no++)
+    {
+      gint parent_no;
+      gboolean found=FALSE;
+      for (parent_no=0;parent_no<n_parent;parent_no++)
+        if (self[prop_no]==parent[parent_no])
+          found=TRUE;
+      /* only print properties if we are an addition compared to
+       * GeglOperation
+       */
+
+      /* Removing pads */
+      if (!strcmp(g_param_spec_get_name (self[prop_no]), "input") ||
+          !strcmp(g_param_spec_get_name (self[prop_no]), "output") ||
+          !strcmp(g_param_spec_get_name (self[prop_no]), "aux"))
+        continue;
+
+      if (!found)
+        {
+          if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_DOUBLE))
+            {
+              gdouble value;
+              cl_float v;
+
+              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
+
+              v = value;
+              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v);
+            }
+          else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_FLOAT))
+            {
+              gfloat value;
+              cl_float v;
+
+              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
+
+              v = value;
+              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v);
+            }
+          else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_INT))
+            {
+              gint value;
+              cl_int v;
+
+              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
+
+              v = value;
+              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_int), (void*)&v);
+            }
+          else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_BOOLEAN))
+            {
+              gboolean value;
+              cl_bool v;
+
+              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
+
+              v = value;
+              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_bool), (void*)&v);
+            }
+          else
+            {
+              g_error ("Unsupported OpenCL kernel argument");
+              return FALSE;
+            }
+        }
+    }
+
+  if (self)
+    g_free (self);
+  if (parent)
+    g_free (parent);
+
+  return TRUE;
+}
 
 gchar **
 gegl_operation_list_keys (const gchar *operation_name,
diff --git a/gegl/operation/gegl-operation.h b/gegl/operation/gegl-operation.h
index 0679500..609b368 100644
--- a/gegl/operation/gegl-operation.h
+++ b/gegl/operation/gegl-operation.h
@@ -25,6 +25,7 @@
 
 
 #include "gegl-buffer.h"
+#include "opencl/gegl-cl.h"
 
 G_BEGIN_DECLS
 
@@ -139,7 +140,10 @@ struct _GeglOperationClass
   GeglNode*     (*detect)                    (GeglOperation       *operation,
                                               gint                 x,
                                               gint                 y);
-  gpointer      pad[10];
+
+  gegl_cl_run_data *cl_data;
+
+  gpointer      pad[9];
 };
 
 
@@ -246,6 +250,11 @@ void     gegl_operation_invalidate       (GeglOperation       *operation,
                                           const GeglRectangle *roi,
                                           gboolean             clear_cache);
 
+gboolean gegl_operation_cl_set_kernel_args (GeglOperation *operation,
+                                            cl_kernel      kernel,
+                                            gint          *p,
+                                            cl_int        *err);
+
 /* internal utility functions used by gegl, these should not be used
  * externally */
 gboolean gegl_operation_calc_need_rects      (GeglOperation       *operation,
diff --git a/operations/common/brightness-contrast.c b/operations/common/brightness-contrast.c
index 1081e82..0f932eb 100644
--- a/operations/common/brightness-contrast.c
+++ b/operations/common/brightness-contrast.c
@@ -107,11 +107,11 @@ process (GeglOperation       *op,
 
 #include "opencl/gegl-cl.h"
 
-static const char* kernel_source =
+static const gchar* kernel_source =
 "__kernel void kernel_bc(__global const float4     *in,         \n"
 "                        __global       float4     *out,        \n"
-"                        float brightness,                      \n"
-"                        float contrast)                        \n"
+"                        float contrast,                        \n"
+"                        float brightness)                      \n"
 "{                                                              \n"
 "  int gid = get_global_id(0);                                  \n"
 "  float4 in_v  = in[gid];                                      \n"
@@ -121,51 +121,6 @@ static const char* kernel_source =
 "  out[gid]  =  out_v;                                          \n"
 "}                                                              \n";
 
-static gegl_cl_run_data *cl_data = NULL;
-
-/* OpenCL processing function */
-static cl_int
-cl_process (GeglOperation       *op,
-            cl_mem               in_tex,
-            cl_mem               out_tex,
-            size_t               global_worksize,
-            const GeglRectangle *roi,
-            int                  level)
-{
-  /* Retrieve a pointer to GeglChantO structure which contains all the
-   * chanted properties
-   */
-
-  GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
-
-  gfloat brightness = o->brightness;
-  gfloat contrast   = o->contrast;
-
-  cl_int cl_err = 0;
-
-  if (!cl_data)
-    {
-      const char *kernel_name[] = {"kernel_bc", NULL};
-      cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
-    }
-
-  if (!cl_data) return 1;
-
-  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
-  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&out_tex);
-  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&brightness);
-  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&contrast);
-  if (cl_err != CL_SUCCESS) return cl_err;
-
-  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
-                                        cl_data->kernel[0], 1,
-                                        NULL, &global_worksize, NULL,
-                                        0, NULL, NULL);
-  if (cl_err != CL_SUCCESS) return cl_err;
-
-  return cl_err;
-}
-
 /*
  * The class init function sets up information needed for this operations class
  * (template) in the GObject OO framework.
@@ -185,7 +140,6 @@ gegl_chant_class_init (GeglChantClass *klass)
    * of our superclasses deal with the handling on their level of abstraction)
    */
   point_filter_class->process = process;
-  point_filter_class->cl_process = cl_process;
 
   /* specify the name this operation is found under in the GUI/when
    * programming/in XML
@@ -196,6 +150,8 @@ gegl_chant_class_init (GeglChantClass *klass)
       "name",       "gegl:brightness-contrast",
       "categories", "color", 
       "description", _("Changes the light level and contrast."),
+      "cl-source"  , kernel_source,
+      "cl-kernel"  , "kernel_bc",
       NULL);
 }