[gegl/soc-2013-opecl-ops] Operations: Add OpenCL support to dot
- From: Carlos Zubieta <czubieta src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl/soc-2013-opecl-ops] Operations: Add OpenCL support to dot
- Date: Thu, 19 Sep 2013 05:03:14 +0000 (UTC)
commit 4ff7bcab9fa87c4e6b54ee52291b756bb7cee758
Author: Carlos Zubieta <czubieta dev gmail com>
Date: Wed Sep 18 23:46:37 2013 -0500
Operations: Add OpenCL support to dot
opencl/dot.cl | 72 ++++++++++++++++++++++
opencl/dot.cl.h | 73 +++++++++++++++++++++++
operations/common/dot.c | 151 ++++++++++++++++++++++++++++++++++++++++++++++-
3 files changed, 294 insertions(+), 2 deletions(-)
---
diff --git a/opencl/dot.cl b/opencl/dot.cl
new file mode 100644
index 0000000..4c504df
--- /dev/null
+++ b/opencl/dot.cl
@@ -0,0 +1,72 @@
+/* This file is an image processing operation for GEGL
+ *
+ * GEGL is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * GEGL is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.
+ *
+ * Copyright 2013 Carlos Zubieta <czubieta dev gmail com>
+ */
+
+__kernel void cl_calc_block_colors (__global const float4 *in,
+ __global float4 *blocks_colors,
+ int cx0,
+ int cy0,
+ int size,
+ float weight,
+ int roi_x,
+ int roi_y,
+ int line_width)
+{
+ int cx = cx0 + get_global_id(0);
+ int cy = cy0 + get_global_id(1);
+ int px = (cx * size) - roi_x + size;
+ int py = (cy * size) - roi_y + size;
+ float4 mean = (float4)(0.0f);
+ float4 tmp;
+ int i, j;
+
+ for( j = py; j < py+size; ++j)
+ {
+ for (i = px; i < px+size; ++i)
+ {
+ mean += in[j * line_width + i];
+ }
+ }
+ block_colors[(cx-cx0) + get_global_size(0) * (cy-cy0)] = mean*weight;
+}
+
+__kernel void cl_dot (__global const float4 *block_colors,
+ __global float4 *out,
+ int cx0,
+ int cy0,
+ int size,
+ float radius2,
+ int roi_x,
+ int roi_y,
+ int block_count_x)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+ int x = gidx + roi_x;
+ int y = gidy + roi_y;
+ int cy = y/size;
+ int cx = x/size;
+ float cellx = convert_float(x - cx * size) - convert_float(size) / 2.0;
+ float celly = convert_float(y - cy * size) - convert_float(size) / 2.0;
+ float4 tmp = (float4)(0.0);
+
+ if((cellx * cellx + celly * celly) <= radius2)
+ tmp = block_colors[(cx-cx0) + block_count_x * (cy-cy0)];
+
+ out[gidx + get_global_size(0) * gidy] = tmp;
+}
+
diff --git a/opencl/dot.cl.h b/opencl/dot.cl.h
new file mode 100644
index 0000000..91777dd
--- /dev/null
+++ b/opencl/dot.cl.h
@@ -0,0 +1,73 @@
+/* This file is an image processing operation for GEGL
+ *
+ * GEGL is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * GEGL is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.
+ *
+ * Copyright 2013 Carlos Zubieta <czubieta dev gmail com>
+ */
+
+static const char* dot_cl_source =
+"__kernel void cl_calc_block_colors (__global const float4 *in, \n"
+" __global float4 *block_colors, \n"
+" int cx0, \n"
+" int cy0, \n"
+" int size, \n"
+" float weight, \n"
+" int roi_x, \n"
+" int roi_y, \n"
+" int line_width) \n"
+"{ \n"
+" int cx = cx0 + get_global_id(0); \n"
+" int cy = cy0 + get_global_id(1); \n"
+" int px = (cx * size) - roi_x + size; \n"
+" int py = (cy * size) - roi_y + size; \n"
+" float4 mean = (float4)(0.0f); \n"
+" float4 tmp; \n"
+" int i, j; \n"
+" \n"
+" for( j = py; j < py+size; ++j) \n"
+" { \n"
+" for (i = px; i < px+size; ++i) \n"
+" { \n"
+" mean += in[j * line_width + i]; \n"
+" } \n"
+" } \n"
+" block_colors[(cx-cx0) + get_global_size(0) * (cy-cy0)] = mean * weight; \n"
+"} \n"
+" \n"
+"__kernel void cl_dot (__global const float4 *block_colors, \n"
+" __global float4 *out, \n"
+" int cx0, \n"
+" int cy0, \n"
+" int size, \n"
+" float radius2, \n"
+" int roi_x, \n"
+" int roi_y, \n"
+" int block_count_x) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" int x = gidx + roi_x; \n"
+" int y = gidy + roi_y; \n"
+" int cy = y/size; \n"
+" int cx = x/size; \n"
+" float cellx = convert_float(x - cx * size) - convert_float(size) / 2.0; \n"
+" float celly = convert_float(y - cy * size) - convert_float(size) / 2.0; \n"
+" float4 tmp = (float4)(0.0); \n"
+" \n"
+" if((cellx * cellx + celly * celly) <= radius2) \n"
+" tmp = block_colors[(cx-cx0) + block_count_x * (cy-cy0)]; \n"
+" \n"
+" out[gidx + get_global_size(0) * gidy] = tmp; \n"
+"} \n"
+;
diff --git a/operations/common/dot.c b/operations/common/dot.c
index 3531273..0812ed1 100644
--- a/operations/common/dot.c
+++ b/operations/common/dot.c
@@ -62,6 +62,142 @@ prepare (GeglOperation *operation)
babl_format ("RGBA float"));
}
+#include "opencl/gegl-cl.h"
+#include "buffer/gegl-buffer-cl-iterator.h"
+#include "opencl/dot.cl.h"
+#include <stdio.h>
+
+GEGL_CL_STATIC
+
+static gboolean
+cl_dot (cl_mem in,
+ cl_mem out,
+ const GeglRectangle *roi,
+ gint size,
+ gfloat ratio)
+{
+ cl_int cl_err = 0;
+ cl_mem block_colors = NULL;
+
+ GEGL_CL_BUILD(dot, "cl_calc_block_colors", "cl_dot")
+
+ {
+ cl_int cl_size = size;
+ cl_float radius2 = (cl_float)(size*ratio*size*ratio)/4.0;
+ cl_float weight = 1.0 / (cl_float)(size*size);
+ cl_int roi_x = roi->x;
+ cl_int roi_y = roi->y;
+ cl_int line_width = roi->width + 2*size;
+ cl_int cx0 = CELL_X(roi->x, size);
+ cl_int cy0 = CELL_Y(roi->y, size);
+ cl_int block_count_x = CELL_X(roi->x + roi->width - 1, size) - cx0 + 1;
+ cl_int block_count_y = CELL_Y(roi->y + roi->height - 1, size) - cy0 + 1;
+ size_t glbl_s[2] = {block_count_x, block_count_y};
+
+ block_colors = gegl_clCreateBuffer(gegl_cl_get_context(),
+ CL_MEM_READ_WRITE,
+ glbl_s[0] * glbl_s[1] * sizeof(cl_float4),
+ NULL, &cl_err);
+ CL_CHECK;
+
+ GEGL_CL_ARG_START(cl_data->kernel[0])
+ GEGL_CL_ARG(cl_mem, in)
+ GEGL_CL_ARG(cl_mem, block_colors)
+ GEGL_CL_ARG(cl_int, cx0)
+ GEGL_CL_ARG(cl_int, cy0)
+ GEGL_CL_ARG(cl_int, cl_size)
+ GEGL_CL_ARG(cl_float, weight)
+ GEGL_CL_ARG(cl_int, roi_x)
+ GEGL_CL_ARG(cl_int, roi_y)
+ GEGL_CL_ARG(cl_int, line_width)
+ GEGL_CL_ARG_END
+
+ /* calculate the average color of all the blocks */
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 2,
+ NULL, glbl_s, NULL,
+ 0, NULL, NULL);
+ CL_CHECK;
+
+ glbl_s[0] = roi->width;
+ glbl_s[1] = roi->height;
+
+ GEGL_CL_ARG_START(cl_data->kernel[1])
+ GEGL_CL_ARG(cl_mem, block_colors)
+ GEGL_CL_ARG(cl_mem, out)
+ GEGL_CL_ARG(cl_int, cx0)
+ GEGL_CL_ARG(cl_int, cy0)
+ GEGL_CL_ARG(cl_int, cl_size)
+ GEGL_CL_ARG(cl_float, radius2)
+ GEGL_CL_ARG(cl_int, roi_x)
+ GEGL_CL_ARG(cl_int, roi_y)
+ GEGL_CL_ARG(cl_int, block_count_x)
+ GEGL_CL_ARG_END
+
+ /* set each pixel to the average color of the block it belongs to */
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[1], 2,
+ NULL, glbl_s, NULL,
+ 0, NULL, NULL);
+ CL_CHECK;
+
+ cl_err = gegl_clFinish(gegl_cl_get_command_queue ());
+ CL_CHECK;
+
+ GEGL_CL_RELEASE(block_colors)
+ }
+
+ return FALSE;
+
+error:
+ if(block_colors)
+ GEGL_CL_RELEASE(block_colors)
+ return TRUE;
+}
+
+static gboolean
+cl_process (GeglOperation *operation,
+ GeglBuffer *input,
+ GeglBuffer *output,
+ const GeglRectangle *result)
+{
+ const Babl *in_format = gegl_operation_get_format (operation, "input");
+ const Babl *out_format = gegl_operation_get_format (operation, "output");
+
+ gint err;
+
+ GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+
+ 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);
+
+ GEGL_CL_BUFFER_ITERATE_START(i, err)
+ {
+ err = cl_dot(i->tex[read],
+ i->tex[0],
+ &i->roi[0],
+ o->size,
+ o->ratio);
+ }
+ GEGL_CL_BUFFER_ITERATE_END(err)
+
+ return TRUE;
+}
+
static void
calc_block_colors (gfloat* block_colors,
const gfloat* input,
@@ -156,6 +292,14 @@ process (GeglOperation *operation,
GeglOperationAreaFilter *op_area;
gfloat* buf;
+ if (gegl_cl_is_accelerated ())
+ {
+ if (cl_process (operation, input, output, roi))
+ return TRUE;
+ else
+ gegl_cl_disable();
+ }
+
op_area = GEGL_OPERATION_AREA_FILTER (operation);
src_rect = *roi;
src_rect.x -= op_area->left;
@@ -165,10 +309,12 @@ process (GeglOperation *operation,
buf = g_new0 (gfloat, src_rect.width * src_rect.height * 4);
- gegl_buffer_get (input, &src_rect, 1.0, babl_format ("RGBA float"), buf, GEGL_AUTO_ROWSTRIDE,
GEGL_ABYSS_NONE);
+ gegl_buffer_get (input, &src_rect, 1.0, babl_format ("RGBA float"), buf,
+ GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);
dot(buf, roi, o);
- gegl_buffer_set (output, roi, 0, babl_format ("RGBA float"), buf, GEGL_AUTO_ROWSTRIDE);
+ gegl_buffer_set (output, roi, 0, babl_format ("RGBA float"), buf,
+ GEGL_AUTO_ROWSTRIDE);
g_free (buf);
@@ -187,6 +333,7 @@ gegl_chant_class_init (GeglChantClass *klass)
filter_class->process = process;
operation_class->prepare = prepare;
+ operation_class->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name", "gegl:dot",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
