From 7c63aa5f1ca9e80013e0db438807c5188d4e0d77 Mon Sep 17 00:00:00 2001
From: Harisankar Sadasivan <hari.sadasivan@amd.com>
Date: Wed, 31 Jul 2024 19:07:53 -0400
Subject: [PATCH] changes for debug

---
 .../01_gemm/run_gemm_example_streamk_v2.inc   | 28 +++++++++----------
 include/ck/host_utility/kernel_launch.hpp     |  2 +-
 include/ck/stream_config.hpp                  |  4 +--
 .../gridwise_gemm_xdl_cshuffle_streamk_v3.hpp |  2 +-
 script/cmake-ck-dev.sh                        |  2 +-
 5 files changed, 19 insertions(+), 19 deletions(-)

diff --git a/example/01_gemm/run_gemm_example_streamk_v2.inc b/example/01_gemm/run_gemm_example_streamk_v2.inc
index fbde664cb2..bfff7d4a25 100644
--- a/example/01_gemm/run_gemm_example_streamk_v2.inc
+++ b/example/01_gemm/run_gemm_example_streamk_v2.inc
@@ -267,7 +267,7 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
 
         auto ref_argument = ref_gemm.MakeArgument(
             a_m_k, b_k_n, c_m_n_host_result, PassThrough{}, PassThrough{}, PassThrough{});
-
+        printf("inside do verification\n");
         ref_invoker.Run(ref_argument);
 
         ave_time = invoker.Run(argument, StreamConfig{nullptr, false, 1});
@@ -281,19 +281,19 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
         return ck::utils::check_err(c_m_n_device_result_converted, c_m_n_host_result);
 #else
         printf("device copy initiated\n"); // HS
-        if((workspace_size != 0) && (Streamk_sel > 0))
-        {
-            printf("entered if\n");
-            workspace.FromDevice(c_m_n_device_result.mData.data());
-        }
-        else
-            c_m_n_device_buf.FromDevice(c_m_n_device_result.mData.data());
-        printf("device copy finished\n"); // HS
-        pass &= ck::utils::check_err(c_m_n_device_result,
-                                     c_m_n_host_result,
-                                     "Error: Incorrect results!",
-                                     get_rtol<CDataType>(),
-                                     get_atol<CDataType>());
+        // if((workspace_size != 0) && (Streamk_sel > 0))
+        // {
+        //     printf("entered if\n");
+        //     workspace.FromDevice(c_m_n_device_result.mData.data());
+        // }
+        // else
+        //     c_m_n_device_buf.FromDevice(c_m_n_device_result.mData.data());
+        // printf("device copy finished\n"); // HS
+        // pass &= ck::utils::check_err(c_m_n_device_result,
+        //                              c_m_n_host_result,
+        //                              "Error: Incorrect results!",
+        //                              get_rtol<CDataType>(),
+        //                              get_atol<CDataType>());
 #endif
     }
 
diff --git a/include/ck/host_utility/kernel_launch.hpp b/include/ck/host_utility/kernel_launch.hpp
index a616433ac9..24c91929a1 100644
--- a/include/ck/host_utility/kernel_launch.hpp
+++ b/include/ck/host_utility/kernel_launch.hpp
@@ -109,7 +109,7 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
             printf("Warm up %d times\n", stream_config.cold_niters_);
         }
         // warm up
-        preprocess();
+        //HS preprocess();
         for(int i = 0; i < stream_config.cold_niters_; ++i)
         {
             kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
diff --git a/include/ck/stream_config.hpp b/include/ck/stream_config.hpp
index 37ba250cf5..7073584940 100644
--- a/include/ck/stream_config.hpp
+++ b/include/ck/stream_config.hpp
@@ -11,8 +11,8 @@ struct StreamConfig
     hipStream_t stream_id_ = nullptr;
     bool time_kernel_      = false;
     int log_level_         = 0;
-    int cold_niters_       = 5;
-    int nrepeat_           = 50;
+    int cold_niters_       = 0;//HS 
+    int nrepeat_           = 1;//HS
 
     bool flush_cache   = false;
     int rotating_count = 1;
diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp
index 21a0e8e9a7..b17a816ad5 100644
--- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp
+++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp
@@ -1186,7 +1186,6 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3
                                void* p_workspace,
                                Block2CTileMap_streamk block_2_ctile_map_streamk)
     {
-
         const AElementwiseOperation a_element_op{};
         const BElementwiseOperation b_element_op{};
         const CElementwiseOperation c_element_op{};
@@ -1899,6 +1898,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3
                 }
             }
         } // for loop
+        if(threadIdx.x==0)printf("kernel ends");
     }
 
     template <bool HasMainKBlockLoop,
diff --git a/script/cmake-ck-dev.sh b/script/cmake-ck-dev.sh
index 26326523f4..d1e8385f6c 100755
--- a/script/cmake-ck-dev.sh
+++ b/script/cmake-ck-dev.sh
@@ -8,7 +8,7 @@ MY_PROJECT_SOURCE=$1
 if [ $# -ge 2 ] ; then
     GPU_TARGETS=$2
 else
-    GPU_TARGETS="gfx908;gfx90a;gfx940"
+    GPU_TARGETS="gfx90a"
 fi
 
 cmake                                                                                             \