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

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

commit 9d8e6db7f9c937aa2ee72e934d9bfb2b973c9266
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 |  184 +++++++++++++++++++++++++++++++++++++++++
 gegl/opencl/gegl-cl-texture.h |   56 +++++++++++++
 3 files changed, 243 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..5aa3073
--- /dev/null
+++ b/gegl/opencl/gegl-cl-texture.c
@@ -0,0 +1,184 @@
+#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_warning("OpenCL Alloc Error: %s", gegl_cl_errstring(errcode));
+    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,
+                      GeglRectangle        *src_rect,
+                      GeglClTexture        *dst,
+                      gint                  dst_x,
+                      gint                  dst_y)
+{
+  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
+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;
+
+  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) );
+}
+
+void
+gegl_cl_texture_clear (const GeglClTexture *texture,
+                       GeglRectangle       *_rect)
+{
+  gfloat zero[4] = {0.0f, 0.0f, 0.0f, 0.0f};
+  gegl_cl_texture_fill (texture, _rect, zero);
+}
diff --git a/gegl/opencl/gegl-cl-texture.h b/gegl/opencl/gegl-cl-texture.h
new file mode 100644
index 0000000..e916324
--- /dev/null
+++ b/gegl/opencl/gegl-cl-texture.h
@@ -0,0 +1,56 @@
+#ifndef __GEGL_CL_TEXTURE_H__
+#define __GEGL_CL_TEXTURE_H__
+
+#include <glib-object.h>
+
+#include "gegl.h"
+#include "gegl-cl-types.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,
+                                          GeglRectangle       *src_rect,
+                                          GeglClTexture       *dst,
+                                          gint                 dst_x,
+                                          gint                 dst_y);
+
+void            gegl_cl_texture_fill     (const GeglClTexture *texture,
+                                          GeglRectangle       *rect,
+                                          gfloat               color[4]);
+
+void            gegl_cl_texture_clear    (const GeglClTexture *texture,
+                                          GeglRectangle       *rect);
+
+#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__ */