[gegl/gsoc2011-opencl: 16/46] added clFinish to ensure synchronization

[Victor Matheus de Araujo Oliveira <vmaolive src gnome org>]

commit 66f996598e439c4572755c464f1a07522da694be
Author: Victor Oliveira <victormatheus gmail com>
Date:   Thu May 19 13:18:59 2011 -0300

    added clFinish to ensure synchronization

 gegl/opencl/gegl-cl-init.c  |    1 +
 gegl/opencl/gegl-cl-init.h  |    2 ++
 gegl/opencl/gegl-cl-types.h |    1 +
 operations/common/over.c    |   20 +++++++++++---------
 4 files changed, 15 insertions(+), 9 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-init.c b/gegl/opencl/gegl-cl-init.c
index 6ada529..36fc48b 100644
--- a/gegl/opencl/gegl-cl-init.c
+++ b/gegl/opencl/gegl-cl-init.c
@@ -126,6 +126,7 @@ gegl_cl_init (GError **error)
       CL_LOAD_FUNCTION (clEnqueueWriteBuffer)
       CL_LOAD_FUNCTION (clEnqueueReadBuffer)
       CL_LOAD_FUNCTION (clEnqueueNDRangeKernel)
+      CL_LOAD_FUNCTION (clFinish)
 
       CL_LOAD_FUNCTION (clReleaseKernel)
       CL_LOAD_FUNCTION (clReleaseProgram)
diff --git a/gegl/opencl/gegl-cl-init.h b/gegl/opencl/gegl-cl-init.h
index 347a0dc..440bf75 100644
--- a/gegl/opencl/gegl-cl-init.h
+++ b/gegl/opencl/gegl-cl-init.h
@@ -29,6 +29,7 @@ t_clCreateBuffer            gegl_clCreateBuffer            = NULL;
 t_clEnqueueWriteBuffer      gegl_clEnqueueWriteBuffer      = NULL;
 t_clEnqueueReadBuffer       gegl_clEnqueueReadBuffer       = NULL;
 t_clEnqueueNDRangeKernel    gegl_clEnqueueNDRangeKernel    = NULL;
+t_clFinish                  gegl_clFinish                  = NULL;
 
 t_clReleaseKernel       gegl_clReleaseKernel       = NULL;
 t_clReleaseProgram      gegl_clReleaseProgram      = NULL;
@@ -56,6 +57,7 @@ extern t_clCreateBuffer            gegl_clCreateBuffer;
 extern t_clEnqueueWriteBuffer      gegl_clEnqueueWriteBuffer;
 extern t_clEnqueueReadBuffer       gegl_clEnqueueReadBuffer;
 extern t_clEnqueueNDRangeKernel    gegl_clEnqueueNDRangeKernel;
+extern t_clFinish                  gegl_clFinish;
 
 extern t_clReleaseKernel       gegl_clReleaseKernel;
 extern t_clReleaseProgram      gegl_clReleaseProgram;
diff --git a/gegl/opencl/gegl-cl-types.h b/gegl/opencl/gegl-cl-types.h
index 16d0580..ad04f6e 100644
--- a/gegl/opencl/gegl-cl-types.h
+++ b/gegl/opencl/gegl-cl-types.h
@@ -26,6 +26,7 @@ typedef cl_mem            (*t_clCreateBuffer           ) (cl_context, cl_mem_fla
 typedef cl_int            (*t_clEnqueueWriteBuffer     ) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, cl_event *);
 typedef cl_int            (*t_clEnqueueReadBuffer      ) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, void *, cl_uint, const cl_event *, cl_event *);
 typedef cl_int            (*t_clEnqueueNDRangeKernel   ) (cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *);
+typedef cl_int            (*t_clFinish                 ) (cl_command_queue);
 
 typedef cl_int (*t_clReleaseKernel      ) (cl_kernel);
 typedef cl_int (*t_clReleaseProgram     ) (cl_program);
