[Gegl-developer] [PATCH] Add opencl implementation of operation channel-mixer



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




[Date Prev][Date Next]   [Thread Prev][Thread Next]   [Thread Index] [Date Index] [Author Index]