TESTING only: enforce larger time span for kernel execution

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,

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

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 {

src/stream/TRIAD.hpp

@@ -26,6 +26,7 @@
 #define TRIAD_BODY  \
   a[i] = b[i] + alpha * c[i] ;
 
+constexpr int extra_kernel_reps = 20;
 
 #include "common/KernelBase.hpp"