Cool!
Because double-precision is not supported (or very slow) in most GPUs,
we try to not use it in our OpenCL code. Maybe this is something we
should be more careful though.
Your point that gegl_cl_compile_and_build should have compiler flags
is pretty good and I'll work on that.
About your code, no complains except that you could use clamp instead
of fmin followed by fmax, but I can do that though.
Victor Oliveira
On Wed, Jan 22, 2014 at 12:26 AM, Yongjia Zhang <zhang_yong_jia@xxxxxxx> wrote:
> From: Yongjia Zhang <Zhang_Yong_jia@xxxxxxx>
>
> 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@xxxxxxxxx>
> ---
> 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@xxxxxxxxx
> List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list
>
_______________________________________________
gegl-developer-list mailing list
List address: gegl-developer-list@xxxxxxxxx
List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list
