SOCL: dot not use StarPU's profiling info

socl/src/cl_enqueuecopybuffer.c

@@ -22,6 +22,10 @@ static void soclEnqueueCopyBuffer_opencl_task(void *descr[], void *args) {
    cl_event ev;
    command_copy_buffer cmd = (command_copy_buffer)args;;
 
+  cl_event event = command_event_get(cmd);
+  event->prof_start = _socl_nanotime();
+  gc_entity_release(event);
+
    wid = starpu_worker_get_id();
    starpu_opencl_get_queue(wid, &cq);
 
@@ -38,6 +42,10 @@ static void soclEnqueueCopyBuffer_opencl_task(void *descr[], void *args) {
 static void soclEnqueueCopyBuffer_cpu_task(void *descr[], void *args) {
    command_copy_buffer cmd = (command_copy_buffer)args;;
 
+  cl_event ev = command_event_get(cmd);
+  ev->prof_start = _socl_nanotime();
+  gc_entity_release(ev);
+
    void * src = (void*)STARPU_VARIABLE_GET_PTR(descr[0]);
    void * dst = (void*)STARPU_VARIABLE_GET_PTR(descr[1]);
 

socl/src/cl_enqueuemapbuffer.c

@@ -19,6 +19,10 @@
 static void mapbuffer_task(void *args) {
 	command_map_buffer cmd = (command_map_buffer)args;
 
+  cl_event ev = command_event_get(cmd);
+  ev->prof_start = _socl_nanotime();
+  gc_entity_release(ev);
+
 	enum starpu_access_mode mode = (cmd->map_flags == CL_MAP_READ ? STARPU_R : STARPU_RW);
 
 	starpu_data_acquire_cb(cmd->buffer->handle, mode, command_completed_task_callback, cmd);

socl/src/cl_enqueuendrangekernel.c

@@ -24,6 +24,10 @@ void soclEnqueueNDRangeKernel_task(void *descr[], void *args) {
    int wid;
    cl_int err;
 
+  cl_event ev = command_event_get(cmd);
+  ev->prof_start = _socl_nanotime();
+  gc_entity_release(ev);
+
    wid = starpu_worker_get_id();
    starpu_opencl_get_queue(wid, &cq);
 

socl/src/cl_enqueuereadbuffer.c

@@ -19,6 +19,10 @@
 static void soclEnqueueReadBuffer_cpu_task(void *descr[], void *args) {
    command_read_buffer cmd = (command_read_buffer)args;
 
+  cl_event ev = command_event_get(cmd);
+  ev->prof_start = _socl_nanotime();
+  gc_entity_release(ev);
+
    void * ptr = (void*)STARPU_VARIABLE_GET_PTR(descr[0]);
    DEBUG_MSG("[Buffer %d] Reading %ld bytes from %p to %p\n", cmd->buffer->id, cmd->cb, ptr+cmd->offset, cmd->ptr);
 
@@ -33,6 +37,10 @@ static void soclEnqueueReadBuffer_cpu_task(void *descr[], void *args) {
 static void soclEnqueueReadBuffer_opencl_task(void *descr[], void *args) {
    command_read_buffer cmd = (command_read_buffer)args;
 
+  cl_event event = command_event_get(cmd);
+  event->prof_start = _socl_nanotime();
+  gc_entity_release(event);
+
    cl_mem mem = (cl_mem)STARPU_VARIABLE_GET_PTR(descr[0]);
 
    DEBUG_MSG("[Buffer %d] Reading %ld bytes from offset %ld into %p\n", cmd->buffer->id, cmd->cb, cmd->offset, cmd->ptr);

socl/src/cl_enqueuewritebuffer.c

@@ -20,6 +20,10 @@
 static void soclEnqueueWriteBuffer_cpu_task(void *descr[], void *args) {
    command_write_buffer cmd = (command_write_buffer)args;
 
+  cl_event ev = command_event_get(cmd);
+  ev->prof_start = _socl_nanotime();
+  gc_entity_release(ev);
+
    void * ptr = (void*)STARPU_VARIABLE_GET_PTR(descr[0]);
    DEBUG_MSG("[Buffer %d] Writing %ld bytes from %p to %p\n", cmd->buffer->id, cmd->cb, cmd->ptr, ptr+cmd->offset);
 
@@ -36,6 +40,10 @@ static void soclEnqueueWriteBuffer_cpu_task(void *descr[], void *args) {
 static void soclEnqueueWriteBuffer_opencl_task(void *descr[], void *args) {
    command_write_buffer cmd = (command_write_buffer)args;
 
+  cl_event event = command_event_get(cmd);
+  event->prof_start = _socl_nanotime();
+  gc_entity_release(event);
+
    cl_mem mem = (cl_mem)STARPU_VARIABLE_GET_PTR(descr[0]);
 
    DEBUG_MSG("[Buffer %d] Writing %ld bytes to offset %ld from %p\n", cmd->buffer->id, cmd->cb, cmd->offset, cmd->ptr);

socl/src/cl_geteventprofilinginfo.c

@@ -24,23 +24,14 @@ soclGetEventProfilingInfo(cl_event          event,
                         void *              param_value,
                         size_t *            param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
 {
-   struct starpu_task_profiling_info * prof = event->profiling_info;
-
-   if (prof == NULL)
-      return CL_PROFILING_INFO_NOT_AVAILABLE;
-
-   #define TONANO(t) ((cl_ulong)t.tv_nsec + (cl_ulong)(t.tv_sec)*1e9)
-
    switch (param_name) {
-      case CL_PROFILING_COMMAND_QUEUED:
-      INFO_CASE_VALUE(CL_PROFILING_COMMAND_SUBMIT, cl_ulong, TONANO(prof->submit_time));
-      INFO_CASE_VALUE(CL_PROFILING_COMMAND_START, cl_ulong, TONANO(prof->start_time));
-      INFO_CASE_VALUE(CL_PROFILING_COMMAND_END, cl_ulong, TONANO(prof->end_time));
+      INFO_CASE_VALUE(CL_PROFILING_COMMAND_QUEUED, cl_ulong, event->prof_queued);
+      INFO_CASE_VALUE(CL_PROFILING_COMMAND_SUBMIT, cl_ulong, event->prof_submit);
+      INFO_CASE_VALUE(CL_PROFILING_COMMAND_START, cl_ulong, event->prof_start);
+      INFO_CASE_VALUE(CL_PROFILING_COMMAND_END, cl_ulong, event->prof_end);
       default:
          return CL_INVALID_VALUE;
    }
 
-   #undef TONANO
-
    return CL_SUCCESS;
 }

socl/src/command_queue.c

@@ -29,6 +29,10 @@
 
 void command_queue_enqueue_ex(cl_command_queue cq, cl_command cmd, cl_uint num_events, const cl_event * events) {
 
+  cl_event ev = command_event_get_ex(cmd);
+  ev->prof_queued = _socl_nanotime();
+  gc_entity_release(ev);
+
 	/* Check if the command is a barrier */
 	int is_barrier = (cmd->typ == CL_COMMAND_BARRIER || !(cq->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE));
 

socl/src/event.c

@@ -38,7 +38,10 @@ cl_event event_create(void) {
    ev->id = event_unique_id();
    ev->status = CL_SUBMITTED;
    ev->command = NULL;
-   ev->profiling_info = NULL;
+   ev->prof_queued = 0L;
+   ev->prof_submit = 0L;
+   ev->prof_start = 0L;
+   ev->prof_end = 0L;
    ev->cq = NULL;
 
    return ev;
@@ -52,9 +55,5 @@ static void release_callback_event(void * e) {
   /* Destruct object */
   //FIXME
   //starpu_tag_remove(event->id);
-  if (event->profiling_info != NULL) {
-    free(event->profiling_info);
-    event->profiling_info = NULL;
-  }
 }
 

socl/src/socl.h

@@ -178,8 +178,8 @@ struct _cl_event {
    */
   int id;
 
-  /* Profiling info are copied here */
-  struct starpu_task_profiling_info *profiling_info;
+  /* Profiling info */
+  cl_ulong prof_queued, prof_submit, prof_start, prof_end;
 };
 
 struct _cl_mem {

socl/src/task.c

@@ -23,16 +23,17 @@ void command_completed(cl_command cmd) {
   
   cl_event ev = command_event_get_ex(cmd);
   ev->status = CL_COMPLETE;
+  
+  ev->prof_end = _socl_nanotime();
+
+  /* Commands without codelets (marker, barrier, unmap...) take no time */
+  if (task->cl == NULL) 
+    ev->prof_start = ev->prof_end;
 
   /* Trigger the tag associated to the command event */
   DEBUG_MSG("Trigger event %d\n", ev->id);
   starpu_tag_notify_from_apps(ev->id);
 
-  if (task->profiling_info != NULL && (intptr_t)task->profiling_info != -ENOSYS) {
-    ev->profiling_info = malloc(sizeof(*task->profiling_info));
-    memcpy(ev->profiling_info, task->profiling_info, sizeof(*task->profiling_info));
-  }
-
   gc_entity_release(ev);
 }
 
@@ -99,6 +100,10 @@ cl_int task_submit_ex(starpu_task task, cl_command cmd) {
   task->callback_func = command_completed_task_callback;
   gc_entity_store(&task->callback_arg, cmd);
 
+  cl_event ev = command_event_get_ex(cmd);
+  ev->prof_submit = _socl_nanotime();
+  gc_entity_release(ev);
+
   /* Submit task */
   int ret = (task->cl != NULL && task->cl->where == STARPU_OPENCL ?
         starpu_task_submit_to_ctx(task, cmd->event->cq->context->sched_ctx) :

socl/src/util.c

@@ -52,3 +52,10 @@ void ** memdup_deep_varsize_safea(const void **p, unsigned n, size_t * size) {
 	}
 	return s;
 }
+
+cl_ulong _socl_nanotime() {
+   struct timespec ts;
+   clock_gettime(CLOCK_MONOTONIC, &ts);
+
+   return (ts.tv_sec * 1e9 + ts.tv_nsec);
+}

socl/src/util.h

@@ -44,4 +44,6 @@ void ** memdup_deep_varsize_safea(const void **p, unsigned n, size_t * size);
 
 #define memdup_deep_varsize_safe(p,n,size) ((typeof(p))memdup_deep_varsize_safea((const void **)p,n,size))
 
+cl_ulong _socl_nanotime();
+
 #endif /* SOCL_UTIL_H */