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>

5 files changed, 144 insertions, 186 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) {
diff --git a/src/cl_api_mem.c b/src/cl_api_mem.c
index de186841..09f9a141 100644
--- a/src/cl_api_mem.c
+++ b/src/cl_api_mem.c
@@ -107,7 +107,7 @@ clGetMemObjectInfo(cl_mem memobj,
     } else if (memobj->type == CL_MEM_IMAGE_TYPE) {
       parent = memobj;
     } else if (memobj->type == CL_MEM_BUFFER1D_IMAGE_TYPE) {
-      struct _cl_mem_buffer1d_image* image_buffer = (struct _cl_mem_buffer1d_image*)memobj;
+      struct _cl_mem_buffer1d_image *image_buffer = (struct _cl_mem_buffer1d_image *)memobj;
       parent = image_buffer->descbuffer;
     } else
       parent = NULL;
@@ -309,31 +309,21 @@ clEnqueueMapBuffer(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      ptr = data->ptr;
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
-      err = cl_enqueue_handle(data, CL_SUBMITTED); // Submit to get the address.
+      err = cl_event_exec(e, CL_SUBMITTED, CL_TRUE); // Submit to get the address.
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
 
-      e->status = CL_SUBMITTED;
-      ptr = data->ptr;
-      assert(ptr);
-
       cl_command_queue_enqueue_event(command_queue, e);
     }
 
+    ptr = data->ptr;
+    assert(ptr);
     err = cl_mem_record_map_mem(buffer, ptr, &mem_ptr, offset, size, NULL, NULL);
     assert(err == CL_SUCCESS);
   } while (0);
@@ -403,16 +393,15 @@ clEnqueueUnmapMemObject(cl_command_queue command_queue,
     data->ptr = mapped_ptr;
 
     if (e_status == CL_COMPLETE) { // No need to wait
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      e->status = CL_COMPLETE;
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else { // May need to wait some event to complete.
+      err = cl_event_exec(e, CL_QUEUED, CL_FALSE);
+      if (err != CL_SUCCESS) {
+        break;
+      }
       cl_command_queue_enqueue_event(command_queue, e);
     }
   } while (0);
@@ -507,16 +496,15 @@ clEnqueueReadBuffer(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
+      err = cl_event_exec(e, CL_QUEUED, CL_FALSE);
+      if (err != CL_SUCCESS) {
+        break;
+      }
       cl_command_queue_enqueue_event(command_queue, e);
     }
   } while (0);
@@ -611,16 +599,15 @@ clEnqueueWriteBuffer(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
+      err = cl_event_exec(e, CL_QUEUED, CL_FALSE);
+      if (err != CL_SUCCESS) {
+        break;
+      }
       cl_command_queue_enqueue_event(command_queue, e);
     }
   } while (0);
@@ -761,16 +748,15 @@ clEnqueueReadBufferRect(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
+      err = cl_event_exec(e, CL_QUEUED, CL_FALSE);
+      if (err != CL_SUCCESS) {
+        break;
+      }
       cl_command_queue_enqueue_event(command_queue, e);
     }
   } while (0);
@@ -913,16 +899,15 @@ clEnqueueWriteBufferRect(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
+      err = cl_event_exec(e, CL_QUEUED, CL_FALSE);
+      if (err != CL_SUCCESS) {
+        break;
+      }
       cl_command_queue_enqueue_event(command_queue, e);
     }
   } while (0);
