[gegl] opencl: Replace u8 -> float gamma conversions with a small LUT
- From: Daniel Sabo <daniels src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl: Replace u8 -> float gamma conversions with a small LUT
- Date: Sat, 2 Nov 2013 04:41:35 +0000 (UTC)
commit 252bd5d111a860f368d67f7f45f678d6019abe93
Author: Daniel Sabo <DanielSabo gmail com>
Date: Fri Nov 1 16:38:52 2013 -0700
opencl: Replace u8 -> float gamma conversions with a small LUT
No performance testing done yet, but fixes some accuracy issues.
gegl/opencl/gegl-cl-color.c | 93 ++++++++---
gegl/opencl/gegl-cl-color.h | 2 +-
opencl/colors-8bit-lut.cl | 382 ++++++++++++++++++++++++++++++++++++++++++
opencl/colors-8bit-lut.cl.h | 384 +++++++++++++++++++++++++++++++++++++++++++
opencl/colors.cl | 128 --------------
opencl/colors.cl.h | 128 --------------
6 files changed, 837 insertions(+), 280 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index f28b52e..804dc53 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 fd0a731..35d14d5 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 0000000..93c00a7
--- /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 0000000..0567a78
--- /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 c311b22..3ff0820 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 db59ed4..3d8c882 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"
;
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
