[gegl/gsoc2011-opencl: 16/21] GeglClTexture uses kernel cache and minor changes

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

commit d8fc1a8a4a2b45a1bda5ecb0f4d969a52cdc4923
Author: Victor Oliveira <victormatheus gmail com>
Date:   Thu Aug 18 13:32:52 2011 -0300

    GeglClTexture uses kernel cache and minor changes

 gegl/opencl/gegl-cl-texture.c |  129 +++++++++++++++++++----------------------
 1 files changed, 59 insertions(+), 70 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-texture.c b/gegl/opencl/gegl-cl-texture.c
index d73d74f..6308d82 100644
--- a/gegl/opencl/gegl-cl-texture.c
+++ b/gegl/opencl/gegl-cl-texture.c
@@ -103,15 +103,18 @@ gegl_cl_texture_copy (const GeglClTexture  *src,
   if (_src_rect)
     src_rect = *_src_rect;
 
-  const size_t src_origin[3] = {src_rect.x, src_rect.y, 0};
-  const size_t dst_origin[3] = {dst_x, dst_y, 0};
-  const size_t region[3] = {src_rect.width, src_rect.height, 1};
-
-  gegl_clEnqueueCopyImage (gegl_cl_get_command_queue(),
-                           src->data, dst->data,
-                           src_origin, dst_origin, region,
-                           0, NULL, NULL);
-  gegl_clFinish(gegl_cl_get_command_queue());
+  {
+    const size_t src_origin[3] = {src_rect.x, src_rect.y, 0};
+    const size_t dst_origin[3] = {dst_x, dst_y, 0};
+    const size_t region[3] = {src_rect.width, src_rect.height, 1};
+
+    gegl_clEnqueueCopyImage (gegl_cl_get_command_queue(),
+                             src->data, dst->data,
+                             src_origin, dst_origin, region,
+                             0, NULL, NULL);
+    gegl_clFinish(gegl_cl_get_command_queue());
+  }
+
 }
 
 void
@@ -119,71 +122,57 @@ gegl_cl_texture_fill (const GeglClTexture *texture,
                       GeglRectangle       *_rect,
                       gfloat               color[4])
 {
-  cl_int errcode;
-  char buffer[10000];
-  GeglRectangle rect = {0, 0, texture->width, texture->height};
-
-  cl_float cl_color[4];
-
-  size_t global_worksize[2];
-
-  const char* kernel_source[] =
-  {
-  "__kernel void kernel_clear (__write_only image2d_t img,  \n",
-  "                            int offset_x, int offset_y,  \n",
-  "                            float4 c)                    \n",
-  "{                                                        \n",
-  "  int2 offset = (int2)(offset_x, offset_y);              \n",
-  "  int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n",
-  "  float4 v = c;                                          \n",
-  "  write_imagef(img,gid+offset,v);                        \n",
-  "}                                                        \n",
-  NULL,
-  };
-
-  cl_program program;
-  cl_kernel kernel;
 
+  GeglRectangle rect = {0, 0, texture->width, texture->height};
   if (_rect)
     rect = *_rect;
 
-  cl_color[0] = color[0];
-  cl_color[1] = color[1];
-  cl_color[2] = color[2];
-  cl_color[3] = color[3];
-
-  global_worksize[0] = rect.width;
-  global_worksize[1] = rect.height;
-
-  CL_SAFE_CALL( program = gegl_clCreateProgramWithSource(gegl_cl_get_context(),
-                                                         gegl_cl_count_lines(kernel_source),
-                                                         (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,
-                                                         gegl_cl_get_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);
-    }
-
-  CL_SAFE_CALL( kernel = gegl_clCreateKernel(program, "kernel_clear", &errcode) );
-  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 0, sizeof(cl_mem),    (void*)&texture->data) );
-  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 1, sizeof(cl_int),    (void*)&rect.x) );
-  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 2, sizeof(cl_int),    (void*)&rect.y) );
-  CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 3, sizeof(cl_float4), (void*)&cl_color) );
-
-  CL_SAFE_CALL( errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), kernel, 2,
-                                                      NULL, global_worksize, NULL,
-                                                      0, NULL, NULL) );
-  CL_SAFE_CALL( errcode = gegl_clFinish(gegl_cl_get_command_queue()) );
-
-  CL_SAFE_CALL( errcode = gegl_clReleaseProgram(program) );
-  CL_SAFE_CALL( errcode = gegl_clReleaseKernel(kernel) );
+  {
+    cl_int errcode;
+    char buffer[10000];
+
+    cl_float cl_color[4];
+
+    size_t global_worksize[2];
+
+    const char* kernel_source =
+    "__kernel void kernel_clear (__write_only image2d_t img,  "
+    "                            int offset_x, int offset_y,  "
+    "                            float4 c)                    "
+    "{                                                        "
+    "  int2 offset = (int2)(offset_x, offset_y);              "
+    "  int2 gid = (int2)(get_global_id(0), get_global_id(1)); "
+    "  float4 v = c;                                          "
+    "  write_imagef(img,gid+offset,v);                        "
+    "}                                                        ";
+
+    gegl_cl_run_data *cl_data;
+    const char *kernel_name[] = {"kernel_clear", NULL};
+
+    cl_color[0] = color[0];
+    cl_color[1] = color[1];
+    cl_color[2] = color[2];
+    cl_color[3] = color[3];
+
+    global_worksize[0] = rect.width;
+    global_worksize[1] = rect.height;
+
+    cl_data = gegl_cl_compile_and_build (kernel_source,
+                                         gegl_cl_get_context (),
+                                         gegl_cl_get_device (),
+                                         kernel_name);
+
+    CL_SAFE_CALL( errcode = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),    (void*)&texture->data) );
+    CL_SAFE_CALL( errcode = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int),    (void*)&rect.x) );
+    CL_SAFE_CALL( errcode = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int),    (void*)&rect.y) );
+    CL_SAFE_CALL( errcode = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float4), (void*)&cl_color) );
+
+    CL_SAFE_CALL( errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
+                                                        cl_data->kernel[0], 2,
+                                                        NULL, global_worksize, NULL,
+                                                        0, NULL, NULL) );
+    CL_SAFE_CALL( errcode = gegl_clFinish(gegl_cl_get_command_queue()) );
+  }
 }
 
 void