[gegl/opencl_colors: 2/4] Use a hash table to store OpenCL color conversions



commit 7495bce7ea18501424a1d9ff31093bab82660068
Author: Daniel Sabo <DanielSabo gmail com>
Date:   Fri Nov 1 00:52:15 2013 -0700

    Use a hash table to store OpenCL color conversions

 gegl/opencl/gegl-cl-color.c |  284 +++++++++++++++++--------------------------
 1 files changed, 114 insertions(+), 170 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 3aa6c65..f28b52e 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -31,108 +31,120 @@
 
 #include "opencl/colors.cl.h"
 
-static GeglClRunData *kernels_color = NULL;
-
 #define CL_FORMAT_N 11
 
 static const Babl *format[CL_FORMAT_N];
 
-enum
+static GHashTable  *color_kernels_hash = NULL;
+
+typedef struct
 {
-  CL_RGBAU8_TO_RGBAF          = 0,
-  CL_RGBAF_TO_RGBAU8          = 1,
+  const Babl *from_fmt;
+  const Babl *to_fmt;
+  const char *kernel_name;
+  cl_kernel   kernel;
+} ColorConversionInfo;
+
+static guint
+color_kernels_hash_hashfunc (gconstpointer key)
+{
+  const ColorConversionInfo *ekey = key;
+  return GPOINTER_TO_INT (ekey->from_fmt) ^ GPOINTER_TO_INT (ekey->to_fmt);
+}
 
-  CL_RGBAF_TO_RAGABAF         = 2,
-  CL_RAGABAF_TO_RGBAF         = 3,
-  CL_RGBAU8_TO_RAGABAF        = 4,
-  CL_RAGABAF_TO_RGBAU8        = 5,
+static gboolean
+color_kernels_hash_equalfunc (gconstpointer a,
+                              gconstpointer b)
+{
+  const ColorConversionInfo *ea = a;
+  const ColorConversionInfo *eb = b;
 
-  CL_RGBAF_TO_RGBA_GAMMA_F    = 6,
-  CL_RGBA_GAMMA_F_TO_RGBAF    = 7,
-  CL_RGBAU8_TO_RGBA_GAMMA_F   = 8,
-  CL_RGBA_GAMMA_F_TO_RGBAU8   = 9,
+  if (ea->from_fmt == eb->from_fmt &&
+      ea->to_fmt   == eb->to_fmt)
+    {
+      return TRUE;
+    }
+  return FALSE;
+}
 
