From: Yongjia Zhang <Zhang_Yong_jia 126 com>
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<yongjia zhang intel com>
---
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
_______________________________________________
gegl-developer-list mailing list
List address: gegl-developer-list gnome org
List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list
