Add opencl implementation of operation texturize-canvas.
Signed-off-by:Yongjia Zhang<yongjia.zhang@xxxxxxxxx>
---
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
_______________________________________________
gegl-developer-list mailing list
List address: gegl-developer-list@xxxxxxxxx
List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list