-  CL_RGBAF_TO_YCBCRAF         = 10,
-  CL_YCBCRAF_TO_RGBAF         = 11,
-  CL_RGBAU8_TO_YCBCRAF        = 12,
-  CL_YCBCRAF_TO_RGBAU8        = 13,
+void
+gegl_cl_color_compile_kernels (void)
+{
+  static GeglClRunData *float_kernels = NULL;
 
-  CL_RGBU8_TO_RGBAF           = 14,
-  CL_RGBAF_TO_RGBU8           = 15,
+  ColorConversionInfo float_conversions[] = {
+    { babl_format ("RGBA u8"), babl_format ("RGBA float"), "rgbau8_to_rgbaf", NULL },
+    { babl_format ("RGBA float"), babl_format ("RGBA u8"), "rgbaf_to_rgbau8", NULL },
 
-  CL_YU8_TO_YF                = 16,
+    { babl_format ("RGBA float"), babl_format("RaGaBaA float"), "rgbaf_to_ragabaf", NULL },
+    { babl_format ("RaGaBaA float"), babl_format("RGBA float"), "ragabaf_to_rgbaf", NULL },
+    { babl_format ("RGBA u8"), babl_format("RaGaBaA float"), "rgbau8_to_ragabaf", NULL },
+    { babl_format ("RaGaBaA float"), babl_format("RGBA u8"), "ragabaf_to_rgbau8", NULL },
 
-  CL_RGBAF_TO_YAF             = 17,
-  CL_YAF_TO_RGBAF             = 18,
-  CL_RGBAU8_TO_YAF            = 19,
-  CL_YAF_TO_RGBAU8            = 20,
+    { babl_format ("RGBA float"), babl_format("R'G'B'A float"), "rgbaf_to_rgba_gamma_f", NULL },
+    { babl_format ("R'G'B'A float"), babl_format("RGBA float"), "rgba_gamma_f_to_rgbaf", NULL },
+    { babl_format ("RGBA u8"), babl_format("R'G'B'A float"), "rgbau8_to_rgba_gamma_f", NULL },
+    { babl_format ("R'G'B'A float"), babl_format("RGBA u8"), "rgba_gamma_f_to_rgbau8", NULL },
 
-  CL_RGBAF_TO_RGBA_GAMMA_U8   = 21,
-  CL_RGBA_GAMMA_U8_TO_RGBAF   = 22,
+    { babl_format ("RGBA float"), babl_format("Y'CbCrA float"), "rgbaf_to_ycbcraf", NULL },
+    { babl_format ("Y'CbCrA float"), babl_format("RGBA float"), "ycbcraf_to_rgbaf", NULL },
+    { babl_format ("RGBA u8"), babl_format("Y'CbCrA float"), "rgbau8_to_ycbcraf", NULL },
+    { babl_format ("Y'CbCrA float"), babl_format("RGBA u8"), "ycbcraf_to_rgbau8", NULL },
 
-  CL_RGBAF_TO_RGB_GAMMA_U8    = 23,
-  CL_RGB_GAMMA_U8_TO_RGBAF    = 24,
+    { babl_format ("RGB u8"), babl_format("RGBA float"), "rgbu8_to_rgbaf", NULL },
+    { babl_format ("RGBA float"), babl_format("RGB u8"), "rgbaf_to_rgbu8", NULL },
 
-  CL_RGBA_GAMMA_U8_TO_RAGABAF = 25,
-  CL_RAGABAF_TO_RGBA_GAMMA_U8 = 26,
-  CL_RGB_GAMMA_U8_TO_RAGABAF  = 27,
-  CL_RAGABAF_TO_RGB_GAMMA_U8  = 28,
+    { babl_format ("Y u8"), babl_format("Y float"), "yu8_to_yf", NULL },
 
-  CL_RGBA_GAMMA_U8_TO_YAF     = 29,
-  CL_YAF_TO_RGBA_GAMMA_U8     = 30,
-  CL_RGB_GAMMA_U8_TO_YAF      = 31,
-  CL_YAF_TO_RGB_GAMMA_U8      = 32,
-};
+    { babl_format ("RGBA float"), babl_format("YA float"), "rgbaf_to_yaf", NULL },
+    { babl_format ("YA float"), babl_format("RGBA float"), "yaf_to_rgbaf", NULL },
+    { babl_format ("RGBA u8"), babl_format("YA float"), "rgbau8_to_yaf", NULL },
+    { babl_format ("YA float"), babl_format("RGBA u8"), "yaf_to_rgbau8", NULL },
 
-void
-gegl_cl_color_compile_kernels(void)
-{
-  const char *kernel_name[] = {"rgbau8_to_rgbaf",         /* 0  */
-                               "rgbaf_to_rgbau8",         /* 1  */
+    { babl_format ("RGBA float"), babl_format("R'G'B'A u8"), "rgbaf_to_rgba_gamma_u8", NULL },
+    { babl_format ("R'G'B'A u8"), babl_format("RGBA float"), "rgba_gamma_u8_to_rgbaf", NULL },
 
-                               "rgbaf_to_ragabaf",        /* 2  */
-                               "ragabaf_to_rgbaf",        /* 3  */
-                               "rgbau8_to_ragabaf",       /* 4  */
-                               "ragabaf_to_rgbau8",       /* 5  */
+    { babl_format ("RGBA float"), babl_format("R'G'B' u8"), "rgbaf_to_rgb_gamma_u8", NULL },
+    { babl_format ("R'G'B' u8"), babl_format("RGBA float"), "rgb_gamma_u8_to_rgbaf", NULL },
 
-                               "rgbaf_to_rgba_gamma_f",   /* 6  */
-                               "rgba_gamma_f_to_rgbaf",   /* 7  */
-                               "rgbau8_to_rgba_gamma_f",  /* 8  */
-                               "rgba_gamma_f_to_rgbau8",  /* 9  */
+    { babl_format ("R'G'B'A u8"), babl_format("RaGaBaA float"), "rgba_gamma_u8_to_ragabaf", NULL },
+    { babl_format ("RaGaBaA float"), babl_format("R'G'B'A u8"), "ragabaf_to_rgba_gamma_u8", NULL },
+    { babl_format ("R'G'B' u8"), babl_format("RaGaBaA float"), "rgb_gamma_u8_to_ragabaf", NULL },
+    { babl_format ("RaGaBaA float"), babl_format("R'G'B' u8"), "ragabaf_to_rgb_gamma_u8", NULL },
 
-                               "rgbaf_to_ycbcraf",        /* 10 */
-                               "ycbcraf_to_rgbaf",        /* 11 */
-                               "rgbau8_to_ycbcraf",       /* 12 */
-                               "ycbcraf_to_rgbau8",       /* 13 */
+    { babl_format ("R'G'B'A u8"), babl_format("YA float"), "rgba_gamma_u8_to_yaf", NULL },
+    { babl_format ("YA float"), babl_format("R'G'B'A u8"), "yaf_to_rgba_gamma_u8", NULL },
+    { babl_format ("R'G'B' u8"), babl_format("YA float"), "rgb_gamma_u8_to_yaf", NULL },
+    { babl_format ("YA float"), babl_format("R'G'B' u8"), "yaf_to_rgb_gamma_u8", NULL }
+  };
 
-                               "rgbu8_to_rgbaf",          /* 14 */
-                               "rgbaf_to_rgbu8",          /* 15 */
+  const char *kernel_name[G_N_ELEMENTS (float_conversions) + 1];
 
-                               "yu8_to_yf",               /* 16 */
+  for (int i = 0; i < G_N_ELEMENTS (float_conversions); ++i)
+    {
+      kernel_name[i] = float_conversions[i].kernel_name;
+    }
 
-                               "rgbaf_to_yaf",            /* 17 */
-                               "yaf_to_rgbaf",            /* 18 */
-                               "rgbau8_to_yaf",           /* 19 */
-                               "yaf_to_rgbau8",           /* 20 */
+  kernel_name[G_N_ELEMENTS (float_conversions)] = NULL;
 
-                               "rgbaf_to_rgba_gamma_u8",  /* 21  */
-                               "rgba_gamma_u8_to_rgbaf",  /* 22  */
+  float_kernels = gegl_cl_compile_and_build (colors_cl_source, kernel_name);
 
-                               "rgbaf_to_rgb_gamma_u8",   /* 23  */
-                               "rgb_gamma_u8_to_rgbaf",   /* 24  */
+  color_kernels_hash = g_hash_table_new_full (color_kernels_hash_hashfunc,
+                                              color_kernels_hash_equalfunc,
+                                              NULL,
+                                              NULL);
 
-                               "rgba_gamma_u8_to_ragabaf", /* 25 */
-                               "ragabaf_to_rgba_gamma_u8", /* 26 */
-                               "rgb_gamma_u8_to_ragabaf",  /* 27 */
-                               "ragabaf_to_rgb_gamma_u8",  /* 28 */
+  for (int i = 0; i < G_N_ELEMENTS (float_conversions); ++i)
+    {
+      ColorConversionInfo *info = g_new (ColorConversionInfo, 1);
 
-                               "rgba_gamma_u8_to_yaf",     /* 29 */
-                               "yaf_to_rgba_gamma_u8",     /* 30 */
-                               "rgb_gamma_u8_to_yaf",      /* 31 */
-                               "yaf_to_rgb_gamma_u8",      /* 32 */
+      float_conversions[i].kernel = float_kernels->kernel[i];
+      *info = float_conversions[i];
 
-                               NULL};
+      g_hash_table_insert (color_kernels_hash, info, info);
+    }
 
+  /* FIXME: This concept of supported formats it not useful */
   format[0] = babl_format ("RGBA u8");
   format[1] = babl_format ("RGBA float");
   format[2] = babl_format ("RaGaBaA float");
@@ -144,83 +156,18 @@ gegl_cl_color_compile_kernels(void)
   format[8] = babl_format ("YA float");
   format[9] = babl_format ("R'G'B'A u8");
   format[10] = babl_format ("R'G'B' u8");
-
-  kernels_color = gegl_cl_compile_and_build (colors_cl_source, kernel_name);
 }
 
