[Beignet] [PATCH] Improve event execute function.
Yang, Rong R
rong.r.yang at intel.com
Wed Dec 28 08:08:53 UTC 2016
LGTM, pushed, thanks.
> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> junyan.he at inbox.com
> Sent: Tuesday, December 27, 2016 18:45
> To: beignet at lists.freedesktop.org
> Subject: [Beignet] [PATCH] Improve event execute function.
>
> 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
>
>
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list