[gegl/soc-2013-opecl-ops: 5/6] Added OpenCL support to box-min



commit 46ecc6ba44c4aed99c918dea21880f2064025e0e
Author: Carlos Zubieta <czubieta dev gmail com>
Date:   Tue Jul 23 18:47:59 2013 -0500

    Added OpenCL support to box-min

 opencl/box-min.cl             |   49 ++++++++++++++
 opencl/box-min.cl.h           |   51 ++++++++++++++
 operations/workshop/box-min.c |  144 ++++++++++++++++++++++++++++++++++++++++-
 3 files changed, 242 insertions(+), 2 deletions(-)
---
diff --git a/opencl/box-min.cl b/opencl/box-min.cl
new file mode 100644
index 0000000..f0978de
--- /dev/null
+++ b/opencl/box-min.cl
@@ -0,0 +1,49 @@
+__kernel void kernel_min_hor (__global const float4     *in,
+                              __global       float4     *aux,
+                              int width, int radius)
+{
+  const int in_index = get_global_id(0) * (width + 2 * radius)
+                       + (radius + get_global_id (1));
+
+  const int aux_index = get_global_id(0) * width + get_global_id (1);
+  int i;
+  float4 min;
+  float4 in_v;
+
+  min = (float4)(1000000000.0f);
+
+  if (get_global_id(1) < width)
+    {
+      for (i=-radius; i <= radius; i++)
+        {
+          in_v = in[in_index + i];
+          min = min > in_v ? in_v : min;
+        }
+        aux[aux_index] = min;
+    }
+}
+
+__kernel void kernel_max_ver (__global const float4     *aux,
+                              __global       float4     *out,
+                              int width, int radius)
+{
+
+  const int out_index = get_global_id(0) * width + get_global_id (1);
+  int aux_index = out_index;
+  int i;
+  float4 min;
+  float4 aux_v;
+
+  min = (float4)(1000000000.0f);
+
+  if(get_global_id(1) < width)
+    {
+      for (i=-radius; i <= radius; i++)
+        {
+          aux_v = aux[aux_index];
+          min = min > aux_v ? aux_v : min;
+          aux_index += width;
+        }
+        out[out_index] = min;
+    }
+}
diff --git a/opencl/box-min.cl.h b/opencl/box-min.cl.h
new file mode 100644
index 0000000..1233258
--- /dev/null
+++ b/opencl/box-min.cl.h
@@ -0,0 +1,51 @@
+static const char* box_min_cl_source = 
+"__kernel void kernel_min_hor (__global const float4     *in,                  \n"
+"                              __global       float4     *aux,                 \n"
+"                              int width, int radius)                          \n"
+"{                                                                             \n"
+"  const int in_index = get_global_id(0) * (width + 2 * radius)                \n"
+"                       + (radius + get_global_id (1));                        \n"
+"                                                                              \n"
+"  const int aux_index = get_global_id(0) * width + get_global_id (1);         \n"
+"  int i;                                                                      \n"
+"  float4 min;                                                                 \n"
+"  float4 in_v;                                                                \n"
+"                                                                              \n"
+"  min = (float4)(1000000000.0f);                                              \n"
+"                                                                              \n"
+"  if (get_global_id(1) < width)                                               \n"
+"    {                                                                         \n"
+"      for (i=-radius; i <= radius; i++)                                       \n"
+"        {                                                                     \n"
+"          in_v = in[in_index + i];                                            \n"
+"          min = min > in_v ? in_v : min;                                      \n"
+"        }                                                                     \n"
+"        aux[aux_index] = min;                                                 \n"
+"    }                                                                         \n"
+"}                                                                             \n"
+"                                                                              \n"
+"__kernel void kernel_min_ver (__global const float4     *aux,                 \n"
+"                              __global       float4     *out,                 \n"
+"                              int width, int radius)                          \n"
+"{                                                                             \n"
+"                                                                              \n"
+"  const int out_index = get_global_id(0) * width + get_global_id (1);         \n"
+"  int aux_index = out_index;                                                  \n"
+"  int i;                                                                      \n"
+"  float4 min;                                                                 \n"
+"  float4 aux_v;                                                               \n"
+"                                                                              \n"
+"  min = (float4)(1000000000.0f);                                              \n"
+"                                                                              \n"
+"  if(get_global_id(1) < width)                                                \n"
+"    {                                                                         \n"
+"      for (i=-radius; i <= radius; i++)                                       \n"
+"        {                                                                     \n"
+"          aux_v = aux[aux_index];                                             \n"
+"          min = min > aux_v ? aux_v : min;                                    \n"
+"          aux_index += width;                                                 \n"
+"        }                                                                     \n"
+"        out[out_index] = min;                                                 \n"
+"    }                                                                         \n"
+"}                                                                             \n"
+;
diff --git a/operations/workshop/box-min.c b/operations/workshop/box-min.c
index 529a411..16aec0d 100644
--- a/operations/workshop/box-min.c
+++ b/operations/workshop/box-min.c
@@ -162,6 +162,136 @@ static void prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "output", babl_format ("RGBA float"));
 }
 
