Improve event execute function.

Modify the event exec function, make it as the uniformal entry
for all event command execution. This will help the timestamp
record and profiling feature a lot.

V2:
1. Set event init state to bigger than CL_QUEUED.
Event state should be set to CL_QUEUED exactly when it is to be queued.
Profiling feature make this requirement clearer. We need to record the
timestamp exactly when it it to be queued. So we need to add a additional
state beyond CL_QUEUED.

2. Fix cl_event_update_timestamp_gen bugi, the CL_SUMITTED time may be less.
GPU may record the timestamp of CL_RUNNING before CPU record timestamp of
CL_SUMITTED. It is a async process and it is hard for us to control.
According to SPEC, we need to record timestamp after some state is done.
We can just now set CL_SUMITTED to CL_RUNNING timestamp if the CL_SUBMITTED
timestamp is the bigger one.

Signed-off-by: Junyan He <junyan.he@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

1 files changed, 9 insertions, 17 deletions

diff --git a/src/cl_api_kernel.c b/src/cl_api_kernel.c
index 723152f7..c7d73313 100644
--- a/src/cl_api_kernel.c
+++ b/src/cl_api_kernel.c
@@ -226,13 +226,11 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
     if (event_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (event_status == CL_COMPLETE) {
-      err = cl_enqueue_handle(&e->exec_data, CL_SUBMITTED);
-      if (err != CL_SUCCESS) {
-        break;
-      }
+    }
 
-      e->status = CL_SUBMITTED;
+    err = cl_event_exec(e, (event_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED), CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -349,19 +347,13 @@ clEnqueueNativeKernel(cl_command_queue command_queue,
     new_mem_list = NULL;
     new_args_mem_loc = NULL; // Event delete will free them.
 
-    if (e_status == CL_COMPLETE) {
-      // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
-      if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
-        break;
-      }
+    err = cl_event_exec(e, (e_status == CL_COMPLETE ? CL_COMPLETE : CL_QUEUED), CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
+    }
 
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-    } else {
+    if (e_status != CL_COMPLETE)
       cl_command_queue_enqueue_event(command_queue, e);
-    }
   } while (0);
 
   if (err != CL_SUCCESS) {