SOCL: commands are entities now

Barrier command has been fixed as it has been introduced in OpenCL 1.2.

socl/src/cl_enqueuebarrier.c

@@ -19,9 +19,16 @@
 CL_API_ENTRY cl_int CL_API_CALL
 soclEnqueueBarrier(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
 {
-	command_marker cmd = command_barrier_create();
+	command_barrier cmd = command_barrier_create();
 
 	command_queue_enqueue(cq, cmd, 0, NULL);
 
 	return CL_SUCCESS;
 }
+
+cl_int command_barrier_submit(command_barrier cmd) {
+	struct starpu_task *task;
+	task = task_create(CL_COMMAND_BARRIER);
+
+	return task_submit(task, cmd);
+}

socl/src/cl_finish.c

@@ -19,7 +19,7 @@
 CL_API_ENTRY cl_int CL_API_CALL
 soclFinish(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0 {
 
-	command_marker cmd = command_barrier_create();
+	command_barrier cmd = command_barrier_create();
 
 	command_queue_enqueue(cq, cmd, 0, NULL);
 

socl/src/command.c

@@ -20,7 +20,8 @@
 /* Forward extern declaration */
 extern void soclEnqueueNDRangeKernel_task(void *descr[], void *args);
 
-void command_init_ex(cl_command cmd, cl_command_type typ) {
+void command_init_ex(cl_command cmd, cl_command_type typ, void (*cb)(void*)) {
+	gc_entity_init(&cmd->_entity, cb);
 	cmd->typ = typ;
 	cmd->num_events = 0;
 	cmd->events = NULL;
@@ -48,6 +49,7 @@ void command_submit_ex(cl_command cmd) {
 		SUBMIT(CL_COMMAND_MAP_BUFFER, command_map_buffer)
 		SUBMIT(CL_COMMAND_UNMAP_MEM_OBJECT, command_unmap_mem_object)
 		SUBMIT(CL_COMMAND_MARKER, command_marker)
+		SUBMIT(CL_COMMAND_BARRIER, command_barrier)
 		default:
 			ERROR_STOP("Trying to submit unknown command (type %x)", cmd->typ);
 	}
@@ -88,8 +90,23 @@ void command_graph_dump_ex(cl_command cmd) {
 }
 
 #define nullOrDup(name,size) cmd->name = memdup_safe(name,size)
+#define nullOrFree(name) if (cmd->name != NULL) free((void*)cmd->name)
 #define dup(name) cmd->name = name
-#define dupEntity(name) do { cmd->name = name; gc_entity_retain(name); } while (0);
+
+void command_ndrange_kernel_release(void * arg) {
+	command_ndrange_kernel cmd = (command_ndrange_kernel)arg;
+
+	gc_entity_unstore(&cmd->kernel);
+	nullOrFree(global_work_offset);
+	nullOrFree(global_work_size);
+	nullOrFree(local_work_size);
+	free(cmd->arg_sizes);
+	free(cmd->arg_types);
+	unsigned int i;
+	for (i=0; i<cmd->num_args; i++)
+		free(cmd->args[i]);
+	free(cmd->args);
+}
 
 command_ndrange_kernel command_ndrange_kernel_create (
 		cl_kernel        kernel,
@@ -99,9 +116,10 @@ command_ndrange_kernel command_ndrange_kernel_create (
 		const size_t *   local_work_size)
 {
 	command_ndrange_kernel cmd = calloc(1, sizeof(struct command_ndrange_kernel_t));
-	command_init(cmd, CL_COMMAND_NDRANGE_KERNEL);
+	command_init(cmd, CL_COMMAND_NDRANGE_KERNEL, command_ndrange_kernel_release);
+
+	gc_entity_store(&cmd->kernel, kernel);
 
-	dupEntity(kernel);
 	dup(work_dim);
 	nullOrDup(global_work_offset, work_dim*sizeof(size_t));
 	nullOrDup(global_work_size, work_dim*sizeof(size_t));
@@ -139,10 +157,10 @@ command_ndrange_kernel command_task_create (cl_kernel kernel) {
 	return cmd;
 }
 
-command_marker command_barrier_create () {
+command_barrier command_barrier_create () {
 
-	command_marker cmd = malloc(sizeof(struct command_marker_t));
-	command_init(cmd, CL_COMMAND_BARRIER);
+	command_barrier cmd = malloc(sizeof(struct command_barrier_t));
+	command_init(cmd, CL_COMMAND_BARRIER, NULL);
 
 	return cmd;
 }
@@ -150,11 +168,20 @@ command_marker command_barrier_create () {
 command_marker command_marker_create () {
 
 	command_marker cmd = malloc(sizeof(struct command_marker_t));
-	command_init(cmd, CL_COMMAND_MARKER);
+	command_init(cmd, CL_COMMAND_MARKER, NULL);
 
 	return cmd;
 }
 
+void command_map_buffer_release(void * arg) {
+	command_map_buffer cmd = (command_map_buffer)arg;
+
+	/* We DO NOT unstore (release) the buffer as unmap will do it
+	  gc_entity_unstore(&cmd->buffer); */
+
+	gc_entity_unstore(&cmd->event);
+}
+
 command_map_buffer command_map_buffer_create(
 		cl_mem buffer,
 		cl_map_flags map_flags,
@@ -164,33 +191,46 @@ command_map_buffer command_map_buffer_create(
 		) {
 
 	command_map_buffer cmd = malloc(sizeof(struct command_map_buffer_t));
-	command_init(cmd, CL_COMMAND_MAP_BUFFER);
+	command_init(cmd, CL_COMMAND_MAP_BUFFER, command_map_buffer_release);
 
-	dupEntity(buffer);
+	gc_entity_store(&cmd->buffer, buffer);
 	dup(map_flags);
 	dup(offset);
 	dup(cb);
-	dupEntity(event);
+	gc_entity_store(&cmd->event, event);
 
 	return cmd;
 }
 
+void command_unmap_mem_object_release(void * arg) {
+	command_unmap_mem_object cmd = (command_unmap_mem_object)arg;
+
+	/* We release the buffer twice because map buffer command did not */
+	gc_entity_release(cmd->buffer);
+	gc_entity_unstore(&cmd->buffer);
+}
+
 command_unmap_mem_object command_unmap_mem_object_create(cl_mem buffer, void * ptr) {
 	command_unmap_mem_object cmd = malloc(sizeof(struct command_unmap_mem_object_t));
-	command_init(cmd, CL_COMMAND_UNMAP_MEM_OBJECT);
+	command_init(cmd, CL_COMMAND_UNMAP_MEM_OBJECT, command_unmap_mem_object_release);
 
-	dupEntity(buffer);
+	gc_entity_store(&cmd->buffer, buffer);
 	dup(ptr);
 
 	return cmd;
 }
 
+void command_read_buffer_release(void *arg) {
+	command_read_buffer cmd = (command_read_buffer)arg;
+	gc_entity_unstore(&cmd->buffer);
+}
+
 command_read_buffer command_read_buffer_create(cl_mem buffer, size_t offset, size_t cb, void * ptr) {
 
 	command_read_buffer cmd = malloc(sizeof(struct command_read_buffer_t));
-	command_init(cmd, CL_COMMAND_READ_BUFFER);
+	command_init(cmd, CL_COMMAND_READ_BUFFER, command_read_buffer_release);
 
-	dupEntity(buffer);
+	gc_entity_store(&cmd->buffer, buffer);
 	dup(offset);
 	dup(cb);
 	dup(ptr);
@@ -198,12 +238,17 @@ command_read_buffer command_read_buffer_create(cl_mem buffer, size_t offset, siz
 	return cmd;
 }
 
+void command_write_buffer_release(void *arg) {
+	command_write_buffer cmd = (command_write_buffer)arg;
+	gc_entity_unstore(&cmd->buffer);
+}
+
 command_write_buffer command_write_buffer_create(cl_mem buffer, size_t offset, size_t cb, const void * ptr) {
 
 	command_write_buffer cmd = malloc(sizeof(struct command_write_buffer_t));
-	command_init(cmd, CL_COMMAND_WRITE_BUFFER);
+	command_init(cmd, CL_COMMAND_WRITE_BUFFER, command_write_buffer_release);
 
-	dupEntity(buffer);
+	gc_entity_store(&cmd->buffer, buffer);
 	dup(offset);
 	dup(cb);
 	dup(ptr);
@@ -211,14 +256,20 @@ command_write_buffer command_write_buffer_create(cl_mem buffer, size_t offset, s
 	return cmd;
 }
 
+void command_copy_buffer_release(void *arg) {
+	command_copy_buffer cmd = (command_copy_buffer)arg;
+	gc_entity_unstore(&cmd->src_buffer);
+	gc_entity_unstore(&cmd->dst_buffer);
+}
+
 command_copy_buffer command_copy_buffer_create( cl_mem src_buffer, cl_mem dst_buffer,
 		size_t src_offset, size_t dst_offset, size_t cb)
 {
 	command_copy_buffer cmd = malloc(sizeof(struct command_copy_buffer_t));
-	command_init(cmd, CL_COMMAND_COPY_BUFFER);
+	command_init(cmd, CL_COMMAND_COPY_BUFFER, command_copy_buffer_release);
 
-	dupEntity(src_buffer);
-	dupEntity(dst_buffer);
+	gc_entity_store(&cmd->src_buffer, src_buffer);
+	gc_entity_store(&cmd->dst_buffer, dst_buffer);
 	dup(src_offset);
 	dup(dst_offset);
 	dup(cb);
@@ -229,7 +280,6 @@ command_copy_buffer command_copy_buffer_create( cl_mem src_buffer, cl_mem dst_bu
 #undef nullOrDup
 #undef nodeNullOrDup
 #undef dup
-#undef dupEntity
 #undef nodeDup
 #undef memdup
 

socl/src/command.h

@@ -14,6 +14,8 @@
  * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  */
 
+#include "socl.h"
+
 #ifndef SOCL_COMMANDS_H
 #define SOCL_COMMANDS_H
 
@@ -25,9 +27,9 @@ typedef struct cl_command_t * cl_command;
  * Command constructors for each kind of command use this method
  * Implicit and explicit dependencies must be passed as parameters
  */
-void command_init_ex(cl_command cmd, cl_command_type typ);
-#define command_init(cmd,typ) \
-	command_init_ex((cl_command)cmd,typ)
+void command_init_ex(cl_command cmd, cl_command_type typ, void (*cb)(void*));
+#define command_init(cmd,typ,cb) \
+	command_init_ex((cl_command)cmd,typ,cb)
 
 /** Submit a command for execution */
 void command_submit_ex(cl_command cmd);
@@ -45,6 +47,7 @@ void command_graph_dump_ex(cl_command cmd);
  * OpenCL Commands
  **************************/
 struct cl_command_t {
+	CL_ENTITY;
 	cl_command_type	typ;	 	/* Command type */
 	cl_uint 	num_events;	/* Number of dependencies */
 	cl_event * 	events;		/* Dependencies */
@@ -135,6 +138,10 @@ typedef struct command_marker_t {
 	CL_COMMAND
 } * command_marker;
 
+typedef struct command_barrier_t {
+	CL_COMMAND
+} * command_barrier;
+
 /*************************
  * Constructor functions
  *************************/
@@ -148,7 +155,7 @@ command_ndrange_kernel command_ndrange_kernel_create (
 
 command_ndrange_kernel command_task_create (cl_kernel kernel);
 
-command_marker command_barrier_create ();
+command_barrier command_barrier_create ();
 
 command_marker command_marker_create ();
 
@@ -192,6 +199,7 @@ cl_int command_copy_buffer_submit(command_copy_buffer cmd);
 cl_int command_map_buffer_submit(command_map_buffer cmd);
 cl_int command_unmap_mem_object_submit(command_unmap_mem_object cmd);
 cl_int command_marker_submit(command_marker cmd);
+cl_int command_barrier_submit(command_barrier cmd);
 
 
 #endif /* SOCL_COMMANDS_H */

socl/src/socl.h

@@ -47,17 +47,6 @@ typedef struct starpu_task * starpu_task;
  */
 typedef struct entity * entity;
 
-#include "command.h"
-#include "command_list.h"
-#include "command_queue.h"
-#include "debug.h"
-#include "event.h"
-#include "gc.h"
-#include "mem_objects.h"
-#include "task.h"
-#include "util.h"
-
-
 struct entity {
   struct _cl_icd_dispatch * dispatch;
   /* Reference count */
@@ -76,6 +65,18 @@ struct entity {
 #define CL_ENTITY struct entity _entity;
 
 
+#include "command.h"
+#include "command_list.h"
+#include "command_queue.h"
+#include "debug.h"
+#include "event.h"
+#include "gc.h"
+#include "mem_objects.h"
+#include "task.h"
+#include "util.h"
+
+
+
 struct _cl_platform_id {struct _cl_icd_dispatch *dispatch;};
 struct _cl_device_id {struct _cl_icd_dispatch *dispatch; int device_id; int worker_id;};