summaryrefslogtreecommitdiff
path: root/opencl
diff options
context:
space:
mode:
Diffstat (limited to 'opencl')
-rw-r--r--opencl/colors.cl43
-rw-r--r--opencl/colors.cl.h43
2 files changed, 86 insertions, 0 deletions
diff --git a/opencl/colors.cl b/opencl/colors.cl
index 9448a6d5..2f74d66d 100644
--- a/opencl/colors.cl
+++ b/opencl/colors.cl
@@ -176,6 +176,34 @@ __kernel void ragabaf_to_rgb_gamma_u8 (__global const float4 * in,
#endif
}
+/* -- R'G'B' float -- */
+
+/* R'G'B' float -> RGBA float */
+__kernel void rgb_gamma_f_to_rgbaf (__global const float * in,
+ __global float4 * out)
+{
+ int gid = get_global_id(0);
+ float4 out_v;
+
+#if (__OPENCL_VERSION__ != CL_VERSION_1_0)
+ float3 in_v;
+ in_v = vload3 (gid, in);
+ out_v = (float4)(gamma_2_2_to_linear(in_v.x),
+ gamma_2_2_to_linear(in_v.y),
+ gamma_2_2_to_linear(in_v.z),
+ 1.0f);
+#else
+ float r = in[3 * gid + 0];
+ float g = in[3 * gid + 1];
+ float b = in[3 * gid + 2];
+ out_v = (float4)(gamma_2_2_to_linear(r),
+ gamma_2_2_to_linear(g),
+ gamma_2_2_to_linear(b),
+ 1.0f);
+#endif
+ out[gid] = out_v;
+}
+
/* -- R'G'B'A float -- */
/* rgba float -> r'g'b'a float */
@@ -409,6 +437,21 @@ __kernel void rgbaf_to_yaf (__global const float4 * in,
out[gid] = out_v;
}
+/* RaGaBaA float -> YA float */
+__kernel void ragabaf_to_yaf (__global const float4 * in,
+ __global float2 * out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 tmp_v = in_v / (float4) (in_v.w);
+
+ float luminance = tmp_v.x * RGB_LUMINANCE_RED +
+ tmp_v.y * RGB_LUMINANCE_GREEN +
+ tmp_v.z * RGB_LUMINANCE_BLUE;
+
+ out[gid] = (float2) (luminance, in_v.w);
+}
+
/* YA float -> RGBA float */
__kernel void yaf_to_rgbaf (__global const float2 * in,
__global float4 * out)
diff --git a/opencl/colors.cl.h b/opencl/colors.cl.h
index 43dd1a2f..b5bc65ac 100644
--- a/opencl/colors.cl.h
+++ b/opencl/colors.cl.h
@@ -177,6 +177,34 @@ static const char* colors_cl_source =
"#endif \n"
"} \n"
" \n"
+"/* -- R'G'B' float -- */ \n"
+" \n"
+"/* R'G'B' float -> RGBA float */ \n"
+"__kernel void rgb_gamma_f_to_rgbaf (__global const float * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 out_v; \n"
+" \n"
+"#if (__OPENCL_VERSION__ != CL_VERSION_1_0) \n"
+" float3 in_v; \n"
+" in_v = vload3 (gid, in); \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"
+" 1.0f); \n"
+"#else \n"
+" float r = in[3 * gid + 0]; \n"
+" float g = in[3 * gid + 1]; \n"
+" float b = in[3 * gid + 2]; \n"
+" out_v = (float4)(gamma_2_2_to_linear(r), \n"
+" gamma_2_2_to_linear(g), \n"
+" gamma_2_2_to_linear(b), \n"
+" 1.0f); \n"
+"#endif \n"
+" out[gid] = out_v; \n"
+"} \n"
+" \n"
"/* -- R'G'B'A float -- */ \n"
" \n"
"/* rgba float -> r'g'b'a float */ \n"
@@ -410,6 +438,21 @@ static const char* colors_cl_source =
" out[gid] = out_v; \n"
"} \n"
" \n"
+"/* RaGaBaA float -> YA float */ \n"
+"__kernel void ragabaf_to_yaf (__global const float4 * in, \n"
+" __global float2 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 tmp_v = in_v / (float4) (in_v.w); \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[gid] = (float2) (luminance, in_v.w); \n"
+"} \n"
+" \n"
"/* YA float -> RGBA float */ \n"
"__kernel void yaf_to_rgbaf (__global const float2 * in, \n"
" __global float4 * out) \n"