diff --git a/include/ck/host_utility/kernel_launch.hpp b/include/ck/host_utility/kernel_launch.hpp
index b4d5a440f3..4ce05035fa 100755
--- a/include/ck/host_utility/kernel_launch.hpp
+++ b/include/ck/host_utility/kernel_launch.hpp
@@ -20,6 +20,7 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
 #if CK_TIME_KERNEL
     if(stream_config.time_kernel_)
     {
+#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
         if(ck::get_device_name() == "gfx940" || ck::get_device_name() == "gfx941" ||
            ck::get_device_name() == "gfx942")
         {
@@ -48,6 +49,7 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
                 (1000.0 / total_time); // we need longer runtime to ramp up the clk on MI300s
             stream_config.nrepeat_ = stream_config.cold_niters_;
         }
+#endif
 #if DEBUG_LOG
         printf("%s: grid_dim {%d, %d, %d}, block_dim {%d, %d, %d} \n",
                __func__,
@@ -121,6 +123,36 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
 #if CK_TIME_KERNEL
     if(stream_config.time_kernel_)
     {
+#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+        if(ck::get_device_name() == "gfx940" || ck::get_device_name() == "gfx941" ||
+           ck::get_device_name() == "gfx942")
+        {
+            hipEvent_t start, stop;
+
+            hip_check_error(hipEventCreate(&start));
+            hip_check_error(hipEventCreate(&stop));
+
+            hip_check_error(hipDeviceSynchronize());
+            hip_check_error(hipEventRecord(start, stream_config.stream_id_));
+
+            for(int i = 0; i < stream_config.nrepeat_; ++i)
+            {
+                kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
+                hip_check_error(hipGetLastError());
+            }
+
+            hip_check_error(hipEventRecord(stop, stream_config.stream_id_));
+            hip_check_error(hipEventSynchronize(stop));
+
+            float total_time = 0;
+
+            hip_check_error(hipEventElapsedTime(&total_time, start, stop));
+            total_time /= 10;
+            stream_config.cold_niters_ =
+                (1000.0 / total_time); // we need longer runtime to ramp up the clk on MI300s
+            stream_config.nrepeat_ = stream_config.cold_niters_;
+        }
+#endif
 #if DEBUG_LOG
         printf("%s: grid_dim {%d, %d, %d}, block_dim {%d, %d, %d} \n",
                __func__,