From 1f74bccd4adca68d95965c199d5a0a7f49334042 Mon Sep 17 00:00:00 2001 From: Carlos Zubieta Date: Tue, 23 Jul 2013 18:08:06 -0500 Subject: Added OpenCL support to color-to-alpha --- opencl/color-to-alpha.cl | 63 ++++++++++++++++++++++++++++++++++++ opencl/color-to-alpha.cl.h | 65 ++++++++++++++++++++++++++++++++++++++ operations/common/color-to-alpha.c | 55 ++++++++++++++++++++++++++++++++ 3 files changed, 183 insertions(+) create mode 100644 opencl/color-to-alpha.cl create mode 100644 opencl/color-to-alpha.cl.h diff --git a/opencl/color-to-alpha.cl b/opencl/color-to-alpha.cl new file mode 100644 index 00000000..9ca08949 --- /dev/null +++ b/opencl/color-to-alpha.cl @@ -0,0 +1,63 @@ +__kernel void cl_color_to_alpha(__global const float4 *in, + __global float4 *out, + float4 color) +{ + int gid = get_global_id(0); + float4 in_v = in[gid]; + float4 out_v = in_v; + float4 alpha; + + alpha.w = in_v.w; + + /*First component*/ + if ( color.x < 0.00001f ) + alpha.x = in_v.x; + else if ( in_v.x > color.x + 0.00001f ) + alpha.x = (in_v.x - color.x) / (1.0f - color.x); + else if ( in_v.x < color.x - 0.00001f ) + alpha.x = (color.x - in_v.x) / color.x; + else + alpha.x = 0.0f; + /*Second component*/ + if ( color.y < 0.00001f ) + alpha.y = in_v.y; + else if ( in_v.y > color.y + 0.00001f ) + alpha.y = (in_v.y - color.y) / (1.0f - color.y); + else if ( in_v.y < color.y - 0.00001f ) + alpha.y = (color.y - in_v.y) / color.y; + else + alpha.y = 0.0f; + /*Third component*/ + if ( color.z < 0.00001f ) + alpha.z = in_v.z; + else if ( in_v.z > color.z + 0.00001f ) + alpha.z = (in_v.z - color.z) / (1.0f - color.z); + else if ( in_v.z < color.z - 0.00001f ) + alpha.z = (color.z - in_v.z) / color.z; + else + alpha.z = 0.0f; + + if (alpha.x > alpha.y) + { + if (alpha.x > alpha.z) + out_v.w = alpha.x; + else + out_v.w = alpha.z; + } + else if (alpha.y > alpha.z) + { + out_v.w = alpha.y; + } + else + { + out_v.w = alpha.z; + } + + if (out_v.w >= 0.00001f) + { + out_v.xyz = (out_v.xyz - color.xyz) / out_v.www + color.xyz; + out_v.w *= alpha.w; + } + + out[gid] = out_v; +} diff --git a/opencl/color-to-alpha.cl.h b/opencl/color-to-alpha.cl.h new file mode 100644 index 00000000..409e0bc6 --- /dev/null +++ b/opencl/color-to-alpha.cl.h @@ -0,0 +1,65 @@ +static const char* color_to_alpha_cl_source = +"__kernel void cl_color_to_alpha(__global const float4 *in, \n" +" __global float4 *out, \n" +" float4 color) \n" +"{ \n" +" int gid = get_global_id(0); \n" +" float4 in_v = in[gid]; \n" +" float4 out_v = in_v; \n" +" float4 alpha; \n" +" \n" +" alpha.w = in_v.w; \n" +" \n" +" /*First component*/ \n" +" if ( color.x < 0.00001f ) \n" +" alpha.x = in_v.x; \n" +" else if ( in_v.x > color.x + 0.00001f ) \n" +" alpha.x = (in_v.x - color.x) / (1.0f - color.x); \n" +" else if ( in_v.x < color.x - 0.00001f ) \n" +" alpha.x = (color.x - in_v.x) / color.x; \n" +" else \n" +" alpha.x = 0.0f; \n" +" /*Second component*/ \n" +" if ( color.y < 0.00001f ) \n" +" alpha.y = in_v.y; \n" +" else if ( in_v.y > color.y + 0.00001f ) \n" +" alpha.y = (in_v.y - color.y) / (1.0f - color.y); \n" +" else if ( in_v.y < color.y - 0.00001f ) \n" +" alpha.y = (color.y - in_v.y) / color.y; \n" +" else \n" +" alpha.y = 0.0f; \n" +" /*Third component*/ \n" +" if ( color.z < 0.00001f ) \n" +" alpha.z = in_v.z; \n" +" else if ( in_v.z > color.z + 0.00001f ) \n" +" alpha.z = (in_v.z - color.z) / (1.0f - color.z); \n" +" else if ( in_v.z < color.z - 0.00001f ) \n" +" alpha.z = (color.z - in_v.z) / color.z; \n" +" else \n" +" alpha.z = 0.0f; \n" +" \n" +" if (alpha.x > alpha.y) \n" +" { \n" +" if (alpha.x > alpha.z) \n" +" out_v.w = alpha.x; \n" +" else \n" +" out_v.w = alpha.z; \n" +" } \n" +" else if (alpha.y > alpha.z) \n" +" { \n" +" out_v.w = alpha.y; \n" +" } \n" +" else \n" +" { \n" +" out_v.w = alpha.z; \n" +" } \n" +" \n" +" if (out_v.w >= 0.00001f) \n" +" { \n" +" out_v.xyz = (out_v.xyz - color.xyz) / out_v.www + color.xyz; \n" +" out_v.w *= alpha.w; \n" +" } \n" +" \n" +" out[gid] = out_v; \n" +"} \n" +; diff --git a/operations/common/color-to-alpha.c b/operations/common/color-to-alpha.c index 2d61c242..a8a5904a 100644 --- a/operations/common/color-to-alpha.c +++ b/operations/common/color-to-alpha.c @@ -130,6 +130,58 @@ color_to_alpha (const gfloat *color, dst[3] *= alpha[3]; } +#include "opencl/gegl-cl.h" +#include "opencl/color-to-alpha.cl.h" + +static GeglClRunData * cl_data = NULL; + +static gboolean +cl_process (GeglOperation *operation, + cl_mem in, + cl_mem out, + size_t global_worksize, + const GeglRectangle *roi, + gint level) +{ + GeglChantO *o = GEGL_CHANT_PROPERTIES (operation); + gfloat color[4]; + gegl_color_get_pixel (o->color, babl_format ("R'G'B'A float"), color); + + if (!cl_data) + { + const char *kernel_name[] = {"cl_color_to_alpha",NULL}; + cl_data = gegl_cl_compile_and_build (color_to_alpha_cl_source, kernel_name); + } + if (!cl_data) + return TRUE; + else + { + cl_int cl_err = 0; + cl_float4 f_color; + f_color.s[0] = color[0]; + f_color.s[1] = color[1]; + f_color.s[2] = color[2]; + f_color.s[3] = color[3]; + + cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in); + CL_CHECK; + cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out); + CL_CHECK; + cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float4),(void*)&f_color); + CL_CHECK; + + cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), + cl_data->kernel[0], 1, + NULL, &global_worksize, NULL, + 0, NULL, NULL); + CL_CHECK; + } + + return FALSE; + + error: + return TRUE; +} static gboolean @@ -192,7 +244,10 @@ gegl_chant_class_init (GeglChantClass *klass) filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass); filter_class->process = process; + filter_class->cl_process = cl_process; + operation_class->prepare = prepare; + operation_class->opencl_support = TRUE; gegl_operation_class_set_keys (operation_class, "name", "gegl:color-to-alpha", -- cgit v1.2.3