[gegl] opencl: Use the right roi for pre_edgelaplace, still not enough

[Téo Mazars <teom src gnome org>]

commit b0009d05349a997dd58a36f22fe99fd4700a29f2
Author: Téo Mazars <teo mazars ensimag fr>
Date:   Tue Oct 15 12:25:29 2013 +0200

    opencl: Use the right roi for pre_edgelaplace, still not enough
    
    but getting closer

 opencl/edge-laplace.cl           |    4 ++--
 opencl/edge-laplace.cl.h         |    4 ++--
 operations/common/edge-laplace.c |   22 +++++++++++++---------
 3 files changed, 17 insertions(+), 13 deletions(-)
---
diff --git a/opencl/edge-laplace.cl b/opencl/edge-laplace.cl
index d038836..80d3767 100644
--- a/opencl/edge-laplace.cl
+++ b/opencl/edge-laplace.cl
@@ -46,10 +46,10 @@ kernel void pre_edgelaplace (global float4 *in,
     int gidx = get_global_id(0);
     int gidy = get_global_id(1);
 
-    int src_width  = get_global_size(0) + LAPLACE_RADIUS * 2;
+    int src_width  = get_global_size(0) + LAPLACE_RADIUS;
     int src_height = get_global_size(1);
 
-    int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;
+    int i = gidx + LAPLACE_RADIUS - 1, j = gidy + LAPLACE_RADIUS - 1;
     int gid1d = i + j * src_width;
 
     float pix_fl[4] = {
diff --git a/opencl/edge-laplace.cl.h b/opencl/edge-laplace.cl.h
index 5682583..f0f7a0c 100644
--- a/opencl/edge-laplace.cl.h
+++ b/opencl/edge-laplace.cl.h
@@ -47,10 +47,10 @@ static const char* edge_laplace_cl_source =
 "    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_width  = get_global_size(0) + LAPLACE_RADIUS;                     \n"
 "    int src_height = get_global_size(1);                                      \n"
 "                                                                              \n"
-"    int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;                 \n"
+"    int i = gidx + LAPLACE_RADIUS - 1, j = gidy + LAPLACE_RADIUS - 1;         \n"
 "    int gid1d = i + j * src_width;                                            \n"
 "                                                                              \n"
 "    float pix_fl[4] = {                                                       \n"
diff --git a/operations/common/edge-laplace.c b/operations/common/edge-laplace.c
index f73f8b8..e627cc4 100644
--- a/operations/common/edge-laplace.c
+++ b/operations/common/edge-laplace.c
@@ -286,7 +286,8 @@ cl_edge_laplace (cl_mem                in_tex,
                  gint                  radius)
 {
   cl_int cl_err = 0;
-  size_t global_ws[2];
+  size_t global_ws_in[2];
+  size_t global_ws_aux[2];
 
   if (!cl_data)
     {
@@ -295,8 +296,11 @@ cl_edge_laplace (cl_mem                in_tex,
     }
   if (!cl_data) return TRUE;
 
-  global_ws[0] = roi->width;
-  global_ws[1] = roi->height;
+  global_ws_in[0] = roi->width  + LAPLACE_RADIUS;
+  global_ws_in[1] = roi->height + LAPLACE_RADIUS;
+
+  global_ws_aux[0] = roi->width;
+  global_ws_aux[1] = roi->height;
 
   cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 0, sizeof (cl_mem),
                                 (void*) &in_tex);
@@ -307,7 +311,7 @@ cl_edge_laplace (cl_mem                in_tex,
 
   cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                         cl_data->kernel[0], 2,
-                                        NULL, global_ws, NULL,
+                                        NULL, global_ws_in, NULL,
                                         0, NULL, NULL);
   CL_CHECK;
 
@@ -320,7 +324,7 @@ cl_edge_laplace (cl_mem                in_tex,
 
   cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                         cl_data->kernel[1], 2,
-                                        NULL, global_ws, NULL,
+                                        NULL, global_ws_aux, NULL,
                                         0, NULL, NULL);
   CL_CHECK;
 
@@ -363,10 +367,10 @@ cl_process (GeglOperation       *operation,
                                              result,
                                              in_format,
                                              GEGL_CL_BUFFER_AUX,
-                                             op_area->left,
-                                             op_area->right,
-                                             op_area->top,
-                                             op_area->bottom,
+                                             op_area->left   - 1,
+                                             op_area->right  - 1,
+                                             op_area->top    - 1,
+                                             op_area->bottom - 1,
                                              GEGL_ABYSS_NONE);
 
   while (gegl_buffer_cl_iterator_next (i, &err))