@@ -1029,13 +1014,11 @@ clEnqueueCopyBuffer(cl_command_queue command_queue,
     if (e_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (e_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, e_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED, CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -1224,12 +1207,10 @@ clEnqueueCopyBufferRect(cl_command_queue command_queue,
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
     } else if (e_status == CL_COMPLETE) {
-      err = cl_enqueue_handle(&e->exec_data, CL_SUBMITTED);
+      err = cl_event_exec(e, CL_SUBMITTED, CL_FALSE);
       if (err != CL_SUCCESS) {
         break;
       }
-
-      e->status = CL_SUBMITTED;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -1324,13 +1305,11 @@ clEnqueueFillBuffer(cl_command_queue command_queue,
     if (e_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (e_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, e_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED, CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -1413,13 +1392,11 @@ clEnqueueMigrateMemObjects(cl_command_queue command_queue,
     if (e_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (e_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, e_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED, CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -1598,31 +1575,22 @@ clEnqueueMapImage(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      ptr = data->ptr;
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
-      err = cl_enqueue_handle(data, CL_SUBMITTED); // Submit to get the address.
+      err = cl_event_exec(e, CL_SUBMITTED, CL_TRUE); // Submit to get the address.
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
 
-      e->status = CL_SUBMITTED;
-      ptr = data->ptr;
-      assert(ptr);
-
       cl_command_queue_enqueue_event(command_queue, e);
     }
 
+    ptr = data->ptr;
+    assert(ptr);
+
     /* Store and write back map info. */
     if (mem->flags & CL_MEM_USE_HOST_PTR) {
       if (image_slice_pitch)
@@ -1797,16 +1765,15 @@ clEnqueueReadImage(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
+      err = cl_event_exec(e, CL_QUEUED, CL_FALSE);
+      if (err != CL_SUCCESS) {
+        break;
+      }
       cl_command_queue_enqueue_event(command_queue, e);
     }
   } while (0);
@@ -1950,16 +1917,15 @@ clEnqueueWriteImage(cl_command_queue command_queue,
 
     if (e_status == CL_COMPLETE) {
       // Sync mode, no need to queue event.
-      err = cl_enqueue_handle(data, CL_COMPLETE);
+      err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
-        assert(err < 0);
-        e->status = err;
         break;
       }
-
-      e->status = CL_COMPLETE; // Just set the status, no notify. No one depend on us now.
-      cl_event_update_timestamp(e, CL_QUEUED, CL_COMPLETE);
     } else {
+      err = cl_event_exec(e, CL_QUEUED, CL_FALSE);
+      if (err != CL_SUCCESS) {
+        break;
+      }
       cl_command_queue_enqueue_event(command_queue, e);
     }
   } while (0);
@@ -2093,13 +2059,11 @@ clEnqueueCopyImage(cl_command_queue command_queue,
     if (e_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (e_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, e_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED, CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -2206,13 +2170,11 @@ clEnqueueCopyImageToBuffer(cl_command_queue command_queue,
     if (e_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (e_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, e_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED, CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -2320,13 +2282,11 @@ clEnqueueCopyBufferToImage(cl_command_queue command_queue,
     if (e_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (e_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, e_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED, CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
@@ -2432,13 +2392,11 @@ clEnqueueFillImage(cl_command_queue command_queue,
     if (e_status < CL_COMPLETE) { // Error happend, cancel.
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
-    } else if (e_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, e_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED, CL_FALSE);
+    if (err != CL_SUCCESS) {
+      break;
     }
 
     cl_command_queue_enqueue_event(command_queue, e);
diff --git a/src/cl_command_queue_enqueue.c b/src/cl_command_queue_enqueue.c
index 7070effe..91fabf97 100644
--- a/src/cl_command_queue_enqueue.c
+++ b/src/cl_command_queue_enqueue.c
@@ -84,7 +84,7 @@ worker_thread_function(void *Arg)
     list_for_each_safe(pos, n, &ready_list)
     {
       e = list_entry(pos, _cl_event, enqueue_node);
-      cl_event_exec(e, exec_status);
+      cl_event_exec(e, exec_status, CL_FALSE);
     }
 
     /* Notify all waiting for flush. */
@@ -93,12 +93,10 @@ worker_thread_function(void *Arg)
     CL_OBJECT_NOTIFY_COND(queue);
     CL_OBJECT_UNLOCK(queue);
 
-    for (exec_status = CL_RUNNING; exec_status >= CL_COMPLETE; exec_status--) {
-      list_for_each_safe(pos, n, &ready_list)
-      {
-        e = list_entry(pos, _cl_event, enqueue_node);
-        cl_event_exec(e, exec_status);
-      }
+    list_for_each_safe(pos, n, &ready_list)
+    {
+      e = list_entry(pos, _cl_event, enqueue_node);
+      cl_event_exec(e, CL_COMPLETE, CL_FALSE);
     }
 
     /* Clear and delete all the events. */
@@ -135,8 +133,6 @@ LOCAL void
 cl_command_queue_enqueue_event(cl_command_queue queue, cl_event event)
 {
   CL_OBJECT_INC_REF(event);
-  cl_event_update_timestamp(event, CL_QUEUED, event->status);
-
   assert(CL_OBJECT_IS_COMMAND_QUEUE(queue));
   CL_OBJECT_LOCK(queue);
   assert(queue->worker.quit == CL_FALSE);
diff --git a/src/cl_event.c b/src/cl_event.c
index f849fd7b..e5a122b1 100644
--- a/src/cl_event.c
+++ b/src/cl_event.c
@@ -58,6 +58,12 @@ cl_event_update_timestamp_gen(cl_event event, cl_int status)
       if (ts == CL_EVENT_INVALID_TIMESTAMP)
         ts++;
       event->timestamp[3] = ts;
+
+      /* Set the submit time the same as running time if it is later. */
+      if (event->timestamp[1] > event->timestamp[2] ||
+          event->timestamp[2] - event->timestamp[1] > 0x0FFFFFFFFFF /*Overflowed */)
+        event->timestamp[1] = event->timestamp[2];
+
       return;
     }
   } else {
@@ -70,15 +76,13 @@ cl_event_update_timestamp_gen(cl_event event, cl_int status)
 }
 
 LOCAL void
-cl_event_update_timestamp(cl_event event, cl_int from, cl_int to)
+cl_event_update_timestamp(cl_event event, cl_int state)
 {
   int i;
   cl_bool re_cal = CL_FALSE;
   cl_ulong ts[4];
 
-  assert(from >= to);
-  assert(from >= CL_COMPLETE || from <= CL_QUEUED);
-  assert(to >= CL_COMPLETE || to <= CL_QUEUED);
+  assert(state >= CL_COMPLETE || state <= CL_QUEUED);
 
   if (event->event_type == CL_COMMAND_USER)
     return;
@@ -87,16 +91,11 @@ cl_event_update_timestamp(cl_event event, cl_int from, cl_int to)
   if ((event->queue->props & CL_QUEUE_PROFILING_ENABLE) == 0)
     return;
 
-  i = CL_QUEUED - from;
-  if (event->timestamp[i] == CL_EVENT_INVALID_TIMESTAMP)
-    cl_event_update_timestamp_gen(event, from);
-  i++;
+  /* Should not record the timestamp twice. */
+  assert(event->timestamp[CL_QUEUED - state] == CL_EVENT_INVALID_TIMESTAMP);
+  cl_event_update_timestamp_gen(event, state);
 
-  for (; i <= CL_QUEUED - to; i++) {
-    cl_event_update_timestamp_gen(event, CL_QUEUED - i);
-  }
-
-  if (to == CL_COMPLETE) {
+  if (state == CL_COMPLETE) {
     // TODO: Need to set the CL_PROFILING_COMMAND_COMPLETE when enable child enqueue.
     // Just a duplicate of event complete time now.
     event->timestamp[4] = event->timestamp[3];
@@ -168,7 +167,7 @@ cl_event_new(cl_context ctx, cl_command_queue queue, cl_command_type type,
   if (type == CL_COMMAND_USER) {
     e->status = CL_SUBMITTED;
   } else {
-    e->status = CL_QUEUED;
+    e->status = CL_EVENT_STATE_UNKNOWN;
   }
 
   if (type == CL_COMMAND_USER) {
@@ -383,16 +382,6 @@ cl_event_set_status(cl_event event, cl_int status)
     return CL_INVALID_OPERATION;
   }
 
-  if (status >= CL_COMPLETE && !CL_EVENT_IS_USER(event) &&
-      (event->queue->props & CL_QUEUE_PROFILING_ENABLE) != 0) {
-    // Call update_timestamp without event lock.
-    CL_OBJECT_TAKE_OWNERSHIP_WITHLOCK(event, 1);
-    CL_OBJECT_UNLOCK(event);
-    cl_event_update_timestamp(event, event->status, status);
-    CL_OBJECT_LOCK(event);
-    CL_OBJECT_RELEASE_OWNERSHIP_WITHLOCK(event);
-  }
-
   event->status = status;
 
   /* Call all the callbacks. */
@@ -572,39 +561,60 @@ cl_event_check_waitlist(cl_uint num_events_in_wait_list, const cl_event *event_w
   return err;
 }
 
-LOCAL void
-cl_event_exec(cl_event event, cl_int exec_status)
+/* When we call this function, all the events it depends
+   on should already be ready, unless ignore_depends is set. */
+LOCAL cl_uint
+cl_event_exec(cl_event event, cl_int exec_to_status, cl_bool ignore_depends)
 {
   /* We are MT safe here, no one should call this
      at the same time. No need to lock */
   cl_int ret = CL_SUCCESS;
-  cl_int status = cl_event_get_status(event);
+  cl_int cur_status = cl_event_get_status(event);
   cl_int depend_status;
+  cl_int s;
 
-  if (status < CL_COMPLETE || status <= exec_status) {
-    return;
+  assert(exec_to_status >= CL_COMPLETE);
+  assert(exec_to_status <= CL_QUEUED);
+  if (cur_status < CL_COMPLETE) {
+    return cur_status;
   }
 
   depend_status = cl_event_is_ready(event);
-  assert(depend_status <= CL_COMPLETE);
+  assert(depend_status <= CL_COMPLETE || ignore_depends || exec_to_status == CL_QUEUED);
   if (depend_status < CL_COMPLETE) { // Error happend, cancel exec.
     ret = cl_event_set_status(event, depend_status);
-    return;
+    return depend_status;
   }
 
-  /* Do the according thing based on event type. */
-  ret = cl_enqueue_handle(&event->exec_data, exec_status);
+  if (cur_status <= exec_to_status) {
+    return ret;
+  }
 
-  if (ret != CL_SUCCESS) {
-    assert(ret < 0);
-    DEBUGP(DL_WARNING, "Exec event %p error, type is %d, error staus is %d",
-           event, event->event_type, ret);
-    ret = cl_event_set_status(event, ret);
-    assert(ret == CL_SUCCESS);
-  } else {
-    ret = cl_event_set_status(event, exec_status);
-    assert(ret == CL_SUCCESS);
+  /* Exec to the target status. */
+  for (s = cur_status - 1; s >= exec_to_status; s--) {
+    assert(s >= CL_COMPLETE);
+    ret = cl_enqueue_handle(&event->exec_data, s);
+
+    if (ret != CL_SUCCESS) {
+      assert(ret < 0);
+      DEBUGP(DL_WARNING, "Exec event %p error, type is %d, error staus is %d",
+             event, event->event_type, ret);
+      ret = cl_event_set_status(event, ret);
+      assert(ret == CL_SUCCESS);
+      return ret; // Failed and we never do further.
+    } else {
+      assert(!CL_EVENT_IS_USER(event));
+      if ((event->queue->props & CL_QUEUE_PROFILING_ENABLE) != 0) {
+        /* record the timestamp before actually doing something. */
+        cl_event_update_timestamp(event, s);
+      }
+
+      ret = cl_event_set_status(event, s);
+      assert(ret == CL_SUCCESS);
+    }
   }
+
+  return ret;
 }
 
 /* 0 means ready, >0 means not ready, <0 means error. */
diff --git a/src/cl_event.h b/src/cl_event.h
index cb43244f..f28844af 100644
--- a/src/cl_event.h
+++ b/src/cl_event.h
@@ -57,6 +57,8 @@ typedef struct _cl_event {
          ((cl_base_object)obj)->magic == CL_OBJECT_EVENT_MAGIC &&  \
          CL_OBJECT_GET_REF(obj) >= 1))
 
+#define CL_EVENT_STATE_UNKNOWN 0x4
+
 #define CL_EVENT_IS_MARKER(E) (E->event_type == CL_COMMAND_MARKER)
 #define CL_EVENT_IS_BARRIER(E) (E->event_type == CL_COMMAND_BARRIER)
 #define CL_EVENT_IS_USER(E) (E->event_type == CL_COMMAND_USER)
@@ -68,7 +70,7 @@ extern cl_event cl_event_create(cl_context ctx, cl_command_queue queue, cl_uint
                                 const cl_event *event_list, cl_command_type type, cl_int *errcode_ret);
 extern cl_int cl_event_check_waitlist(cl_uint num_events_in_wait_list, const cl_event *event_wait_list,
                                       cl_event* event, cl_context ctx);
-extern void cl_event_exec(cl_event event, cl_int exec_status);
+extern cl_uint cl_event_exec(cl_event event, cl_int exec_to_status, cl_bool ignore_depends);
 /* 0 means ready, >0 means not ready, <0 means error. */
 extern cl_int cl_event_is_ready(cl_event event);
 extern cl_int cl_event_get_status(cl_event event);
@@ -82,5 +84,5 @@ extern cl_int cl_event_wait_for_event_ready(cl_event event);
 extern cl_event cl_event_create_marker_or_barrier(cl_command_queue queue, cl_uint num_events_in_wait_list,
                                                   const cl_event *event_wait_list, cl_bool is_barrier,
                                                   cl_int* error);
-extern void cl_event_update_timestamp(cl_event event, cl_int from_status, cl_int to_status);
+extern void cl_event_update_timestamp(cl_event event, cl_int status);
 #endif /* __CL_EVENT_H__ */