[gegl] operations: Clean up OpenCL calls

[Daniel Sabo <daniels src gnome org>]

commit b9c357b3d3aeed701eb80ac26942111ab082fceb
Author: Daniel Sabo <DanielSabo gmail com>
Date:   Mon Oct 21 19:41:54 2013 -0700

    operations: Clean up OpenCL calls

 operations/common/box-blur.c        |   64 ++++++++++++++++--------------
 operations/common/edge-laplace.c    |   33 ++++++++--------
 operations/common/gaussian-blur.c   |   72 ++++++++++++++++-------------------
 operations/common/noise-reduction.c |   30 ++++++++-------
 4 files changed, 100 insertions(+), 99 deletions(-)
---
diff --git a/operations/common/box-blur.c b/operations/common/box-blur.c
index de21183..d44d6f4 100644
--- a/operations/common/box-blur.c
+++ b/operations/common/box-blur.c
@@ -200,7 +200,9 @@ cl_box_blur (cl_mem                in_tex,
       const char *kernel_name[] = {"kernel_blur_hor", "kernel_blur_ver", NULL};
       cl_data = gegl_cl_compile_and_build (box_blur_cl_source, kernel_name);
     }
-  if (!cl_data) return TRUE;
+
+  if (!cl_data)
+    return TRUE;
 
   local_ws_hor[0] = 1;
   local_ws_hor[1] = 256;
@@ -212,31 +214,31 @@ cl_box_blur (cl_mem                in_tex,
   global_ws_ver[0] = roi->height;
   global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1];
 
-  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_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int),   (void*)&roi->width);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int),   (void*)&radius);
+
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
+                                    sizeof(cl_mem), (void*)&in_tex,
+                                    sizeof(cl_mem), (void*)&aux_tex,
+                                    sizeof(cl_int), (void*)&roi->width,
+                                    sizeof(cl_int), (void*)&radius,
+                                    NULL);
   CL_CHECK;
 
-  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                         cl_data->kernel[0], 2,
                                         NULL, global_ws_hor, local_ws_hor,
                                         0, NULL, NULL);
   CL_CHECK;
 
-  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_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_int),   (void*)&roi->width);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_int),   (void*)&radius);
+
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
+                                    sizeof(cl_mem), (void*)&aux_tex,
+                                    sizeof(cl_mem), (void*)&out_tex,
+                                    sizeof(cl_int), (void*)&roi->width,
+                                    sizeof(cl_int), (void*)&radius,
+                                    NULL);
   CL_CHECK;
 
-  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                         cl_data->kernel[1], 2,
                                         NULL, global_ws_ver, local_ws_ver,
                                         0, NULL, NULL);
@@ -257,7 +259,7 @@ 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 err = 0;
 
   GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
   GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
@@ -286,21 +288,23 @@ cl_process (GeglOperation       *operation,
                                               op_area->top,
                                               op_area->bottom);
 
-  while (gegl_buffer_cl_iterator_next (i, &err))
+  while (gegl_buffer_cl_iterator_next (i, &err) && !err)
     {
-      if (err) 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;
+      err = cl_box_blur (i->tex[read],
+                         i->tex[aux],
+                         i->tex[0],
+                         i->size[0],
+                         &i->roi[0],
+                         ceil (o->radius));
+
+      if (err)
+        {
+          gegl_buffer_cl_iterator_stop (i);
+          break;
+        }
     }
 
-  return TRUE;
+  return !err;
 }
 
 static gboolean
diff --git a/operations/common/edge-laplace.c b/operations/common/edge-laplace.c
index 4eed257..5e79ba6 100644
--- a/operations/common/edge-laplace.c
+++ b/operations/common/edge-laplace.c
@@ -302,11 +302,10 @@ cl_edge_laplace (cl_mem                in_tex,
   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);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 1, sizeof (cl_mem),
-                                (void*) &aux_tex);
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
+                                    sizeof (cl_mem), &in_tex,
+                                    sizeof (cl_mem), &aux_tex,
+                                    NULL);
   CL_CHECK;
 
   cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
@@ -315,11 +314,11 @@ cl_edge_laplace (cl_mem                in_tex,
                                         0, NULL, NULL);
   CL_CHECK;
 
-  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_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
+                                    sizeof (cl_mem), &aux_tex,
+                                    sizeof (cl_mem), &out_tex,
+                                    NULL);
   CL_CHECK;
 
   cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
@@ -342,7 +341,7 @@ 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 err = 0;
 
   GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
 
@@ -370,10 +369,8 @@ cl_process (GeglOperation       *operation,
                                               op_area->top    - 1,
                                               op_area->bottom - 1);
 
-  while (gegl_buffer_cl_iterator_next (i, &err))
+  while (gegl_buffer_cl_iterator_next (i, &err) && !err)
     {
-      if (err) return FALSE;
-
       err = cl_edge_laplace (i->tex[read],
                              i->tex[aux],
                              i->tex[0],
@@ -381,10 +378,14 @@ cl_process (GeglOperation       *operation,
                              &i->roi[0],
                              LAPLACE_RADIUS);
 
-      if (err) return FALSE;
+      if (err)
+        {
+          gegl_buffer_cl_iterator_stop (i);
+          break;
+        }
     }
 
-  return TRUE;
+  return !err;
 }
 
 static void
