[gegl] colors.cl: Add some conversions
- From: Daniel Sabo <daniels src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] colors.cl: Add some conversions
- Date: Sun, 3 Nov 2013 00:43:17 +0000 (UTC)
commit e90df9bc4d11087c6089746fd22739a71912d71d
Author: Daniel Sabo <DanielSabo gmail com>
Date: Sat Nov 2 15:56:19 2013 -0700
colors.cl: Add some conversions
gegl/opencl/gegl-cl-color.c | 16 ++++++++++++++-
opencl/colors.cl | 44 +++++++++++++++++++++++++++++++++++++++++++
opencl/colors.cl.h | 44 +++++++++++++++++++++++++++++++++++++++++++
3 files changed, 103 insertions(+), 1 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 401fa23..ae5fcf7 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -142,7 +142,17 @@ gegl_cl_color_compile_kernels (void)
{ babl_format ("RaGaBaA float"), babl_format("R'G'B' u8"), "ragabaf_to_rgb_gamma_u8", NULL },
{ babl_format ("YA float"), babl_format("R'G'B'A u8"), "yaf_to_rgba_gamma_u8", NULL },
- { babl_format ("YA float"), babl_format("R'G'B' u8"), "yaf_to_rgb_gamma_u8", NULL }
+ { babl_format ("YA float"), babl_format("R'G'B' u8"), "yaf_to_rgb_gamma_u8", NULL },
+
+ { babl_format ("YA float"), babl_format ("RaGaBaA float"), "yaf_to_ragabaf", NULL },
+
+ { babl_format ("Y float"), babl_format ("RaGaBaA float"), "yf_to_ragabaf", NULL },
+
+ { babl_format ("RGBA float"), babl_format ("RGB float"), "rgbaf_to_rgbf", NULL },
+
+ /* Reuse some conversions */
+ { babl_format ("R'G'B' u8"), babl_format ("R'G'B'A float"), "rgbu8_to_rgbaf", NULL },
+ { babl_format ("R'G'B'A u8"), babl_format ("R'G'B'A float"), "rgbau8_to_rgbaf", NULL },
};
ColorConversionInfo lut8_conversions[] = {
@@ -220,6 +230,10 @@ gegl_cl_color_supported (const Babl *in_format,
if (color_kernels_hash && find_color_kernel (in_format, out_format))
return GEGL_CL_COLOR_CONVERT;
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Missing OpenCL conversion for %s -> %s",
+ babl_get_name (in_format),
+ babl_get_name (out_format));
+
return GEGL_CL_COLOR_NOT_SUPPORTED;
}
diff --git a/opencl/colors.cl b/opencl/colors.cl
index 382cee9..9448a6d 100644
--- a/opencl/colors.cl
+++ b/opencl/colors.cl
@@ -66,6 +66,23 @@ __kernel void rgbaf_to_rgbau8 (__global const float4 * in,
out[gid] = convert_uchar4_sat_rte(255.0f * out_v);
}
+
+/* RGBA float -> RGB float */
+__kernel void rgbaf_to_rgbf (__global const float4 * in,
+ __global float * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ vstore3 (in_v.xyz, gid, out);
+#else
+ out[3 * gid] = in_v.x;
+ out[3 * gid + 1] = in_v.y;
+ out[3 * gid + 2] = in_v.z;
+#endif
+}
+
/* -- RaGaBaA float -- */
/* RGBA float -> RaGaBaA float */
@@ -354,6 +371,19 @@ __kernel void yu8_to_yf (__global const uchar * in,
out[gid] = out_v;
}
+/* -- Y float -- */
+
+/* Y float -> RaGaBaA float */
+__kernel void yf_to_ragabaf (__global const float * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float y = in[gid];
+ float4 out_v = (float4) (y, y, y, 1.0f);
+
+ out[gid] = out_v;
+}
+
/* -- YA float -- */
/* babl reference file: babl/base/rgb-constants.h */
@@ -390,6 +420,20 @@ __kernel void yaf_to_rgbaf (__global const float2 * in,
out[gid] = out_v;
}
+/* YA float -> RaGaBaA float */
+__kernel void yaf_to_ragabaf (__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.y,
+ in_v.x * in_v.y,
+ in_v.x * in_v.y,
+ in_v.y);
+
+ out[gid] = out_v;
+}
+
/* RGBA u8 -> YA float */
__kernel void rgbau8_to_yaf (__global const uchar4 * in,
__global float2 * out)
diff --git a/opencl/colors.cl.h b/opencl/colors.cl.h
index 1acbbc6..43dd1a2 100644
--- a/opencl/colors.cl.h
+++ b/opencl/colors.cl.h
@@ -67,6 +67,23 @@ static const char* colors_cl_source =
" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
"} \n"
" \n"
+" \n"
+"/* RGBA float -> RGB float */ \n"
+"__kernel void rgbaf_to_rgbf (__global const float4 * in, \n"
+" __global float * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" \n"
+"#if (__OPENCL_VERSION__ != CL_VERSION_1_0) \n"
+" vstore3 (in_v.xyz, gid, out); \n"
+"#else \n"
+" out[3 * gid] = in_v.x; \n"
+" out[3 * gid + 1] = in_v.y; \n"
+" out[3 * gid + 2] = in_v.z; \n"
+"#endif \n"
+"} \n"
+" \n"
"/* -- RaGaBaA float -- */ \n"
" \n"
"/* RGBA float -> RaGaBaA float */ \n"
@@ -355,6 +372,19 @@ static const char* colors_cl_source =
" out[gid] = out_v; \n"
"} \n"
" \n"
+"/* -- Y float -- */ \n"
+" \n"
+"/* Y float -> RaGaBaA float */ \n"
+"__kernel void yf_to_ragabaf (__global const float * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float y = in[gid]; \n"
+" float4 out_v = (float4) (y, y, y, 1.0f); \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
"/* -- YA float -- */ \n"
" \n"
"/* babl reference file: babl/base/rgb-constants.h */ \n"
@@ -391,6 +421,20 @@ static const char* colors_cl_source =
" out[gid] = out_v; \n"
"} \n"
" \n"
+"/* YA float -> RaGaBaA float */ \n"
+"__kernel void yaf_to_ragabaf (__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.y, \n"
+" in_v.x * in_v.y, \n"
+" in_v.x * in_v.y, \n"
+" 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"
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]