[gegl/gsoc2011-opencl: 5/14] gegl:over uses opencl

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

commit ecf727ea00bd9358ef6daa27404306ce5003d409
Author: Victor Oliveira <victormatheus gmail com>
Date:   Sat Jul 2 00:47:11 2011 -0300

    gegl:over uses opencl
    
    this is an example/test of opencl functionalities

 operations/common/over.c |  131 +++++++++++++++++++++++++++++++++++++++++++++-
 1 files changed, 129 insertions(+), 2 deletions(-)
---
diff --git a/operations/common/over.c b/operations/common/over.c
index 1519ec2..c7e5e07 100644
--- a/operations/common/over.c
+++ b/operations/common/over.c
@@ -17,8 +17,8 @@
  */
 
 #include "config.h"
-#include <glib/gi18n-lib.h>
 
+#include <glib/gi18n-lib.h>
 
 #ifdef GEGL_CHANT_PROPERTIES
 
@@ -30,6 +30,7 @@
 #define GEGL_CHANT_C_FILE        "over.c"
 
 #include "gegl-chant.h"
+#include "gegl-cl-init.h"
 
 static void prepare (GeglOperation *operation)
 {
@@ -40,8 +41,115 @@ static void prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "output", format);
 }
 
+#define CL_SAFE_CALL(func)                                          \
+func;                                                               \
+if (errcode != CL_SUCCESS)                                          \
+{                                                                   \
+  g_warning("OpenCL error in %s, Line %u in file %s\nError:%s",     \
+            #func, __LINE__, __FILE__, gegl_cl_errstring(errcode)); \
+  return FALSE;                                                     \
+}
+
 static gboolean
-process (GeglOperation        *op,
+cl_process (GeglOperation        *op,
+            void                *in_buf,
+            void                *aux_buf,
+            void                *out_buf,
+            glong                n_pixels,
+            const GeglRectangle *roi)
+{
+  gfloat * GEGL_ALIGNED in = in_buf;
+  gfloat * GEGL_ALIGNED aux = aux_buf;
+  gfloat * GEGL_ALIGNED out = out_buf;
+  cl_uint _n_pixels = (cl_uint)n_pixels;
+
+  const char* kernel_source[] =
+  {
+  "__kernel void kernel_over (__global float4* in,               \n",
+  "                           __global float4* aux,              \n",
+  "                           __global float4* out,              \n",
+  "                           uint n_pixels)                     \n",
+  "{                                                             \n",
+  "    size_t gid = get_global_id(0);                            \n",
+  "    if (gid < n_pixels)                                       \n",
+  "    {                                                         \n",
+  "        float4 _in  = in[gid];                                \n",
+  "        float4 _aux = aux[gid];                               \n",
+  "        out[gid] = (float4)(_aux.x + _in.x * (1.0f - _aux.w), \n",
+  "                            _aux.y + _in.y * (1.0f - _aux.w), \n",
+  "                            _aux.z + _in.z * (1.0f - _aux.w), \n",
+  "                            _aux.w + _in.w - _aux.w * _in.w); \n",
+  "    }                                                         \n",
+  "}                                                             \n",
+  };
+
+  char buffer[16384];
+
+  cl_int errcode;
+  cl_command_queue cq;
+  cl_program program;
+  cl_kernel kernel;
+
+  size_t local_worksize;
+  size_t global_worksize;
+
+  cl_mem d_in;
+  cl_mem d_aux;
+  cl_mem d_out;
+
+  if (aux==NULL)
+    return TRUE;
+
+  /* -- Configuration -- */
+
+  CL_SAFE_CALL( cq      = gegl_clCreateCommandQueue(gegl_cl_get_context(), gegl_cl_get_device(), 0, &errcode) );
+  CL_SAFE_CALL( program = gegl_clCreateProgramWithSource(gegl_cl_get_context(), 16, (const char **)&kernel_source, NULL, &errcode) );
+  errcode = gegl_clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+  if (errcode != CL_SUCCESS)
+    {
+      CL_SAFE_CALL( errcode = gegl_clGetProgramBuildInfo(program, gegl_cl_get_device(), CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, NULL) );
+      g_warning("OpenCL Build Error in Line %u in file %s\nError:%s\n%s",
+                __LINE__, __FILE__, gegl_cl_errstring(errcode), buffer);
+      return FALSE;
+    }
+
+  CL_SAFE_CALL( d_in  = gegl_clCreateBuffer(gegl_cl_get_context(), CL_MEM_READ_ONLY,  sizeof(cl_float4) * _n_pixels, NULL, &errcode) );
+  CL_SAFE_CALL( d_aux = gegl_clCreateBuffer(gegl_cl_get_context(), CL_MEM_READ_ONLY,  sizeof(cl_float4) * _n_pixels, NULL, &errcode) );
+  CL_SAFE_CALL( d_out = gegl_clCreateBuffer(gegl_cl_get_context(), CL_MEM_WRITE_ONLY, sizeof(cl_float4) * _n_pixels, NULL, &errcode) );
+
+  CL_SAFE_CALL( kernel  = gegl_clCreateKernel(program, "kernel_over", &errcode) );
+  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&d_in) );
+  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&d_aux) );
+  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&d_out) );
+  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 3, sizeof(cl_uint), (void*)&_n_pixels) );
+
+  /* -- Running -- */
+
+  CL_SAFE_CALL( errcode = gegl_clGetKernelWorkGroupInfo(kernel, gegl_cl_get_device(), CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &local_worksize, NULL) );
+  global_worksize = MAX( ((_n_pixels+local_worksize-1) / local_worksize) * local_worksize, local_worksize );
+
+  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_in,  CL_FALSE, 0, sizeof(cl_float4) * _n_pixels, in,  0, NULL, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_aux, CL_FALSE, 0, sizeof(cl_float4) * _n_pixels, aux, 0, NULL, NULL) );
+
+  CL_SAFE_CALL( errcode = gegl_clFinish(cq) );
+  CL_SAFE_CALL( errcode = gegl_clEnqueueNDRangeKernel(cq, kernel, 1, NULL, &global_worksize, &local_worksize, 0, NULL, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clFinish(cq) );
+
+  CL_SAFE_CALL( errcode = gegl_clEnqueueReadBuffer(cq, d_out, CL_TRUE, 0, sizeof(cl_float4) * _n_pixels, out, 0, NULL, NULL) );
+
+  CL_SAFE_CALL( errcode = gegl_clReleaseProgram(program) );
+  CL_SAFE_CALL( errcode = gegl_clReleaseCommandQueue(cq) );
+
+  CL_SAFE_CALL( errcode = gegl_clReleaseKernel(kernel) );
+  CL_SAFE_CALL( errcode = gegl_clReleaseMemObject(d_in) );
+  CL_SAFE_CALL( errcode = gegl_clReleaseMemObject(d_aux) );
+  CL_SAFE_CALL( errcode = gegl_clReleaseMemObject(d_out));
+
+  return TRUE;
+}
+
+static gboolean
+_process (GeglOperation        *op,
           void                *in_buf,
           void                *aux_buf,
           void                *out_buf,
@@ -67,9 +175,28 @@ process (GeglOperation        *op,
       aux += 4;
       out += 4;
     }
+
   return TRUE;
 }
 
+static gboolean
+process (GeglOperation        *op,
+          void                *in_buf,
+          void                *aux_buf,
+          void                *out_buf,
+          glong                n_pixels,
+          const GeglRectangle *roi)
+{
+  if (gegl_cl_init(NULL))
+  {
+    return cl_process(op, in_buf, aux_buf, out_buf, n_pixels, roi);
+  }
+  else
+  {
+    return _process(op, in_buf, aux_buf, out_buf, n_pixels, roi);
+  }
+}
+
 /* Fast paths */
 static gboolean operation_process (GeglOperation        *operation,
                                    GeglOperationContext *context,