diff --git a/operations/common/over.c b/operations/common/over.c
index 0422fb2..0fb9167 100644
--- a/operations/common/over.c
+++ b/operations/common/over.c
@@ -107,10 +107,10 @@ cl_process (GeglOperation        *op,
   /* -- Configuration -- */
 
   CL_SAFE_CALL( errcode = gegl_clGetPlatformIDs (1, &platform, NULL) );
-  CL_SAFE_CALL( errcode = gegl_clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, NULL) );
   CL_SAFE_CALL( ctx     = gegl_clCreateContext(0, 1, &device, NULL, NULL, &errcode) );
   CL_SAFE_CALL( cq      = gegl_clCreateCommandQueue(ctx, device, 0, &errcode) );
-  CL_SAFE_CALL( program = gegl_clCreateProgramWithSource(ctx, 16, (const char **)&kernel_source, NULL, &errcode) );
+  CL_SAFE_CALL( program = gegl_clCreateProgramWithSource(ctx, 11, (const char **)&kernel_source, NULL, &errcode) );
   errcode = gegl_clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
   if (errcode != CL_SUCCESS)
     {
@@ -120,9 +120,9 @@ cl_process (GeglOperation        *op,
       return FALSE;
     }
 
-  CL_SAFE_CALL( d_in  = gegl_clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sizeof(cl_float) * _n_pixels, NULL, &errcode) );
-  CL_SAFE_CALL( d_aux = gegl_clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sizeof(cl_float) * _n_pixels, NULL, &errcode) );
-  CL_SAFE_CALL( d_out = gegl_clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, sizeof(cl_float) * _n_pixels, NULL, &errcode) );
+  CL_SAFE_CALL( d_in  = gegl_clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sizeof(cl_float4) * _n_pixels, NULL, &errcode) );
+  CL_SAFE_CALL( d_aux = gegl_clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sizeof(cl_float4) * _n_pixels, NULL, &errcode) );
+  CL_SAFE_CALL( d_out = gegl_clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, sizeof(cl_float4) * _n_pixels, NULL, &errcode) );
 
   CL_SAFE_CALL( kernel  = gegl_clCreateKernel(program, "kernel_over", &errcode) );
   CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&d_in) );
@@ -133,14 +133,16 @@ cl_process (GeglOperation        *op,
   /* -- Running -- */
 
   CL_SAFE_CALL( errcode = gegl_clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &local_worksize, NULL) );
-  global_worksize = ((_n_pixels+local_worksize-1) / local_worksize) * local_worksize;
+  global_worksize = MAX( (((_n_pixels+local_worksize-1) / local_worksize) * local_worksize) / 4, local_worksize );
 
-  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_in,  CL_FALSE, 0, sizeof(cl_float) * _n_pixels, in,  0, NULL, NULL) );
-  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_aux, CL_FALSE, 0, sizeof(cl_float) * _n_pixels, aux, 0, NULL, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_in,  CL_FALSE, 0, sizeof(cl_float4) * _n_pixels, in,  0, NULL, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clEnqueueWriteBuffer(cq, d_aux, CL_FALSE, 0, sizeof(cl_float4) * _n_pixels, aux, 0, NULL, NULL) );
 
+  CL_SAFE_CALL( errcode = gegl_clFinish(cq) );
   CL_SAFE_CALL( errcode = gegl_clEnqueueNDRangeKernel(cq, kernel, 1, NULL, &global_worksize, &local_worksize, 0, NULL, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clFinish(cq) );
 
-  CL_SAFE_CALL( errcode = gegl_clEnqueueReadBuffer(cq, d_out, CL_TRUE, 0, sizeof(cl_float) * _n_pixels, out, 0, NULL, NULL) );
+  CL_SAFE_CALL( errcode = gegl_clEnqueueReadBuffer(cq, d_out, CL_TRUE, 0, sizeof(cl_float4) * _n_pixels, out, 0, NULL, NULL) );
 
   CL_SAFE_CALL( errcode = gegl_clReleaseProgram(program) );
   CL_SAFE_CALL( errcode = gegl_clReleaseCommandQueue(cq) );