opencl: Replace u8 -> float gamma conversions with a small LUT

No performance testing done yet, but fixes some accuracy issues.

6 files changed, 837 insertions, 280 deletions

diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index f28b52eb..804dc532 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -30,6 +30,7 @@
 #include "gegl-cl-color.h"
 
 #include "opencl/colors.cl.h"
+#include "opencl/colors-8bit-lut.cl.h"
 
 #define CL_FORMAT_N 11
 
@@ -67,10 +68,47 @@ color_kernels_hash_equalfunc (gconstpointer a,
   return FALSE;
 }
 
-void
+static gboolean
+gegl_cl_color_load_conversion_set (ColorConversionInfo  *conversions,
+                                   gint                  num_conversions,
+                                   GeglClRunData       **kernels,
+                                   const gchar          *source)
+{
+  const char *kernel_names[num_conversions + 1];
+
+  for (int i = 0; i < num_conversions; ++i)
+    {
+      kernel_names[i] = conversions[i].kernel_name;
+    }
+
+  kernel_names[num_conversions] = NULL;
+
+  *kernels = gegl_cl_compile_and_build (source, kernel_names);
+
+  if (!(*kernels))
+    {
+      return FALSE;
+    }
+
+  for (int i = 0; i < num_conversions; ++i)
+    {
+      ColorConversionInfo *info = g_new (ColorConversionInfo, 1);
+
+      conversions[i].kernel = (*kernels)->kernel[i];
+      *info = conversions[i];
+
+      g_hash_table_insert (color_kernels_hash, info, info);
+    }
+
+  return TRUE;
+}
+
+gboolean
 gegl_cl_color_compile_kernels (void)
 {
   static GeglClRunData *float_kernels = NULL;
+  static GeglClRunData *lut8_kernels = NULL;
+  gboolean result = TRUE;
 
   ColorConversionInfo float_conversions[] = {
     { babl_format ("RGBA u8"), babl_format ("RGBA float"), "rgbau8_to_rgbaf", NULL },
@@ -102,49 +140,56 @@ gegl_cl_color_compile_kernels (void)
     { babl_format ("YA float"), babl_format("RGBA u8"), "yaf_to_rgbau8", NULL },
 
     { 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 },
-
     { 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 },
 
-    { 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 },
 
-    { 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 }
   };
 
-  const char *kernel_name[G_N_ELEMENTS (float_conversions) + 1];
-
-  for (int i = 0; i < G_N_ELEMENTS (float_conversions); ++i)
-    {
-      kernel_name[i] = float_conversions[i].kernel_name;
-    }
-
-  kernel_name[G_N_ELEMENTS (float_conversions)] = NULL;
+  ColorConversionInfo lut8_conversions[] = {
+    { babl_format ("R'G'B'A u8"), babl_format("RGBA float"), "rgba_gamma_u8_to_rgbaf", NULL },
+    { babl_format ("R'G'B'A u8"), babl_format("RaGaBaA float"), "rgba_gamma_u8_to_ragabaf", NULL },
+    { babl_format ("R'G'B'A u8"), babl_format("YA float"), "rgba_gamma_u8_to_yaf", NULL },
+    { babl_format ("R'G'B' u8"), babl_format("RGBA float"), "rgb_gamma_u8_to_rgbaf", NULL },
+    { babl_format ("R'G'B' u8"), babl_format("RaGaBaA float"), "rgb_gamma_u8_to_ragabaf", NULL },
+    { babl_format ("R'G'B' u8"), babl_format("YA float"), "rgb_gamma_u8_to_yaf", NULL },
+  };
 
-  float_kernels = gegl_cl_compile_and_build (colors_cl_source, kernel_name);
+  /* Fail if the kernels are already compiled, meaning this must have been called twice */
+  g_return_val_if_fail (!float_kernels, FALSE);
+  g_return_val_if_fail (!color_kernels_hash, FALSE);
 
   color_kernels_hash = g_hash_table_new_full (color_kernels_hash_hashfunc,
                                               color_kernels_hash_equalfunc,
                                               NULL,
                                               NULL);
 
-  for (int i = 0; i < G_N_ELEMENTS (float_conversions); ++i)
+  gegl_cl_color_load_conversion_set (float_conversions,
+                                     G_N_ELEMENTS (float_conversions),
+                                     &float_kernels,
+                                     colors_cl_source);
+
+  if (!float_kernels)
     {
-      ColorConversionInfo *info = g_new (ColorConversionInfo, 1);
+      g_warning ("Failed to compile coloer conversions (float_kernels)");
+      result = FALSE;
+    }
 
-      float_conversions[i].kernel = float_kernels->kernel[i];
-      *info = float_conversions[i];
+  gegl_cl_color_load_conversion_set (lut8_conversions,
+                                     G_N_ELEMENTS (lut8_conversions),
+                                     &lut8_kernels,
+                                     colors_8bit_lut_cl_source);
 
-      g_hash_table_insert (color_kernels_hash, info, info);
+  if (!lut8_kernels)
+    {
+      g_warning ("Failed to compile coloer conversions (lut8_kernels)");
+      result = FALSE;
     }
 
-  /* FIXME: This concept of supported formats it not useful */
+  /* FIXME: Remove this and just use the babl bbp/components */
   format[0] = babl_format ("RGBA u8");
   format[1] = babl_format ("RGBA float");
   format[2] = babl_format ("RaGaBaA float");
@@ -156,6 +201,8 @@ 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");
+
+  return result;
 }
 
 static cl_kernel
diff --git a/gegl/opencl/gegl-cl-color.h b/gegl/opencl/gegl-cl-color.h
index fd0a7316..35d14d56 100644
--- a/gegl/opencl/gegl-cl-color.h
+++ b/gegl/opencl/gegl-cl-color.h
@@ -30,7 +30,7 @@ typedef enum
 } GeglClColorOp;
 
 /** Compile and register OpenCL kernel for color conversion */
