Commit 252bd5d1 authored by Daniel Sabo's avatar Daniel Sabo

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

No performance testing done yet, but fixes some accuracy issues.
parent a40cb7fd
......@@ -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
......
......@@ -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
......
/* 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
This diff is collapsed.
......@@ -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;
}
......@@ -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"