[gegl/soc-2013-opecl-ops] Added OpenCL support to noise-hurl



commit 8416910e6ef5c30a13da94ce308a7afb27f20906
Author: Carlos Zubieta <czubieta dev gmail com>
Date:   Wed Sep 11 05:51:18 2013 -0500

    Added OpenCL support to noise-hurl

 opencl/noise-hurl.cl           |   56 ++++++++++++++++++
 opencl/noise-hurl.cl.h         |   58 ++++++++++++++++++
 operations/common/noise-hurl.c |  125 +++++++++++++++++++++++++++++++++++----
 3 files changed, 226 insertions(+), 13 deletions(-)
---
diff --git a/opencl/noise-hurl.cl b/opencl/noise-hurl.cl
new file mode 100644
index 0000000..6cc0412
--- /dev/null
+++ b/opencl/noise-hurl.cl
@@ -0,0 +1,56 @@
+/* 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_noise_hurl(__global       float4 *src,
+                            __global const int    *random_data,
+                            __global const long   *random_primes,
+                                           int     x_offset,
+                                           int     y_offset,
+                                           int     roi_width,
+                                           int     whole_region_width,
+                                           int     seed,
+                                           float   pct_random,
+                                           int     offset)
+{
+  int gid  = get_global_id(0);
+  int gidy = gid / roi_width;
+  int gidx = gid - gidy*roi_width;
+
+  int x = gidx + x_offset;
+  int y = gidy + y_offset;
+  int n = 4 * (x + whole_region_width * y + offset);
+
+  float4 src_v = src[gid];
+
+  float pc          = gegl_cl_random_float_range (random_data, random_primes,
+                                                  seed, x, y, 0, n, 0, 100);
+  float red_noise   = gegl_cl_random_float (random_data, random_primes,
+                                            seed, x, y, 0, n+1);
+  float green_noise = gegl_cl_random_float (random_data, random_primes,
+                                            seed, x, y, 0, n+2);
+  float blue_noise  = gegl_cl_random_float (random_data, random_primes,
+                                            seed, x, y, 0, n+3);
+
+  if(pc <= pct_random)
+    {
+      src_v.x = red_noise;
+      src_v.y = green_noise;
+      src_v.z = blue_noise;
+    }
+  src[gid] = src_v;
+}
diff --git a/opencl/noise-hurl.cl.h b/opencl/noise-hurl.cl.h
new file mode 100644
index 0000000..89d5dc1
--- /dev/null
+++ b/opencl/noise-hurl.cl.h
@@ -0,0 +1,58 @@
+/* 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* noise_hurl_cl_source =
+"__kernel void cl_noise_hurl(__global       float4 *src,                       \n"
+"                            __global const int    *random_data,               \n"
+"                            __global const long   *random_primes,             \n"
+"                                           int     x_offset,                  \n"
+"                                           int     y_offset,                  \n"
+"                                           int     roi_width,                 \n"
+"                                           int     whole_region_width,        \n"
+"                                           int     seed,                      \n"
+"                                           float   pct_random,                \n"
+"                                           int     offset)                    \n"
+"{                                                                             \n"
+"  int gid  = get_global_id(0);                                                \n"
+"  int gidy = gid / roi_width;                                                 \n"
+"  int gidx = gid - gidy*roi_width;                                            \n"
+"                                                                              \n"
+"  int x = gidx + x_offset;                                                    \n"
+"  int y = gidy + y_offset;                                                    \n"
+"  int n = 4 * (x + whole_region_width * y + offset);                          \n"
+"                                                                              \n"
+"  float4 src_v = src[gid];                                                    \n"
+"                                                                              \n"
+"  float pc          = gegl_cl_random_float_range (random_data, random_primes, \n"
+"                                                  seed, x, y, 0, n, 0, 100);  \n"
+"  float red_noise   = gegl_cl_random_float (random_data, random_primes,       \n"
+"                                            seed, x, y, 0, n+1);              \n"
+"  float green_noise = gegl_cl_random_float (random_data, random_primes,       \n"
+"                                            seed, x, y, 0, n+2);              \n"
+"  float blue_noise  = gegl_cl_random_float (random_data, random_primes,       \n"
+"                                            seed, x, y, 0, n+3);              \n"
+"                                                                              \n"
+"  if(pc <= pct_random)                                                        \n"
+"    {                                                                         \n"
+"      src_v.x = red_noise;                                                    \n"
+"      src_v.y = green_noise;                                                  \n"
+"      src_v.z = blue_noise;                                                   \n"
+"    }                                                                         \n"
+"  src[gid] = src_v;                                                           \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/noise-hurl.c b/operations/common/noise-hurl.c
index b047ca0..f2a88a3 100644
--- a/operations/common/noise-hurl.c
+++ b/operations/common/noise-hurl.c
@@ -61,12 +61,14 @@ process (GeglOperation       *operation,
          const GeglRectangle *roi,
          gint                 level)
 {
-  GeglChantO *o  = GEGL_CHANT_PROPERTIES (operation);
-  gfloat     *GEGL_ALIGNED in_pixel;
-  gfloat     *GEGL_ALIGNED out_pixel;
-  gfloat     *out_pix;
-  gint        i, cnt;
-  gint x, y, n;
+  GeglChantO    *o  = GEGL_CHANT_PROPERTIES (operation);
+  gfloat        *GEGL_ALIGNED in_pixel;
+  gfloat        *GEGL_ALIGNED out_pixel;
+  gfloat        *out_pix;
+  GeglRectangle *whole_region;
+  gint          total_size;
+  gint          i, cnt;
+  gint          offset;
 
   in_pixel  = in_buf;
   out_pixel = out_buf;
@@ -79,32 +81,37 @@ process (GeglOperation       *operation,
       in_pixel += 1;
     }
 
-  n = 0;
+  whole_region = gegl_operation_source_get_bounding_box (operation, "input");
+  total_size   = whole_region->width*whole_region->height;
+  offset = 0;
 
   for (cnt = 0; cnt < o->repeat; cnt++)
     {
-      x = roi->x;
-      y = roi->y;
+      gint x = roi->x;
+      gint y = roi->y;
 
       out_pix = out_pixel;
 
       for (i = 0; i < n_pixels; i++)
         {
           gfloat red, green, blue, alpha;
+          gint idx, n;
 
           red   = out_pix[0];
           green = out_pix[1];
           blue  = out_pix[2];
           alpha = out_pix[3];
 
+          idx = x + whole_region->width * y;
+          n = 4 * (idx + offset);
+
           if (gegl_random_float_range (o->seed, x, y, 0, n, 0.0, 100.0) <=
               o->pct_random)
             {
-              red   = gegl_random_float_range (o->seed, x, y, 0, n+1, 0.0, 1.0);
-              green = gegl_random_float_range (o->seed, x, y, 0, n+2, 0.0, 1.0);
-              blue  = gegl_random_float_range (o->seed, x, y, 0, n+3, 0.0, 1.0);
+              red   = gegl_random_float (o->seed, x, y, 0, n+1);
+              green = gegl_random_float (o->seed, x, y, 0, n+2);
+              blue  = gegl_random_float (o->seed, x, y, 0, n+3);
             }
-          n += 4;
 
           out_pix[0] = red;
           out_pix[1] = green;
@@ -120,11 +127,101 @@ process (GeglOperation       *operation,
               y++;
             }
         }
+      offset += total_size;
     }
 
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/noise-hurl.cl.h"
+
+GEGL_CL_STATIC
+
+static gboolean
+cl_process (GeglOperation       *operation,
+            cl_mem              in,
+            cl_mem              out,
+            size_t              global_worksize,
+            const GeglRectangle *roi,
+            gint                level)
+{
+  GeglChantO    *o          = GEGL_CHANT_PROPERTIES (operation);
+  GeglRectangle *wr         = gegl_operation_source_get_bounding_box (operation,
+                                                                      "input");
+  cl_int   cl_err           = 0;
+  cl_mem   cl_random_data   = NULL;
+  cl_mem   cl_random_primes = NULL;
+  cl_float pct_random       = o->pct_random;
+  cl_int   seed             = o->seed;
+  cl_int   x_offset         = roi->x;
+  cl_int   y_offset         = roi->y;
+  cl_int   roi_width        = roi->width;
+  cl_int   wr_width         = wr->width;
+  int      total_size       = wr->width*wr->height;
+  cl_int   offset;
+  int      it;
+
+  GEGL_CL_BUILD(noise_hurl, "cl_noise_hurl")
+
+  {
+  cl_random_data = gegl_cl_load_random_data(&cl_err);
+  CL_CHECK;
+  cl_random_primes = gegl_cl_load_random_primes(&cl_err);
+  CL_CHECK;
+
+  cl_err = gegl_clEnqueueCopyBuffer(gegl_cl_get_command_queue(),
+                                    in , out , 0 , 0 ,
+                                    global_worksize * sizeof(cl_float4),
+                                    0, NULL, NULL);
+  CL_CHECK;
+
+  GEGL_CL_ARG_START(cl_data->kernel[0])
+  GEGL_CL_ARG(cl_mem,   out)
+  GEGL_CL_ARG(cl_mem,   cl_random_data)
+  GEGL_CL_ARG(cl_mem,   cl_random_primes)
+  GEGL_CL_ARG(cl_int,   x_offset)
+  GEGL_CL_ARG(cl_int,   y_offset)
+  GEGL_CL_ARG(cl_int,   roi_width)
+  GEGL_CL_ARG(cl_int,   wr_width)
+  GEGL_CL_ARG(cl_int,   seed)
+  GEGL_CL_ARG(cl_float, pct_random)
+  GEGL_CL_ARG_END
+
+  offset = 0;
+
+  for(it = 0; it < o->repeat; ++it)
+    {
+      cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_int),
+                                   (void*)&offset);
+      CL_CHECK;
+      cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                         cl_data->kernel[0], 1,
+                                         NULL, &global_worksize, NULL,
+                                         0, NULL, NULL);
+      CL_CHECK;
+
+      offset += total_size;
+    }
+
+  cl_err = gegl_clFinish(gegl_cl_get_command_queue ());
+  CL_CHECK;
+
+  GEGL_CL_RELEASE(cl_random_data)
+  GEGL_CL_RELEASE(cl_random_primes)
+  }
+
+  return  FALSE;
+
+error:
+  if(cl_random_data)
+    GEGL_CL_RELEASE(cl_random_data)
+  if(cl_random_primes)
+    GEGL_CL_RELEASE(cl_random_primes)
+
+  return TRUE;
+}
+
 static void
 gegl_chant_class_init (GeglChantClass *klass)
 {
@@ -135,7 +232,9 @@ gegl_chant_class_init (GeglChantClass *klass)
   point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
   operation_class->prepare    = prepare;
+  operation_class->opencl_support = TRUE;
   point_filter_class->process = process;
+  point_filter_class->cl_process = cl_process;
 
   gegl_operation_class_set_keys (operation_class,
     "name",       "gegl:noise-hurl",


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