From 04510ebe52c4f616e2cf29fb19484f8e094f6881 Mon Sep 17 00:00:00 2001
From: Amr Akmal Moustafa Abouelmagd <abouelmagd1@dane6.llnl.gov>
Date: Wed, 13 Aug 2025 12:18:09 -0700
Subject: [PATCH 1/3] TESTING only: enforce larger time span for kernel
 execution

---
 src/stream/TRIAD-Hip.cpp | 9 +++++++--
 src/stream/TRIAD.hpp     | 1 +
 2 files changed, 8 insertions(+), 2 deletions(-)

diff --git a/src/stream/TRIAD-Hip.cpp b/src/stream/TRIAD-Hip.cpp
index e9654d5f0..8c414e979 100644
--- a/src/stream/TRIAD-Hip.cpp
+++ b/src/stream/TRIAD-Hip.cpp
@@ -84,14 +84,19 @@ void TRIAD::runHipVariantImpl(VariantID vid)
   } else if ( vid == RAJA_HIP ) {
 
     startTimer();
+
     for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
 
       RAJA::forall< RAJA::hip_exec<block_size, true /*async*/> >( res,
         RAJA::RangeSegment(ibegin, iend), [=] __device__ (Index_type i) {
-        TRIAD_BODY;
-      });
 
+          for (RepIndex_type extra_rep = 0; extra_rep < extra_kernel_reps; ++extra_rep) {
+            TRIAD_BODY;
+          }
+
+      });
     }
+
     stopTimer();
 
   } else {
diff --git a/src/stream/TRIAD.hpp b/src/stream/TRIAD.hpp
index f901314a4..8f31380f1 100644
--- a/src/stream/TRIAD.hpp
+++ b/src/stream/TRIAD.hpp
@@ -26,6 +26,7 @@
 #define TRIAD_BODY  \
   a[i] = b[i] + alpha * c[i] ;
 
+constexpr int extra_kernel_reps = 5;
 
 #include "common/KernelBase.hpp"
 

From dd132de3e263108bdb1cfde84dad8683aec8851a Mon Sep 17 00:00:00 2001
From: Amr Akmal Moustafa Abouelmagd <abouelmagd1@tuolumne2151.llnl.gov>
Date: Wed, 13 Aug 2025 15:16:38 -0700
Subject: [PATCH 2/3] Basic_MAT_MAT_SHARED

---
 src/basic/MAT_MAT_SHARED-Hip.cpp | 3 ++-
 src/basic/MAT_MAT_SHARED.hpp     | 1 +
 2 files changed, 3 insertions(+), 1 deletion(-)

diff --git a/src/basic/MAT_MAT_SHARED-Hip.cpp b/src/basic/MAT_MAT_SHARED-Hip.cpp
index 9c58d9267..9ccc011da 100644
--- a/src/basic/MAT_MAT_SHARED-Hip.cpp
+++ b/src/basic/MAT_MAT_SHARED-Hip.cpp
@@ -86,7 +86,7 @@ void MAT_MAT_SHARED::runHipVariantImpl(VariantID vid)
     for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
 
       auto mat_mat_shared_lambda = [=] __device__() {
-
+      for (RepIndex_type extra_rep = 0; extra_rep < extra_kernel_reps; ++extra_rep) {
         auto outer_y = [&](Index_type by) {
           auto outer_x = [&](Index_type bx) {
             MAT_MAT_SHARED_BODY_0(tile_size)
@@ -173,6 +173,7 @@ void MAT_MAT_SHARED::runHipVariantImpl(VariantID vid)
           Index_type by = blockIdx.y;
           if(by < Ny) outer_y(by);
         }
+      }
       };
 
       RPlaunchHipKernel( (lambda_hip<tile_size*tile_size,
diff --git a/src/basic/MAT_MAT_SHARED.hpp b/src/basic/MAT_MAT_SHARED.hpp
index db448860b..85f17b270 100644
--- a/src/basic/MAT_MAT_SHARED.hpp
+++ b/src/basic/MAT_MAT_SHARED.hpp
@@ -75,6 +75,7 @@ constexpr rajaperf::Index_type TL_SZ = 16;
   Real_ptr A = m_A;                                                            \
   Real_ptr B = m_B;                                                            \
   Real_ptr C = m_C;
+  // constexpr int extra_kernel_reps = 5;
 
 /*
  When doing the device compile pass hipcc/clang will put in the device

From d58cd3ed80a5a6c0ed0fd98b302c087499d4fce2 Mon Sep 17 00:00:00 2001
From: Amr Akmal Moustafa Abouelmagd <abouelmagd1@dane4.llnl.gov>
Date: Mon, 25 Aug 2025 08:31:29 -0700
Subject: [PATCH 3/3] Increased outer loop to be 20 iterations

---
 src/stream/TRIAD.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/stream/TRIAD.hpp b/src/stream/TRIAD.hpp
index 8f31380f1..c555db7bb 100644
--- a/src/stream/TRIAD.hpp
+++ b/src/stream/TRIAD.hpp
@@ -26,7 +26,7 @@
 #define TRIAD_BODY  \
   a[i] = b[i] + alpha * c[i] ;
 
-constexpr int extra_kernel_reps = 5;
+constexpr int extra_kernel_reps = 20;
 
 #include "common/KernelBase.hpp"