Hi Zhang,

I've pushed your commit, thanks for the contribution.

Besides the other patch where I had problems in my nvidia card, is there another patch you sent I haven't reviewed?

On Sun, Feb 9, 2014 at 5:58 PM, Yongjia Zhang wrote:

Add opencl implementation of operation texturize-canvas. Signed-off-by:Yongjia Zhang
---
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; idepth * 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@gnome.org List membership:
https://mail.gnome.org/mailman/listinfo/gegl-developer-list