[Gegl-developer] [PATCH] Add opencl implementation of operation texturize-canvas

[Yongjia Zhang <zhang_yong_jia 126 com>]

Add opencl implementation of operation texturize-canvas.
Signed-off-by:Yongjia Zhang<yongjia zhang intel com>
---
 opencl/texturize-canvas.cl           | 31 ++++++++++++
 opencl/texturize-canvas.cl.h         | 32 ++++++++++++
 operations/common/texturize-canvas.c | 97 ++++++++++++++++++++++++++++++++++++
 3 files changed, 160 insertions(+)
 create mode 100644 opencl/texturize-canvas.cl
 create mode 100644 opencl/texturize-canvas.cl.h

diff --git a/opencl/texturize-canvas.cl b/opencl/texturize-canvas.cl
new file mode 100644
index 0000000..36de13b
--- /dev/null
+++ b/opencl/texturize-canvas.cl
@@ -0,0 +1,31 @@
+#define CLAMP(val,lo,hi) (val < lo) ? lo : ((hi < val) ? hi : val )
+__kernel cl_texturize_canvas(__global const float * in,
+                             __global float * out,
+                             __global float * sdata,
+                             const int x,
+                             const int y,
+                             const int xm,
+                             const int ym,
+                             const int offs,
+                             const float mult,
+                             const int components,
+                             const int has_alpha)
+{
+    int col = get_global_id(0);
+    int row = get_global_id(1);
+    int step = components + has_alpha;
+    int index = step * (row * get_global_size(0) + col);
+    int canvas_index = ((x + col) & 127) * xm +
+                       ((y + row) & 127) * ym + offs;
+    float color;
+    int i;
+    float tmp = mult * sdata[canvas_index];
+    for(i=0; i<components; ++i)
+    {
+       color = tmp + src[index];
+       out[index++] = CLAMP(color,0.0f,1.0f);
+    }
+    if(has_alpha)
+       out[index] = in[index]; 
+}
+
diff --git a/opencl/texturize-canvas.cl.h b/opencl/texturize-canvas.cl.h
new file mode 100644
index 0000000..52ac243
--- /dev/null
+++ b/opencl/texturize-canvas.cl.h
@@ -0,0 +1,32 @@
+static const char *texturize_canvas_cl_source =
+"#define CLAMP(val,lo,hi) (val < lo) ? lo : ((hi < val) ? hi : val )        \n"
+"__kernel void cl_texturize_canvas(__global const float * in,               \n"
+"                                  __global float * out,                    \n"
+"                                  __global const float * sdata,            \n"
+"                                  const int x,                             \n"
+"                                  const int y,                             \n"
+"                                  const int xm,                            \n"
+"                                  const int ym,                            \n"
+"                                  const int offs,                          \n"
+"                                  const float mult,                        \n"
+"                                  const int components,                    \n"
+"                                  const int has_alpha)                     \n"
+"{                                                                          \n"
+"    int col = get_global_id(0);                                            \n"
+"    int row = get_global_id(1);                                            \n"
+"    int step = components + has_alpha;                                     \n"
+"    int index = step * (row * get_global_size(0) + col);                   \n"
+"    int canvas_index = ((x + col) & 127) * xm +                            \n"
+"                       ((y + row) & 127) * ym + offs;                      \n"
+"    float color;                                                           \n"
+"    int i;                                                                 \n"
+"    float tmp = mult * sdata[canvas_index];                                \n"
+"    for(i=0; i<components; ++i)                                            \n"
+"    {                                                                      \n"
+"       color = tmp + in[index];                                            \n"
+"       out[index++] = CLAMP(color,0.0f,1.0f);                              \n"
+"    }                                                                      \n"
+"    if(has_alpha)                                                          \n"
+"       out[index] = in[index];                                             \n"
+"}                                                                          \n"
+;
diff --git a/operations/common/texturize-canvas.c b/operations/common/texturize-canvas.c
index f552509..d475201 100644
--- a/operations/common/texturize-canvas.c
+++ b/operations/common/texturize-canvas.c
@@ -4183,6 +4183,101 @@ prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "output", new_format);
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/texturize-canvas.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 *opt = GEGL_CHANT_PROPERTIES(op);
+   float mult = (float)opt->depth * 0.25;
+   const Babl *format = gegl_operation_get_format(op, "input");
+   int has_alpha = babl_format_has_alpha(format);
+   int components = babl_format_get_n_components(format) - has_alpha;
+   int xm, ym, offs;
+   size_t global_ws[] = {roi->width, roi->height};
+   
+   switch(opt->direction)
+   {
+      default:
+      case GEGL_TEXTURIZE_CANVAS_DIRECTION_TOP_RIGHT:
+          xm = 1;
+          ym = 128;
+          offs = 0;
+          break;
+      case GEGL_TEXTURIZE_CANVAS_DIRECTION_TOP_LEFT:
+          xm = -1;
+          ym=128;
+          offs = 127;
+          break;
+      case GEGL_TEXTURIZE_CANVAS_DIRECTION_BOTTOM_LEFT:
+          xm = 128;
+          ym = 1;
+          offs = 0;
+          break;
+      case GEGL_TEXTURIZE_CANVAS_DIRECTION_BOTTOM_RIGHT:
+          xm = 128;
+          ym = -1;
+          offs = 127;
+          break;
+   }
+
+   if(!cl_data)
+   {
+      const char *kernel_name[] = {"cl_texturize_canvas", NULL};
+      cl_data = gegl_cl_compile_and_build(texturize_canvas_cl_source, kernel_name);
+   }
+   if(!cl_data)
+      return TRUE;
+   else
+   {
+      cl_int cl_err = 0;
+       
+      cl_mem sdata_tex = gegl_clCreateBuffer(gegl_cl_get_context(),
+                                             CL_MEM_USE_HOST_PTR|CL_MEM_READ_ONLY,
+                                             sizeof(cl_float)*128*128,(void *)sdata , &cl_err);
+      CL_CHECK;
+
+      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_mem), (void *)&sdata_tex,
+                                       sizeof(cl_int), (void *)&roi->x,
+                                       sizeof(cl_int), (void *)&roi->y,
+                                       sizeof(cl_int), (void *)&xm,
+                                       sizeof(cl_int), (void *)&ym,
+                                       sizeof(cl_int), (void *)&offs,
+                                       sizeof(cl_float), (void *)&mult,
+                                       sizeof(cl_int), (void *)&components,
+                                       sizeof(cl_int), (void *)&has_alpha, NULL);
+      CL_CHECK;
+
+      cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
+                                           cl_data->kernel[0], 2, NULL,
+                                           global_ws, NULL, 0, NULL, NULL);
+      CL_CHECK;
+   
+      cl_err = gegl_clFinish(gegl_cl_get_command_queue());
+      CL_CHECK;
+
+      cl_err = gegl_clReleaseMemObject(sdata_tex);
+      CL_CHECK_ONLY(cl_err);
+
+      return FALSE;
+
+error:
+      return TRUE;             
+   }
+}
+
+
 static gboolean
 process (GeglOperation       *operation,
          void                *in_buf,
@@ -4271,8 +4366,10 @@ 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;
 
+  operation_class->opencl_support = TRUE;
   gegl_operation_class_set_keys (operation_class,
     "name"       , "gegl:texturize-canvas",
     "categories" , "artistic",
-- 
1.8.3.2