Signed-off-by: Yongjia Zhang
---
opencl/motion-blur-zoom.cl | 78 ++++++++++++++++++++++++++ opencl/motion-blur-zoom.cl.h | 80 ++++++++++++++++++++++++++ operations/common/motion-blur-zoom.c | 105 +++++++++++++++++++++++++++++++++++ 3 files changed, 263 insertions(+)
create mode 100644 opencl/motion-blur-zoom.cl create mode 100644 opencl/motion-blur-zoom.cl.h

diff --git a/opencl/motion-blur-zoom.cl b/opencl/motion-blur-zoom.cl new file mode 100644
index 0000000..085c70c
--- /dev/null
+++ b/opencl/motion-blur-zoom.cl
@@ -0,0 +1,78 @@
+#define CLAMP(val,lo,hi) ((val) NOMINAL_NUM_IT) + xy_len = min(NOMINAL_NUM_IT + (int)sqrt((float)(xy_len - NOMINAL_NUM_IT)), + MAX_NUM_IT);
+ inv_xy_len = 1.0f / xy_len;
+ dxx = (x_end - x_start) * inv_xy_len; + dyy = (y_end - y_start) * inv_xy_len; + ix = x_start;
+ iy = y_start;
+
+ for( int i=0; iwidth,roi->height}; + if(!cl_data)
+ {
+ const char *kernel_name[] = {"motion_blur_zoom", NULL}; + cl_data = gegl_cl_compile_and_build(motion_blur_zoom_cl_source, kernel_name); + }
+ if(!cl_data)
+ return TRUE;
+
+ 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_int), (void *)&src_rect->width, + sizeof(cl_int), (void *)&src_rect->height, + sizeof(cl_int), (void *)&src_rect->x, + sizeof(cl_int), (void *)&src_rect->y, + sizeof(cl_int), (void *)&roi->x, + sizeof(cl_int), (void *)&roi->y, + sizeof(cl_float), (void *)¢er_x, + sizeof(cl_float), (void *)¢er_y, + sizeof(cl_float), (void *)&factor, 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;
+
+ return FALSE;
+
+error:
+ return TRUE;
+}
+
+
+static gboolean
+cl_process(GeglOperation *operation, + GeglBuffer *input,
+ GeglBuffer *output,
+ const GeglRectangle *result, + const GeglRectangle *src_rect) +{
+ GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER(operation); + GeglChantO *o = GEGL_CHANT_PROPERTIES(operation); +
+ const Babl *in_format = gegl_operation_get_format(operation,"input"); + const Babl *out_format = gegl_operation_get_format(operation, "output"); +
+ gint err;
+
+ GeglBufferClIterator *i = gegl_buffer_cl_iterator_new(output, + result, + out_format, + GEGL_CL_BUFFER_WRITE); + gint read = gegl_buffer_cl_iterator_add_2(i, + input, + result, + in_format, + GEGL_CL_BUFFER_READ, + op_area->left, + op_area->right, + op_area->top, + op_area->bottom, + GEGL_ABYSS_NONE); + while(gegl_buffer_cl_iterator_next(i,&err)) + {
+ if(err) return FALSE;
+ err = cl_motion_blur_zoom(i->tex[read], + i->tex[0], + &i->roi[0], + &i->roi[read], + o->center_x, + o->center_y, + o->factor); +
+ if(err) return FALSE;
+ }
+
+ return TRUE;
+
+
+}
+
+
static inline gfloat *
get_pixel_color (gfloat *in_buf, const GeglRectangle *rect, @@ -130,6 +231,9 @@ process (GeglOperation *operation, src_rect.width += op_area->left + op_area->right; src_rect.height += op_area->top + op_area->bottom;
+ if(gegl_operation_use_opencl(operation)) + if(cl_process(operation, input, output, roi, &src_rect)) + return TRUE;
in_buf = g_new (gfloat, src_rect.width * src_rect.height * 4); out_buf = g_new0 (gfloat, roi->width * roi->height * 4); out_pixel = out_buf;
@@ -218,6 +322,7 @@ gegl_chant_class_init (GeglChantClass *klass) filter_class = GEGL_OPERATION_FILTER_CLASS (klass);
operation_class->prepare = prepare; + operation_class->opencl_support = TRUE; filter_class->process = process;
gegl_operation_class_set_keys (operation_class,

1.8.3.2