-void          gegl_cl_color_compile_kernels(void);
+gboolean      gegl_cl_color_compile_kernels (void);
 
 /** Return TRUE if the Babl format is supported with OpenCL.
  *  If present, returns the byte per pixel in *bytes
diff --git a/opencl/colors-8bit-lut.cl b/opencl/colors-8bit-lut.cl
new file mode 100644
index 00000000..93c00a7b
--- /dev/null
+++ b/opencl/colors-8bit-lut.cl
@@ -0,0 +1,382 @@
+/* This file is part of GEGL
+ *
+ * GEGL is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * GEGL is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.
+ *
+ * Copyright 2012 Victor Oliveira (victormatheus@gmail.com)
+ *           2013 Daniel Sabo
+ */
+
+__constant float u8_gamma_to_linear_lut[] = {
+0.0000000000000000f,
+0.0003035269835488f,
+0.0006070539670977f,
+0.0009105809506465f,
+0.0012141079341954f,
+0.0015176349177442f,
+0.0018211619012930f,
+0.0021246888848419f,
+0.0024282158683907f,
+0.0027317428519395f,
+0.0030352698354884f,
+0.0033465357638992f,
+0.0036765073240474f,
+0.0040247170184963f,
+0.0043914420374103f,
+0.0047769534806937f,
+0.0051815167023384f,
+0.0056053916242027f,
+0.0060488330228571f,
+0.0065120907925945f,
+0.0069954101872654f,
+0.0074990320432262f,
+0.0080231929853850f,
+0.0085681256180693f,
+0.0091340587022208f,
+0.0097212173202378f,
+0.0103298230296269f,
+0.0109600940064882f,
+0.0116122451797439f,
+0.0122864883569159f,
+0.0129830323421730f,
+0.0137020830472897f,
+0.0144438435960925f,
+0.0152085144229127f,
+0.0159962933655096f,
+0.0168073757528874f,
+0.0176419544883841f,
+0.0185002201283797f,
+0.0193823609569357f,
+0.0202885630566524f,
+0.0212190103760036f,
+0.0221738847933874f,
+0.0231533661781104f,
+0.0241576324485048f,
+0.0251868596273616f,
+0.0262412218948499f,
+0.0273208916390749f,
+0.0284260395044208f,
+0.0295568344378088f,
+0.0307134437329936f,
+0.0318960330730115f,
+0.0331047665708851f,
+0.0343398068086822f,
+0.0356013148750203f,
+0.0368894504011000f,
+0.0382043715953465f,
+0.0395462352767328f,
+0.0409151969068532f,
+0.0423114106208097f,
+0.0437350292569735f,
+0.0451862043856755f,
+0.0466650863368801f,
+0.0481718242268894f,
+0.0497065659841272f,
+0.0512694583740432f,
+0.0528606470231802f,
+0.0544802764424424f,
+0.0561284900496001f,
+0.0578054301910672f,
+0.0595112381629812f,
+0.0612460542316176f,
+0.0630100176531677f,
+0.0648032666929058f,
+0.0666259386437729f,
+0.0684781698444002f,
+0.0703600956965959f,
+0.0722718506823175f,
+0.0742135683801496f,
+0.0761853814813079f,
+0.0781874218051863f,
+0.0802198203144683f,
+0.0822827071298148f,
+0.0843762115441488f,
+0.0865004620365498f,
+0.0886555862857729f,
+0.0908417111834077f,
+0.0930589628466875f,
+0.0953074666309647f,
+0.0975873471418625f,
+0.0998987282471139f,
+0.1022417330881013f,
+0.1046164840911042f,
+0.1070231029782676f,
+0.1094617107782993f,
+0.1119324278369056f,
+0.1144353738269737f,
+0.1169706677585108f,
+0.1195384279883456f,
+0.1221387722296019f,
+0.1247718175609505f,
+0.1274376804356474f,
+0.1301364766903643f,
+0.1328683215538180f,
+0.1356333296552057f,
+0.1384316150324518f,
+0.1412632911402716f,
+0.1441284708580578f,
+0.1470272664975950f,
+0.1499597898106086f,
+0.1529261519961502f,
+0.1559264637078274f,
+0.1589608350608804f,
+0.1620293756391110f,
+0.1651321945016676f,
+0.1682694001896907f,
+0.1714411007328226f,
+0.1746474036555850f,
+0.1778884159836291f,
+0.1811642442498602f,
+0.1844749945004410f,
+0.1878207723006779f,
+0.1912016827407914f,
+0.1946178304415758f,
+0.1980693195599489f,
+0.2015562537943971f,
+0.2050787363903169f,
+0.2086368701452558f,
+0.2122307574140552f,
+0.2158605001138993f,
+0.2195261997292692f,
+0.2232279573168085f,
+0.2269658735100984f,
+0.2307400485243492f,
+0.2345505821610052f,
+0.2383975738122710f,
+0.2422811224655549f,
+0.2462013267078355f,
+0.2501582847299534f,
+0.2541520943308267f,
+0.2581828529215958f,
+0.2622506575296962f,
+0.2663556048028625f,
+0.2704977910130658f,
+0.2746773120603846f,
+0.2788942634768104f,
+0.2831487404299921f,
+0.2874408377269175f,
+0.2917706498175359f,
+0.2961382707983211f,
+0.3005437944157765f,
+0.3049873140698863f,
+0.3094689228175085f,
+0.3139887133757175f,
+0.3185467781250919f,
+0.3231432091129507f,
+0.3277780980565422f,
+0.3324515363461794f,
+0.3371636150483304f,
+0.3419144249086609f,
+0.3467040563550296f,
+0.3515325995004394f,
+0.3564001441459435f,
+0.3613067797835095f,
+0.3662525955988395f,
+0.3712376804741491f,
+0.3762621229909065f,
+0.3813260114325301f,
+0.3864294337870490f,
+0.3915724777497233f,
+0.3967552307256269f,
+0.4019777798321958f,
+0.4072402119017367f,
+0.4125426134839038f,
+0.4178850708481375f,
+0.4232676699860717f,
+0.4286904966139066f,
+0.4341536361747489f,
+0.4396571738409188f,
+0.4452011945162279f,
+0.4507857828382235f,
+0.4564110231804047f,
+0.4620769996544071f,
+0.4677837961121590f,
+0.4735314961480095f,
+0.4793201831008268f,
+0.4851499400560704f,
+0.4910208498478356f,
+0.4969329950608704f,
+0.5028864580325687f,
+0.5088813208549338f,
+0.5149176653765214f,
+0.5209955732043543f,
+0.5271151257058131f,
+0.5332764040105052f,
+0.5394794890121072f,
+0.5457244613701866f,
+0.5520114015120001f,
+0.5583403896342679f,
+0.5647115057049292f,
+0.5711248294648731f,
+0.5775804404296506f,
+0.5840784178911641f,
+0.5906188409193369f,
+0.5972017883637634f,
+0.6038273388553378f,
+0.6104955708078648f,
+0.6172065624196511f,
+0.6239603916750761f,
+0.6307571363461468f,
+0.6375968739940326f,
+0.6444796819705821f,
+0.6514056374198242f,
+0.6583748172794485f,
+0.6653872982822721f,
+0.6724431569576875f,
+0.6795424696330938f,
+0.6866853124353135f,
+0.6938717612919899f,
+0.7011018919329731f,
+0.7083757798916868f,
+0.7156935005064807f,
+0.7230551289219693f,
+0.7304607400903537f,
+0.7379104087727308f,
+0.7454042095403874f,
+0.7529422167760779f,
+0.7605245046752924f,
+0.7681511472475070f,
+0.7758222183174236f,
+0.7835377915261935f,
+0.7912979403326302f,
+0.7991027380144090f,
+0.8069522576692516f,
+0.8148465722161012f,
+0.8227857543962835f,
+0.8307698767746546f,
+0.8387990117407400f,
+0.8468732315098580f,
+0.8549926081242338f,
+0.8631572134541023f,
+0.8713671191987972f,
+0.8796223968878317f,
+0.8879231178819663f,
+0.8962693533742664f,
+0.9046611743911496f,
+0.9130986517934192f,
+0.9215818562772946f,
+0.9301108583754237f,
+0.9386857284578880f,
+0.9473065367331999f,
+0.9559733532492861f,
+0.9646862478944651f,
+0.9734452903984125f,
+0.9822505503331171f,
+0.9911020971138298f,
+1.0000000000000000f
+};
+
+/* babl reference file: babl/base/rgb-constants.h */
+#define RGB_LUMINANCE_RED    (0.222491f)
+#define RGB_LUMINANCE_GREEN  (0.716888f)
+#define RGB_LUMINANCE_BLUE   (0.060621f)
+
+/* R'G'B' u8 -> RGBA float */
+__kernel void rgb_gamma_u8_to_rgbaf (__global const uchar  * in,
+                                     __global       float4 * out)
+{
+  int gid = get_global_id(0);
+
+  float r = u8_gamma_to_linear_lut[(int)in[3 * gid + 0]];
+  float g = u8_gamma_to_linear_lut[(int)in[3 * gid + 1]];
+  float b = u8_gamma_to_linear_lut[(int)in[3 * gid + 2]];
+
+  out[gid] = (float4) (r, g, b, 1.0f);
+}
+
+/* R'G'B' u8 -> RaGaBaA float */
+__kernel void rgb_gamma_u8_to_ragabaf (__global const uchar  * in,
+                                       __global       float4 * out)
+{
+  int gid = get_global_id(0);
+
+  float r = u8_gamma_to_linear_lut[(int)in[3 * gid + 0]];
+  float g = u8_gamma_to_linear_lut[(int)in[3 * gid + 1]];
+  float b = u8_gamma_to_linear_lut[(int)in[3 * gid + 2]];
+
+  out[gid] = (float4) (r, g, b, 1.0f);
+}
+
+/* R'G'B' u8 -> YA float */
+__kernel void rgb_gamma_u8_to_yaf (__global const uchar  * in,
+                                   __global       float2 * out)
+{
+  int gid = get_global_id(0);
+
+  float r = u8_gamma_to_linear_lut[(int)in[3 * gid + 0]];
+  float g = u8_gamma_to_linear_lut[(int)in[3 * gid + 1]];
+  float b = u8_gamma_to_linear_lut[(int)in[3 * gid + 2]];
+
+  float luminance = r * RGB_LUMINANCE_RED +
+                    g * RGB_LUMINANCE_GREEN +
+                    b * RGB_LUMINANCE_BLUE;
+
+  out[gid] = (float2) (luminance, 1.0f);
+}
+
+/* R'G'B'A u8 -> RGBA float */
+__kernel void rgba_gamma_u8_to_rgbaf (__global const uchar4 * in,
+                                      __global       float4 * out)
+{
+  int gid = get_global_id(0);
+  uchar4 in_v = in[gid];
+
+  float r = u8_gamma_to_linear_lut[(int)in_v.x];
+  float g = u8_gamma_to_linear_lut[(int)in_v.y];
+  float b = u8_gamma_to_linear_lut[(int)in_v.z];
+
+  out[gid] = (float4) (r, g, b, in_v.w / 255.0f);
+}
+
+/* R'G'B'A u8 -> RaGaBaA float */
+__kernel void rgba_gamma_u8_to_ragabaf (__global const uchar4 * in,
+                                        __global       float4 * out)
+{
+  int gid = get_global_id(0);
+  uchar4 in_v  = in[gid];
+  float4 out_v;
+  float4 tmp_v;
+
+  tmp_v = (float4)(u8_gamma_to_linear_lut[(int)in_v.x],
+                   u8_gamma_to_linear_lut[(int)in_v.x],
+                   u8_gamma_to_linear_lut[(int)in_v.x],
+                   in_v.w / 255.0f);
+
+  out_v   = tmp_v * tmp_v.w;
+  out_v.w = tmp_v.w;
+
+  out[gid] = out_v;
+}
+
+/* R'G'B'A u8 -> YA float */
+__kernel void rgba_gamma_u8_to_yaf (__global const uchar4 * in,
+                                    __global       float2 * out)
+{
+  int gid = get_global_id(0);
+  uchar4 in_v  = in[gid];
+  float4 tmp_v;
+  tmp_v = (float4)(u8_gamma_to_linear_lut[(int)in_v.x],
+                   u8_gamma_to_linear_lut[(int)in_v.x],
+                   u8_gamma_to_linear_lut[(int)in_v.x],
+                   in_v.w / 255.0f);
+  float2 out_v;
+
+  float luminance = tmp_v.x * RGB_LUMINANCE_RED +
+                    tmp_v.y * RGB_LUMINANCE_GREEN +
+                    tmp_v.z * RGB_LUMINANCE_BLUE;
+
+  out_v.x = luminance;
+  out_v.y = tmp_v.w;
+
+  out[gid] = out_v;
+}
\ No newline at end of file
diff --git a/opencl/colors-8bit-lut.cl.h b/opencl/colors-8bit-lut.cl.h
new file mode 100644
index 00000000..0567a78d
--- /dev/null
+++ b/opencl/colors-8bit-lut.cl.h
@@ -0,0 +1,384 @@
+static const char* colors_8bit_lut_cl_source =
+"/* This file is part of GEGL                                                  \n"
+" *                                                                            \n"
+" * GEGL is free software; you can redistribute it and/or                      \n"
+" * modify it under the terms of the GNU Lesser General Public                 \n"
+" * License as published by the Free Software Foundation; either               \n"
+" * version 3 of the License, or (at your option) any later version.           \n"
+" *                                                                            \n"
+" * GEGL is distributed in the hope that it will be useful,                    \n"
+" * but WITHOUT ANY WARRANTY; without even the implied warranty of             \n"
+" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU          \n"
+" * Lesser General Public License for more details.                            \n"
+" *                                                                            \n"
+" * You should have received a copy of the GNU Lesser General Public           \n"
+" * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.       \n"
+" *                                                                            \n"
+" * Copyright 2012 Victor Oliveira (victormatheus@gmail.com)                   \n"
+" *           2013 Daniel Sabo                                                 \n"
+" */                                                                           \n"
+"                                                                              \n"
+"__constant float u8_gamma_to_linear_lut[] = {                                 \n"
+"0.0000000000000000f,                                                          \n"
+"0.0003035269835488f,                                                          \n"
+"0.0006070539670977f,                                                          \n"
+"0.0009105809506465f,                                                          \n"
+"0.0012141079341954f,                                                          \n"
+"0.0015176349177442f,                                                          \n"
+"0.0018211619012930f,                                                          \n"
+"0.0021246888848419f,                                                          \n"
+"0.0024282158683907f,                                                          \n"
+"0.0027317428519395f,                                                          \n"
+"0.0030352698354884f,                                                          \n"
+"0.0033465357638992f,                                                          \n"
+"0.0036765073240474f,                                                          \n"
+"0.0040247170184963f,                                                          \n"
+"0.0043914420374103f,                                                          \n"
+"0.0047769534806937f,                                                          \n"
+"0.0051815167023384f,                                                          \n"
+"0.0056053916242027f,                                                          \n"
+"0.0060488330228571f,                                                          \n"
+"0.0065120907925945f,                                                          \n"
+"0.0069954101872654f,                                                          \n"
+"0.0074990320432262f,                                                          \n"
+"0.0080231929853850f,                                                          \n"
+"0.0085681256180693f,                                                          \n"
+"0.0091340587022208f,                                                          \n"
+"0.0097212173202378f,                                                          \n"
+"0.0103298230296269f,                                                          \n"
+"0.0109600940064882f,                                                          \n"
+"0.0116122451797439f,                                                          \n"
+"0.0122864883569159f,                                                          \n"
+"0.0129830323421730f,                                                          \n"
+"0.0137020830472897f,                                                          \n"
+"0.0144438435960925f,                                                          \n"
+"0.0152085144229127f,                                                          \n"
+"0.0159962933655096f,                                                          \n"
+"0.0168073757528874f,                                                          \n"
+"0.0176419544883841f,                                                          \n"
+"0.0185002201283797f,                                                          \n"
+"0.0193823609569357f,                                                          \n"
+"0.0202885630566524f,                                                          \n"
+"0.0212190103760036f,                                                          \n"
+"0.0221738847933874f,                                                          \n"
+"0.0231533661781104f,                                                          \n"
+"0.0241576324485048f,                                                          \n"
+"0.0251868596273616f,                                                          \n"
+"0.0262412218948499f,                                                          \n"
+"0.0273208916390749f,                                                          \n"
+"0.0284260395044208f,                                                          \n"
+"0.0295568344378088f,                                                          \n"
+"0.0307134437329936f,                                                          \n"
+"0.0318960330730115f,                                                          \n"
+"0.0331047665708851f,                                                          \n"
+"0.0343398068086822f,                                                          \n"
+"0.0356013148750203f,                                                          \n"
+"0.0368894504011000f,                                                          \n"
+"0.0382043715953465f,                                                          \n"
+"0.0395462352767328f,                                                          \n"
+"0.0409151969068532f,                                                          \n"
+"0.0423114106208097f,                                                          \n"
+"0.0437350292569735f,                                                          \n"
+"0.0451862043856755f,                                                          \n"
+"0.0466650863368801f,                                                          \n"
+"0.0481718242268894f,                                                          \n"
+"0.0497065659841272f,                                                          \n"
+"0.0512694583740432f,                                                          \n"
+"0.0528606470231802f,                                                          \n"
+"0.0544802764424424f,                                                          \n"
+"0.0561284900496001f,                                                          \n"
+"0.0578054301910672f,                                                          \n"
+"0.0595112381629812f,                                                          \n"
+"0.0612460542316176f,                                                          \n"
+"0.0630100176531677f,                                                          \n"
+"0.0648032666929058f,                                                          \n"
+"0.0666259386437729f,                                                          \n"
+"0.0684781698444002f,                                                          \n"
+"0.0703600956965959f,                                                          \n"
+"0.0722718506823175f,                                                          \n"
+"0.0742135683801496f,                                                          \n"
+"0.0761853814813079f,                                                          \n"
+"0.0781874218051863f,                                                          \n"
+"0.0802198203144683f,                                                          \n"
+"0.0822827071298148f,                                                          \n"
+"0.0843762115441488f,                                                          \n"
+"0.0865004620365498f,                                                          \n"
+"0.0886555862857729f,                                                          \n"
+"0.0908417111834077f,                                                          \n"
+"0.0930589628466875f,                                                          \n"
+"0.0953074666309647f,                                                          \n"
+"0.0975873471418625f,                                                          \n"
+"0.0998987282471139f,                                                          \n"
+"0.1022417330881013f,                                                          \n"
+"0.1046164840911042f,                                                          \n"
+"0.1070231029782676f,                                                          \n"
+"0.1094617107782993f,                                                          \n"
+"0.1119324278369056f,                                                          \n"
+"0.1144353738269737f,                                                          \n"
+"0.1169706677585108f,                                                          \n"
+"0.1195384279883456f,                                                          \n"
+"0.1221387722296019f,                                                          \n"
+"0.1247718175609505f,                                                          \n"
+"0.1274376804356474f,                                                          \n"
+"0.1301364766903643f,                                                          \n"
+"0.1328683215538180f,                                                          \n"
+"0.1356333296552057f,                                                          \n"
+"0.1384316150324518f,                                                          \n"
+"0.1412632911402716f,                                                          \n"
+"0.1441284708580578f,                                                          \n"
+"0.1470272664975950f,                                                          \n"
+"0.1499597898106086f,                                                          \n"
+"0.1529261519961502f,                                                          \n"
+"0.1559264637078274f,                                                          \n"
+"0.1589608350608804f,                                                          \n"
+"0.1620293756391110f,                                                          \n"
+"0.1651321945016676f,                                                          \n"
+"0.1682694001896907f,                                                          \n"
+"0.1714411007328226f,                                                          \n"
+"0.1746474036555850f,                                                          \n"
+"0.1778884159836291f,                                                          \n"
+"0.1811642442498602f,                                                          \n"
+"0.1844749945004410f,                                                          \n"
+"0.1878207723006779f,                                                          \n"
+"0.1912016827407914f,                                                          \n"
+"0.1946178304415758f,                                                          \n"
+"0.1980693195599489f,                                                          \n"
+"0.2015562537943971f,                                                          \n"
+"0.2050787363903169f,                                                          \n"
+"0.2086368701452558f,                                                          \n"
+"0.2122307574140552f,                                                          \n"
+"0.2158605001138993f,                                                          \n"
+"0.2195261997292692f,                                                          \n"
+"0.2232279573168085f,                                                          \n"
+"0.2269658735100984f,                                                          \n"
+"0.2307400485243492f,                                                          \n"
+"0.2345505821610052f,                                                          \n"
+"0.2383975738122710f,                                                          \n"
+"0.2422811224655549f,                                                          \n"
+"0.2462013267078355f,                                                          \n"
+"0.2501582847299534f,                                                          \n"
+"0.2541520943308267f,                                                          \n"
+"0.2581828529215958f,                                                          \n"
+"0.2622506575296962f,                                                          \n"
+"0.2663556048028625f,                                                          \n"
+"0.2704977910130658f,                                                          \n"
+"0.2746773120603846f,                                                          \n"
+"0.2788942634768104f,                                                          \n"
+"0.2831487404299921f,                                                          \n"
+"0.2874408377269175f,                                                          \n"
+"0.2917706498175359f,                                                          \n"
+"0.2961382707983211f,                                                          \n"
+"0.3005437944157765f,                                                          \n"
+"0.3049873140698863f,                                                          \n"
+"0.3094689228175085f,                                                          \n"
+"0.3139887133757175f,                                                          \n"
+"0.3185467781250919f,                                                          \n"
+"0.3231432091129507f,                                                          \n"
+"0.3277780980565422f,                                                          \n"
+"0.3324515363461794f,                                                          \n"
+"0.3371636150483304f,                                                          \n"
+"0.3419144249086609f,                                                          \n"
+"0.3467040563550296f,                                                          \n"
+"0.3515325995004394f,                                                          \n"
+"0.3564001441459435f,                                                          \n"
+"0.3613067797835095f,                                                          \n"
+"0.3662525955988395f,                                                          \n"
+"0.3712376804741491f,                                                          \n"
+"0.3762621229909065f,                                                          \n"
+"0.3813260114325301f,                                                          \n"
+"0.3864294337870490f,                                                          \n"
+"0.3915724777497233f,                                                          \n"
+"0.3967552307256269f,                                                          \n"
+"0.4019777798321958f,                                                          \n"
+"0.4072402119017367f,                                                          \n"
+"0.4125426134839038f,                                                          \n"
+"0.4178850708481375f,                                                          \n"
+"0.4232676699860717f,                                                          \n"
+"0.4286904966139066f,                                                          \n"
+"0.4341536361747489f,                                                          \n"
+"0.4396571738409188f,                                                          \n"
+"0.4452011945162279f,                                                          \n"
+"0.4507857828382235f,                                                          \n"
+"0.4564110231804047f,                                                          \n"
+"0.4620769996544071f,                                                          \n"
+"0.4677837961121590f,                                                          \n"
+"0.4735314961480095f,                                                          \n"
+"0.4793201831008268f,                                                          \n"
+"0.4851499400560704f,                                                          \n"
+"0.4910208498478356f,                                                          \n"
+"0.4969329950608704f,                                                          \n"
+"0.5028864580325687f,                                                          \n"
+"0.5088813208549338f,                                                          \n"
+"0.5149176653765214f,                                                          \n"
+"0.5209955732043543f,                                                          \n"
+"0.5271151257058131f,                                                          \n"
+"0.5332764040105052f,                                                          \n"
+"0.5394794890121072f,                                                          \n"
+"0.5457244613701866f,                                                          \n"
+"0.5520114015120001f,                                                          \n"
+"0.5583403896342679f,                                                          \n"
+"0.5647115057049292f,                                                          \n"
+"0.5711248294648731f,                                                          \n"
+"0.5775804404296506f,                                                          \n"
+"0.5840784178911641f,                                                          \n"
+"0.5906188409193369f,                                                          \n"
+"0.5972017883637634f,                                                          \n"
+"0.6038273388553378f,                                                          \n"
+"0.6104955708078648f,                                                          \n"
+"0.6172065624196511f,                                                          \n"
+"0.6239603916750761f,                                                          \n"
+"0.6307571363461468f,                                                          \n"
+"0.6375968739940326f,                                                          \n"
+"0.6444796819705821f,                                                          \n"
+"0.6514056374198242f,                                                          \n"
+"0.6583748172794485f,                                                          \n"
+"0.6653872982822721f,                                                          \n"
+"0.6724431569576875f,                                                          \n"
+"0.6795424696330938f,                                                          \n"
+"0.6866853124353135f,                                                          \n"
+"0.6938717612919899f,                                                          \n"
+"0.7011018919329731f,                                                          \n"
+"0.7083757798916868f,                                                          \n"
+"0.7156935005064807f,                                                          \n"
+"0.7230551289219693f,                                                          \n"
+"0.7304607400903537f,                                                          \n"
+"0.7379104087727308f,                                                          \n"
+"0.7454042095403874f,                                                          \n"
+"0.7529422167760779f,                                                          \n"
+"0.7605245046752924f,                                                          \n"
+"0.7681511472475070f,                                                          \n"
+"0.7758222183174236f,                                                          \n"
+"0.7835377915261935f,                                                          \n"
+"0.7912979403326302f,                                                          \n"
+"0.7991027380144090f,                                                          \n"
+"0.8069522576692516f,                                                          \n"
+"0.8148465722161012f,                                                          \n"
+"0.8227857543962835f,                                                          \n"
+"0.8307698767746546f,                                                          \n"
+"0.8387990117407400f,                                                          \n"
+"0.8468732315098580f,                                                          \n"
+"0.8549926081242338f,                                                          \n"
+"0.8631572134541023f,                                                          \n"
+"0.8713671191987972f,                                                          \n"
+"0.8796223968878317f,                                                          \n"
+"0.8879231178819663f,                                                          \n"
+"0.8962693533742664f,                                                          \n"
+"0.9046611743911496f,                                                          \n"
+"0.9130986517934192f,                                                          \n"
+"0.9215818562772946f,                                                          \n"
+"0.9301108583754237f,                                                          \n"
+"0.9386857284578880f,                                                          \n"
+"0.9473065367331999f,                                                          \n"
+"0.9559733532492861f,                                                          \n"
+"0.9646862478944651f,                                                          \n"
+"0.9734452903984125f,                                                          \n"
+"0.9822505503331171f,                                                          \n"
+"0.9911020971138298f,                                                          \n"
+"1.0000000000000000f                                                           \n"
+"};                                                                            \n"
+"                                                                              \n"
+"/* babl reference file: babl/base/rgb-constants.h */                          \n"
+"#define RGB_LUMINANCE_RED    (0.222491f)                                      \n"
+"#define RGB_LUMINANCE_GREEN  (0.716888f)                                      \n"
+"#define RGB_LUMINANCE_BLUE   (0.060621f)                                      \n"
+"                                                                              \n"
+"/* R'G'B' u8 -> RGBA float */                                                 \n"
+"__kernel void rgb_gamma_u8_to_rgbaf (__global const uchar  * in,              \n"
+"                                     __global       float4 * out)             \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"                                                                              \n"
+"  float r = u8_gamma_to_linear_lut[(int)in[3 * gid + 0]];                     \n"
+"  float g = u8_gamma_to_linear_lut[(int)in[3 * gid + 1]];                     \n"
+"  float b = u8_gamma_to_linear_lut[(int)in[3 * gid + 2]];                     \n"
+"                                                                              \n"
+"  out[gid] = (float4) (r, g, b, 1.0f);                                        \n"
+"}                                                                             \n"
+"                                                                              \n"
+"/* R'G'B' u8 -> RaGaBaA float */                                              \n"
+"__kernel void rgb_gamma_u8_to_ragabaf (__global const uchar  * in,            \n"
+"                                       __global       float4 * out)           \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"                                                                              \n"
+"  float r = u8_gamma_to_linear_lut[(int)in[3 * gid + 0]];                     \n"
+"  float g = u8_gamma_to_linear_lut[(int)in[3 * gid + 1]];                     \n"
+"  float b = u8_gamma_to_linear_lut[(int)in[3 * gid + 2]];                     \n"
+"                                                                              \n"
+"  out[gid] = (float4) (r, g, b, 1.0f);                                        \n"
+"}                                                                             \n"
+"                                                                              \n"
+"/* R'G'B' u8 -> YA float */                                                   \n"
+"__kernel void rgb_gamma_u8_to_yaf (__global const uchar  * in,                \n"
+"                                   __global       float2 * out)               \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"                                                                              \n"
+"  float r = u8_gamma_to_linear_lut[(int)in[3 * gid + 0]];                     \n"
+"  float g = u8_gamma_to_linear_lut[(int)in[3 * gid + 1]];                     \n"
+"  float b = u8_gamma_to_linear_lut[(int)in[3 * gid + 2]];                     \n"
+"                                                                              \n"
+"  float luminance = r * RGB_LUMINANCE_RED +                                   \n"
+"                    g * RGB_LUMINANCE_GREEN +                                 \n"
+"                    b * RGB_LUMINANCE_BLUE;                                   \n"
+"                                                                              \n"
+"  out[gid] = (float2) (luminance, 1.0f);                                      \n"
+"}                                                                             \n"
+"                                                                              \n"
+"/* R'G'B'A u8 -> RGBA float */                                                \n"
+"__kernel void rgba_gamma_u8_to_rgbaf (__global const uchar4 * in,             \n"
+"                                      __global       float4 * out)            \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  uchar4 in_v = in[gid];                                                      \n"
+"                                                                              \n"
+"  float r = u8_gamma_to_linear_lut[(int)in_v.x];                              \n"
+"  float g = u8_gamma_to_linear_lut[(int)in_v.y];                              \n"
+"  float b = u8_gamma_to_linear_lut[(int)in_v.z];                              \n"
+"                                                                              \n"
+"  out[gid] = (float4) (r, g, b, in_v.w / 255.0f);                             \n"
+"}                                                                             \n"
+"                                                                              \n"
+"/* R'G'B'A u8 -> RaGaBaA float */                                             \n"
+"__kernel void rgba_gamma_u8_to_ragabaf (__global const uchar4 * in,           \n"
+"                                        __global       float4 * out)          \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  uchar4 in_v  = in[gid];                                                     \n"
+"  float4 out_v;                                                               \n"
+"  float4 tmp_v;                                                               \n"
+"                                                                              \n"
+"  tmp_v = (float4)(u8_gamma_to_linear_lut[(int)in_v.x],                       \n"
+"                   u8_gamma_to_linear_lut[(int)in_v.x],                       \n"
+"                   u8_gamma_to_linear_lut[(int)in_v.x],                       \n"
+"                   in_v.w / 255.0f);                                          \n"
+"                                                                              \n"
+"  out_v   = tmp_v * tmp_v.w;                                                  \n"
+"  out_v.w = tmp_v.w;                                                          \n"
+"                                                                              \n"
+"  out[gid] = out_v;                                                           \n"
+"}                                                                             \n"
+"                                                                              \n"
+"/* R'G'B'A u8 -> YA float */                                                  \n"
+"__kernel void rgba_gamma_u8_to_yaf (__global const uchar4 * in,               \n"
+"                                    __global       float2 * out)              \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  uchar4 in_v  = in[gid];                                                     \n"
+"  float4 tmp_v;                                                               \n"
+"  tmp_v = (float4)(u8_gamma_to_linear_lut[(int)in_v.x],                       \n"
+"                   u8_gamma_to_linear_lut[(int)in_v.x],                       \n"
+"                   u8_gamma_to_linear_lut[(int)in_v.x],                       \n"
+"                   in_v.w / 255.0f);                                          \n"
+"  float2 out_v;                                                               \n"
+"                                                                              \n"
+"  float luminance = tmp_v.x * RGB_LUMINANCE_RED +                             \n"
+"                    tmp_v.y * RGB_LUMINANCE_GREEN +                           \n"
+"                    tmp_v.z * RGB_LUMINANCE_BLUE;                             \n"
+"                                                                              \n"
+"  out_v.x = luminance;                                                        \n"
+"  out_v.y = tmp_v.w;                                                          \n"
+"                                                                              \n"
+"  out[gid] = out_v;                                                           \n"
+"}                                                                             \n"
+;
diff --git a/opencl/colors.cl b/opencl/colors.cl
index c311b22f..3ff0820c 100644
--- a/opencl/colors.cl
+++ b/opencl/colors.cl
@@ -118,24 +118,6 @@ __kernel void ragabaf_to_rgbau8 (__global const float4 * in,
   out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
 }
 
