[gegl] core: manage cl_process functions on POINT_COMPOSER operations

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

commit 4de532063f6715316ddfbe76164b7e0783c55c91
Author: Thomas Manni <thomas manni free fr>
Date:   Thu Jun 29 11:16:42 2017 +0200

    core: manage cl_process functions on POINT_COMPOSER operations

 gegl/operation/gegl-operation-point-composer.c |  102 ++++++++++++++++++++++++
 1 files changed, 102 insertions(+), 0 deletions(-)
---
diff --git a/gegl/operation/gegl-operation-point-composer.c b/gegl/operation/gegl-operation-point-composer.c
index e6a86f9..e68f28c 100644
--- a/gegl/operation/gegl-operation-point-composer.c
+++ b/gegl/operation/gegl-operation-point-composer.c
@@ -22,6 +22,7 @@
 #include <glib-object.h>
 
 #include "gegl.h"
+#include "gegl-debug.h"
 #include "gegl-operation-point-composer.h"
 #include "gegl-operation-context.h"
 #include "gegl-config.h"
@@ -30,6 +31,9 @@
 #include <unistd.h>
 #include <string.h>
 
+#include "opencl/gegl-cl.h"
+#include "gegl-buffer-cl-iterator.h"
+
 typedef struct ThreadData
 {
   GeglOperationPointComposerClass *klass;
@@ -171,6 +175,97 @@ static gboolean gegl_operation_point_composer_process
 
 G_DEFINE_TYPE (GeglOperationPointComposer, gegl_operation_point_composer, GEGL_TYPE_OPERATION_COMPOSER)
 
+static gboolean
+gegl_operation_point_composer_cl_process (GeglOperation       *operation,
+                                          GeglBuffer          *input,
+                                          GeglBuffer          *aux,
+                                          GeglBuffer          *output,
+                                          const GeglRectangle *result,
+                                          gint                 level)
+{
+  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);
+  GeglOperationPointComposerClass *point_composer_class = GEGL_OPERATION_POINT_COMPOSER_GET_CLASS 
(operation);
+
+  GeglBufferClIterator *iter = NULL;
+
+  cl_int cl_err = 0;
+  gboolean err;
+
+  GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_COMPOSER: %s", operation_class->name);
+
+  /* Process */
+  iter = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
+
+  gegl_buffer_cl_iterator_add (iter, input, result, in_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
+
+  if (aux)
+    {
+      const Babl *aux_format = gegl_operation_get_format (operation, "aux");
+      gegl_buffer_cl_iterator_add (iter, aux, result, aux_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
+    }
+
+  while (gegl_buffer_cl_iterator_next (iter, &err))
+    {
+      if (err)
+        return FALSE;
+
+      if (point_composer_class->cl_process)
+        {
+          err = point_composer_class->cl_process (operation,
+                                                  iter->tex[1],
+                                                  (aux) ? iter->tex[2] : NULL,
+                                                  iter->tex[0],
+                                                  iter->size[0],
+                                                  &iter->roi[0],
+                                                  level);
+          if (err)
+            {
+              GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name);
+              gegl_buffer_cl_iterator_stop (iter);
+              return FALSE;
+            }
+        }
+      else if (operation_class->cl_data)
+        {
+          gint p = 0;
+          GeglClRunData *cl_data = operation_class->cl_data;
+
+          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[1]);
+          CL_CHECK;
+          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (aux) ? 
(void*)&iter->tex[2] : NULL);
+          CL_CHECK;
+          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[0]);
+          CL_CHECK;
+
+          gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
+          CL_CHECK;
+
+          cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+                                                cl_data->kernel[0], 1,
+                                                NULL, &iter->size[0], NULL,
+                                                0, NULL, NULL);
+          CL_CHECK;
+        }
+      else
+        {
+          g_warning ("OpenCL support enabled, but no way to execute");
+          gegl_buffer_cl_iterator_stop (iter);
+          return FALSE;
+        }
+    }
+
+  return TRUE;
+
+error:
+  GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring (cl_err));
+  if (iter)
+    gegl_buffer_cl_iterator_stop (iter);
+  return FALSE;
+}
+
 static void prepare (GeglOperation *operation)
 {
   const Babl *format = gegl_babl_rgba_linear_float ();
@@ -207,6 +302,7 @@ gegl_operation_point_composer_process (GeglOperation       *operation,
                                        const GeglRectangle *result,
                                        gint                 level)
 {
+  GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
   GeglOperationPointComposerClass *point_composer_class = GEGL_OPERATION_POINT_COMPOSER_GET_CLASS 
(operation);
   const Babl *in_format   = gegl_operation_get_format (operation, "input");
   const Babl *aux_format  = gegl_operation_get_format (operation, "aux");
@@ -219,6 +315,12 @@ gegl_operation_point_composer_process (GeglOperation       *operation,
       const Babl *aux_buf_format = aux?gegl_buffer_get_format(aux):NULL;
       const Babl *output_buf_format = output?gegl_buffer_get_format(output):NULL;
 
+      if (gegl_operation_use_opencl (operation) && (operation_class->cl_data || 
point_composer_class->cl_process))
+        {
+          if (gegl_operation_point_composer_cl_process (operation, input, aux, output, result, level))
+              return TRUE;
+        }
+
       if (gegl_operation_use_threading (operation, result) && result->height > 1)
       {
         gint threads = gegl_config_threads ();