debug add cache

include/ck_tile/host/device_memory.hpp

@@ -6,6 +6,7 @@
 #include <hip/hip_runtime.h>
 #include <stdint.h>
 #include <stdexcept>
+#include <iostream>
 #include "ck_tile/host/hip_check_error.hpp"
 
 namespace ck_tile {

include/ck_tile/host/kernel_launch.hpp

@@ -78,15 +78,21 @@ CK_TILE_HOST float launch_kernel(const stream_config& s, Callables... callables)
     }
     if(s.is_gpu_timer_) {
         gpu_timer timer {};
-
+        float total_time = 0;
         // warmup
         for(int i = 0; i < s.cold_niters_; i++) { (callables(s),...); } HIP_CHECK_ERROR(hipGetLastError());
 
-        timer.start(s.stream_id_);
-        for(int i = 0; i < s.nrepeat_; i++) { (callables(s),...); } HIP_CHECK_ERROR(hipGetLastError());
-        timer.stop(s.stream_id_);
+        for(int i = 0; i < s.nrepeat_; i++) { 
+            if (s.clear_cache) {
+                s.cache_buf.SetValue<int>(i);
+            }
+            timer.start(s.stream_id_);
+            (callables(s),...);
+            timer.stop(s.stream_id_);
+            total_time += timer.duration();
+        } HIP_CHECK_ERROR(hipGetLastError());
 
-        return timer.duration() / s.nrepeat_;
+        return total_time / s.nrepeat_;
     }
     else {
         cpu_timer timer {};

include/ck_tile/host/stream_config.hpp

@@ -4,6 +4,7 @@
 #pragma once
 
 #include <hip/hip_runtime.h>
+#include "device_memory.hpp"
 
 namespace ck_tile {
 /*
@@ -30,5 +31,8 @@ struct stream_config
     int cold_niters_       = 3;
     int nrepeat_           = 10;
     bool is_gpu_timer_     = true; // keep compatible
+    bool clear_cache = false;
+    size_t buf_size = 0;
+    DeviceMem cache_buf{buf_size};
 };
 } // namespace ck_tile