From: Yongjia Zhang

Although function gegl_chant_class_init had set operation_class->opencl_support=yes in source file channel-mixer.c, it didn't have a opencl implementation of operation channel-mixer.
In the cpu version of channel-mixer, all needed information is calculated using variable type double. Since gegl_cl_compile_and_build does not take opencl kernel build options, this implementation uses type float instead of double. If build options could be taken, query the device capability of cl_khr_fp64 extension and then select the proper variable type would be the best.

Signed-off-by: Yongjia Zhang ---
opencl/channel-mixer.cl | 40 ++++++++++++++++++++ opencl/channel-mixer.cl.h | 41 ++++++++++++++++++++ operations/common/channel-mixer.c | 80 +++++++++++++++++++++++++++++++++++++++ 3 files changed, 161 insertions(+)
create mode 100644 opencl/channel-mixer.cl create mode 100644 opencl/channel-mixer.cl.h

diff --git a/opencl/channel-mixer.cl b/opencl/channel-mixer.cl new file mode 100644
index 0000000..2fb899c
--- /dev/null
+++ b/opencl/channel-mixer.cl
@@ -0,0 +1,40 @@
+#define CM_MIX_PIXEL( ch, r, g, b, norm ) \ + c = ch.x * r + ch.y * g + ch.z * b; \ + c *= norm; \ + mix_return = fmin( 1.0f, fmax( 0.0f, c ) ); + +__kernel void cl_channel_mixer(__global const float *in, + __global float *out, + float4 ch_red, + float4 ch_green, + float4 ch_blue, + float4 ch_black, + float red_norm, + float green_norm, + float blue_norm, + float black_norm, + int monochrome, + int has_alpha) +{ + const int step = (has_alpha == 0 ? 3 : 4 ); + const int offset = get_global_id(0) * step; + float mix_return = 0.0f; + float c = 0.0f; + if( monochrome ) + { + CM_MIX_PIXEL( ch_black, in[offset], in[offset+1], in[offset+2], black_norm ); + out[offset] = out[offset+1] = out[offset+2] = mix_return; + } + else + { + CM_MIX_PIXEL( ch_red, in[offset], in[offset+1], in[offset+2], red_norm ); + out[offset] = mix_return; + CM_MIX_PIXEL( ch_green, in[offset], in[offset+1], in[offset+2], green_norm ); + out[offset+1] = mix_return; + CM_MIX_PIXEL( ch_blue, in[offset], in[offset+1], in[offset+2], blue_norm ); + out[offset+2] = mix_return; + } + if( 4==step ) + out[offset+3] = in[offset+3]; +} +
diff --git a/opencl/channel-mixer.cl.h b/opencl/channel-mixer.cl.h new file mode 100644
index 0000000..32e12d6
--- /dev/null
+++ b/opencl/channel-mixer.cl.h
@@ -0,0 +1,41 @@
+static const char* channel_mixer_cl_source = +"#define CM_MIX_PIXEL( ch, r, g, b, norm ) \\\n" +" c = ch.x * r + ch.y * g + ch.z * b; \\\n" +" c *= norm; \\\n" +" mix_return = fmin( 1.0f, fmax( 0.0f, c ) ); \n" +" \n" +"__kernel void cl_channel_mixer(__global const float *in, \n" +" __global float *out, \n" +" float4 ch_red, \n" +" float4 ch_green, \n" +" float4 ch_blue, \n" +" float4 ch_black, \n" +" float red_norm, \n" +" float green_norm, \n" +" float blue_norm, \n" +" float black_norm, \n" +" int monochrome, \n" +" int has_alpha) \n" +"{ \n" +" const int step = (has_alpha == 0 ? 3 : 4 ); \n" +" const int offset = get_global_id(0) * step; \n" +" float mix_return = 0.0f; \n" +" float c = 0.0f; \n" +" if( monochrome ) \n" +" { \n" +" CM_MIX_PIXEL( ch_black, in[offset], in[offset+1], in[offset+2], black_norm ); \n" +" out[offset] = out[offset+1] = out[offset+2] = mix_return; \n" +" } \n" +" else \n" +" { \n" +" CM_MIX_PIXEL( ch_red, in[offset], in[offset+1], in[offset+2], red_norm ); \n" +" out[offset] = mix_return; \n" +" CM_MIX_PIXEL( ch_green, in[offset], in[offset+1], in[offset+2], green_norm ); \n" +" out[offset+1] = mix_return; \n" +" CM_MIX_PIXEL( ch_blue, in[offset], in[offset+1], in[offset+2], blue_norm ); \n" +" out[offset+2] = mix_return; \n" +" } \n" +" if( 4==step ) \n" +" out[offset+3] = in[offset+3]; \n" +"} \n" +;
diff --git a/operations/common/channel-mixer.c b/operations/common/channel-mixer.c index cfa7858..ac317de 100644
--- a/operations/common/channel-mixer.c +++ b/operations/common/channel-mixer.c @@ -275,6 +275,85 @@ process (GeglOperation *op, return TRUE;
}

