diff options
| -rw-r--r-- | src/runtime/CL/CLTuner.cpp | 8 | ||||
| -rw-r--r-- | tests/framework/instruments/OpenCLTimer.cpp | 12 |
2 files changed, 14 insertions, 6 deletions
diff --git a/src/runtime/CL/CLTuner.cpp b/src/runtime/CL/CLTuner.cpp index 56baad82c3..a262d6b95c 100644 --- a/src/runtime/CL/CLTuner.cpp +++ b/src/runtime/CL/CLTuner.cpp @@ -161,8 +161,6 @@ cl::NDRange CLTuner::find_optimal_lws(ICLKernel &kernel) auto interceptor = [this](cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *gwo, const size_t *gws, const size_t *lws, cl_uint num_events_in_wait_list, const cl_event * event_wait_list, cl_event * event) { - ARM_COMPUTE_ERROR_ON_MSG(event != nullptr, "Not supported"); - ARM_COMPUTE_UNUSED(event); if(this->kernel_event_is_set()) { // If the event is already set it means the kernel enqueue is sliced: given that we only time the first slice we can save time by skipping the other enqueues. @@ -174,6 +172,12 @@ cl::NDRange CLTuner::find_optimal_lws(ICLKernel &kernel) // Set OpenCL event this->set_cl_kernel_event(tmp); + if(event != nullptr) + { + //return cl_event from the intercepted call + clRetainEvent(tmp); + *event = tmp; + } return retval; }; CLSymbols::get().clEnqueueNDRangeKernel_ptr = interceptor; diff --git a/tests/framework/instruments/OpenCLTimer.cpp b/tests/framework/instruments/OpenCLTimer.cpp index 2fd5714f65..ca859b6fd9 100644 --- a/tests/framework/instruments/OpenCLTimer.cpp +++ b/tests/framework/instruments/OpenCLTimer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -107,9 +107,6 @@ void OpenCLClock<output_timestamps>::test_start() { if(this->_timer_enabled) { - ARM_COMPUTE_ERROR_ON_MSG(event != nullptr, "Not supported"); - ARM_COMPUTE_UNUSED(event); - kernel_info info; cl::Kernel cpp_kernel(kernel, true); std::stringstream ss; @@ -127,6 +124,13 @@ void OpenCLClock<output_timestamps>::test_start() cl_int retval = this->_real_function(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp); info.event = tmp; this->_kernels.push_back(std::move(info)); + + if(event != nullptr) + { + //return cl_event from the intercepted call + clRetainEvent(tmp); + *event = tmp; + } return retval; } else |