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.
Signed-off-by: Junyan He <[email protected]> --- src/cl_api_kernel.c | 10 +--- src/cl_api_mem.c | 120 +++++++++-------------------------------- src/cl_command_queue_enqueue.c | 14 ++--- src/cl_event.c | 50 ++++++++++------- src/cl_event.h | 2 +- 5 files changed, 65 insertions(+), 131 deletions(-) diff --git a/src/cl_api_kernel.c b/src/cl_api_kernel.c index 723152f..ca7a5f8 100644 --- a/src/cl_api_kernel.c +++ b/src/cl_api_kernel.c @@ -227,12 +227,10 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, 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); + 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); @@ -351,14 +349,10 @@ clEnqueueNativeKernel(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. } else { cl_command_queue_enqueue_event(command_queue, e); } diff --git a/src/cl_api_mem.c b/src/cl_api_mem.c index de18684..b714926 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,15 +393,10 @@ 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. cl_command_queue_enqueue_event(command_queue, e); } @@ -507,15 +492,10 @@ 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 { cl_command_queue_enqueue_event(command_queue, e); } @@ -611,15 +591,10 @@ 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 { cl_command_queue_enqueue_event(command_queue, e); } @@ -761,15 +736,10 @@ 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 { cl_command_queue_enqueue_event(command_queue, e); } @@ -913,15 +883,10 @@ 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 { cl_command_queue_enqueue_event(command_queue, e); } @@ -1030,12 +995,10 @@ clEnqueueCopyBuffer(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); @@ -1224,12 +1187,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); @@ -1325,12 +1286,10 @@ clEnqueueFillBuffer(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); @@ -1414,12 +1373,10 @@ clEnqueueMigrateMemObjects(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); @@ -1598,31 +1555,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,15 +1745,10 @@ 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 { cl_command_queue_enqueue_event(command_queue, e); } @@ -1950,15 +1893,10 @@ 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 { cl_command_queue_enqueue_event(command_queue, e); } @@ -2094,12 +2032,10 @@ clEnqueueCopyImage(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); @@ -2207,12 +2143,10 @@ clEnqueueCopyImageToBuffer(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); @@ -2321,12 +2255,10 @@ clEnqueueCopyBufferToImage(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); @@ -2433,12 +2365,10 @@ clEnqueueFillImage(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); 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..8173578 100644 --- a/src/cl_event.c +++ b/src/cl_event.c @@ -573,39 +573,53 @@ 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); 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--) { + 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 { + 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..ae38800 100644 --- a/src/cl_event.h +++ b/src/cl_event.h @@ -68,7 +68,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); -- 2.7.4 _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
