From 27199ba5ef2a99f43716c7983daeda737a230831 Mon Sep 17 00:00:00 2001
From: dummycoderfe <noplydummmycoder@163.com>
Date: Tue, 29 Oct 2024 09:47:07 +0000
Subject: [PATCH] debug add cache

---
 include/ck_tile/host/device_memory.hpp |  1 +
 include/ck_tile/host/kernel_launch.hpp | 16 +++++++++++-----
 include/ck_tile/host/stream_config.hpp |  4 ++++
 3 files changed, 16 insertions(+), 5 deletions(-)

diff --git a/include/ck_tile/host/device_memory.hpp b/include/ck_tile/host/device_memory.hpp
index 7c8549f74f..bbaacecfe8 100644
--- a/include/ck_tile/host/device_memory.hpp
+++ b/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 {
diff --git a/include/ck_tile/host/kernel_launch.hpp b/include/ck_tile/host/kernel_launch.hpp
index 5c7bf12bfc..c2bae890c4 100644
--- a/include/ck_tile/host/kernel_launch.hpp
+++ b/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 {};
diff --git a/include/ck_tile/host/stream_config.hpp b/include/ck_tile/host/stream_config.hpp
index 47cf0fd5e4..c3a7262f7b 100644
--- a/include/ck_tile/host/stream_config.hpp
+++ b/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