SWDEV-435054 - Improves robustness of HIP test

Makes the test stronger by hiding
dispatch scheduling latency on MI300

Change-Id: Ic9724f4f1b16f1e707060129327248bbb353df45

catch/unit/cooperativeGrps/cg_common_kernels.hh

@@ -0,0 +1,43 @@
+/*
+Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
+
+static __device__ void busy_wait(unsigned long long wait_period) {
+  unsigned long long time_diff = 0;
+#if HT_AMD
+  unsigned long long last_clock = wall_clock64();
+#else
+  unsigned long long last_clock = clock64();
+#endif
+  while (time_diff < wait_period) {
+#if HT_AMD
+  unsigned long long cur_clock = wall_clock64();
+#else
+  unsigned long long cur_clock = clock64();
+#endif
+    if (cur_clock > last_clock) {
+      time_diff += (cur_clock - last_clock);
+    }
+    last_clock = cur_clock;
+  }
+}

catch/unit/cooperativeGrps/coalesced_group.cc

@@ -17,6 +17,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 #include "cooperative_groups_common.hh"
+#include "cg_common_kernels.hh"
 
 #include <cmd_options.hh>
 #include <cpu_grid.h>

catch/unit/cooperativeGrps/coalesced_group_tiled_partition.cc

@@ -18,6 +18,7 @@ THE SOFTWARE.
 */
 
 #include "cooperative_groups_common.hh"
+#include "cg_common_kernels.hh"
 
 #include <bitset>
 #include <optional>

catch/unit/cooperativeGrps/cooperative_groups_common.hh

@@ -43,19 +43,6 @@ __device__ inline unsigned int thread_rank_in_grid() {
   return block_rank_in_grid * block_size + thread_rank_in_block;
 }
 
-static __device__ void busy_wait(unsigned long long wait_period) {
-  unsigned long long time_diff = 0;
-  unsigned long long last_clock = clock64();
-  while (time_diff < wait_period) {
-    unsigned long long cur_clock = clock64();
-    if (cur_clock > last_clock) {
-      time_diff += (cur_clock - last_clock);
-    }
-    last_clock = cur_clock;
-  }
-}
-
-
 template <class T> bool CheckDimensions(unsigned int device, T kernel, dim3 blocks, dim3 threads) {
   hipDeviceProp_t props;
   int max_blocks_per_sm = 0;
@@ -74,5 +61,4 @@ template <class T> bool CheckDimensions(unsigned int device, T kernel, dim3 bloc
   }
 
   return true;
-}
-
+}

catch/unit/cooperativeGrps/grid_group.cc

@@ -51,26 +51,45 @@ static __global__ void grid_group_non_member_thread_rank_getter(unsigned int* th
   thread_ranks[thread_rank_in_grid()] = cg::thread_rank(cg::this_grid());
 }
 
-static __global__ void sync_kernel(unsigned int* atomic_val, unsigned int* array,
-                                   unsigned int loops) {
+static __global__ void sync_kernel(unsigned int* atomic_val, unsigned int *per_loop_atomic,
+                                   unsigned int* array, unsigned int loops) {
   cg::grid_group grid = cg::this_grid();
   unsigned rank = grid.thread_rank();
 
-  int offset = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x;
-  for (int i = 0; i < loops; i++) {
+  int blocks_seen = 0;
+  int grid_blocks = gridDim.x * gridDim.y * gridDim.z;
+
+  for (int iter = 0; iter < loops; iter++) {
     // Make the last thread run way behind everyone else.
     // If the sync below fails, then the other threads may hit the
     // atomicInc instruction many times before the last thread ever gets to it.
-    // If the sync works, then it will likely contain "total number of blocks"*i
+    // If the sync works, then it will likely contain "total number of blocks"*iter
     if (rank == (grid.size() - 1)) {
-      busy_wait(100000);
+      // The last wavefront should spin on this loop's atomic value
+      // until all of the other wavefronts have incremented the
+      // per-loop atomic and hit the grid.sync()
+#if HT_AMD
+      while(__hip_atomic_load(&per_loop_atomic[iter], __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT) <
+            (grid_blocks - 1)) {
+        __builtin_amdgcn_s_sleep(127);
+      }
+
+      // Give the other waves time to maybe go around the loop again
+      // if the barrier has failed
+      __builtin_amdgcn_s_sleep(127);
+#else // CUDA does not seem to need an ordered atomic load
+      while(per_loop_atomic[iter] < (grid_blocks - 1)) {
+      }
+#endif
     }
     if (threadIdx.x == blockDim.x - 1 && threadIdx.y == blockDim.y - 1 &&
         threadIdx.z == blockDim.z - 1) {
-      array[offset] = atomicInc(&atomic_val[0], UINT_MAX);
+      atomicInc(&per_loop_atomic[iter], UINT_MAX);
+      array[((blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x) + blocks_seen] =
+            atomicInc(&atomic_val[0], UINT_MAX);
     }
     grid.sync();
-    offset += gridDim.x * gridDim.y * gridDim.z;
+    blocks_seen += grid_blocks;
   }
 }
 
@@ -259,15 +278,19 @@ TEST_CASE("Unit_Grid_Group_Sync_Positive_Basic") {
   LinearAllocGuard<unsigned int> uint_arr(LinearAllocs::hipHostMalloc,
                                           array_len * sizeof(unsigned int));
   LinearAllocGuard<unsigned int> atomic_val(LinearAllocs::hipMalloc, sizeof(unsigned int));
+  LinearAllocGuard<unsigned int> per_loop_atomic_val(LinearAllocs::hipMalloc, loops * sizeof(unsigned int));
   HIP_CHECK(hipMemset(atomic_val.ptr(), 0, sizeof(unsigned int)));
+  HIP_CHECK(hipMemset(per_loop_atomic_val.ptr(), 0, loops * sizeof(unsigned int)));
 
   // Launch Kernel
   unsigned int* uint_arr_dev_ptr = uint_arr_dev.ptr();
   unsigned int* atomic_val_ptr = atomic_val.ptr();
-  void* params[3];
+  unsigned int* per_loop_atomic_val_ptr = per_loop_atomic_val.ptr();
+  void* params[4];
   params[0] = reinterpret_cast<void*>(&atomic_val_ptr);
-  params[1] = reinterpret_cast<void*>(&uint_arr_dev_ptr);
-  params[2] = reinterpret_cast<void*>(&loops);
+  params[1] = reinterpret_cast<void*>(&per_loop_atomic_val_ptr);
+  params[2] = reinterpret_cast<void*>(&uint_arr_dev_ptr);
+  params[3] = reinterpret_cast<void*>(&loops);
 
   HIP_CHECK(hipLaunchCooperativeKernel(sync_kernel, blocks, threads, params, 0, 0));
 

catch/unit/cooperativeGrps/multi_grid_group.cc

@@ -17,6 +17,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 #include "cooperative_groups_common.hh"
+#include "cg_common_kernels.hh"
 
 #include <cpu_grid.h>
 #include <resource_guards.hh>

catch/unit/cooperativeGrps/thread_block.cc

@@ -18,6 +18,7 @@ THE SOFTWARE.
 */
 
 #include "cooperative_groups_common.hh"
+#include "cg_common_kernels.hh"
 
 #include <cpu_grid.h>
 #include <optional>

catch/unit/cooperativeGrps/thread_block_tile.cc

@@ -18,6 +18,7 @@ THE SOFTWARE.
 */
 
 #include "cooperative_groups_common.hh"
+#include "cg_common_kernels.hh"
 
 #include <bitset>
 #include <array>