Do not #include C source file anymore in SOCL.

Inclusions of C source files have been replaced by a more conventional approach using .h headers.

.gitignore

@@ -20,6 +20,7 @@ Makefile.in
 *.o
 *.lo
 *.la
+*.swp
 .dirstamp
 stamp-h[0-9]*
 starpu.log

configure.ac

@@ -1115,15 +1115,12 @@ AM_CONDITIONAL([HAVE_GUILE], [test "x$GUILE" != "x"])
 
 AC_ARG_ENABLE([socl],
   [AS_HELP_STRING([--enable-socl],
-    [build the OpenCL interfacce (SOCL)])],
+    [build the OpenCL interface (SOCL)])],
   [enable_socl="$enableval"],
   [enable_socl="no"])
 
-#TODO: check that OpenCL is enabled
-
 if test "x$enable_socl" = "xyes"; then
    STARPU_SOCL_SUPPORT
-
    build_socl="yes"
 else
    build_socl="no"

socl/src/Makefile.am

@@ -24,5 +24,83 @@ SUBDIRS =
 lib_LTLIBRARIES = libsocl.la
 
 libsocl_la_SOURCES = 						\
-  opencl.c
+  socl.c \
+  gc.c \
+  graph.c \
+  event.c \
+  init.c \
+  task.c \
+  command_queue.c \
+  mem_objects.c \
+  util.c \
+  devices.c \
+  cl_getplatformids.c \
+  cl_getplatforminfo.c \
+  cl_getdeviceids.c \
+  cl_getdeviceinfo.c \
+  cl_releasecontext.c \
+  cl_createcontext.c \
+  cl_createcontextfromtype.c \
+  cl_retaincontext.c \
+  cl_getcontextinfo.c \
+  cl_releasecommandqueue.c \
+  cl_createcommandqueue.c \
+  cl_retaincommandqueue.c \
+  cl_getcommandqueueinfo.c \
+  cl_setcommandqueueproperty.c \
+  cl_releaseevent.c \
+  cl_waitforevents.c \
+  cl_geteventinfo.c \
+  cl_retainevent.c \
+  cl_enqueuemarker.c \
+  cl_enqueuewaitforevents.c \
+  cl_enqueuebarrier.c \
+  cl_flush.c \
+  cl_finish.c \
+  cl_releasememobject.c \
+  cl_createbuffer.c \
+  cl_createimage2d.c \
+  cl_createimage3d.c \
+  cl_retainmemobject.c \
+  cl_getsupportedimageformats.c \
+  cl_getmemobjectinfo.c \
+  cl_getimageinfo.c \
+  cl_createsampler.c \
+  cl_retainsampler.c \
+  cl_releasesampler.c \
+  cl_getsamplerinfo.c \
+  cl_releaseprogram.c \
+  cl_createprogramwithsource.c \
+  cl_createprogramwithbinary.c \
+  cl_retainprogram.c \
+  cl_buildprogram.c \
+  cl_unloadcompiler.c \
+  cl_getprograminfo.c \
+  cl_getprogrambuildinfo.c \
+  cl_releasekernel.c \
+  cl_createkernel.c \
+  cl_createkernelsinprogram.c \
+  cl_retainkernel.c \
+  cl_setkernelarg.c \
+  cl_getkernelinfo.c \
+  cl_getkernelworkgroupinfo.c \
+  cl_enqueuereadbuffer.c \
+  cl_enqueuewritebuffer.c \
+  cl_enqueuecopybuffer.c \
+  cl_enqueuereadimage.c \
+  cl_enqueuewriteimage.c \
+  cl_enqueuecopyimage.c \
+  cl_enqueuecopyimagetobuffer.c \
+  cl_enqueuecopybuffertoimage.c \
+  cl_enqueuemapbuffer.c \
+  cl_enqueuemapimage.c \
+  cl_enqueueunmapmemobject.c \
+  cl_enqueuetask.c \
+  cl_enqueuendrangekernel.c \
+  cl_enqueuenativekernel.c \
+  cl_geteventprofilinginfo.c \
+  cl_getextensionfunctionaddress.c
+
+
+
 

