From 75d4122405e29505c1d97f55f4ae6da775140ed9 Mon Sep 17 00:00:00 2001 From: Vidhya Sudhan Loganathan Date: Thu, 7 Feb 2019 11:14:42 +0000 Subject: COMPMID-1919: Opencl - Tuning with Instrumentation does not work Change-Id: Iee6a07d5bf6a35af04071865682bcc4a615c14f9 Signed-off-by: Vidhya Sudhan Loganathan Reviewed-on: https://review.mlplatform.org/638 Reviewed-by: Giuseppe Rossini Tested-by: Arm Jenkins Reviewed-by: Pablo Marquez --- src/runtime/CL/CLTuner.cpp | 8 ++++++-- 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::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::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 -- cgit v1.2.1