diff options
| author | Vidhya Sudhan Loganathan <vidhyasudhan.loganathan@arm.com> | 2019-02-07 11:14:42 +0000 |
|---|---|---|
| committer | Isabella Gottardi <isabella.gottardi@arm.com> | 2019-02-12 17:06:05 +0000 |
| commit | 75d4122405e29505c1d97f55f4ae6da775140ed9 (patch) | |
| tree | b2ba18f599d48d668844e19b75e5677b72b67d55 | |
| parent | 6ea4d782c08ff76a8861dccb2ee8d378d2e2345b (diff) | |
| download | ComputeLibrary-75d4122405e29505c1d97f55f4ae6da775140ed9.tar.gz | |
COMPMID-1919: Opencl - Tuning with Instrumentation does not work
Change-Id: Iee6a07d5bf6a35af04071865682bcc4a615c14f9
Signed-off-by: Vidhya Sudhan Loganathan <vidhyasudhan.loganathan@arm.com>
Reviewed-on: https://review.mlplatform.org/638
Reviewed-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Pablo Marquez <pablo.tello@arm.com>
| -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 |