[gegl] opencl: missing color kernels
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl: missing color kernels
- Date: Sun, 13 Jan 2013 02:22:13 +0000 (UTC)
commit ecad72c46949dd7abaec437960f159edd5d73899
Author: Victor Oliveira <victormatheus gmail com>
Date: Sun Jan 13 00:21:16 2013 -0200
opencl: missing color kernels
opencl/colors.cl | 635 +++++++++++++++++++++++++++++++++++++++++++++++++++
opencl/colors.cl.h | 637 ++++++++++++++++++++++++++++++++++++++++++++++++++++
2 files changed, 1272 insertions(+), 0 deletions(-)
---
diff --git a/opencl/colors.cl b/opencl/colors.cl
new file mode 100644
index 0000000..8f43a61
--- /dev/null
+++ b/opencl/colors.cl
@@ -0,0 +1,635 @@
+/* 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)
+ */
+
+/* This is almost a copy-paste from babl/base color conversion functions
+
+ XXX: This code is very hard to maintain and keep in sync with BABL, we should
+ think in something better
+*/
+
+/* Alpha threshold used in the reference implementation for
+ * un-pre-multiplication of color data:
+ *
+ * 0.01 / (2^16 - 1)
+ */
+#define BABL_ALPHA_THRESHOLD 0.0f
+
+/* babl reference file: babl/base/util.h */
+float linear_to_gamma_2_2 (float value)
+{
+ if (value > 0.003130804954f)
+ return 1.055f * native_powr (value, (1.0f/2.4f)) - 0.055f;
+ return 12.92f * value;
+}
+
+float gamma_2_2_to_linear (float value)
+{
+ if (value > 0.04045f)
+ return native_powr ((value + 0.055f) / 1.055f, 2.4f);
+ return value / 12.92f;
+}
+
+/* -- RGBA float/u8 -- */
+
+/* RGBA u8 -> RGBA float */
+__kernel void rgbau8_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 = in_v;
+ out[gid] = out_v;
+}
+
+/* RGBA float -> RGBA u8 */
+__kernel void rgbaf_to_rgbau8 (__global const float4 * in,
+ __global uchar4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v = in_v;
+ out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
+}
+
+/* -- RaGaBaA float -- */
+
+/* RGBA float -> RaGaBaA float */
+__kernel void rgbaf_to_ragabaf (__global const float4 * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+ out_v = in_v * in_v.w;
+ out_v.w = in_v.w;
+ out[gid] = out_v;
+}
+
+
+/* RaGaBaA float -> RGBA float */
+__kernel void ragabaf_to_rgbaf (__global const float4 * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+ out_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f);
+ out_v.w = in_v.w;
+ out[gid] = out_v;
+}
+
+/* RGBA u8 -> RaGaBaA float */
+__kernel void rgbau8_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 out_v;
+ out_v = in_v * in_v.w;
+ out_v.w = in_v.w;
+ out[gid] = out_v;
+}
+
+
+/* RaGaBaA float -> RGBA u8 */
+__kernel void ragabaf_to_rgbau8 (__global const float4 * in,
+ __global uchar4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+ out_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f);
+ out_v.w = in_v.w;
+ 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)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 tmp_v;
+ tmp_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f);
+ tmp_v.w = in_v.w;
+ float4 out_v;
+ out_v = (float4)(linear_to_gamma_2_2(tmp_v.x),
+ linear_to_gamma_2_2(tmp_v.y),
+ linear_to_gamma_2_2(tmp_v.z),
+ tmp_v.w);
+ 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)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 tmp_v;
+ tmp_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f);
+ tmp_v.w = in_v.w;
+ float4 out_v;
+ out_v = (float4)(linear_to_gamma_2_2(tmp_v.x),
+ linear_to_gamma_2_2(tmp_v.y),
+ linear_to_gamma_2_2(tmp_v.z),
+ tmp_v.w);
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ vstore3 (convert_uchar3_sat_rte(255.0f * out_v.xyz), gid, out);
+#else
+ uchar4 sat = convert_uchar4_sat_rte (255.0f * out_v);
+ out[3 * gid] = sat.x;
+ out[3 * gid + 1] = sat.y;
+ out[3 * gid + 2] = sat.z;
+#endif
+}
+
+/* -- R'G'B'A float -- */
+
+/* rgba float -> r'g'b'a float */
+__kernel void rgbaf_to_rgba_gamma_f (__global const float4 * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+ out_v = (float4)(linear_to_gamma_2_2(in_v.x),
+ linear_to_gamma_2_2(in_v.y),
+ linear_to_gamma_2_2(in_v.z),
+ in_v.w);
+ out[gid] = out_v;
+}
+
+/* r'g'b'a float -> rgba float */
+__kernel void rgba_gamma_f_to_rgbaf (__global const float4 * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ 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;
+}
+
+/* rgba u8 -> r'g'b'a float */
+__kernel void rgbau8_to_rgba_gamma_f (__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)(linear_to_gamma_2_2(in_v.x),
+ linear_to_gamma_2_2(in_v.y),
+ linear_to_gamma_2_2(in_v.z),
+ in_v.w);
+ out[gid] = out_v;
+}
+
+/* r'g'b'a float -> rgba u8 */
+__kernel void rgba_gamma_f_to_rgbau8 (__global const float4 * in,
+ __global uchar4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ 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] = convert_uchar4_sat_rte(255.0f * out_v);
+}
+
+/* -- Y'CbCrA float -- */
+
+/* RGBA float -> Y'CbCrA float */
+__kernel void rgbaf_to_ycbcraf (__global const float4 * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ float3 rgb = (float3)(linear_to_gamma_2_2(in_v.x),
+ linear_to_gamma_2_2(in_v.y),
+ linear_to_gamma_2_2(in_v.z));
+#else
+ float4 rgb = (float4)(linear_to_gamma_2_2(in_v.x),
+ linear_to_gamma_2_2(in_v.y),
+ linear_to_gamma_2_2(in_v.z), 1.0f);
+#endif
+
+ out_v = (float4)( 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z,
+ -0.168736f * rgb.x - 0.331264f * rgb.y + 0.5f * rgb.z,
+ 0.5f * rgb.x - 0.418688f * rgb.y - 0.081312f * rgb.z,
+ in_v.w);
+ out[gid] = out_v;
+}
+
+/* Y'CbCrA float -> RGBA float */
+__kernel void ycbcraf_to_rgbaf (__global const float4 * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+
+ float4 rgb = (float4)(1.0f * in_v.x + 0.0f * in_v.y + 1.40200f * in_v.z,
+ 1.0f * in_v.x - 0.344136f * in_v.y - 0.71414136f * in_v.z,
+ 1.0f * in_v.x + 1.772f * in_v.y + 0.0f * in_v.z,
+ 0.0f);
+
+ out_v = (float4)(gamma_2_2_to_linear(rgb.x),
+ gamma_2_2_to_linear(rgb.y),
+ gamma_2_2_to_linear(rgb.z),
+ in_v.w);
+ out[gid] = out_v;
+}
+
+/* RGBA u8 -> Y'CbCrA float */
+__kernel void rgbau8_to_ycbcraf (__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;
+
+ float4 rgb = (float4)(linear_to_gamma_2_2(in_v.x),
+ linear_to_gamma_2_2(in_v.y),
+ linear_to_gamma_2_2(in_v.z),
+ 0.0f);
+
+ out_v = (float4)( 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z,
+ -0.168736f * rgb.x - 0.331264f * rgb.y + 0.5f * rgb.z,
+ 0.5f * rgb.x - 0.418688f * rgb.y - 0.081312f * rgb.z,
+ in_v.w);
+ out[gid] = out_v;
+}
+
+/* Y'CbCrA float -> RGBA u8 */
+__kernel void ycbcraf_to_rgbau8 (__global const float4 * in,
+ __global uchar4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+
+ float4 rgb = (float4)(1.0f * in_v.x + 0.0f * in_v.y + 1.40200f * in_v.z,
+ 1.0f * in_v.x - 0.344136f * in_v.y - 0.71414136f * in_v.z,
+ 1.0f * in_v.x + 1.772f * in_v.y + 0.0f * in_v.z,
+ 0.0f);
+
+ out_v = (float4)(gamma_2_2_to_linear(rgb.x),
+ gamma_2_2_to_linear(rgb.y),
+ gamma_2_2_to_linear(rgb.z),
+ in_v.w);
+ out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
+}
+
+/* -- RGB u8 -- */
+
+/* RGB u8 -> RGBA float */
+__kernel void rgbu8_to_rgbaf (__global const uchar * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ uchar3 in_v;
+ float4 out_v;
+ in_v = vload3 (gid, in);
+ out_v.xyz = convert_float3(in_v) / 255.0f;
+ out_v.w = 1.0f;
+#else
+ uchar4 in_v = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255);
+ float4 out_v = convert_float4 (in_v) / 255.0f;
+#endif
+ out[gid] = out_v;
+}
+
+/* RGBA float -> RGB u8 */
+__kernel void rgbaf_to_rgbu8 (__global const float4 * in,
+ __global uchar * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ uchar3 out_v = convert_uchar3_sat_rte(255.0f * in_v.xyz);
+ vstore3 (out_v, gid, out);
+#else
+ uchar4 out_v = convert_uchar4_sat_rte(255.0f * in_v);
+ out[3 * gid] = out_v.x;
+ out[3 * gid + 1] = out_v.y;
+ out[3 * gid + 2] = out_v.z;
+#endif
+}
+
+/* -- Y u8 -- */
+
+/* Y u8 -> Y float */
+__kernel void yu8_to_yf (__global const uchar * in,
+ __global float * out)
+{
+ int gid = get_global_id(0);
+ float in_v = convert_float (in[gid]) / 255.0f;
+ float out_v;
+ out_v = in_v;
+ out[gid] = out_v;
+}
+
+/* -- YA float -- */
+
+/* babl reference file: babl/base/rgb-constants.h */
+#if 0
+#define CONTEMPORARY_MONITOR
+#endif
+
+#ifdef CONTEMPORARY_MONITOR
+ /* source: http://www.poynton.com/ColorFAQ.html */
+ #define RGB_LUMINANCE_RED (0.212671f)
+ #define RGB_LUMINANCE_GREEN (0.715160f)
+ #define RGB_LUMINANCE_BLUE (0.072169f)
+#else
+ /* this is not correct, but the constants are kept around */
+ #define RGB_LUMA_RED (0.299)
+ #define RGB_LUMA_GREEN (0.587)
+ #define RGB_LUMA_BLUE (0.114)
+ #define RGB_LUMINANCE_RED RGB_LUMA_RED
+ #define RGB_LUMINANCE_GREEN RGB_LUMA_GREEN
+ #define RGB_LUMINANCE_BLUE RGB_LUMA_BLUE
+#endif
+
+/* RGBA float -> YA float */
+__kernel void rgbaf_to_yaf (__global const float4 * in,
+ __global float2 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float2 out_v;
+
+ float luminance = in_v.x * RGB_LUMINANCE_RED +
+ in_v.y * RGB_LUMINANCE_GREEN +
+ in_v.z * RGB_LUMINANCE_BLUE;
+
+ out_v.x = luminance;
+ out_v.y = in_v.w;
+
+ out[gid] = out_v;
+}
+
+/* YA float -> RGBA float */
+__kernel void yaf_to_rgbaf (__global const float2 * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float2 in_v = in[gid];
+ float4 out_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y);
+
+ out[gid] = out_v;
+}
+
+/* RGBA u8 -> YA float */
+__kernel void rgbau8_to_yaf (__global const uchar4 * in,
+ __global float2 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = convert_float4(in[gid]) / 255.0f;
+ float2 out_v;
+
+ float luminance = in_v.x * RGB_LUMINANCE_RED +
+ in_v.y * RGB_LUMINANCE_GREEN +
+ in_v.z * RGB_LUMINANCE_BLUE;
+
+ out_v.x = luminance;
+ out_v.y = in_v.w;
+
+ out[gid] = out_v;
+}
+
+/* YA float -> RGBA u8 */
+__kernel void yaf_to_rgbau8 (__global const float2 * in,
+ __global uchar4 * out)
+{
+ int gid = get_global_id(0);
+ float2 in_v = in[gid];
+ float4 out_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y);
+
+ 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)
+{
+ int gid = get_global_id(0);
+ float2 in_v = in[gid];
+ float4 tmp_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y);
+
+ float4 out_v;
+ out_v = (float4)(linear_to_gamma_2_2(tmp_v.x),
+ linear_to_gamma_2_2(tmp_v.y),
+ linear_to_gamma_2_2(tmp_v.z),
+ tmp_v.w);
+ 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)
+{
+ int gid = get_global_id(0);
+ float2 in_v = in[gid];
+ uchar tmp = convert_uchar_sat_rte (255.0f * linear_to_gamma_2_2 (in_v.x));
+
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ vstore3 ((uchar3)tmp, gid, out);
+#else
+ out[3 * gid] = out[3 * gid + 1] = out[3 * gid + 2] = tmp;
+#endif
+}
+
+/* R'G'B'A u8 */
+
+/* rgba float -> r'g'b'a u8 */
+__kernel void rgbaf_to_rgba_gamma_u8 (__global const float4 * in,
+ __global uchar4 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+ out_v = (float4)(linear_to_gamma_2_2(in_v.x),
+ linear_to_gamma_2_2(in_v.y),
+ linear_to_gamma_2_2(in_v.z),
+ in_v.w);
+ 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 */
+__kernel void rgbaf_to_rgb_gamma_u8 (__global const float4 * in,
+ __global uchar * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 tmp_v;
+ tmp_v = (float4)(linear_to_gamma_2_2(in_v.x),
+ linear_to_gamma_2_2(in_v.y),
+ linear_to_gamma_2_2(in_v.z),
+ in_v.w);
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ uchar3 out_v;
+ out_v = convert_uchar3_sat_rte(255.0f * tmp_v.xyz);
+ vstore3 (out_v, gid, out);
+#else
+ uchar4 out_v = convert_uchar4_sat_rte (255.0f * tmp_v);
+ out[3 * gid] = out_v.x;
+ out[3 * gid + 1] = out_v.y;
+ 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
new file mode 100644
index 0000000..60dd747
--- /dev/null
+++ b/opencl/colors.cl.h
@@ -0,0 +1,637 @@
+static const char* colors_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"
+" */ \n"
+" \n"
+"/* This is almost a copy-paste from babl/base color conversion functions \n"
+" \n"
+" XXX: This code is very hard to maintain and keep in sync with BABL, we should\n"
+" think in something better \n"
+"*/ \n"
+" \n"
+"/* Alpha threshold used in the reference implementation for \n"
+" * un-pre-multiplication of color data: \n"
+" * \n"
+" * 0.01 / (2^16 - 1) \n"
+" */ \n"
+"#define BABL_ALPHA_THRESHOLD 0.0f \n"
+" \n"
+"/* babl reference file: babl/base/util.h */ \n"
+"float linear_to_gamma_2_2 (float value) \n"
+"{ \n"
+" if (value > 0.003130804954f) \n"
+" return 1.055f * native_powr (value, (1.0f/2.4f)) - 0.055f; \n"
+" return 12.92f * value; \n"
+"} \n"
+" \n"
+"float gamma_2_2_to_linear (float value) \n"
+"{ \n"
+" if (value > 0.04045f) \n"
+" return native_powr ((value + 0.055f) / 1.055f, 2.4f); \n"
+" return value / 12.92f; \n"
+"} \n"
+" \n"
+"/* -- RGBA float/u8 -- */ \n"
+" \n"
+"/* RGBA u8 -> RGBA float */ \n"
+"__kernel void rgbau8_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 = in_v; \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* RGBA float -> RGBA u8 */ \n"
+"__kernel void rgbaf_to_rgbau8 (__global const float4 * in, \n"
+" __global uchar4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v = in_v; \n"
+" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
+"} \n"
+" \n"
+"/* -- RaGaBaA float -- */ \n"
+" \n"
+"/* RGBA float -> RaGaBaA float */ \n"
+"__kernel void rgbaf_to_ragabaf (__global const float4 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v = in_v * in_v.w; \n"
+" out_v.w = in_v.w; \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+" \n"
+"/* RaGaBaA float -> RGBA float */ \n"
+"__kernel void ragabaf_to_rgbaf (__global const float4 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f); \n"
+" out_v.w = in_v.w; \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* RGBA u8 -> RaGaBaA float */ \n"
+"__kernel void rgbau8_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 out_v; \n"
+" out_v = in_v * in_v.w; \n"
+" out_v.w = in_v.w; \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+" \n"
+"/* RaGaBaA float -> RGBA u8 */ \n"
+"__kernel void ragabaf_to_rgbau8 (__global const float4 * in, \n"
+" __global uchar4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f); \n"
+" out_v.w = in_v.w; \n"
+" 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"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 tmp_v; \n"
+" tmp_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f); \n"
+" tmp_v.w = in_v.w; \n"
+" float4 out_v; \n"
+" out_v = (float4)(linear_to_gamma_2_2(tmp_v.x), \n"
+" linear_to_gamma_2_2(tmp_v.y), \n"
+" linear_to_gamma_2_2(tmp_v.z), \n"
+" tmp_v.w); \n"
+" 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"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 tmp_v; \n"
+" tmp_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f); \n"
+" tmp_v.w = in_v.w; \n"
+" float4 out_v; \n"
+" out_v = (float4)(linear_to_gamma_2_2(tmp_v.x), \n"
+" linear_to_gamma_2_2(tmp_v.y), \n"
+" linear_to_gamma_2_2(tmp_v.z), \n"
+" tmp_v.w); \n"
+"#if (__OPENCL_VERSION__ != CL_VERSION_1_0) \n"
+" vstore3 (convert_uchar3_sat_rte(255.0f * out_v.xyz), gid, out); \n"
+"#else \n"
+" uchar4 sat = convert_uchar4_sat_rte (255.0f * out_v); \n"
+" out[3 * gid] = sat.x; \n"
+" out[3 * gid + 1] = sat.y; \n"
+" out[3 * gid + 2] = sat.z; \n"
+"#endif \n"
+"} \n"
+" \n"
+"/* -- R'G'B'A float -- */ \n"
+" \n"
+"/* rgba float -> r'g'b'a float */ \n"
+"__kernel void rgbaf_to_rgba_gamma_f (__global const float4 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v = (float4)(linear_to_gamma_2_2(in_v.x), \n"
+" linear_to_gamma_2_2(in_v.y), \n"
+" linear_to_gamma_2_2(in_v.z), \n"
+" in_v.w); \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* r'g'b'a float -> rgba float */ \n"
+"__kernel void rgba_gamma_f_to_rgbaf (__global const float4 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \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"
+"/* rgba u8 -> r'g'b'a float */ \n"
+"__kernel void rgbau8_to_rgba_gamma_f (__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)(linear_to_gamma_2_2(in_v.x), \n"
+" linear_to_gamma_2_2(in_v.y), \n"
+" linear_to_gamma_2_2(in_v.z), \n"
+" in_v.w); \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* r'g'b'a float -> rgba u8 */ \n"
+"__kernel void rgba_gamma_f_to_rgbau8 (__global const float4 * in, \n"
+" __global uchar4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \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] = convert_uchar4_sat_rte(255.0f * out_v); \n"
+"} \n"
+" \n"
+"/* -- Y'CbCrA float -- */ \n"
+" \n"
+"/* RGBA float -> Y'CbCrA float */ \n"
+"__kernel void rgbaf_to_ycbcraf (__global const float4 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" \n"
+"#if (__OPENCL_VERSION__ != CL_VERSION_1_0) \n"
+" float3 rgb = (float3)(linear_to_gamma_2_2(in_v.x), \n"
+" linear_to_gamma_2_2(in_v.y), \n"
+" linear_to_gamma_2_2(in_v.z)); \n"
+"#else \n"
+" float4 rgb = (float4)(linear_to_gamma_2_2(in_v.x), \n"
+" linear_to_gamma_2_2(in_v.y), \n"
+" linear_to_gamma_2_2(in_v.z), 1.0f); \n"
+"#endif \n"
+" \n"
+" out_v = (float4)( 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z,\n"
+" -0.168736f * rgb.x - 0.331264f * rgb.y + 0.5f * rgb.z,\n"
+" 0.5f * rgb.x - 0.418688f * rgb.y - 0.081312f * rgb.z,\n"
+" in_v.w); \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* Y'CbCrA float -> RGBA float */ \n"
+"__kernel void ycbcraf_to_rgbaf (__global const float4 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" \n"
+" float4 rgb = (float4)(1.0f * in_v.x + 0.0f * in_v.y + 1.40200f * in_v.z,\n"
+" 1.0f * in_v.x - 0.344136f * in_v.y - 0.71414136f * in_v.z,\n"
+" 1.0f * in_v.x + 1.772f * in_v.y + 0.0f * in_v.z,\n"
+" 0.0f); \n"
+" \n"
+" out_v = (float4)(gamma_2_2_to_linear(rgb.x), \n"
+" gamma_2_2_to_linear(rgb.y), \n"
+" gamma_2_2_to_linear(rgb.z), \n"
+" in_v.w); \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* RGBA u8 -> Y'CbCrA float */ \n"
+"__kernel void rgbau8_to_ycbcraf (__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"
+" \n"
+" float4 rgb = (float4)(linear_to_gamma_2_2(in_v.x), \n"
+" linear_to_gamma_2_2(in_v.y), \n"
+" linear_to_gamma_2_2(in_v.z), \n"
+" 0.0f); \n"
+" \n"
+" out_v = (float4)( 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z,\n"
+" -0.168736f * rgb.x - 0.331264f * rgb.y + 0.5f * rgb.z,\n"
+" 0.5f * rgb.x - 0.418688f * rgb.y - 0.081312f * rgb.z,\n"
+" in_v.w); \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* Y'CbCrA float -> RGBA u8 */ \n"
+"__kernel void ycbcraf_to_rgbau8 (__global const float4 * in, \n"
+" __global uchar4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" \n"
+" float4 rgb = (float4)(1.0f * in_v.x + 0.0f * in_v.y + 1.40200f * in_v.z,\n"
+" 1.0f * in_v.x - 0.344136f * in_v.y - 0.71414136f * in_v.z,\n"
+" 1.0f * in_v.x + 1.772f * in_v.y + 0.0f * in_v.z,\n"
+" 0.0f); \n"
+" \n"
+" out_v = (float4)(gamma_2_2_to_linear(rgb.x), \n"
+" gamma_2_2_to_linear(rgb.y), \n"
+" gamma_2_2_to_linear(rgb.z), \n"
+" in_v.w); \n"
+" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
+"} \n"
+" \n"
+"/* -- RGB u8 -- */ \n"
+" \n"
+"/* RGB u8 -> RGBA float */ \n"
+"__kernel void rgbu8_to_rgbaf (__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"
+" uchar3 in_v; \n"
+" float4 out_v; \n"
+" in_v = vload3 (gid, in); \n"
+" out_v.xyz = convert_float3(in_v) / 255.0f; \n"
+" out_v.w = 1.0f; \n"
+"#else \n"
+" uchar4 in_v = (uchar4) (in[3 * gid], in[3 * gid + 1], in[3 * gid + 2], 255);\n"
+" float4 out_v = convert_float4 (in_v) / 255.0f; \n"
+"#endif \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* RGBA float -> RGB u8 */ \n"
+"__kernel void rgbaf_to_rgbu8 (__global const float4 * in, \n"
+" __global uchar * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+"#if (__OPENCL_VERSION__ != CL_VERSION_1_0) \n"
+" uchar3 out_v = convert_uchar3_sat_rte(255.0f * in_v.xyz); \n"
+" vstore3 (out_v, gid, out); \n"
+"#else \n"
+" uchar4 out_v = convert_uchar4_sat_rte(255.0f * in_v); \n"
+" out[3 * gid] = out_v.x; \n"
+" out[3 * gid + 1] = out_v.y; \n"
+" out[3 * gid + 2] = out_v.z; \n"
+"#endif \n"
+"} \n"
+" \n"
+"/* -- Y u8 -- */ \n"
+" \n"
+"/* Y u8 -> Y float */ \n"
+"__kernel void yu8_to_yf (__global const uchar * in, \n"
+" __global float * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float in_v = convert_float (in[gid]) / 255.0f; \n"
+" float out_v; \n"
+" out_v = in_v; \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* -- YA float -- */ \n"
+" \n"
+"/* babl reference file: babl/base/rgb-constants.h */ \n"
+"#if 0 \n"
+"#define CONTEMPORARY_MONITOR \n"
+"#endif \n"
+" \n"
+"#ifdef CONTEMPORARY_MONITOR \n"
+" /* source: http://www.poynton.com/ColorFAQ.html */ \n"
+" #define RGB_LUMINANCE_RED (0.212671f) \n"
+" #define RGB_LUMINANCE_GREEN (0.715160f) \n"
+" #define RGB_LUMINANCE_BLUE (0.072169f) \n"
+"#else \n"
+" /* this is not correct, but the constants are kept around */ \n"
+" #define RGB_LUMA_RED (0.299) \n"
+" #define RGB_LUMA_GREEN (0.587) \n"
+" #define RGB_LUMA_BLUE (0.114) \n"
+" #define RGB_LUMINANCE_RED RGB_LUMA_RED \n"
+" #define RGB_LUMINANCE_GREEN RGB_LUMA_GREEN \n"
+" #define RGB_LUMINANCE_BLUE RGB_LUMA_BLUE \n"
+"#endif \n"
+" \n"
+"/* RGBA float -> YA float */ \n"
+"__kernel void rgbaf_to_yaf (__global const float4 * in, \n"
+" __global float2 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float2 out_v; \n"
+" \n"
+" float luminance = in_v.x * RGB_LUMINANCE_RED + \n"
+" in_v.y * RGB_LUMINANCE_GREEN + \n"
+" in_v.z * RGB_LUMINANCE_BLUE; \n"
+" \n"
+" out_v.x = luminance; \n"
+" out_v.y = in_v.w; \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* YA float -> RGBA float */ \n"
+"__kernel void yaf_to_rgbaf (__global const float2 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float2 in_v = in[gid]; \n"
+" float4 out_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y); \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* RGBA u8 -> YA float */ \n"
+"__kernel void rgbau8_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"
+" float2 out_v; \n"
+" \n"
+" float luminance = in_v.x * RGB_LUMINANCE_RED + \n"
+" in_v.y * RGB_LUMINANCE_GREEN + \n"
+" in_v.z * RGB_LUMINANCE_BLUE; \n"
+" \n"
+" out_v.x = luminance; \n"
+" out_v.y = in_v.w; \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
+"/* YA float -> RGBA u8 */ \n"
+"__kernel void yaf_to_rgbau8 (__global const float2 * in, \n"
+" __global uchar4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float2 in_v = in[gid]; \n"
+" float4 out_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y); \n"
+" \n"
+" 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"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float2 in_v = in[gid]; \n"
+" float4 tmp_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y); \n"
+" \n"
+" float4 out_v; \n"
+" out_v = (float4)(linear_to_gamma_2_2(tmp_v.x), \n"
+" linear_to_gamma_2_2(tmp_v.y), \n"
+" linear_to_gamma_2_2(tmp_v.z), \n"
+" tmp_v.w); \n"
+" 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"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float2 in_v = in[gid]; \n"
+" uchar tmp = convert_uchar_sat_rte (255.0f * linear_to_gamma_2_2 (in_v.x)); \n"
+" \n"
+"#if (__OPENCL_VERSION__ != CL_VERSION_1_0) \n"
+" vstore3 ((uchar3)tmp, gid, out); \n"
+"#else \n"
+" out[3 * gid] = out[3 * gid + 1] = out[3 * gid + 2] = tmp; \n"
+"#endif \n"
+"} \n"
+" \n"
+"/* R'G'B'A u8 */ \n"
+" \n"
+"/* rgba float -> r'g'b'a u8 */ \n"
+"__kernel void rgbaf_to_rgba_gamma_u8 (__global const float4 * in, \n"
+" __global uchar4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v = (float4)(linear_to_gamma_2_2(in_v.x), \n"
+" linear_to_gamma_2_2(in_v.y), \n"
+" linear_to_gamma_2_2(in_v.z), \n"
+" in_v.w); \n"
+" 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"
+"__kernel void rgbaf_to_rgb_gamma_u8 (__global const float4 * in, \n"
+" __global uchar * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 tmp_v; \n"
+" tmp_v = (float4)(linear_to_gamma_2_2(in_v.x), \n"
+" linear_to_gamma_2_2(in_v.y), \n"
+" linear_to_gamma_2_2(in_v.z), \n"
+" in_v.w); \n"
+"#if (__OPENCL_VERSION__ != CL_VERSION_1_0) \n"
+" uchar3 out_v; \n"
+" out_v = convert_uchar3_sat_rte(255.0f * tmp_v.xyz); \n"
+" vstore3 (out_v, gid, out); \n"
+"#else \n"
+" uchar4 out_v = convert_uchar4_sat_rte (255.0f * tmp_v); \n"
+" out[3 * gid] = out_v.x; \n"
+" out[3 * gid + 1] = out_v.y; \n"
+" 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]
