COMPMID-1777: Add option to make instruments output timestamps instead of duration

Change-Id: Iafc1d6cd8003de64a3439ad807f4002036c73a73

9 files changed, 240 insertions, 50 deletions

diff --git a/tests/framework/Framework.cpp b/tests/framework/Framework.cpp
index ffdecf3aa4..cc7852906c 100644
--- a/tests/framework/Framework.cpp
+++ b/tests/framework/Framework.cpp
@@ -42,9 +42,19 @@ namespace framework
 {
 Framework::Framework()
 {
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMESTAMPS, ScaleFactor::NONE), Instrument::make_instrument<WallClockTimestamps, ScaleFactor::NONE>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMESTAMPS, ScaleFactor::TIME_MS),
+                                   Instrument::make_instrument<WallClockTimestamps, ScaleFactor::TIME_MS>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMESTAMPS, ScaleFactor::TIME_S),
+                                   Instrument::make_instrument<WallClockTimestamps, ScaleFactor::TIME_S>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMER, ScaleFactor::NONE), Instrument::make_instrument<WallClockTimer, ScaleFactor::NONE>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMER, ScaleFactor::TIME_MS), Instrument::make_instrument<WallClockTimer, ScaleFactor::TIME_MS>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMER, ScaleFactor::TIME_S), Instrument::make_instrument<WallClockTimer, ScaleFactor::TIME_S>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMESTAMPS, ScaleFactor::NONE), Instrument::make_instrument<SchedulerTimestamps, ScaleFactor::NONE>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMESTAMPS, ScaleFactor::TIME_MS),
+                                   Instrument::make_instrument<SchedulerTimestamps, ScaleFactor::TIME_MS>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMESTAMPS, ScaleFactor::TIME_S),
+                                   Instrument::make_instrument<SchedulerTimestamps, ScaleFactor::TIME_S>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMER, ScaleFactor::NONE), Instrument::make_instrument<SchedulerTimer, ScaleFactor::NONE>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMER, ScaleFactor::TIME_MS), Instrument::make_instrument<SchedulerTimer, ScaleFactor::TIME_MS>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMER, ScaleFactor::TIME_S), Instrument::make_instrument<SchedulerTimer, ScaleFactor::TIME_S>);
@@ -59,6 +69,10 @@ Framework::Framework()
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::MALI, ScaleFactor::SCALE_1M), Instrument::make_instrument<MaliCounter, ScaleFactor::SCALE_1M>);
 #endif /* MALI_ENABLED */
 #ifdef ARM_COMPUTE_CL
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::NONE), Instrument::make_instrument<OpenCLTimestamps, ScaleFactor::NONE>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::TIME_US), Instrument::make_instrument<OpenCLTimestamps, ScaleFactor::TIME_US>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::TIME_MS), Instrument::make_instrument<OpenCLTimestamps, ScaleFactor::TIME_MS>);
+    _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::TIME_S), Instrument::make_instrument<OpenCLTimestamps, ScaleFactor::TIME_S>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMER, ScaleFactor::NONE), Instrument::make_instrument<OpenCLTimer, ScaleFactor::NONE>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMER, ScaleFactor::TIME_US), Instrument::make_instrument<OpenCLTimer, ScaleFactor::TIME_US>);
     _available_instruments.emplace(std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMER, ScaleFactor::TIME_MS), Instrument::make_instrument<OpenCLTimer, ScaleFactor::TIME_MS>);
diff --git a/tests/framework/instruments/Instruments.cpp b/tests/framework/instruments/Instruments.cpp
index 6d65b014e3..2288124a21 100644
--- a/tests/framework/instruments/Instruments.cpp
+++ b/tests/framework/instruments/Instruments.cpp
@@ -44,6 +44,12 @@ InstrumentsDescription instrument_type_from_name(const std::string &name)
         { "wall_clock_timer", std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMER, ScaleFactor::NONE) },
         { "wall_clock_timer_ms", std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMER, ScaleFactor::TIME_MS) },
         { "wall_clock_timer_s", std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMER, ScaleFactor::TIME_S) },
