[gegl] Simplying opencl buffer iterators
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] Simplying opencl buffer iterators
- Date: Wed, 5 Jun 2013 00:41:43 +0000 (UTC)
commit 33bb648b5b19607d6b8e62c9dc8a7ffcc2e511f0
Author: Victor Oliveira <victormatheus gmail com>
Date: Tue Jun 4 21:13:01 2013 -0300
Simplying opencl buffer iterators
opencl buffer iterators now iterate over just
one region at a time, instead of possibly many.
This change is because the overhead of many
clFinish calls is not that great and it was
already happening in many places because of the
gpu caching code.
gegl/buffer/gegl-buffer-cl-iterator.c | 208 +++++++++----------
gegl/buffer/gegl-buffer-cl-iterator.h | 8 +-
gegl/opencl/gegl-cl.h | 13 +-
gegl/operation/gegl-operation-point-composer.c | 19 +-
gegl/operation/gegl-operation-point-filter.c | 17 +-
opencl/edge-laplace.cl | 180 ++++++++++++++++
opencl/edge-sobel.cl | 71 ++++++
operations/common/bilateral-filter-fast.c | 25 ++-
operations/common/bilateral-filter.c | 46 +++--
operations/common/box-blur.c | 60 ++++--
operations/common/c2g.c | 67 ++++---
operations/common/edge-laplace.c | 273 ++++++------------------
operations/common/edge-sobel.c | 150 +++++---------
operations/common/gaussian-blur.c | 69 ++++---
operations/common/motion-blur-linear.c | 41 +++--
operations/common/noise-reduction.c | 64 ++++--
operations/common/oilify.c | 46 +++--
operations/common/pixelize.c | 51 ++++--
operations/common/snn-mean.c | 38 +++-
operations/common/write-buffer.c | 50 +++--
20 files changed, 851 insertions(+), 645 deletions(-)
---
diff --git a/gegl/buffer/gegl-buffer-cl-iterator.c b/gegl/buffer/gegl-buffer-cl-iterator.c
index 814d6ff..9832e0d 100644
--- a/gegl/buffer/gegl-buffer-cl-iterator.c
+++ b/gegl/buffer/gegl-buffer-cl-iterator.c
@@ -39,17 +39,16 @@
typedef struct GeglBufferClIterators
{
/* current region of interest */
- gint n;
- size_t size [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX]; /* length of current data in pixels */
- cl_mem tex [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
- GeglRectangle roi [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
+ size_t size [GEGL_CL_BUFFER_MAX_ITERATORS]; /* length of current data in pixels */
+ cl_mem tex [GEGL_CL_BUFFER_MAX_ITERATORS];
+ GeglRectangle roi [GEGL_CL_BUFFER_MAX_ITERATORS];
/* the following is private: */
- cl_mem tex_buf [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
- cl_mem tex_op [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
+ cl_mem tex_buf [GEGL_CL_BUFFER_MAX_ITERATORS];
+ cl_mem tex_op [GEGL_CL_BUFFER_MAX_ITERATORS];
/* don't free textures loaded from cache */
- gboolean tex_buf_from_cache [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
+ gboolean tex_buf_from_cache [GEGL_CL_BUFFER_MAX_ITERATORS];
gint iterators;
gint iteration_no;
@@ -218,7 +217,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
GeglBufferClIterators *i = (gpointer)iterator;
gboolean result = FALSE;
- gint no, j;
+ gint no;
cl_int cl_err = 0;
int color_err = 0;
@@ -259,53 +258,51 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
/* color conversion in the GPU (output) */
if (i->conv[no] == GEGL_CL_COLOR_CONVERT)
- for (j=0; j < i->n; j++)
{
- color_err = gegl_cl_color_conv (i->tex_op[no][j], i->tex_buf[no][j], i->size[no][j],
+ color_err = gegl_cl_color_conv (i->tex_op[no], i->tex_buf[no], i->size[no],
i->format[no], i->buffer[no]->soft_format);
if (color_err) goto error;
}
/* GPU -> CPU */
- for (j=0; j < i->n; j++)
{
gpointer data;
/* tile-ize */
if (i->conv[no] == GEGL_CL_COLOR_NOT_SUPPORTED)
{
- data = g_malloc(i->size[no][j] * i->op_cl_format_size [no]);
+ data = g_malloc(i->size[no] * i->op_cl_format_size [no]);
cl_err = gegl_clEnqueueReadBuffer(gegl_cl_get_command_queue(),
- i->tex_op[no][j], CL_TRUE,
- 0, i->size[no][j] * i->op_cl_format_size [no], data,
+ i->tex_op[no], CL_TRUE,
+ 0, i->size[no] * i->op_cl_format_size[no], data,
0, NULL, NULL);
CL_CHECK;
/* color conversion using BABL */
- gegl_buffer_set (i->buffer[no], &i->roi[no][j], 0, i->format[no], data,
GEGL_AUTO_ROWSTRIDE);
+ gegl_buffer_set (i->buffer[no], &i->roi[no], 0, i->format[no], data,
GEGL_AUTO_ROWSTRIDE);
g_free(data);
}
else
#ifdef OPENCL_USE_CACHE
{
- gegl_buffer_cl_cache_new (i->buffer[no], &i->roi[no][j], i->tex_buf[no][j]);
+ gegl_buffer_cl_cache_new (i->buffer[no], &i->roi[no], i->tex_buf[no]);
/* don't release this texture */
- i->tex_buf[no][j] = NULL;
+ i->tex_buf[no] = NULL;
}
#else
{
- data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
+ data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no], CL_TRUE,
CL_MAP_READ,
- 0, i->size[no][j] * i->buf_cl_format_size [no],
+ 0, i->size[no] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
CL_CHECK;
/* color conversion using BABL */
- gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->format[no], data,
GEGL_AUTO_ROWSTRIDE);
+ gegl_buffer_set (i->buffer[no], &i->roi[no], i->format[no], data, GEGL_AUTO_ROWSTRIDE);
- cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j],
data,
+ cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no],
data,
0, NULL, NULL);
CL_CHECK;
}
@@ -319,47 +316,42 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
CL_CHECK;
for (no=0; no < i->iterators; no++)
- for (j=0; j < i->n; j++)
{
- if (i->tex_buf_from_cache [no][j])
+ if (i->tex_buf_from_cache [no])
{
- gboolean ok = gegl_buffer_cl_cache_release (i->tex_buf[no][j]);
+ gboolean ok = gegl_buffer_cl_cache_release (i->tex_buf[no]);
g_assert (ok);
}
- if (i->tex_buf[no][j] && !i->tex_buf_from_cache [no][j])
- gegl_clReleaseMemObject (i->tex_buf[no][j]);
+ if (i->tex_buf[no] && !i->tex_buf_from_cache [no])
+ gegl_clReleaseMemObject (i->tex_buf[no]);
- if (i->tex_op [no][j])
- gegl_clReleaseMemObject (i->tex_op [no][j]);
+ if (i->tex_op [no])
+ gegl_clReleaseMemObject (i->tex_op [no]);
- i->tex [no][j] = NULL;
- i->tex_buf[no][j] = NULL;
- i->tex_op [no][j] = NULL;
+ i->tex [no] = NULL;
+ i->tex_buf[no] = NULL;
+ i->tex_op [no] = NULL;
}
}
g_assert (i->iterators > 0);
result = (i->roi_no >= i->rois)? FALSE : TRUE;
- i->n = MIN(GEGL_CL_NTEX, i->rois - i->roi_no);
-
/* then we iterate all */
for (no=0; no<i->iterators;no++)
{
- for (j = 0; j < i->n; j++)
{
- GeglRectangle r = {i->rect[no].x + i->roi_all[i->roi_no+j].x - i->area[no][0],
- i->rect[no].y + i->roi_all[i->roi_no+j].y - i->area[no][2],
- i->roi_all[i->roi_no+j].width + i->area[no][0] + i->area[no][1],
- i->roi_all[i->roi_no+j].height + i->area[no][2] + i->area[no][3]};
- i->roi [no][j] = r;
- i->size[no][j] = r.width * r.height;
+ GeglRectangle r = {i->rect[no].x + i->roi_all[i->roi_no].x - i->area[no][0],
+ i->rect[no].y + i->roi_all[i->roi_no].y - i->area[no][2],
+ i->roi_all[i->roi_no].width + i->area[no][0] + i->area[no][1],
+ i->roi_all[i->roi_no].height + i->area[no][2] + i->area[no][3]};
+ i->roi [no] = r;
+ i->size[no] = r.width * r.height;
}
if (i->flags[no] == GEGL_CL_BUFFER_READ)
{
- for (j=0; j < i->n; j++)
{
gpointer data;
@@ -369,30 +361,30 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
case GEGL_CL_COLOR_NOT_SUPPORTED:
{
- gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]);
+ gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no]);
- g_assert (i->tex_op[no][j] == NULL);
- i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_op[no] == NULL);
+ i->tex_op[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
- i->size[no][j] * i->op_cl_format_size [no],
+ i->size[no] * i->op_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
/* pre-pinned memory */
- data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
+ data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no], CL_TRUE,
CL_MAP_WRITE,
- 0, i->size[no][j] * i->op_cl_format_size [no],
+ 0, i->size[no] * i->op_cl_format_size [no],
0, NULL, NULL, &cl_err);
CL_CHECK;
/* color conversion using BABL */
- gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->format[no], data,
GEGL_AUTO_ROWSTRIDE, i->abyss_policy[no]);
+ gegl_buffer_get (i->buffer[no], &i->roi[no], 1.0, i->format[no], data,
GEGL_AUTO_ROWSTRIDE, i->abyss_policy[no]);
- cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no][j],
data,
+ cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no], data,
0, NULL, NULL);
CL_CHECK;
- i->tex[no][j] = i->tex_op[no][j];
+ i->tex[no] = i->tex_op[no];
break;
}
@@ -400,37 +392,37 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
case GEGL_CL_COLOR_EQUAL:
{
- i->tex_buf[no][j] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no][j]);
+ i->tex_buf[no] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no]);
- if (i->tex_buf[no][j])
- i->tex_buf_from_cache [no][j] = TRUE; /* don't free texture from cache */
+ if (i->tex_buf[no])
+ i->tex_buf_from_cache [no] = TRUE; /* don't free texture from cache */
else
{
- gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]);
+ gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no]);
- g_assert (i->tex_buf[no][j] == NULL);
- i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_buf[no] == NULL);
+ i->tex_buf[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
- i->size[no][j] * i->buf_cl_format_size [no],
+ i->size[no] * i->buf_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
/* pre-pinned memory */
- data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j],
CL_TRUE,
+ data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no], CL_TRUE,
CL_MAP_WRITE,
- 0, i->size[no][j] * i->buf_cl_format_size [no],
+ 0, i->size[no] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
CL_CHECK;
/* color conversion will be performed in the GPU later */
- gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->buffer[no]->soft_format,
data, GEGL_AUTO_ROWSTRIDE, i->abyss_policy[no]);
+ gegl_buffer_get (i->buffer[no], &i->roi[no], 1.0, i->buffer[no]->soft_format, data,
GEGL_AUTO_ROWSTRIDE, i->abyss_policy[no]);
- cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(),
i->tex_buf[no][j], data,
+ cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no],
data,
0, NULL, NULL);
CL_CHECK;
}
- i->tex[no][j] = i->tex_buf[no][j];
+ i->tex[no] = i->tex_buf[no];
break;
}
@@ -438,50 +430,50 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
case GEGL_CL_COLOR_CONVERT:
{
- i->tex_buf[no][j] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no][j]);
+ i->tex_buf[no] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no]);
- if (i->tex_buf[no][j])
- i->tex_buf_from_cache [no][j] = TRUE; /* don't free texture from cache */
+ if (i->tex_buf[no])
+ i->tex_buf_from_cache [no] = TRUE; /* don't free texture from cache */
else
{
- gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]);
+ gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no]);
- g_assert (i->tex_buf[no][j] == NULL);
- i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_buf[no] == NULL);
+ i->tex_buf[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
- i->size[no][j] * i->buf_cl_format_size [no],
+ i->size[no] * i->buf_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
/* pre-pinned memory */
- data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j],
CL_TRUE,
+ data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no], CL_TRUE,
CL_MAP_WRITE,
- 0, i->size[no][j] * i->buf_cl_format_size [no],
+ 0, i->size[no] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
CL_CHECK;
/* color conversion will be performed in the GPU later */
- gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->buffer[no]->soft_format,
data, GEGL_AUTO_ROWSTRIDE, i->abyss_policy[no]);
+ gegl_buffer_get (i->buffer[no], &i->roi[no], 1.0, i->buffer[no]->soft_format, data,
GEGL_AUTO_ROWSTRIDE, i->abyss_policy[no]);
- cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(),
i->tex_buf[no][j], data,
+ cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no],
data,
0, NULL, NULL);
CL_CHECK;
}
- g_assert (i->tex_op[no][j] == NULL);
- i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_op[no] == NULL);
+ i->tex_op[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
- i->size[no][j] * i->op_cl_format_size [no],
+ i->size[no] * i->op_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
/* color conversion in the GPU (input) */
- g_assert (i->tex_buf[no][j] && i->tex_op[no][j]);
- color_err = gegl_cl_color_conv (i->tex_buf[no][j], i->tex_op[no][j], i->size[no][j],
+ g_assert (i->tex_buf[no] && i->tex_op[no]);
+ color_err = gegl_cl_color_conv (i->tex_buf[no], i->tex_op[no], i->size[no],
i->buffer[no]->soft_format, i->format[no]);
if (color_err) goto error;
- i->tex[no][j] = i->tex_op[no][j];
+ i->tex[no] = i->tex_op[no];
break;
}
@@ -490,21 +482,20 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
}
else if (i->flags[no] == GEGL_CL_BUFFER_WRITE)
{
- for (j=0; j < i->n; j++)
{
switch (i->conv[no])
{
case GEGL_CL_COLOR_NOT_SUPPORTED:
{
- g_assert (i->tex_op[no][j] == NULL);
- i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_op[no] == NULL);
+ i->tex_op[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_WRITE_ONLY,
- i->size[no][j] * i->op_cl_format_size [no],
+ i->size[no] * i->op_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
- i->tex[no][j] = i->tex_op[no][j];
+ i->tex[no] = i->tex_op[no];
break;
}
@@ -512,14 +503,14 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
case GEGL_CL_COLOR_EQUAL:
{
- g_assert (i->tex_buf[no][j] == NULL);
- i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_buf[no] == NULL);
+ i->tex_buf[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE, /* cache */
- i->size[no][j] * i->buf_cl_format_size [no],
+ i->size[no] * i->buf_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
- i->tex[no][j] = i->tex_buf[no][j];
+ i->tex[no] = i->tex_buf[no];
break;
}
@@ -527,21 +518,21 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
case GEGL_CL_COLOR_CONVERT:
{
- g_assert (i->tex_buf[no][j] == NULL);
- i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_buf[no] == NULL);
+ i->tex_buf[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE, /* cache */
- i->size[no][j] * i->buf_cl_format_size [no],
+ i->size[no] * i->buf_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
- g_assert (i->tex_op[no][j] == NULL);
- i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_op[no] == NULL);
+ i->tex_op[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
- i->size[no][j] * i->op_cl_format_size [no],
+ i->size[no] * i->op_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
- i->tex[no][j] = i->tex_op[no][j];
+ i->tex[no] = i->tex_op[no];
break;
}
@@ -550,22 +541,20 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
}
else if (i->flags[no] == GEGL_CL_BUFFER_AUX)
{
- for (j=0; j < i->n; j++)
{
- g_assert (i->tex_op[no][j] == NULL);
- i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
+ g_assert (i->tex_op[no] == NULL);
+ i->tex_op[no] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
- i->size[no][j] * i->op_cl_format_size [no],
+ i->size[no] * i->op_cl_format_size [no],
NULL, &cl_err);
CL_CHECK;
- i->tex[no][j] = i->tex_op[no][j];
+ i->tex[no] = i->tex_op[no];
}
}
}
- i->roi_no += i->n;
-
+ i->roi_no ++;
i->iteration_no++;
if (result == FALSE)
@@ -601,20 +590,15 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
error:
for (no=0; no<i->iterators;no++)
- for (j=0; j < i->n; j++)
{
- if (i->tex_buf[no][j]) gegl_clReleaseMemObject (i->tex_buf[no][j]);
- if (i->tex_op [no][j]) gegl_clReleaseMemObject (i->tex_op [no][j]);
+ if (i->tex_buf[no]) gegl_clReleaseMemObject (i->tex_buf[no]);
+ if (i->tex_op [no]) gegl_clReleaseMemObject (i->tex_op [no]);
- i->tex [no][j] = NULL;
- i->tex_buf[no][j] = NULL;
- i->tex_op [no][j] = NULL;
+ i->tex [no] = NULL;
+ i->tex_buf[no] = NULL;
+ i->tex_op [no] = NULL;
}
- /* something pretty bad happened, so it's better to just disable opencl at all for next operations */
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: Disabling OpenCL!");
- gegl_cl_disable();
-
*err = TRUE;
return FALSE;
}
diff --git a/gegl/buffer/gegl-buffer-cl-iterator.h b/gegl/buffer/gegl-buffer-cl-iterator.h
index 152d367..039d325 100644
--- a/gegl/buffer/gegl-buffer-cl-iterator.h
+++ b/gegl/buffer/gegl-buffer-cl-iterator.h
@@ -22,7 +22,6 @@
#include "gegl-buffer.h"
#include "opencl/gegl-cl.h"
-#define GEGL_CL_NTEX 16
#define GEGL_CL_BUFFER_MAX_ITERATORS 6
enum
@@ -34,10 +33,9 @@ enum
typedef struct GeglBufferClIterator
{
- gint n;
- size_t size [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX]; /* length of current data in pixels */
- cl_mem tex [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
- GeglRectangle roi [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
+ size_t size [GEGL_CL_BUFFER_MAX_ITERATORS]; /* length of current data in pixels */
+ cl_mem tex [GEGL_CL_BUFFER_MAX_ITERATORS];
+ GeglRectangle roi [GEGL_CL_BUFFER_MAX_ITERATORS];
} GeglBufferClIterator;
gint gegl_buffer_cl_iterator_add (GeglBufferClIterator *iterator,
diff --git a/gegl/opencl/gegl-cl.h b/gegl/opencl/gegl-cl.h
index ab67230..2137567 100644
--- a/gegl/opencl/gegl-cl.h
+++ b/gegl/opencl/gegl-cl.h
@@ -51,20 +51,13 @@
{ cl_err = gegl_clReleaseMemObject(obj); \
CL_CHECK; }
-#define GEGL_CL_BUFFER_ITERATE_START(I, J, ERR) \
- while (gegl_buffer_cl_iterator_next (I, & ERR)) \
+#define GEGL_CL_BUFFER_ITERATE_START(I, ERR) \
+ while (gegl_buffer_cl_iterator_next (I, & ERR)) \
{ \
if (ERR) return FALSE; \
- for (J=0; J < I ->n; J++) \
- {
#define GEGL_CL_BUFFER_ITERATE_END(ERR) \
- if (ERR) \
- { \
- g_warning("[OpenCL] Error"); \
- return FALSE; \
- } \
- } \
+ if (ERR) return FALSE; \
}
diff --git a/gegl/operation/gegl-operation-point-composer.c b/gegl/operation/gegl-operation-point-composer.c
index 4e606b4..79e16e9 100644
--- a/gegl/operation/gegl-operation-point-composer.c
+++ b/gegl/operation/gegl-operation-point-composer.c
@@ -172,19 +172,20 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
{
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format,
GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
+
if (aux)
foo = gegl_buffer_cl_iterator_add (i, aux, result, aux_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
- for (j=0; j < i->n; j++)
+
{
if (point_composer_class->cl_process)
{
- err = point_composer_class->cl_process(operation, i->tex[read][j],
- (aux)? i->tex[foo][j] : NULL,
- i->tex[0][j], i->size[0][j], &i->roi[0][j], level);
+ err = point_composer_class->cl_process(operation, i->tex[read],
+ (aux)? i->tex[foo] : NULL,
+ i->tex[0], i->size[0], &i->roi[0], level);
if (err)
{
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name);
@@ -196,16 +197,16 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
gint p = 0;
GeglClRunData *cl_data = operation_class->cl_data;
- cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem),
(void*)&i->tex[read][j]);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read]);
CL_CHECK;
if (aux)
- cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem),
(void*)&i->tex[foo][j]);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[foo]);
else
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), NULL);
CL_CHECK;
- cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[0][j]);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[0]);
CL_CHECK;
gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
@@ -213,7 +214,7 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
- NULL, &i->size[0][j], NULL,
+ NULL, &i->size[0], NULL,
0, NULL, NULL);
CL_CHECK;
}
@@ -229,7 +230,7 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
return TRUE;
error:
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring(cl_err));
return FALSE;
}
diff --git a/gegl/operation/gegl-operation-point-filter.c b/gegl/operation/gegl-operation-point-filter.c
index 1460560..9203dea 100644
--- a/gegl/operation/gegl-operation-point-filter.c
+++ b/gegl/operation/gegl-operation-point-filter.c
@@ -106,18 +106,20 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
{
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format,
GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
+
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
- for (j=0; j < i->n; j++)
+
{
if (point_filter_class->cl_process)
{
- err = point_filter_class->cl_process(operation, i->tex[read][j], i->tex[0][j],
- i->size[0][j], &i->roi[0][j], level);
+ err = point_filter_class->cl_process(operation, i->tex[read], i->tex[0],
+ i->size[0], &i->roi[0], level);
if (err)
{
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name);
+ gegl_cl_disable();
return FALSE;
}
}
@@ -126,9 +128,9 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
gint p = 0;
GeglClRunData *cl_data = operation_class->cl_data;
- cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem),
(void*)&i->tex[read][j]);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read]);
CL_CHECK;
- cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[ 0
][j]);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[ 0 ]);
CL_CHECK;
gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
@@ -136,7 +138,7 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
- NULL, &i->size[0][j], NULL,
+ NULL, &i->size[0], NULL,
0, NULL, NULL);
CL_CHECK;
}
@@ -152,7 +154,8 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
return TRUE;
error:
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring(cl_err));
+ gegl_cl_disable();
return FALSE;
}
diff --git a/opencl/edge-laplace.cl b/opencl/edge-laplace.cl
new file mode 100644
index 0000000..2259a5f
--- /dev/null
+++ b/opencl/edge-laplace.cl
@@ -0,0 +1,180 @@
+#define LAPLACE_RADIUS 1
+void minmax(float x1, float x2, float x3,
+ float x4, float x5,
+ float *min_result,
+ float *max_result)
+{
+ float min1, min2, max1, max2;
+
+ if (x1 > x2)
+ {
+ max1 = x1;
+ min1 = x2;
+ }
+ else
+ {
+ max1 = x2;
+ min1 = x1;
+ }
+
+ if (x3 > x4)
+ {
+ max2 = x3;
+ min2 = x4;
+ }
+ else
+ {
+ max2 = x4;
+ min2 = x3;
+ }
+
+ if (min1 < min2)
+ *min_result = fmin(min1, x5);
+ else
+ *min_result = fmin(min2, x5);
+ if (max1 > max2)
+ *max_result = fmax(max1, x5);
+ else
+ *max_result = fmax(max2, x5);
+}
+
+kernel void pre_edgelaplace (global float4 *in,
+ global float4 *out)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+
+ int src_width = get_global_size(0) + LAPLACE_RADIUS * 2;
+ int src_height = get_global_size(1);
+
+ int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;
+ int gid1d = i + j * src_width;
+
+ float pix_fl[4] = {
+ in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y,
+ in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w
+ };
+ float pix_fm[4] = {
+ in[gid1d - src_width].x, in[gid1d - src_width].y,
+ in[gid1d - src_width].z, in[gid1d - src_width].w
+ };
+ float pix_fr[4] = {
+ in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y,
+ in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w
+ };
+ float pix_ml[4] = {
+ in[gid1d - 1 ].x, in[gid1d - 1 ].y,
+ in[gid1d - 1 ].z, in[gid1d - 1 ].w
+ };
+ float pix_mm[4] = {
+ in[gid1d ].x, in[gid1d ].y,
+ in[gid1d ].z, in[gid1d ].w
+ };
+ float pix_mr[4] = {
+ in[gid1d + 1 ].x, in[gid1d + 1 ].y,
+ in[gid1d + 1 ].z, in[gid1d + 1 ].w
+ };
+ float pix_bl[4] = {
+ in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y,
+ in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w
+ };
+ float pix_bm[4] = {
+ in[gid1d + src_width].x, in[gid1d + src_width].y,
+ in[gid1d + src_width].z, in[gid1d + src_width].w
+ };
+ float pix_br[4] = {
+ in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y,
+ in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w
+ };
+
+ int c;
+ float minval, maxval;
+ float gradient[4];
+
+ for (c = 0;c < 3; ++c)
+ {
+ minmax(pix_fm[c], pix_bm[c], pix_ml[c], pix_mr[c],
+ pix_mm[c], &minval, &maxval);
+ gradient[c] = 0.5f *
+ fmax((maxval - pix_mm[c]),(pix_mm[c] - minval));
+ gradient[c] =
+ (pix_fl[c] + pix_fm[c] + pix_fr[c] +
+ pix_ml[c] + pix_mr[c] + pix_bl[c] +
+ pix_bm[c] + pix_br[c] - 8.0f * pix_mm[c]) >
+ 0.0f ? gradient[c] : -1.0f * gradient[c];
+ }
+ gradient[3] = pix_mm[3];
+
+ out[gid1d] = (float4)
+ (gradient[0], gradient[1], gradient[2], gradient[3]);
+}
+
+kernel void knl_edgelaplace (global float4 *in,
+ global float4 *out)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+
+ int src_width = get_global_size(0) + LAPLACE_RADIUS * 2;
+ int src_height = get_global_size(1);
+
+ int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;
+ int gid1d = i + j * src_width;
+
+ float pix_fl[4] = {
+ in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y,
+ in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w
+ };
+ float pix_fm[4] = {
+ in[gid1d - src_width].x, in[gid1d - src_width].y,
+ in[gid1d - src_width].z, in[gid1d - src_width].w
+ };
+ float pix_fr[4] = {
+ in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y,
+ in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w
+ };
+ float pix_ml[4] = {
+ in[gid1d - 1 ].x, in[gid1d - 1 ].y,
+ in[gid1d - 1 ].z, in[gid1d - 1 ].w
+ };
+ float pix_mm[4] = {
+ in[gid1d ].x, in[gid1d ].y,
+ in[gid1d ].z, in[gid1d ].w
+ };
+ float pix_mr[4] = {
+ in[gid1d + 1 ].x, in[gid1d + 1 ].y,
+ in[gid1d + 1 ].z, in[gid1d + 1 ].w
+ };
+ float pix_bl[4] = {
+ in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y,
+ in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w
+ };
+ float pix_bm[4] = {
+ in[gid1d + src_width].x, in[gid1d + src_width].y,
+ in[gid1d + src_width].z, in[gid1d + src_width].w
+ };
+ float pix_br[4] = {
+ in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y,
+ in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w
+ };
+
+ int c;
+ float value[4];
+
+ for (c = 0;c < 3; ++c)
+ {
+ float current = pix_mm[c];
+ current =
+ ((current > 0.0f) &&
+ (pix_fl[c] < 0.0f || pix_fm[c] < 0.0f ||
+ pix_fr[c] < 0.0f || pix_ml[c] < 0.0f ||
+ pix_mr[c] < 0.0f || pix_bl[c] < 0.0f ||
+ pix_bm[c] < 0.0f || pix_br[c] < 0.0f )
+ ) ? current : 0.0f;
+ value[c] = current;
+ }
+ value[3] = pix_mm[3];
+
+ out[gidx + gidy * get_global_size(0)] = (float4)
+ (value[0], value[1], value[2], value[3]);
+}
diff --git a/opencl/edge-sobel.cl b/opencl/edge-sobel.cl
new file mode 100644
index 0000000..ed8aed6
--- /dev/null
+++ b/opencl/edge-sobel.cl
@@ -0,0 +1,71 @@
+#define SOBEL_RADIUS 1
+kernel void kernel_edgesobel(global float4 *in,
+ global float4 *out,
+ const int horizontal,
+ const int vertical,
+ const int keep_signal,
+ const int has_alpha)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+
+ float4 hor_grad = 0.0f;
+ float4 ver_grad = 0.0f;
+ float4 gradient = 0.0f;
+
+ int dst_width = get_global_size(0);
+ int src_width = dst_width + SOBEL_RADIUS * 2;
+
+ int i = gidx + SOBEL_RADIUS, j = gidy + SOBEL_RADIUS;
+ int gid1d = i + j * src_width;
+
+ float4 pix_fl = in[gid1d - 1 - src_width];
+ float4 pix_fm = in[gid1d - src_width];
+ float4 pix_fr = in[gid1d + 1 - src_width];
+ float4 pix_ml = in[gid1d - 1 ];
+ float4 pix_mm = in[gid1d ];
+ float4 pix_mr = in[gid1d + 1 ];
+ float4 pix_bl = in[gid1d - 1 + src_width];
+ float4 pix_bm = in[gid1d + src_width];
+ float4 pix_br = in[gid1d + 1 + src_width];
+
+ if (horizontal)
+ {
+ hor_grad +=
+ - 1.0f * pix_fl + 1.0f * pix_fr
+ - 2.0f * pix_ml + 2.0f * pix_mr
+ - 1.0f * pix_bl + 1.0f * pix_br;
+ }
+ if (vertical)
+ {
+ ver_grad +=
+ - 1.0f * pix_fl - 2.0f * pix_fm
+ - 1.0f * pix_fr + 1.0f * pix_bl
+ + 2.0f * pix_bm + 1.0f * pix_br;
+ }
+
+ if (horizontal && vertical)
+ {
+ gradient = sqrt(
+ hor_grad * hor_grad +
+ ver_grad * ver_grad) / 1.41f;
+ }
+ else
+ {
+ if (keep_signal)
+ gradient = hor_grad + ver_grad;
+ else
+ gradient = fabs(hor_grad + ver_grad);
+ }
+
+ if (has_alpha)
+ {
+ gradient.w = pix_mm.w;
+ }
+ else
+ {
+ gradient.w = 1.0f;
+ }
+
+ out[gidx + gidy * dst_width] = gradient;
+}
diff --git a/operations/common/bilateral-filter-fast.c b/operations/common/bilateral-filter-fast.c
index d2fea54..1266a58 100644
--- a/operations/common/bilateral-filter-fast.c
+++ b/operations/common/bilateral-filter-fast.c
@@ -430,18 +430,25 @@ bilateral_cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- gint j;
- cl_int cl_err;
- GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
- gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format, GEGL_CL_BUFFER_READ,
GEGL_ABYSS_NONE);
+ GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
+ result,
+ out_format,
+ GEGL_CL_BUFFER_WRITE);
- GEGL_CL_BUFFER_ITERATE_START(i, j, err)
+ gint read = gegl_buffer_cl_iterator_add (i,
+ input,
+ result,
+ in_format,
+ GEGL_CL_BUFFER_READ,
+ GEGL_ABYSS_NONE);
+
+ GEGL_CL_BUFFER_ITERATE_START(i, err)
{
- err = cl_bilateral(i->tex[read][j],
- i->tex[0][j],
- &i->roi[0][j],
- &i->roi[read][j],
+ err = cl_bilateral(i->tex[read],
+ i->tex[0],
+ &i->roi[0],
+ &i->roi[read],
s_sigma,
r_sigma);
}
diff --git a/operations/common/bilateral-filter.c b/operations/common/bilateral-filter.c
index de4f775..d2f4392 100644
--- a/operations/common/bilateral-filter.c
+++ b/operations/common/bilateral-filter.c
@@ -76,10 +76,10 @@ cl_bilateral_filter (cl_mem in_tex,
size_t global_ws[2];
if (!cl_data)
- {
- const char *kernel_name[] = {"bilateral_filter", NULL};
- cl_data = gegl_cl_compile_and_build (bilateral_filter_cl_source, kernel_name);
- }
+ {
+ const char *kernel_name[] = {"bilateral_filter", NULL};
+ cl_data = gegl_cl_compile_and_build (bilateral_filter_cl_source, kernel_name);
+ }
if (!cl_data) return TRUE;
global_ws[0] = roi->width;
@@ -115,26 +115,38 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- gint j;
- cl_int cl_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);
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_bilateral_filter(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j],
ceil(o->blur_radius), o->edge_preservation);
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:bilateral-filter");
- return FALSE;
- }
- }
+
+ err = cl_bilateral_filter(i->tex[read],
+ i->tex[0],
+ i->size[0],
+ &i->roi[0],
+ ceil(o->blur_radius),
+ o->edge_preservation);
+
+ if (err) return FALSE;
}
return TRUE;
diff --git a/operations/common/box-blur.c b/operations/common/box-blur.c
index ca5563a..9ea43b8 100644
--- a/operations/common/box-blur.c
+++ b/operations/common/box-blur.c
@@ -256,29 +256,53 @@ cl_process (GeglOperation *operation,
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
+
gint err;
- gint j;
- cl_int cl_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);
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_box_blur(i->tex[read][j], i->tex[aux][j], i->tex[0][j], i->size[0][j], &i->roi[0][j],
ceil (o->radius));
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:box-blur");
- return FALSE;
- }
- }
+
+ err = cl_box_blur(i->tex[read],
+ i->tex[aux],
+ i->tex[0],
+ i->size[0],
+ &i->roi[0],
+ ceil (o->radius));
+
+ if (err) return FALSE;
}
+
return TRUE;
}
@@ -297,8 +321,12 @@ process (GeglOperation *operation,
op_area = GEGL_OPERATION_AREA_FILTER (operation);
if (gegl_cl_is_accelerated ())
- if (cl_process (operation, input, output, result))
- return TRUE;
+ {
+ if (cl_process (operation, input, output, result))
+ return TRUE;
+ else
+ gegl_cl_disable();
+ }
rect = *result;
tmprect = *result;
diff --git a/operations/common/c2g.c b/operations/common/c2g.c
index 14df6f3..b330bee 100644
--- a/operations/common/c2g.c
+++ b/operations/common/c2g.c
@@ -167,14 +167,14 @@ static GeglClRunData *cl_data = NULL;
static gboolean
cl_c2g (cl_mem in_tex,
- cl_mem out_tex,
- size_t global_worksize,
- const GeglRectangle *src_roi,
- const GeglRectangle *roi,
- gint radius,
- gint samples,
- gint iterations,
- gdouble rgamma)
+ cl_mem out_tex,
+ size_t global_worksize,
+ const GeglRectangle *src_roi,
+ const GeglRectangle *roi,
+ gint radius,
+ gint samples,
+ gint iterations,
+ gdouble rgamma)
{
cl_int cl_err = 0;
cl_mem cl_lut_cos, cl_lut_sin, cl_radiuses;
@@ -254,36 +254,51 @@ error:
}
static gboolean
-cl_process (GeglOperation *operation,
- GeglBuffer *input,
- GeglBuffer *output,
- const GeglRectangle *result)
+cl_process (GeglOperation *operation,
+ GeglBuffer *input,
+ GeglBuffer *output,
+ const GeglRectangle *result)
{
const Babl *in_format = babl_format("RGBA float");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- cl_int cl_err;
- gint j;
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);
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_c2g(i->tex[read][j], i->tex[0][j],i->size[0][j], &i->roi[read][j], &i->roi[0][j],
o->radius, o->samples, o->iterations, RGAMMA);
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:c2g");
- return FALSE;
- }
- }
+
+ err = cl_c2g(i->tex[read],
+ i->tex[0],
+ i->size[0],
+ &i->roi[read],
+ &i->roi[0],
+ o->radius,
+ o->samples,
+ o->iterations,
+ RGAMMA);
+
+ if (err) return FALSE;
}
+
return TRUE;
}
diff --git a/operations/common/edge-laplace.c b/operations/common/edge-laplace.c
index a8232fa..5dabff7 100644
--- a/operations/common/edge-laplace.c
+++ b/operations/common/edge-laplace.c
@@ -239,191 +239,11 @@ edge_laplace (GeglBuffer *src,
#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"
-static const char* kernel_source =
-"#define LAPLACE_RADIUS 1 \n"
-"void minmax(float x1, float x2, float x3, \n"
-" float x4, float x5, \n"
-" float *min_result, \n"
-" float *max_result) \n"
-"{ \n"
-" float min1, min2, max1, max2; \n"
-" \n"
-" if (x1 > x2) \n"
-" { \n"
-" max1 = x1; \n"
-" min1 = x2; \n"
-" } \n"
-" else \n"
-" { \n"
-" max1 = x2; \n"
-" min1 = x1; \n"
-" } \n"
-" \n"
-" if (x3 > x4) \n"
-" { \n"
-" max2 = x3; \n"
-" min2 = x4; \n"
-" } \n"
-" else \n"
-" { \n"
-" max2 = x4; \n"
-" min2 = x3; \n"
-" } \n"
-" \n"
-" if (min1 < min2) \n"
-" *min_result = fmin(min1, x5); \n"
-" else \n"
-" *min_result = fmin(min2, x5); \n"
-" if (max1 > max2) \n"
-" *max_result = fmax(max1, x5); \n"
-" else \n"
-" *max_result = fmax(max2, x5); \n"
-"} \n"
-" \n"
-"kernel void pre_edgelaplace (global float4 *in, \n"
-" global float4 *out) \n"
-"{ \n"
-" int gidx = get_global_id(0); \n"
-" int gidy = get_global_id(1); \n"
-" \n"
-" int src_width = get_global_size(0) + LAPLACE_RADIUS * 2; \n"
-" int src_height = get_global_size(1); \n"
-" \n"
-" int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS; \n"
-" int gid1d = i + j * src_width; \n"
-" \n"
-" float pix_fl[4] = { \n"
-" in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y, \n"
-" in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w \n"
-" }; \n"
-" float pix_fm[4] = { \n"
-" in[gid1d - src_width].x, in[gid1d - src_width].y, \n"
-" in[gid1d - src_width].z, in[gid1d - src_width].w \n"
-" }; \n"
-" float pix_fr[4] = { \n"
-" in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y, \n"
-" in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w \n"
-" }; \n"
-" float pix_ml[4] = { \n"
-" in[gid1d - 1 ].x, in[gid1d - 1 ].y, \n"
-" in[gid1d - 1 ].z, in[gid1d - 1 ].w \n"
-" }; \n"
-" float pix_mm[4] = { \n"
-" in[gid1d ].x, in[gid1d ].y, \n"
-" in[gid1d ].z, in[gid1d ].w \n"
-" }; \n"
-" float pix_mr[4] = { \n"
-" in[gid1d + 1 ].x, in[gid1d + 1 ].y, \n"
-" in[gid1d + 1 ].z, in[gid1d + 1 ].w \n"
-" }; \n"
-" float pix_bl[4] = { \n"
-" in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y, \n"
-" in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w \n"
-" }; \n"
-" float pix_bm[4] = { \n"
-" in[gid1d + src_width].x, in[gid1d + src_width].y, \n"
-" in[gid1d + src_width].z, in[gid1d + src_width].w \n"
-" }; \n"
-" float pix_br[4] = { \n"
-" in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y, \n"
-" in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w \n"
-" }; \n"
-" \n"
-" int c; \n"
-" float minval, maxval; \n"
-" float gradient[4]; \n"
-" \n"
-" for (c = 0;c < 3; ++c) \n"
-" { \n"
-" minmax(pix_fm[c], pix_bm[c], pix_ml[c], pix_mr[c], \n"
-" pix_mm[c], &minval, &maxval); \n"
-" gradient[c] = 0.5f * \n"
-" fmax((maxval - pix_mm[c]),(pix_mm[c] - minval)); \n"
-" gradient[c] = \n"
-" (pix_fl[c] + pix_fm[c] + pix_fr[c] + \n"
-" pix_ml[c] + pix_mr[c] + pix_bl[c] + \n"
-" pix_bm[c] + pix_br[c] - 8.0f * pix_mm[c]) > \n"
-" 0.0f ? gradient[c] : -1.0f * gradient[c]; \n"
-" } \n"
-" gradient[3] = pix_mm[3]; \n"
-" \n"
-" out[gid1d] = (float4) \n"
-" (gradient[0], gradient[1], gradient[2], gradient[3]); \n"
-"} \n"
-" \n"
-"kernel void knl_edgelaplace (global float4 *in, \n"
-" global float4 *out) \n"
-"{ \n"
-" int gidx = get_global_id(0); \n"
-" int gidy = get_global_id(1); \n"
-" \n"
-" int src_width = get_global_size(0) + LAPLACE_RADIUS * 2; \n"
-" int src_height = get_global_size(1); \n"
-" \n"
-" int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS; \n"
-" int gid1d = i + j * src_width; \n"
-" \n"
-" float pix_fl[4] = { \n"
-" in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y, \n"
-" in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w \n"
-" }; \n"
-" float pix_fm[4] = { \n"
-" in[gid1d - src_width].x, in[gid1d - src_width].y, \n"
-" in[gid1d - src_width].z, in[gid1d - src_width].w \n"
-" }; \n"
-" float pix_fr[4] = { \n"
-" in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y, \n"
-" in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w \n"
-" }; \n"
-" float pix_ml[4] = { \n"
-" in[gid1d - 1 ].x, in[gid1d - 1 ].y, \n"
-" in[gid1d - 1 ].z, in[gid1d - 1 ].w \n"
-" }; \n"
-" float pix_mm[4] = { \n"
-" in[gid1d ].x, in[gid1d ].y, \n"
-" in[gid1d ].z, in[gid1d ].w \n"
-" }; \n"
-" float pix_mr[4] = { \n"
-" in[gid1d + 1 ].x, in[gid1d + 1 ].y, \n"
-" in[gid1d + 1 ].z, in[gid1d + 1 ].w \n"
-" }; \n"
-" float pix_bl[4] = { \n"
-" in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y, \n"
-" in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w \n"
-" }; \n"
-" float pix_bm[4] = { \n"
-" in[gid1d + src_width].x, in[gid1d + src_width].y, \n"
-" in[gid1d + src_width].z, in[gid1d + src_width].w \n"
-" }; \n"
-" float pix_br[4] = { \n"
-" in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y, \n"
-" in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w \n"
-" }; \n"
-" \n"
-" int c; \n"
-" float value[4]; \n"
-" \n"
-" for (c = 0;c < 3; ++c) \n"
-" { \n"
-" float current = pix_mm[c]; \n"
-" current = \n"
-" ((current > 0.0f) && \n"
-" (pix_fl[c] < 0.0f || pix_fm[c] < 0.0f || \n"
-" pix_fr[c] < 0.0f || pix_ml[c] < 0.0f || \n"
-" pix_mr[c] < 0.0f || pix_bl[c] < 0.0f || \n"
-" pix_bm[c] < 0.0f || pix_br[c] < 0.0f ) \n"
-" ) ? current : 0.0f; \n"
-" value[c] = current; \n"
-" } \n"
-" value[3] = pix_mm[3]; \n"
-" \n"
-" out[gidx + gidy * get_global_size(0)] = (float4) \n"
-" (value[0], value[1], value[2], value[3]); \n"
-"} \n";
+#include "opencl/edge-laplace.cl.h"
static GeglClRunData *cl_data = NULL;
-static cl_int
+static gboolean
cl_edge_laplace (cl_mem in_tex,
cl_mem aux_tex,
cl_mem out_tex,
@@ -433,37 +253,43 @@ cl_edge_laplace (cl_mem in_tex,
{
cl_int cl_err = 0;
size_t global_ws[2];
- if (!cl_data)
- {
- const char *kernel_name[] = {"pre_edgelaplace", "knl_edgelaplace", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
- }
- if (!cl_data) return 1;
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"pre_edgelaplace", "knl_edgelaplace", NULL};
+ cl_data = gegl_cl_compile_and_build (edge_laplace_cl_source, kernel_name);
+ }
+ if (!cl_data) return TRUE;
global_ws[0] = roi->width;
global_ws[1] = roi->height;
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&aux_tex);
- if (cl_err != CL_SUCCESS) return cl_err;
+ 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_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
+ CL_CHECK;
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&aux_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&out_tex);
- if (cl_err != CL_SUCCESS) return cl_err;
+ 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_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[1], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
- return cl_err;
+ CL_CHECK;
+
+ return FALSE;
+
+error:
+ return TRUE;
}
static gboolean
@@ -475,27 +301,50 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- gint j;
- cl_int cl_err;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (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, op_area->left,
op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE);
+ 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,
+ 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;
- for (j=0; j < i->n; j++)
{
- cl_err = cl_edge_laplace(i->tex[read][j], i->tex[aux][j], i->tex[0][j], &i->roi[read][j],
&i->roi[0][j], LAPLACE_RADIUS);
- if (cl_err != CL_SUCCESS)
- {
- g_warning("[OpenCL] Error in gegl:edge-laplace: %s", gegl_cl_errstring(cl_err));
- return FALSE;
- }
+ if (err) return FALSE;
+
+ err = cl_edge_laplace(i->tex[read],
+ i->tex[aux],
+ i->tex[0],
+ &i->roi[read],
+ &i->roi[0],
+ LAPLACE_RADIUS);
+
+ if (err) return FALSE;
}
- }
+
return TRUE;
}
diff --git a/operations/common/edge-sobel.c b/operations/common/edge-sobel.c
index 2ebc65f..3007e60 100644
--- a/operations/common/edge-sobel.c
+++ b/operations/common/edge-sobel.c
@@ -79,82 +79,11 @@ static void prepare (GeglOperation *operation)
#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"
-static const char* kernel_source =
-"#define SOBEL_RADIUS 1 \n"
-"kernel void kernel_edgesobel(global float4 *in, \n"
-" global float4 *out, \n"
-" const int horizontal, \n"
-" const int vertical, \n"
-" const int keep_signal, \n"
-" const int has_alpha) \n"
-"{ \n"
-" int gidx = get_global_id(0); \n"
-" int gidy = get_global_id(1); \n"
-" \n"
-" float4 hor_grad = 0.0f; \n"
-" float4 ver_grad = 0.0f; \n"
-" float4 gradient = 0.0f; \n"
-" \n"
-" int dst_width = get_global_size(0); \n"
-" int src_width = dst_width + SOBEL_RADIUS * 2; \n"
-" \n"
-" int i = gidx + SOBEL_RADIUS, j = gidy + SOBEL_RADIUS; \n"
-" int gid1d = i + j * src_width; \n"
-" \n"
-" float4 pix_fl = in[gid1d - 1 - src_width]; \n"
-" float4 pix_fm = in[gid1d - src_width]; \n"
-" float4 pix_fr = in[gid1d + 1 - src_width]; \n"
-" float4 pix_ml = in[gid1d - 1 ]; \n"
-" float4 pix_mm = in[gid1d ]; \n"
-" float4 pix_mr = in[gid1d + 1 ]; \n"
-" float4 pix_bl = in[gid1d - 1 + src_width]; \n"
-" float4 pix_bm = in[gid1d + src_width]; \n"
-" float4 pix_br = in[gid1d + 1 + src_width]; \n"
-" \n"
-" if (horizontal) \n"
-" { \n"
-" hor_grad += \n"
-" - 1.0f * pix_fl + 1.0f * pix_fr \n"
-" - 2.0f * pix_ml + 2.0f * pix_mr \n"
-" - 1.0f * pix_bl + 1.0f * pix_br; \n"
-" } \n"
-" if (vertical) \n"
-" { \n"
-" ver_grad += \n"
-" - 1.0f * pix_fl - 2.0f * pix_fm \n"
-" - 1.0f * pix_fr + 1.0f * pix_bl \n"
-" + 2.0f * pix_bm + 1.0f * pix_br; \n"
-" } \n"
-" \n"
-" if (horizontal && vertical) \n"
-" { \n"
-" gradient = sqrt( \n"
-" hor_grad * hor_grad + \n"
-" ver_grad * ver_grad) / 1.41f; \n"
-" } \n"
-" else \n"
-" { \n"
-" if (keep_signal) \n"
-" gradient = hor_grad + ver_grad; \n"
-" else \n"
-" gradient = fabs(hor_grad + ver_grad); \n"
-" } \n"
-" \n"
-" if (has_alpha) \n"
-" { \n"
-" gradient.w = pix_mm.w; \n"
-" } \n"
-" else \n"
-" { \n"
-" gradient.w = 1.0f; \n"
-" } \n"
-" \n"
-" out[gidx + gidy * dst_width] = gradient; \n"
-"} \n";
+#include "opencl/edge-sobel.cl.h"
static GeglClRunData *cl_data = NULL;
-static cl_int
+static gboolean
cl_edge_sobel (cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
@@ -174,25 +103,33 @@ cl_edge_sobel (cl_mem in_tex,
if (!cl_data)
{
const char *kernel_name[] = {"kernel_edgesobel", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ cl_data = gegl_cl_compile_and_build (edge_sobel_cl_source, kernel_name);
}
- if (!cl_data) return 0;
-
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&n_horizontal);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&n_vertical);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&n_keep_signal);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&n_has_alpha);
- if (cl_err != CL_SUCCESS) return cl_err;
+ if (!cl_data) return TRUE;
+
+ 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*)&out_tex);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&n_horizontal);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&n_vertical);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&n_keep_signal);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&n_has_alpha);
+ CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
cl_data->kernel[0], 2,
NULL, gbl_size, NULL,
0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
+ CL_CHECK;
- return CL_SUCCESS;
+ return FALSE;
+
+error:
+ return TRUE;
}
static gboolean
@@ -205,27 +142,42 @@ cl_process (GeglOperation *operation,
const Babl *in_format = babl_format ("RGBA float");
const Babl *out_format = babl_format ("RGBA float");
gint err;
- gint j;
- cl_int cl_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);
+ 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;
- for (j=0; j < i->n; j++)
{
- cl_err = cl_edge_sobel(i->tex[read][j], i->tex[0][j], i->size[0][j],&i->roi[0][j], o->horizontal,
o->vertical, o->keep_signal, has_alpha);
- if (cl_err != CL_SUCCESS)
- {
- g_warning("[OpenCL] Error in gegl:edge-sobel: %s", gegl_cl_errstring(cl_err));
- return FALSE;
- }
+ if (err) return FALSE;
+
+ err = cl_edge_sobel(i->tex[read],
+ i->tex[0],
+ i->size[0],
+ &i->roi[0],
+ o->horizontal,
+ o->vertical,
+ o->keep_signal,
+ has_alpha);
+
+ if (err) return FALSE;
}
- }
+
return TRUE;
}
diff --git a/operations/common/gaussian-blur.c b/operations/common/gaussian-blur.c
index b9f22b7..14f73db 100644
--- a/operations/common/gaussian-blur.c
+++ b/operations/common/gaussian-blur.c
@@ -543,7 +543,6 @@ cl_process (GeglOperation *operation,
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
gint j;
- cl_int cl_err;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
@@ -566,36 +565,52 @@ cl_process (GeglOperation *operation,
fmatrix_y[j] = (gfloat) cmatrix_y[j];
{
- 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);
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_gaussian_blur(i->tex[read][j],
- i->tex[0][j],
- i->tex[aux][j],
- i->size[0][j],
- &i->roi[0][j],
- &i->roi[read][j],
- &i->roi[aux][j],
- fmatrix_x,
- cmatrix_len_x,
- op_area->left,
- fmatrix_y,
- cmatrix_len_y,
- op_area->top);
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:gaussian-blur");
- return FALSE;
- }
- }
+
+ err = cl_gaussian_blur(i->tex[read],
+ i->tex[0],
+ i->tex[aux],
+ i->size[0],
+ &i->roi[0],
+ &i->roi[read],
+ &i->roi[aux],
+ fmatrix_x,
+ cmatrix_len_x,
+ op_area->left,
+ fmatrix_y,
+ cmatrix_len_y,
+ op_area->top);
+
+ if (err) return FALSE;
}
}
diff --git a/operations/common/motion-blur-linear.c b/operations/common/motion-blur-linear.c
index 792619a..7c187c0 100644
--- a/operations/common/motion-blur-linear.c
+++ b/operations/common/motion-blur-linear.c
@@ -137,29 +137,42 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- cl_int cl_err;
- gint j;
gdouble theta = o->angle * G_PI / 180.0;
gfloat offset_x = (gfloat)(o->length * cos(theta));
gfloat offset_y = (gfloat)(o->length * sin(theta));
gint num_steps = (gint)ceil(o->length) + 1;
- 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_CLAMP);
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_motion_blur_linear(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j],
&i->roi[read][j], num_steps, offset_x, offset_y);
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:motion-blur-linear");
- return FALSE;
- }
- }
+
+ err = cl_motion_blur_linear(i->tex[read],
+ i->tex[0],
+ i->size[0],
+ &i->roi[0],
+ &i->roi[read],
+ num_steps,
+ offset_x,
+ offset_y);
+
+ if (err) return FALSE;
}
return TRUE;
diff --git a/operations/common/noise-reduction.c b/operations/common/noise-reduction.c
index b0f1063..500adbf 100644
--- a/operations/common/noise-reduction.c
+++ b/operations/common/noise-reduction.c
@@ -257,36 +257,52 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- gint j;
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,
- op_area->left, op_area->right, op_area->top, op_area->bottom,
GEGL_ABYSS_NONE);
+ 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,
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_noise_reduction(i->tex[read][j],
- i->tex[aux][j],
- i->tex[0][j],
- i->size[0][j],
- &i->roi[read][j],
- &i->roi[0][j],
- o->iterations);
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:noise-reduction");
- return FALSE;
- }
- }
- }
+ {
+ if (err) return FALSE;
+
+ err = cl_noise_reduction(i->tex[read],
+ i->tex[aux],
+ i->tex[0],
+ i->size[0],
+ &i->roi[read],
+ &i->roi[0],
+ o->iterations);
+
+ if (err) return FALSE;
+ }
+
return TRUE;
}
diff --git a/operations/common/oilify.c b/operations/common/oilify.c
index aee83fc..77fd317 100644
--- a/operations/common/oilify.c
+++ b/operations/common/oilify.c
@@ -340,31 +340,39 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- gint j;
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,
- o->mask_radius, o->mask_radius, o->mask_radius,
o->mask_radius, GEGL_ABYSS_CLAMP);
+ 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,
+ o->mask_radius,
+ o->mask_radius,
+ o->mask_radius,
+ o->mask_radius,
+ GEGL_ABYSS_CLAMP);
+
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
- for (j=0; j < i->n; j++)
- {
- err = cl_oilify(i->tex[read][j],
- i->tex[0][j],
- i->size[0][j],&i->roi[0][j],
- o->mask_radius,
- o->intensities,
- o->exponent,
- o->use_inten);
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:oilify");
- return FALSE;
- }
- }
+
+ err = cl_oilify(i->tex[read],
+ i->tex[0],
+ i->size[0],
+ &i->roi[0],
+ o->mask_radius,
+ o->intensities,
+ o->exponent,
+ o->use_inten);
+
+ if (err) return FALSE;
}
return TRUE;
diff --git a/operations/common/pixelize.c b/operations/common/pixelize.c
index e3a27e2..4756dc9 100644
--- a/operations/common/pixelize.c
+++ b/operations/common/pixelize.c
@@ -235,27 +235,52 @@ cl_process (GeglOperation *operation,
gboolean has_alpha = babl_format_has_alpha (gegl_operation_get_format (operation, "output"));
GeglAbyssPolicy read_abyss = has_alpha ? GEGL_ABYSS_NONE : GEGL_ABYSS_BLACK;
gint err;
- gint j;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
- GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, roi, out_format, GEGL_CL_BUFFER_WRITE);
- gint read = gegl_buffer_cl_iterator_add_2 (i, input, roi, in_format, GEGL_CL_BUFFER_READ,
op_area->left, op_area->right, op_area->top, op_area->bottom, read_abyss);
- gint aux = gegl_buffer_cl_iterator_add_2 (i, NULL, roi, in_format, GEGL_CL_BUFFER_AUX,
op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE);
+ GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
+ roi,
+ out_format,
+ GEGL_CL_BUFFER_WRITE);
+
+ gint read = gegl_buffer_cl_iterator_add_2 (i,
+ input,
+ roi,
+ in_format,
+ GEGL_CL_BUFFER_READ,
+ op_area->left,
+ op_area->right,
+ op_area->top,
+ op_area->bottom,
+ read_abyss);
+
+ gint aux = gegl_buffer_cl_iterator_add_2 (i,
+ NULL,
+ roi,
+ in_format,
+ GEGL_CL_BUFFER_AUX,
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_pixelise(i->tex[read][j], i->tex[aux][j], i->tex[0][j],&i->roi[read][j], &i->roi[0][j],
o->size_x, o->size_y);
- if (err != CL_SUCCESS)
- {
- g_warning("[OpenCL] Error in gegl:pixelize");
- return FALSE;
- }
- }
+
+ err = cl_pixelise(i->tex[read],
+ i->tex[aux],
+ i->tex[0],
+ &i->roi[read],
+ &i->roi[0],
+ o->size_x,
+ o->size_y);
+
+ if (err) return FALSE;
}
+
return TRUE;
}
diff --git a/operations/common/snn-mean.c b/operations/common/snn-mean.c
index 52d320c..8fff871 100644
--- a/operations/common/snn-mean.c
+++ b/operations/common/snn-mean.c
@@ -282,26 +282,40 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
- gint j;
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);
+ 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;
- for (j=0; j < i->n; j++)
- {
- err = cl_snn_mean(i->tex[read][j], i->tex[0][j], &i->roi[read][j], &i->roi[0][j], ceil(o->radius),
o->pairs);
- if (err)
- {
- g_warning("[OpenCL] Error in gegl:snn-mean");
- return FALSE;
- }
- }
+
+ err = cl_snn_mean(i->tex[read],
+ i->tex[0],
+ &i->roi[read],
+ &i->roi[0],
+ ceil(o->radius),
+ o->pairs);
+
+ if (err) return FALSE;
}
+
return TRUE;
}
diff --git a/operations/common/write-buffer.c b/operations/common/write-buffer.c
index eb86bdf..2986199 100644
--- a/operations/common/write-buffer.c
+++ b/operations/common/write-buffer.c
@@ -57,29 +57,51 @@ process (GeglOperation *operation,
size_t size;
gboolean err;
cl_int cl_err = 0;
- gint j;
- GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, output->soft_format,
GEGL_CL_BUFFER_WRITE);
- gint read = gegl_buffer_cl_iterator_add (i, input, result, output->soft_format,
GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
+ GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
+ result,
+ output->soft_format,
+ GEGL_CL_BUFFER_WRITE);
+
+ gint read = gegl_buffer_cl_iterator_add (i,
+ input,
+ result,
+ output->soft_format,
+ GEGL_CL_BUFFER_READ,
+ GEGL_ABYSS_NONE);
gegl_cl_color_babl (output->soft_format, &size);
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "write-buffer: %p %p %s %s {%d %d %d %d}", input, output,
babl_get_name(input->soft_format), babl_get_name(output->soft_format),
- result->x, result->y,
result->width, result->height);
+ GEGL_NOTE (GEGL_DEBUG_OPENCL,
+ "write-buffer: "
+ "%p %p %s %s {%d %d %d %d}",
+ input,
+ output,
+ babl_get_name(input->soft_format),
+ babl_get_name(output->soft_format),
+ result->x,
+ result->y,
+ result->width,
+ result->height);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) break;
- for (j=0; j < i->n; j++)
+
+ cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue (),
+ i->tex[read],
+ i->tex[0],
+ 0,
+ 0,
+ i->size[0] * size,
+ 0,
+ NULL,
+ NULL);
+
+ if (cl_err != CL_SUCCESS)
{
- cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue (),
- i->tex[read][j], i->tex[0][j], 0, 0, i->size[0][j] *
size,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS)
- {
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in gegl_buffer_copy: %s",
gegl_cl_errstring(cl_err));
- break;
- }
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring(cl_err));
+ break;
}
}
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
