use hip launch ext

library/include/ck/library/host_tensor/device.hpp

@@ -6,6 +6,7 @@
 #include <chrono>
 #include <hip/hip_runtime.h>
 #include <hip/hip_fp16.h>
+#include <hip/hip_ext.h>
 
 #include "stream_config.hpp"
 #include "ck/options.hpp"
@@ -97,17 +98,34 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
 
         printf("Start running %d times...\n", nrepeat);
 
-        KernelTimer timer;
-        timer.Start();
+        //KernelTimer timer;
+        //timer.Start();
+
+        hipEvent_t start, stop;
+        hip_check_error(hipEventCreate(&start));
+        hip_check_error(hipEventCreate(&stop));
+
+        float total_time = 0.0f;
+
 
         for(int i = 0; i < nrepeat; ++i)
         {
-            kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
+            hipExtLaunchKernelGGL(kernel, grid_dim, block_dim, 0, nullptr, start, stop, 0,
+                                            args...);
+
+            //kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
+            hip_check_error(hipEventSynchronize(stop));
+
+            float time = 0.0f;
+            hip_check_error(hipEventElapsedTime(&time, start, stop));
+            total_time += time;
+
         }
 
-        timer.End();
+        //timer.End();
 
-        return timer.GetElapsedTime() / nrepeat;
+        //return timer.GetElapsedTime() / nrepeat;
+        return total_time/nrepeat;
     }
     else
     {

library/src/host_tensor/device.cpp

@@ -39,13 +39,13 @@ struct KernelTimerImpl
 
     void Start()
     {
-        hip_check_error(hipDeviceSynchronize());
-        hip_check_error(hipEventRecord(mStart, nullptr));
+        //hip_check_error(hipDeviceSynchronize());
+        //hip_check_error(hipEventRecord(mStart, nullptr));
     }
 
     void End()
     {
-        hip_check_error(hipEventRecord(mEnd, nullptr));
+        //hip_check_error(hipEventRecord(mEnd, nullptr));
         hip_check_error(hipEventSynchronize(mEnd));
     }