-/* RGBA_GAMMA_U8 -> RaGaBaA float */
-__kernel void rgba_gamma_u8_to_ragabaf (__global const uchar4 * in,
-                                        __global       float4 * out)
-{
-  int gid = get_global_id(0);
-  float4 in_v  = convert_float4(in[gid]) / 255.0f;
-  float4 tmp_v;
-  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),
-                   gamma_2_2_to_linear(in_v.y),
-                   gamma_2_2_to_linear(in_v.z),
-                   in_v.w);
-  float4 out_v;
-  out_v   = tmp_v * tmp_v.w;
-  out_v.w = tmp_v.w;
-  out[gid] = out_v;
-}
-
-
 /* RaGaBaA float -> RGBA_GAMMA_U8 */
 __kernel void ragabaf_to_rgba_gamma_u8 (__global const float4 * in,
                                         __global       uchar4 * out)
@@ -153,29 +135,6 @@ __kernel void ragabaf_to_rgba_gamma_u8 (__global const float4 * in,
   out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
 }
 
-/* RGB_GAMMA_U8 -> RaGaBaA float */
-__kernel void rgb_gamma_u8_to_ragabaf (__global const uchar  * in,
-                                       __global       float4 * out)
-{
-  int gid = get_global_id(0);
-#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
-  float3 in_v  = convert_float3(vload3 (gid, in)) / 255.0f;
-#else
-  uchar4 i4 = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255);
-  float4 in_v = convert_float4 (i4) / 255.0f;
-#endif
-  float4 tmp_v;
-  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),
-                   gamma_2_2_to_linear(in_v.y),
-                   gamma_2_2_to_linear(in_v.z),
-                   1.0f);
-  float4 out_v;
-  out_v   = tmp_v * tmp_v.w;
-  out_v.w = tmp_v.w;
-  out[gid] = out_v;
-}
-
-
 /* RaGaBaA float -> RGB_GAMMA_U8 */
 __kernel void ragabaf_to_rgb_gamma_u8 (__global const float4 * in,
                                        __global       uchar  * out)