+#include "opencl/gegl-cl.h"
+#include "buffer/gegl-buffer-cl-iterator.h"
+
+#include "opencl/box-min.cl.h"
+
+static GeglClRunData *cl_data = NULL;
+
+static gboolean
+cl_box_min (cl_mem                in_tex,
+            cl_mem                aux_tex,
+            cl_mem                out_tex,
+            size_t                global_worksize,
+            const GeglRectangle  *roi,
+            gint                  radius)
+{
+  cl_int cl_err = 0;
+  size_t global_ws_hor[2], global_ws_ver[2];
+  size_t local_ws_hor[2], local_ws_ver[2];
+
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"kernel_min_hor", "kernel_min_ver", NULL};
+      cl_data = gegl_cl_compile_and_build (box_min_cl_source, kernel_name);
+    }
+  if (!cl_data) return TRUE;
+
+  local_ws_hor[0] = 1;
+  local_ws_hor[1] = 256;
+  global_ws_hor[0] = roi->height + 2 * radius;
+  global_ws_hor[1] = ((roi->width + local_ws_hor[1] -1)/local_ws_hor[1]) * local_ws_hor[1];
+
+  local_ws_ver[0] = 1;
+  local_ws_ver[1] = 256;
+  global_ws_ver[0] = roi->height;
+  global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1];
+
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
+  CL_CHECK;
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&aux_tex);
+  CL_CHECK;
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int),   (void*)&roi->width);
+  CL_CHECK;
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int),   (void*)&radius);
+  CL_CHECK;
+
+  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                        cl_data->kernel[0], 2,
+                                        NULL, global_ws_hor, local_ws_hor,
+                                        0, NULL, NULL);
+  CL_CHECK;
+
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem),   (void*)&aux_tex);
+  CL_CHECK;
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem),   (void*)&out_tex);
+  CL_CHECK;
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_int),   (void*)&roi->width);
+  CL_CHECK;
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_int),   (void*)&radius);
+  CL_CHECK;
+
+  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                        cl_data->kernel[1], 2,
+                                        NULL, global_ws_ver, local_ws_ver,
+                                        0, NULL, NULL);
+  CL_CHECK;
+
+  return FALSE;
+
+error:
+  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);
+
+  gint aux  = gegl_buffer_cl_iterator_add_2 (i,
+                                             NULL,
+                                             result,
+                                             in_format,
+                                             GEGL_CL_BUFFER_AUX,
+                                             0,
+                                             0,
+                                             op_area->top,
+                                             op_area->bottom,
+                                             GEGL_ABYSS_NONE);
+
+  while (gegl_buffer_cl_iterator_next (i, &err))
+    {
+      if (err) return FALSE;
+
+      err = cl_box_min(i->tex[read],
+                       i->tex[aux],
+                       i->tex[0],
+                       i->size[0],
+                       &i->roi[0],
+                       ceil (o->radius));
+
+      if (err) return FALSE;
+    }
+
+  return TRUE;
+}
+
 static gboolean
 process (GeglOperation       *operation,
          GeglBuffer          *input,
@@ -172,8 +302,16 @@ process (GeglOperation       *operation,
   GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
   GeglRectangle input_rect = gegl_operation_get_required_for_output (operation, "input", result);
 
-  hor_min ( input, &input_rect, output, result, o->radius);
-  ver_min (output,      result, output, result, o->radius);
+  if (gegl_cl_is_accelerated ())
+    {
+      if (cl_process (operation, input, output, result))
+        return TRUE;
+      else
+        gegl_cl_disable();
+    }
+
+  hor_mim ( input, &input_rect, output, result, o->radius);
+  ver_mim (output,      result, output, result, o->radius);
 
   return  TRUE;
 }
@@ -191,6 +329,8 @@ 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:box-min",
   "categories"  , "misc",


[Date Prev][Date Next]   [Thread Prev][Thread Next]   [Thread Index] [Date Index] [Author Index]