+        { "wall_clock_timestamps", std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMESTAMPS, ScaleFactor::NONE) },
+        { "wall_clock_timestamps_ms", std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMESTAMPS, ScaleFactor::TIME_MS) },
+        { "wall_clock_timestamps_s", std::pair<InstrumentType, ScaleFactor>(InstrumentType::WALL_CLOCK_TIMESTAMPS, ScaleFactor::TIME_S) },
+        { "scheduler_timestamps", std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMESTAMPS, ScaleFactor::NONE) },
+        { "scheduler_timestamps_ms", std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMESTAMPS, ScaleFactor::TIME_MS) },
+        { "scheduler_timestamps_s", std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMESTAMPS, ScaleFactor::TIME_S) },
         { "scheduler_timer", std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMER, ScaleFactor::NONE) },
         { "scheduler_timer_ms", std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMER, ScaleFactor::TIME_MS) },
         { "scheduler_timer_s", std::pair<InstrumentType, ScaleFactor>(InstrumentType::SCHEDULER_TIMER, ScaleFactor::TIME_S) },
@@ -55,6 +61,10 @@ InstrumentsDescription instrument_type_from_name(const std::string &name)
         { "mali", std::pair<InstrumentType, ScaleFactor>(InstrumentType::MALI, ScaleFactor::NONE) },
         { "mali_k", std::pair<InstrumentType, ScaleFactor>(InstrumentType::MALI, ScaleFactor::SCALE_1K) },
         { "mali_m", std::pair<InstrumentType, ScaleFactor>(InstrumentType::MALI, ScaleFactor::SCALE_1M) },
+        { "opencl_timestamps", std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::NONE) },
+        { "opencl_timestamps_us", std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::TIME_US) },
+        { "opencl_timestamps_ms", std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::TIME_MS) },
+        { "opencl_timestamps_s", std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMESTAMPS, ScaleFactor::TIME_S) },
         { "opencl_timer", std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMER, ScaleFactor::NONE) },
         { "opencl_timer_us", std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMER, ScaleFactor::TIME_US) },
         { "opencl_timer_ms", std::pair<InstrumentType, ScaleFactor>(InstrumentType::OPENCL_TIMER, ScaleFactor::TIME_MS) },
diff --git a/tests/framework/instruments/Instruments.h b/tests/framework/instruments/Instruments.h
index 77c74b7b3e..370db8d4dc 100644
--- a/tests/framework/instruments/Instruments.h
+++ b/tests/framework/instruments/Instruments.h
@@ -54,6 +54,9 @@ enum class InstrumentType : unsigned int
     OPENCL_TIMER            = 0x0400,
     SCHEDULER_TIMER         = 0x0500,
     OPENCL_MEMORY_USAGE     = 0x0600,
+    WALL_CLOCK_TIMESTAMPS   = 0x0700,
+    OPENCL_TIMESTAMPS       = 0x0800,
+    SCHEDULER_TIMESTAMPS    = 0x0900,
 };
 
 using InstrumentsDescription = std::pair<InstrumentType, ScaleFactor>;
