From 72f4ae5a53fe24226ff16ed9c339171887d74874 Mon Sep 17 00:00:00 2001
From: Anthony Barbier <anthony.barbier@arm.com>
Date: Wed, 7 Nov 2018 17:33:54 +0000
Subject: COMPMID-1777: Add option to make instruments output timestamps
 instead of duration

Change-Id: Iafc1d6cd8003de64a3439ad807f4002036c73a73
---
 tests/framework/instruments/OpenCLTimer.cpp | 55 +++++++++++++++++++++--------
 1 file changed, 41 insertions(+), 14 deletions(-)

(limited to 'tests/framework/instruments/OpenCLTimer.cpp')

diff --git a/tests/framework/instruments/OpenCLTimer.cpp b/tests/framework/instruments/OpenCLTimer.cpp
index 4391c433e5..b23b8a8878 100644
--- a/tests/framework/instruments/OpenCLTimer.cpp
+++ b/tests/framework/instruments/OpenCLTimer.cpp
@@ -39,12 +39,21 @@ namespace test
 {
 namespace framework
 {
-std::string OpenCLTimer::id() const
+template <bool output_timestamps>
+std::string    OpenCLClock<output_timestamps>::id() const
 {
-    return "OpenCLTimer";
+    if(output_timestamps)
+    {
+        return "OpenCLTimestamps";
+    }
+    else
+    {
+        return "OpenCLTimer";
+    }
 }
 
-OpenCLTimer::OpenCLTimer(ScaleFactor scale_factor)
+template <bool output_timestamps>
+OpenCLClock<output_timestamps>::OpenCLClock(ScaleFactor scale_factor)
     : _kernels(), _real_function(nullptr), _real_graph_function(nullptr), _prefix(), _timer_enabled(false)
 {
     auto                        q     = CLScheduler::get().queue();
@@ -77,7 +86,8 @@ OpenCLTimer::OpenCLTimer(ScaleFactor scale_factor)
     }
 }
 
-void OpenCLTimer::test_start()
+template <bool output_timestamps>
+void           OpenCLClock<output_timestamps>::test_start()
 {
     // Start intercepting enqueues:
     ARM_COMPUTE_ERROR_ON(_real_function != nullptr);
@@ -100,9 +110,9 @@ void OpenCLTimer::test_start()
             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;
+            kernel_info       info;
+            cl::Kernel        cpp_kernel(kernel, true);
+            std::stringstream ss;
             ss << this->_prefix << cpp_kernel.getInfo<CL_KERNEL_FUNCTION_NAME>();
             if(gws != nullptr)
             {
@@ -144,17 +154,20 @@ void OpenCLTimer::test_start()
     graph::TaskExecutor::get().execute_function = task_interceptor;
 }
 
-void OpenCLTimer::start()
+template <bool output_timestamps>
+void           OpenCLClock<output_timestamps>::start()
 {
     _kernels.clear();
     _timer_enabled = true;
 }
-void OpenCLTimer::stop()
+template <bool output_timestamps>
+void           OpenCLClock<output_timestamps>::stop()
 {
     _timer_enabled = false;
 }
 
-void OpenCLTimer::test_stop()
+template <bool output_timestamps>
+void           OpenCLClock<output_timestamps>::test_stop()
 {
     // Restore real function
     CLSymbols::get().clEnqueueNDRangeKernel_ptr = _real_function;
@@ -163,20 +176,34 @@ void OpenCLTimer::test_stop()
     _real_function                              = nullptr;
 }
 
-Instrument::MeasurementsMap OpenCLTimer::measurements() const
+template <bool              output_timestamps>
+Instrument::MeasurementsMap OpenCLClock<output_timestamps>::measurements() const
 {
     MeasurementsMap measurements;
     unsigned int    kernel_number = 0;
     for(auto kernel : _kernels)
     {
-        cl_ulong start = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
-        cl_ulong end   = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
+        cl_ulong start, end;
+        kernel.event.getProfilingInfo(CL_PROFILING_COMMAND_START, &start);
+        kernel.event.getProfilingInfo(CL_PROFILING_COMMAND_END, &end);
+        std::string name = kernel.name + " #" + support::cpp11::to_string(kernel_number++);
 
-        measurements.emplace(kernel.name + " #" + support::cpp11::to_string(kernel_number++), Measurement((end - start) / _scale_factor, _unit));
+        if(output_timestamps)
+        {
+            measurements.emplace("[start]" + name, Measurement(start / static_cast<cl_ulong>(_scale_factor), _unit));
+            measurements.emplace("[end]" + name, Measurement(end / static_cast<cl_ulong>(_scale_factor), _unit));
+        }
+        else
+        {
+            measurements.emplace(name, Measurement((end - start) / _scale_factor, _unit));
+        }
     }
 
     return measurements;
 }
+
+template class OpenCLClock<true>;
+template class OpenCLClock<false>;
 } // namespace framework
 } // namespace test
 } // namespace arm_compute
-- 
cgit v1.2.1