From 3c22a826a6ebec4b59c31634e440bed05d3df006 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Tue, 1 Jul 2025 14:21:39 +0000 Subject: [PATCH 1/7] SWDEV-493792 - add split barriers for grid_group --- .../hip/amd_detail/amd_hip_cooperative_groups.h | 13 +++++++++++++ .../include/hip/amd_detail/device_library_decls.h | 2 ++ .../hip/amd_detail/hip_cooperative_groups_helper.h | 4 +++- 3 files changed, 18 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index 62cb66fb908b..e5c268830c23 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h @@ -212,6 +212,19 @@ class grid_group : public thread_group { //! @copydoc thread_group::sync __CG_QUALIFIER__ void sync() const { internal::grid::sync(); } __CG_QUALIFIER__ dim3 group_dim() const { return internal::grid::grid_dim(); } + struct arrival_token { + unsigned int signal; + }; + //! Arrive at a barrier + __CG_QUALIFIER__ arrival_token barrier_arrive() const { + arrival_token t; + t.signal = internal::grid::barrier_signal(); + return t; + } + //! Arrive at a barrier + __CG_QUALIFIER__ void barrier_wait(arrival_token&& t) const { + internal::grid::barrier_wait(t.signal); + } }; /** \ingroup CooperativeGConstruct diff --git a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h index 39c2e5968639..791d2e7137d7 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h @@ -97,6 +97,8 @@ extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_size(void); extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank(void); extern "C" __device__ __attribute__((const)) int __ockl_multi_grid_is_valid(void); extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(void); +extern "C" __device__ __attribute__((const)) uint __ockl_grid_bar_arrive(void); +extern "C" __device__ __attribute__((convergent)) void __ockl_grid_bar_wait(uint); extern "C" __device__ void __ockl_atomic_add_noret_f32(float*, float); diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index d9671eb34296..18c23491c3ed 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -196,6 +196,9 @@ __CG_STATIC_QUALIFIER__ dim3 grid_dim() { static_cast<__hip_uint32_t>(gridDim.z))); } +__CG_STATIC_QUALIFIER__ unsigned int barrier_signal() { return __ockl_grid_bar_arrive(); } + +__CG_STATIC_QUALIFIER__ void barrier_wait(unsigned int s) { __ockl_grid_bar_wait(s); } } // namespace grid /** @@ -237,7 +240,6 @@ __CG_STATIC_QUALIFIER__ dim3 block_dim() { return (dim3(static_cast<__hip_uint32_t>(blockDim.x), static_cast<__hip_uint32_t>(blockDim.y), static_cast<__hip_uint32_t>(blockDim.z))); } - } // namespace workgroup namespace tiled_group { From d8dae8e4f88b55a51a59355c64cf1b4136c1101f Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Thu, 11 Sep 2025 23:44:43 +0100 Subject: [PATCH 2/7] add tests --- .../hip_cooperative_groups_helper.h | 9 ++ .../catch/unit/cooperativeGrps/CMakeLists.txt | 1 + .../unit/cooperativeGrps/split_barrier.cc | 114 ++++++++++++++++++ 3 files changed, 124 insertions(+) create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index 18c23491c3ed..325f2ea0861c 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -240,6 +240,15 @@ __CG_STATIC_QUALIFIER__ dim3 block_dim() { return (dim3(static_cast<__hip_uint32_t>(blockDim.x), static_cast<__hip_uint32_t>(blockDim.y), static_cast<__hip_uint32_t>(blockDim.z))); } + +__CG_STATIC_QUALIFIER__ void barrier_arrive() { + __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup"); +} + +__CG_STATIC_QUALIFIER__ void barrier_wait() { + __builtin_amdgcn_s_barrier(); + __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup"); +} } // namespace workgroup namespace tiled_group { diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt index f22d40f94487..c0d2a00e3c50 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt @@ -20,6 +20,7 @@ set(TEST_SRC binary_partition.cc cg_ballot.cc cg_any_all.cc + split_barrier.cc ) if(HIP_PLATFORM STREQUAL "nvidia") diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc b/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc new file mode 100644 index 000000000000..ca732e21d124 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc @@ -0,0 +1,114 @@ +/* +Copyright (c) 2025 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. +*/ + +#include +#include + +static __global__ void wg_split_barrier(float* out, float* in) { + namespace cg = cooperative_groups; + + __shared__ float mid[32]; + size_t i = threadIdx.x; + auto tb = cg::this_thread_block(); + + out[i] = in[i] * 2.0f; + + auto tok = tb.barrier_arrive(); + + // use tid 0 to populate shared mem + if (i == 0) { + for (size_t j = 0; j < 32; j++) { + mid[j] = in[j]; + } + } + + tb.barrier_wait(std::move(tok)); + + out[i] += mid[i]; +} + +TEST_CASE("Unit_coop_thread_block_split_barrier") { + constexpr size_t size = 32; + float *d_out, *d_in; + + HIP_CHECK(hipMalloc(&d_out, sizeof(float) * size)); + HIP_CHECK(hipMalloc(&d_in, sizeof(float) * size)); + + std::vector in(size, 0.0f), out = in; + for (size_t i = 0; i < size; i++) { + in[i] = i + 1; + } + + HIP_CHECK(hipMemset(d_out, 0, sizeof(float) * size)); + HIP_CHECK(hipMemcpy(d_in, in.data(), sizeof(float) * size, hipMemcpyHostToDevice)); + wg_split_barrier<<<1, size>>>(d_out, d_in); + HIP_CHECK(hipMemcpy(out.data(), d_out, sizeof(float) * size, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipFree(d_out)); + HIP_CHECK(hipFree(d_in)); + + for (size_t i = 0; i < size; i++) { + INFO("Index: " << i << " in: " << in[i] << " out: " << out[i]); + REQUIRE((in[i] * 3.0f) == Approx(out[i])); + } +} + +static __global__ void grid_split_barrier(int* data, int* result, int N) { + namespace cg = cooperative_groups; + cg::grid_group grid = cg::this_grid(); + + int gid = blockIdx.x * blockDim.x + threadIdx.x; + auto tok = grid.barrier_arrive(); + if (gid < N) { + data[gid] = gid + 1; + } + + grid.barrier_wait(std::move(tok)); + + if (grid.thread_rank() == 0) { + int sum = 0; + for (int i = 0; i < N; i++) sum += data[i]; + *result = sum; + } +} + +TEST_CASE("Unit_coop_grids_split_barrier") { + int N = 1024; + const int threads = 128; + const int blocks = (N + threads - 1) / threads; + + int *d_in, *d_out; + HIP_CHECK(hipMalloc(&d_in, N * sizeof(int))); + HIP_CHECK(hipMalloc(&d_out, sizeof(int))); + + void* args[] = {&d_in, &d_out, &N}; + + dim3 grid(blocks); + dim3 block(threads); + + HIP_CHECK(hipLaunchCooperativeKernel((void*)grid_split_barrier, grid, block, args, 0, 0)); + HIP_CHECK(hipDeviceSynchronize()); + + int out = 0; + HIP_CHECK(hipMemcpy(&out, d_out, sizeof(int), hipMemcpyDeviceToHost)); + + HIP_CHECK(hipFree(d_in)); + HIP_CHECK(hipFree(d_out)); + REQUIRE(out == ((N * (N + 1)) / 2)); +} From 601fec226057e5a04186d2096f5bbb2e9cf2d2ce Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Thu, 11 Sep 2025 23:56:49 +0100 Subject: [PATCH 3/7] Update change log --- projects/clr/CHANGELOG.md | 1 + .../include/hip/amd_detail/amd_hip_cooperative_groups.h | 8 ++++++++ 2 files changed, 9 insertions(+) diff --git a/projects/clr/CHANGELOG.md b/projects/clr/CHANGELOG.md index 292125904dad..80756af55e4e 100644 --- a/projects/clr/CHANGELOG.md +++ b/projects/clr/CHANGELOG.md @@ -8,6 +8,7 @@ Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs * New HIP APIs - `hipKernelGetParamInfo` returns the offset and size of a kernel parameter +* Support for `barrier_arrive` and `barrier_wait` for `grid_group` and `thread_block`. * New HIP supports - `grid_group::block_rank()` returns the rank of the block in the calling thread diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index e5c268830c23..80b7f71c2378 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h @@ -308,6 +308,14 @@ class thread_block : public thread_group { __CG_STATIC_QUALIFIER__ void sync() { internal::workgroup::sync(); } //! Returns the group dimensions. __CG_QUALIFIER__ dim3 group_dim() { return internal::workgroup::block_dim(); } + struct arrival_token {}; + //! Arrive at a barrier + __CG_QUALIFIER__ arrival_token barrier_arrive() const { + internal::workgroup::barrier_arrive(); + return arrival_token{}; + } + //! Arrive at a barrier + __CG_QUALIFIER__ void barrier_wait(arrival_token&&) const { internal::workgroup::barrier_wait(); } }; /** \ingroup CooperativeGConstruct From 9f12fc0fdf5602a1ad023ca335061e4433964b86 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 15 Sep 2025 16:51:12 +0100 Subject: [PATCH 4/7] Add Navi4 split barrier --- .../hip/amd_detail/hip_cooperative_groups_helper.h | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index 325f2ea0861c..771e8419ffd6 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -243,10 +243,19 @@ __CG_STATIC_QUALIFIER__ dim3 block_dim() { __CG_STATIC_QUALIFIER__ void barrier_arrive() { __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup"); +#if __has_builtin(__builtin_amdgcn_s_barrier_signal) && \ + __has_builtin(__builtin_amdgcn_s_barrier_wait) + __builtin_amdgcn_s_barrier_signal(-1); +#endif // __builtin_amdgcn_s_barrier_signal && __builtin_amdgcn_s_barrier_wait } __CG_STATIC_QUALIFIER__ void barrier_wait() { +#if __has_builtin(__builtin_amdgcn_s_barrier_signal) && \ + __has_builtin(__builtin_amdgcn_s_barrier_wait) + __builtin_amdgcn_s_barrier_wait(-1); +#else __builtin_amdgcn_s_barrier(); +#endif // __builtin_amdgcn_s_barrier_signal && __builtin_amdgcn_s_barrier_wait __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup"); } } // namespace workgroup From 01e4c60399b69c6dd717e0680e334013b22ea566 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 24 Sep 2025 10:32:24 +0100 Subject: [PATCH 5/7] Update docs --- projects/hip/docs/how-to/hip_runtime_api/cooperative_groups.rst | 1 - 1 file changed, 1 deletion(-) diff --git a/projects/hip/docs/how-to/hip_runtime_api/cooperative_groups.rst b/projects/hip/docs/how-to/hip_runtime_api/cooperative_groups.rst index a3e32cd294a0..bf34ab7583fe 100644 --- a/projects/hip/docs/how-to/hip_runtime_api/cooperative_groups.rst +++ b/projects/hip/docs/how-to/hip_runtime_api/cooperative_groups.rst @@ -494,7 +494,6 @@ HIP doesn't support the following CUDA functions/operators in ``cooperative_grou * ``synchronize`` * ``memcpy_async`` * ``wait`` and ``wait_prior`` -* ``barrier_arrive`` and ``barrier_wait`` * ``invoke_one`` and ``invoke_one_broadcast`` * ``reduce`` * ``reduce_update_async`` and ``reduce_store_async`` From e6e013e0f294a7349576558dd6d7fa90d6a59ad3 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 5 Dec 2025 16:12:53 +0000 Subject: [PATCH 6/7] Use new Catch2 Approx macro --- projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc b/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc index ca732e21d124..4a2b91a2fc60 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc @@ -65,7 +65,7 @@ TEST_CASE("Unit_coop_thread_block_split_barrier") { for (size_t i = 0; i < size; i++) { INFO("Index: " << i << " in: " << in[i] << " out: " << out[i]); - REQUIRE((in[i] * 3.0f) == Approx(out[i])); + REQUIRE((in[i] * 3.0f) == Catch::Approx(out[i])); } } From 3657f6a81ad06e728a66832f3ca2e1a16b2cf388 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com> Date: Mon, 12 Jan 2026 16:59:03 +0000 Subject: [PATCH 7/7] Update split_barrier.cc to check for coop groups --- .../unit/cooperativeGrps/split_barrier.cc | 53 +++++++++++-------- 1 file changed, 31 insertions(+), 22 deletions(-) diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc b/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc index 4a2b91a2fc60..f7376b6605c6 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/split_barrier.cc @@ -17,10 +17,10 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include #include +#include -static __global__ void wg_split_barrier(float* out, float* in) { +static __global__ void wg_split_barrier(float *out, float *in) { namespace cg = cooperative_groups; __shared__ float mid[32]; @@ -56,9 +56,11 @@ TEST_CASE("Unit_coop_thread_block_split_barrier") { } HIP_CHECK(hipMemset(d_out, 0, sizeof(float) * size)); - HIP_CHECK(hipMemcpy(d_in, in.data(), sizeof(float) * size, hipMemcpyHostToDevice)); + HIP_CHECK( + hipMemcpy(d_in, in.data(), sizeof(float) * size, hipMemcpyHostToDevice)); wg_split_barrier<<<1, size>>>(d_out, d_in); - HIP_CHECK(hipMemcpy(out.data(), d_out, sizeof(float) * size, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(out.data(), d_out, sizeof(float) * size, + hipMemcpyDeviceToHost)); HIP_CHECK(hipFree(d_out)); HIP_CHECK(hipFree(d_in)); @@ -69,7 +71,7 @@ TEST_CASE("Unit_coop_thread_block_split_barrier") { } } -static __global__ void grid_split_barrier(int* data, int* result, int N) { +static __global__ void grid_split_barrier(int *data, int *result, int N) { namespace cg = cooperative_groups; cg::grid_group grid = cg::this_grid(); @@ -83,32 +85,39 @@ static __global__ void grid_split_barrier(int* data, int* result, int N) { if (grid.thread_rank() == 0) { int sum = 0; - for (int i = 0; i < N; i++) sum += data[i]; + for (int i = 0; i < N; i++) + sum += data[i]; *result = sum; } } TEST_CASE("Unit_coop_grids_split_barrier") { - int N = 1024; - const int threads = 128; - const int blocks = (N + threads - 1) / threads; + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); - int *d_in, *d_out; - HIP_CHECK(hipMalloc(&d_in, N * sizeof(int))); - HIP_CHECK(hipMalloc(&d_out, sizeof(int))); + if (prop.cooperativeLaunch != 0) { + int N = 1024; + const int threads = 128; + const int blocks = (N + threads - 1) / threads; - void* args[] = {&d_in, &d_out, &N}; + int *d_in, *d_out; + HIP_CHECK(hipMalloc(&d_in, N * sizeof(int))); + HIP_CHECK(hipMalloc(&d_out, sizeof(int))); - dim3 grid(blocks); - dim3 block(threads); + void *args[] = {&d_in, &d_out, &N}; - HIP_CHECK(hipLaunchCooperativeKernel((void*)grid_split_barrier, grid, block, args, 0, 0)); - HIP_CHECK(hipDeviceSynchronize()); + dim3 grid(blocks); + dim3 block(threads); - int out = 0; - HIP_CHECK(hipMemcpy(&out, d_out, sizeof(int), hipMemcpyDeviceToHost)); + HIP_CHECK(hipLaunchCooperativeKernel((void *)grid_split_barrier, grid, + block, args, 0, 0)); + HIP_CHECK(hipDeviceSynchronize()); - HIP_CHECK(hipFree(d_in)); - HIP_CHECK(hipFree(d_out)); - REQUIRE(out == ((N * (N + 1)) / 2)); + int out = 0; + HIP_CHECK(hipMemcpy(&out, d_out, sizeof(int), hipMemcpyDeviceToHost)); + + HIP_CHECK(hipFree(d_in)); + HIP_CHECK(hipFree(d_out)); + REQUIRE(out == ((N * (N + 1)) / 2)); + } }