COMPMID-663 Adding instrument to return timestamps of clGetEventProfilingInfo

Change-Id: I1037054615593205f07e25fb9b16fecd13407c2c
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/95142
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>

6 files changed, 226 insertions, 1 deletions

diff --git a/tests/framework/Framework.cpp b/tests/framework/Framework.cpp
index 82ea0fa93d..a5c665c458 100644
--- a/tests/framework/Framework.cpp
+++ b/tests/framework/Framework.cpp
@@ -50,6 +50,9 @@ Framework::Framework()
 #ifdef MALI_ENABLED
     _available_instruments.emplace(InstrumentType::MALI, Instrument::make_instrument<MaliCounter>);
 #endif /* MALI_ENABLED */
+#ifdef OPENCL_TIMER_ENABLED
+    _available_instruments.emplace(InstrumentType::OPENCL_TIMER, Instrument::make_instrument<OpenCLTimer>);
+#endif /* OPENCL_TIMER_ENABLED */
 }
 
 std::set<InstrumentType> Framework::available_instruments() const
diff --git a/tests/framework/SConscript b/tests/framework/SConscript
index 92cbd552bb..52b8bed10f 100644
--- a/tests/framework/SConscript
+++ b/tests/framework/SConscript
@@ -28,7 +28,8 @@ Import('vars')
 # vars is imported from arm_compute:
 variables = [
     BoolVariable("pmu", "Enable PMU counters", False),
-    BoolVariable("mali", "Enable Mali hardware counters", False)
+    BoolVariable("mali", "Enable Mali hardware counters", False),
+    BoolVariable("opencl_timer", "Enable OpenCL timers", False)
 ]
 
 # We need a separate set of Variables for the Help message (Otherwise the global variables will get displayed twice)
@@ -62,6 +63,15 @@ if not framework_env['pmu']:
 else:
     framework_env.Append(CPPDEFINES = ['PMU_ENABLED'])
 
+if not framework_env['opencl_timer']:
+    # Remove OpenCLTimer files
+    files = [f for f in files if "OpenCLTimer" not in os.path.basename(str(f))]
+else:
+    if not framework_env["opencl"]:
+        print("ERROR: You need opencl=1 to be able to use opencl_timer=1")
+        Exit(1)
+    framework_env.Append(CPPDEFINES = ['OPENCL_TIMER_ENABLED'])
+
 if not framework_env['mali']:
     # Remove MALI files
     files = [f for f in files if "MaliCounter" not in os.path.basename(str(f))]
diff --git a/tests/framework/instruments/Instruments.cpp b/tests/framework/instruments/Instruments.cpp
index 1c21f0773d..eca6d6f437 100644
--- a/tests/framework/instruments/Instruments.cpp
+++ b/tests/framework/instruments/Instruments.cpp
@@ -45,6 +45,7 @@ InstrumentType instrument_type_from_name(const std::string &name)
         { "pmu_cycles", InstrumentType::PMU_CYCLE_COUNTER },
         { "pmu_instructions", InstrumentType::PMU_INSTRUCTION_COUNTER },
         { "mali", InstrumentType::MALI },
+        { "opencl_timer", InstrumentType::OPENCL_TIMER },
     };
 
     try
diff --git a/tests/framework/instruments/Instruments.h b/tests/framework/instruments/Instruments.h
index df6aa62ec7..ffe7cb681d 100644
--- a/tests/framework/instruments/Instruments.h
+++ b/tests/framework/instruments/Instruments.h
@@ -25,6 +25,7 @@
 #define ARM_COMPUTE_TEST_INSTRUMENTS
 
 #include "MaliCounter.h"
+#include "OpenCLTimer.h"
 #include "PMUCounter.h"
 #include "WallClockTimer.h"
 
@@ -46,6 +47,7 @@ enum class InstrumentType : unsigned int
     PMU_CYCLE_COUNTER       = 0x0201,
     PMU_INSTRUCTION_COUNTER = 0x0202,
     MALI                    = 0x0300,
+    OPENCL_TIMER            = 0x0400,
 };
 
 InstrumentType instrument_type_from_name(const std::string &name);