-
-
-static gint
-choose_kernel (const Babl *in_format,
-               const Babl *out_format)
+static cl_kernel
+find_color_kernel (const Babl *in_format,
+                   const Babl *out_format)
 {
-  gint kernel = -1;
-
-  if      (in_format == babl_format ("RGBA float"))
-    {
-      if      (out_format == babl_format ("RGBA u8"))          kernel = CL_RGBAF_TO_RGBAU8;
-      else if (out_format == babl_format ("RaGaBaA float"))    kernel = CL_RGBAF_TO_RAGABAF;
-      else if (out_format == babl_format ("R'G'B'A u8"))       kernel = CL_RGBAF_TO_RGBA_GAMMA_U8;
-      else if (out_format == babl_format ("R'G'B' u8"))        kernel = CL_RGBAF_TO_RGB_GAMMA_U8;
-      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 (out_format == babl_format ("YA float"))         kernel = CL_RGBAF_TO_YAF;
-    }
-  else if (in_format == babl_format ("RGBA u8"))
-    {
-      if      (out_format == babl_format ("RGBA float"))       kernel = CL_RGBAU8_TO_RGBAF;
-      else if (out_format == babl_format ("RaGaBaA float"))    kernel = CL_RGBAU8_TO_RAGABAF;
-      else if (out_format == babl_format ("R'G'B'A float"))    kernel = CL_RGBAU8_TO_RGBA_GAMMA_F;
-      else if (out_format == babl_format ("Y'CbCrA float"))    kernel = CL_RGBAU8_TO_YCBCRAF;
-      else if (out_format == babl_format ("YA float"))         kernel = CL_RGBAU8_TO_YAF;
-    }
-  else if (in_format == babl_format ("RaGaBaA float"))
-    {
-      if      (out_format == babl_format ("RGBA float"))       kernel = CL_RAGABAF_TO_RGBAF;
-      else if (out_format == babl_format ("RGBA u8"))          kernel = CL_RAGABAF_TO_RGBAU8;
-      else if (out_format == babl_format ("R'G'B'A u8"))       kernel = CL_RAGABAF_TO_RGBA_GAMMA_U8;
-      else if (out_format == babl_format ("R'G'B' u8"))        kernel = CL_RAGABAF_TO_RGB_GAMMA_U8;
-    }
-  else if (in_format == babl_format ("R'G'B'A float"))
-    {
-      if      (out_format == babl_format ("RGBA float"))       kernel = CL_RGBA_GAMMA_F_TO_RGBAF;
-      else if (out_format == babl_format ("RGBA u8"))          kernel = CL_RGBA_GAMMA_F_TO_RGBAU8;
-    }
-  else if (in_format == babl_format ("Y'CbCrA float"))
-    {
-      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;
-    }
-  else if (in_format == babl_format ("Y u8"))
-    {
-      if      (out_format == babl_format ("Y float"))          kernel = CL_YU8_TO_YF;
-    }
-  else if (in_format == babl_format ("YA float"))
-    {
-      if      (out_format == babl_format ("RGBA float"))       kernel = CL_YAF_TO_RGBAF;
-      else if (out_format == babl_format ("RGBA u8"))          kernel = CL_YAF_TO_RGBAU8;
-      else if (out_format == babl_format ("R'G'B'A u8"))       kernel = CL_YAF_TO_RGBA_GAMMA_U8;
-      else if (out_format == babl_format ("R'G'B' u8"))        kernel = CL_YAF_TO_RGB_GAMMA_U8;
-    }
-  else if (in_format == babl_format ("R'G'B'A u8"))
-    {
-      if      (out_format == babl_format ("RGBA float"))       kernel = CL_RGBA_GAMMA_U8_TO_RGBAF;
-      else if (out_format == babl_format ("RaGaBaA float"))    kernel = CL_RGBA_GAMMA_U8_TO_RAGABAF;
-      else if (out_format == babl_format ("YA float"))         kernel = CL_RGBA_GAMMA_U8_TO_YAF;
-    }
-  else if (in_format == babl_format ("R'G'B' u8"))
-    {
-      if      (out_format == babl_format ("RGBA float"))       kernel = CL_RGB_GAMMA_U8_TO_RGBAF;
-      else if (out_format == babl_format ("RaGaBaA float"))    kernel = CL_RGB_GAMMA_U8_TO_RAGABAF;
-      else if (out_format == babl_format ("YA float"))         kernel = CL_RGB_GAMMA_U8_TO_YAF;
-    }
-
-  return kernel;
+  ColorConversionInfo *info;
+  ColorConversionInfo search = { in_format, out_format, NULL, NULL };
+  info = g_hash_table_lookup (color_kernels_hash, &search);
+  if (info)
+    return info->kernel;
+  return NULL;
 }
 
 gboolean