@@ -460,29 +419,6 @@ __kernel void yaf_to_rgbau8 (__global const float2 * in,
   out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
 }
 
-/* R'G'B'A u8 -> YA float */
-__kernel void rgba_gamma_u8_to_yaf (__global const uchar4 * in,
-                                    __global       float2 * out)
-{
-  int gid = get_global_id(0);
-  float4 in_v  = convert_float4(in[gid]) / 255.0f;
-  float4 tmp_v;
-  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),
-                   gamma_2_2_to_linear(in_v.y),
-                   gamma_2_2_to_linear(in_v.z),
-                   in_v.w);
-  float2 out_v;
-
-  float luminance = tmp_v.x * RGB_LUMINANCE_RED +
-                    tmp_v.y * RGB_LUMINANCE_GREEN +
-                    tmp_v.z * RGB_LUMINANCE_BLUE;
-
-  out_v.x = luminance;
-  out_v.y = tmp_v.w;
-
-  out[gid] = out_v;
-}
-
 /* YA float -> R'G'B'A u8 */
 __kernel void yaf_to_rgba_gamma_u8 (__global const float2 * in,
                                     __global       uchar4 * out)
@@ -499,34 +435,6 @@ __kernel void yaf_to_rgba_gamma_u8 (__global const float2 * in,
   out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
 }
 
