IVGCVSW-7222 Fix incorrect kernel measurements in profiling output

* Some CL kernels are not run after the first inference and this breaks
  the profiler which is expecting a measurement for every kernel each run
* Add a function HasKernelMeasurements() to ascertain if the Event is
  returning kernel measurements and if so insert 0.0 values for any missing
  kernel measurements.
* Fix ExecuteNetwork to only print a json object after all inferences
  have completed


Signed-off-by: Kevin May <kevin.may@arm.com>
Change-Id: I99f2bb0db847f5a52ab4c5705b072155c6b6f333

9 files changed, 48 insertions, 1 deletions

diff --git a/src/armnn/Instrument.hpp b/src/armnn/Instrument.hpp
index 8dde2ecefa..11505d1b46 100644
--- a/src/armnn/Instrument.hpp
+++ b/src/armnn/Instrument.hpp
@@ -61,6 +61,8 @@ public:
 
     virtual const char* GetName() const = 0;
 
+    virtual bool HasKernelMeasurements() const { return false;}
+
 };
 
 } //namespace armnn
diff --git a/src/armnn/Profiling.cpp b/src/armnn/Profiling.cpp
index 805b61e8ef..db2962ecad 100644
--- a/src/armnn/Profiling.cpp
+++ b/src/armnn/Profiling.cpp
@@ -316,6 +316,20 @@ void ExtractJsonObjects(unsigned int inferenceIndex,
     }
     std::vector<Measurement> instrumentMeasurements = parentEvent->GetMeasurements();
     unsigned int childIdx = 0;
+    unsigned int numSkippedKernels = 0;
+    if (inferenceIndex > 0)
+    {
+        for (auto &i: parentEvent->GetInstruments())
+        {
+            if (i->HasKernelMeasurements())
+            {
+                numSkippedKernels = static_cast<unsigned int>(parentObject.m_Children.size() -
+                                                               instrumentMeasurements.size());
+                childIdx = numSkippedKernels;
+            }
+        }
+    }
+
     for (size_t measurementIndex = 0; measurementIndex < instrumentMeasurements.size(); ++measurementIndex, ++childIdx)
     {
         if (inferenceIndex == 0)
@@ -328,6 +342,13 @@ void ExtractJsonObjects(unsigned int inferenceIndex,
             ARMNN_ASSERT(parentObject.NumChildren() == childIdx);
             parentObject.AddChild(measurementObject);
         }
+        else
+        {
+            if (numSkippedKernels > 0)
+            {
+                parentObject.GetChild(--numSkippedKernels).AddMeasurement(0.0);
+            }
+        }
 
         parentObject.GetChild(childIdx).AddMeasurement(instrumentMeasurements[measurementIndex].m_Value);
     }
diff --git a/src/armnn/ProfilingEvent.cpp b/src/armnn/ProfilingEvent.cpp
index e341344544..9f045ffabb 100644
--- a/src/armnn/ProfilingEvent.cpp
+++ b/src/armnn/ProfilingEvent.cpp
@@ -66,6 +66,11 @@ const std::vector<Measurement> Event::GetMeasurements() const
     return measurements;
 }
 
+const std::vector<Event::InstrumentPtr>& Event::GetInstruments() const
+{
+    return m_Instruments;
+}
+
 const std::string& Event::GetName() const
 {
     return m_EventName;
diff --git a/src/armnn/ProfilingEvent.hpp b/src/armnn/ProfilingEvent.hpp
index 3b9d5261ef..e4e26ec235 100644
--- a/src/armnn/ProfilingEvent.hpp
+++ b/src/armnn/ProfilingEvent.hpp
@@ -55,6 +55,10 @@ public:
     /// \return Recorded measurements of the event
     const std::vector<Measurement> GetMeasurements() const;
 
+    /// Get the Instruments used by this Event
+    /// \return Return a reference to the collection of Instruments
+    const std::vector<InstrumentPtr>& GetInstruments() const;
+
     /// Get the name of the event
     /// \return Name of the event
     const std::string& GetName() const;
diff --git a/src/backends/cl/OpenClTimer.cpp b/src/backends/cl/OpenClTimer.cpp
index a4958c10e0..a39377505d 100644
--- a/src/backends/cl/OpenClTimer.cpp
+++ b/src/backends/cl/OpenClTimer.cpp
@@ -85,6 +85,11 @@ void OpenClTimer::Stop()
     CLSymbols::get().clEnqueueNDRangeKernel_ptr = m_OriginalEnqueueFunction;
 }
 
+bool OpenClTimer::HasKernelMeasurements() const
+{
+    return m_Kernels.size() > 0;
+}
+
 std::vector<Measurement> OpenClTimer::GetMeasurements() const
 {
     std::vector<Measurement> measurements;
diff --git a/src/backends/cl/OpenClTimer.hpp b/src/backends/cl/OpenClTimer.hpp
index 5539e885d8..e517fa4b8a 100644
--- a/src/backends/cl/OpenClTimer.hpp
+++ b/src/backends/cl/OpenClTimer.hpp
@@ -29,6 +29,9 @@ public:
     /// Stop the OpenCl timer
     void Stop() override;
 
+    /// Return true if this Instrument has kernels for recording measurements
+    bool HasKernelMeasurements() const override;
+
     /// Get the name of the timer
     /// \return Name of the timer
     const char* GetName() const override { return "OpenClKernelTimer"; }
diff --git a/src/backends/neon/NeonTimer.cpp b/src/backends/neon/NeonTimer.cpp
index a7d3032b13..dbb1503d24 100644
--- a/src/backends/neon/NeonTimer.cpp
+++ b/src/backends/neon/NeonTimer.cpp
@@ -42,6 +42,11 @@ void NeonTimer::Stop()
     m_RealScheduler = nullptr;
 }
 
+bool NeonTimer::HasKernelMeasurements() const
+{
+    return m_Kernels.size() > 0;
+}
+
 std::vector<Measurement> NeonTimer::GetMeasurements() const
 {
     std::vector<Measurement> measurements = m_Kernels;
diff --git a/src/backends/neon/NeonTimer.hpp b/src/backends/neon/NeonTimer.hpp
index 31d3e85a7c..f2be7dc326 100644
--- a/src/backends/neon/NeonTimer.hpp
+++ b/src/backends/neon/NeonTimer.hpp
@@ -30,6 +30,8 @@ public:
 
     void Stop() override;
 
+    bool HasKernelMeasurements() const override;
+
     std::vector<Measurement> GetMeasurements() const override;
 
     const char* GetName() const override;
diff --git a/tests/ExecuteNetwork/ArmNNExecutor.cpp b/tests/ExecuteNetwork/ArmNNExecutor.cpp
index b894db8049..b655ef8bc3 100644
--- a/tests/ExecuteNetwork/ArmNNExecutor.cpp
+++ b/tests/ExecuteNetwork/ArmNNExecutor.cpp
@@ -184,7 +184,7 @@ void ArmNNExecutor::ExecuteSync()
         const auto inferenceDuration = armnn::GetTimeDuration(start_time);
 
         // If profiling is enabled print out the results
-        if(profiler && profiler->IsProfilingEnabled())
+        if(profiler && profiler->IsProfilingEnabled() && x == (m_Params.m_Iterations - 1))
         {
             profiler->Print(std::cout);
         }