[gegl] cl: More color conversion from and to R'G'B'A u8 and R'G'B' u8
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] cl: More color conversion from and to R'G'B'A u8 and R'G'B' u8
- Date: Sun, 1 Apr 2012 19:51:42 +0000 (UTC)
commit 0142c23b03f370b3a6d4b4b9897c3cb0f4ef9dde
Author: Victor Oliveira <victormatheus gmail com>
Date: Sun Apr 1 16:50:36 2012 -0300
cl: More color conversion from and to R'G'B'A u8 and R'G'B' u8
gegl/opencl/gegl-cl-color-kernel.h | 150 +++++++++++++++++++++++++++++++++++-
gegl/opencl/gegl-cl-color.c | 31 +++++++-
2 files changed, 179 insertions(+), 2 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color-kernel.h b/gegl/opencl/gegl-cl-color-kernel.h
index 6159080..056545b 100644
--- a/gegl/opencl/gegl-cl-color-kernel.h
+++ b/gegl/opencl/gegl-cl-color-kernel.h
@@ -96,6 +96,76 @@ static const char* kernel_color_source =
" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
"} \n"
+/* RGBA_GAMMA_U8 -> RaGaBaA float */
+"__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"
+
+/* RaGaBaA float -> RGBA_GAMMA_U8 */
+"__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"
+
+/* RGB_GAMMA_U8 -> RaGaBaA float */
+"__kernel void rgb_gamma_u8_to_ragabaf (__global const uchar * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float3 in_v = convert_float3(vload3 (gid, in)) / 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"
+" 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"
+
+/* RaGaBaA float -> RGB_GAMMA_U8 */
+"__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"
+" vstore3 (convert_uchar3_sat_rte(255.0f * out_v.xyz * out_v.w), gid, out); \n"
+"} \n"
+
/* -- R'G'B'A float -- */
/* rgba float -> r'g'b'a float */
@@ -309,7 +379,6 @@ static const char* kernel_color_source =
" out[gid] = out_v; \n"
"} \n"
-
/* RGBA u8 -> YA float */
"__kernel void rgbau8_to_yaf (__global const uchar4 * in, \n"
" __global float2 * out) \n"
@@ -339,6 +408,85 @@ static const char* kernel_color_source =
" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
"} \n"
+/* R'G'B'A u8 -> YA float */
+"__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"
+
+/* YA float -> R'G'B'A u8 */
+"__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"
+
+/* R'G'B' u8 -> YA float */
+"__kernel void rgb_gamma_u8_to_yaf (__global const uchar * in, \n"
+" __global float2 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float3 in_v = convert_float3(vload3 (gid, in)) / 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"
+" 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"
+
+/* YA float -> R'G'B' u8 */
+"__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"
+" 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"
+" vstore3 (convert_uchar3_sat_rte(255.0f * out_v.xyz * out_v.w), gid, out); \n"
+"} \n"
+
+
/* R'G'B'A u8 */
/* rgba float -> r'g'b'a u8 */
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 2f6c4a8..379a1ac 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -47,6 +47,16 @@ CL_RGBA_GAMMA_U8_TO_RGBAF = 22,
CL_RGBAF_TO_RGB_GAMMA_U8 = 23,
CL_RGB_GAMMA_U8_TO_RGBAF = 24,
+
+CL_RGBA_GAMMA_U8_TO_RAGABAF = 25,
+CL_RAGABAF_TO_RGBA_GAMMA_U8 = 26,
+CL_RGB_GAMMA_U8_TO_RAGABAF = 27,
+CL_RAGABAF_TO_RGB_GAMMA_U8 = 28,
+
+CL_RGBA_GAMMA_U8_TO_YAF = 29,
+CL_YAF_TO_RGBA_GAMMA_U8 = 30,
+CL_RGB_GAMMA_U8_TO_YAF = 31,
+CL_YAF_TO_RGB_GAMMA_U8 = 32,
};
void
@@ -86,6 +96,16 @@ gegl_cl_color_compile_kernels(void)
"rgbaf_to_rgb_gamma_u8", /* 23 */
"rgb_gamma_u8_to_rgbaf", /* 24 */
+ "rgba_gamma_u8_to_ragabaf", /* 25 */
+ "ragabaf_to_rgba_gamma_u8", /* 26 */
+ "rgb_gamma_u8_to_ragabaf", /* 27 */
+ "ragabaf_to_rgb_gamma_u8", /* 28 */
+
+ "rgba_gamma_u8_to_yaf", /* 29 */
+ "yaf_to_rgba_gamma_u8", /* 30 */
+ "rgb_gamma_u8_to_yaf", /* 31 */
+ "yaf_to_rgb_gamma_u8", /* 32 */
+
NULL};
format[0] = babl_format ("RGBA u8");
@@ -104,6 +124,7 @@ gegl_cl_color_compile_kernels(void)
}
+
static gint
choose_kernel (const Babl *in_format, const Babl *out_format)
{
@@ -113,8 +134,8 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
{
if (out_format == babl_format ("RGBA u8")) kernel = CL_RGBAF_TO_RGBAU8;
else if (out_format == babl_format ("RaGaBaA float")) kernel = CL_RGBAF_TO_RAGABAF;
- else if (out_format == babl_format ("R'G'B' u8")) kernel = CL_RGBAF_TO_RGB_GAMMA_U8;
else if (out_format == babl_format ("R'G'B'A u8")) kernel = CL_RGBAF_TO_RGBA_GAMMA_U8;
+ else if (out_format == babl_format ("R'G'B' u8")) kernel = CL_RGBAF_TO_RGB_GAMMA_U8;
else if (out_format == babl_format ("R'G'B'A float")) kernel = CL_RGBAF_TO_RGBA_GAMMA_F;
else if (out_format == babl_format ("Y'CbCrA float")) kernel = CL_RGBAF_TO_YCBCRAF;
else if (out_format == babl_format ("RGB u8")) kernel = CL_RGBAF_TO_RGBU8;
@@ -132,6 +153,8 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
{
if (out_format == babl_format ("RGBA float")) kernel = CL_RAGABAF_TO_RGBAF;
else if (out_format == babl_format ("RGBA u8")) kernel = CL_RAGABAF_TO_RGBAU8;
+ else if (out_format == babl_format ("R'G'B'A u8")) kernel = CL_RAGABAF_TO_RGBA_GAMMA_U8;
+ else if (out_format == babl_format ("R'G'B' u8")) kernel = CL_RAGABAF_TO_RGB_GAMMA_U8;
}
else if (in_format == babl_format ("R'G'B'A float"))
{
@@ -155,14 +178,20 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
{
if (out_format == babl_format ("RGBA float")) kernel = CL_YAF_TO_RGBAF;
else if (out_format == babl_format ("RGBA u8")) kernel = CL_YAF_TO_RGBAU8;
+ else if (out_format == babl_format ("R'G'B'A u8")) kernel = CL_YAF_TO_RGBA_GAMMA_U8;
+ else if (out_format == babl_format ("R'G'B' u8")) kernel = CL_YAF_TO_RGB_GAMMA_U8;
}
else if (in_format == babl_format ("R'G'B'A u8"))
{
if (out_format == babl_format ("RGBA float")) kernel = CL_RGBA_GAMMA_U8_TO_RGBAF;
+ else if (out_format == babl_format ("RaGaBaA float")) kernel = CL_RGBA_GAMMA_U8_TO_RAGABAF;
+ else if (out_format == babl_format ("YA float")) kernel = CL_RGBA_GAMMA_U8_TO_YAF;
}
else if (in_format == babl_format ("R'G'B' u8"))
{
if (out_format == babl_format ("RGBA float")) kernel = CL_RGB_GAMMA_U8_TO_RGBAF;
+ else if (out_format == babl_format ("RaGaBaA float")) kernel = CL_RGB_GAMMA_U8_TO_RAGABAF;
+ else if (out_format == babl_format ("YA float")) kernel = CL_RGB_GAMMA_U8_TO_YAF;
}
return kernel;
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
