[gegl/gsoc2011-opencl: 7/10] OpenCL Texture Class

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

commit 34dc3bfee0daf8d3945c6694928c8606b9e71ff8
Author: Victor Oliveira <victormatheus gmail com>
Date:   Wed Jun 1 14:25:16 2011 -0300

    OpenCL Texture Class

 gegl/opencl/Makefile.am       |    4 +-
 gegl/opencl/gegl-cl-texture.c |  152 +++++++++++++++++++++++++++++++++++++++++
 gegl/opencl/gegl-cl-texture.h |   46 ++++++++++++
 3 files changed, 201 insertions(+), 1 deletions(-)
---
diff --git a/gegl/opencl/Makefile.am b/gegl/opencl/Makefile.am
index fc122dc..90df443 100644
--- a/gegl/opencl/Makefile.am
+++ b/gegl/opencl/Makefile.am
@@ -19,8 +19,10 @@ noinst_LTLIBRARIES = libcl.la
 #libcl_public_HEADERS = gegl-cl-init.h
 
 libcl_la_SOURCES = \
+	gegl-cl-types.h \
 	gegl-cl-init.c \
 	gegl-cl-init.h \
-	gegl-cl-types.h
+	gegl-cl-texture.c \
+	gegl-cl-texture.h
 
 #libcl_la_SOURCES = $(libcl_sources) $(libcl_public_HEADERS)
diff --git a/gegl/opencl/gegl-cl-texture.c b/gegl/opencl/gegl-cl-texture.c
new file mode 100644
index 0000000..e3ebea9
--- /dev/null
+++ b/gegl/opencl/gegl-cl-texture.c
@@ -0,0 +1,152 @@
+#include "gegl.h"
+#include "gegl-cl-types.h"
+#include "gegl-cl-init.h"
+
+#include "gegl-cl-texture.h"
+
+GeglClTexture *
+gegl_cl_texture_new (const gint width, const gint height)
+{
+  cl_int errcode;
+
+  GeglClTexture *texture = g_new (GeglClTexture, 1);
+  texture->width = width;
+  texture->height = height;
+  texture->format.image_channel_order     = CL_RGBA;
+  texture->format.image_channel_data_type = CL_FLOAT;
+  texture->data  = gegl_clCreateImage2D (gegl_cl_get_context(), CL_MEM_READ_WRITE,
+                                         &texture->format,
+                                         texture->width,
+                                         texture->height,
+                                         0,  NULL, &errcode);
+  if (errcode != CL_SUCCESS)
+  {
+    g_free(texture);
+    return NULL;
+  }
+
+  return texture;
+}
+
+void
+gegl_cl_texture_free (GeglClTexture *texture)
+{
+  gegl_clReleaseMemObject (texture->data);
+  g_free (texture);
+}
+
+void
+gegl_cl_texture_get (const GeglClTexture *texture,
+                     gpointer             dst)
+{
+  const size_t origin[3] = {0,0,0};
+  const size_t region[3] = {texture->width,
+                            texture->height,
+                            1};
+  gegl_clEnqueueReadImage(gegl_cl_get_command_queue(),
+                          texture->data, CL_TRUE, origin, region, 0, 0, dst,
+                          0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+}
+
+void
+gegl_cl_texture_set (GeglClTexture  *texture,
+                     const gpointer  src)
+{
+  const size_t origin[3] = {0,0,0};
+  const size_t region[3] = {texture->width,
+                            texture->height,
+                            1};
+  gegl_clEnqueueWriteImage(gegl_cl_get_command_queue(),
+                          texture->data, CL_TRUE, origin, region, 0, 0, src,
+                          0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+}
+
+GeglClTexture *
+gegl_cl_texture_dup (const GeglClTexture *texture)
+{
+  const size_t origin[3] = {0,0,0};
+  const size_t region[3] = {texture->width,
+                            texture->height,
+                            1};
+
+  GeglClTexture *new_texture = gegl_cl_texture_new (texture->width,
+                                                    texture->height);
+
+  gegl_clEnqueueCopyImage(gegl_cl_get_command_queue(),
+                          texture->data, new_texture->data,
+                          origin, origin, region,
+                          0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+
+  return new_texture;
+}
+
+void
+gegl_cl_texture_copy (const GeglClTexture *src,
+                      GeglClTexture *dst)
+{
+  const size_t origin[3] = {0,0,0};
+  const size_t region[3] = {src->width,
+                            src->height,
+                            1};
+
+  gegl_clEnqueueCopyImage (gegl_cl_get_command_queue(),
+                           src->data, dst->data,
+                           origin, origin, region,
+                           0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+}
+
+void
+gegl_cl_texture_clear (const GeglClTexture *texture)
+{
+  cl_int errcode;
+  char buffer[10000];
+
+  size_t global_worksize[2] = {texture->width, texture->height};
+
+  const char* kernel_source[] =
+  {
+  "__kernel void kernel_clear (__write_only image2d_t img)  \n",
+  "{                                                        \n",
+  "  int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n",
+  "  float4 v = (float4) 0.0f;                              \n",
+  "  write_imagef(img,gid,v);                               \n",
+  "}                                                        \n",
+  NULL,
+  };
+
+  cl_program program;
+  cl_kernel kernel;
+
+  g_log(G_LOG_DOMAIN, G_LOG_LEVEL_DEBUG, "lines: %u\n", gegl_cl_count_lines(kernel_source));
+
+  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_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) );
+}
diff --git a/gegl/opencl/gegl-cl-texture.h b/gegl/opencl/gegl-cl-texture.h
new file mode 100644
index 0000000..26c187c
--- /dev/null
+++ b/gegl/opencl/gegl-cl-texture.h
@@ -0,0 +1,46 @@
+#ifndef __GEGL_CL_TEXTURE_H__
+#define __GEGL_CL_TEXTURE_H__
+
+#include <glib-object.h>
+
+#include "gegl-cl-init.h"
+
+G_BEGIN_DECLS
+
+struct _GeglClTexture
+{
+  cl_mem data;
+  cl_image_format format;
+  gint width;
+  gint height;
+};
+
+typedef struct _GeglClTexture GeglClTexture;
+
+GType           gegl_cl_texture_get_type (void) G_GNUC_CONST;
+
+GeglClTexture  *gegl_cl_texture_new      (const gint           width,
+                                          const gint           height);
+
+void            gegl_cl_texture_free     (GeglClTexture       *texture);
+
+void            gegl_cl_texture_get      (const GeglClTexture *texture,
+                                          gpointer             dest);
+
+void            gegl_cl_texture_set      (GeglClTexture       *texture,
+                                          const gpointer       src);
+
+GeglClTexture  *gegl_cl_texture_dup      (const GeglClTexture *texture);
+
+void            gegl_cl_texture_copy     (const GeglClTexture *src,
+                                          GeglClTexture *dst);
+
+void            gegl_cl_texture_clear    (const GeglClTexture *texture);
+
+#define GEGL_TYPE_CL_TEXTURE   (gegl_cl_texture_get_type())
+
+#define gegl_cl_texture_get_data(texture) (texture->data)
+
+G_END_DECLS
+
+#endif  /* __GEGL_CL_TEXTURE_H__ */