commit 2b003e23b68a4f9ec6797e7def54712cc01eebe7
Author: Massimo Valentini <mvalentini src gnome org>
Date: Tue Nov 20 18:44:56 2012 +0100
opencl: make color conversions work with OpenCL1.0
using the runtime preprocessor
gegl/opencl/gegl-cl-color-kernel.h | 58 ++++++++++++++++++++++++++++++-----
1 files changed, 49 insertions(+), 9 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color-kernel.h b/gegl/opencl/gegl-cl-color-kernel.h
index e32d884..e485ecd 100644
--- a/gegl/opencl/gegl-cl-color-kernel.h
+++ b/gegl/opencl/gegl-cl-color-kernel.h
@@ -159,7 +159,12 @@ static const char* kernel_color_source =
" __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"
@@ -186,7 +191,14 @@ static const char* kernel_color_source =
" 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"
/* -- R'G'B'A float -- */
@@ -335,11 +347,16 @@ static const char* kernel_color_source =
" __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"
@@ -349,8 +366,15 @@ static const char* kernel_color_source =
"{ \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"
/* -- Y u8 -- */
@@ -490,7 +514,12 @@ static const char* kernel_color_source =
" __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"
@@ -514,14 +543,13 @@ static const char* kernel_color_source =
"{ \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"
+" uchar tmp = convert_uchar_sat_rte (255.0f * linear_to_gamma_2_2 (in_v.x)); \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), gid, out); \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"
@@ -564,13 +592,20 @@ static const char* kernel_color_source =
" 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"
+"#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"
/* r'g'b' u8 -> rgba float */
@@ -578,11 +613,16 @@ static const char* kernel_color_source =
" __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"
-" float4 out_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"