diff --git a/catch/unit/cooperativeGrps/cg_common_kernels.hh b/catch/unit/cooperativeGrps/cg_common_kernels.hh new file mode 100644 index 0000000000..bcfddf68e3 --- /dev/null +++ b/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 +#include + +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; + } +} diff --git a/catch/unit/cooperativeGrps/coalesced_group.cc b/catch/unit/cooperativeGrps/coalesced_group.cc index 4ff7e7d510..b24d0b0aa4 100644 --- a/catch/unit/cooperativeGrps/coalesced_group.cc +++ b/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 #include diff --git a/catch/unit/cooperativeGrps/coalesced_group_tiled_partition.cc b/catch/unit/cooperativeGrps/coalesced_group_tiled_partition.cc index c354bbe3a9..2353d20470 100644 --- a/catch/unit/cooperativeGrps/coalesced_group_tiled_partition.cc +++ b/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 #include diff --git a/catch/unit/cooperativeGrps/cooperative_groups_common.hh b/catch/unit/cooperativeGrps/cooperative_groups_common.hh index aeec9e5fcc..2eb1314d4f 100644 --- a/catch/unit/cooperativeGrps/cooperative_groups_common.hh +++ b/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 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 bool CheckDimensions(unsigned int device, T kernel, dim3 bloc } return true; -} - +} \ No newline at end of file diff --git a/catch/unit/cooperativeGrps/grid_group.cc b/catch/unit/cooperativeGrps/grid_group.cc index eb97359456..11da111469 100644 --- a/catch/unit/cooperativeGrps/grid_group.cc +++ b/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 uint_arr(LinearAllocs::hipHostMalloc, array_len * sizeof(unsigned int)); LinearAllocGuard atomic_val(LinearAllocs::hipMalloc, sizeof(unsigned int)); + LinearAllocGuard 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(&atomic_val_ptr); - params[1] = reinterpret_cast(&uint_arr_dev_ptr); - params[2] = reinterpret_cast(&loops); + params[1] = reinterpret_cast(&per_loop_atomic_val_ptr); + params[2] = reinterpret_cast(&uint_arr_dev_ptr); + params[3] = reinterpret_cast(&loops); HIP_CHECK(hipLaunchCooperativeKernel(sync_kernel, blocks, threads, params, 0, 0)); diff --git a/catch/unit/cooperativeGrps/multi_grid_group.cc b/catch/unit/cooperativeGrps/multi_grid_group.cc index a7490bb360..0a4a767a4e 100644 --- a/catch/unit/cooperativeGrps/multi_grid_group.cc +++ b/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 #include diff --git a/catch/unit/cooperativeGrps/thread_block.cc b/catch/unit/cooperativeGrps/thread_block.cc index a5018a9cec..d7d3a0dde9 100644 --- a/catch/unit/cooperativeGrps/thread_block.cc +++ b/catch/unit/cooperativeGrps/thread_block.cc @@ -18,6 +18,7 @@ THE SOFTWARE. */ #include "cooperative_groups_common.hh" +#include "cg_common_kernels.hh" #include #include diff --git a/catch/unit/cooperativeGrps/thread_block_tile.cc b/catch/unit/cooperativeGrps/thread_block_tile.cc index 8ae266b0e4..a3ac01aaf4 100644 --- a/catch/unit/cooperativeGrps/thread_block_tile.cc +++ b/catch/unit/cooperativeGrps/thread_block_tile.cc @@ -18,6 +18,7 @@ THE SOFTWARE. */ #include "cooperative_groups_common.hh" +#include "cg_common_kernels.hh" #include #include