diff --git a/operations/common/gaussian-blur.c b/operations/common/gaussian-blur.c
index f42beec..007cb42 100644
--- a/operations/common/gaussian-blur.c
+++ b/operations/common/gaussian-blur.c
@@ -463,67 +463,61 @@ cl_gaussian_blur (cl_mem                in_tex,
   if (!cl_data)
     return TRUE;
 
-  cl_matrix_x = gegl_clCreateBuffer(gegl_cl_get_context(),
-                                    CL_MEM_COPY_HOST_PTR | CL_MEM_READ_ONLY,
-                                    matrix_length_x * sizeof(cl_float), dmatrix_x, &cl_err);
+  cl_matrix_x = gegl_clCreateBuffer (gegl_cl_get_context(),
+                                     CL_MEM_COPY_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_COPY_HOST_PTR | CL_MEM_READ_ONLY,
-                                    matrix_length_y * sizeof(cl_float), dmatrix_y, &cl_err);
+  cl_matrix_y = gegl_clCreateBuffer (gegl_cl_get_context(),
+                                     CL_MEM_COPY_HOST_PTR | CL_MEM_READ_ONLY,
+                                     matrix_length_y * sizeof(cl_float), dmatrix_y, &cl_err);
   CL_CHECK;
 
   global_ws[0] = aux_rect->width;
   global_ws[1] = aux_rect->height;
 
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&in_tex);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_int), (void*)&src_rect->width);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem), (void*)&aux_tex);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_mem), (void*)&cl_matrix_x);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 4, sizeof(cl_int), (void*)&matrix_length_x);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 5, sizeof(cl_int), (void*)&xoff);
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
+                                    sizeof(cl_mem), (void*)&in_tex,
+                                    sizeof(cl_int), (void*)&src_rect->width,
+                                    sizeof(cl_mem), (void*)&aux_tex,
+                                    sizeof(cl_mem), (void*)&cl_matrix_x,
+                                    sizeof(cl_int), (void*)&matrix_length_x,
+                                    sizeof(cl_int), (void*)&xoff,
+                                    NULL);
   CL_CHECK;
 
-  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
-                                       cl_data->kernel[1], 2,
-                                       NULL, global_ws, NULL,
-                                       0, NULL, NULL);
+  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+                                        cl_data->kernel[1], 2,
+                                        NULL, global_ws, NULL,
+                                        0, NULL, NULL);
   CL_CHECK;
 
 
   global_ws[0] = roi->width;
   global_ws[1] = roi->height;
 
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&aux_tex);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int), (void*)&aux_rect->width);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_mem), (void*)&cl_matrix_y);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&matrix_length_y);
-  CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&yoff);
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
+                                    sizeof(cl_mem), (void*)&aux_tex,
+                                    sizeof(cl_int), (void*)&aux_rect->width,
+                                    sizeof(cl_mem), (void*)&out_tex,
+                                    sizeof(cl_mem), (void*)&cl_matrix_y,
+                                    sizeof(cl_int), (void*)&matrix_length_y,
+                                    sizeof(cl_int), (void*)&yoff,
+                                    NULL);
   CL_CHECK;
 
-  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
-                                       cl_data->kernel[0], 2,
-                                       NULL, global_ws, NULL,
-                                       0, NULL, NULL);
+  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+                                        cl_data->kernel[0], 2,
+                                        NULL, global_ws, NULL,
+                                        0, NULL, NULL);
   CL_CHECK;
 
-  cl_err = gegl_clFinish(gegl_cl_get_command_queue ());
+  cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
   CL_CHECK;
 
-  cl_err = gegl_clReleaseMemObject(cl_matrix_x);
+  cl_err = gegl_clReleaseMemObject (cl_matrix_x);
   CL_CHECK;
-  cl_err = gegl_clReleaseMemObject(cl_matrix_y);
+  cl_err = gegl_clReleaseMemObject (cl_matrix_y);
   CL_CHECK;
 
   return FALSE;
diff --git a/operations/common/noise-reduction.c b/operations/common/noise-reduction.c
index e12655d..8e66e2d 100644
--- a/operations/common/noise-reduction.c
+++ b/operations/common/noise-reduction.c
@@ -256,7 +256,7 @@ 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 err = 0;
 
   GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
   GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
@@ -285,22 +285,24 @@ cl_process (GeglOperation       *operation,
                                               op_area->top,
                                               op_area->bottom);
 
-  while (gegl_buffer_cl_iterator_next (i, &err))
+  while (gegl_buffer_cl_iterator_next (i, &err) && !err)
     {
-      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;
+      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)
+        {
+          gegl_buffer_cl_iterator_stop (i);
+          break;
+        }
     }
 
-  return TRUE;
+  return !err;
 }
 
 #define INPLACE 1