[gegl] cl: support for R'G'B' u8 and R'G'B'A u8
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] cl: support for R'G'B' u8 and R'G'B'A u8
- Date: Fri, 30 Mar 2012 20:14:08 +0000 (UTC)
commit 4ac1e6a086348865beb295a425bbcd659d1719dc
Author: Victor Oliveira <victormatheus gmail com>
Date: Fri Mar 30 16:16:51 2012 -0300
cl: support for R'G'B' u8 and R'G'B'A u8
gegl/opencl/gegl-cl-color-kernel.h | 65 ++++++++++++++++++++++++++++++++++++
gegl/opencl/gegl-cl-color.c | 30 ++++++++++++++++-
2 files changed, 94 insertions(+), 1 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color-kernel.h b/gegl/opencl/gegl-cl-color-kernel.h
index 34c2522..6159080 100644
--- a/gegl/opencl/gegl-cl-color-kernel.h
+++ b/gegl/opencl/gegl-cl-color-kernel.h
@@ -337,4 +337,69 @@ static const char* kernel_color_source =
" 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"
+
+/* R'G'B'A u8 */
+
+/* rgba float -> r'g'b'a u8 */
+"__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"
+
+/* r'g'b'a u8 -> rgba float */
+"__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"
+
+/* R'G'B' u8 */
+
+/* rgba float -> r'g'b' u8 */
+"__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"
+" uchar3 out_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"
+" out_v = convert_uchar3_sat_rte(255.0f * tmp_v.w * tmp_v.xyz); \n"
+" vstore3 (out_v, gid, out); \n"
+"} \n"
+
+/* r'g'b' u8 -> rgba float */
+"__kernel void rgb_gamma_u8_to_rgbaf (__global const uchar * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" uchar3 in_v; \n"
+" float3 tmp_v; \n"
+" float4 out_v; \n"
+" in_v = vload3 (gid, in); \n"
+" tmp_v = convert_float3(in_v) / 255.0f; \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";
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 20b91db..fc0fa7f 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -8,7 +8,7 @@
static gegl_cl_run_data *kernels_color = NULL;
-#define CL_FORMAT_N 9
+#define CL_FORMAT_N 11
static const Babl *format[CL_FORMAT_N];
@@ -41,6 +41,12 @@ CL_RGBAF_TO_YAF = 17,
CL_YAF_TO_RGBAF = 18,
CL_RGBAU8_TO_YAF = 19,
CL_YAF_TO_RGBAU8 = 20,
+
+CL_RGBAF_TO_RGBA_GAMMA_U8 = 21,
+CL_RGBA_GAMMA_U8_TO_RGBAF = 22,
+
+CL_RGBAF_TO_RGB_GAMMA_U8 = 23,
+CL_RGB_GAMMA_U8_TO_RGBAF = 24,
};
void
@@ -74,6 +80,12 @@ gegl_cl_color_compile_kernels(void)
"rgbau8_to_yaf", /* 19 */
"yaf_to_rgbau8", /* 20 */
+ "rgbaf_to_rgba_gamma_u8", /* 21 */
+ "rgba_gamma_u8_to_rgbaf", /* 22 */
+
+ "rgbaf_to_rgb_gamma_u8", /* 23 */
+ "rgb_gamma_u8_to_rgbaf", /* 24 */
+
NULL};
format[0] = babl_format ("RGBA u8");
@@ -85,6 +97,8 @@ gegl_cl_color_compile_kernels(void)
format[6] = babl_format ("Y u8");
format[7] = babl_format ("Y float");
format[8] = babl_format ("YA float");
+ format[9] = babl_format ("R'G'B'A u8");
+ format[10] = babl_format ("R'G'B' u8");
kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
}
@@ -99,6 +113,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'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;
@@ -140,6 +156,14 @@ 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 (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 (in_format == babl_format ("R'G'B' u8"))
+ {
+ if (out_format == babl_format ("RGBA float")) kernel = CL_RGB_GAMMA_U8_TO_RGBAF;
+ }
return kernel;
}
@@ -171,6 +195,10 @@ gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes)
*bytes = sizeof (cl_float);
else if (buffer_format == babl_format ("YA float"))
*bytes = sizeof (cl_float2);
+ else if (buffer_format == babl_format("R'G'B'A u8"))
+ *bytes = sizeof (cl_uchar4);
+ else if (buffer_format == babl_format("R'G'B' u8"))
+ *bytes = sizeof (cl_uchar3);
else
*bytes = sizeof (cl_float4);
}
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
