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

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

commit 0859b84924f498697b7dfc66f4c075803cf547a6
Author: Victor Matheus de A. Oliveira <victormatheus gmail com>
Date:   Sun May 15 20:31:15 2011 -0300

    gegl:over uses opencl
    
    Signed-off-by: Victor Oliveira <victormatheus gmail com>

 operations/common/over.c |  134 +++++++++++++++++++++++++++++++++++++++++++--
 1 files changed, 128 insertions(+), 6 deletions(-)
---
diff --git a/operations/common/over.c b/operations/common/over.c
index 18c6515..633daad 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)
 {
@@ -38,10 +39,18 @@ static void prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "input", format);
   gegl_operation_set_format (operation, "aux", format);
   gegl_operation_set_format (operation, "output", format);
-}
 
+  gegl_cl_init(NULL);
+}
 
-#include <gegl-cl-init.h>
+#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
 cl_process (GeglOperation        *op,
@@ -51,13 +60,109 @@ cl_process (GeglOperation        *op,
             glong                n_pixels,
             const GeglRectangle *roi)
 {
-  gegl_cl_init(NULL);
+  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)(_in.x + _aux.x * (1.0f - _in.w),   \n"
+  "                            _in.x + _aux.x * (1.0f - _in.w),   \n"
+  "                            _in.x + _aux.x * (1.0f - _in.w),   \n"
+  "                            _in.w + _aux.w - _in.w * _aux.w);  \n"
+  "    }                                                          \n"
+  "}                                                              \n";
+
+  size_t kernel_size = sizeof(kernel_source);
+
+  cl_int errcode;
+
+  cl_platform_id platform;
+  cl_device_id device;
+
+  cl_context_properties ctx_properties[3];
+
+  cl_context ctx;
+  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( errcode = gegl_clGetPlatformIDs (1, &platform, NULL) );
+  //CL_SAFE_CALL( errcode = gegl_clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL) );
+
+  /*
+  ctx_properties[0] = CL_CONTEXT_PLATFORM;
+  ctx_properties[1] = (cl_context_properties) platform;
+  ctx_properties[2] = 0;
+  */
+
+  CL_SAFE_CALL( ctx = gegl_clCreateContext(0, 1, &device, NULL, NULL, &errcode) );
+  CL_SAFE_CALL( cq  = gegl_clCreateCommandQueue(ctx, device, 0, &errcode) );
+  CL_SAFE_CALL( program = gegl_clCreateProgramWithSource(ctx, 1, (const char **)&kernel_source, &kernel_size, &errcode) );
+  CL_SAFE_CALL( errcode = gegl_clBuildProgram(program, 0, NULL, NULL, NULL, NULL) );
+
+/*
+  CL_SAFE_CALL( d_in  = gegl_clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sizeof(cl_float) * _n_pixels, NULL, &errcode) );
+  CL_SAFE_CALL( d_aux = gegl_clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sizeof(cl_float) * _n_pixels, NULL, &errcode) );
+  CL_SAFE_CALL( d_out = gegl_clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, sizeof(cl_float) * _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, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &local_worksize, NULL) );
+  global_worksize = (_n_pixels+local_worksize-1) / local_worksize;
+
+  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_in,  CL_FALSE, 0, sizeof(cl_float) * _n_pixels, in,  0, NULL, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_aux, CL_FALSE, 0, sizeof(cl_float) * _n_pixels, aux, 0, NULL, NULL) );
+
+  CL_SAFE_CALL( errcode = gegl_clEnqueueNDRangeKernel(cq, kernel, 1, NULL, &global_worksize, &local_worksize, 0, NULL, NULL) );
+
+  CL_SAFE_CALL( errcode = gegl_clEnqueueReadBuffer(cq, d_out, CL_TRUE, 0, sizeof(cl_float) * _n_pixels, out, 0, NULL, NULL) );
+
+  CL_SAFE_CALL( errcode = gegl_clReleaseKernel(kernel) );
+  CL_SAFE_CALL( errcode = gegl_clReleaseProgram(program) );
+  CL_SAFE_CALL( errcode = gegl_clReleaseCommandQueue(cq) );
+  CL_SAFE_CALL( errcode = gegl_clReleaseContext(ctx) );
+  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,
+_process (GeglOperation        *op,
           void                *in_buf,
           void                *aux_buf,
           void                *out_buf,
@@ -84,10 +189,27 @@ process (GeglOperation        *op,
       out += 4;
     }
 
-  cl_process(op, in_buf, aux_buf, out_buf, n_pixels, roi);
   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_is_accelerated())
+    {
+      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,