From e28c562a20a46c9c60533faaf39a3b9ab3117179 Mon Sep 17 00:00:00 2001
From: Cloud Han <cloudhan@outlook.com>
Date: Tue, 18 Jan 2022 13:51:27 +0800
Subject: add log and simple statistics for async kernel

---
 examples/OpenCLVectorAdd/OpenCLVectorAdd.cpp | 26 +++++++++++++++++++++++++-
 1 file changed, 25 insertions(+), 1 deletion(-)

(limited to 'examples')

diff --git a/examples/OpenCLVectorAdd/OpenCLVectorAdd.cpp b/examples/OpenCLVectorAdd/OpenCLVectorAdd.cpp
index 36d4a12f..71184e00 100644
--- a/examples/OpenCLVectorAdd/OpenCLVectorAdd.cpp
+++ b/examples/OpenCLVectorAdd/OpenCLVectorAdd.cpp
@@ -1,3 +1,4 @@
+#include <algorithm>
 #include <iostream>
 #include <cassert>
 #include <string>
@@ -137,7 +138,10 @@ int main()
         }
     }
 
-    for (int i = 0; i < 100; ++i)
+    const int numLaunchs = 100;
+    std::vector<cl_event> kernelLaunchEvts;
+    kernelLaunchEvts.reserve(numLaunchs);
+    for (int i = 0; i < numLaunchs; ++i)
     {
         int n_value = static_cast<int>(N);
         ZoneScopedN("VectorAdd Kernel Launch");
@@ -151,6 +155,8 @@ int main()
         cl_event vectorAddKernelEvent;
         CL_ASSERT(clEnqueueNDRangeKernel(commandQueue, vectorAddKernel, 1, nullptr, &N, nullptr, 0, nullptr, &vectorAddKernelEvent));
         TracyCLZoneSetEvent(vectorAddKernelEvent);
+        CL_ASSERT(clRetainEvent(vectorAddKernelEvent));
+        kernelLaunchEvts.push_back(vectorAddKernelEvent);
         std::cout << "VectorAdd Kernel Enqueued" << std::endl;
     }
 
@@ -164,6 +170,24 @@ int main()
     }
 
     CL_ASSERT(clFinish(commandQueue));
+    std::vector<float> durations(kernelLaunchEvts.size());
+    for (int i=0; i<kernelLaunchEvts.size(); i++) {
+        cl_event evt = kernelLaunchEvts[i];
+        cl_ulong start;
+        cl_ulong end;
+        CL_ASSERT(clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, nullptr));
+        CL_ASSERT(clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, nullptr));
+        CL_ASSERT(clReleaseEvent(evt));
+        durations[i] = (end - start) * 0.001;
+        std::cout << "VectorAdd Kernel " << i << " tooks " <<  static_cast<int>(durations[i]) << "us" << std::endl;
+    };
+    float avg = std::accumulate(durations.cbegin(), durations.cend(), 0.0f) / durations.size();
+    float stddev2 = std::accumulate(durations.cbegin(), durations.cend(), 0.0f, [avg](const float& acc, const float& v) {
+        auto d = v - avg;
+        return acc + d*d;
+    }) / (durations.size() - 1.0f);
+    std::cout << "VectorAdd runtime avg: " << avg << "us, std: " << sqrt(stddev2) << "us over " << numLaunchs << " runs." << std::endl;
+
 
     TracyCLCollect(tracyCLCtx);
 
-- 
cgit v1.2.3