[gegl] opencl support for gegl:opacity and Y u8/Y float color format

[Ãyvind KolÃs <ok src gnome org>]

commit eafe6369adf03973eef786a8bd0d98ec4adf8eb6
Author: Victor Oliveira <victormatheus gmail com>
Date:   Thu Feb 23 18:02:07 2012 -0200

    opencl support for gegl:opacity and Y u8/Y float color format

 gegl/opencl/gegl-cl-color-kernel.h |   14 ++++++-
 gegl/opencl/gegl-cl-color.c        |   28 +++++++++---
 operations/common/opacity.c        |   85 ++++++++++++++++++++++++++++++++++++
 3 files changed, 119 insertions(+), 8 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color-kernel.h b/gegl/opencl/gegl-cl-color-kernel.h
index 4353593..f54d97b 100644
--- a/gegl/opencl/gegl-cl-color-kernel.h
+++ b/gegl/opencl/gegl-cl-color-kernel.h
@@ -257,4 +257,16 @@ static const char* kernel_color_source =
 "  float4 in_v  = in[gid];                                                                \n"
 "  float4 out_v = in_v;                                                                   \n"
 "  out[gid] = convert_uchar3_sat_rte(255.0f * out_v.w * out_v.xyz);                       \n"
-"}                                                                                        \n";
+"}                                                                                        \n"
+
+/* -- Y u8 -- */
+
+"__kernel void yu8_to_yf (__global const uchar * in,                                      \n"
+"                         __global       float * out)                                     \n"
+"{                                                                                        \n"
+"  int gid = get_global_id(0);                                                            \n"
+"  float in_v  = convert_float (in[gid]) / 255.0f;                                        \n"
+"  float out_v;                                                                           \n"
+"  out_v = in_v;                                                                          \n"
+"  out[gid] = out_v;                                                                      \n"
+"}                                                                                        \n";
\ No newline at end of file
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 13054cc..83a7d20 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -6,7 +6,7 @@
 
 static gegl_cl_run_data *kernels_color = NULL;
 
-#define CL_FORMAT_N 6
+#define CL_FORMAT_N 8
 
 static const Babl *format[CL_FORMAT_N];
 
@@ -32,6 +32,8 @@ CL_YCBCRAF_TO_RGBAU8      = 13,
 
 CL_RGBU8_TO_RGBAF         = 14,
 CL_RGBAF_TO_RGBU8         = 15,
+
+CL_YU8_TO_YF              = 16,
 };
 
 void
@@ -58,14 +60,18 @@ gegl_cl_color_compile_kernels(void)
                                "rgbu8_to_rgbaf",          /* 14 */
                                "rgbaf_to_rgbu8",          /* 15 */
 
+                               "yu8_to_yf",               /* 16 */
+
                                NULL};
 
-  format[0] = babl_format ("RGBA u8"),
-  format[1] = babl_format ("RGBA float"),
-  format[2] = babl_format ("RaGaBaA float"),
-  format[3] = babl_format ("R'G'B'A float"),
-  format[4] = babl_format ("Y'CbCrA float"),
-  format[5] = babl_format ("RGB u8"),
+  format[0] = babl_format ("RGBA u8");
+  format[1] = babl_format ("RGBA float");
+  format[2] = babl_format ("RaGaBaA float");
+  format[3] = babl_format ("R'G'B'A float");
+  format[4] = babl_format ("Y'CbCrA float");
+  format[5] = babl_format ("RGB u8");
+  format[6] = babl_format ("Y u8");
+  format[7] = babl_format ("Y float");
 
   kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
 }
@@ -110,6 +116,10 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
     {
       if      (out_format == babl_format ("RGBA float"))       kernel = CL_RGBU8_TO_RGBAF;
     }
+  else if (in_format == babl_format ("Y u8"))
+    {
+      if      (out_format == babl_format ("Y float"))          kernel = CL_YU8_TO_YF;
+    }
 
   return kernel;
 }
@@ -132,6 +142,10 @@ gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes)
         *bytes = sizeof (cl_uchar4);
       else if (buffer_format == babl_format ("RGB u8"))
         *bytes = sizeof (cl_uchar3);
+      else if (buffer_format == babl_format ("Y u8"))
+        *bytes = sizeof (cl_uchar);
+      else if (buffer_format == babl_format ("Y float"))
+        *bytes = sizeof (cl_float);
       else
         *bytes = sizeof (cl_float4);
     }
diff --git a/operations/common/opacity.c b/operations/common/opacity.c
index a63d22f..19f5f2a 100644
--- a/operations/common/opacity.c
+++ b/operations/common/opacity.c
@@ -88,6 +88,88 @@ process (GeglOperation       *op,
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+
+static const char* kernel_source =
+"__kernel void kernel_op_3 (__global const float4     *in,      \n"
+"                           __global const float      *aux,     \n"
+"                           __global       float4     *out,     \n"
+"                           float value)                        \n"
+"{                                                              \n"
+"  int gid = get_global_id(0);                                  \n"
+"  float4 in_v  = in [gid];                                     \n"
+"  float  aux_v = aux[gid];                                     \n"
+"  float4 out_v;                                                \n"
+"  out_v = in_v * aux_v * value;                                \n"
+"  out[gid]  =  out_v;                                          \n"
+"}                                                              \n"
+
+"__kernel void kernel_op_2 (__global const float4     *in,      \n"
+"                           __global       float4     *out,     \n"
+"                           float value)                        \n"
+"{                                                              \n"
+"  int gid = get_global_id(0);                                  \n"
+"  float4 in_v  = in [gid];                                     \n"
+"  float4 out_v;                                                \n"
+"  out_v = in_v * value;                                        \n"
+"  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              aux_tex,
+            cl_mem              out_tex,
+            size_t              global_worksize,
+            const GeglRectangle *roi)
+{
+  gfloat value = GEGL_CHANT_PROPERTIES (op)->value;
+
+  cl_int cl_err = 0;
+
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"kernel_op_3", "kernel_op_2", NULL};
+      cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+    }
+
+  if (!cl_data) return 1;
+
+  if (aux_tex)
+    {
+      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*)&aux_tex);
+      cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem),   (void*)&out_tex);
+      cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&value);
+      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;
+    }
+  else
+    {
+      cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem),   (void*)&in_tex);
+      cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem),   (void*)&out_tex);
+      cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_float), (void*)&value);
+      if (cl_err != CL_SUCCESS) return cl_err;
+
+      cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                            cl_data->kernel[1], 1,
+                                            NULL, &global_worksize, NULL,
+                                            0, NULL, NULL);
+      if (cl_err != CL_SUCCESS) return cl_err;
+    }
+
+  return cl_err;
+}
+
 /* Fast path when opacity is a no-op
  */
 static gboolean operation_process (GeglOperation        *operation,
@@ -128,8 +210,11 @@ gegl_chant_class_init (GeglChantClass *klass)
   operation_class->prepare = prepare;
   operation_class->process = operation_process;
   point_composer_class->process = process;
+  point_composer_class->cl_process = cl_process;
 
   operation_class->name        = "gegl:opacity";
+  operation_class->opencl_support = TRUE;
+
   operation_class->categories  = "transparency";
   operation_class->description =
         _("Weights the opacity of the input both the value of the aux"