@@ -269,7 +216,7 @@ gegl_cl_color_supported (const Babl *in_format,
   if (in_format == out_format)
     return GEGL_CL_COLOR_EQUAL;
 
-  if (kernels_color && choose_kernel (in_format, out_format) >= 0)
+  if (color_kernels_hash && find_color_kernel (in_format, out_format))
     return GEGL_CL_COLOR_CONVERT;
 
   return GEGL_CL_COLOR_NOT_SUPPORTED;
@@ -282,39 +229,36 @@ gegl_cl_color_conv (cl_mem         in_tex,
                     const Babl    *in_format,
                     const Babl    *out_format)
 {
-  cl_int cl_err;
-
-  if (gegl_cl_color_supported (in_format, out_format) == GEGL_CL_COLOR_NOT_SUPPORTED)
-    return FALSE;
+  cl_kernel kernel = NULL;
+  cl_int cl_err = 0;
 
   if (in_format == out_format)
     {
-      size_t s;
-      gegl_cl_color_babl (in_format, &s);
+      size_t s = babl_format_get_bytes_per_pixel (in_format);
 
       /* just copy in_tex to out_tex */
       cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
                                          in_tex, out_tex, 0, 0, size * s,
                                          0, NULL, NULL);
       CL_CHECK;
+      return FALSE;
     }
-  else
-    {
-      gint k = choose_kernel (in_format, out_format);
 
-      cl_err = gegl_clSetKernelArg(kernels_color->kernel[k], 0, sizeof(cl_mem), (void*)&in_tex);
-      CL_CHECK;
-
-      cl_err = gegl_clSetKernelArg(kernels_color->kernel[k], 1, sizeof(cl_mem), (void*)&out_tex);
-      CL_CHECK;
-
-      cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
-                                           kernels_color->kernel[k], 1,
-                                           NULL, &size, NULL,
-                                           0, NULL, NULL);
-      CL_CHECK;
-    }
+  kernel = find_color_kernel (in_format, out_format);
+  if (!kernel)
+    return FALSE;
 
+  cl_err = gegl_cl_set_kernel_args (kernel,
+                                    sizeof(cl_mem), &in_tex,
+                                    sizeof(cl_mem), &out_tex,
+                                    NULL);
+  CL_CHECK;
+
+  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+                                        kernel, 1,
+                                        NULL, &size, NULL,
+                                        0, NULL, NULL);
+  CL_CHECK;
   return FALSE;
 
 error:


[Date Prev][Date Next]   [Thread Prev][Thread Next]   [Thread Index] [Date Index] [Author Index]