[Beignet] [PATCH 12/18] Backend: Add the profiling output with source.

[junyan.he at inbox.com]

From: Junyan He <junyan.he at linux.intel.com>

We provide a more direct way for profiling with source.
The output would be:

Format: Average Timestamp  Exec number Source
                             |  __kernel void
TS:       49  Num: 256 ----> |  test_fill_image_2d_array(__write_only image2d_array_t dst)
                             |  {
                             |    int coordx;
TS:       59  Num: 256 ----> |    int coordy;
                             |    int coordz;
                             |    coordx = (int)get_global_id(0);
                             |    coordy = (int)get_global_id(1);
                             |    coordz = (int)get_global_id(2);
                             |    uint4 color4 = {0, 1, 2 ,3};
                             |    if (coordz < 7)
                             |      write_imageui(dst, (int3)(coordx, coordy, coordz), color4);
TS:      893  Num: 256 ----> |  }

We also support 3 output mode,
level 1: just brief timestamp with line number.
level 2: timestamp with source
level 3: output the detail logs.

Signed-off-by: Junyan He <junyan.he at linux.intel.com>
---
 backend/src/ir/profiling.cpp | 72 +++++++++++++++++++++++++++++++++++++++++++-
 backend/src/ir/profiling.hpp |  1 +
 2 files changed, 72 insertions(+), 1 deletion(-)

diff --git a/backend/src/ir/profiling.cpp b/backend/src/ir/profiling.cpp
index 09537fa..075945a 100644
--- a/backend/src/ir/profiling.cpp
+++ b/backend/src/ir/profiling.cpp
@@ -22,6 +22,10 @@
 
 #include <stdio.h>
 #include <stdlib.h>
+#include <fstream>
+#include <dlfcn.h>
+#include <sstream>
+#include <iostream>
 #include "ir/profiling.hpp"
 #include "src/cl_device_data.h"
 
@@ -31,12 +35,78 @@ namespace ir
 {
   pthread_mutex_t ProfilingInfo::lock = PTHREAD_MUTEX_INITIALIZER;
 
+  void ProfilingInfo::outputResultInSource(ProfilingReportItem* log, uint32_t num)
+  {
+    uint32_t tsPoints[MaxTimestampProfilingPoints];
+    uint32_t tsPointsValid[MaxTimestampProfilingPoints];
+    for (uint32_t i = 0; i < MaxTimestampProfilingPoints; i++) {
+      uint32_t ts = 0;
+      tsPointsValid[i] = 0;
+      for (uint32_t j = 0; j < num; j++) {
+        ts += log[j].userTimestamp[i];
+        if (log[j].userTimestamp[i])
+          tsPointsValid[i]++;
+      }
+      if (tsPointsValid[i]) {
+        tsPoints[i] = ts/tsPointsValid[i];
+      } else {
+        tsPoints[i] = 0;
+      }
+    }
+
+    for (uint32_t i = 0; i < MaxTimestampProfilingPoints; i++) {
+      if (profPoints.find(i) != profPoints.end()) {
+        printf("Line %8d:    Timestamp:%10d  Thread Exec:%d\n",
+            profPoints[i], tsPoints[i], tsPointsValid[i]);
+      }
+    }
+
+    if (profilingType <= ProfilingSimple)
+      return;
+
+    if (source.length() == 0)
+      return;
+
+    std::map<uint32_t, uint32_t> lineProfMap;
+    for (auto &e : profPoints) {
+      lineProfMap.insert(std::pair<uint32_t, uint32_t>(e.second, e.first));
+    }
+
+    uint32_t l = 1;
+    std::string oneLine;
+    std::istringstream srcStream(source);
+    printf("\n\nFormat: Average Timestamp  \tExec number \tSource\n");
+    while (getline(srcStream, oneLine, '\n')) {
+      if (lineProfMap.find(l) == lineProfMap.end()) {
+        printf("                             |  %s\n", oneLine.c_str());
+      } else {
+        uint32_t p = lineProfMap[l];
+        printf("TS:%9d  Num:%4d ----> |  %s\n",
+            tsPoints[p], tsPointsValid[p], oneLine.c_str());
+      }
+      l++;
+    }
+  }
+
   void ProfilingInfo::outputProfilingInfo(void * logBuf)
   {
     LockOutput lock;
+
     uint32_t logNum = *reinterpret_cast<uint32_t*>(logBuf);
-    printf("Total log number is %u\n", logNum);
     ProfilingReportItem* log = reinterpret_cast<ProfilingReportItem*>((char*)logBuf + 4);
+
+    if (logNum == 0) {
+      printf("Warning: No profiling log found\n");
+      return;
+    }
+
+    printf("Total log number is %u\n", logNum);
+
+    outputResultInSource(log, logNum);
+
+    if (profilingType <= ProfilingNormal)
+      return;
+
     for (int i = 0; i < (int)logNum; i++) {
       GBE_ASSERT(log->simdType == ProfilingSimdType8 || log->simdType == ProfilingSimdType16);
       uint32_t simd = log->simdType == ProfilingSimdType16 ? 16 : 8;
diff --git a/backend/src/ir/profiling.hpp b/backend/src/ir/profiling.hpp
index d36f978..57eb08f 100644
--- a/backend/src/ir/profiling.hpp
+++ b/backend/src/ir/profiling.hpp
@@ -143,6 +143,7 @@ namespace gbe
       }
 
     private:
+      void outputResultInSource(ProfilingReportItem* log, uint32_t num);
       uint32_t bti;
       uint32_t profilingType;
       uint32_t deviceID;
-- 
1.9.1