[gegl/gsoc2011-opencl: 18/19] Brightness-Contrast op with OpenCL support

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

commit fa5d9ad8034ec9bc8b9bcdf1b5f1e7f1bd7a72fd
Author: Victor Oliveira <victormatheus gmail com>
Date:   Thu Aug 18 13:29:21 2011 -0300

    Brightness-Contrast op with OpenCL support

 operations/common/brightness-contrast.c |   71 ++++++++++++++++++++++++++++++-
 1 files changed, 70 insertions(+), 1 deletions(-)
---
diff --git a/operations/common/brightness-contrast.c b/operations/common/brightness-contrast.c
index 81cf8e3..5b00b17 100644
--- a/operations/common/brightness-contrast.c
+++ b/operations/common/brightness-contrast.c
@@ -105,6 +105,74 @@ process (GeglOperation       *op,
   return TRUE;
 }
 
+/* BEGIN OpenCL */
+
+#include "gegl-cl-init.h"
+
+static const char* kernel_source =
+"sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |              \n"
+"                    CLK_ADDRESS_NONE            |              \n"
+"                    CLK_FILTER_NEAREST;                        \n"
+"__kernel void kernel_bc(__read_only  image2d_t in,             \n"
+"                        __write_only image2d_t out,            \n"
+"                         float brightness,                     \n"
+"                         float contrast)                       \n"
+"{                                                              \n"
+"  int2 gid = (int2)(get_global_id(0), get_global_id(1));       \n"
+"  float4 in_v  = read_imagef(in, sampler, gid);                \n"
+"  float4 out_v;                                                \n"
+"  out_v.xyz = (in_v.xyz - 0.5f) * contrast + brightness + 0.5f;\n"
+"  out_v.w   =  in_v.w;                                         \n"
+"  write_imagef(out, gid, out_v);                               \n"
+"}                                                              \n";
+
+static gegl_cl_run_data *cl_data = NULL;
+
+/* OpenCL processing function */
+static gboolean
+cl_process (GeglOperation       *op,
+            GeglClTexture       *in_tex,
+            GeglClTexture       *out_tex,
+            glong                n_pixels,
+            const GeglRectangle *roi)
+{
+  /* 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;
+
+  size_t global_worksize[2] = {roi->width, roi->height};
+
+  cl_int errcode;
+
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"kernel_bc", NULL};
+      cl_data = gegl_cl_compile_and_build (kernel_source,
+                                           gegl_cl_get_context (),
+                                           gegl_cl_get_device (),
+                                           kernel_name);
+    }
+
+  CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex->data));
+  CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&out_tex->data));
+  CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&brightness));
+  CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&contrast));
+
+  CL_SAFE_CALL(errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                                     cl_data->kernel[0], 2,
+                                                     NULL, global_worksize, NULL,
+                                                     0, NULL, NULL) );
+
+  CL_SAFE_CALL(errcode = gegl_clFinish(gegl_cl_get_command_queue ()));
+  return TRUE;
+}
+
+/* END OpenCL */
+
 /*
  * The class init function sets up information needed for this operations class
  * (template) in the GObject OO framework.
@@ -123,7 +191,8 @@ gegl_chant_class_init (GeglChantClass *klass)
   /* override the process method of the point filter class (the process methods
    * of our superclasses deal with the handling on their level of abstraction)
    */
-  point_filter_class->process = process;
+  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