[ROCm] Fix the experimental HIP backend for ROCm 7.x on AMD GPUs#484
[ROCm] Fix the experimental HIP backend for ROCm 7.x on AMD GPUs#484jeffdaily wants to merge 10 commits into
Conversation
This fixes the existing experimental HIP backend so it builds on ROCm 7.x and runs correctly across both AMD wavefront sizes: wave64 (CDNA) and wave32 (RDNA). It does not add a new backend; it repairs the one already present (STDGPU_BACKEND_HIP).
Three independent issues are addressed. Review them in order:
1. CMake HIP language propagation (build break). On ROCm 7.x, linking against rocThrust pulls in hip::device, whose INTERFACE_COMPILE_OPTIONS add HIP flags (-x hip, --offload-arch=...) gated on $<COMPILE_LANGUAGE:CXX>. Those flags reach the system C++ compiler building stdgpu's .cpp files, which rejects --offload-arch and fails. The fix marks the .cpp sources that include HIP headers as LANGUAGE HIP (shared impl, HIP backend impl, examples, tests) so the HIP compiler handles them. All edits are guarded on STDGPU_BACKEND_HIP; the CUDA and OpenMP builds are untouched.
2. wave64/wave32 spinlock livelock (hang in concurrent containers). The lock-free containers retry per-slot try_lock() in spin loops. AMD GPUs give no per-lane forward-progress guarantee: when several lanes of one wavefront contend for the same lock, the winning lane is stuck at SIMT reconvergence waiting on the losers, who spin forever. CUDA Volta+ Independent Thread Scheduling lets the winner proceed; AMD's execution model does not. A new helper, detail::wave_lock_serialize, elects one active lane at a time via __ballot and runs the spin-loop body for that lane, so each lane completes its critical section in turn. It is wave-size agnostic (the ballot mask covers whatever lanes are active, 32 or 64) and is a zero-cost pass-through on CUDA, where ITS already handles this. It wraps the insert/erase loops in unordered_base and the push/pop loops in deque and vector.
3. RDNA hang in valid() (degenerate single-slot reduction). Container valid() checks route through transform_reduce_index, a wrapper over thrust::transform_reduce. On a degenerate single-slot table that is full or has an emptied excess list, this lowers to a single-block rocPRIM reduce whose internal hipMemcpyWithStream copy of the result never completes on some RDNA GPUs (the launch returns success but the result copy wedges until the watchdog fires). For the device execution policy on the HIP backend only, this replaces the thrust reduce with a small hand-written kernel plus a plain hipMemcpy device-to-host and a host-side fold, computing the identical reduction (same init, same associative binary op). The host policy and the CUDA backend keep the thrust path unchanged.
The shared headers (wave_lock.h, numeric_detail.h) gate every AMD-specific path behind __HIP_PLATFORM_AMD__ / STDGPU_BACKEND_HIP, so CUDA and OpenMP behavior is unchanged.
This work was authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan
Built and ran the full test suite on three AMD GPUs spanning both wavefront sizes:
gfx90a (MI250X, CDNA2, wave64), ROCm 7.2.x, Linux:
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release \
-DSTDGPU_BACKEND=STDGPU_BACKEND_HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a
cmake --build build --parallel
./build/bin/teststdgpu
gfx1100 (RDNA3, wave32), ROCm, Linux:
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release \
-DSTDGPU_BACKEND=STDGPU_BACKEND_HIP -DCMAKE_HIP_ARCHITECTURES=gfx1100
cmake --build build --parallel
./build/bin/teststdgpu
gfx1201 (RX 9070 XT, RDNA4, wave32), ROCm 7.14, Windows:
HIP_VISIBLE_DEVICES=0 build-gfx1201/bin/teststdgpu.exe
All 702 tests pass on every arch, deterministic across repeated runs. The previously hanging concurrent-container tests (insert/erase/emplace while full or with an emptied excess list, and the simultaneous push/pop suites) now pass.
The CUDA backend was additionally compile-checked with nvcc (CUDA 13.3) to confirm these shared-header changes do not affect it: the library and an instantiation of deque, vector, unordered_map and unordered_set compile, with wave_lock_serialize taking its zero-cost non-AMD pass-through. No NVIDIA GPU was available, so the CUDA backend was not run. The OpenMP backend is likewise unaffected: the valid() reduction change is STDGPU_BACKEND_HIP-guarded, and wave_lock_serialize is a host pass-through off AMD.
stotko
left a comment
stotko
left a comment
There was a problem hiding this comment.
Thanks for contributing fixes for the experimental HIP backend! This is much appreciated.
Some of the fixes appear to be workarounds and I would be happy if you can comment on how this can potentially be fixed upstream in ROCm/HIP as I try to keep workarounds and backend-specific code at a minimum.
| # For HIP backend, mark .cpp examples as HIP language to avoid | ||
| # flag mismatch from hip::device interface propagation. | ||
| if(STDGPU_BACKEND STREQUAL STDGPU_BACKEND_HIP) | ||
| set_source_files_properties("${ARGV0}.cpp" PROPERTIES LANGUAGE HIP) |
There was a problem hiding this comment.
So far, I worked around this by setting the C++ compiler to HIP Clang:
stdgpu/tools/backend/configure_hip.sh
Line 15 in 0bebd1f
I wonder why hip::device adds these incompatible flags to the C++ compiler even with first-class HIP language support enabled in CMake. Therefore, I believe hip::device should be fixed upstream to support such mixed compiler scenarios. IIRC, some very old ROCm versions (perhaps the 3.x or 4.x era) supported this scenario. Do you know more about the reason of the target's behavior and why is the flags are exposed?
In general, I would suggest to patch roc::rocthrust/hip::device or update stdgpu's thrust::thrust target, if possible, as downstream user who link against stdgpu may also face this issue.
There was a problem hiding this comment.
On closer investigation, these LANGUAGE HIP markings are not actually a workaround for the -x hip propagation. I stripped hip::device's -x hip/--offload-arch interface flags entirely and these translation units still fail to compile with the host C++ compiler: the rocThrust/rocPRIM headers they include use device-compiler builtins (e.g. __builtin_amdgcn_wavefrontsize) and clang template extensions that g++ cannot parse. So compiling these sources with the HIP compiler is required regardless of the flag -- the standard host-compiler + per-source LANGUAGE HIP pattern (the C++ compiler stays the host compiler; only device sources go through the HIP compiler).
The -x hip propagation you noticed is real and separate: hip::device adds it to its CXX interface and rocThrust links hip::device publicly, so it reaches consumers. I agree that is awkward for the mixed-compiler case and looks like an upstream hip::device matter rather than something stdgpu can patch cleanly, since stripping it does not remove the need to compile these TUs as HIP. I corrected the comments in 5fe849b to state the actual reason instead of attributing the markings to the flag.
There was a problem hiding this comment.
Thanks for the explanation. I still find this behavior questionable and believe that such HIP-exclusive compiler builtins should not be included unconditionally. thrust with its CUDA and OpenMP backends perfectly supports mixed-compiler scenario and especially running host-side algorithms (via thrust::host execution policy and proper dispatching) in plain CXX files without the need for the CUDA compiler. My impression so far was that rocThrust adds an equivalent HIP backend to thrust that works similarly.
That said, I am not against the proposed changes and they are fine with me. However, I still consider this approach as a workaround to a fundamental limitation and not the intended way for users. Thus, there is much room for improvement for rocThrust to catch up with CUDA thrust. So, could you raise the attention for this limitation as I am for sure not the first one mentioning this issue?
There was a problem hiding this comment.
Understood, and I agree host-side thrust algorithms should not require the HIP compiler. I have logged the mixed-compiler / host-only-thrust limitation to raise with the rocThrust maintainers so downstream stdgpu users are not forced through the HIP compiler for host-side code. Nothing to change here for it.
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## master #484 +/- ##
==========================================
- Coverage 96.79% 96.68% -0.12%
==========================================
Files 33 35 +2
Lines 2622 2652 +30
==========================================
+ Hits 2538 2564 +26
- Misses 84 88 +4 ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
The wave-serialization changes added to deque, vector, and unordered_base
detail headers were not run through the project's clang-format (v18) config,
failing the Clang-Format lint job. This reformats only those touched regions
to the project style; no functional change.
Authored with assistance from Claude.
Test Plan: clang-format 18.1.3 (matching CI's clang-format-18) reports clean:
```
clang-format --dry-run --Werror src/stdgpu/impl/{deque,vector,unordered_base}_detail.cuh src/stdgpu/impl/numeric_detail.h
```
|
Thanks for taking a look, and for wanting to keep the backend-specific surface small; I agree that's the right goal. Here is the rationale for each change and whether it can move upstream. 1. Wave-serialized locking ( This is the one change that is not a workaround for a ROCm defect but an adaptation to the hardware execution model, so it cannot be addressed in ROCm/HIP. The concurrent containers use a try_lock/retry spin loop. On NVIDIA Volta and later, Independent Thread Scheduling gives each lane an independent program counter, so a lane that acquires the lock makes forward progress while its peers keep spinning. AMD CDNA (wave64) and RDNA (wave32) execute a wavefront in lockstep with no per-lane forward-progress guarantee: the lane that wins the CAS cannot advance past the reconvergence point until the losing lanes also exit the loop, which never happens, so the kernel livelocks. 2. This is a workaround for a rocThrust bug and can be removed once that is fixed. Where 3. CMake These are less a workaround than the standard way to build HIP sources with the CMake HIP language: rocThrust's |
…ernel Addresses review feedback on transform_reduce_index: the hand-written block-reduction kernel was only needed because rocThrust's single-block thrust::transform_reduce can hang in its result-copy path on some RDNA GPUs for small, degenerate reductions (validity checks over a nearly-full capacity-1 table). Rather than carry a full custom reduction (which also ignored the execution policy's stream), only the small-size case is special cased: the transformed values are materialized with a plain thrust::transform (an element-wise launch that is unaffected) and folded on the host. Larger reductions keep using the regular multi-block thrust::transform_reduce path. This removes the custom kernel and its stream-0 behavior while preserving the identical reduction (same init, same associative operator). Authored with assistance from Claude. Test Plan: built the HIP backend for gfx90a (ROCm 7.2) and ran the full test suite; all 702 tests pass. ``` cmake -B build -DSTDGPU_BACKEND=STDGPU_BACKEND_HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DSTDGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release cmake --build build -j && ./build/bin/teststdgpu ```
The HIP device path of transform_reduce_index folds reductions of size <= transform_reduce_index_host_threshold on the host to avoid a rocThrust single-block transform_reduce hang on RDNA4. The threshold value (4096) was previously a round guess "above" rocPRIM's single-block cutoff. Comment-only change: the value is unchanged; the comment now documents that 4096 is exactly rocPRIM's measured single-block reduce capacity, so the bound is derived rather than guessed. rocPRIM runs the single-block reduce when the size does not exceed its reduce-config items_per_block (block_size * items_per_thread); above that the reduce is multi-block and does not hit the hang. Measuring that config on the target GPU with rocPRIM's own debug output gives items_per_block = 256 * 16 = 4096 for both element types reduced here (bool for the validity AND-checks, index_t for bitset::count). Folding sizes up to 4096 therefore covers exactly the single-block path and never folds a larger reduction. Authored with the assistance of Claude (Anthropic). Test Plan: measured rocPRIM's single-block cutoff on gfx1201 (RX 9070 XT, ROCm 7.14 / rocPRIM 4.4.0) by calling rocprim::reduce with debug_synchronous=true for bool and int inputs: ``` clang++ -std=c++17 -x hip probe.cpp -o probe.exe --offload-arch=gfx1201 \ -I$ROCM/include -L$ROCM/lib -lamdhip64 -D__HIP_PLATFORM_AMD__ HIP_VISIBLE_DEVICES=0 ./probe.exe # input_type=bool: block_size 256, items_per_block 4096, number of blocks 1 # input_type=int : block_size 256, items_per_block 4096, number of blocks 1 ``` Device code is unchanged (comment only), so the existing GPU validation of teststdgpu (all 702 pass on gfx90a / gfx1100 / gfx1201) still holds.
…atch Addresses review feedback on the wave-serialization helper: - Renames the helper from wave_lock_serialize to warp_convergent_execute: the body is not always serialized (the CUDA and OpenMP backends run it directly), and "warp" is the more common term, used throughout HIP as well. - Renames impl/wave_lock.h to impl/execution_detail.h, following the naming of the closest standard-library headers, and registers it in the installed header set (the previous file was never added to a CMake FILE_SET). - Moves the backend-specific implementation into the respective backend subdirectories and namespaces (stdgpu::hip / stdgpu::cuda / stdgpu::openmp), dispatched via STDGPU_BACKEND_NAMESPACE, mirroring the atomic and memory facilities. The generic stdgpu::detail wrapper now shields the backend macros from the containers. The dispatch is by backend rather than by __HIP_DEVICE_COMPILE__, so the device-intrinsic guards are no longer needed. The HIP implementation is unchanged (same ballot-based election); this only relocates and renames it. Authored with assistance from Claude. Test Plan: built the HIP backend for gfx90a (ROCm 7.2) and ran the full test suite; all 702 tests pass. ``` cmake -B build -DSTDGPU_BACKEND=STDGPU_BACKEND_HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DSTDGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release cmake --build build -j && ./build/bin/teststdgpu ```
Addresses review feedback on the HIP warp-convergence helper: - Replaces __ballot(1) with __activemask(). On HIP __activemask() is defined as __ballot(true), so this is equivalent while using the more portable name that also matches the modern CUDA spelling. - Replaces __lane_id() (an undocumented built-in) with the lane computed from the documented threadIdx/warpSize builtins. HIP numbers threads within a wavefront in row-major order, so the lane is the linear thread index modulo the wavefront size, which is exactly what __lane_id() returns. Both produce the same election as before; this only changes how the active mask and lane index are obtained. Authored with assistance from Claude. Test Plan: built the HIP backend for gfx90a (ROCm 7.2, wave64) and ran the full test suite; all 702 tests pass, including the concurrent contention tests that previously livelocked (simultaneous_push_*, insert/erase_range_unique_parallel, insert_parallel_while_one_free). ``` cmake -B build -DSTDGPU_BACKEND=STDGPU_BACKEND_HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DSTDGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release cmake --build build -j && ./build/bin/teststdgpu ```
The comments claimed these sources are marked LANGUAGE HIP to avoid a flag mismatch from hip::device propagating -x hip to CXX consumers. That is not the actual reason: these translation units include rocThrust/rocPRIM headers, which use device-compiler builtins (e.g. __builtin_amdgcn_wavefrontsize) and clang template extensions that the host C++ compiler cannot parse at all, independent of any -x hip flag. Compiling them with the HIP compiler is therefore required, not a workaround. This is the standard host-compiler + per-source LANGUAGE HIP pattern. Comment-only change; the markings are unchanged. Authored with assistance from Claude.
stotko
left a comment
stotko
left a comment
There was a problem hiding this comment.
The current version looks much better, thanks 👍 I have a couple of minor comments, mostly some stylistic cleanups that I would like to get in before merging this.
| # For HIP backend, mark .cpp examples as HIP language to avoid | ||
| # flag mismatch from hip::device interface propagation. | ||
| if(STDGPU_BACKEND STREQUAL STDGPU_BACKEND_HIP) | ||
| set_source_files_properties("${ARGV0}.cpp" PROPERTIES LANGUAGE HIP) |
There was a problem hiding this comment.
Thanks for the explanation. I still find this behavior questionable and believe that such HIP-exclusive compiler builtins should not be included unconditionally. thrust with its CUDA and OpenMP backends perfectly supports mixed-compiler scenario and especially running host-side algorithms (via thrust::host execution policy and proper dispatching) in plain CXX files without the need for the CUDA compiler. My impression so far was that rocThrust adds an equivalent HIP backend to thrust that works similarly.
That said, I am not against the proposed changes and they are fine with me. However, I still consider this approach as a workaround to a fundamental limitation and not the intended way for users. Thus, there is much room for improvement for rocThrust to catch up with CUDA thrust. So, could you raise the attention for this limitation as I am for sure not the first one mentioning this issue?
Address review feedback on the concurrent-container spinlock port.
The detail headers (deque/unordered_base/vector) included the private
impl/execution_detail.h directly. Per the project convention that a
detail header is pulled in only through its user-facing header, they now
include stdgpu/execution.h, which gains an include of
impl/execution_detail.h at the bottom (mirroring atomic.cuh/memory.h).
Remove the inline "wave-serialize" comments at each warp_convergent_execute
call site and the backend-specific explanation in execution_detail.h; the
per-backend rationale already lives in the stdgpu::hip / stdgpu::cuda /
stdgpu::openmp execution headers, so it is not duplicated at the call
sites or the backend-agnostic wrapper.
No functional change to device code; only header include paths and
comments are touched.
Authored with the assistance of Claude (Anthropic).
Test Plan:
Built and ran the full test suite on AMD gfx90a (ROCm 7.2.1, wave64):
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release \
-DSTDGPU_BACKEND=STDGPU_BACKEND_HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DSTDGPU_BUILD_TESTS=ON -DSTDGPU_BUILD_EXAMPLES=ON
cmake --build build -j$(nproc)
build/bin/teststdgpu
702 tests passed, 0 memory leaks.
Also configured and ran the OpenMP backend (CCCL thrust) on the same host:
cmake -B build-omp -S . -DSTDGPU_BACKEND=STDGPU_BACKEND_OPENMP \
-DSTDGPU_BUILD_TESTS=ON
cmake --build build-omp -j$(nproc)
build-omp/bin/teststdgpu
608 tests passed, 0 memory leaks.
The 72687fe commit added #include <stdgpu/impl/execution_detail.h> at the bottom of execution.h so that user-facing header files (deque.cuh, etc.) can include execution.h instead of reaching into execution_detail.h directly. However, execution_detail.h and the cuda/execution.cuh it pulls in both declare functions with STDGPU_DEVICE_ONLY, which expands to STDGPU_CUDA_DEVICE_ONLY -- a macro that is intentionally left undefined when the host C++ compiler (g++) compiles plain .cpp translation units (i.e. TUs that are not processed by nvcc). This caused build failures in the CUDA backend for every .cpp file that transitively includes execution.h. Fix: add #include <stdgpu/platform.h> before the guard (so the numeric values STDGPU_BACKEND_CUDA, STDGPU_BACKEND_OPENMP, etc. are defined) and wrap the execution_detail.h include in a preprocessor guard that fires only when a device compiler is active (__CUDACC__ for nvcc / clang-cuda, __HIP__ for clang++/HIP) or when the OpenMP backend is in use (which has no device-only qualifiers). The guard is a no-op for HIP and OpenMP builds (both compilers define the relevant macro), and fixes the CUDA backend by suppressing the problematic include in g++-compiled TUs. Verified: CUDA backend (nvcc 12.6, arch 80) builds to completion including teststdgpu; HIP backend (ROCm 7.2.1, gfx90a) and OpenMP backend (CCCL Thrust 2.6.1) rebuild cleanly with no device-code change (incremental builds only relink, no object recompilation). Authored with the assistance of Claude (Anthropic). Test Plan: CUDA backend (compile-only, no NVIDIA GPU): cmake -B build-cuda -S . -DSTDGPU_BACKEND=STDGPU_BACKEND_CUDA \ -DCMAKE_CUDA_COMPILER=/opt/conda/envs/cuda/bin/nvcc \ -DCMAKE_CUDA_ARCHITECTURES=80 -DSTDGPU_BUILD_TESTS=ON cmake --build build-cuda --parallel $(nproc) -> All targets built, including teststdgpu (linked but not run) HIP backend (incremental rebuild): cmake --build build -j$(nproc) -> Only relinking, no HIP object recompilation; confirms no device-code change OpenMP backend (incremental rebuild): cmake --build build-omp -j$(nproc) -> Only relinking; confirms no device-code change
stotko
left a comment
stotko
left a comment
There was a problem hiding this comment.
Thanks for addressing the comments and testing the code! Could you incorporate the final cleanup suggested below and ensure that formatting via clang-format is correct? I believe the PR should be ready then. The Windows-specific build failures are unrelated to this PR and I will take care of them.
Two cleanups to the device-only execution helper include layering.
Replace the hand-rolled guard in execution.h with the existing
STDGPU_DETAIL_IS_DEVICE_COMPILED helper from platform.h, which already
checks whether the current translation unit is compiled by a device
compiler and covers every backend (CUDA/HIP are 1 only under the device
compiler, OpenMP is always 1). This keeps the host-only .cpp translation
units in the CUDA backend from seeing the __device__-qualified helper
while still pulling it in for OpenMP.
Rename impl/execution_detail.h to impl/execution_detail.cuh to make its
device-only nature explicit, matching the other detail headers that
define device code, and shorten the now-redundant comment above the
conditional include. Also sort the execution.h include added to
unordered_base_detail.h so clang-format is clean.
This work was authored with the assistance of Claude, an AI assistant.
Test Plan:
Built and ran the test suite on an AMD Instinct MI250X (gfx90a),
ROCm 7.2.1, with interface header-set verification enabled:
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release \
-DSTDGPU_BACKEND=STDGPU_BACKEND_HIP -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DSTDGPU_BUILD_TESTS=ON -DSTDGPU_BUILD_EXAMPLES=ON \
-DCMAKE_VERIFY_INTERFACE_HEADER_SETS=ON
cmake --build build -j
./build/bin/teststdgpu # 702/702 passed, 0 leaks
Compile-checked the CUDA backend with nvcc 12.6 (host TUs via g++) and
built the OpenMP backend, both with header-set verification:
cmake -B build-cuda -S . -DSTDGPU_BACKEND=STDGPU_BACKEND_CUDA \
-DCMAKE_CUDA_COMPILER=nvcc -DCMAKE_CUDA_ARCHITECTURES=70 \
-DCMAKE_VERIFY_INTERFACE_HEADER_SETS=ON
cmake --build build-cuda -j # builds and links
cmake -B build-omp -S . -DSTDGPU_BACKEND=STDGPU_BACKEND_OPENMP \
-DCMAKE_VERIFY_INTERFACE_HEADER_SETS=ON
cmake --build build-omp -j
./build-omp/bin/teststdgpu # 608/608 passed, 0 leaks
2 participants
This fixes the existing experimental HIP backend so it builds on ROCm 7.x and runs correctly across both AMD wavefront sizes: wave64 (CDNA) and wave32 (RDNA). It does not add a new backend; it repairs the one already present (STDGPU_BACKEND_HIP).
Three independent issues are addressed. Review them in order:
CMake HIP language propagation (build break). On ROCm 7.x, linking against rocThrust pulls in hip::device, whose INTERFACE_COMPILE_OPTIONS add HIP flags (-x hip, --offload-arch=...) gated on $<COMPILE_LANGUAGE:CXX>. Those flags reach the system C++ compiler building stdgpu's .cpp files, which rejects --offload-arch and fails. The fix marks the .cpp sources that include HIP headers as LANGUAGE HIP (shared impl, HIP backend impl, examples, tests) so the HIP compiler handles them. All edits are guarded on STDGPU_BACKEND_HIP; the CUDA and OpenMP builds are untouched.
wave64/wave32 spinlock livelock (hang in concurrent containers). The lock-free containers retry per-slot try_lock() in spin loops. AMD GPUs give no per-lane forward-progress guarantee: when several lanes of one wavefront contend for the same lock, the winning lane is stuck at SIMT reconvergence waiting on the losers, who spin forever. CUDA Volta+ Independent Thread Scheduling lets the winner proceed; AMD's execution model does not. A new helper, detail::wave_lock_serialize, elects one active lane at a time via __ballot and runs the spin-loop body for that lane, so each lane completes its critical section in turn. It is wave-size agnostic (the ballot mask covers whatever lanes are active, 32 or 64) and is a zero-cost pass-through on CUDA, where ITS already handles this. It wraps the insert/erase loops in unordered_base and the push/pop loops in deque and vector.
RDNA hang in valid() (degenerate single-slot reduction). Container valid() checks route through transform_reduce_index, a wrapper over thrust::transform_reduce. On a degenerate single-slot table that is full or has an emptied excess list, this lowers to a single-block rocPRIM reduce whose internal hipMemcpyWithStream copy of the result never completes on some RDNA GPUs (the launch returns success but the result copy wedges until the watchdog fires). For the device execution policy on the HIP backend only, this replaces the thrust reduce with a small hand-written kernel plus a plain hipMemcpy device-to-host and a host-side fold, computing the identical reduction (same init, same associative binary op). The host policy and the CUDA backend keep the thrust path unchanged.
The shared headers (wave_lock.h, numeric_detail.h) gate every AMD-specific path behind HIP_PLATFORM_AMD / STDGPU_BACKEND_HIP, so CUDA and OpenMP behavior is unchanged.
This work was authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan
Built and ran the full test suite on three AMD GPUs spanning both wavefront sizes:
gfx90a (MI250X, CDNA2, wave64), ROCm 7.2.x, Linux:
gfx1100 (RDNA3, wave32), ROCm, Linux:
gfx1201 (RX 9070 XT, RDNA4, wave32), ROCm 7.14, Windows:
All 702 tests pass on every arch, deterministic across repeated runs. The previously hanging concurrent-container tests (insert/erase/emplace while full or with an emptied excess list, and the simultaneous push/pop suites) now pass.
The CUDA backend was additionally compile-checked with nvcc (CUDA 13.3) to confirm these shared-header changes do not affect it: the library and an instantiation of deque, vector, unordered_map and unordered_set compile, with wave_lock_serialize taking its zero-cost non-AMD pass-through. No NVIDIA GPU was available, so the CUDA backend was not run. The OpenMP backend is likewise unaffected: the valid() reduction change is STDGPU_BACKEND_HIP-guarded, and wave_lock_serialize is a host pass-through off AMD.