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

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

commit 700e2716f7d5328802f4f1b4fa19eed1dbc7a87b
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 |  118 +++++++++++++++++++++++++++++++++++++++++
 gegl/opencl/gegl-cl-texture.h |   47 ++++++++++++++++
 3 files changed, 168 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..6259f62
--- /dev/null
+++ b/gegl/opencl/gegl-cl-texture.c
@@ -0,0 +1,118 @@
+#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 size)
+{
+  cl_int errcode;
+
+  GeglClTexture *texture = g_new (GeglClTexture, 1);
+  texture->data  = gegl_clCreateBuffer (gegl_cl_get_context(), CL_MEM_READ_ONLY,
+                                        size, 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 gint           size)
+{
+  gegl_clEnqueueReadBuffer(gegl_cl_get_command_queue(),
+                           texture->data, CL_TRUE, 0, size, dst,
+                           0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+}
+
+void
+gegl_cl_texture_set (GeglClTexture  *texture,
+                     const gpointer  src,
+                     const gint      size)
+{
+  gegl_clEnqueueWriteBuffer(gegl_cl_get_command_queue(),
+                            texture->data, CL_TRUE, 0, size, src,
+                            0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+}
+
+GeglClTexture *
+gegl_cl_texture_dup (const GeglClTexture *texture,
+                     const gint           size)
+{
+  GeglClTexture *new_texture = gegl_cl_texture_new (size);
+
+  gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
+                            texture->data, new_texture->data,
+                            0, 0, size, 0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+
+  return new_texture;
+}
+
+void
+gegl_cl_texture_copy (const GeglClTexture *src,
+                      GeglClTexture *dst,
+                      const gint           size)
+{
+  gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
+                            src->data, dst->data,
+                            0, 0, size, 0, NULL, NULL);
+  gegl_clFinish(gegl_cl_get_command_queue());
+}
+
+void
+gegl_cl_texture_clear (const GeglClTexture *texture,
+                       const gint           size)
+{
+  cl_int errcode;
+  char buffer[16384];
+
+  size_t global_worksize;
+
+  const char* kernel_source[] =
+  {
+  "__kernel void kernel_clear (__global float* dst) \n",
+  "{                                                \n",
+  "  dst[get_global_id(0)] = 0.0f;                  \n",
+  "}                                                \n",
+  };
+
+  cl_program program;
+  cl_kernel kernel;
+
+  CL_SAFE_CALL( program = gegl_clCreateProgramWithSource(gegl_cl_get_context(), 4, (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) );
+
+  // THIS ASSUMES TEXTURE TYPE IS FLOAT !!
+  global_worksize = (size/4);
+
+  CL_SAFE_CALL( errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), kernel, 1, 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..761f115
--- /dev/null
+++ b/gegl/opencl/gegl-cl-texture.h
@@ -0,0 +1,47 @@
+#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;
+};
+
+typedef struct _GeglClTexture GeglClTexture;
+
+GType           gegl_cl_texture_get_type (void) G_GNUC_CONST;
+
+GeglClTexture  *gegl_cl_texture_new      (const gint           size);
+
+void            gegl_cl_texture_free     (GeglClTexture       *texture);
+
+void            gegl_cl_texture_get      (const GeglClTexture *texture,
+                                          gpointer             dest,
+                                          const gint           size);
+
+void            gegl_cl_texture_set      (GeglClTexture       *texture,
+                                          const gpointer       src,
+                                          const gint           size);
+
+GeglClTexture  *gegl_cl_texture_dup      (const GeglClTexture *texture,
+                                          const gint           size);
+
+void            gegl_cl_texture_copy     (const GeglClTexture *src,
+                                          GeglClTexture *dst,
+                                          const gint           size);
+
+void            gegl_cl_texture_clear    (const GeglClTexture *texture,
+                                          const gint           size);
+
+#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__ */