LGTM, pushed, thanks.
> -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf Of > [email protected] > Sent: Tuesday, December 27, 2016 18:45 > To: [email protected] > Subject: [Beignet] [PATCH] Improve event execute function. > > From: Junyan He <[email protected]> > > 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 <[email protected]> > --- > 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 > [email protected] > https://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
