From 64b8b046f2f844d292a692e788201acecb586359 Mon Sep 17 00:00:00 2001 From: Junyan He Date: Tue, 27 Dec 2016 18:45:22 +0800 Subject: 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 Reviewed-by: Yang Rong --- src/cl_api_kernel.c | 26 +++++++++----------------- 1 file changed, 9 insertions(+), 17 deletions(-) (limited to 'src/cl_api_kernel.c') 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) { -- cgit v1.2.1