[gegl] Support for RGBu8

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

commit f51dcd2c8c5b181dae4c7f1bccba874818dc3fb5
Author: Victor Oliveira <victormatheus gmail com>
Date:   Wed Feb 15 13:21:26 2012 -0200

    Support for RGBu8
    
    And minor changes in color conversion

 gegl/buffer/gegl-buffer-cl-cache.c           |    6 ++--
 gegl/buffer/gegl-buffer-cl-iterator.c        |    4 +-
 gegl/opencl/gegl-cl-color-kernel.h           |   24 +++++++++++++++++
 gegl/opencl/gegl-cl-color.c                  |   35 ++++++++++++++-----------
 gegl/opencl/gegl-cl-color.h                  |    2 +-
 gegl/operation/gegl-operation-point-filter.c |    4 +-
 6 files changed, 52 insertions(+), 23 deletions(-)
---
diff --git a/gegl/buffer/gegl-buffer-cl-cache.c b/gegl/buffer/gegl-buffer-cl-cache.c
index aa5a34d..5568d4f 100644
--- a/gegl/buffer/gegl-buffer-cl-cache.c
+++ b/gegl/buffer/gegl-buffer-cl-cache.c
@@ -88,7 +88,7 @@ gegl_buffer_cl_cache_merge (GeglBuffer          *buffer,
   if (!roi)
     roi = &buffer->extent;
 
-  gegl_cl_color_babl (buffer->format, NULL, &size);
+  gegl_cl_color_babl (buffer->format, &size);
 
   if (G_UNLIKELY (!cache_entries))
     {
@@ -213,8 +213,8 @@ gegl_buffer_cl_cache_from (GeglBuffer          *buffer,
   gint i;
 
   gegl_cl_color_op conv = gegl_cl_color_supported (buffer->format, format);
-  gegl_cl_color_babl (buffer->format, NULL, &buf_size);
-  gegl_cl_color_babl (format,         NULL, &dest_size);
+  gegl_cl_color_babl (buffer->format, &buf_size);
+  gegl_cl_color_babl (format,         &dest_size);
 
   if (G_UNLIKELY (!cache_entries))
     {
diff --git a/gegl/buffer/gegl-buffer-cl-iterator.c b/gegl/buffer/gegl-buffer-cl-iterator.c
index 86162f8..037972b 100644
--- a/gegl/buffer/gegl-buffer-cl-iterator.c
+++ b/gegl/buffer/gegl-buffer-cl-iterator.c
@@ -94,8 +94,8 @@ gegl_buffer_cl_iterator_add (GeglBufferClIterator  *iterator,
   else
     i->conv[self] = gegl_cl_color_supported (buffer->format, format);
 
-  gegl_cl_color_babl (buffer->format, NULL, &i->buf_cl_format_size[self]);
-  gegl_cl_color_babl (format,         NULL, &i->op_cl_format_size [self]);
+  gegl_cl_color_babl (buffer->format, &i->buf_cl_format_size[self]);
+  gegl_cl_color_babl (format,         &i->op_cl_format_size [self]);
 
   if (self!=0)
     {
diff --git a/gegl/opencl/gegl-cl-color-kernel.h b/gegl/opencl/gegl-cl-color-kernel.h
index 0b5b1f7..5bfa4b9 100644
--- a/gegl/opencl/gegl-cl-color-kernel.h
+++ b/gegl/opencl/gegl-cl-color-kernel.h
@@ -234,4 +234,28 @@ static const char* kernel_color_source =
 "                   linear_to_gamma_2_2(rgb.z),                                           \n"
 "                   in_v.w);                                                              \n"
 "  out[gid] = convert_uchar4_sat_rte(255.0f * out_v);                                     \n"
+"}                                                                                        \n"
+
+/* -- RGB u8 -- */
+
+/* RGB u8 -> RGBA float */
+"__kernel void rgbu8_to_rgbaf (__global const uchar3 * in,                                \n"
+"                              __global       float4 * out)                               \n"
+"{                                                                                        \n"
+"  int gid = get_global_id(0);                                                            \n"
+"  float3 in_v  = convert_float3(in[gid]) / 255.0f;                                       \n"
+"  float4 out_v;                                                                          \n"
+"  out_v.xyz = in_v;                                                                      \n"
+"  out_v.w   = 1.0f;                                                                      \n"
+"  out[gid] = out_v;                                                                      \n"
+"}                                                                                        \n"
+
+/* RGBA float -> RGB u8 */
+"__kernel void rgbaf_to_rgbu8 (__global const float4 * in,                                \n"
+"                              __global       uchar3 * out)                               \n"
+"{                                                                                        \n"
+"  int gid = get_global_id(0);                                                            \n"
+"  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";
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 8b4aba4..13054cc 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 5
+#define CL_FORMAT_N 6
 
 static const Babl *format[CL_FORMAT_N];
 
@@ -29,6 +29,9 @@ CL_RGBAF_TO_YCBCRAF       = 10,
 CL_YCBCRAF_TO_RGBAF       = 11,
 CL_RGBAU8_TO_YCBCRAF      = 12,
 CL_YCBCRAF_TO_RGBAU8      = 13,
+
+CL_RGBU8_TO_RGBAF         = 14,
+CL_RGBAF_TO_RGBU8         = 15,
 };
 
 void
@@ -52,6 +55,9 @@ gegl_cl_color_compile_kernels(void)
                                "rgbau8_to_ycbcraf",       /* 12 */
                                "ycbcraf_to_rgbau8",       /* 13 */
 
+                               "rgbu8_to_rgbaf",          /* 14 */
+                               "rgbaf_to_rgbu8",          /* 15 */
+
                                NULL};
 
   format[0] = babl_format ("RGBA u8"),
@@ -59,6 +65,7 @@ gegl_cl_color_compile_kernels(void)
   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"),
 
   kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
 }
@@ -75,6 +82,7 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
       else if (out_format == babl_format ("RaGaBaA float"))    kernel = CL_RGBAF_TO_RAGABAF;
       else if (out_format == babl_format ("R'G'B'A float"))    kernel = CL_RGBAF_TO_RGBA_GAMMA_F;
       else if (out_format == babl_format ("Y'CbCrA float"))    kernel = CL_RGBAF_TO_YCBCRAF;
+      else if (out_format == babl_format ("RGB u8"))           kernel = CL_RGBAF_TO_RGBU8;
     }
   else if (in_format == babl_format ("RGBA u8"))
     {
@@ -98,12 +106,16 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
       if      (out_format == babl_format ("RGBA float"))       kernel = CL_YCBCRAF_TO_RGBAF;
       else if (out_format == babl_format ("RGBA u8"))          kernel = CL_YCBCRAF_TO_RGBAU8;
     }
+  else if (in_format == babl_format ("RGB u8"))
+    {
+      if      (out_format == babl_format ("RGBA float"))       kernel = CL_RGBU8_TO_RGBAF;
+    }
 
   return kernel;
 }
 
 gboolean
-gegl_cl_color_babl (const Babl *buffer_format, cl_image_format *cl_format, size_t *bytes)
+gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes)
 {
   int i;
   gboolean supported_format = FALSE;
@@ -114,23 +126,16 @@ gegl_cl_color_babl (const Babl *buffer_format, cl_image_format *cl_format, size_
   if (!supported_format)
     return FALSE;
 
-  if (cl_format)
+  if (bytes)
     {
       if (buffer_format == babl_format ("RGBA u8"))
-        {
-          cl_format->image_channel_order     = CL_RGBA;
-          cl_format->image_channel_data_type = CL_UNORM_INT8;
-        }
+        *bytes = sizeof (cl_uchar4);
+      else if (buffer_format == babl_format ("RGB u8"))
+        *bytes = sizeof (cl_uchar3);
       else
-        {
-          cl_format->image_channel_order      = CL_RGBA;
-          cl_format->image_channel_data_type  = CL_FLOAT;
-        }
+        *bytes = sizeof (cl_float4);
     }
 
-  if (bytes)
-    *bytes  = (buffer_format == babl_format ("RGBA u8"))? sizeof (cl_uchar4) : sizeof (cl_float4);
-
   return TRUE;
 }
 
@@ -160,7 +165,7 @@ gegl_cl_color_conv (cl_mem in_tex, cl_mem out_tex, const size_t size,
   if (in_format == out_format)
     {
       size_t s;
-      gegl_cl_color_babl (in_format, NULL, &s);
+      gegl_cl_color_babl (in_format, &s);
 
       /* just copy in_tex to out_tex */
       errcode = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
diff --git a/gegl/opencl/gegl-cl-color.h b/gegl/opencl/gegl-cl-color.h
index 797ed6c..d9ca87f 100644
--- a/gegl/opencl/gegl-cl-color.h
+++ b/gegl/opencl/gegl-cl-color.h
@@ -13,7 +13,7 @@ typedef enum
 
 void gegl_cl_color_compile_kernels(void);
 
-gboolean gegl_cl_color_babl (const Babl *buffer_format, cl_image_format *cl_format, size_t *bytes);
+gboolean gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes);
 
 gegl_cl_color_op gegl_cl_color_supported (const Babl *in_format, const Babl *out_format);
 
diff --git a/gegl/operation/gegl-operation-point-filter.c b/gegl/operation/gegl-operation-point-filter.c
index c3eea28..fb19c55 100644
--- a/gegl/operation/gegl-operation-point-filter.c
+++ b/gegl/operation/gegl-operation-point-filter.c
@@ -88,8 +88,8 @@ gegl_operation_point_filter_cl_process (GeglOperation       *operation,
   gboolean err;
 
   /* non-texturizable format! */
-  if (!gegl_cl_color_babl (in_format,  NULL, NULL) ||
-      !gegl_cl_color_babl (out_format, NULL, NULL))
+  if (!gegl_cl_color_babl (in_format,  NULL) ||
+      !gegl_cl_color_babl (out_format, NULL))
     {
       g_warning ("[OpenCL] Non-texturizable input of output format!");
       return FALSE;