-/* R'G'B' u8 -> YA float */
-__kernel void rgb_gamma_u8_to_yaf (__global const uchar  * in,
-                                   __global       float2 * out)
-{
-  int gid = get_global_id(0);
-#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
-  float3 in_v  = convert_float3(vload3 (gid, in)) / 255.0f;
-#else
-  uchar4 u_v = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255);
-  float4 in_v = convert_float4 (u_v) / 255.0f;
-#endif
-  float4 tmp_v;
-  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),
-                   gamma_2_2_to_linear(in_v.y),
-                   gamma_2_2_to_linear(in_v.z),
-                   1.0f);
-  float2 out_v;
-
-  float luminance = tmp_v.x * RGB_LUMINANCE_RED +
-                    tmp_v.y * RGB_LUMINANCE_GREEN +
-                    tmp_v.z * RGB_LUMINANCE_BLUE;
-
-  out_v.x = luminance;
-  out_v.y = tmp_v.w;
-
-  out[gid] = out_v;
-}
-
 /* YA float -> R'G'B' u8 */
 __kernel void yaf_to_rgb_gamma_u8 (__global const float2 * in,
                                    __global       uchar  * out)
@@ -558,20 +466,6 @@ __kernel void rgbaf_to_rgba_gamma_u8 (__global const float4 * in,
   out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
 }
 
