commit fe6a4711f45cd2b4925def560ec271ea1d30e047
Author: Victor Oliveira <victormatheus gmail com>
Date: Mon Feb 13 19:20:43 2012 -0200
Changing color conversion to linear buffers
And organizing a little
gegl/opencl/gegl-cl-color-kernel.h | 186 +++++++++++++++++++++++++-----------
gegl/opencl/gegl-cl-color.c | 97 +++++++++++--------
2 files changed, 186 insertions(+), 97 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color-kernel.h b/gegl/opencl/gegl-cl-color-kernel.h
index 898be52..0b5b1f7 100644
--- a/gegl/opencl/gegl-cl-color-kernel.h
+++ b/gegl/opencl/gegl-cl-color-kernel.h
@@ -21,11 +21,34 @@ static const char* kernel_color_source =
" return native_powr ((value + 0.055f) / 1.055f, 2.4f); \n"
" return value / 12.92f; \n"
"} \n"
-" \n"
-" \n"
-"/* RGBA float -> RaGaBaA float */ \n"
-"__kernel void non_premultiplied_to_premultiplied (__global const float4 * in, \n"
-" __global float4 * out) \n"
+
+/* -- RGBA float/u8 -- */
+
+/* RGBA u8 -> RGBA float */
+"__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"
+
+/* RGBA float -> RGBA u8 */
+"__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"
+
+/* -- RaGaBaA float -- */
+
+/* RGBA float -> RaGaBaA float */
+"__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"
@@ -35,9 +58,10 @@ static const char* kernel_color_source =
" out[gid] = out_v; \n"
"} \n"
" \n"
-"/* RaGaBaA float -> RGBA float */ \n"
-"__kernel void premultiplied_to_non_premultiplied (__global const float4 * in, \n"
-" __global float4 * out) \n"
+
+/* RaGaBaA float -> RGBA float */
+"__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"
@@ -46,10 +70,37 @@ static const char* kernel_color_source =
" out_v.w = in_v.w; \n"
" out[gid] = out_v; \n"
"} \n"
+
+/* RGBA u8 -> RaGaBaA float */
+"__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"
-"/* RGBA float -> R'G'B'A float */ \n"
-"__kernel void rgba2rgba_gamma_2_2 (__global const float4 * in, \n"
-" __global float4 * out) \n"
+
+/* RaGaBaA float -> RGBA u8 */
+"__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(out_v); \n"
+"} \n"
+
+/* -- R'G'B'A float -- */
+
+/* rgba float -> r'g'b'a float */
+"__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"
@@ -60,10 +111,10 @@ static const char* kernel_color_source =
" in_v.w); \n"
" out[gid] = out_v; \n"
"} \n"
-" \n"
-"/* R'G'B'A float -> RGBA float */ \n"
-"__kernel void rgba_gamma_2_22rgba (__global const float4 * in, \n"
-" __global float4 * out) \n"
+
+/* r'g'b'a float -> rgba float */
+"__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"
@@ -74,63 +125,83 @@ static const char* kernel_color_source =
" in_v.w); \n"
" out[gid] = out_v; \n"
"} \n"
-" \n"
-"/* RGBA float -> R'aG'aB'aA float */ \n"
-"__kernel void rgba2rgba_gamma_2_2_premultiplied (__global const float4 * in, \n"
-" __global float4 * out) \n"
+
+/* rgba u8 -> r'g'b'a float */
+"__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 = in[gid]; \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) * in_v.w, \n"
-" linear_to_gamma_2_2(in_v.y) * in_v.w, \n"
-" linear_to_gamma_2_2(in_v.z) * in_v.w, \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'aG'aB'aA float -> RGBA float */ \n"
-"__kernel void rgba_gamma_2_2_premultiplied2rgba (__global const float4 * in, \n"
-" __global float4 * out) \n"
+
+/* r'g'b'a float -> rgba u8 */
+"__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 = (in_v.w > BABL_ALPHA_THRESHOLD)? (float4)(linear_to_gamma_2_2(in_v.x) / in_v.w,\n"
-" linear_to_gamma_2_2(in_v.y) / in_v.w,\n"
-" linear_to_gamma_2_2(in_v.z) / in_v.w,\n"
-" in_v.w) : \n"
-" (float4)(0.0f); \n"
-" out[gid] = 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"
-"/* RGBA float -> RGBA u8 */ \n"
-"__kernel void rgbaf_to_rgbau8 (__global const float4 * in, \n"
-" __global uchar4 * out) \n"
+
+/* -- Y'CbCrA float -- */
+
+/* RGBA float -> Y'CbCrA float */
+"__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 = in_v * 255.0f; \n"
-" out[gid] = convert_uchar4_sat_rte(out_v); \n"
-"} \n"
+" float4 out_v; \n"
" \n"
-"/* RGBAu8 -> 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]); \n"
-" float4 out_v = in_v / 255.0f; \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"
+
+/* Y'CbCrA float -> RGBA float */
+"__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"
-"/* RGBA float -> Y'CbCrA float */ \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"
-"__kernel void rgba_to_ycbcra (__global const float4 * in, \n"
-" __global float4 * out) \n"
+" out_v = (float4)(linear_to_gamma_2_2(rgb.x), \n"
+" linear_to_gamma_2_2(rgb.y), \n"
+" linear_to_gamma_2_2(rgb.z), \n"
+" in_v.w); \n"
+" out[gid] = out_v; \n"
+"} \n"
+
+/* RGBA u8 -> Y'CbCrA float */
+"__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 = in[gid]; \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"
@@ -144,11 +215,10 @@ static const char* kernel_color_source =
" in_v.w); \n"
" out[gid] = out_v; \n"
"} \n"
-" \n"
-"/* Y'CbCrA float -> RGBA float */ \n"
-" \n"
-"__kernel void ycbcra_to_rgba (__global const float4 * in, \n"
-" __global float4 * out) \n"
+
+/* Y'CbCrA float -> RGBA u8 */
+"__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"
@@ -163,5 +233,5 @@ static const char* kernel_color_source =
" linear_to_gamma_2_2(rgb.y), \n"
" linear_to_gamma_2_2(rgb.z), \n"
" in_v.w); \n"
-" out[gid] = out_v; \n"
-"} \n";
\ No newline at end of file
+" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
+"} \n";
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 40fc574..8b4aba4 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -6,31 +6,59 @@
static gegl_cl_run_data *kernels_color = NULL;
-#define CL_FORMAT_N 6
+#define CL_FORMAT_N 5
static const Babl *format[CL_FORMAT_N];
+enum
+{
+CL_RGBAU8_TO_RGBAF = 0,
+CL_RGBAF_TO_RGBAU8 = 1,
+
+CL_RGBAF_TO_RAGABAF = 2,
+CL_RAGABAF_TO_RGBAF = 3,
+CL_RGBAU8_TO_RAGABAF = 4,
+CL_RAGABAF_TO_RGBAU8 = 5,
+
+CL_RGBAF_TO_RGBA_GAMMA_F = 6,
+CL_RGBA_GAMMA_F_TO_RGBAF = 7,
+CL_RGBAU8_TO_RGBA_GAMMA_F = 8,
+CL_RGBA_GAMMA_F_TO_RGBAU8 = 9,
+
+CL_RGBAF_TO_YCBCRAF = 10,
+CL_YCBCRAF_TO_RGBAF = 11,
+CL_RGBAU8_TO_YCBCRAF = 12,
+CL_YCBCRAF_TO_RGBAU8 = 13,
+};
+
void
gegl_cl_color_compile_kernels(void)
{
- const char *kernel_name[] = {"non_premultiplied_to_premultiplied", /* 0 */
- "premultiplied_to_non_premultiplied", /* 1 */
- "rgba2rgba_gamma_2_2", /* 2 */
- "rgba_gamma_2_22rgba", /* 3 */
- "rgba2rgba_gamma_2_2_premultiplied", /* 4 */
- "rgba_gamma_2_2_premultiplied2rgba", /* 5 */
- "rgbaf_to_rgbau8", /* 6 */
- "rgbau8_to_rgbaf", /* 7 */
- "rgba_to_ycbcra", /* 8 */
- "ycbcra_to_rgba", /* 9 */
+ const char *kernel_name[] = {"rgbau8_to_rgbaf", /* 0 */
+ "rgbaf_to_rgbau8", /* 1 */
+
+ "rgbaf_to_ragabaf", /* 2 */
+ "ragabaf_to_rgbaf", /* 3 */
+ "rgbau8_to_ragabaf", /* 4 */
+ "ragabaf_to_rgbau8", /* 5 */
+
+ "rgbaf_to_rgba_gamma_f", /* 6 */
+ "rgba_gamma_f_to_rgbaf", /* 7 */
+ "rgbau8_to_rgba_gamma_f", /* 8 */
+ "rgba_gamma_f_to_rgbau8", /* 9 */
+
+ "rgbaf_to_ycbcraf", /* 10 */
+ "ycbcraf_to_rgbaf", /* 11 */
+ "rgbau8_to_ycbcraf", /* 12 */
+ "ycbcraf_to_rgbau8", /* 13 */
+
NULL};
format[0] = babl_format ("RGBA u8"),
format[1] = babl_format ("RGBA float"),
format[2] = babl_format ("RaGaBaA float"),
format[3] = babl_format ("R'G'B'A float"),
- format[4] = babl_format ("R'aG'aB'aA float"),
- format[5] = babl_format ("Y'CbCrA float"),
+ format[4] = babl_format ("Y'CbCrA float"),
kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
}
@@ -43,41 +71,32 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
if (in_format == babl_format ("RGBA float"))
{
- if (out_format == babl_format ("RaGaBaA float")) kernel = 0;
- else if (out_format == babl_format ("R'G'B'A float")) kernel = 2;
- else if (out_format == babl_format ("R'aG'aB'aA float")) kernel = 4;
- else if (out_format == babl_format ("RGBA u8")) kernel = 6;
- else if (out_format == babl_format ("Y'CbCrA float")) kernel = 8;
+ 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'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 (in_format == babl_format ("RaGaBaA float"))
+ else if (in_format == babl_format ("RGBA u8"))
{
- if (out_format == babl_format ("RGBA float")) kernel = 1;
- else if (out_format == babl_format ("RGBA u8")) kernel = 1;
+ if (out_format == babl_format ("RGBA float")) kernel = CL_RGBAU8_TO_RGBAF;
+ else if (out_format == babl_format ("RaGaBaA float")) kernel = CL_RGBAU8_TO_RAGABAF;
+ else if (out_format == babl_format ("R'G'B'A float")) kernel = CL_RGBAU8_TO_RGBA_GAMMA_F;
+ else if (out_format == babl_format ("Y'CbCrA float")) kernel = CL_RGBAU8_TO_YCBCRAF;
}
- else if (in_format == babl_format ("R'G'B'A float"))
- {
- if (out_format == babl_format ("RGBA float")) kernel = 3;
- else if (out_format == babl_format ("RGBA u8")) kernel = 3;
- }
- else if (in_format == babl_format ("R'aG'aB'aA float"))
+ else if (in_format == babl_format ("RaGaBaA float"))
{
- if (out_format == babl_format ("RGBA float")) kernel = 5;
- else if (out_format == babl_format ("RGBA u8")) kernel = 5;
+ 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 (in_format == babl_format ("RGBA u8")) /* read_imagef and write_imagef abstract texture format for us,
- * so I can use the same functions as RGBA float
- */
+ else if (in_format == babl_format ("R'G'B'A float"))
{
- if (out_format == babl_format ("RGBA float")) kernel = 7;
- else if (out_format == babl_format ("RaGaBaA float")) kernel = 0;
- else if (out_format == babl_format ("R'G'B'A float")) kernel = 2;
- else if (out_format == babl_format ("R'aG'aB'aA float")) kernel = 4;
- else if (out_format == babl_format ("Y'CbCrA float")) kernel = 8;
+ if (out_format == babl_format ("RGBA float")) kernel = CL_RGBA_GAMMA_F_TO_RGBAF;
+ else if (out_format == babl_format ("RGBA u8")) kernel = CL_RGBA_GAMMA_F_TO_RGBAU8;
}
else if (in_format == babl_format ("Y'CbCrA float"))
{
- if (out_format == babl_format ("RGBA float")) kernel = 9;
- else if (out_format == babl_format ("RGBA u8")) kernel = 9;
+ if (out_format == babl_format ("RGBA float")) kernel = CL_YCBCRAF_TO_RGBAF;
+ else if (out_format == babl_format ("RGBA u8")) kernel = CL_YCBCRAF_TO_RGBAU8;
}
return kernel;