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>

1 files changed, 6 insertions, 2 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;