Added OpenCL support to color-to-alpha

3 files changed, 183 insertions, 0 deletions

diff --git a/opencl/color-to-alpha.cl b/opencl/color-to-alpha.cl
new file mode 100644
index 00000000..9ca08949
--- /dev/null
+++ b/opencl/color-to-alpha.cl
@@ -0,0 +1,63 @@
+__kernel void cl_color_to_alpha(__global const float4 *in,
+                                __global       float4 *out,
+                                float4                color)
+{
+  int gid = get_global_id(0);
+  float4 in_v = in[gid];
+  float4 out_v = in_v;
+  float4 alpha;
+
+  alpha.w = in_v.w;
+
+  /*First component*/
+  if ( color.x < 0.00001f )
+    alpha.x = in_v.x;
+  else if ( in_v.x > color.x + 0.00001f )
+    alpha.x = (in_v.x - color.x) / (1.0f - color.x);
+  else if ( in_v.x < color.x - 0.00001f )
+    alpha.x = (color.x - in_v.x) / color.x;
+  else
+    alpha.x = 0.0f;
+  /*Second component*/
+  if ( color.y < 0.00001f )
+    alpha.y = in_v.y;
+  else if ( in_v.y > color.y + 0.00001f )
+    alpha.y = (in_v.y - color.y) / (1.0f - color.y);
+  else if ( in_v.y < color.y - 0.00001f )
+    alpha.y = (color.y - in_v.y) / color.y;
+  else
+    alpha.y = 0.0f;
+  /*Third component*/
+  if ( color.z < 0.00001f )
+    alpha.z = in_v.z;
+  else if ( in_v.z > color.z + 0.00001f )
+    alpha.z = (in_v.z - color.z) / (1.0f - color.z);
+  else if ( in_v.z < color.z - 0.00001f )
+    alpha.z = (color.z - in_v.z) / color.z;
+  else
+    alpha.z = 0.0f;
+
+  if (alpha.x > alpha.y)
+    {
+      if (alpha.x > alpha.z)
+        out_v.w = alpha.x;
+      else
+        out_v.w = alpha.z;
+    }
+  else if (alpha.y > alpha.z)
+    {
+      out_v.w = alpha.y;
+    }
+  else
+    {
+      out_v.w = alpha.z;
+    }
+
+  if (out_v.w >= 0.00001f)
+    {
+      out_v.xyz = (out_v.xyz - color.xyz) / out_v.www + color.xyz;
+      out_v.w *= alpha.w;
+    }
+
+    out[gid] = out_v;
+}
diff --git a/opencl/color-to-alpha.cl.h b/opencl/color-to-alpha.cl.h
new file mode 100644
index 00000000..409e0bc6
--- /dev/null
+++ b/opencl/color-to-alpha.cl.h
@@ -0,0 +1,65 @@
+static const char* color_to_alpha_cl_source =
+"__kernel void cl_color_to_alpha(__global const float4 *in,                    \n"
+"                                __global       float4 *out,                   \n"
+"                                float4                color)                  \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  float4 in_v = in[gid];                                                      \n"
+"  float4 out_v = in_v;                                                        \n"
+"  float4 alpha;                                                               \n"
+"                                                                              \n"
+"  alpha.w = in_v.w;                                                           \n"
+"                                                                              \n"
+"  /*First component*/                                                         \n"
+"  if ( color.x < 0.00001f )                                                   \n"
+"    alpha.x = in_v.x;                                                         \n"
+"  else if ( in_v.x > color.x + 0.00001f )                                     \n"
+"    alpha.x = (in_v.x - color.x) / (1.0f - color.x);                          \n"
+"  else if ( in_v.x < color.x - 0.00001f )                                     \n"
+"    alpha.x = (color.x - in_v.x) / color.x;                                   \n"
+"  else                                                                        \n"
+"    alpha.x = 0.0f;                                                           \n"
+"  /*Second component*/                                                        \n"
+"  if ( color.y < 0.00001f )                                                   \n"
+"    alpha.y = in_v.y;                                                         \n"
+"  else if ( in_v.y > color.y + 0.00001f )                                     \n"
+"    alpha.y = (in_v.y - color.y) / (1.0f - color.y);                          \n"
+"  else if ( in_v.y < color.y - 0.00001f )                                     \n"
+"    alpha.y = (color.y - in_v.y) / color.y;                                   \n"
+"  else                                                                        \n"
+"    alpha.y = 0.0f;                                                           \n"
+"  /*Third component*/                                                         \n"
+"  if ( color.z < 0.00001f )                                                   \n"
+"    alpha.z = in_v.z;                                                         \n"
+"  else if ( in_v.z > color.z + 0.00001f )                                     \n"
+"    alpha.z = (in_v.z - color.z) / (1.0f - color.z);                          \n"
+"  else if ( in_v.z < color.z - 0.00001f )                                     \n"
+"    alpha.z = (color.z - in_v.z) / color.z;                                   \n"
+"  else                                                                        \n"
+"    alpha.z = 0.0f;                                                           \n"
+"                                                                              \n"
+"  if (alpha.x > alpha.y)                                                      \n"
+"    {                                                                         \n"
+"      if (alpha.x > alpha.z)                                                  \n"
+"        out_v.w = alpha.x;                                                    \n"
+"      else                                                                    \n"
+"        out_v.w = alpha.z;                                                    \n"
+"    }                                                                         \n"
+"  else if (alpha.y > alpha.z)                                                 \n"
+"    {                                                                         \n"
+"      out_v.w = alpha.y;                                                      \n"
+"    }                                                                         \n"
+"  else                                                                        \n"
+"    {                                                                         \n"
+"      out_v.w = alpha.z;                                                      \n"
+"    }                                                                         \n"
+"                                                                              \n"
+"  if (out_v.w >= 0.00001f)                                                    \n"
+"    {                                                                         \n"
+"      out_v.xyz = (out_v.xyz - color.xyz) / out_v.www + color.xyz;            \n"
+"      out_v.w *= alpha.w;                                                     \n"
+"    }                                                                         \n"
+"                                                                              \n"
+"    out[gid] = out_v;                                                         \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/color-to-alpha.c b/operations/common/color-to-alpha.c
index 2d61c242..a8a5904a 100644
--- a/operations/common/color-to-alpha.c
+++ b/operations/common/color-to-alpha.c
@@ -130,6 +130,58 @@ color_to_alpha (const gfloat *color,
   dst[3] *= alpha[3];
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/color-to-alpha.cl.h"
+
+static GeglClRunData * cl_data = NULL;
+
+static gboolean
+cl_process (GeglOperation       *operation,
+            cl_mem              in,
+            cl_mem              out,
+            size_t              global_worksize,
+            const GeglRectangle *roi,
+            gint                level)
+{
+  GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+  gfloat      color[4];
+  gegl_color_get_pixel (o->color, babl_format ("R'G'B'A float"), color);
+
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"cl_color_to_alpha",NULL};
+      cl_data = gegl_cl_compile_and_build (color_to_alpha_cl_source, kernel_name);
+    }
+  if (!cl_data)
+    return TRUE;
+  else
+  {
+    cl_int cl_err = 0;
+    cl_float4 f_color;
+    f_color.s[0] = color[0];
+    f_color.s[1] = color[1];
+    f_color.s[2] = color[2];
+    f_color.s[3] = color[3];
+
+    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0,  sizeof(cl_mem),   (void*)&in);
+    CL_CHECK;
+    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1,  sizeof(cl_mem),   (void*)&out);
+    CL_CHECK;
+    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2,  sizeof(cl_float4),(void*)&f_color);
+    CL_CHECK;
+
+    cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                         cl_data->kernel[0], 1,
+                                         NULL, &global_worksize, NULL,
+                                         0, NULL, NULL);
+    CL_CHECK;
+  }
+
+  return  FALSE;
+
+ error:
+  return TRUE;
+}
 
 
 static gboolean
@@ -192,7 +244,10 @@ gegl_chant_class_init (GeglChantClass *klass)
   filter_class    = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
   filter_class->process    = process;
+  filter_class->cl_process = cl_process;
+
   operation_class->prepare = prepare;
+  operation_class->opencl_support = TRUE;
 
   gegl_operation_class_set_keys (operation_class,
     "name",        "gegl:color-to-alpha",