socl/src/cl_buildprogram.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 struct bp_data {
    cl_program program;
    char * options;

socl/src/cl_createbuffer.c.inc

@@ -14,6 +14,24 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
+static void release_callback_memobject(void * e) {
+  cl_mem mem = (cl_mem)e;
+
+  /* Release references */
+  gc_entity_unstore(&mem->context);
+
+  //Delete this mem_object from the mem_object list
+  mem_object_release(mem);
+
+  /* Destruct object */
+  starpu_data_unregister_no_coherency(mem->handle);
+
+  if (!(mem->flags & CL_MEM_USE_HOST_PTR))
+    free(mem->ptr);
+}
+
 
 /**
  * \brief Create a buffer

socl/src/cl_createcommandqueue.c.inc

@@ -14,6 +14,27 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
+static void release_callback_command_queue(void * e) {
+  cl_command_queue cq = (cl_command_queue)e;
+
+  //Disable StarPU profiling if necessary
+  if (cq->properties & CL_QUEUE_PROFILING_ENABLE) {
+    profiling_queue_count -= 1;
+    if (profiling_queue_count == 0)
+      starpu_profiling_status_set(STARPU_PROFILING_DISABLE);
+  }
+
+  /* Release references */
+  gc_entity_unstore(&cq->context);
+
+  /* Destruct object */
+  pthread_spin_destroy(&cq->spin);
+  free(cq->events);
+}
+
+
 
 CL_API_ENTRY cl_command_queue CL_API_CALL
 soclCreateCommandQueue(cl_context                   context, 

socl/src/cl_createcontext.c.inc

@@ -14,6 +14,19 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
+static void release_callback_context(void * e) {
+  cl_context context = (cl_context)e;
+
+  /* Destruct object */
+  if (context->properties != NULL)
+    free(context->properties);
+
+  free(context->devices);
+}
+
+
 
 CL_API_ENTRY cl_context CL_API_CALL
 soclCreateContext(const cl_context_properties * properties,

socl/src/cl_createcontextfromtype.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_context CL_API_CALL
 soclCreateContextFromType(const cl_context_properties * properties,
                         cl_device_type                UNUSED(device_type),

socl/src/cl_createimage2d.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_mem CL_API_CALL
 soclCreateImage2D(cl_context              UNUSED(context),
                 cl_mem_flags            UNUSED(flags),

socl/src/cl_createimage3d.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_mem CL_API_CALL
 soclCreateImage3D(cl_context              UNUSED(context),
                 cl_mem_flags            UNUSED(flags),

socl/src/cl_createkernel.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 
 static void soclCreateKernel_task(void *data) {
    struct _cl_kernel *k = (struct _cl_kernel *)data;
@@ -53,6 +55,50 @@ static void soclCreateKernel_task(void *data) {
    }
 }
 
+static void rk_task(void *data) {
+   cl_kernel k = (cl_kernel)data;
+
+   int range = starpu_worker_get_range();
+
+   cl_int err = clReleaseKernel(k->cl_kernels[range]);
+   if (err != CL_SUCCESS)
+      DEBUG_CL("clReleaseKernel", err);
+}
+
+static void release_callback_kernel(void * e) {
+  cl_kernel kernel = (cl_kernel)e;
+
+  //Free args
+  unsigned int j;
+  for (j=0; j<kernel->arg_count; j++) {
+    switch (kernel->arg_type[j]) {
+      case Null:
+        break;
+      case Buffer:
+        gc_entity_unstore((cl_mem*)&kernel->arg_value[j]);
+        break;
+      case Immediate:
+        free(kernel->arg_value[j]);
+        break;
+    }
+  }
+  if (kernel->arg_size != NULL)
+    free(kernel->arg_size);
+  if (kernel->arg_value != NULL)
+    free(kernel->arg_value);
+  if (kernel->arg_type != NULL)
+    free(kernel->arg_type);
+
+  //Release real kernels...
+  starpu_execute_on_each_worker(rk_task, kernel, STARPU_OPENCL);
+
+  gc_entity_unstore(&kernel->program);
+
+  free(kernel->kernel_name);
+  free(kernel->cl_kernels);
+  free(kernel->errcodes);
+}
+
 CL_API_ENTRY cl_kernel CL_API_CALL
 soclCreateKernel(cl_program    program,
                const char *    kernel_name,

socl/src/cl_createkernelsinprogram.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclCreateKernelsInProgram(cl_program   UNUSED(program),
                          cl_uint        UNUSED(num_kernels),

socl/src/cl_createprogramwithbinary.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_program CL_API_CALL
 soclCreateProgramWithBinary(cl_context                     UNUSED(context),
                           cl_uint                        UNUSED(num_devices),

socl/src/cl_createprogramwithsource.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 struct cpws_data {
    struct _cl_program *program;
    cl_int *errcodes;
@@ -37,6 +39,31 @@ static void soclCreateProgramWithSource_task(void *data) {
 
 }
 
+static void rp_task(void *data) {
+   struct _cl_program *d = (struct _cl_program*)data;
+
+   int range = starpu_worker_get_range();
+
+   cl_int err = clReleaseProgram(d->cl_programs[range]);
+   if (err != CL_SUCCESS)
+      DEBUG_CL("clReleaseProgram", err);
+}
+
+static void release_callback_program(void * e) {
+  cl_program program = (cl_program)e;
+
+  /* Destruct object */
+  starpu_execute_on_each_worker(rp_task, program, STARPU_OPENCL);
+
+  /* Release references */
+  gc_entity_unstore(&program->context);
+
+  free(program->cl_programs);
+
+  if (program->options != NULL)
+    free(program->options);
+}
+
 
 CL_API_ENTRY cl_program CL_API_CALL
 soclCreateProgramWithSource(cl_context      context,

socl/src/cl_createsampler.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_sampler CL_API_CALL
 soclCreateSampler(cl_context          UNUSED(context),
                 cl_bool             UNUSED(normalized_coords), 

socl/src/cl_enqueuebarrier.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueBarrier(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_enqueuecopybuffer.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 struct arg_copybuffer {
    size_t src_offset, dst_offset;
    cl_mem src_buffer, dst_buffer;

socl/src/cl_enqueuecopybuffertoimage.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueCopyBufferToImage(cl_command_queue UNUSED(command_queue),
                            cl_mem           UNUSED(src_buffer),

socl/src/cl_enqueuecopyimage.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueCopyImage(cl_command_queue   UNUSED(command_queue),
                    cl_mem               UNUSED(src_image),

socl/src/cl_enqueuecopyimagetobuffer.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueCopyImageToBuffer(cl_command_queue UNUSED(command_queue),
                            cl_mem           UNUSED(src_image),

socl/src/cl_enqueuemapbuffer.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 struct mb_data {
   cl_event ev;
   cl_mem buffer;

socl/src/cl_enqueuemapimage.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY void * CL_API_CALL
 soclEnqueueMapImage(cl_command_queue  UNUSED(command_queue),
                   cl_mem            UNUSED(image), 

socl/src/cl_enqueuemarker.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueMarker(cl_command_queue  cq,
                 cl_event *          event) CL_API_SUFFIX__VERSION_1_0

socl/src/cl_enqueuenativekernel.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueNativeKernel(cl_command_queue  UNUSED(command_queue),
 					       __attribute__((unused)) void (*user_func)(void *), 

socl/src/cl_enqueuendrangekernel.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 typedef struct running_kernel * running_kernel;
 
 struct running_kernel {
@@ -127,25 +129,30 @@ static struct starpu_perfmodel_t perf_model = {
   .symbol = "perf_model"
 };
 
-CL_API_ENTRY cl_int CL_API_CALL
-soclEnqueueNDRangeKernel(cl_command_queue cq,
-                       cl_kernel        kernel,
-                       cl_uint          work_dim,
-                       const size_t *   global_work_offset,
-                       const size_t *   global_work_size,
-                       const size_t *   local_work_size,
-                       cl_uint          num_events,
-                       const cl_event * events,
-                       cl_event *       event) CL_API_SUFFIX__VERSION_1_0
-{
+/**
+ * Real kernel enqueuing command
+ */
+cl_int node_play_enqueue_kernel(node_enqueue_kernel n) {
+
    struct starpu_task *task;
    running_kernel arg;
    starpu_codelet *codelet;
    cl_event ev;
-
-   /***********************
-    * Allocate structures *
-    ***********************/
+   
+   /* Alias struc fields */
+   cl_command_queue cq = n->cq;
+   cl_kernel        kernel = n->kernel;
+   cl_uint          work_dim = n->work_dim;
+   const size_t *   global_work_offset = n->global_work_offset;
+   const size_t *   global_work_size = n->global_work_size;
+   const size_t *   local_work_size = n->local_work_size;
+   cl_uint          num_events = n->num_events;
+   const cl_event * events = n->events;
+   cl_event *       event = n->event;
+   char 	    is_task = n->is_task;
+
+
+   /* Allocate structures */
 
    /* Codelet */
    codelet = (starpu_codelet*)malloc(sizeof(starpu_codelet));
@@ -160,7 +167,7 @@ soclEnqueueNDRangeKernel(cl_command_queue cq,
    }
 
    /* StarPU task */
-   task = task_create(&work_dim != &et_work_dim ? CL_COMMAND_NDRANGE_KERNEL: CL_COMMAND_TASK);
+   task = task_create(is_task ? CL_COMMAND_TASK : CL_COMMAND_NDRANGE_KERNEL);
    ev = task_event(task);
 
    /*******************
@@ -281,3 +288,30 @@ soclEnqueueNDRangeKernel(cl_command_queue cq,
 
    return ret;
 }
+
+/**
+ * Virtual kernel enqueueing command
+ */
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNDRangeKernel(cl_command_queue cq,
+                       cl_kernel        kernel,
+                       cl_uint          work_dim,
+                       const size_t *   global_work_offset,
+                       const size_t *   global_work_size,
+                       const size_t *   local_work_size,
+                       cl_uint          num_events,
+                       const cl_event * events,
+                       cl_event *       event) CL_API_SUFFIX__VERSION_1_1
+{
+	node_enqueue_kernel n;
+
+	n = graph_create_enqueue_kernel(0, cq, kernel, work_dim, global_work_offset, global_work_size,
+		local_work_size, num_events, events, event, kernel->arg_count, kernel->arg_size,
+		kernel->arg_type, kernel->arg_value);
+	
+	//FIXME: temporarily, we execute the node directly. In the future, we will postpone this.
+	node_play_enqueue_kernel(n);
+
+	//graph_store(n);
+	return CL_SUCCESS;
+}

socl/src/cl_enqueuereadbuffer.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 struct arg_readbuffer {
    size_t offset;
    size_t cb;

socl/src/cl_enqueuereadimage.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueReadImage(cl_command_queue   UNUSED(command_queue),
                    cl_mem               UNUSED(image),

socl/src/cl_enqueuetask.c.inc

@@ -14,10 +14,12 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
-static cl_uint et_work_dim = 3;
-static const size_t et_global_work_offset[3] = {0,0,0};
-static const size_t et_global_work_size[3] = {1,1,1};
-static const size_t * et_local_work_size = NULL;
+#include "socl.h"
+
+static cl_uint work_dim = 3;
+static const size_t global_work_offset[3] = {0,0,0};
+static const size_t global_work_size[3] = {1,1,1};
+static const size_t * local_work_size = NULL;
 
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueNDRangeKernel(cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *) CL_API_SUFFIX__VERSION_1_0;
@@ -29,5 +31,15 @@ soclEnqueueTask(cl_command_queue cq,
               const cl_event *  events,
               cl_event *        event) CL_API_SUFFIX__VERSION_1_0
 {
-   return soclEnqueueNDRangeKernel(cq, kernel, et_work_dim, et_global_work_offset, et_global_work_size, et_local_work_size, num_events, events, event); 
+	node_enqueue_kernel n;
+
+	n = graph_create_enqueue_kernel(1, cq, kernel, work_dim, global_work_offset, global_work_size,
+		local_work_size, num_events, events, event, kernel->arg_count, kernel->arg_size,
+		kernel->arg_type, kernel->arg_value);
+	
+	//FIXME: temporarily, we execute the node directly. In the future, we will postpone this.
+	node_play_enqueue_kernel(n);
+
+	//graph_store(n);
+	return CL_SUCCESS;
 }

socl/src/cl_enqueueunmapmemobject.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueUnmapMemObject(cl_command_queue cq,
                         cl_mem            memobj,

socl/src/cl_enqueuewaitforevents.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueWaitForEvents(cl_command_queue cq,
                        cl_uint          num_events,

socl/src/cl_enqueuewritebuffer.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 struct arg_writebuffer {
    size_t offset;
    size_t cb;

socl/src/cl_enqueuewriteimage.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueWriteImage(cl_command_queue  UNUSED(command_queue),
                     cl_mem              UNUSED(image),

socl/src/cl_finish.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclFinish(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_flush.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclFlush(cl_command_queue UNUSED(command_queue)) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_getcommandqueueinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetCommandQueueInfo(cl_command_queue    cq,
                       cl_command_queue_info param_name,

socl/src/cl_getcontextinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetContextInfo(cl_context       context, 
                  cl_context_info    param_name, 

socl/src/cl_getdeviceids.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 
 /**
  * \brief Return one device of each kind

socl/src/cl_getdeviceinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 
 /**
  * \brief Return dummy infos

socl/src/cl_geteventinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetEventInfo(cl_event       event,
                cl_event_info    param_name,

socl/src/cl_geteventprofilinginfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetEventProfilingInfo(cl_event          event,
                         cl_profiling_info   param_name,

socl/src/cl_getextensionfunctionaddress.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY void * CL_API_CALL
 soclGetExtensionFunctionAddress(const char * UNUSED(func_name)) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_getimageinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetImageInfo(cl_mem           UNUSED(image),
                cl_image_info    UNUSED(param_name), 

socl/src/cl_getkernelinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetKernelInfo(cl_kernel       kernel,
                 cl_kernel_info  param_name,

socl/src/cl_getkernelworkgroupinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 struct gkwgi_data {
    cl_kernel_work_group_info param_name;
    cl_kernel kernel;

socl/src/cl_getmemobjectinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetMemObjectInfo(cl_mem           mem,
                    cl_mem_info      param_name, 

socl/src/cl_getplatformids.c.inc

@@ -14,6 +14,7 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
 
 /**
  * \brief Get StarPU platform ID

socl/src/cl_getplatforminfo.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
 
 /**
  * \brief Get information about StarPU platform

socl/src/cl_getprogrambuildinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetProgramBuildInfo(cl_program          program,
                       cl_device_id          UNUSED(device),

socl/src/cl_getprograminfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetProgramInfo(cl_program       program,
                  cl_program_info    param_name,

socl/src/cl_getsamplerinfo.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "getinfo.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetSamplerInfo(cl_sampler       UNUSED(sampler),
                  cl_sampler_info    UNUSED(param_name),

socl/src/cl_getsupportedimageformats.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclGetSupportedImageFormats(cl_context           UNUSED(context),
                            cl_mem_flags         UNUSED(flags),

socl/src/cl_releasecommandqueue.c

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "socl.h"
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseCommandQueue(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
+{
+
+  gc_entity_release(cq);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasecommandqueue.c.inc

@@ -1,43 +0,0 @@
-/* StarPU --- Runtime system for heterogeneous multicore architectures.
- *
- * Copyright (C) 2010,2011 University of Bordeaux
- *
- * StarPU is free software; you can redistribute it and/or modify
- * it under the terms of the GNU Lesser General Public License as published by
- * the Free Software Foundation; either version 2.1 of the License, or (at
- * your option) any later version.
- *
- * StarPU is distributed in the hope that it will be useful, but
- * WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
- *
- * See the GNU Lesser General Public License in COPYING.LGPL for more details.
- */
-
-static void release_callback_command_queue(void * e) {
-  cl_command_queue cq = (cl_command_queue)e;
-
-  //Disable StarPU profiling if necessary
-  if (cq->properties & CL_QUEUE_PROFILING_ENABLE) {
-    profiling_queue_count -= 1;
-    if (profiling_queue_count == 0)
-      starpu_profiling_status_set(STARPU_PROFILING_DISABLE);
-  }
-
-  /* Release references */
-  gc_entity_unstore(&cq->context);
-
-  /* Destruct object */
-  pthread_spin_destroy(&cq->spin);
-  free(cq->events);
-}
-
-
-CL_API_ENTRY cl_int CL_API_CALL
-soclReleaseCommandQueue(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
-{
-
-  gc_entity_release(cq);
-
-  return CL_SUCCESS;
-}

socl/src/cl_releasecontext.c.inc

@@ -14,16 +14,7 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
-static void release_callback_context(void * e) {
-  cl_context context = (cl_context)e;
-
-  /* Destruct object */
-  if (context->properties != NULL)
-    free(context->properties);
-
-  free(context->devices);
-}
-
+#include "socl.h"
 
 CL_API_ENTRY cl_int CL_API_CALL
 soclReleaseContext(cl_context context) CL_API_SUFFIX__VERSION_1_0

socl/src/cl_releaseevent.c

@@ -0,0 +1,28 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "socl.h"
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseEvent(cl_event event) CL_API_SUFFIX__VERSION_1_0
+{
+  if (event == NULL)
+    return CL_INVALID_EVENT;
+
+  gc_entity_release(event);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasekernel.c

@@ -0,0 +1,28 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "socl.h"
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseKernel(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0
+{
+  if (kernel == NULL)
+    return CL_INVALID_KERNEL;
+
+  gc_entity_release(kernel);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasekernel.c.inc

@@ -1,70 +0,0 @@
-/* StarPU --- Runtime system for heterogeneous multicore architectures.
- *
- * Copyright (C) 2010,2011 University of Bordeaux
- *
- * StarPU is free software; you can redistribute it and/or modify
- * it under the terms of the GNU Lesser General Public License as published by
- * the Free Software Foundation; either version 2.1 of the License, or (at
- * your option) any later version.
- *
- * StarPU is distributed in the hope that it will be useful, but
- * WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
- *
- * See the GNU Lesser General Public License in COPYING.LGPL for more details.
- */
-
-static void rk_task(void *data) {
-   cl_kernel k = (cl_kernel)data;
-
-   int range = starpu_worker_get_range();
-
-   cl_int err = clReleaseKernel(k->cl_kernels[range]);
-   if (err != CL_SUCCESS)
-      DEBUG_CL("clReleaseKernel", err);
-}
-
-static void release_callback_kernel(void * e) {
-  cl_kernel kernel = (cl_kernel)e;
-
-  //Free args
-  unsigned int j;
-  for (j=0; j<kernel->arg_count; j++) {
-    switch (kernel->arg_type[j]) {
-      case Null:
-        break;
-      case Buffer:
-        gc_entity_unstore((cl_mem*)&kernel->arg_value[j]);
-        break;
-      case Immediate:
-        free(kernel->arg_value[j]);
-        break;
-    }
-  }
-  if (kernel->arg_size != NULL)
-    free(kernel->arg_size);
-  if (kernel->arg_value != NULL)
-    free(kernel->arg_value);
-  if (kernel->arg_type != NULL)
-    free(kernel->arg_type);
-
-  //Release real kernels...
-  starpu_execute_on_each_worker(rk_task, kernel, STARPU_OPENCL);
-
-  gc_entity_unstore(&kernel->program);
-
-  free(kernel->kernel_name);
-  free(kernel->cl_kernels);
-  free(kernel->errcodes);
-}
-
-CL_API_ENTRY cl_int CL_API_CALL
-soclReleaseKernel(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0
-{
-  if (kernel == NULL)
-    return CL_INVALID_KERNEL;
-
-  gc_entity_release(kernel);
-
-  return CL_SUCCESS;
-}

socl/src/cl_releasememobject.c

@@ -0,0 +1,25 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "socl.h"
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseMemObject(cl_mem mem) CL_API_SUFFIX__VERSION_1_0
+{
+  gc_entity_release(mem);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasememobject.c.inc

@@ -1,39 +0,0 @@
-/* StarPU --- Runtime system for heterogeneous multicore architectures.
- *
- * Copyright (C) 2010,2011 University of Bordeaux
- *
- * StarPU is free software; you can redistribute it and/or modify
- * it under the terms of the GNU Lesser General Public License as published by
- * the Free Software Foundation; either version 2.1 of the License, or (at
- * your option) any later version.
- *
- * StarPU is distributed in the hope that it will be useful, but
- * WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
- *
- * See the GNU Lesser General Public License in COPYING.LGPL for more details.
- */
-
-static void release_callback_memobject(void * e) {
-  cl_mem mem = (cl_mem)e;
-
-  /* Release references */
-  gc_entity_unstore(&mem->context);
-
-  //Delete this mem_object from the mem_object list
-  mem_object_release(mem);
-
-  /* Destruct object */
-  starpu_data_unregister_no_coherency(mem->handle);
-
-  if (!(mem->flags & CL_MEM_USE_HOST_PTR))
-    free(mem->ptr);
-}
-
-CL_API_ENTRY cl_int CL_API_CALL
-soclReleaseMemObject(cl_mem mem) CL_API_SUFFIX__VERSION_1_0
-{
-  gc_entity_release(mem);
-
-  return CL_SUCCESS;
-}

socl/src/cl_releaseprogram.c

@@ -0,0 +1,28 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "socl.h"
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseProgram(cl_program program) CL_API_SUFFIX__VERSION_1_0
+{
+  if (program == NULL)
+    return CL_INVALID_PROGRAM;
+
+  gc_entity_release(program);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releaseprogram.c.inc

@@ -1,51 +0,0 @@
-/* StarPU --- Runtime system for heterogeneous multicore architectures.
- *
- * Copyright (C) 2010,2011 University of Bordeaux
- *
- * StarPU is free software; you can redistribute it and/or modify
- * it under the terms of the GNU Lesser General Public License as published by
- * the Free Software Foundation; either version 2.1 of the License, or (at
- * your option) any later version.
- *
- * StarPU is distributed in the hope that it will be useful, but
- * WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
- *
- * See the GNU Lesser General Public License in COPYING.LGPL for more details.
- */
-
-static void rp_task(void *data) {
-   struct _cl_program *d = (struct _cl_program*)data;
-
-   int range = starpu_worker_get_range();
-
-   cl_int err = clReleaseProgram(d->cl_programs[range]);
-   if (err != CL_SUCCESS)
-      DEBUG_CL("clReleaseProgram", err);
-}
-
-static void release_callback_program(void * e) {
-  cl_program program = (cl_program)e;
-
-  /* Destruct object */
-  starpu_execute_on_each_worker(rp_task, program, STARPU_OPENCL);
-
-  /* Release references */
-  gc_entity_unstore(&program->context);
-
-  free(program->cl_programs);
-
-  if (program->options != NULL)
-    free(program->options);
-}
-
-CL_API_ENTRY cl_int CL_API_CALL
-soclReleaseProgram(cl_program program) CL_API_SUFFIX__VERSION_1_0
-{
-  if (program == NULL)
-    return CL_INVALID_PROGRAM;
-
-  gc_entity_release(program);
-
-  return CL_SUCCESS;
-}

socl/src/cl_releasesampler.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclReleaseSampler(cl_sampler UNUSED(sampler)) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_retaincommandqueue.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclRetainCommandQueue(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_retaincontext.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclRetainContext(cl_context context) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_retainevent.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclRetainEvent(cl_event event) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_retainkernel.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclRetainKernel(cl_kernel    kernel) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_retainmemobject.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclRetainMemObject(cl_mem mem) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_retainprogram.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclRetainProgram(cl_program program) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_retainsampler.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclRetainSampler(cl_sampler UNUSED(sampler)) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_setcommandqueueproperty.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclSetCommandQueueProperty(cl_command_queue            command_queue,
                           cl_command_queue_properties   properties, 

socl/src/cl_setkernelarg.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclSetKernelArg(cl_kernel  kernel,
                cl_uint      arg_index,

socl/src/cl_unloadcompiler.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclUnloadCompiler(void) CL_API_SUFFIX__VERSION_1_0
 {

socl/src/cl_waitforevents.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 CL_API_ENTRY cl_int CL_API_CALL
 soclWaitForEvents(cl_uint           num_events,
                 const cl_event *    event_list) CL_API_SUFFIX__VERSION_1_0

socl/src/helper_command_queue.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "task.h"
+#include "gc.h"
 
 /**
  * WARNING: command queues do NOT hold references on events. Only events hold references
@@ -25,7 +28,7 @@
  * Enqueue the given task but put fake_event into the command queue.
  * This is used when a tag notified by application is used (cf clEnqueueMapBuffer, etc.)
  */
-static cl_int command_queue_enqueue_fakeevent(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events, cl_event fake_event) {
+cl_int command_queue_enqueue_fakeevent(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events, cl_event fake_event) {
 
   int in_order = !(cq->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
 
@@ -77,12 +80,12 @@ static cl_int command_queue_enqueue_fakeevent(cl_command_queue cq, starpu_task *
    return CL_SUCCESS;
 }
 
-static cl_int command_queue_enqueue(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events) {
+cl_int command_queue_enqueue(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events) {
   return command_queue_enqueue_fakeevent(cq, task, barrier, num_events, events, NULL);
 }
 
 
-static cl_event enqueueBarrier(cl_command_queue cq) {
+cl_event enqueueBarrier(cl_command_queue cq) {
 
    //CL_COMMAND_MARKER has been chosen as CL_COMMAND_BARRIER doesn't exist
    starpu_task * task = task_create(CL_COMMAND_MARKER);

socl/src/command_queue.h

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_COMMAND_QUEUE_H
+#define SOCl_COMMAND_QUEUE_H
+
+cl_int command_queue_enqueue(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events);
+
+cl_int command_queue_enqueue_fakeevent(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events, cl_event fake_event);
+
+cl_event enqueueBarrier(cl_command_queue cq);
+
+#endif /* SOCl_COMMAND_QUEUE_H */

socl/src/helper_debug.c.inc

@@ -14,9 +14,12 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
-#include <stdio.h>
+#ifndef SOCL_DEBUG_H
+#define SOCL_DEBUG_H
 
-#ifdef DEBUG
+#ifdef STARPU_VERBOSE
+#define DEBUG
+#include <stdio.h>
    #define DEBUG_MSG(...) do { fprintf(stderr, "[SOCL] [%s] ", __func__); fprintf(stderr, __VA_ARGS__); } while (0);
    #define DEBUG_MSG_NOHEAD(...) fprintf(stderr, __VA_ARGS__)
    #define DEBUG_ERROR(...) do { fprintf(stderr, "[SOCL] ERROR: "__VA_ARGS__); exit(1); } while (0);
@@ -31,7 +34,7 @@
 #define ERROR_MSG_NOHEAD(...) fprintf(stderr, __VA_ARGS__)
 #define ERROR_STOP(...) do { ERROR_MSG(__VA_ARGS__); exit(1); } while(0);
 
-#ifdef DEBUG
+#ifdef STARPU_VERBOSE
 void DEBUG_CL(char *s, cl_int err) {
    #define ERR_CASE(a) case a: DEBUG_MSG("[OpenCL] %s CL error: %s\n", s, #a); break;
    switch(err) {
@@ -91,3 +94,5 @@ void DEBUG_CL(char *s, cl_int err) {
 #else
    #define DEBUG_CL(...) while(0);
 #endif
+
+#endif /* SOCL_DEBUG_H */

socl/src/device_descriptions.c.inc

@@ -14,48 +14,15 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
-// OpenCL 1.0 : Mandatory format: major_number.minor_number
-static const char SOCL_DRIVER_VERSION[] = "0.1";
+#include "socl.h"
+#include "devices.h"
 
-static const cl_uint SOCL_DEVICE_VENDOR_ID = 666;
+// OpenCL 1.0 : Mandatory format: major_number.minor_number
+const char * SOCL_DRIVER_VERSION = "0.1";
 
-struct _cl_device_id {
-   cl_device_type    type;
-   cl_uint           max_compute_units;        //OpenCL 1.0: minimum value is 1
-   cl_uint           max_work_item_dimensions; //OpenCL 1.0: minimum value is 3
-   size_t            max_work_item_sizes[3];   //array size should be set accordingly to the maximum max_work_item_dimensions
-   size_t            max_work_group_size;      //OpenCL 1.0: minimum value is 1
-   cl_uint           preferred_vector_widths[6]; //Char, Short, Int, Long, Float, Double
-   cl_uint           max_clock_frequency;
-   cl_uint           address_bits;             //OpenCL 1.0: 32 or 64
-   cl_ulong          max_mem_alloc_size;       //OpenCL 1.0: minimum value is max(CL_DEVICE_GLOBAL_MEM_SIZE/4, 128*1024*1024)
-   cl_bool           image_support;
-   //image fields not present
-   size_t            max_parameter_size;       //OpenCL 1.0: minimum is 256
-   cl_uint           mem_base_addr_align;
-   cl_uint           min_data_type_align_size;
-   cl_device_fp_config single_fp_config;       //OpenCL 1.0: CL_FP_ROUND_TO_NEAREST and CL_FP_INF_NAN are mandatory
-   cl_device_mem_cache_type global_mem_cache_type;
-   cl_uint           global_mem_cacheline_size;
-   cl_ulong          global_mem_cache_size;
-   cl_ulong          global_mem_size;
-   cl_ulong          max_constant_buffer_size; //OpenCL 1.0: minimum value is 64KB
-   cl_uint           max_constant_args;        //OpenCL 1.0: minimum value is 8
-   cl_device_local_mem_type local_mem_type;
-   cl_ulong          local_mem_size;           //OpenCL 1.0: minimum value is 16KB
-   cl_bool           error_correction_support;
-   size_t            profiling_timer_resolution;
-   cl_bool           endian_little;
-   cl_bool           available;
-   cl_bool           compiler_available;       //OpenCL 1.0: mandatory for FULL_PROFILE platforms
-   cl_device_exec_capabilities execution_capabilities; //OpenCL 1.0: CL_EXEC_KERNEL is mandatory
-   cl_command_queue_properties queue_properties; //OpenCL 1.0: CL_QUEUE_PROFILING_ENABLE is mandatory
-   char              name[40];                 //Array size has been arbitrarily defined
-   //versions, profile and vendor are statically defined for all devices
-   char              extensions[100];          //Array size has been arbitrarily defined
-};
+const cl_uint SOCL_DEVICE_VENDOR_ID = 666;
 
-static const struct _cl_device_id socl_devices[] = {
+const struct _cl_device_id socl_devices[] = {
    { 
       .type = CL_DEVICE_TYPE_CPU,
       .max_compute_units = 1,
@@ -315,7 +282,7 @@ static const struct _cl_device_id socl_devices[] = {
 
 };
 
-static const int socl_device_count = sizeof(socl_devices) / sizeof(struct _cl_device_id);
+const int socl_device_count = sizeof(socl_devices) / sizeof(struct _cl_device_id);
 
 
 

socl/src/devices.h

@@ -0,0 +1,65 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_DEVICES_H
+#define SOCL_DEVICES_H
+
+// OpenCL 1.0 : Mandatory format: major_number.minor_number
+const char * SOCL_DRIVER_VERSION;
+
+const cl_uint SOCL_DEVICE_VENDOR_ID;
+
+struct _cl_device_id {
+   cl_device_type    type;
+   cl_uint           max_compute_units;        //OpenCL 1.0: minimum value is 1
+   cl_uint           max_work_item_dimensions; //OpenCL 1.0: minimum value is 3
+   size_t            max_work_item_sizes[3];   //array size should be set accordingly to the maximum max_work_item_dimensions
+   size_t            max_work_group_size;      //OpenCL 1.0: minimum value is 1
+   cl_uint           preferred_vector_widths[6]; //Char, Short, Int, Long, Float, Double
+   cl_uint           max_clock_frequency;
+   cl_uint           address_bits;             //OpenCL 1.0: 32 or 64
+   cl_ulong          max_mem_alloc_size;       //OpenCL 1.0: minimum value is max(CL_DEVICE_GLOBAL_MEM_SIZE/4, 128*1024*1024)
+   cl_bool           image_support;
+   //image fields not present
+   size_t            max_parameter_size;       //OpenCL 1.0: minimum is 256
+   cl_uint           mem_base_addr_align;
+   cl_uint           min_data_type_align_size;
+   cl_device_fp_config single_fp_config;       //OpenCL 1.0: CL_FP_ROUND_TO_NEAREST and CL_FP_INF_NAN are mandatory
+   cl_device_mem_cache_type global_mem_cache_type;
+   cl_uint           global_mem_cacheline_size;
+   cl_ulong          global_mem_cache_size;
+   cl_ulong          global_mem_size;
+   cl_ulong          max_constant_buffer_size; //OpenCL 1.0: minimum value is 64KB
+   cl_uint           max_constant_args;        //OpenCL 1.0: minimum value is 8
+   cl_device_local_mem_type local_mem_type;
+   cl_ulong          local_mem_size;           //OpenCL 1.0: minimum value is 16KB
+   cl_bool           error_correction_support;
+   size_t            profiling_timer_resolution;
+   cl_bool           endian_little;
+   cl_bool           available;
+   cl_bool           compiler_available;       //OpenCL 1.0: mandatory for FULL_PROFILE platforms
+   cl_device_exec_capabilities execution_capabilities; //OpenCL 1.0: CL_EXEC_KERNEL is mandatory
+   cl_command_queue_properties queue_properties; //OpenCL 1.0: CL_QUEUE_PROFILING_ENABLE is mandatory
+   char              name[40];                 //Array size has been arbitrarily defined
+   //versions, profile and vendor are statically defined for all devices
+   char              extensions[100];          //Array size has been arbitrarily defined
+};
+
+const struct _cl_device_id socl_devices[100];
+
+const int socl_device_count;
+
+#endif /* SOCL_DEVICES_H */

socl/src/cl_releaseevent.c.inc

@@ -14,6 +14,33 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "event.h"
+#include "gc.h"
+
+static void release_callback_event(void * e);
+
+/**
+ * Create a new event
+ *
+ * Events have one-to-one relation with tag. Tag number is event ID
+ */
+cl_event event_create(void) {
+   static int id = 1;
+   cl_event ev;
+   ev = gc_entity_alloc(sizeof(struct _cl_event), release_callback_event);
+
+   ev->next = NULL;
+   ev->prev = NULL;
+   ev->id = __sync_fetch_and_add(&id,1);
+   ev->status = CL_SUBMITTED;
+   ev->type = 0;
+   ev->profiling_info = NULL;
+   ev->cq = NULL;
+
+   return ev;
+}
+
 static void release_callback_event(void * e) {
   cl_event event = (cl_event)e;
 
@@ -47,13 +74,3 @@ static void release_callback_event(void * e) {
   //starpu_tag_remove(event->id);
 }
 
-CL_API_ENTRY cl_int CL_API_CALL
-soclReleaseEvent(cl_event event) CL_API_SUFFIX__VERSION_1_0
-{
-  if (event == NULL)
-    return CL_INVALID_EVENT;
-
-  gc_entity_release(event);
-
-  return CL_SUCCESS;
-}

socl/src/event.h

@@ -0,0 +1,29 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_EVENT_H
+#define SOCL_EVENT_H
+
+#include "socl.h"
+
+/**
+ * Create a new event
+ *
+ * Events have one-to-one relation with tag. Tag number is event ID
+ */
+cl_event event_create(void);
+
+#endif /* SOCL_EVENT_H */

socl/src/gc.c.inc

@@ -14,6 +14,12 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "gc.h"
+#include "event.h"
+#include "socl.h"
+
+#include <stdlib.h>
+
 /**
  * Garbage collection thread
  */
@@ -83,12 +89,12 @@ static void * gc_thread_routine(void *UNUSED(arg)) {
 static pthread_t gc_thread;
 
 /* Start garbage collection */
-static void gc_start(void) {
+void gc_start(void) {
   pthread_create(&gc_thread, NULL, gc_thread_routine, NULL);
 }
 
 /* Stop garbage collection */
-static void gc_stop(void) {
+void gc_stop(void) {
   GC_LOCK;
 
   gc_stop_required = 1;
@@ -98,12 +104,7 @@ static void gc_stop(void) {
   pthread_join(gc_thread, NULL);
 }
 
-/**
- * Decrement reference counter and release entity if applicable
- */
-#define gc_entity_release(a) gc_entity_release_ex(&(a)->_entity)
-
-static void gc_entity_release_ex(entity e) {
+void gc_entity_release_ex(entity e) {
 
   /* Decrement reference count */
   int refs = __sync_sub_and_fetch(&e->refs, 1);
@@ -133,7 +134,7 @@ static void gc_entity_release_ex(entity e) {
 /**
  * Initialize entity
  */
-static void gc_entity_init(void *arg, void (*release_callback)(void*)) {
+void gc_entity_init(void *arg, void (*release_callback)(void*)) {
   struct entity * e = (entity)arg;
 
   e->refs = 1;
@@ -153,34 +154,30 @@ static void gc_entity_init(void *arg, void (*release_callback)(void*)) {
 /**
  * Allocate and initialize entity
  */
-static void * gc_entity_alloc(unsigned int size, void (*release_callback)(void*)) {
+void * gc_entity_alloc(unsigned int size, void (*release_callback)(void*)) {
   void * e = malloc(size);
   gc_entity_init(e, release_callback);
   return e;
 }
 
-/**
- * Retain entity
- */
-static void gc_entity_retain(void *arg) {
+/** Retain entity */
+void gc_entity_retain(void *arg) {
+	struct entity * e = (entity)arg;
 
-  struct entity * e = (entity)arg;
-
-  __sync_fetch_and_add(&e->refs, 1);
+	__sync_fetch_and_add(&e->refs, 1);
 }
 
+int gc_active_entity_count(void) {
+	int i = 0;
 
-#define gc_entity_store(dest,e) \
-  do {\
-    gc_entity_retain(e); \
-    *dest = e;\
-  } while(0);
+	entity e = entities;
+	while (e != NULL) {
+		i++;
+		e = e->next;
+	}
 
-#define gc_entity_unstore(dest) \
-  do {\
-    gc_entity_release(*dest); \
-    *dest = NULL;\
-  } while(0);
+	return i;
+}
 
 #undef GC_LOCK
 #undef GC_UNLOCK

socl/src/gc.h

@@ -0,0 +1,50 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_GC_H
+#define SOCL_GC_H
+
+#include "socl.h"
+
+void gc_start(void);
+void gc_stop(void);
+
+void gc_entity_init(void *arg, void (*release_callback)(void*));
+void * gc_entity_alloc(unsigned int size, void (*release_callback)(void*));
+void gc_entity_retain(void *arg);
+
+/** Decrement reference counter and release entity if applicable */
+void gc_entity_release_ex(entity e);
+
+int gc_active_entity_count(void);
+
+#define gc_entity_release(a) gc_entity_release_ex(&(a)->_entity)
+
+#define gc_entity_store(dest,e) \
+  do {\
+    gc_entity_retain(e); \
+    *dest = e;\
+  } while(0);
+
+#define gc_entity_unstore(dest) \
+  do {\
+    gc_entity_release(*dest); \
+    *dest = NULL;\
+  } while(0);
+
+
+
+#endif

socl/src/helper_getinfo.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#ifndef SOCL_GETINFO_H
+#define SOCL_GETINFO_H
 
 #define INFO_CASE_EX2(var) if (param_value != NULL) { \
       if (param_value_size < sizeof(var)) \
@@ -42,3 +44,5 @@
    if (param_value_size_ret != NULL) \
       *param_value_size_ret = size; \
    break;
+
+#endif /* SOCL_GETINFO_H */

socl/src/graph.c

@@ -0,0 +1,123 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "socl.h"
+#include "graph.h"
+#include "event.h"
+
+static pthread_spinlock_t graph_lock;
+static graph_node graph_nodes = NULL;
+
+
+/**
+ * Initialize graph structure
+ */
+void graph_init(void) {
+	pthread_spin_init(&graph_lock, PTHREAD_PROCESS_PRIVATE);
+}
+
+/**
+ * Release graph structure
+ */
+void graph_destroy(void) {
+	pthread_spin_destroy(&graph_lock);
+}
+
+/**
+ * Initialize a graph node
+ */
+void graph_node_init(graph_node node) {
+	node->id = -1;
+	node->next = NULL;
+}
+
+/**
+ * Store a node in the graph
+ */
+void graph_store(void * node) {
+	pthread_spin_lock(&graph_lock);
+
+	graph_node n = (graph_node)node;
+	n->next = graph_nodes;
+	graph_nodes = n;
+
+	pthread_spin_unlock(&graph_lock);
+}
+
+
+
+/**
+ * Duplicate a memory area into a fresh allocated buffer
+ */
+static void * memdupa(const void *p, size_t size) {
+	void * s = malloc(size);
+	memcpy(s,p,size);
+	return s;
+}
+
+#define memdup(p, size) ((typeof(p))memdupa(p,size))
+#define nullOrDup(name,size) s->name = (name == NULL ? NULL : memdup(name,size))
+#define dup(name) s->name = name
+
+
+node_enqueue_kernel graph_create_enqueue_kernel(char is_task,
+		cl_command_queue cq,
+		cl_kernel        kernel,
+		cl_uint          work_dim,
+		const size_t *   global_work_offset,
+		const size_t *   global_work_size,
+		const size_t *   local_work_size,
+		cl_uint          num_events,
+		const cl_event * events,
+		cl_event *       event,
+		cl_uint 		num_args,
+		size_t *		arg_sizes,
+		enum kernel_arg_type * arg_types,
+		void **		args)
+{
+	node_enqueue_kernel s = malloc(sizeof(struct node_enqueue_kernel_t));
+	graph_node_init(&s->node);
+	s->node.id = NODE_ENQUEUE_KERNEL;
+
+	dup(is_task);
+	dup(cq);
+	dup(kernel);
+	dup(work_dim);
+	nullOrDup(global_work_offset, work_dim*sizeof(size_t));
+	nullOrDup(global_work_size, work_dim*sizeof(size_t));
+	nullOrDup(local_work_size, work_dim*sizeof(size_t));
+	dup(num_events);
+	nullOrDup(events, num_events * sizeof(cl_event));
+	dup(num_args);
+	nullOrDup(arg_sizes, num_args * sizeof(size_t));
+	nullOrDup(arg_types, num_args * sizeof(enum kernel_arg_type));
+	nullOrDup(args, num_args * sizeof(void*));
+
+	
+	if (event != NULL) {
+		*event = event_create();
+		s->event = event;
+	}
+	else {
+		s->event = NULL;
+	}
+
+	return s;
+}
+
+#undef nullOrDup
+#undef memdup
+#undef dup

socl/src/graph.h

@@ -0,0 +1,73 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_GRAPH_H
+#define SOCL_GRAPH_H
+
+#include "socl.h"
+
+typedef struct graph_node_t * graph_node;
+
+struct graph_node_t {
+	int id; /* Kind of node */
+	graph_node next; /* Linked-list of nodes... */
+};
+
+void graph_init(void);
+void graph_destroy(void);
+void graph_node_init(graph_node node);
+void graph_store(void * node);
+
+#define NODE_ENQUEUE_KERNEL 1
+
+
+typedef struct node_enqueue_kernel_t {
+	struct graph_node_t node;
+
+	char 		 is_task; /* Set if clEnqueueTask is used */
+	cl_command_queue cq;
+	cl_kernel        kernel;
+	cl_uint          work_dim;
+	const size_t *   global_work_offset;
+	const size_t *   global_work_size;
+	const size_t *   local_work_size;
+	cl_uint          num_events;
+	const cl_event * events;
+	cl_event * 	 event;
+	cl_uint 	 num_args;
+	size_t *	 arg_sizes;
+	enum kernel_arg_type * arg_types;
+	void **		 args;
+} * node_enqueue_kernel;
+
+node_enqueue_kernel graph_create_enqueue_kernel(char is_task,
+		cl_command_queue cq,
+		cl_kernel        kernel,
+		cl_uint          work_dim,
+		const size_t *   global_work_offset,
+		const size_t *   global_work_size,
+		const size_t *   local_work_size,
+		cl_uint          num_events,
+		const cl_event * events,
+		cl_event *       event,
+		cl_uint 		num_args,
+		size_t *		arg_sizes,
+		enum kernel_arg_type * arg_types,
+		void **		args);
+
+cl_int node_play_enqueue_kernel(node_enqueue_kernel n);
+
+#endif /* SOCL_GRAPH_H */

socl/src/init.c.inc

@@ -14,16 +14,22 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+#include "graph.h"
+#include "gc.h"
+#include "mem_objects.h"
+
 /**
  * Initialize SOCL
  */
 __attribute__((constructor)) static void socl_init() {
   
   mem_object_init();
+  graph_init();
 
   starpu_init(NULL);
   
-  //Disable dataflow implicit dependencies
+  /* Disable dataflow implicit dependencies */
   starpu_data_set_default_sequential_consistency_flag(0);
 
   gc_start();
@@ -40,14 +46,12 @@ __attribute__((destructor)) static void socl_shutdown() {
 
   starpu_task_wait_for_all();
 
-  int i = 0;
-  while (entities != NULL) {
-    i++;
-    entities = entities->next;
-  }
+  int active_entities = gc_active_entity_count();
+
+  if (active_entities != 0)
+    fprintf(stderr, "Unreleased entities: %d\n", active_entities);
 
-  if (i != 0)
-    fprintf(stderr, "Unreleased entities: %d\n", i);
+  graph_destroy();
 
   starpu_shutdown();
 }

socl/src/helper_mem_objects.c.inc

@@ -14,6 +14,9 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
+
 #define mem_object_hash_key 257
 
 static cl_mem p_mem_objects[mem_object_hash_key] = {NULL};
@@ -22,7 +25,7 @@ static volatile pthread_spinlock_t p_mem_objects_spinlock[mem_object_hash_key];
 #define LOCK(i) pthread_spin_lock(&p_mem_objects_spinlock[i]);
 #define UNLOCK(i) pthread_spin_unlock(&p_mem_objects_spinlock[i]);
 
-static void mem_object_init(void) {
+void mem_object_init(void) {
   int i;
   for (i=0; i<mem_object_hash_key; i++) {
     pthread_spin_init(&p_mem_objects_spinlock[i], 0);
@@ -36,7 +39,7 @@ static int mem_object_hash(const void * addr) {
   return (int)t3;
 }
 
-static void mem_object_store(cl_mem m) {
+void mem_object_store(cl_mem m) {
    int hash = mem_object_hash(m);
 
    LOCK(hash);
@@ -50,7 +53,7 @@ static void mem_object_store(cl_mem m) {
    UNLOCK(hash);
 }
 
-static void mem_object_release(cl_mem m) {
+void mem_object_release(cl_mem m) {
 
    int hash = mem_object_hash(m);
 
@@ -68,7 +71,7 @@ static void mem_object_release(cl_mem m) {
    UNLOCK(hash)
 }
 
-static cl_mem mem_object_fetch(const void * addr) {
+cl_mem mem_object_fetch(const void * addr) {
    int hash = mem_object_hash(*(cl_mem*)addr);
 
    LOCK(hash);

socl/src/mem_objects.h

@@ -0,0 +1,25 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_MEM_OBJECTS_H
+#define SOCL_MEM_OBJECTS_H
+
+void mem_object_init(void);
+void mem_object_store(cl_mem m);
+void mem_object_release(cl_mem m);
+cl_mem mem_object_fetch(const void * addr);
+
+#endif /* SOCL_MEM_OBJECTS_H */

socl/src/opencl.c

@@ -1,343 +0,0 @@
-/* StarPU --- Runtime system for heterogeneous multicore architectures.
- *
- * Copyright (C) 2010,2011 University of Bordeaux
- *
- * StarPU is free software; you can redistribute it and/or modify
- * it under the terms of the GNU Lesser General Public License as published by
- * the Free Software Foundation; either version 2.1 of the License, or (at
- * your option) any later version.
- *
- * StarPU is distributed in the hope that it will be useful, but
- * WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
- *
- * See the GNU Lesser General Public License in COPYING.LGPL for more details.
- */
-
-
-#ifndef CL_HEADERS
-#include "CL/cl.h"
-#else
-#include CL_HEADERS "CL/cl.h"
-#endif
-
-
-#include <string.h>
-#include <stdlib.h>
-#include <stdint.h>
-#include <unistd.h>
-#include <pthread.h>
-
-#include <starpu.h>
-#include <starpu_opencl.h>
-#include <starpu_data_interfaces.h>
-#include <starpu_profiling.h>
-#include <starpu_task.h>
-
-typedef struct starpu_task starpu_task;
-
-#ifdef UNUSED
-#elif defined(__GNUC__)
-   #define UNUSED(x) UNUSED_ ## x __attribute__((unused))
-#else
-   #define UNUSED(x) x
-#endif
-
-#define RETURN_EVENT(ev, event) \
-   if (event != NULL) \
-      *event = ev; \
-   else\
-      gc_entity_release(ev);
-
-#include "helper_debug.c.inc"
-#include "helper_getinfo.c.inc"
-
-/**
- * Entity that can be managed by the garbage collector
- */
-typedef struct entity * entity;
-
-struct entity {
-  /* Reference count */
-  size_t refs;
-
-  /* Callback called on release */
-  void (*release_callback)(void*entity);
-
-  /* Next entity in garbage collector queue */
-  entity prev;
-  entity next;
-};
-
-/* OpenCL entities (context, command queues, buffers...) must use
- * this macro as their first field */
-#define CL_ENTITY struct entity _entity;
-
-
-struct _cl_platform_id {};
-
-static struct _cl_platform_id socl_platform = {};
-
-static const char SOCL_PROFILE[] = "FULL_PROFILE";
-static const char SOCL_VERSION[] = "OpenCL 1.0 StarPU Edition (0.0.1)";
-static const char SOCL_PLATFORM_NAME[]    = "StarPU Platform";
-static const char SOCL_VENDOR[]  = "INRIA";
-static const char SOCL_PLATFORM_EXTENSIONS[] = "";
-
-struct _cl_context {
-  CL_ENTITY;
-
-  void (*pfn_notify)(const char *, const void *, size_t, void *);
-  void *user_data;
-
-  /* Associated devices */
-  cl_device_id * devices;
-  cl_uint num_devices;
-
-  /* Properties */
-  cl_context_properties * properties;
-  cl_uint num_properties;
-
-  /* ID  */
-#ifdef DEBUG
-  int id;
-#endif
-};
-
-
-struct _cl_command_queue {
-  CL_ENTITY;
-
-  cl_command_queue_properties properties;
-  cl_device_id device;
-  cl_context context;
-
-  /* Stored command events */
-  cl_event events;
-
-  /* Last enqueued barrier-like event */
-  cl_event barrier;
-
-  /* Mutex */
-  pthread_spinlock_t spin;
-
-  /* ID  */
-#ifdef DEBUG
-  int id;
-#endif
-};
-
-struct _cl_event {
-  CL_ENTITY;
-
-  /* Command queue */
-  cl_command_queue cq;
-
-  /* Command type */
-  cl_command_type type;
-
-  /* Command queue list */
-  cl_event prev;
-  cl_event next;
-
-  /* Event status */
-  cl_int status;
-
-  /* ID  
-   * This ID is used as a tag for StarPU dependencies
-   */
-  int id;
-
-  /* Profiling info are copied here */
-  struct starpu_task_profiling_info *profiling_info;
-};
-
-struct _cl_mem {
-  CL_ENTITY;
-
-  /* StarPU handle */
-  starpu_data_handle handle;
-
-  /* Pointer to data in host memory */
-  void *ptr;    
-
-  /* Buffer size */
-  size_t size;
-
-  /* Indicates how many references (mapping, MEM_USE_HOST_PTR...) require
-   * coherence in host memory. If set to zero, no coherency is maintained
-   * (this is the most efficient) */
-  int map_count; 
-
-  /* Creation flags */
-  cl_mem_flags flags;
-
-  /* Creation context */
-  cl_context context;
-
-  /* Access mode */
-  int mode;
-
-  /* Host ptr */
-  void * host_ptr;
-
-  /* Fields used to store cl_mems in mem_objects list */
-  cl_mem prev;
-  cl_mem next;
-
-  /* Indicates if a buffer may contain meaningful data. Otherwise
-     we don't have to transfer it */
-  int scratch;
-
-  /* ID  */
-#ifdef DEBUG
-  int id;
-#endif
-};
-
-struct _cl_program {
-  CL_ENTITY;
-
-  /* Real OpenCL Programs
-   * There is one entry for each device (even non OpenCL ones)
-   * in order to index this array with dev_id
-   */
-  cl_program *cl_programs;
-
-  /* Context used to create this program */
-  cl_context context;
-
-  /* Options  */
-  char * options;
-  unsigned int options_size;
-
-  /* ID  */
-#ifdef DEBUG
-  int id;
-#endif
-};
-
-enum kernel_arg_type { Null, Buffer, Immediate };
-
-struct _cl_kernel {
-  CL_ENTITY;
-
-  /* Associated program */
-  cl_program program;
-
-  /* Kernel name */
-  char * kernel_name;
-
-  /* Real OpenCL kernels */
-  cl_kernel *cl_kernels;
-
-  /* clCreateKernel return codes */
-  cl_int *errcodes;
-
-  /* Arguments */
-  unsigned int arg_count;
-  size_t *arg_size;
-  enum kernel_arg_type  *arg_type;
-  void  **arg_value;
-
-  /* ID  */
-#ifdef DEBUG
-  int id;
-#endif
-};
-
-/* Command queues with profiling enabled
- * This allows us to disable StarPU profiling it
- * is equal to 0
- */
-static int profiling_queue_count = 0;
-
-#include "helper_workerid.c.inc"
-
-#include "gc.c.inc"
-
-#include "cl_getplatformids.c.inc"
-#include "cl_getplatforminfo.c.inc"
-
-#include "device_descriptions.c.inc"
-#include "cl_getdeviceids.c.inc"
-#include "cl_getdeviceinfo.c.inc"
-
-#include "cl_releasecontext.c.inc"
-#include "cl_createcontext.c.inc"
-#include "cl_createcontextfromtype.c.inc"
-#include "cl_retaincontext.c.inc"
-#include "cl_getcontextinfo.c.inc"
-
-#include "cl_releasecommandqueue.c.inc"
-#include "cl_createcommandqueue.c.inc"
-#include "cl_retaincommandqueue.c.inc"
-#include "cl_getcommandqueueinfo.c.inc"
-#include "cl_setcommandqueueproperty.c.inc"
-
-#include "cl_releaseevent.c.inc"
-#include "helper_event.c.inc"
-#include "helper_task.c.inc"
-#include "cl_waitforevents.c.inc"
-#include "cl_geteventinfo.c.inc"
-#include "cl_retainevent.c.inc"
-
-#include "helper_command_queue.c.inc"
-
-#include "cl_enqueuemarker.c.inc"
-#include "cl_enqueuewaitforevents.c.inc"
-#include "cl_enqueuebarrier.c.inc"
-#include "cl_flush.c.inc"
-#include "cl_finish.c.inc"
-
-#include "helper_mem_objects.c.inc"
-#include "cl_releasememobject.c.inc"
-#include "cl_createbuffer.c.inc"
-#include "cl_createimage2d.c.inc"
-#include "cl_createimage3d.c.inc"
-#include "cl_retainmemobject.c.inc"
-#include "cl_getsupportedimageformats.c.inc"
-#include "cl_getmemobjectinfo.c.inc"
-#include "cl_getimageinfo.c.inc"
-
-#include "cl_createsampler.c.inc"
-#include "cl_retainsampler.c.inc"
-#include "cl_releasesampler.c.inc"
-#include "cl_getsamplerinfo.c.inc"
-
-#include "cl_releaseprogram.c.inc"
-#include "cl_createprogramwithsource.c.inc"
-#include "cl_createprogramwithbinary.c.inc"
-#include "cl_retainprogram.c.inc"
-#include "cl_buildprogram.c.inc"
-#include "cl_unloadcompiler.c.inc"
-#include "cl_getprograminfo.c.inc"
-#include "cl_getprogrambuildinfo.c.inc"
-
-#include "cl_releasekernel.c.inc"
-#include "cl_createkernel.c.inc"
-#include "cl_createkernelsinprogram.c.inc"
-#include "cl_retainkernel.c.inc"
-#include "cl_setkernelarg.c.inc"
-#include "cl_getkernelinfo.c.inc"
-#include "cl_getkernelworkgroupinfo.c.inc"
-
-#include "cl_enqueuereadbuffer.c.inc"
-#include "cl_enqueuewritebuffer.c.inc"
-#include "cl_enqueuecopybuffer.c.inc"
-#include "cl_enqueuereadimage.c.inc"
-#include "cl_enqueuewriteimage.c.inc"
-#include "cl_enqueuecopyimage.c.inc"
-#include "cl_enqueuecopyimagetobuffer.c.inc"
-#include "cl_enqueuecopybuffertoimage.c.inc"
-#include "cl_enqueuemapbuffer.c.inc"
-#include "cl_enqueuemapimage.c.inc"
-#include "cl_enqueueunmapmemobject.c.inc"
-#include "cl_enqueuetask.c.inc"
-#include "cl_enqueuendrangekernel.c.inc"
-#include "cl_enqueuenativekernel.c.inc"
-
-#include "cl_geteventprofilinginfo.c.inc"
-#include "cl_getextensionfunctionaddress.c.inc"
-
-#include "init.c.inc"

socl/src/helper_event.c.inc

@@ -14,24 +14,19 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
-/**
- * Create a new event
- *
- * Events have one-to-one relation with tag. Tag number is event ID
- */
-static cl_event event_create(void) {
-   static int id = 1;
-   cl_event ev;
-   ev = gc_entity_alloc(sizeof(struct _cl_event), release_callback_event);
+#include "socl.h"
 
-   ev->next = NULL;
-   ev->prev = NULL;
-   ev->id = __sync_fetch_and_add(&id,1);
-   ev->status = CL_SUBMITTED;
-   ev->type = 0;
-   ev->profiling_info = NULL;
-   ev->cq = NULL;
+struct _cl_platform_id socl_platform = {};
 
-   return ev;
-}
+const char * SOCL_PROFILE = "FULL_PROFILE";
+const char * SOCL_VERSION = "OpenCL 1.0 StarPU Edition (0.0.1)";
+const char * SOCL_PLATFORM_NAME    = "StarPU Platform";
+const char * SOCL_VENDOR  = "INRIA";
+const char * SOCL_PLATFORM_EXTENSIONS = "";
 
+
+/* Command queues with profiling enabled
+ * This allows us to disable StarPU profiling it
+ * is equal to 0
+ */
+int profiling_queue_count = 0;

socl/src/socl.h

@@ -0,0 +1,730 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_H
+#define SOCL_H
+
+#ifndef CL_HEADERS
+#include "CL/cl.h"
+#else
+#include CL_HEADERS "CL/cl.h"
+#endif
+
+#include <string.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <unistd.h>
+#include <pthread.h>
+
+#include <starpu.h>
+#include <starpu_opencl.h>
+#include <starpu_data_interfaces.h>
+#include <starpu_profiling.h>
+#include <starpu_task.h>
+
+typedef struct starpu_task starpu_task;
+
+#ifdef UNUSED
+#elif defined(__GNUC__)
+   #define UNUSED(x) UNUSED_ ## x __attribute__((unused))
+#else
+   #define UNUSED(x) x
+#endif
+
+
+/**
+ * Entity that can be managed by the garbage collector
+ */
+typedef struct entity * entity;
+
+#include "command_queue.h"
+#include "debug.h"
+#include "devices.h"
+#include "event.h"
+#include "gc.h"
+#include "graph.h"
+#include "mem_objects.h"
+#include "task.h"
+#include "util.h"
+
+
+struct entity {
+  /* Reference count */
+  size_t refs;
+
+  /* Callback called on release */
+  void (*release_callback)(void*entity);
+
+  /* Next entity in garbage collector queue */
+  entity prev;
+  entity next;
+};
+
+/* OpenCL entities (context, command queues, buffers...) must use
+ * this macro as their first field */
+#define CL_ENTITY struct entity _entity;
+
+struct _cl_platform_id {};
+
+#define RETURN_EVENT(ev, event) \
+   if (event != NULL) \
+      *event = ev; \
+   else\
+      gc_entity_release(ev);
+
+/* Constants */
+struct _cl_platform_id socl_platform;
+const char * SOCL_PROFILE;
+const char * SOCL_VERSION;
+const char * SOCL_PLATFORM_NAME;
+const char * SOCL_VENDOR;
+const char * SOCL_PLATFORM_EXTENSIONS;
+
+struct _cl_context {
+  CL_ENTITY;
+
+  void (*pfn_notify)(const char *, const void *, size_t, void *);
+  void *user_data;
+
+  /* Associated devices */
+  cl_device_id * devices;
+  cl_uint num_devices;
+
+  /* Properties */
+  cl_context_properties * properties;
+  cl_uint num_properties;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+
+struct _cl_command_queue {
+  CL_ENTITY;
+
+  cl_command_queue_properties properties;
+  cl_device_id device;
+  cl_context context;
+
+  /* Stored command events */
+  cl_event events;
+
+  /* Last enqueued barrier-like event */
+  cl_event barrier;
+
+  /* Mutex */
+  pthread_spinlock_t spin;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+struct _cl_event {
+  CL_ENTITY;
+
+  /* Command queue */
+  cl_command_queue cq;
+
+  /* Command type */
+  cl_command_type type;
+
+  /* Command queue list */
+  cl_event prev;
+  cl_event next;
+
+  /* Event status */
+  cl_int status;
+
+  /* ID  
+   * This ID is used as a tag for StarPU dependencies
+   */
+  int id;
+
+  /* Profiling info are copied here */
+  struct starpu_task_profiling_info *profiling_info;
+};
+
+struct _cl_mem {
+  CL_ENTITY;
+
+  /* StarPU handle */
+  starpu_data_handle handle;
+
+  /* Pointer to data in host memory */
+  void *ptr;    
+
+  /* Buffer size */
+  size_t size;
+
+  /* Indicates how many references (mapping, MEM_USE_HOST_PTR...) require
+   * coherence in host memory. If set to zero, no coherency is maintained
+   * (this is the most efficient) */
+  int map_count; 
+
+  /* Creation flags */
+  cl_mem_flags flags;
+
+  /* Creation context */
+  cl_context context;
+
+  /* Access mode */
+  int mode;
+
+  /* Host ptr */
+  void * host_ptr;
+
+  /* Fields used to store cl_mems in mem_objects list */
+  cl_mem prev;
+  cl_mem next;
+
+  /* Indicates if a buffer may contain meaningful data. Otherwise
+     we don't have to transfer it */
+  int scratch;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+struct _cl_program {
+  CL_ENTITY;
+
+  /* Real OpenCL Programs
+   * There is one entry for each device (even non OpenCL ones)
+   * in order to index this array with dev_id
+   */
+  cl_program *cl_programs;
+
+  /* Context used to create this program */
+  cl_context context;
+
+  /* Options  */
+  char * options;
+  unsigned int options_size;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+enum kernel_arg_type { Null, Buffer, Immediate };
+
+struct _cl_kernel {
+  CL_ENTITY;
+
+  /* Associated program */
+  cl_program program;
+
+  /* Kernel name */
+  char * kernel_name;
+
+  /* Real OpenCL kernels */
+  cl_kernel *cl_kernels;
+
+  /* clCreateKernel return codes */
+  cl_int *errcodes;
+
+  /* Arguments */
+  unsigned int arg_count;
+  size_t *arg_size;
+  enum kernel_arg_type  *arg_type;
+  void  **arg_value;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+/* Global vars */
+
+/* Command queues with profiling enabled
+ * This allows us to disable StarPU profiling it
+ * is equal to 0
+ */
+int profiling_queue_count;
+
+/***************************************************************************/
+
+/* Platform API */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetPlatformIDs(cl_uint          /* num_entries */,
+                 cl_platform_id * /* platforms */,
+                 cl_uint *        /* num_platforms */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL 
+soclGetPlatformInfo(cl_platform_id   /* platform */, 
+                  cl_platform_info /* param_name */,
+                  size_t           /* param_value_size */, 
+                  void *           /* param_value */,
+                  size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Device APIs */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetDeviceIDs(cl_platform_id   /* platform */,
+               cl_device_type   /* device_type */, 
+               cl_uint          /* num_entries */, 
+               cl_device_id *   /* devices */, 
+               cl_uint *        /* num_devices */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetDeviceInfo(cl_device_id    /* device */,
+                cl_device_info  /* param_name */, 
+                size_t          /* param_value_size */, 
+                void *          /* param_value */,
+                size_t *        /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Context APIs  */
+extern CL_API_ENTRY cl_context CL_API_CALL
+soclCreateContext(const cl_context_properties * /* properties */,
+                cl_uint                       /* num_devices */,
+                const cl_device_id *          /* devices */,
+                void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */,
+                void *                        /* user_data */,
+                cl_int *                      /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_context CL_API_CALL
+soclCreateContextFromType(const cl_context_properties * /* properties */,
+                        cl_device_type                /* device_type */,
+                        void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */,
+                        void *                        /* user_data */,
+                        cl_int *                      /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainContext(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseContext(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetContextInfo(cl_context         /* context */, 
+                 cl_context_info    /* param_name */, 
+                 size_t             /* param_value_size */, 
+                 void *             /* param_value */, 
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Command Queue APIs */
+extern CL_API_ENTRY cl_command_queue CL_API_CALL
+soclCreateCommandQueue(cl_context                     /* context */, 
+                     cl_device_id                   /* device */, 
+                     cl_command_queue_properties    /* properties */,
+                     cl_int *                       /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainCommandQueue(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseCommandQueue(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetCommandQueueInfo(cl_command_queue      /* command_queue */,
+                      cl_command_queue_info /* param_name */,
+                      size_t                /* param_value_size */,
+                      void *                /* param_value */,
+                      size_t *              /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclSetCommandQueueProperty(cl_command_queue              /* command_queue */,
+                          cl_command_queue_properties   /* properties */, 
+                          cl_bool                        /* enable */,
+                          cl_command_queue_properties * /* old_properties */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Memory Object APIs  */
+extern CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateBuffer(cl_context   /* context */,
+               cl_mem_flags /* flags */,
+               size_t       /* size */,
+               void *       /* host_ptr */,
+               cl_int *     /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateImage2D(cl_context              /* context */,
+                cl_mem_flags            /* flags */,
+                const cl_image_format * /* image_format */,
+                size_t                  /* image_width */,
+                size_t                  /* image_height */,
+                size_t                  /* image_row_pitch */, 
+                void *                  /* host_ptr */,
+                cl_int *                /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+                        
+extern CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateImage3D(cl_context              /* context */,
+                cl_mem_flags            /* flags */,
+                const cl_image_format * /* image_format */,
+                size_t                  /* image_width */, 
+                size_t                  /* image_height */,
+                size_t                  /* image_depth */, 
+                size_t                  /* image_row_pitch */, 
+                size_t                  /* image_slice_pitch */, 
+                void *                  /* host_ptr */,
+                cl_int *                /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+                        
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainMemObject(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseMemObject(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetSupportedImageFormats(cl_context           /* context */,
+                           cl_mem_flags         /* flags */,
+                           cl_mem_object_type   /* image_type */,
+                           cl_uint              /* num_entries */,
+                           cl_image_format *    /* image_formats */,
+                           cl_uint *            /* num_image_formats */) CL_API_SUFFIX__VERSION_1_0;
+                                    
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetMemObjectInfo(cl_mem           /* memobj */,
+                   cl_mem_info      /* param_name */, 
+                   size_t           /* param_value_size */,
+                   void *           /* param_value */,
+                   size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetImageInfo(cl_mem           /* image */,
+               cl_image_info    /* param_name */, 
+               size_t           /* param_value_size */,
+               void *           /* param_value */,
+               size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Sampler APIs  */
+extern CL_API_ENTRY cl_sampler CL_API_CALL
+soclCreateSampler(cl_context          /* context */,
+                cl_bool             /* normalized_coords */, 
+                cl_addressing_mode  /* addressing_mode */, 
+                cl_filter_mode      /* filter_mode */,
+                cl_int *            /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainSampler(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseSampler(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetSamplerInfo(cl_sampler         /* sampler */,
+                 cl_sampler_info    /* param_name */,
+                 size_t             /* param_value_size */,
+                 void *             /* param_value */,
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                            
+/* Program Object APIs  */
+extern CL_API_ENTRY cl_program CL_API_CALL
+soclCreateProgramWithSource(cl_context        /* context */,
+                          cl_uint           /* count */,
+                          const char **     /* strings */,
+                          const size_t *    /* lengths */,
+                          cl_int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_program CL_API_CALL
+soclCreateProgramWithBinary(cl_context                     /* context */,
+                          cl_uint                        /* num_devices */,
+                          const cl_device_id *           /* device_list */,
+                          const size_t *                 /* lengths */,
+                          const unsigned char **         /* binaries */,
+                          cl_int *                       /* binary_status */,
+                          cl_int *                       /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainProgram(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseProgram(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclBuildProgram(cl_program           /* program */,
+               cl_uint              /* num_devices */,
+               const cl_device_id * /* device_list */,
+               const char *         /* options */, 
+               void (*pfn_notify)(cl_program /* program */, void * /* user_data */),
+               void *               /* user_data */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclUnloadCompiler(void) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetProgramInfo(cl_program         /* program */,
+                 cl_program_info    /* param_name */,
+                 size_t             /* param_value_size */,
+                 void *             /* param_value */,
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetProgramBuildInfo(cl_program            /* program */,
+                      cl_device_id          /* device */,
+                      cl_program_build_info /* param_name */,
+                      size_t                /* param_value_size */,
+                      void *                /* param_value */,
+                      size_t *              /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                            
+/* Kernel Object APIs */
+extern CL_API_ENTRY cl_kernel CL_API_CALL
+soclCreateKernel(cl_program      /* program */,
+               const char *    /* kernel_name */,
+               cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclCreateKernelsInProgram(cl_program     /* program */,
+                         cl_uint        /* num_kernels */,
+                         cl_kernel *    /* kernels */,
+                         cl_uint *      /* num_kernels_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainKernel(cl_kernel    /* kernel */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseKernel(cl_kernel   /* kernel */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclSetKernelArg(cl_kernel    /* kernel */,
+               cl_uint      /* arg_index */,
+               size_t       /* arg_size */,
+               const void * /* arg_value */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetKernelInfo(cl_kernel       /* kernel */,
+                cl_kernel_info  /* param_name */,
+                size_t          /* param_value_size */,
+                void *          /* param_value */,
+                size_t *        /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetKernelWorkGroupInfo(cl_kernel                  /* kernel */,
+                         cl_device_id               /* device */,
+                         cl_kernel_work_group_info  /* param_name */,
+                         size_t                     /* param_value_size */,
+                         void *                     /* param_value */,
+                         size_t *                   /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Event Object APIs  */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclWaitForEvents(cl_uint             /* num_events */,
+                const cl_event *    /* event_list */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetEventInfo(cl_event         /* event */,
+               cl_event_info    /* param_name */,
+               size_t           /* param_value_size */,
+               void *           /* param_value */,
+               size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainEvent(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseEvent(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Profiling APIs  */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetEventProfilingInfo(cl_event            /* event */,
+                        cl_profiling_info   /* param_name */,
+                        size_t              /* param_value_size */,
+                        void *              /* param_value */,
+                        size_t *            /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                                
+/* Flush and Finish APIs */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclFlush(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclFinish(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Enqueued Commands APIs */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueReadBuffer(cl_command_queue    /* command_queue */,
+                    cl_mem              /* buffer */,
+                    cl_bool             /* blocking_read */,
+                    size_t              /* offset */,
+                    size_t              /* cb */, 
+                    void *              /* ptr */,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWriteBuffer(cl_command_queue   /* command_queue */, 
+                     cl_mem             /* buffer */, 
+                     cl_bool            /* blocking_write */, 
+                     size_t             /* offset */, 
+                     size_t             /* cb */, 
+                     const void *       /* ptr */, 
+                     cl_uint            /* num_events_in_wait_list */, 
+                     const cl_event *   /* event_wait_list */, 
+                     cl_event *         /* event */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyBuffer(cl_command_queue    /* command_queue */, 
+                    cl_mem              /* src_buffer */,
+                    cl_mem              /* dst_buffer */, 
+                    size_t              /* src_offset */,
+                    size_t              /* dst_offset */,
+                    size_t              /* cb */, 
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueReadImage(cl_command_queue     /* command_queue */,
+                   cl_mem               /* image */,
+                   cl_bool              /* blocking_read */, 
+                   const size_t *       /* origin[3] */,
+                   const size_t *       /* region[3] */,
+                   size_t               /* row_pitch */,
+                   size_t               /* slice_pitch */, 
+                   void *               /* ptr */,
+                   cl_uint              /* num_events_in_wait_list */,
+                   const cl_event *     /* event_wait_list */,
+                   cl_event *           /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWriteImage(cl_command_queue    /* command_queue */,
+                    cl_mem              /* image */,
+                    cl_bool             /* blocking_write */, 
+                    const size_t *      /* origin[3] */,
+                    const size_t *      /* region[3] */,
+                    size_t              /* input_row_pitch */,
+                    size_t              /* input_slice_pitch */, 
+                    const void *        /* ptr */,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyImage(cl_command_queue     /* command_queue */,
+                   cl_mem               /* src_image */,
+                   cl_mem               /* dst_image */, 
+                   const size_t *       /* src_origin[3] */,
+                   const size_t *       /* dst_origin[3] */,
+                   const size_t *       /* region[3] */, 
+                   cl_uint              /* num_events_in_wait_list */,
+                   const cl_event *     /* event_wait_list */,
+                   cl_event *           /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyImageToBuffer(cl_command_queue /* command_queue */,
+                           cl_mem           /* src_image */,
+                           cl_mem           /* dst_buffer */, 
+                           const size_t *   /* src_origin[3] */,
+                           const size_t *   /* region[3] */, 
+                           size_t           /* dst_offset */,
+                           cl_uint          /* num_events_in_wait_list */,
+                           const cl_event * /* event_wait_list */,
+                           cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyBufferToImage(cl_command_queue /* command_queue */,
+                           cl_mem           /* src_buffer */,
+                           cl_mem           /* dst_image */, 
+                           size_t           /* src_offset */,
+                           const size_t *   /* dst_origin[3] */,
+                           const size_t *   /* region[3] */, 
+                           cl_uint          /* num_events_in_wait_list */,
+                           const cl_event * /* event_wait_list */,
+                           cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY void * CL_API_CALL
+soclEnqueueMapBuffer(cl_command_queue /* command_queue */,
+                   cl_mem           /* buffer */,
+                   cl_bool          /* blocking_map */, 
+                   cl_map_flags     /* map_flags */,
+                   size_t           /* offset */,
+                   size_t           /* cb */,
+                   cl_uint          /* num_events_in_wait_list */,
+                   const cl_event * /* event_wait_list */,
+                   cl_event *       /* event */,
+                   cl_int *         /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY void * CL_API_CALL
+soclEnqueueMapImage(cl_command_queue  /* command_queue */,
+                  cl_mem            /* image */, 
+                  cl_bool           /* blocking_map */, 
+                  cl_map_flags      /* map_flags */, 
+                  const size_t *    /* origin[3] */,
+                  const size_t *    /* region[3] */,
+                  size_t *          /* image_row_pitch */,
+                  size_t *          /* image_slice_pitch */,
+                  cl_uint           /* num_events_in_wait_list */,
+                  const cl_event *  /* event_wait_list */,
+                  cl_event *        /* event */,
+                  cl_int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueUnmapMemObject(cl_command_queue /* command_queue */,
+                        cl_mem           /* memobj */,
+                        void *           /* mapped_ptr */,
+                        cl_uint          /* num_events_in_wait_list */,
+                        const cl_event *  /* event_wait_list */,
+                        cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNDRangeKernel(cl_command_queue /* command_queue */,
+                       cl_kernel        /* kernel */,
+                       cl_uint          /* work_dim */,
+                       const size_t *   /* global_work_offset */,
+                       const size_t *   /* global_work_size */,
+                       const size_t *   /* local_work_size */,
+                       cl_uint          /* num_events_in_wait_list */,
+                       const cl_event * /* event_wait_list */,
+                       cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueTask(cl_command_queue  /* command_queue */,
+              cl_kernel         /* kernel */,
+              cl_uint           /* num_events_in_wait_list */,
+              const cl_event *  /* event_wait_list */,
+              cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNativeKernel(cl_command_queue  /* command_queue */,
+					  void (*user_func)(void *), 
+                      void *            /* args */,
+                      size_t            /* cb_args */, 
+                      cl_uint           /* num_mem_objects */,
+                      const cl_mem *    /* mem_list */,
+                      const void **     /* args_mem_loc */,
+                      cl_uint           /* num_events_in_wait_list */,
+                      const cl_event *  /* event_wait_list */,
+                      cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueMarker(cl_command_queue    /* command_queue */,
+                cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWaitForEvents(cl_command_queue /* command_queue */,
+                       cl_uint          /* num_events */,
+                       const cl_event * /* event_list */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueBarrier(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Extension function access
+ *
+ * Returns the extension function address for the given function name,
+ * or NULL if a valid function can not be found.  The client must
+ * check to make sure the address is not NULL, before using or 
+ * calling the returned function address.
+ */
+extern CL_API_ENTRY void * CL_API_CALL
+soclGetExtensionFunctionAddress(const char * /* func_name */) CL_API_SUFFIX__VERSION_1_0;
+
+#endif /* SOCL_H */

socl/src/helper_task.c.inc

@@ -14,7 +14,11 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
-static inline cl_event task_event(starpu_task *task) {
+#include "socl.h"
+#include "gc.h"
+#include "event.h"
+
+cl_event task_event(starpu_task *task) {
   return (cl_event)task->callback_arg;
 }
 
@@ -39,7 +43,7 @@ static void task_release_callback(void *arg) {
  * Task's callback_arg is event
  * Task's tag is set to event ID
  */
-static starpu_task * task_create(cl_command_type type) {
+starpu_task * task_create(cl_command_type type) {
    cl_event event;
    struct starpu_task * task;
 
@@ -64,7 +68,7 @@ static starpu_task * task_create(cl_command_type type) {
 }
 
 
-static void task_dependency_add(starpu_task * task, cl_uint num, const cl_event *events) {
+void task_dependency_add(starpu_task * task, cl_uint num, const cl_event *events) {
    unsigned int i;
 
    for (i=0; i<num; i++) {
@@ -102,7 +106,7 @@ static starpu_codelet cputask_codelet = {
    .cpu_func = &cputask_task
 };
 
-static starpu_task * task_create_cpu(cl_command_type type, void (*callback)(void*), void *arg, int free_arg) {
+starpu_task * task_create_cpu(cl_command_type type, void (*callback)(void*), void *arg, int free_arg) {
   
   struct cputask_arg * a = malloc(sizeof(struct cputask_arg));
   a->callback = callback;

socl/src/task.h

@@ -0,0 +1,27 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_TASK_H
+#define SOCL_TASK_H
+
+#include "socl.h"
+
+starpu_task * task_create(cl_command_type type);
+void task_dependency_add(starpu_task * task, cl_uint num, const cl_event *events);
+starpu_task * task_create_cpu(cl_command_type type, void (*callback)(void*), void *arg, int free_arg);
+cl_event task_event(starpu_task *task);
+
+#endif /* SOCL_TASK_H */

socl/src/helper_workerid.c.inc

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 int starpu_worker_get_range() {
    int id = starpu_worker_get_id();
    int i, oid = 0;

socl/src/util.h

@@ -0,0 +1,22 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef SOCL_UTIL_H
+#define SOCL_UTIL_H
+
+int starpu_worker_get_range();
+
+#endif /* SOCL_UTIL_H */