@@ -77,6 +79,9 @@ inline ::std::stringstream &operator<<(::std::stringstream &stream, InstrumentTy
         case InstrumentType::MALI:
             stream << "MALI";
             break;
+        case InstrumentType::OPENCL_TIMER:
+            stream << "OPENCL_TIMER";
+            break;
         case InstrumentType::ALL:
             stream << "ALL";
             break;
diff --git a/tests/framework/instruments/OpenCLTimer.cpp b/tests/framework/instruments/OpenCLTimer.cpp
new file mode 100644
index 0000000000..819ecdb743
--- /dev/null
+++ b/tests/framework/instruments/OpenCLTimer.cpp
@@ -0,0 +1,143 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "OpenCLTimer.h"
+
+#include "../Framework.h"
+#include "../Utils.h"
+
+#include "arm_compute/runtime/CL/CLScheduler.h"
+
+#ifndef ARM_COMPUTE_CL
+#error "You can't use OpenCLTimer without OpenCL"
+#endif /* ARM_COMPUTE_CL */
+
+namespace arm_compute
+{
+namespace test
+{
+namespace framework
+{
+std::string OpenCLTimer::id() const
+{
+    return "OpenCLTimer";
+}
+
+/* Function to be used to intercept kernel enqueues and store their OpenCL Event */
+class Interceptor
+{
+public:
+    explicit Interceptor(OpenCLTimer &timer)
+        : _timer(timer)
+    {
+    }
+
+    cl_int operator()(
+        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);
+
+        OpenCLTimer::kernel_info info;
+        cl::Kernel               cpp_kernel(kernel, true);
+        std::stringstream        ss;
+        ss << cpp_kernel.getInfo<CL_KERNEL_FUNCTION_NAME>();
+        if(gws != nullptr)
+        {
+            ss << " GWS[" << gws[0] << "," << gws[1] << "," << gws[2] << "]";
+        }
+        if(lws != nullptr)
+        {
+            ss << " LWS[" << lws[0] << "," << lws[1] << "," << lws[2] << "]";
+        }
+        info.name = ss.str();
+        cl_event tmp;
+        cl_int   retval = _timer.real_function(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp);
+        info.event      = tmp;
+        _timer.kernels.push_back(std::move(info));
+        return retval;
+    }
+
+private:
+    OpenCLTimer &_timer;
+};
+
+OpenCLTimer::OpenCLTimer()
+    : real_function(CLSymbols::get().clEnqueueNDRangeKernel_ptr)
+{
+    auto                        q     = CLScheduler::get().queue();
+    cl_command_queue_properties props = q.getInfo<CL_QUEUE_PROPERTIES>();
+    if((props & CL_QUEUE_PROFILING_ENABLE) == 0)
+    {
+        CLScheduler::get().set_queue(cl::CommandQueue(CLScheduler::get().context(), props | CL_QUEUE_PROFILING_ENABLE));
+    }
+}
+
+void OpenCLTimer::start()
+{
+    kernels.clear();
+    // Start intercepting enqueues:
+    CLSymbols::get().clEnqueueNDRangeKernel_ptr = Interceptor(*this);
+}
+
+void OpenCLTimer::stop()
+{
+    // Restore real function
+    CLSymbols::get().clEnqueueNDRangeKernel_ptr = real_function;
+}
+
+Instrument::MeasurementsMap OpenCLTimer::measurements() const
+{
+    MeasurementsMap measurements;
+    unsigned int    kernel_number = 0;
+    for(auto kernel : kernels)
+    {
+        //cl_int status = kernel.event.getInfo(<CL_EVENT_COMMAND_EXECUTION_STATUS>();
+        cl_ulong queued = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_QUEUED>();
+        cl_ulong submit = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_SUBMIT>();
+        cl_ulong start  = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
+        cl_ulong end    = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
+
+        std::list<std::string> raw_data =
+        {
+            "queued", support::cpp11::to_string(queued),
+            "submit", support::cpp11::to_string(submit),
+            "start", support::cpp11::to_string(start),
+            "end", support::cpp11::to_string(end),
+        };
+        measurements.emplace(kernel.name + " #" + support::cpp11::to_string(kernel_number++), Measurement(end - start, "ns", raw_data));
+    }
+
+    return measurements;
+}
+} // namespace framework
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/framework/instruments/OpenCLTimer.h b/tests/framework/instruments/OpenCLTimer.h
new file mode 100644
index 0000000000..d00a2e8bf2
--- /dev/null
+++ b/tests/framework/instruments/OpenCLTimer.h
@@ -0,0 +1,63 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_OPENCL_TIMER
+#define ARM_COMPUTE_TEST_OPENCL_TIMER
+
+#include "Instrument.h"
+
+#ifdef ARM_COMPUTE_CL
+#include "arm_compute/core/CL/OpenCL.h"
+#endif /* ARM_COMPUTE_CL */
+
+#include <list>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace framework
+{
+/** Instrument creating measurements based on the information returned by clGetEventProfilingInfo for each OpenCL kernel executed*/
+class OpenCLTimer : public Instrument
+{
+public:
+    OpenCLTimer();
+    std::string     id() const override;
+    void            start() override;
+    void            stop() override;
+    MeasurementsMap measurements() const override;
+#ifdef ARM_COMPUTE_CL
+    struct kernel_info
+    {
+        cl::Event   event{}; /**< OpenCL event associated to the kernel enqueue */
+        std::string name{};  /**< OpenCL Kernel name */
+    };
+    std::list<kernel_info>                          kernels{};
+    std::function<decltype(clEnqueueNDRangeKernel)> real_function;
+#endif /* ARM_COMPUTE_CL */
+};
+} // namespace framework
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_OPENCL_TIMER */