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