+#include "opencl/gegl-cl.h"
+#include "opencl/channel-mixer.cl.h" +
+static GeglClRunData *cl_data = NULL; +
+static gboolean
+cl_process(GeglOperation *op,
+ cl_mem in_tex,
+ cl_mem out_tex,
+ glong samples,
+ const GeglRectangle *roi, + gint level)
+{
+ GeglChantO *o = GEGL_CHANT_PROPERTIES(op); + CmParamsType *mix = (CmParamsType *)o->chant_data; + float red_norm, green_norm, blue_norm, black_norm; + cl_float4 ch_red, ch_green, ch_blue, ch_black; + int global_ws[]={samples};
+
+ g_assert( mix != NULL );
+
+ red_norm = (float)cm_calculate_norm(mix, &mix->red); + green_norm = (float)cm_calculate_norm(mix, &mix->green); + blue_norm = (float)cm_calculate_norm(mix, &mix->blue); + black_norm = (float)cm_calculate_norm(mix, &mix->black); +
+ /*Convert double to float*/
+ ch_red.s[0] = (float)mix->red.red_gain; + ch_red.s[1] = (float)mix->red.green_gain; + ch_red.s[2] = (float)mix->red.blue_gain; + ch_green.s[0] = (float)mix->green.red_gain; + ch_green.s[1] = (float)mix->green.green_gain; + ch_green.s[2] = (float)mix->green.blue_gain; + ch_blue.s[0] = (float)mix->blue.red_gain; + ch_blue.s[1] = (float)mix->blue.green_gain; + ch_blue.s[2] = (float)mix->blue.blue_gain; + ch_black.s[0] = (float)mix->black.red_gain; + ch_black.s[1] = (float)mix->black.green_gain; + ch_black.s[2] = (float)mix->black.blue_gain; +
+ if( !cl_data )
+ {
+ const char *kernel_name[] = {"cl_channel_mixer", NULL}; + cl_data = gegl_cl_compile_and_build( channel_mixer_cl_source, kernel_name ); + }
+ if( !cl_data )
+ return TRUE;
+ else
+ {
+ cl_int cl_err = 0;
+
+ cl_err = gegl_cl_set_kernel_args(cl_data->kernel[0], + sizeof(cl_mem),(void*)&in_tex, + sizeof(cl_mem),(void*)&out_tex, + sizeof(cl_float4),(void*)&ch_red, + sizeof(cl_float4),(void*)&ch_green, + sizeof(cl_float4),(void*)&ch_blue, + sizeof(cl_float4),(void*)&ch_black, + sizeof(cl_float),(void*)&red_norm, + sizeof(cl_float),(void*)&green_norm, + sizeof(cl_float),(void*)&blue_norm, + sizeof(cl_float),(void*)&black_norm, + sizeof(cl_int),(void*)&mix->monochrome, + sizeof(cl_int),(void*)&mix->has_alpha, NULL); + CL_CHECK;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), + cl_data->kernel[0], 1, + NULL, global_ws, NULL, + 0, NULL, NULL); + CL_CHECK;
+
+ return FALSE;
+
+error:
+ return TRUE;
+ }
+
+}
static void
gegl_chant_class_init (GeglChantClass *klass) {
@@ -285,6 +364,7 @@ gegl_chant_class_init (GeglChantClass *klass) point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
point_filter_class->process = process; + point_filter_class->cl_process = cl_process; operation_class->prepare = prepare; G_OBJECT_CLASS (klass)->finalize = finalize;

1.8.3.2