diff --git a/HIP-Basic/CMakeLists.txt b/HIP-Basic/CMakeLists.txt index 58f03dee5..ed29ae528 100644 --- a/HIP-Basic/CMakeLists.txt +++ b/HIP-Basic/CMakeLists.txt @@ -93,6 +93,11 @@ endif() if(NOT CMAKE_SYSTEM_NAME MATCHES "Windows") add_subdirectory(cooperative_groups) + # This cooperative groups example uses AMD-specific extensions + # (cooperative_groups::memcpy_async) that are not available on the CUDA backend. + if(NOT "${ROCM_EXAMPLES_GPU_LANGUAGE}" STREQUAL "CUDA") + add_subdirectory(cooperative_groups_double_buffered_tile) + endif() endif() add_subdirectory(bandwidth) diff --git a/HIP-Basic/Makefile b/HIP-Basic/Makefile index 42fd31d0b..03e14bbd1 100644 --- a/HIP-Basic/Makefile +++ b/HIP-Basic/Makefile @@ -57,6 +57,7 @@ endif ifneq ($(GPU_RUNTIME), CUDA) EXAMPLES += \ assembly_to_executable \ + cooperative_groups_double_buffered_tile \ llvm_ir_to_executable \ module_api \ static_device_library diff --git a/HIP-Basic/cooperative_groups_double_buffered_tile/.gitignore b/HIP-Basic/cooperative_groups_double_buffered_tile/.gitignore new file mode 100644 index 000000000..70992bcd3 --- /dev/null +++ b/HIP-Basic/cooperative_groups_double_buffered_tile/.gitignore @@ -0,0 +1 @@ +hip_cooperative_groups_double_buffered_tile diff --git a/HIP-Basic/cooperative_groups_double_buffered_tile/CMakeLists.txt b/HIP-Basic/cooperative_groups_double_buffered_tile/CMakeLists.txt new file mode 100644 index 000000000..6152f9dae --- /dev/null +++ b/HIP-Basic/cooperative_groups_double_buffered_tile/CMakeLists.txt @@ -0,0 +1,60 @@ +# MIT License +# +# Copyright (c) 2026 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. + +set(example_name hip_cooperative_groups_double_buffered_tile) + +cmake_minimum_required(VERSION 3.21 FATAL_ERROR) +project(${example_name} LANGUAGES CXX) + +include("${CMAKE_CURRENT_LIST_DIR}/../../Common/HipPlatform.cmake") +select_gpu_language() +enable_language(${ROCM_EXAMPLES_GPU_LANGUAGE}) +select_hip_platform() + +enable_language(${ROCM_EXAMPLES_GPU_LANGUAGE}) +set(CMAKE_${ROCM_EXAMPLES_GPU_LANGUAGE}_STANDARD 17) +set(CMAKE_${ROCM_EXAMPLES_GPU_LANGUAGE}_EXTENSIONS OFF) +set(CMAKE_${ROCM_EXAMPLES_GPU_LANGUAGE}_STANDARD_REQUIRED ON) + +include("${CMAKE_CURRENT_LIST_DIR}/../../Common/ROCmPath.cmake") + +add_executable(${example_name} main.hip) +# Make example runnable using ctest +add_test(NAME ${example_name} COMMAND ${example_name}) + +set(include_dirs "../../Common" "../../External") +if(ROCM_EXAMPLES_GPU_LANGUAGE STREQUAL "CUDA") + list(APPEND include_dirs "${ROCM_PATH}/include") +else() + # Add NDEBUG for HIP version >= 5.5 and < 6.0 due to a known bug in the cooperative groups header + if( + ${hip-lang_VERSION} VERSION_GREATER_EQUAL 5.5 + AND ${hip-lang_VERSION} VERSION_LESS 6 + ) + add_compile_definitions(NDEBUG) + endif() +endif() + +target_include_directories(${example_name} PRIVATE ${include_dirs}) +set_source_files_properties(main.hip PROPERTIES LANGUAGE ${ROCM_EXAMPLES_GPU_LANGUAGE}) + +install(TARGETS ${example_name}) diff --git a/HIP-Basic/cooperative_groups_double_buffered_tile/Makefile b/HIP-Basic/cooperative_groups_double_buffered_tile/Makefile new file mode 100644 index 000000000..4d75a4113 --- /dev/null +++ b/HIP-Basic/cooperative_groups_double_buffered_tile/Makefile @@ -0,0 +1,63 @@ +# MIT License +# +# Copyright (c) 2026 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. + +EXAMPLE := hip_cooperative_groups_double_buffered_tile +COMMON_INCLUDE_DIR := ../../Common +GPU_RUNTIME ?= HIP + +# HIP variables +ROCM_PATH ?= /opt/rocm +HIP_INCLUDE_DIR := $(ROCM_PATH)/include + +HIPCXX ?= $(ROCM_PATH)/bin/hipcc + +# Common variables and flags +CXX_STD := c++17 +ICXXFLAGS := -std=$(CXX_STD) +ICPPFLAGS := -I $(COMMON_INCLUDE_DIR) +ILDFLAGS := +ILDLIBS := + +ifeq ($(GPU_RUNTIME), CUDA) + ICXXFLAGS += -x cu + ICPPFLAGS += -isystem $(HIP_INCLUDE_DIR) +else ifeq ($(GPU_RUNTIME), HIP) + CXXFLAGS ?= -Wall -Wextra +else + $(error GPU_RUNTIME is set to "$(GPU_RUNTIME)". GPU_RUNTIME must be either CUDA or HIP) +endif + +ICXXFLAGS += $(CXXFLAGS) +ICPPFLAGS += $(CPPFLAGS) +ILDFLAGS += $(LDFLAGS) +ILDLIBS += $(LDLIBS) + +$(EXAMPLE): main.hip $(COMMON_INCLUDE_DIR)/example_utils.hpp + $(HIPCXX) $(ICXXFLAGS) $(ICPPFLAGS) $(ILDFLAGS) -o $@ $< $(ILDLIBS) + +test: $(EXAMPLE) + ./$(EXAMPLE) $(TEST_ARGS) + +clean: + $(RM) $(EXAMPLE) + +.PHONY: clean test diff --git a/HIP-Basic/cooperative_groups_double_buffered_tile/README.md b/HIP-Basic/cooperative_groups_double_buffered_tile/README.md new file mode 100644 index 000000000..23f25161d --- /dev/null +++ b/HIP-Basic/cooperative_groups_double_buffered_tile/README.md @@ -0,0 +1,96 @@ +# Cooperative Groups Double-Buffered Tile Example + +## Description + +This program showcases a double-buffered tile load pipeline built from two cooperative groups +APIs: the group-collective `cooperative_groups::memcpy_async` and the split barrier +(`barrier_arrive` / `barrier_wait`) of a `thread_block`. + +A single block streams a 1D array through two LDS (shared memory) buffers, one tile at a time. +The async load of the next tile is issued into the other buffer while the current tile is consumed, +and a split barrier separates the moment a thread has finished reading a buffer from the moment the +block guarantees that every thread is done. The kernel applies the element-wise operation +`out[i] = scale * in[i] + bias`, which is trivial to validate against a CPU reference. + +`cooperative_groups::memcpy_async` is an **asynchronous**, group-collective copy (typically +global <-> LDS). HIP does not expose a separate wait handle (there is no `cg::wait()`), so its +completion must be enforced by a following group barrier - either a `block.sync()` (as the official +reference test does) or, as in this example, the `barrier_wait` of a split barrier whose +`barrier_arrive` is issued *after* the copy. Ordering matters: the prefetch of the next tile is +issued **before** `barrier_arrive`, so the release fence in `barrier_arrive` orders the copy's +completion and the acquire fence in `barrier_wait` makes the prefetched buffer visible to every +thread by the next iteration. The split barrier additionally lets the current tile's computation +run as independent work between `barrier_arrive` and `barrier_wait`. Correctness (no data races, +validated output) is the top priority. + +This example targets the AMD/HIP (ROCm) backend: the NVIDIA path of +`` is not yet implemented (the header carries a TODO for it), +and it requires a ROCm version recent enough to ship that public header. + +### Application flow + +1. A number of compile-time constants define the tile size (also the block size), the number of + tiles, the total element count, and the constants of the element-wise operation. +2. The input array is set up in host memory and the output array is allocated. +3. The input is copied to the device. +4. The double-buffered pipeline kernel is launched in a single block. + 1. The first tile is loaded into the first LDS buffer with `memcpy_async`, followed by a + block-wide `sync` that completes the async load and makes the tile visible to all threads. + 2. For each tile the block issues the async load of the next tile into the other buffer, calls + `barrier_arrive`, then - as independent work between arrive and wait - consumes the current + buffer (applies the element-wise operation and writes the result to global memory), and finally + calls `barrier_wait` to complete the barrier and the in-flight prefetch. +5. The result array is copied back to the host and all device memory is freed. +6. The elements of the result are compared with the CPU reference. The result of the comparison is + printed to the standard output. + +## Key APIs and Concepts + +- `cooperative_groups::this_thread_block` returns the `thread_block` group that represents all + threads of the block. The block is used both as the group for the collective copies and as the + group that owns the split barrier. +- `cooperative_groups::memcpy_async(group, dst, src, count_in_bytes)` is an asynchronous + group-collective copy that is designed for global <-> LDS transfers. Every thread of the group + must call it collectively. **In HIP the `count` argument is expressed in bytes** (here `tile_size * + sizeof(float)`). The copy is asynchronous and HIP exposes no separate wait handle, so its + completion is enforced by a following group barrier (`block.sync()` or the split barrier's + `barrier_wait`). On hardware or compilers without the asynchronous LDS builtins it falls back to a + correct traditional per-thread copy, so it always produces correct results. +- The split barrier decomposes a block barrier into two phases. `thread_block::barrier_arrive` + signals that a thread has reached the barrier and returns an `arrival_token`; it emits a release + fence and does not block, which exposes a window for independent work. `thread_block::barrier_wait` + consumes the moved token, blocks until every thread of the block has arrived, and emits an acquire + fence. Together they act as a full block barrier whose release/acquire fences order the prefetch + issued just before `barrier_arrive`. +- Two LDS buffers (`__shared__ float buf[2][tile_size]`) are alternated between iterations so that + the buffer being consumed is never the buffer being prefetched. +- Race-freedom: (a) the buffer prefetched during an iteration (`buf[(t + 1) & 1]`) is always + different from the buffer read as independent work (`buf[cur]`), so the load and the reads target + disjoint memory; (b) the prefetch is issued before `barrier_arrive`, so the barrier's + release/acquire fences order its completion and visibility before the next iteration consumes it; + (c) `buf[cur]` is only overwritten by the prefetch issued in the next iteration, which cannot begin + until every thread has passed this iteration's `barrier_wait` - i.e. after every thread has + finished reading `buf[cur]`. In this configuration the tile size equals the block size, so every + thread also copies exactly the element it later reads. + +## Demonstrated API Calls + +### HIP runtime + +#### Device symbols + +- `cooperative_groups::this_thread_block` +- `thread_block` +- `cooperative_groups::memcpy_async` +- `thread_block::barrier_arrive` +- `thread_block::barrier_wait` +- `thread_block::sync` +- All above from the [`cooperative_groups` namespace](https://github.com/ROCm/clr/blob/develop/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h) + +#### Host symbols + +- `hipMalloc` +- `hipMemcpy` +- `hipStreamDefault` +- `hipGetLastError` +- `hipFree` diff --git a/HIP-Basic/cooperative_groups_double_buffered_tile/main.hip b/HIP-Basic/cooperative_groups_double_buffered_tile/main.hip new file mode 100644 index 000000000..d20d1c525 --- /dev/null +++ b/HIP-Basic/cooperative_groups_double_buffered_tile/main.hip @@ -0,0 +1,218 @@ +// MIT License +// +// Copyright (c) 2026 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 "example_utils.hpp" + +#include +#include +#include + +#include +#include + +#include +#include +#include + +/// \brief Number of `float` elements in a single tile. +/// The tile size is also used as the block size, so every thread of the block +/// is responsible for exactly one element per tile. +constexpr unsigned int tile_size = 256; + +/// \brief Number of tiles streamed through the double-buffered pipeline. +constexpr unsigned int num_tiles = 8; + +/// \brief Total number of elements of the input and output arrays. +constexpr unsigned int size = num_tiles * tile_size; + +/// \brief Scale factor of the element-wise operation `out[i] = scale * in[i] + bias`. +constexpr float scale = 2.0f; + +/// \brief Bias term of the element-wise operation `out[i] = scale * in[i] + bias`. +constexpr float bias = 3.0f; + +/// \brief Double-buffered tile pipeline over a 1D array. +/// +/// The kernel streams `num_tiles` tiles of `tile_size` elements through two LDS +/// (shared memory) buffers. While one buffer is being consumed, the next tile is +/// prefetched into the other buffer. A split barrier (`barrier_arrive` / +/// `barrier_wait`) separates the point where a thread has finished reading a +/// buffer from the point where the block guarantees that every thread is done, +/// which is what makes reusing the two buffers race-free. +/// +/// Note: `cooperative_groups::memcpy_async` is an asynchronous, group-collective +/// copy (typically global <-> LDS). HIP does not expose a separate wait handle +/// (there is no `cg::wait()`), so its completion must be enforced by a following +/// group barrier. Here the copy of the next tile is issued *before* +/// `barrier_arrive`, so the release fence of `barrier_arrive` and the acquire +/// fence of `barrier_wait` order the copy's completion and make the prefetched +/// buffer visible to every thread by the next iteration. The split barrier also +/// lets the current tile's computation run as independent work between +/// `barrier_arrive` and `barrier_wait`. +__global__ void double_buffered_tile_kernel(const float* in, float* out) +{ + namespace cg = cooperative_groups; + + // The group used for the collective copies and the split barrier: every + // thread of the block participates. + cg::thread_block block = cg::this_thread_block(); + + // The two LDS buffers that are alternated between iterations. + __shared__ float buf[2][tile_size]; + + // Rank of this thread inside the block; also the element index within a tile. + const unsigned int tid = block.thread_rank(); + + // Load the first tile into buf[0]. `memcpy_async` is an asynchronous, + // collective call: every thread of the group must call it. Its completion is + // enforced here with a block-wide barrier before the tile is consumed. In + // this configuration each thread copies exactly the element it later reads. + cg::memcpy_async(block, buf[0], in + 0, tile_size * sizeof(float)); + + // Complete the async load and make the tile visible to every thread. + block.sync(); + + for(unsigned int t = 0; t < num_tiles; ++t) + { + // Index of the buffer holding the tile consumed during this iteration. + const unsigned int cur = t & 1u; + + // Kick off the async load of the next tile into the other buffer. It is + // issued *before* `barrier_arrive` so that the release fence in + // `barrier_arrive` orders its completion and the acquire fence in + // `barrier_wait` makes it visible. The target buf[(t + 1) & 1] is never + // the buffer being consumed (buf[cur]), so the load and the reads below + // target disjoint memory. + if(t + 1 < num_tiles) + { + cg::memcpy_async(block, + buf[(t + 1) & 1u], + in + (t + 1) * tile_size, + tile_size * sizeof(float)); + } + + // Split barrier, part 1: `barrier_arrive` publishes the prefetch (release + // fence) and returns a token without blocking, opening a window for + // independent work. + auto token = block.barrier_arrive(); + + // Independent work performed while the barrier is in progress: consume + // the current tile (already valid) and write this tile's results. This + // reads buf[cur] only, so it is independent of the in-flight prefetch + // into the other buffer. + const unsigned int global_index = t * tile_size + tid; + out[global_index] = scale * buf[cur][tid] + bias; + + // Split barrier, part 2: `barrier_wait` completes the barrier (block + // barrier + acquire fence). After it: + // - the prefetch into the other buffer is complete and visible, ready to + // be consumed by the next iteration, and + // - every thread has finished reading buf[cur], so it may safely be + // overwritten later. + // + // Race-freedom argument: + // - The prefetch targets the *other* buffer, buf[(t + 1) & 1], which is + // disjoint from buf[cur] read as independent work this iteration. + // - The prefetch is issued before `barrier_arrive`, so the barrier's + // release/acquire fences order its completion and visibility: when the + // next iteration consumes that buffer, the copy is guaranteed done. + // - buf[cur] is only overwritten by the prefetch issued in the *next* + // iteration (whose target buf[((t + 1) + 1) & 1] == buf[cur]). That + // prefetch cannot begin until every thread has passed this iteration's + // `barrier_wait`, i.e. after every thread finished reading buf[cur]. + block.barrier_wait(std::move(token)); + } +} + +/// \brief Host reference: applies the same element-wise operation over the array. +std::vector ref_transform(const std::vector& input) +{ + std::vector result(input.size()); + for(std::size_t i = 0; i < input.size(); ++i) + { + result[i] = scale * input[i] + bias; + } + return result; +} + +int main() +{ + // Total size (in bytes) of the input and output arrays. + constexpr std::size_t size_bytes = sizeof(float) * size; + + // The tile size is used as the block size, so every thread handles exactly + // one element per tile. + static_assert(size == num_tiles * tile_size, "size must equal num_tiles * tile_size"); + + // Set up the input data on the host. + std::vector h_in(size); + for(unsigned int i = 0; i < size; ++i) + { + h_in[i] = static_cast(i); + } + std::vector h_out(size); + + // Allocate device memory for the input and output arrays. + float* d_in{}; + float* d_out{}; + HIP_CHECK(hipMalloc(&d_in, size_bytes)); + HIP_CHECK(hipMalloc(&d_out, size_bytes)); + + // Transfer the input array to the device. + HIP_CHECK(hipMemcpy(d_in, h_in.data(), size_bytes, hipMemcpyHostToDevice)); + + // Launch the double-buffered pipeline in a single block. No cooperative + // launch is required because the block-level split barrier and the + // block-collective copies only synchronize threads within one block. + double_buffered_tile_kernel<<>>(d_in, d_out); + + // Check that the kernel launch was successful. + HIP_CHECK(hipGetLastError()); + + // Transfer the result back to the host. + HIP_CHECK(hipMemcpy(h_out.data(), d_out, size_bytes, hipMemcpyDeviceToHost)); + + // Free the resources on the device. + HIP_CHECK(hipFree(d_in)); + HIP_CHECK(hipFree(d_out)); + + // Perform the reference (CPU) calculation. + const std::vector ref = ref_transform(h_in); + + // Check the results' validity. The chosen constants and inputs make every + // result exactly representable as a `float`, so an exact comparison is used. + unsigned int errors{}; + for(unsigned int i = 0; i < size; ++i) + { + errors += (h_out[i] != ref[i]); + } + + if(errors) + { + std::cout << "Validation failed. Errors: " << errors << std::endl; + return error_exit_code; + } + else + { + std::cout << "Validation passed." << std::endl; + } +} diff --git a/README.md b/README.md index 606dd9fc2..55a9427bf 100644 --- a/README.md +++ b/README.md @@ -159,6 +159,8 @@ The following options are available when building with CMake. - [assembly_to_executable](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/assembly_to_executable): Program and accompanying build systems that show how to manually compile and link a HIP application from host and device code. - [bandwidth](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/bandwidth): Program that measures memory bandwidth from host to device, device to host, and device to device. - [bit_extract](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/bit_extract): Program that showcases how to use HIP built-in bit extract. + - [cooperative_groups](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/cooperative_groups): Showcases the cooperative groups programming model with a block- and partition-level reduction. + - [cooperative_groups_double_buffered_tile](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/cooperative_groups_double_buffered_tile): Double-buffered tile pipeline over a 1D array combining `memcpy_async` with a split barrier. - [device_globals](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/device_globals): Show cases how to set global variables on the device from the host. - [device_query](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/device_query): Program that showcases how properties from the device may be queried. - [dynamic_shared](https://github.com/ROCm/rocm-examples/tree/amd-staging/HIP-Basic/dynamic_shared): Program that showcases how to use dynamic shared memory with the help of a simple matrix transpose kernel.