[gegl/gsoc2011-opencl: 12/14] OpenCL is running, but still there are some bugs

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

commit c1ea1ee0c7c39e673c2b00f0146ada25f02e517c
Author: Victor Oliveira <victormatheus gmail com>
Date:   Tue May 17 17:28:42 2011 -0300

    OpenCL is running, but still there are some bugs

 operations/common/over.c |   97 +++++++++++++++++++++------------------------
 1 files changed, 45 insertions(+), 52 deletions(-)
---
diff --git a/operations/common/over.c b/operations/common/over.c
index 633daad..0422fb2 100644
--- a/operations/common/over.c
+++ b/operations/common/over.c
@@ -39,8 +39,6 @@ static void prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "input", format);
   gegl_operation_set_format (operation, "aux", format);
   gegl_operation_set_format (operation, "output", format);
-
-  gegl_cl_init(NULL);
 }
 
 #define CL_SAFE_CALL(func)                                          \
@@ -65,38 +63,37 @@ cl_process (GeglOperation        *op,
   gfloat * GEGL_ALIGNED out = out_buf;
   cl_uint _n_pixels = (cl_uint)n_pixels;
 
-  const char kernel_source[] =
-  "__kernel void kernel_over (__global float4* in,                \n"
-  "                           __global float4* aux,               \n"
-  "                           __global float4* out,               \n"
-  "                           uint n_pixels)                      \n"
-  "{                                                              \n"
-  "    size_t gid = get_global_id(0);                             \n"
-  "    if (gid < n_pixels)                                        \n"
-  "    {                                                          \n"
-  "        float4 _in  = in[gid];                                 \n"
-  "        float4 _aux = aux[gid];                                \n"
-  "        out[gid] = (float4)(_in.x + _aux.x * (1.0f - _in.w),   \n"
-  "                            _in.x + _aux.x * (1.0f - _in.w),   \n"
-  "                            _in.x + _aux.x * (1.0f - _in.w),   \n"
-  "                            _in.w + _aux.w - _in.w * _aux.w);  \n"
-  "    }                                                          \n"
-  "}                                                              \n";
-
-  size_t kernel_size = sizeof(kernel_source);
+  const char* kernel_source[] =
+  {
+  "__kernel void kernel_over (__global float4* in,                \n",
+  "                           __global float4* aux,               \n",
+  "                           __global float4* out,               \n",
+  "                           uint n_pixels)                      \n",
+  "{                                                              \n",
+  "    size_t gid = get_global_id(0);                             \n",
+  "    if (gid < n_pixels)                                        \n",
+  "    {                                                          \n",
+  "        float4 _in  = in[gid];                                 \n",
+  "        float4 _aux = aux[gid];                                \n",
+  "        out[gid] = (float4)(_in.x + _aux.x * (1.0f - _in.w),   \n",
+  "                            _in.x + _aux.x * (1.0f - _in.w),   \n",
+  "                            _in.x + _aux.x * (1.0f - _in.w),   \n",
+  "                            _in.w + _aux.w - _in.w * _aux.w);  \n",
+  "    }                                                          \n",
+  "}                                                              \n",
+  };
+
+  size_t len;
+  char buffer[16384];
 
   cl_int errcode;
-
   cl_platform_id platform;
   cl_device_id device;
-
-  cl_context_properties ctx_properties[3];
-
   cl_context ctx;
   cl_command_queue cq;
-
   cl_program program;
   cl_kernel kernel;
+
   size_t local_worksize;
   size_t global_worksize;
 
@@ -110,37 +107,33 @@ cl_process (GeglOperation        *op,
   /* -- Configuration -- */
 
   CL_SAFE_CALL( errcode = gegl_clGetPlatformIDs (1, &platform, NULL) );
-  //CL_SAFE_CALL( errcode = gegl_clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, NULL) );
   CL_SAFE_CALL( errcode = gegl_clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 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) );
+  errcode = gegl_clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+  if (errcode != CL_SUCCESS)
+    {
+      CL_SAFE_CALL( errcode = gegl_clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, NULL) );
+      g_warning("OpenCL Build Error in Line %u in file %s\nError:%s\n%s",
+                __LINE__, __FILE__, gegl_cl_errstring(errcode), buffer);
+      return FALSE;
+    }
 
-  /*
-  ctx_properties[0] = CL_CONTEXT_PLATFORM;
-  ctx_properties[1] = (cl_context_properties) platform;
-  ctx_properties[2] = 0;
-  */
-
-  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, 1, (const char **)&kernel_source, &kernel_size, &errcode) );
-  CL_SAFE_CALL( errcode = gegl_clBuildProgram(program, 0, NULL, NULL, NULL, NULL) );
-
-/*
   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( kernel = gegl_clCreateKernel(program, "kernel_over", &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) );
   CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&d_aux) );
   CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&d_out) );
   CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 3, sizeof(cl_uint), (void*)&_n_pixels) );
-*/
 
   /* -- 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;
+  global_worksize = ((_n_pixels+local_worksize-1) / local_worksize) * 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) );
@@ -149,14 +142,14 @@ cl_process (GeglOperation        *op,
 
   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_clReleaseKernel(kernel) );
   CL_SAFE_CALL( errcode = gegl_clReleaseProgram(program) );
   CL_SAFE_CALL( errcode = gegl_clReleaseCommandQueue(cq) );
   CL_SAFE_CALL( errcode = gegl_clReleaseContext(ctx) );
+
+  CL_SAFE_CALL( errcode = gegl_clReleaseKernel(kernel) );
   CL_SAFE_CALL( errcode = gegl_clReleaseMemObject(d_in) );
   CL_SAFE_CALL( errcode = gegl_clReleaseMemObject(d_aux) );
   CL_SAFE_CALL( errcode = gegl_clReleaseMemObject(d_out));
-*/
 
   return TRUE;
 }
@@ -200,14 +193,14 @@ process (GeglOperation        *op,
           glong                n_pixels,
           const GeglRectangle *roi)
 {
-  if (gegl_cl_is_accelerated())
-    {
-      return cl_process(op, in_buf, aux_buf, out_buf, n_pixels, roi);
-    }
+  if (gegl_cl_init(NULL))
+  {
+    return cl_process(op, in_buf, aux_buf, out_buf, n_pixels, roi);
+  }
   else
-    {
-      return _process(op, in_buf, aux_buf, out_buf, n_pixels, roi);
-    }
+  {
+    return _process(op, in_buf, aux_buf, out_buf, n_pixels, roi);
+  }
 }
 
 /* Fast paths */