[gegl] new automated testing system for OpenCL showed a lot of bugs in filters
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] new automated testing system for OpenCL showed a lot of bugs in filters
- Date: Tue, 19 Feb 2013 18:10:02 +0000 (UTC)
commit 5d68c2cfae6bd91a2515017dfba66e340f56f836
Author: Victor Oliveira <victormatheus gmail com>
Date: Tue Feb 19 15:06:28 2013 -0300
new automated testing system for OpenCL showed a lot of bugs in filters
opencl/gaussian-blur.cl | 43 ++++++++++++++++--------------------
opencl/gaussian-blur.cl.h | 43 ++++++++++++++++--------------------
operations/common/gaussian-blur.c | 18 +++------------
operations/common/motion-blur.c | 2 +-
operations/common/oilify.c | 2 +-
operations/common/opacity.c | 2 +-
6 files changed, 45 insertions(+), 65 deletions(-)
---
diff --git a/opencl/gaussian-blur.cl b/opencl/gaussian-blur.cl
index 03b2502..46d288e 100644
--- a/opencl/gaussian-blur.cl
+++ b/opencl/gaussian-blur.cl
@@ -1,20 +1,3 @@
-float4 fir_get_mean_component_1D(const global float4 *buf,
- int offset,
- const int delta_offset,
- constant float *cmatrix,
- const int matrix_length)
-{
- float4 acc = 0.0f;
- int i;
-
- for(i=0; i<matrix_length; i++)
- {
- acc += buf[offset] * cmatrix[i];
- offset += delta_offset;
- }
- return acc;
-}
-
__kernel void fir_ver_blur(const global float4 *src_buf,
const int src_width,
global float4 *dst_buf,
@@ -27,10 +10,16 @@ __kernel void fir_ver_blur(const global float4 *src_buf,
int gid = gidx + gidy * get_global_size(0);
int radius = matrix_length / 2;
- int src_offset = gidx + (gidy - radius + yoff) * src_width;
+ int src_offset = gidx + (gidy + yoff) * src_width;
- dst_buf[gid] = fir_get_mean_component_1D(
- src_buf, src_offset, src_width, cmatrix, matrix_length);
+ float4 v = 0.0f;
+
+ for (int i=-radius; i <= radius; i++)
+ {
+ v += src_buf[src_offset + i * src_width] * cmatrix[i+radius];
+ }
+
+ dst_buf[gid] = v;
}
__kernel void fir_hor_blur(const global float4 *src_buf,
@@ -38,15 +27,21 @@ __kernel void fir_hor_blur(const global float4 *src_buf,
global float4 *dst_buf,
constant float *cmatrix,
const int matrix_length,
- const int yoff)
+ const int xoff)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
int gid = gidx + gidy * get_global_size(0);
int radius = matrix_length / 2;
- int src_offset = gidy * src_width + (gidx - radius + yoff);
+ int src_offset = gidy * src_width + (gidx + xoff);
+
+ float4 v = 0.0f;
+
+ for (int i=-radius; i <= radius; i++)
+ {
+ v += src_buf[src_offset + i] * cmatrix[i+radius];
+ }
- dst_buf[gid] = fir_get_mean_component_1D(
- src_buf, src_offset, 1, cmatrix, matrix_length);
+ dst_buf[gid] = v;
}
diff --git a/opencl/gaussian-blur.cl.h b/opencl/gaussian-blur.cl.h
index acf7f9d..1bf44ac 100644
--- a/opencl/gaussian-blur.cl.h
+++ b/opencl/gaussian-blur.cl.h
@@ -1,21 +1,4 @@
static const char* gaussian_blur_cl_source =
-"float4 fir_get_mean_component_1D(const global float4 *buf, \n"
-" int offset, \n"
-" const int delta_offset, \n"
-" constant float *cmatrix, \n"
-" const int matrix_length) \n"
-"{ \n"
-" float4 acc = 0.0f; \n"
-" int i; \n"
-" \n"
-" for(i=0; i<matrix_length; i++) \n"
-" { \n"
-" acc += buf[offset] * cmatrix[i]; \n"
-" offset += delta_offset; \n"
-" } \n"
-" return acc; \n"
-"} \n"
-" \n"
"__kernel void fir_ver_blur(const global float4 *src_buf, \n"
" const int src_width, \n"
" global float4 *dst_buf, \n"
@@ -28,10 +11,16 @@ static const char* gaussian_blur_cl_source =
" int gid = gidx + gidy * get_global_size(0); \n"
" \n"
" int radius = matrix_length / 2; \n"
-" int src_offset = gidx + (gidy - radius + yoff) * src_width; \n"
+" int src_offset = gidx + (gidy + yoff) * src_width; \n"
" \n"
-" dst_buf[gid] = fir_get_mean_component_1D( \n"
-" src_buf, src_offset, src_width, cmatrix, matrix_length); \n"
+" float4 v = 0.0f; \n"
+" \n"
+" for (int i=-radius; i <= radius; i++) \n"
+" { \n"
+" v += src_buf[src_offset + i * src_width] * cmatrix[i+radius]; \n"
+" } \n"
+" \n"
+" dst_buf[gid] = v; \n"
"} \n"
" \n"
"__kernel void fir_hor_blur(const global float4 *src_buf, \n"
@@ -39,16 +28,22 @@ static const char* gaussian_blur_cl_source =
" global float4 *dst_buf, \n"
" constant float *cmatrix, \n"
" const int matrix_length, \n"
-" const int yoff) \n"
+" const int xoff) \n"
"{ \n"
" int gidx = get_global_id(0); \n"
" int gidy = get_global_id(1); \n"
" int gid = gidx + gidy * get_global_size(0); \n"
" \n"
" int radius = matrix_length / 2; \n"
-" int src_offset = gidy * src_width + (gidx - radius + yoff); \n"
+" int src_offset = gidy * src_width + (gidx + xoff); \n"
+" \n"
+" float4 v = 0.0f; \n"
+" \n"
+" for (int i=-radius; i <= radius; i++) \n"
+" { \n"
+" v += src_buf[src_offset + i] * cmatrix[i+radius]; \n"
+" } \n"
" \n"
-" dst_buf[gid] = fir_get_mean_component_1D( \n"
-" src_buf, src_offset, 1, cmatrix, matrix_length); \n"
+" dst_buf[gid] = v; \n"
"} \n"
;
diff --git a/operations/common/gaussian-blur.c b/operations/common/gaussian-blur.c
index a158342..b5c0221 100644
--- a/operations/common/gaussian-blur.c
+++ b/operations/common/gaussian-blur.c
@@ -451,23 +451,13 @@ cl_gaussian_blur (cl_mem in_tex,
if (!cl_data) return TRUE;
cl_matrix_x = gegl_clCreateBuffer(gegl_cl_get_context(),
- CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
- matrix_length_x * sizeof(cl_float), NULL, &cl_err);
- CL_CHECK;
-
- cl_err = gegl_clEnqueueWriteBuffer(gegl_cl_get_command_queue(), cl_matrix_x,
- CL_TRUE, 0, matrix_length_x * sizeof(cl_float), dmatrix_x,
- 0, NULL, NULL);
+ CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
+ matrix_length_x * sizeof(cl_float), dmatrix_x, &cl_err);
CL_CHECK;
cl_matrix_y = gegl_clCreateBuffer(gegl_cl_get_context(),
- CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
- matrix_length_y * sizeof(cl_float), NULL, &cl_err);
- CL_CHECK;
-
- cl_err = gegl_clEnqueueWriteBuffer(gegl_cl_get_command_queue(), cl_matrix_y,
- CL_TRUE, 0, matrix_length_y * sizeof(cl_float), dmatrix_y,
- 0, NULL, NULL);
+ CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
+ matrix_length_y * sizeof(cl_float), dmatrix_y, &cl_err);
CL_CHECK;
{
diff --git a/operations/common/motion-blur.c b/operations/common/motion-blur.c
index b6ac612..9428b14 100644
--- a/operations/common/motion-blur.c
+++ b/operations/common/motion-blur.c
@@ -143,7 +143,7 @@ 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_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);
+ op_area->left, op_area->right, op_area->top,
op_area->bottom, GEGL_ABYSS_CLAMP);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
diff --git a/operations/common/oilify.c b/operations/common/oilify.c
index 22d9074..aee83fc 100644
--- a/operations/common/oilify.c
+++ b/operations/common/oilify.c
@@ -346,7 +346,7 @@ 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_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ,
- o->mask_radius, o->mask_radius, o->mask_radius,
o->mask_radius, GEGL_ABYSS_NONE);
+ 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;
diff --git a/operations/common/opacity.c b/operations/common/opacity.c
index 3c454ac..8b850ff 100644
--- a/operations/common/opacity.c
+++ b/operations/common/opacity.c
@@ -219,7 +219,7 @@ cl_process (GeglOperation *op,
value = GEGL_CHANT_PROPERTIES (op)->value;
- kernel = (GEGL_CHANT_PROPERTIES (op)->chant_data != NULL);
+ kernel = (GEGL_CHANT_PROPERTIES (op)->chant_data == NULL);
cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 0, sizeof(cl_mem), (void*)&in_tex);
CL_CHECK;
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