-/* r'g'b'a u8 -> rgba float */
-__kernel void rgba_gamma_u8_to_rgbaf (__global const uchar4 * in,
-                                      __global       float4 * out)
-{
-  int gid = get_global_id(0);
-  float4 in_v  = convert_float4(in[gid]) / 255.0f;
-  float4 out_v;
-  out_v = (float4)(gamma_2_2_to_linear(in_v.x),
-                   gamma_2_2_to_linear(in_v.y),
-                   gamma_2_2_to_linear(in_v.z),
-                   in_v.w);
-  out[gid] = out_v;
-}
-
 /* R'G'B' u8 */
 
 /* rgba float -> r'g'b' u8 */
@@ -596,25 +490,3 @@ __kernel void rgbaf_to_rgb_gamma_u8 (__global const float4 * in,
   out[3 * gid + 2] = out_v.z;
 #endif
 }
-
-/* r'g'b' u8 -> rgba float */
-__kernel void rgb_gamma_u8_to_rgbaf (__global const uchar  * in,
-                                     __global       float4 * out)
-{
-  int gid = get_global_id(0);
-  float4 out_v;
-#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
-  uchar3 in_v;
-  float3 tmp_v;
-  in_v = vload3 (gid, in);
-  tmp_v = convert_float3(in_v) / 255.0f;
-#else
-  uchar4 in_v = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255);
-  float4 tmp_v = convert_float4 (in_v) / 255.0f;
-#endif
-  out_v = (float4)(gamma_2_2_to_linear(tmp_v.x),
-                   gamma_2_2_to_linear(tmp_v.y),
-                   gamma_2_2_to_linear(tmp_v.z),
-                   1.0f);
-  out[gid] = out_v;
-}
diff --git a/opencl/colors.cl.h b/opencl/colors.cl.h
index db59ed47..3d8c882a 100644
--- a/opencl/colors.cl.h
+++ b/opencl/colors.cl.h
@@ -119,24 +119,6 @@ static const char* colors_cl_source =
 "  out[gid] = convert_uchar4_sat_rte(255.0f * out_v);                          \n"
 "}                                                                             \n"
 "                                                                              \n"
