workshop: add the FIR case to the work-in-progress gaussian-blur

Things missing are:

- OpenCL handles only 4-components babl formats,
- The IIR case doesn't make the extent grow

2 files changed, 148 insertions, 0 deletions

diff --git a/opencl/gblur-1d.cl b/opencl/gblur-1d.cl
new file mode 100644
index 00000000..834cc858
--- /dev/null
+++ b/opencl/gblur-1d.cl
@@ -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 Victor Oliveira <victormatheus@gmail.com>
+ * Copyright 2013 Téo Mazars      <teomazars@gmail.com>
+ */
+
+__kernel void fir_ver_blur(const global float4 *src_buf,
+                                 global float4 *dst_buf,
+                           const global float  *cmatrix,
+                           const        int     clen)
+{
+    const int gidx          = get_global_id (0);
+    const int gidy          = get_global_id (1);
+    const int src_rowstride = get_global_size (0);
+    const int dst_rowstride = get_global_size (0);
+
+    const int half_clen     = clen / 2;
+
+    const int src_offset    = gidx + (gidy + half_clen) * src_rowstride;
+    const int dst_offset    = gidx +  gidy              * dst_rowstride;
+
+    const int src_start_ind = src_offset - half_clen * src_rowstride;
+
+    float4 v = 0.0f;
+
+    for (int i = 0; i < clen; i++)
+      {
+        v += src_buf[src_start_ind + i * src_rowstride] * cmatrix[i];
+      }
+
+    dst_buf[dst_offset] = v;
+}
+
+
+__kernel void fir_hor_blur(const global float4 *src_buf,
+                                 global float4 *dst_buf,
+                           const global float  *cmatrix,
+                           const        int     clen)
+{
+    const int gidx          = get_global_id (0);
+    const int gidy          = get_global_id (1);
+    const int src_rowstride = get_global_size (0) + clen - 1; /*== 2*(clen/2) */
+    const int dst_rowstride = get_global_size (0);
+
+    const int half_clen     = clen / 2;
+
+    const int src_offset    = gidx + gidy * src_rowstride + half_clen;
+    const int dst_offset    = gidx + gidy * dst_rowstride;
+
+    const int src_start_ind = src_offset - half_clen;
+
+    float4 v = 0.0f;
+
+    for (int i = 0; i < clen; i++)
+      {
+        v += src_buf[src_start_ind + i] * cmatrix[i];
+      }
+
+    dst_buf[dst_offset] = v;
+}
diff --git a/opencl/gblur-1d.cl.h b/opencl/gblur-1d.cl.h
new file mode 100644
index 00000000..a3391f9d
--- /dev/null
+++ b/opencl/gblur-1d.cl.h
@@ -0,0 +1,75 @@
+static const char* gblur_1d_cl_source =
+"/* This file is an image processing operation for GEGL                        \n"
+" *                                                                            \n"
+" * GEGL is free software; you can redistribute it and/or                      \n"
+" * modify it under the terms of the GNU Lesser General Public                 \n"
+" * License as published by the Free Software Foundation; either               \n"
+" * version 3 of the License, or (at your option) any later version.           \n"
+" *                                                                            \n"
+" * GEGL is distributed in the hope that it will be useful,                    \n"
+" * but WITHOUT ANY WARRANTY; without even the implied warranty of             \n"
+" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU          \n"
+" * Lesser General Public License for more details.                            \n"
+" *                                                                            \n"
+" * You should have received a copy of the GNU Lesser General Public           \n"
+" * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.       \n"
+" *                                                                            \n"
+" * Copyright 2013 Victor Oliveira <victormatheus@gmail.com>                   \n"
+" * Copyright 2013 T\303\251o Mazars      <teomazars@gmail.com>                \n"
+" */                                                                           \n"
+"                                                                              \n"
+"__kernel void fir_ver_blur(const global float4 *src_buf,                      \n"
+"                                 global float4 *dst_buf,                      \n"
+"                           const global float  *cmatrix,                      \n"
+"                           const        int     clen)                         \n"
+"{                                                                             \n"
+"    const int gidx          = get_global_id (0);                              \n"
+"    const int gidy          = get_global_id (1);                              \n"
+"    const int src_rowstride = get_global_size (0);                            \n"
+"    const int dst_rowstride = get_global_size (0);                            \n"
+"                                                                              \n"
+"    const int half_clen     = clen / 2;                                       \n"
+"                                                                              \n"
+"    const int src_offset    = gidx + (gidy + half_clen) * src_rowstride;      \n"
+"    const int dst_offset    = gidx +  gidy              * dst_rowstride;      \n"
+"                                                                              \n"
+"    const int src_start_ind = src_offset - half_clen * src_rowstride;         \n"
+"                                                                              \n"
+"    float4 v = 0.0f;                                                          \n"
+"                                                                              \n"
+"    for (int i = 0; i < clen; i++)                                            \n"
+"      {                                                                       \n"
+"        v += src_buf[src_start_ind + i * src_rowstride] * cmatrix[i];         \n"
+"      }                                                                       \n"
+"                                                                              \n"
+"    dst_buf[dst_offset] = v;                                                  \n"
+"}                                                                             \n"
+"                                                                              \n"
+"                                                                              \n"
+"__kernel void fir_hor_blur(const global float4 *src_buf,                      \n"
+"                                 global float4 *dst_buf,                      \n"
+"                           const global float  *cmatrix,                      \n"
+"                           const        int     clen)                         \n"
+"{                                                                             \n"
+"    const int gidx          = get_global_id (0);                              \n"
+"    const int gidy          = get_global_id (1);                              \n"
+"    const int src_rowstride = get_global_size (0) + clen - 1; /*== 2*(clen/2) */\n"
+"    const int dst_rowstride = get_global_size (0);                            \n"
+"                                                                              \n"
+"    const int half_clen     = clen / 2;                                       \n"
+"                                                                              \n"
+"    const int src_offset    = gidx + gidy * src_rowstride + half_clen;        \n"
+"    const int dst_offset    = gidx + gidy * dst_rowstride;                    \n"
+"                                                                              \n"
+"    const int src_start_ind = src_offset - half_clen;                         \n"
+"                                                                              \n"
+"    float4 v = 0.0f;                                                          \n"
+"                                                                              \n"
+"    for (int i = 0; i < clen; i++)                                            \n"
+"      {                                                                       \n"
+"        v += src_buf[src_start_ind + i] * cmatrix[i];                         \n"
+"      }                                                                       \n"
+"                                                                              \n"
+"    dst_buf[dst_offset] = v;                                                  \n"
+"}                                                                             \n"
+;