@@ -72,6 +75,22 @@ inline ::std::stringstream &operator<<(::std::stringstream &stream, InstrumentsD
 {
     switch(instrument.first)
     {
+        case InstrumentType::WALL_CLOCK_TIMESTAMPS:
+            switch(instrument.second)
+            {
+                case ScaleFactor::NONE:
+                    stream << "WALL_CLOCK_TIMESTAMPS";
+                    break;
+                case ScaleFactor::TIME_MS:
+                    stream << "WALL_CLOCK_TIMESTAMPS_MS";
+                    break;
+                case ScaleFactor::TIME_S:
+                    stream << "WALL_CLOCK_TIMESTAMPS_S";
+                    break;
+                default:
+                    throw std::invalid_argument("Unsupported instrument scale");
+            }
+            break;
         case InstrumentType::WALL_CLOCK_TIMER:
             switch(instrument.second)
             {
@@ -88,6 +107,22 @@ inline ::std::stringstream &operator<<(::std::stringstream &stream, InstrumentsD
                     throw std::invalid_argument("Unsupported instrument scale");
             }
             break;
+        case InstrumentType::SCHEDULER_TIMESTAMPS:
+            switch(instrument.second)
+            {
+                case ScaleFactor::NONE:
+                    stream << "SCHEDULER_TIMESTAMPS";
+                    break;
+                case ScaleFactor::TIME_MS:
+                    stream << "SCHEDULER_TIMESTAMPS_MS";
+                    break;
+                case ScaleFactor::TIME_S:
+                    stream << "SCHEDULER_TIMESTAMPS_S";
+                    break;
+                default:
+                    throw std::invalid_argument("Unsupported instrument scale");
+            }
+            break;
         case InstrumentType::SCHEDULER_TIMER:
             switch(instrument.second)
             {
@@ -142,6 +177,25 @@ inline ::std::stringstream &operator<<(::std::stringstream &stream, InstrumentsD
                     throw std::invalid_argument("Unsupported instrument scale");
             }
             break;
+        case InstrumentType::OPENCL_TIMESTAMPS:
+            switch(instrument.second)
+            {
+                case ScaleFactor::NONE:
+                    stream << "OPENCL_TIMESTAMPS";
+                    break;
+                case ScaleFactor::TIME_US:
+                    stream << "OPENCL_TIMESTAMPS_US";
+                    break;
+                case ScaleFactor::TIME_MS:
+                    stream << "OPENCL_TIMESTAMPS_MS";
+                    break;
+                case ScaleFactor::TIME_S:
+                    stream << "OPENCL_TIMESTAMPS_S";
+                    break;
+                default:
+                    throw std::invalid_argument("Unsupported instrument scale");
+            }
+            break;
         case InstrumentType::OPENCL_TIMER:
             switch(instrument.second)
             {
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
diff --git a/tests/framework/instruments/OpenCLTimer.h b/tests/framework/instruments/OpenCLTimer.h
index c5f3bce92a..8722e6b348 100644
--- a/tests/framework/instruments/OpenCLTimer.h
+++ b/tests/framework/instruments/OpenCLTimer.h
@@ -41,14 +41,15 @@ namespace test
 namespace framework
 {
 /** Instrument creating measurements based on the information returned by clGetEventProfilingInfo for each OpenCL kernel executed*/
-class OpenCLTimer : public Instrument
+template <bool output_timestamps>
+class OpenCLClock : public Instrument
 {
 public:
     /** Construct an OpenCL timer.
      *
      * @param[in] scale_factor Measurement scale factor.
      */
-    OpenCLTimer(ScaleFactor scale_factor);
+    OpenCLClock(ScaleFactor scale_factor);
     std::string     id() const override;
     void            test_start() override;
     void            start() override;
@@ -60,7 +61,7 @@ private:
 #ifdef ARM_COMPUTE_CL
     struct kernel_info
     {
-        cl::Event   event{}; /**< OpenCL event associated to the kernel enqueue */
+        ::cl::Event event{}; /**< OpenCL event associated to the kernel enqueue */
         std::string name{};  /**< OpenCL Kernel name */
     };
     std::list<kernel_info>                          _kernels;
@@ -73,6 +74,10 @@ private:
 private:
     float _scale_factor{};
 };
+
+using OpenCLTimer      = OpenCLClock<false>;
+using OpenCLTimestamps = OpenCLClock<true>;
+
 } // namespace framework
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/framework/instruments/SchedulerTimer.cpp b/tests/framework/instruments/SchedulerTimer.cpp
index 76f1a58e9c..bd84a777a9 100644
--- a/tests/framework/instruments/SchedulerTimer.cpp
+++ b/tests/framework/instruments/SchedulerTimer.cpp
@@ -34,16 +34,25 @@ namespace test
 {
 namespace framework
 {
-std::string SchedulerTimer::id() const
+template <bool output_timestamps>
+std::string    SchedulerClock<output_timestamps>::id() const
 {
-    return "SchedulerTimer";
+    if(output_timestamps)
+    {
+        return "SchedulerTimestamps";
+    }
+    else
+    {
+        return "SchedulerTimer";
+    }
 }
 
+template <bool    output_timestamps>
 class Interceptor final : public IScheduler
 {
 public:
     /** Default constructor. */
-    Interceptor(std::list<SchedulerTimer::kernel_info> &kernels, IScheduler &real_scheduler, ScaleFactor scale_factor)
+    Interceptor(std::list<struct SchedulerClock<output_timestamps>::kernel_info> &kernels, IScheduler &real_scheduler, ScaleFactor scale_factor)
         : _kernels(kernels), _real_scheduler(real_scheduler), _timer(scale_factor), _prefix()
     {
     }
@@ -69,7 +78,7 @@ public:
         _real_scheduler.schedule(kernel, hints.split_dimension());
         _timer.stop();
 
-        SchedulerTimer::kernel_info info;
+        typename SchedulerClock<output_timestamps>::kernel_info info;
         info.name         = kernel->name();
         info.prefix       = _prefix;
         info.measurements = _timer.measurements();
@@ -82,7 +91,7 @@ public:
         _real_scheduler.run_tagged_workloads(workloads, tag);
         _timer.stop();
 
-        SchedulerTimer::kernel_info info;
+        typename SchedulerClock<output_timestamps>::kernel_info info;
         info.name         = tag != nullptr ? tag : "Unknown";
         info.prefix       = _prefix;
         info.measurements = _timer.measurements();
@@ -97,28 +106,30 @@ protected:
     }
 
 private:
-    std::list<SchedulerTimer::kernel_info> &_kernels;
-    IScheduler                             &_real_scheduler;
-    WallClockTimer                          _timer;
-    std::string                             _prefix;
+    std::list<struct SchedulerClock<output_timestamps>::kernel_info> &_kernels;
+    IScheduler                                                       &_real_scheduler;
+    WallClock<output_timestamps>                                      _timer;
+    std::string                                                       _prefix;
 };
 
-SchedulerTimer::SchedulerTimer(ScaleFactor scale_factor)
+template <bool output_timestamps>
+SchedulerClock<output_timestamps>::SchedulerClock(ScaleFactor scale_factor)
     : _kernels(), _real_scheduler(nullptr), _real_scheduler_type(), _real_graph_function(nullptr), _scale_factor(scale_factor), _interceptor(nullptr)
 {
 }
 
-void SchedulerTimer::test_start()
+template <bool output_timestamps>
+void           SchedulerClock<output_timestamps>::test_start()
 {
     // Start intercepting tasks:
     ARM_COMPUTE_ERROR_ON(_real_graph_function != nullptr);
     _real_graph_function  = graph::TaskExecutor::get().execute_function;
     auto task_interceptor = [this](graph::ExecutionTask & task)
     {
-        Interceptor *scheduler = nullptr;
-        if(dynamic_cast<Interceptor *>(this->_interceptor.get()) != nullptr)
+        Interceptor<output_timestamps> *scheduler = nullptr;
+        if(dynamic_cast<Interceptor<output_timestamps> *>(this->_interceptor.get()) != nullptr)
         {
-            scheduler = arm_compute::utils::cast::polymorphic_downcast<Interceptor *>(_interceptor.get());
+            scheduler = arm_compute::utils::cast::polymorphic_downcast<Interceptor<output_timestamps> *>(_interceptor.get());
             if(task.node != nullptr && !task.node->name().empty())
             {
                 scheduler->set_prefix(task.node->name() + "/");
@@ -143,18 +154,20 @@ void SchedulerTimer::test_start()
     if(_real_scheduler_type != Scheduler::Type::CUSTOM)
     {
         _real_scheduler = &Scheduler::get();
-        _interceptor    = std::make_shared<Interceptor>(_kernels, *_real_scheduler, _scale_factor);
+        _interceptor    = std::make_shared<Interceptor<output_timestamps>>(_kernels, *_real_scheduler, _scale_factor);
         Scheduler::set(std::static_pointer_cast<IScheduler>(_interceptor));
         graph::TaskExecutor::get().execute_function = task_interceptor;
     }
 }
 
-void SchedulerTimer::start()
+template <bool output_timestamps>
+void           SchedulerClock<output_timestamps>::start()
 {
     _kernels.clear();
 }
 
-void SchedulerTimer::test_stop()
+template <bool output_timestamps>
+void           SchedulerClock<output_timestamps>::test_stop()
 {
     // Restore real scheduler
     Scheduler::set(_real_scheduler_type);
@@ -164,17 +177,43 @@ void SchedulerTimer::test_stop()
     _real_graph_function                        = nullptr;
 }
 
-Instrument::MeasurementsMap SchedulerTimer::measurements() const
+template <bool              output_timestamps>
+Instrument::MeasurementsMap SchedulerClock<output_timestamps>::measurements() const
 {
     MeasurementsMap measurements;
     unsigned int    kernel_number = 0;
     for(auto kernel : _kernels)
     {
-        measurements.emplace(kernel.prefix + kernel.name + " #" + support::cpp11::to_string(kernel_number++), kernel.measurements.begin()->second);
+        std::string name = kernel.prefix + kernel.name + " #" + support::cpp11::to_string(kernel_number++);
+        if(output_timestamps)
+        {
+            ARM_COMPUTE_ERROR_ON(kernel.measurements.size() != 2);
+            for(auto m : kernel.measurements)
+            {
+                if(m.first.find("[start]") != std::string::npos)
+                {
+                    measurements.emplace("[start]" + name, m.second);
+                }
+                else if(m.first.find("[end]") != std::string::npos)
+                {
+                    measurements.emplace("[end]" + name, m.second);
+                }
+                else
+                {
+                    ARM_COMPUTE_ERROR("Measurement not handled");
+                }
+            }
+        }
+        else
+        {
+            measurements.emplace(name, kernel.measurements.begin()->second);
+        }
     }
 
     return measurements;
 }
+template class SchedulerClock<true>;
+template class SchedulerClock<false>;
 } // namespace framework
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/framework/instruments/SchedulerTimer.h b/tests/framework/instruments/SchedulerTimer.h
index 55d5f25b75..64adb488ae 100644
--- a/tests/framework/instruments/SchedulerTimer.h
+++ b/tests/framework/instruments/SchedulerTimer.h
@@ -37,19 +37,26 @@ namespace test
 namespace framework
 {
 /** Instrument creating measurements based on the information returned by clGetEventProfilingInfo for each OpenCL kernel executed*/
-class SchedulerTimer : public Instrument
+template <bool output_timestamps>
+class SchedulerClock : public Instrument
 {
 public:
     /** Construct a Scheduler timer.
      *
      * @param[in] scale_factor Measurement scale factor.
      */
-    SchedulerTimer(ScaleFactor scale_factor);
+    SchedulerClock(ScaleFactor scale_factor);
 
     /** Prevent instances of this class from being copy constructed */
-    SchedulerTimer(const SchedulerTimer &) = delete;
+    SchedulerClock(const SchedulerClock &) = delete;
     /** Prevent instances of this class from being copied */
-    SchedulerTimer &operator=(const SchedulerTimer &) = delete;
+    SchedulerClock &operator=(const SchedulerClock &) = delete;
+    /** Use the default move assignment operator */
+    SchedulerClock &operator=(SchedulerClock &&) = default;
+    /** Use the default move constructor */
+    SchedulerClock(SchedulerClock &&) = default;
+    /** Use the default destructor */
+    ~SchedulerClock() = default;
 
     std::string                 id() const override;
     void                        test_start() override;
@@ -73,6 +80,10 @@ private:
     ScaleFactor                                  _scale_factor;
     std::shared_ptr<IScheduler>                  _interceptor;
 };
+
+using SchedulerTimer      = SchedulerClock<false>;
+using SchedulerTimestamps = SchedulerClock<true>;
+
 } // namespace framework
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/framework/instruments/WallClockTimer.cpp b/tests/framework/instruments/WallClockTimer.cpp
index 6c69360236..f309efca1c 100644
--- a/tests/framework/instruments/WallClockTimer.cpp
+++ b/tests/framework/instruments/WallClockTimer.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -32,26 +32,52 @@ namespace test
 {
 namespace framework
 {
-std::string WallClockTimer::id() const
+template <bool output_timestamps>
+std::string    WallClock<output_timestamps>::id() const
 {
-    return "Wall clock";
+    if(output_timestamps)
+    {
+        return "Wall clock timestamps";
+    }
+    else
+    {
+        return "Wall clock";
+    }
 }
 
-void WallClockTimer::start()
+template <bool output_timestamps>
+void           WallClock<output_timestamps>::start()
 {
     _start = std::chrono::high_resolution_clock::now();
 }
 
-void WallClockTimer::stop()
+template <bool output_timestamps>
+void           WallClock<output_timestamps>::stop()
 {
     _stop = std::chrono::high_resolution_clock::now();
 }
 
-Instrument::MeasurementsMap WallClockTimer::measurements() const
+template <bool              output_timestamps>
+Instrument::MeasurementsMap WallClock<output_timestamps>::measurements() const
 {
-    const auto delta = std::chrono::duration_cast<std::chrono::microseconds>(_stop - _start);
-    return MeasurementsMap{ { "Wall clock time", Measurement(delta.count() / _scale_factor, _unit) } };
+    MeasurementsMap measurements;
+    if(output_timestamps)
+    {
+        // _start / _stop are in ns, so divide by an extra 1000:
+        measurements.emplace("[start]Wall clock time", Measurement(_start.time_since_epoch().count() / static_cast<uint64_t>(1000 * _scale_factor), _unit));
+        measurements.emplace("[end]Wall clock time", Measurement(_stop.time_since_epoch().count() / static_cast<uint64_t>(1000 * _scale_factor), _unit));
+    }
+    else
+    {
+        const auto delta = std::chrono::duration_cast<std::chrono::microseconds>(_stop - _start);
+        measurements.emplace("Wall clock time", Measurement(delta.count() / _scale_factor, _unit));
+    }
+    return measurements;
 }
+
+template class WallClock<true>;
+template class WallClock<false>;
+
 } // namespace framework
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/framework/instruments/WallClockTimer.h b/tests/framework/instruments/WallClockTimer.h
index c9829aea12..d659ab1a80 100644
--- a/tests/framework/instruments/WallClockTimer.h
+++ b/tests/framework/instruments/WallClockTimer.h
@@ -35,14 +35,15 @@ namespace test
 namespace framework
 {
 /** Implementation of an instrument to measure elapsed wall-clock time in milliseconds. */
-class WallClockTimer : public Instrument
+template <bool output_timestamps>
+class WallClock : public Instrument
 {
 public:
     /** Construct a Wall clock timer.
      *
      * @param[in] scale_factor Measurement scale factor.
      */
-    WallClockTimer(ScaleFactor scale_factor)
+    WallClock(ScaleFactor scale_factor)
     {
         switch(scale_factor)
         {
@@ -73,6 +74,9 @@ private:
     std::chrono::high_resolution_clock::time_point _stop{};
     float                                          _scale_factor{};
 };
+
+using WallClockTimer      = WallClock<false>;
+using WallClockTimestamps = WallClock<true>;
 } // namespace framework
 } // namespace test
 } // namespace arm_compute