-"/* RGBA_GAMMA_U8 -> RaGaBaA float */                                          \n"
-"__kernel void rgba_gamma_u8_to_ragabaf (__global const uchar4 * in,           \n"
-"                                        __global       float4 * out)          \n"
-"{                                                                             \n"
-"  int gid = get_global_id(0);                                                 \n"
-"  float4 in_v  = convert_float4(in[gid]) / 255.0f;                            \n"
-"  float4 tmp_v;                                                               \n"
-"  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),                               \n"
-"                   gamma_2_2_to_linear(in_v.y),                               \n"
-"                   gamma_2_2_to_linear(in_v.z),                               \n"
-"                   in_v.w);                                                   \n"
-"  float4 out_v;                                                               \n"
-"  out_v   = tmp_v * tmp_v.w;                                                  \n"
-"  out_v.w = tmp_v.w;                                                          \n"
-"  out[gid] = out_v;                                                           \n"
-"}                                                                             \n"
-"                                                                              \n"
-"                                                                              \n"
 "/* RaGaBaA float -> RGBA_GAMMA_U8 */                                          \n"
 "__kernel void ragabaf_to_rgba_gamma_u8 (__global const float4 * in,           \n"
 "                                        __global       uchar4 * out)          \n"
@@ -154,29 +136,6 @@ static const char* colors_cl_source =
 "  out[gid] = convert_uchar4_sat_rte(255.0f * out_v);                          \n"
 "}                                                                             \n"
 "                                                                              \n"
-"/* RGB_GAMMA_U8 -> RaGaBaA float */                                           \n"
-"__kernel void rgb_gamma_u8_to_ragabaf (__global const uchar  * in,            \n"
-"                                       __global       float4 * out)           \n"
-"{                                                                             \n"
-"  int gid = get_global_id(0);                                                 \n"
-"#if (__OPENCL_VERSION__ != CL_VERSION_1_0)                                    \n"
-"  float3 in_v  = convert_float3(vload3 (gid, in)) / 255.0f;                   \n"
-"#else                                                                         \n"
-"  uchar4 i4 = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255);  \n"
-"  float4 in_v = convert_float4 (i4) / 255.0f;                                 \n"
-"#endif                                                                        \n"
-"  float4 tmp_v;                                                               \n"
-"  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),                               \n"
-"                   gamma_2_2_to_linear(in_v.y),                               \n"
-"                   gamma_2_2_to_linear(in_v.z),                               \n"
-"                   1.0f);                                                     \n"
-"  float4 out_v;                                                               \n"
-"  out_v   = tmp_v * tmp_v.w;                                                  \n"
-"  out_v.w = tmp_v.w;                                                          \n"
-"  out[gid] = out_v;                                                           \n"
-"}                                                                             \n"
-"                                                                              \n"
-"                                                                              \n"
 "/* RaGaBaA float -> RGB_GAMMA_U8 */                                           \n"
 "__kernel void ragabaf_to_rgb_gamma_u8 (__global const float4 * in,            \n"
 "                                       __global       uchar  * out)           \n"
