[Beignet] [PATCH] Improve event execute function.
junyan.he at inbox.com
junyan.he at inbox.com
Tue Dec 27 10:45:22 UTC 2016
From: Junyan He <junyan.he at intel.com>
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 at intel.com>
---
src/cl_api_kernel.c | 26 ++----
src/cl_api_mem.c | 190 ++++++++++++++++-------------------------
src/cl_command_queue_enqueue.c | 14 ++-
src/cl_event.c | 94 +++++++++++---------
src/cl_event.h | 6 +-
5 files changed, 144 insertions(+), 186 deletions(-)
diff --git a/src/cl_api_kernel.c b/src/cl_api_kernel.c
index 723152f..c7d7331 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 de18684..09f9a14 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 cf9ee3f..9de15aa 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 212f184..519e6c6 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. */
@@ -573,39 +562,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 9df5ab6..6853ce9 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__ */
--
2.7.4
More information about the Beignet
mailing list