@@ -461,29 +420,6 @@ static const char* colors_cl_source =
 "  out[gid] = convert_uchar4_sat_rte(255.0f * out_v);                          \n"
 "}                                                                             \n"
 "                                                                              \n"
-"/* R'G'B'A u8 -> YA float */                                                  \n"
-"__kernel void rgba_gamma_u8_to_yaf (__global const uchar4 * in,               \n"
-"                                    __global       float2 * out)              \n"
-"{                                                                             \n"
-"  int gid = get_global_id(0);                                                 \n"
-"  float4 in_v  = convert_float4(in[gid]) / 255.0f;                            \n"
-"  float4 tmp_v;                                                               \n"
-"  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),                               \n"
-"                   gamma_2_2_to_linear(in_v.y),                               \n"
-"                   gamma_2_2_to_linear(in_v.z),                               \n"
-"                   in_v.w);                                                   \n"
-"  float2 out_v;                                                               \n"
-"                                                                              \n"
-"  float luminance = tmp_v.x * RGB_LUMINANCE_RED +                             \n"
-"                    tmp_v.y * RGB_LUMINANCE_GREEN +                           \n"
-"                    tmp_v.z * RGB_LUMINANCE_BLUE;                             \n"
-"                                                                              \n"
-"  out_v.x = luminance;                                                        \n"
-"  out_v.y = tmp_v.w;                                                          \n"
-"                                                                              \n"
-"  out[gid] = out_v;                                                           \n"
-"}                                                                             \n"
-"                                                                              \n"
 "/* YA float -> R'G'B'A u8 */                                                  \n"
 "__kernel void yaf_to_rgba_gamma_u8 (__global const float2 * in,               \n"
 "                                    __global       uchar4 * out)              \n"
@@ -500,34 +436,6 @@ static const char* colors_cl_source =
 "  out[gid] = convert_uchar4_sat_rte(255.0f * out_v);                          \n"
 "}                                                                             \n"
 "                                                                              \n"
-"/* R'G'B' u8 -> YA float */                                                   \n"
-"__kernel void rgb_gamma_u8_to_yaf (__global const uchar  * in,                \n"
-"                                   __global       float2 * out)               \n"
-"{                                                                             \n"
-"  int gid = get_global_id(0);                                                 \n"
-"#if (__OPENCL_VERSION__ != CL_VERSION_1_0)                                    \n"
-"  float3 in_v  = convert_float3(vload3 (gid, in)) / 255.0f;                   \n"
-"#else                                                                         \n"
-"  uchar4 u_v = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255); \n"
-"  float4 in_v = convert_float4 (u_v) / 255.0f;                                \n"
-"#endif                                                                        \n"
-"  float4 tmp_v;                                                               \n"
-"  tmp_v = (float4)(gamma_2_2_to_linear(in_v.x),                               \n"
-"                   gamma_2_2_to_linear(in_v.y),                               \n"
-"                   gamma_2_2_to_linear(in_v.z),                               \n"
-"                   1.0f);                                                     \n"
-"  float2 out_v;                                                               \n"
-"                                                                              \n"
-"  float luminance = tmp_v.x * RGB_LUMINANCE_RED +                             \n"
-"                    tmp_v.y * RGB_LUMINANCE_GREEN +                           \n"
-"                    tmp_v.z * RGB_LUMINANCE_BLUE;                             \n"
-"                                                                              \n"
-"  out_v.x = luminance;                                                        \n"
-"  out_v.y = tmp_v.w;                                                          \n"
-"                                                                              \n"
-"  out[gid] = out_v;                                                           \n"
-"}                                                                             \n"
-"                                                                              \n"
 "/* YA float -> R'G'B' u8 */                                                   \n"
 "__kernel void yaf_to_rgb_gamma_u8 (__global const float2 * in,                \n"
 "                                   __global       uchar  * out)               \n"
@@ -559,20 +467,6 @@ static const char* colors_cl_source =
 "  out[gid] = convert_uchar4_sat_rte(255.0f * out_v);                          \n"
 "}                                                                             \n"
 "                                                                              \n"
-"/* r'g'b'a u8 -> rgba float */                                                \n"
-"__kernel void rgba_gamma_u8_to_rgbaf (__global const uchar4 * in,             \n"
-"                                      __global       float4 * out)            \n"
-"{                                                                             \n"
-"  int gid = get_global_id(0);                                                 \n"
-"  float4 in_v  = convert_float4(in[gid]) / 255.0f;                            \n"
-"  float4 out_v;                                                               \n"
-"  out_v = (float4)(gamma_2_2_to_linear(in_v.x),                               \n"
-"                   gamma_2_2_to_linear(in_v.y),                               \n"
-"                   gamma_2_2_to_linear(in_v.z),                               \n"
-"                   in_v.w);                                                   \n"
-"  out[gid] = out_v;                                                           \n"
-"}                                                                             \n"
-"                                                                              \n"
 "/* R'G'B' u8 */                                                               \n"
 "                                                                              \n"
 "/* rgba float -> r'g'b' u8 */                                                 \n"
@@ -597,26 +491,4 @@ static const char* colors_cl_source =
 "  out[3 * gid + 2] = out_v.z;                                                 \n"
 "#endif                                                                        \n"
 "}                                                                             \n"
-"                                                                              \n"
-"/* r'g'b' u8 -> rgba float */                                                 \n"
-"__kernel void rgb_gamma_u8_to_rgbaf (__global const uchar  * in,              \n"
-"                                     __global       float4 * out)             \n"
-"{                                                                             \n"
-"  int gid = get_global_id(0);                                                 \n"
-"  float4 out_v;                                                               \n"
-"#if (__OPENCL_VERSION__ != CL_VERSION_1_0)                                    \n"
-"  uchar3 in_v;                                                                \n"
-"  float3 tmp_v;                                                               \n"
-"  in_v = vload3 (gid, in);                                                    \n"
-"  tmp_v = convert_float3(in_v) / 255.0f;                                      \n"
-"#else                                                                         \n"
-"  uchar4 in_v = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255);\n"
-"  float4 tmp_v = convert_float4 (in_v) / 255.0f;                              \n"
-"#endif                                                                        \n"
-"  out_v = (float4)(gamma_2_2_to_linear(tmp_v.x),                              \n"
-"                   gamma_2_2_to_linear(tmp_v.y),                              \n"
-"                   gamma_2_2_to_linear(tmp_v.z),                              \n"
-"                   1.0f);                                                     \n"
-"  out[gid] = out_v;                                                           \n"
-"}                                                                             \n"
 ;