From e19572d0ff2ced186faa0ea6f1888587df19c54b Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Fri, 19 Jun 2026 23:54:40 +0000 Subject: [PATCH 01/10] [ROCm] Fix the experimental HIP backend for ROCm 7.x on AMD GPUs 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 $. 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. --- examples/CMakeLists.txt | 5 + src/stdgpu/CMakeLists.txt | 12 ++ src/stdgpu/hip/CMakeLists.txt | 8 ++ src/stdgpu/impl/deque_detail.cuh | 133 ++++++++++++---------- src/stdgpu/impl/numeric_detail.h | 106 ++++++++++++++++- src/stdgpu/impl/unordered_base_detail.cuh | 45 +++++--- src/stdgpu/impl/vector_detail.cuh | 67 ++++++----- src/stdgpu/impl/wave_lock.h | 87 ++++++++++++++ tests/stdgpu/CMakeLists.txt | 18 +++ 9 files changed, 367 insertions(+), 114 deletions(-) create mode 100644 src/stdgpu/impl/wave_lock.h diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 94192defe..153e370a1 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -18,6 +18,11 @@ endmacro() macro(stdgpu_add_example_cpp) stdgpu_detail_add_example(${ARGV0} "cpp") + # 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) + endif() endmacro() diff --git a/src/stdgpu/CMakeLists.txt b/src/stdgpu/CMakeLists.txt index 2c5ffcf34..98e19d206 100644 --- a/src/stdgpu/CMakeLists.txt +++ b/src/stdgpu/CMakeLists.txt @@ -32,6 +32,18 @@ else() add_library(stdgpu STATIC) endif() +# For the HIP backend, the shared impl/*.cpp files include backend headers +# (via STDGPU_DETAIL_BACKEND_HEADER) which pull in HIP/rocPRIM headers that +# require the HIP compiler. Mark them as HIP language to avoid flag mismatch +# when linking against rocThrust (which propagates -x hip to CXX consumers). +if(STDGPU_BACKEND STREQUAL STDGPU_BACKEND_HIP) + set_source_files_properties( + ${CMAKE_CURRENT_SOURCE_DIR}/impl/device.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/impl/iterator.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/impl/memory.cpp + PROPERTIES LANGUAGE HIP) +endif() + target_sources(stdgpu PRIVATE impl/device.cpp impl/iterator.cpp impl/memory.cpp) diff --git a/src/stdgpu/hip/CMakeLists.txt b/src/stdgpu/hip/CMakeLists.txt index dbb1cdb6a..f4e6cb5b3 100644 --- a/src/stdgpu/hip/CMakeLists.txt +++ b/src/stdgpu/hip/CMakeLists.txt @@ -4,9 +4,17 @@ #find_dependency(hip 7.1 REQUIRED) #" PARENT_SCOPE) +# Mark HIP backend sources as HIP language (these include HIP headers). +# Use TARGET_DIRECTORY to ensure properties apply to sources as seen by the target. target_sources(stdgpu PRIVATE impl/device.cpp impl/memory.cpp) +set_source_files_properties( + ${CMAKE_CURRENT_SOURCE_DIR}/impl/device.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/impl/memory.cpp + TARGET_DIRECTORY stdgpu + PROPERTIES LANGUAGE HIP) + if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.23) target_sources(stdgpu PUBLIC FILE_SET stdgpu_backend_headers TYPE HEADERS diff --git a/src/stdgpu/impl/deque_detail.cuh b/src/stdgpu/impl/deque_detail.cuh index 637953aa5..e309e3082 100644 --- a/src/stdgpu/impl/deque_detail.cuh +++ b/src/stdgpu/impl/deque_detail.cuh @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -218,28 +219,31 @@ deque::push_back(const T& element) { index_t push_position = static_cast(_end.fetch_inc_mod(static_cast(capacity()))); - while (!pushed) - { - if (_locks[push_position].try_lock()) + // Wave-serialize to avoid livelock on AMD wave64/wave32 + detail::wave_lock_serialize([&]() { + while (!pushed) { - // START --- critical section --- START - - if (!occupied(push_position)) + if (_locks[push_position].try_lock()) { - allocator_traits::construct(_allocator, &(_data[push_position]), element); - bool was_occupied = _occupied.set(push_position); - pushed = true; + // START --- critical section --- START - if (was_occupied) + if (!occupied(push_position)) { - printf("stdgpu::deque::push_back : Expected entry to be not occupied but actually was\n"); + allocator_traits::construct(_allocator, &(_data[push_position]), element); + bool was_occupied = _occupied.set(push_position); + pushed = true; + + if (was_occupied) + { + printf("stdgpu::deque::push_back : Expected entry to be not occupied but actually was\n"); + } } - } - // END --- critical section --- END - _locks[push_position].unlock(); + // END --- critical section --- END + _locks[push_position].unlock(); + } } - } + }); } else { @@ -271,28 +275,31 @@ deque::pop_back() index_t pop_position = static_cast(_end.fetch_dec_mod(static_cast(capacity()))); pop_position = (pop_position == 0) ? capacity() - 1 : pop_position - 1; // Manually reconstruct stored value - while (!popped.second) - { - if (_locks[pop_position].try_lock()) + // Wave-serialize to avoid livelock on AMD wave64/wave32 + detail::wave_lock_serialize([&]() { + while (!popped.second) { - // START --- critical section --- START - - if (occupied(pop_position)) + if (_locks[pop_position].try_lock()) { - bool was_occupied = _occupied.reset(pop_position); - allocator_traits::construct(_allocator, &popped, _data[pop_position], true); - allocator_traits::destroy(_allocator, &(_data[pop_position])); + // START --- critical section --- START - if (!was_occupied) + if (occupied(pop_position)) { - printf("stdgpu::deque::pop_back : Expected entry to be occupied but actually was not\n"); + bool was_occupied = _occupied.reset(pop_position); + allocator_traits::construct(_allocator, &popped, _data[pop_position], true); + allocator_traits::destroy(_allocator, &(_data[pop_position])); + + if (!was_occupied) + { + printf("stdgpu::deque::pop_back : Expected entry to be occupied but actually was not\n"); + } } - } - // END --- critical section --- END - _locks[pop_position].unlock(); + // END --- critical section --- END + _locks[pop_position].unlock(); + } } - } + }); } else { @@ -331,28 +338,31 @@ deque::push_front(const T& element) index_t push_position = static_cast(_begin.fetch_dec_mod(static_cast(capacity()))); push_position = (push_position == 0) ? capacity() - 1 : push_position - 1; // Manually reconstruct stored value - while (!pushed) - { - if (_locks[push_position].try_lock()) + // Wave-serialize to avoid livelock on AMD wave64/wave32 + detail::wave_lock_serialize([&]() { + while (!pushed) { - // START --- critical section --- START - - if (!occupied(push_position)) + if (_locks[push_position].try_lock()) { - allocator_traits::construct(_allocator, &(_data[push_position]), element); - bool was_occupied = _occupied.set(push_position); - pushed = true; + // START --- critical section --- START - if (was_occupied) + if (!occupied(push_position)) { - printf("stdgpu::deque::push_front : Expected entry to be not occupied but actually was\n"); + allocator_traits::construct(_allocator, &(_data[push_position]), element); + bool was_occupied = _occupied.set(push_position); + pushed = true; + + if (was_occupied) + { + printf("stdgpu::deque::push_front : Expected entry to be not occupied but actually was\n"); + } } - } - // END --- critical section --- END - _locks[push_position].unlock(); + // END --- critical section --- END + _locks[push_position].unlock(); + } } - } + }); } else { @@ -383,28 +393,31 @@ deque::pop_front() { index_t pop_position = static_cast(_begin.fetch_inc_mod(static_cast(capacity()))); - while (!popped.second) - { - if (_locks[pop_position].try_lock()) + // Wave-serialize to avoid livelock on AMD wave64/wave32 + detail::wave_lock_serialize([&]() { + while (!popped.second) { - // START --- critical section --- START - - if (occupied(pop_position)) + if (_locks[pop_position].try_lock()) { - bool was_occupied = _occupied.reset(pop_position); - allocator_traits::construct(_allocator, &popped, _data[pop_position], true); - allocator_traits::destroy(_allocator, &(_data[pop_position])); + // START --- critical section --- START - if (!was_occupied) + if (occupied(pop_position)) { - printf("stdgpu::deque::pop_front : Expected entry to be occupied but actually was not\n"); + bool was_occupied = _occupied.reset(pop_position); + allocator_traits::construct(_allocator, &popped, _data[pop_position], true); + allocator_traits::destroy(_allocator, &(_data[pop_position])); + + if (!was_occupied) + { + printf("stdgpu::deque::pop_front : Expected entry to be occupied but actually was not\n"); + } } - } - // END --- critical section --- END - _locks[pop_position].unlock(); + // END --- critical section --- END + _locks[pop_position].unlock(); + } } - } + }); } else { diff --git a/src/stdgpu/impl/numeric_detail.h b/src/stdgpu/impl/numeric_detail.h index c99117444..ddff0d94b 100644 --- a/src/stdgpu/impl/numeric_detail.h +++ b/src/stdgpu/impl/numeric_detail.h @@ -18,15 +18,100 @@ #include #include +#include #include #include +#include +#include + +#if STDGPU_BACKEND == STDGPU_BACKEND_HIP + #include + + #include + #include +#endif namespace stdgpu { namespace detail { + +#if STDGPU_BACKEND == STDGPU_BACKEND_HIP +template +__global__ void +transform_reduce_index_kernel(IndexType size, T init, BinaryFunction reduce, UnaryFunction f, T* block_results) +{ + __shared__ T shared[block_size]; + + index_t tid = static_cast(threadIdx.x); + IndexType stride = static_cast(block_size) * static_cast(gridDim.x); + + T value = init; + for (IndexType i = static_cast(blockIdx.x) * static_cast(block_size) + tid; i < size; + i += stride) + { + value = reduce(value, f(i)); + } + shared[tid] = value; + __syncthreads(); + + for (index_t offset = block_size / 2; offset > 0; offset /= 2) + { + if (tid < offset) + { + shared[tid] = reduce(shared[tid], shared[tid + offset]); + } + __syncthreads(); + } + + if (tid == 0) + { + block_results[blockIdx.x] = shared[0]; + } +} + +// Hand-written device reduction used in place of thrust::transform_reduce for the device execution policy. +// rocThrust's single-block transform_reduce wedges its internal hipMemcpyWithStream result copy on some +// RDNA GPUs for degenerate inputs; a plain kernel plus a plain hipMemcpy D2H avoids that path while +// computing the identical reduction (same init, same associative BinaryFunction). +template +T +transform_reduce_index_device(IndexType size, T init, BinaryFunction reduce, UnaryFunction f) +{ + if (size <= 0) + { + return init; + } + + constexpr index_t block_size = 256; + constexpr index_t max_blocks = 64; + index_t blocks = static_cast((static_cast(size) + block_size - 1) / block_size); + blocks = (blocks < 1) ? 1 : ((blocks > max_blocks) ? max_blocks : blocks); + + T* block_results = nullptr; + STDGPU_HIP_SAFE_CALL(hipMalloc(&block_results, static_cast(blocks) * sizeof(T))); + + transform_reduce_index_kernel + <<(blocks), block_size>>>(size, init, reduce, f, block_results); + STDGPU_HIP_SAFE_CALL(hipGetLastError()); + STDGPU_HIP_SAFE_CALL(hipDeviceSynchronize()); + + T host_results[max_blocks]; + STDGPU_HIP_SAFE_CALL( + hipMemcpy(host_results, block_results, static_cast(blocks) * sizeof(T), hipMemcpyDeviceToHost)); + STDGPU_HIP_SAFE_CALL(hipFree(block_results)); + + T result = init; + for (index_t i = 0; i < blocks; ++i) + { + result = reduce(result, host_results[i]); + } + return result; +} +#endif + template class iota_functor { @@ -73,12 +158,21 @@ template (policy), - thrust::counting_iterator(0), - thrust::counting_iterator(size), - f, - init, - reduce); +#if STDGPU_BACKEND == STDGPU_BACKEND_HIP + if constexpr (std::is_same_v, execution::device_policy>) + { + return detail::transform_reduce_index_device(size, init, reduce, f); + } + else +#endif + { + return thrust::transform_reduce(std::forward(policy), + thrust::counting_iterator(0), + thrust::counting_iterator(size), + f, + init, + reduce); + } } } // namespace stdgpu diff --git a/src/stdgpu/impl/unordered_base_detail.cuh b/src/stdgpu/impl/unordered_base_detail.cuh index ad72e0fbb..a9ecbbfd9 100644 --- a/src/stdgpu/impl/unordered_base_detail.cuh +++ b/src/stdgpu/impl/unordered_base_detail.cuh @@ -23,6 +23,7 @@ #include #include #include +#include #include #include #include @@ -971,17 +972,22 @@ inline STDGPU_DEVICE_ONLY { pair result(end(), operation_status::failed_collision); - while (true) - { - if (result.second == operation_status::failed_collision && !full() && !_excess_list_positions.empty()) + // Wave-serialize to avoid livelock on AMD wave64/wave32: multiple lanes + // of the same wavefront contending for the same bucket spin forever + // because the winner cannot make progress while waiting at reconvergence. + wave_lock_serialize([&]() { + while (true) { - result = try_insert(value); - } - else - { - break; + if (result.second == operation_status::failed_collision && !full() && !_excess_list_positions.empty()) + { + result = try_insert(value); + } + else + { + break; + } } - } + }); return result.second == operation_status::success ? pair(result.first, true) : pair(result.first, false); @@ -1017,17 +1023,20 @@ unordered_base::erase( { operation_status result = operation_status::failed_collision; - while (true) - { - if (result == operation_status::failed_collision) + // Wave-serialize to avoid livelock on AMD wave64/wave32 + wave_lock_serialize([&]() { + while (true) { - result = try_erase(key); - } - else - { - break; + if (result == operation_status::failed_collision) + { + result = try_erase(key); + } + else + { + break; + } } - } + }); return result == operation_status::success ? 1 : 0; } diff --git a/src/stdgpu/impl/vector_detail.cuh b/src/stdgpu/impl/vector_detail.cuh index fd84cd4d4..8fa4d3465 100644 --- a/src/stdgpu/impl/vector_detail.cuh +++ b/src/stdgpu/impl/vector_detail.cuh @@ -20,6 +20,7 @@ #include #include +#include #include #include #include @@ -195,28 +196,31 @@ vector::push_back(const T& element) // Check position if (0 <= push_position && push_position < capacity()) { - while (!pushed) - { - if (_locks[push_position].try_lock()) + // Wave-serialize to avoid livelock on AMD wave64/wave32 + detail::wave_lock_serialize([&]() { + while (!pushed) { - // START --- critical section --- START - - if (!occupied(push_position)) + if (_locks[push_position].try_lock()) { - allocator_traits::construct(_allocator, &(_data[push_position]), element); - bool was_occupied = _occupied.set(push_position); - pushed = true; + // START --- critical section --- START - if (was_occupied) + if (!occupied(push_position)) { - printf("stdgpu::vector::push_back : Expected entry to be not occupied but actually was\n"); + allocator_traits::construct(_allocator, &(_data[push_position]), element); + bool was_occupied = _occupied.set(push_position); + pushed = true; + + if (was_occupied) + { + printf("stdgpu::vector::push_back : Expected entry to be not occupied but actually was\n"); + } } - } - // END --- critical section --- END - _locks[push_position].unlock(); + // END --- critical section --- END + _locks[push_position].unlock(); + } } - } + }); } else { @@ -248,28 +252,31 @@ vector::pop_back() // Check position if (0 <= pop_position && pop_position < capacity()) { - while (!popped.second) - { - if (_locks[pop_position].try_lock()) + // Wave-serialize to avoid livelock on AMD wave64/wave32 + detail::wave_lock_serialize([&]() { + while (!popped.second) { - // START --- critical section --- START - - if (occupied(pop_position)) + if (_locks[pop_position].try_lock()) { - bool was_occupied = _occupied.reset(pop_position); - allocator_traits::construct(_allocator, &popped, _data[pop_position], true); - allocator_traits::destroy(_allocator, &(_data[pop_position])); + // START --- critical section --- START - if (!was_occupied) + if (occupied(pop_position)) { - printf("stdgpu::vector::pop_back : Expected entry to be occupied but actually was not\n"); + bool was_occupied = _occupied.reset(pop_position); + allocator_traits::construct(_allocator, &popped, _data[pop_position], true); + allocator_traits::destroy(_allocator, &(_data[pop_position])); + + if (!was_occupied) + { + printf("stdgpu::vector::pop_back : Expected entry to be occupied but actually was not\n"); + } } - } - // END --- critical section --- END - _locks[pop_position].unlock(); + // END --- critical section --- END + _locks[pop_position].unlock(); + } } - } + }); } else { diff --git a/src/stdgpu/impl/wave_lock.h b/src/stdgpu/impl/wave_lock.h new file mode 100644 index 000000000..64c6cb863 --- /dev/null +++ b/src/stdgpu/impl/wave_lock.h @@ -0,0 +1,87 @@ +/* + * Copyright 2019 Patrick Stotko + * Copyright 2026 Advanced Micro Devices, Inc. + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef STDGPU_WAVE_LOCK_H +#define STDGPU_WAVE_LOCK_H + +#include + +namespace stdgpu::detail +{ + +/** + * Wave-serialized locking helper for AMD CDNA/RDNA GPUs. + * + * On CUDA Volta+, Independent Thread Scheduling allows a lane that wins + * a spinlock to make forward progress while other lanes spin. On AMD GPUs + * (CDNA wave64, RDNA wave32), there is no per-lane forward-progress + * guarantee: a lane that wins a CAS is stuck at SIMT reconvergence waiting + * on losers who spin forever -- classic livelock. + * + * The fix: wave-serialize by electing one active lane at a time via ballot, + * letting only that lane attempt the lock. Other lanes wait their turn. + * + * Usage: + * wave_lock_serialize([&]() { + * // Your spinlock-based critical section + * while (!done) { + * if (lock.try_lock()) { + * // critical section + * done = true; + * lock.unlock(); + * } + * } + * }); + */ + +#if defined(__HIP_DEVICE_COMPILE__) && defined(__HIP_PLATFORM_AMD__) + +// AMD HIP: wave-serialize using ballot to avoid wave64/wave32 livelock +template +STDGPU_DEVICE_ONLY void +wave_lock_serialize(F&& body) +{ + int lane = static_cast(__lane_id()); + unsigned long long active = __ballot(1); + + while (active) + { + // Elect the lowest-numbered active lane + int elected = __ffsll(static_cast(active)) - 1; + if (lane == elected) + { + body(); + } + // Clear this lane's bit and move to the next + active &= ~(1ull << elected); + } +} + +#else + +// CUDA or other: Independent Thread Scheduling handles this, no serialization needed +template +STDGPU_DEVICE_ONLY void +wave_lock_serialize(F&& body) +{ + body(); +} + +#endif + +} // namespace stdgpu::detail + +#endif // STDGPU_WAVE_LOCK_H diff --git a/tests/stdgpu/CMakeLists.txt b/tests/stdgpu/CMakeLists.txt index 35fcc36e3..dace20ac3 100644 --- a/tests/stdgpu/CMakeLists.txt +++ b/tests/stdgpu/CMakeLists.txt @@ -13,6 +13,24 @@ target_sources(teststdgpu PRIVATE algorithm.cpp target_sources(teststdgpu PRIVATE ../test_memory_utils.cpp) +# For HIP backend, mark .cpp test sources as HIP language to avoid +# flag mismatch from hip::device interface propagation. +if(STDGPU_BACKEND STREQUAL STDGPU_BACKEND_HIP) + set_source_files_properties( + main.cpp + algorithm.cpp + bit.cpp + contract.cpp + functional.cpp + iterator.cpp + limits.cpp + memory.cpp + numeric.cpp + ranges.cpp + ../test_memory_utils.cpp + PROPERTIES LANGUAGE HIP) +endif() + add_subdirectory(${STDGPU_BACKEND_DIRECTORY}) target_include_directories(teststdgpu PRIVATE From e03b4ab274a86dcfeae34a20a7f5e9d143f9da67 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 22 Jun 2026 18:54:59 +0000 Subject: [PATCH 02/10] [ROCm] Apply clang-format to ported wave-lock critical sections 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 ``` --- src/stdgpu/impl/deque_detail.cuh | 174 ++++++++++++---------- src/stdgpu/impl/numeric_detail.h | 6 +- src/stdgpu/impl/unordered_base_detail.cuh | 53 ++++--- src/stdgpu/impl/vector_detail.cuh | 87 ++++++----- 4 files changed, 180 insertions(+), 140 deletions(-) diff --git a/src/stdgpu/impl/deque_detail.cuh b/src/stdgpu/impl/deque_detail.cuh index e309e3082..773c2f875 100644 --- a/src/stdgpu/impl/deque_detail.cuh +++ b/src/stdgpu/impl/deque_detail.cuh @@ -220,30 +220,35 @@ deque::push_back(const T& element) index_t push_position = static_cast(_end.fetch_inc_mod(static_cast(capacity()))); // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize([&]() { - while (!pushed) - { - if (_locks[push_position].try_lock()) + detail::wave_lock_serialize( + [&]() { - // START --- critical section --- START - - if (!occupied(push_position)) + while (!pushed) { - allocator_traits::construct(_allocator, &(_data[push_position]), element); - bool was_occupied = _occupied.set(push_position); - pushed = true; - - if (was_occupied) + if (_locks[push_position].try_lock()) { - printf("stdgpu::deque::push_back : Expected entry to be not occupied but actually was\n"); + // START --- critical section --- START + + if (!occupied(push_position)) + { + allocator_traits::construct(_allocator, + &(_data[push_position]), + element); + bool was_occupied = _occupied.set(push_position); + pushed = true; + + if (was_occupied) + { + printf("stdgpu::deque::push_back : Expected entry to be not occupied but actually " + "was\n"); + } + } + + // END --- critical section --- END + _locks[push_position].unlock(); } } - - // END --- critical section --- END - _locks[push_position].unlock(); - } - } - }); + }); } else { @@ -276,30 +281,36 @@ deque::pop_back() pop_position = (pop_position == 0) ? capacity() - 1 : pop_position - 1; // Manually reconstruct stored value // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize([&]() { - while (!popped.second) - { - if (_locks[pop_position].try_lock()) + detail::wave_lock_serialize( + [&]() { - // START --- critical section --- START - - if (occupied(pop_position)) + while (!popped.second) { - bool was_occupied = _occupied.reset(pop_position); - allocator_traits::construct(_allocator, &popped, _data[pop_position], true); - allocator_traits::destroy(_allocator, &(_data[pop_position])); - - if (!was_occupied) + if (_locks[pop_position].try_lock()) { - printf("stdgpu::deque::pop_back : Expected entry to be occupied but actually was not\n"); + // START --- critical section --- START + + if (occupied(pop_position)) + { + bool was_occupied = _occupied.reset(pop_position); + allocator_traits::construct(_allocator, + &popped, + _data[pop_position], + true); + allocator_traits::destroy(_allocator, &(_data[pop_position])); + + if (!was_occupied) + { + printf("stdgpu::deque::pop_back : Expected entry to be occupied but actually was " + "not\n"); + } + } + + // END --- critical section --- END + _locks[pop_position].unlock(); } } - - // END --- critical section --- END - _locks[pop_position].unlock(); - } - } - }); + }); } else { @@ -339,30 +350,35 @@ deque::push_front(const T& element) push_position = (push_position == 0) ? capacity() - 1 : push_position - 1; // Manually reconstruct stored value // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize([&]() { - while (!pushed) - { - if (_locks[push_position].try_lock()) + detail::wave_lock_serialize( + [&]() { - // START --- critical section --- START - - if (!occupied(push_position)) + while (!pushed) { - allocator_traits::construct(_allocator, &(_data[push_position]), element); - bool was_occupied = _occupied.set(push_position); - pushed = true; - - if (was_occupied) + if (_locks[push_position].try_lock()) { - printf("stdgpu::deque::push_front : Expected entry to be not occupied but actually was\n"); + // START --- critical section --- START + + if (!occupied(push_position)) + { + allocator_traits::construct(_allocator, + &(_data[push_position]), + element); + bool was_occupied = _occupied.set(push_position); + pushed = true; + + if (was_occupied) + { + printf("stdgpu::deque::push_front : Expected entry to be not occupied but actually " + "was\n"); + } + } + + // END --- critical section --- END + _locks[push_position].unlock(); } } - - // END --- critical section --- END - _locks[push_position].unlock(); - } - } - }); + }); } else { @@ -394,30 +410,36 @@ deque::pop_front() index_t pop_position = static_cast(_begin.fetch_inc_mod(static_cast(capacity()))); // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize([&]() { - while (!popped.second) - { - if (_locks[pop_position].try_lock()) + detail::wave_lock_serialize( + [&]() { - // START --- critical section --- START - - if (occupied(pop_position)) + while (!popped.second) { - bool was_occupied = _occupied.reset(pop_position); - allocator_traits::construct(_allocator, &popped, _data[pop_position], true); - allocator_traits::destroy(_allocator, &(_data[pop_position])); - - if (!was_occupied) + if (_locks[pop_position].try_lock()) { - printf("stdgpu::deque::pop_front : Expected entry to be occupied but actually was not\n"); + // START --- critical section --- START + + if (occupied(pop_position)) + { + bool was_occupied = _occupied.reset(pop_position); + allocator_traits::construct(_allocator, + &popped, + _data[pop_position], + true); + allocator_traits::destroy(_allocator, &(_data[pop_position])); + + if (!was_occupied) + { + printf("stdgpu::deque::pop_front : Expected entry to be occupied but actually was " + "not\n"); + } + } + + // END --- critical section --- END + _locks[pop_position].unlock(); } } - - // END --- critical section --- END - _locks[pop_position].unlock(); - } - } - }); + }); } else { diff --git a/src/stdgpu/impl/numeric_detail.h b/src/stdgpu/impl/numeric_detail.h index ddff0d94b..1911373c2 100644 --- a/src/stdgpu/impl/numeric_detail.h +++ b/src/stdgpu/impl/numeric_detail.h @@ -99,8 +99,10 @@ transform_reduce_index_device(IndexType size, T init, BinaryFunction reduce, Una STDGPU_HIP_SAFE_CALL(hipDeviceSynchronize()); T host_results[max_blocks]; - STDGPU_HIP_SAFE_CALL( - hipMemcpy(host_results, block_results, static_cast(blocks) * sizeof(T), hipMemcpyDeviceToHost)); + STDGPU_HIP_SAFE_CALL(hipMemcpy(host_results, + block_results, + static_cast(blocks) * sizeof(T), + hipMemcpyDeviceToHost)); STDGPU_HIP_SAFE_CALL(hipFree(block_results)); T result = init; diff --git a/src/stdgpu/impl/unordered_base_detail.cuh b/src/stdgpu/impl/unordered_base_detail.cuh index a9ecbbfd9..36aa48d55 100644 --- a/src/stdgpu/impl/unordered_base_detail.cuh +++ b/src/stdgpu/impl/unordered_base_detail.cuh @@ -975,19 +975,22 @@ inline STDGPU_DEVICE_ONLY // Wave-serialize to avoid livelock on AMD wave64/wave32: multiple lanes // of the same wavefront contending for the same bucket spin forever // because the winner cannot make progress while waiting at reconvergence. - wave_lock_serialize([&]() { - while (true) - { - if (result.second == operation_status::failed_collision && !full() && !_excess_list_positions.empty()) - { - result = try_insert(value); - } - else + wave_lock_serialize( + [&]() { - break; - } - } - }); + while (true) + { + if (result.second == operation_status::failed_collision && !full() && + !_excess_list_positions.empty()) + { + result = try_insert(value); + } + else + { + break; + } + } + }); return result.second == operation_status::success ? pair(result.first, true) : pair(result.first, false); @@ -1024,19 +1027,21 @@ unordered_base::erase( operation_status result = operation_status::failed_collision; // Wave-serialize to avoid livelock on AMD wave64/wave32 - wave_lock_serialize([&]() { - while (true) - { - if (result == operation_status::failed_collision) - { - result = try_erase(key); - } - else + wave_lock_serialize( + [&]() { - break; - } - } - }); + while (true) + { + if (result == operation_status::failed_collision) + { + result = try_erase(key); + } + else + { + break; + } + } + }); return result == operation_status::success ? 1 : 0; } diff --git a/src/stdgpu/impl/vector_detail.cuh b/src/stdgpu/impl/vector_detail.cuh index 8fa4d3465..88f64c9cf 100644 --- a/src/stdgpu/impl/vector_detail.cuh +++ b/src/stdgpu/impl/vector_detail.cuh @@ -197,30 +197,35 @@ vector::push_back(const T& element) if (0 <= push_position && push_position < capacity()) { // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize([&]() { - while (!pushed) - { - if (_locks[push_position].try_lock()) + detail::wave_lock_serialize( + [&]() { - // START --- critical section --- START - - if (!occupied(push_position)) + while (!pushed) { - allocator_traits::construct(_allocator, &(_data[push_position]), element); - bool was_occupied = _occupied.set(push_position); - pushed = true; - - if (was_occupied) + if (_locks[push_position].try_lock()) { - printf("stdgpu::vector::push_back : Expected entry to be not occupied but actually was\n"); + // START --- critical section --- START + + if (!occupied(push_position)) + { + allocator_traits::construct(_allocator, + &(_data[push_position]), + element); + bool was_occupied = _occupied.set(push_position); + pushed = true; + + if (was_occupied) + { + printf("stdgpu::vector::push_back : Expected entry to be not occupied but actually " + "was\n"); + } + } + + // END --- critical section --- END + _locks[push_position].unlock(); } } - - // END --- critical section --- END - _locks[push_position].unlock(); - } - } - }); + }); } else { @@ -253,30 +258,36 @@ vector::pop_back() if (0 <= pop_position && pop_position < capacity()) { // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize([&]() { - while (!popped.second) - { - if (_locks[pop_position].try_lock()) + detail::wave_lock_serialize( + [&]() { - // START --- critical section --- START - - if (occupied(pop_position)) + while (!popped.second) { - bool was_occupied = _occupied.reset(pop_position); - allocator_traits::construct(_allocator, &popped, _data[pop_position], true); - allocator_traits::destroy(_allocator, &(_data[pop_position])); - - if (!was_occupied) + if (_locks[pop_position].try_lock()) { - printf("stdgpu::vector::pop_back : Expected entry to be occupied but actually was not\n"); + // START --- critical section --- START + + if (occupied(pop_position)) + { + bool was_occupied = _occupied.reset(pop_position); + allocator_traits::construct(_allocator, + &popped, + _data[pop_position], + true); + allocator_traits::destroy(_allocator, &(_data[pop_position])); + + if (!was_occupied) + { + printf("stdgpu::vector::pop_back : Expected entry to be occupied but actually was " + "not\n"); + } + } + + // END --- critical section --- END + _locks[pop_position].unlock(); } } - - // END --- critical section --- END - _locks[pop_position].unlock(); - } - } - }); + }); } else { From 178879c4e22f0820bd629972b407d8bbea417637 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 22 Jun 2026 21:35:54 +0000 Subject: [PATCH 03/10] [ROCm] Fold small device reductions on the host instead of a custom kernel 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 ``` --- src/stdgpu/impl/numeric_detail.h | 129 +++++++++++-------------------- 1 file changed, 43 insertions(+), 86 deletions(-) diff --git a/src/stdgpu/impl/numeric_detail.h b/src/stdgpu/impl/numeric_detail.h index 1911373c2..5c9610789 100644 --- a/src/stdgpu/impl/numeric_detail.h +++ b/src/stdgpu/impl/numeric_detail.h @@ -1,5 +1,6 @@ /* * Copyright 2022 Patrick Stotko + * Copyright 2026 Advanced Micro Devices, Inc. * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. * You may obtain a copy of the License at @@ -26,10 +27,11 @@ #include #if STDGPU_BACKEND == STDGPU_BACKEND_HIP - #include + #include - #include - #include + #include + #include + #include #endif namespace stdgpu @@ -39,79 +41,11 @@ namespace detail { #if STDGPU_BACKEND == STDGPU_BACKEND_HIP -template -__global__ void -transform_reduce_index_kernel(IndexType size, T init, BinaryFunction reduce, UnaryFunction f, T* block_results) -{ - __shared__ T shared[block_size]; - - index_t tid = static_cast(threadIdx.x); - IndexType stride = static_cast(block_size) * static_cast(gridDim.x); - - T value = init; - for (IndexType i = static_cast(blockIdx.x) * static_cast(block_size) + tid; i < size; - i += stride) - { - value = reduce(value, f(i)); - } - shared[tid] = value; - __syncthreads(); - - for (index_t offset = block_size / 2; offset > 0; offset /= 2) - { - if (tid < offset) - { - shared[tid] = reduce(shared[tid], shared[tid + offset]); - } - __syncthreads(); - } - - if (tid == 0) - { - block_results[blockIdx.x] = shared[0]; - } -} - -// Hand-written device reduction used in place of thrust::transform_reduce for the device execution policy. -// rocThrust's single-block transform_reduce wedges its internal hipMemcpyWithStream result copy on some -// RDNA GPUs for degenerate inputs; a plain kernel plus a plain hipMemcpy D2H avoids that path while -// computing the identical reduction (same init, same associative BinaryFunction). -template -T -transform_reduce_index_device(IndexType size, T init, BinaryFunction reduce, UnaryFunction f) -{ - if (size <= 0) - { - return init; - } - - constexpr index_t block_size = 256; - constexpr index_t max_blocks = 64; - index_t blocks = static_cast((static_cast(size) + block_size - 1) / block_size); - blocks = (blocks < 1) ? 1 : ((blocks > max_blocks) ? max_blocks : blocks); - - T* block_results = nullptr; - STDGPU_HIP_SAFE_CALL(hipMalloc(&block_results, static_cast(blocks) * sizeof(T))); - - transform_reduce_index_kernel - <<(blocks), block_size>>>(size, init, reduce, f, block_results); - STDGPU_HIP_SAFE_CALL(hipGetLastError()); - STDGPU_HIP_SAFE_CALL(hipDeviceSynchronize()); - - T host_results[max_blocks]; - STDGPU_HIP_SAFE_CALL(hipMemcpy(host_results, - block_results, - static_cast(blocks) * sizeof(T), - hipMemcpyDeviceToHost)); - STDGPU_HIP_SAFE_CALL(hipFree(block_results)); - - T result = init; - for (index_t i = 0; i < blocks; ++i) - { - result = reduce(result, host_results[i]); - } - return result; -} +// Reductions at or below this size are folded on the host to sidestep a rocThrust +// single-block transform_reduce hang observed on some RDNA GPUs (see +// transform_reduce_index). It must exceed rocThrust's single-block cutoff so that +// larger reductions always take the unaffected multi-block path. +inline constexpr index_t transform_reduce_index_host_threshold = 4096; #endif template @@ -163,18 +97,41 @@ transform_reduce_index(ExecutionPolicy&& policy, IndexType size, T init, BinaryF #if STDGPU_BACKEND == STDGPU_BACKEND_HIP if constexpr (std::is_same_v, execution::device_policy>) { - return detail::transform_reduce_index_device(size, init, reduce, f); + if (size <= 0) + { + return init; + } + + // Small reductions hit rocThrust's single-block transform_reduce, which can + // hang in its result-copy path on some RDNA GPUs for degenerate inputs. + // Materialize the transformed values with a plain thrust::transform (an + // element-wise launch that is unaffected) and fold them on the host. + if (size <= detail::transform_reduce_index_host_threshold) + { + thrust::device_vector values(static_cast(size)); + thrust::transform(policy, + thrust::counting_iterator(0), + thrust::counting_iterator(size), + values.begin(), + f); + + thrust::host_vector host_values(values); + T result = init; + for (std::size_t i = 0; i < host_values.size(); ++i) + { + result = reduce(result, host_values[i]); + } + return result; + } } - else #endif - { - return thrust::transform_reduce(std::forward(policy), - thrust::counting_iterator(0), - thrust::counting_iterator(size), - f, - init, - reduce); - } + + return thrust::transform_reduce(std::forward(policy), + thrust::counting_iterator(0), + thrust::counting_iterator(size), + f, + init, + reduce); } } // namespace stdgpu From 236953b20feb1c27bfa95d969f2495c0fcc5552d Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 22 Jun 2026 14:53:24 -0700 Subject: [PATCH 04/10] [ROCm] Tie host-fold reduction threshold to rocPRIM single-block cutoff 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. --- src/stdgpu/impl/numeric_detail.h | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/src/stdgpu/impl/numeric_detail.h b/src/stdgpu/impl/numeric_detail.h index 5c9610789..0b4b44534 100644 --- a/src/stdgpu/impl/numeric_detail.h +++ b/src/stdgpu/impl/numeric_detail.h @@ -43,8 +43,13 @@ namespace detail #if STDGPU_BACKEND == STDGPU_BACKEND_HIP // Reductions at or below this size are folded on the host to sidestep a rocThrust // single-block transform_reduce hang observed on some RDNA GPUs (see -// transform_reduce_index). It must exceed rocThrust's single-block cutoff so that -// larger reductions always take the unaffected multi-block path. +// transform_reduce_index). rocPRIM runs the single-block reduce exactly when the +// size does not exceed its reduce-config items_per_block (block_size * +// items_per_thread); above that bound the reduce is multi-block and unaffected. +// This bound was measured with rocPRIM 4.4.0 (debug_synchronous) on gfx1201: 256 * +// 16 == 4096 for both reduced element types used here (bool and index_t). Folding +// sizes up to 4096 therefore covers every size that would take the single-block +// path, and no larger reduction is ever folded onto the host. inline constexpr index_t transform_reduce_index_host_threshold = 4096; #endif From 07529aebc23d264541b9378caa8057aa976f47dd Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 22 Jun 2026 22:13:39 +0000 Subject: [PATCH 05/10] [ROCm] Move warp-convergent execution into the backend namespace dispatch 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 ``` --- src/stdgpu/CMakeLists.txt | 1 + src/stdgpu/cuda/CMakeLists.txt | 1 + src/stdgpu/cuda/execution.cuh | 42 +++++++++++ src/stdgpu/hip/CMakeLists.txt | 1 + src/stdgpu/hip/execution.h | 54 ++++++++++++++ src/stdgpu/impl/deque_detail.cuh | 10 +-- src/stdgpu/impl/execution_detail.h | 51 +++++++++++++ src/stdgpu/impl/unordered_base_detail.cuh | 6 +- src/stdgpu/impl/vector_detail.cuh | 6 +- src/stdgpu/impl/wave_lock.h | 87 ----------------------- src/stdgpu/openmp/CMakeLists.txt | 1 + src/stdgpu/openmp/execution.h | 41 +++++++++++ 12 files changed, 203 insertions(+), 98 deletions(-) create mode 100644 src/stdgpu/cuda/execution.cuh create mode 100644 src/stdgpu/hip/execution.h create mode 100644 src/stdgpu/impl/execution_detail.h delete mode 100644 src/stdgpu/impl/wave_lock.h create mode 100644 src/stdgpu/openmp/execution.h diff --git a/src/stdgpu/CMakeLists.txt b/src/stdgpu/CMakeLists.txt index 98e19d206..479577978 100644 --- a/src/stdgpu/CMakeLists.txt +++ b/src/stdgpu/CMakeLists.txt @@ -86,6 +86,7 @@ if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.23) impl/bit_detail.h impl/bitset_detail.cuh impl/deque_detail.cuh + impl/execution_detail.h impl/functional_detail.h impl/iterator_detail.h impl/limits_detail.h diff --git a/src/stdgpu/cuda/CMakeLists.txt b/src/stdgpu/cuda/CMakeLists.txt index 08ed3c6ec..0f95d8430 100644 --- a/src/stdgpu/cuda/CMakeLists.txt +++ b/src/stdgpu/cuda/CMakeLists.txt @@ -14,6 +14,7 @@ if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.23) BASE_DIRS ${STDGPU_INCLUDE_LOCAL_DIR} FILES atomic.cuh device.h + execution.cuh memory.h platform.h platform_check.h) diff --git a/src/stdgpu/cuda/execution.cuh b/src/stdgpu/cuda/execution.cuh new file mode 100644 index 000000000..c4eb674e6 --- /dev/null +++ b/src/stdgpu/cuda/execution.cuh @@ -0,0 +1,42 @@ +/* + * Copyright 2024 Patrick Stotko + * Copyright 2026 Advanced Micro Devices, Inc. + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef STDGPU_CUDA_EXECUTION_H +#define STDGPU_CUDA_EXECUTION_H + +#include + +namespace stdgpu::cuda +{ + +/** + * \brief Runs a (spinlock-based) critical section + * \param[in] body The callable to execute + * + * NVIDIA Volta and later provide Independent Thread Scheduling, so a lane that + * acquires a lock makes forward progress while its peers spin; no per-lane + * serialization is required. + */ +template +STDGPU_DEVICE_ONLY void +warp_convergent_execute(F&& body) +{ + body(); +} + +} // namespace stdgpu::cuda + +#endif // STDGPU_CUDA_EXECUTION_H diff --git a/src/stdgpu/hip/CMakeLists.txt b/src/stdgpu/hip/CMakeLists.txt index f4e6cb5b3..59a0cb360 100644 --- a/src/stdgpu/hip/CMakeLists.txt +++ b/src/stdgpu/hip/CMakeLists.txt @@ -21,6 +21,7 @@ if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.23) BASE_DIRS ${STDGPU_INCLUDE_LOCAL_DIR} FILES atomic.h device.h + execution.h memory.h platform.h platform_check.h) diff --git a/src/stdgpu/hip/execution.h b/src/stdgpu/hip/execution.h new file mode 100644 index 000000000..c9e743e78 --- /dev/null +++ b/src/stdgpu/hip/execution.h @@ -0,0 +1,54 @@ +/* + * Copyright 2024 Patrick Stotko + * Copyright 2026 Advanced Micro Devices, Inc. + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef STDGPU_HIP_EXECUTION_H +#define STDGPU_HIP_EXECUTION_H + +#include + +namespace stdgpu::hip +{ + +/** + * \brief Runs a (spinlock-based) critical section with one wavefront lane active at a time + * \param[in] body The callable to execute + * + * AMD GPUs (CDNA wave64, RDNA wave32) provide no per-lane forward-progress + * guarantee, so a lane that acquires a lock cannot advance past reconvergence + * while its peers keep spinning -- a livelock. Electing one active lane at a + * time via ballot lets each lane run the body in turn. + */ +template +STDGPU_DEVICE_ONLY void +warp_convergent_execute(F&& body) +{ + int lane = static_cast(__lane_id()); + unsigned long long active = __ballot(1); + + while (active) + { + int elected = __ffsll(static_cast(active)) - 1; + if (lane == elected) + { + body(); + } + active &= ~(1ull << elected); + } +} + +} // namespace stdgpu::hip + +#endif // STDGPU_HIP_EXECUTION_H diff --git a/src/stdgpu/impl/deque_detail.cuh b/src/stdgpu/impl/deque_detail.cuh index 773c2f875..31b2d9a68 100644 --- a/src/stdgpu/impl/deque_detail.cuh +++ b/src/stdgpu/impl/deque_detail.cuh @@ -19,7 +19,7 @@ #include #include -#include +#include #include #include #include @@ -220,7 +220,7 @@ deque::push_back(const T& element) index_t push_position = static_cast(_end.fetch_inc_mod(static_cast(capacity()))); // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize( + detail::warp_convergent_execute( [&]() { while (!pushed) @@ -281,7 +281,7 @@ deque::pop_back() pop_position = (pop_position == 0) ? capacity() - 1 : pop_position - 1; // Manually reconstruct stored value // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize( + detail::warp_convergent_execute( [&]() { while (!popped.second) @@ -350,7 +350,7 @@ deque::push_front(const T& element) push_position = (push_position == 0) ? capacity() - 1 : push_position - 1; // Manually reconstruct stored value // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize( + detail::warp_convergent_execute( [&]() { while (!pushed) @@ -410,7 +410,7 @@ deque::pop_front() index_t pop_position = static_cast(_begin.fetch_inc_mod(static_cast(capacity()))); // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize( + detail::warp_convergent_execute( [&]() { while (!popped.second) diff --git a/src/stdgpu/impl/execution_detail.h b/src/stdgpu/impl/execution_detail.h new file mode 100644 index 000000000..e4d4d54c9 --- /dev/null +++ b/src/stdgpu/impl/execution_detail.h @@ -0,0 +1,51 @@ +/* + * Copyright 2019 Patrick Stotko + * Copyright 2026 Advanced Micro Devices, Inc. + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef STDGPU_EXECUTION_DETAIL_H +#define STDGPU_EXECUTION_DETAIL_H + +#include +#include + +#if STDGPU_BACKEND == STDGPU_BACKEND_CUDA + #include STDGPU_DETAIL_BACKEND_HEADER(execution.cuh) +#else + #include STDGPU_DETAIL_BACKEND_HEADER(execution.h) +#endif + +namespace stdgpu::detail +{ + +/** + * \brief Runs a spinlock-based critical section in a way that is safe on the active backend + * \param[in] body The callable to execute + * + * Containers that retry a try_lock in a loop livelock on AMD wavefronts, which + * have no per-lane forward-progress guarantee. The backend implementation + * (warp_convergent_execute in stdgpu::hip / stdgpu::cuda / stdgpu::openmp) + * serializes the wavefront where required and is a no-op where the hardware or + * host execution already guarantees progress. + */ +template +STDGPU_DEVICE_ONLY void +warp_convergent_execute(F&& body) +{ + stdgpu::STDGPU_BACKEND_NAMESPACE::warp_convergent_execute(stdgpu::forward(body)); +} + +} // namespace stdgpu::detail + +#endif // STDGPU_EXECUTION_DETAIL_H diff --git a/src/stdgpu/impl/unordered_base_detail.cuh b/src/stdgpu/impl/unordered_base_detail.cuh index 36aa48d55..9f645c2d0 100644 --- a/src/stdgpu/impl/unordered_base_detail.cuh +++ b/src/stdgpu/impl/unordered_base_detail.cuh @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include #include @@ -975,7 +975,7 @@ inline STDGPU_DEVICE_ONLY // Wave-serialize to avoid livelock on AMD wave64/wave32: multiple lanes // of the same wavefront contending for the same bucket spin forever // because the winner cannot make progress while waiting at reconvergence. - wave_lock_serialize( + warp_convergent_execute( [&]() { while (true) @@ -1027,7 +1027,7 @@ unordered_base::erase( operation_status result = operation_status::failed_collision; // Wave-serialize to avoid livelock on AMD wave64/wave32 - wave_lock_serialize( + warp_convergent_execute( [&]() { while (true) diff --git a/src/stdgpu/impl/vector_detail.cuh b/src/stdgpu/impl/vector_detail.cuh index 88f64c9cf..a31f2c020 100644 --- a/src/stdgpu/impl/vector_detail.cuh +++ b/src/stdgpu/impl/vector_detail.cuh @@ -20,7 +20,7 @@ #include #include -#include +#include #include #include #include @@ -197,7 +197,7 @@ vector::push_back(const T& element) if (0 <= push_position && push_position < capacity()) { // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize( + detail::warp_convergent_execute( [&]() { while (!pushed) @@ -258,7 +258,7 @@ vector::pop_back() if (0 <= pop_position && pop_position < capacity()) { // Wave-serialize to avoid livelock on AMD wave64/wave32 - detail::wave_lock_serialize( + detail::warp_convergent_execute( [&]() { while (!popped.second) diff --git a/src/stdgpu/impl/wave_lock.h b/src/stdgpu/impl/wave_lock.h deleted file mode 100644 index 64c6cb863..000000000 --- a/src/stdgpu/impl/wave_lock.h +++ /dev/null @@ -1,87 +0,0 @@ -/* - * Copyright 2019 Patrick Stotko - * Copyright 2026 Advanced Micro Devices, Inc. - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#ifndef STDGPU_WAVE_LOCK_H -#define STDGPU_WAVE_LOCK_H - -#include - -namespace stdgpu::detail -{ - -/** - * Wave-serialized locking helper for AMD CDNA/RDNA GPUs. - * - * On CUDA Volta+, Independent Thread Scheduling allows a lane that wins - * a spinlock to make forward progress while other lanes spin. On AMD GPUs - * (CDNA wave64, RDNA wave32), there is no per-lane forward-progress - * guarantee: a lane that wins a CAS is stuck at SIMT reconvergence waiting - * on losers who spin forever -- classic livelock. - * - * The fix: wave-serialize by electing one active lane at a time via ballot, - * letting only that lane attempt the lock. Other lanes wait their turn. - * - * Usage: - * wave_lock_serialize([&]() { - * // Your spinlock-based critical section - * while (!done) { - * if (lock.try_lock()) { - * // critical section - * done = true; - * lock.unlock(); - * } - * } - * }); - */ - -#if defined(__HIP_DEVICE_COMPILE__) && defined(__HIP_PLATFORM_AMD__) - -// AMD HIP: wave-serialize using ballot to avoid wave64/wave32 livelock -template -STDGPU_DEVICE_ONLY void -wave_lock_serialize(F&& body) -{ - int lane = static_cast(__lane_id()); - unsigned long long active = __ballot(1); - - while (active) - { - // Elect the lowest-numbered active lane - int elected = __ffsll(static_cast(active)) - 1; - if (lane == elected) - { - body(); - } - // Clear this lane's bit and move to the next - active &= ~(1ull << elected); - } -} - -#else - -// CUDA or other: Independent Thread Scheduling handles this, no serialization needed -template -STDGPU_DEVICE_ONLY void -wave_lock_serialize(F&& body) -{ - body(); -} - -#endif - -} // namespace stdgpu::detail - -#endif // STDGPU_WAVE_LOCK_H diff --git a/src/stdgpu/openmp/CMakeLists.txt b/src/stdgpu/openmp/CMakeLists.txt index 7002eede7..ac5c3fb4f 100644 --- a/src/stdgpu/openmp/CMakeLists.txt +++ b/src/stdgpu/openmp/CMakeLists.txt @@ -14,6 +14,7 @@ if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.23) BASE_DIRS ${STDGPU_INCLUDE_LOCAL_DIR} FILES atomic.h device.h + execution.h memory.h platform.h platform_check.h) diff --git a/src/stdgpu/openmp/execution.h b/src/stdgpu/openmp/execution.h new file mode 100644 index 000000000..473c8c1c5 --- /dev/null +++ b/src/stdgpu/openmp/execution.h @@ -0,0 +1,41 @@ +/* + * Copyright 2024 Patrick Stotko + * Copyright 2026 Advanced Micro Devices, Inc. + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef STDGPU_OPENMP_EXECUTION_H +#define STDGPU_OPENMP_EXECUTION_H + +#include + +namespace stdgpu::openmp +{ + +/** + * \brief Runs a (spinlock-based) critical section + * \param[in] body The callable to execute + * + * The OpenMP backend executes on the host without wavefronts, so there is no + * per-lane forward-progress hazard and no serialization is required. + */ +template +STDGPU_DEVICE_ONLY void +warp_convergent_execute(F&& body) +{ + body(); +} + +} // namespace stdgpu::openmp + +#endif // STDGPU_OPENMP_EXECUTION_H From 7eb8c3cc38fc113640f7f5563fa4aea4683bba7a Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 22 Jun 2026 23:24:59 +0000 Subject: [PATCH 06/10] [ROCm] Use __activemask and threadIdx for warp-convergent lane election 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 ``` --- src/stdgpu/hip/execution.h | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/src/stdgpu/hip/execution.h b/src/stdgpu/hip/execution.h index c9e743e78..9880ceb81 100644 --- a/src/stdgpu/hip/execution.h +++ b/src/stdgpu/hip/execution.h @@ -29,14 +29,18 @@ namespace stdgpu::hip * AMD GPUs (CDNA wave64, RDNA wave32) provide no per-lane forward-progress * guarantee, so a lane that acquires a lock cannot advance past reconvergence * while its peers keep spinning -- a livelock. Electing one active lane at a - * time via ballot lets each lane run the body in turn. + * time from the active-lane mask lets each lane run the body in turn. */ template STDGPU_DEVICE_ONLY void warp_convergent_execute(F&& body) { - int lane = static_cast(__lane_id()); - unsigned long long active = __ballot(1); + // Lane within the wavefront, expressed with the documented threadIdx/warpSize + // builtins: HIP numbers threads within a wavefront in row-major (x fastest) + // order, so the lane is the linear thread index modulo the wavefront size. + const unsigned int linear_thread_id = (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x; + const int lane = static_cast(linear_thread_id % static_cast(warpSize)); + unsigned long long active = __activemask(); while (active) { From b99630f35e11c38aee03e9cbf9fd3229dc2c4178 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Tue, 23 Jun 2026 04:25:53 +0000 Subject: [PATCH 07/10] [ROCm] Correct the rationale for the HIP language source markings 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. --- examples/CMakeLists.txt | 4 ++-- src/stdgpu/CMakeLists.txt | 8 ++++---- src/stdgpu/hip/CMakeLists.txt | 5 +++-- tests/stdgpu/CMakeLists.txt | 5 +++-- 4 files changed, 12 insertions(+), 10 deletions(-) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 153e370a1..c363a8da3 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -18,8 +18,8 @@ endmacro() macro(stdgpu_add_example_cpp) stdgpu_detail_add_example(${ARGV0} "cpp") - # For HIP backend, mark .cpp examples as HIP language to avoid - # flag mismatch from hip::device interface propagation. + # The examples include stdgpu's device-facing headers, which pull in + # rocThrust/rocPRIM and require the HIP compiler; compile them as HIP. if(STDGPU_BACKEND STREQUAL STDGPU_BACKEND_HIP) set_source_files_properties("${ARGV0}.cpp" PROPERTIES LANGUAGE HIP) endif() diff --git a/src/stdgpu/CMakeLists.txt b/src/stdgpu/CMakeLists.txt index 479577978..2808b016f 100644 --- a/src/stdgpu/CMakeLists.txt +++ b/src/stdgpu/CMakeLists.txt @@ -32,10 +32,10 @@ else() add_library(stdgpu STATIC) endif() -# For the HIP backend, the shared impl/*.cpp files include backend headers -# (via STDGPU_DETAIL_BACKEND_HEADER) which pull in HIP/rocPRIM headers that -# require the HIP compiler. Mark them as HIP language to avoid flag mismatch -# when linking against rocThrust (which propagates -x hip to CXX consumers). +# The shared impl/*.cpp sources include the backend headers (via +# STDGPU_DETAIL_BACKEND_HEADER), which pull in rocThrust/rocPRIM. Those headers +# use device-compiler builtins (e.g. __builtin_amdgcn_wavefrontsize) that the +# host C++ compiler cannot parse, so the sources must be compiled as HIP. if(STDGPU_BACKEND STREQUAL STDGPU_BACKEND_HIP) set_source_files_properties( ${CMAKE_CURRENT_SOURCE_DIR}/impl/device.cpp diff --git a/src/stdgpu/hip/CMakeLists.txt b/src/stdgpu/hip/CMakeLists.txt index 59a0cb360..e0add6b7f 100644 --- a/src/stdgpu/hip/CMakeLists.txt +++ b/src/stdgpu/hip/CMakeLists.txt @@ -4,11 +4,12 @@ #find_dependency(hip 7.1 REQUIRED) #" PARENT_SCOPE) -# Mark HIP backend sources as HIP language (these include HIP headers). -# Use TARGET_DIRECTORY to ensure properties apply to sources as seen by the target. target_sources(stdgpu PRIVATE impl/device.cpp impl/memory.cpp) +# These backend sources include rocThrust/rocPRIM headers that require the HIP +# compiler, so compile them as HIP. TARGET_DIRECTORY applies the property in the +# directory where the stdgpu target is defined (the parent of this one). set_source_files_properties( ${CMAKE_CURRENT_SOURCE_DIR}/impl/device.cpp ${CMAKE_CURRENT_SOURCE_DIR}/impl/memory.cpp diff --git a/tests/stdgpu/CMakeLists.txt b/tests/stdgpu/CMakeLists.txt index dace20ac3..e675f0a05 100644 --- a/tests/stdgpu/CMakeLists.txt +++ b/tests/stdgpu/CMakeLists.txt @@ -13,8 +13,9 @@ target_sources(teststdgpu PRIVATE algorithm.cpp target_sources(teststdgpu PRIVATE ../test_memory_utils.cpp) -# For HIP backend, mark .cpp test sources as HIP language to avoid -# flag mismatch from hip::device interface propagation. +# These tests include stdgpu's device-facing headers, which pull in +# rocThrust/rocPRIM. Those headers require the HIP compiler (the host C++ +# compiler cannot parse their device-compiler builtins), so compile them as HIP. if(STDGPU_BACKEND STREQUAL STDGPU_BACKEND_HIP) set_source_files_properties( main.cpp From 76d78906692a906a80890ee91107cf66fe346ec4 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 2 Jul 2026 18:39:37 +0000 Subject: [PATCH 08/10] [ROCm] Include user-facing execution.h; drop redundant comments 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. --- src/stdgpu/execution.h | 2 ++ src/stdgpu/impl/deque_detail.cuh | 6 +----- src/stdgpu/impl/execution_detail.h | 6 ------ src/stdgpu/impl/unordered_base_detail.cuh | 6 +----- src/stdgpu/impl/vector_detail.cuh | 4 +--- 5 files changed, 5 insertions(+), 19 deletions(-) diff --git a/src/stdgpu/execution.h b/src/stdgpu/execution.h index fc7dafd38..d8e4a6142 100644 --- a/src/stdgpu/execution.h +++ b/src/stdgpu/execution.h @@ -92,4 +92,6 @@ constexpr host_policy host; } // namespace stdgpu::execution +#include + #endif // STDGPU_EXECUTION_H diff --git a/src/stdgpu/impl/deque_detail.cuh b/src/stdgpu/impl/deque_detail.cuh index 31b2d9a68..657c25fa8 100644 --- a/src/stdgpu/impl/deque_detail.cuh +++ b/src/stdgpu/impl/deque_detail.cuh @@ -19,7 +19,7 @@ #include #include -#include +#include #include #include #include @@ -219,7 +219,6 @@ deque::push_back(const T& element) { index_t push_position = static_cast(_end.fetch_inc_mod(static_cast(capacity()))); - // Wave-serialize to avoid livelock on AMD wave64/wave32 detail::warp_convergent_execute( [&]() { @@ -280,7 +279,6 @@ deque::pop_back() index_t pop_position = static_cast(_end.fetch_dec_mod(static_cast(capacity()))); pop_position = (pop_position == 0) ? capacity() - 1 : pop_position - 1; // Manually reconstruct stored value - // Wave-serialize to avoid livelock on AMD wave64/wave32 detail::warp_convergent_execute( [&]() { @@ -349,7 +347,6 @@ deque::push_front(const T& element) index_t push_position = static_cast(_begin.fetch_dec_mod(static_cast(capacity()))); push_position = (push_position == 0) ? capacity() - 1 : push_position - 1; // Manually reconstruct stored value - // Wave-serialize to avoid livelock on AMD wave64/wave32 detail::warp_convergent_execute( [&]() { @@ -409,7 +406,6 @@ deque::pop_front() { index_t pop_position = static_cast(_begin.fetch_inc_mod(static_cast(capacity()))); - // Wave-serialize to avoid livelock on AMD wave64/wave32 detail::warp_convergent_execute( [&]() { diff --git a/src/stdgpu/impl/execution_detail.h b/src/stdgpu/impl/execution_detail.h index e4d4d54c9..590bdfc7a 100644 --- a/src/stdgpu/impl/execution_detail.h +++ b/src/stdgpu/impl/execution_detail.h @@ -32,12 +32,6 @@ namespace stdgpu::detail /** * \brief Runs a spinlock-based critical section in a way that is safe on the active backend * \param[in] body The callable to execute - * - * Containers that retry a try_lock in a loop livelock on AMD wavefronts, which - * have no per-lane forward-progress guarantee. The backend implementation - * (warp_convergent_execute in stdgpu::hip / stdgpu::cuda / stdgpu::openmp) - * serializes the wavefront where required and is a no-op where the hardware or - * host execution already guarantees progress. */ template STDGPU_DEVICE_ONLY void diff --git a/src/stdgpu/impl/unordered_base_detail.cuh b/src/stdgpu/impl/unordered_base_detail.cuh index 9f645c2d0..a69979a42 100644 --- a/src/stdgpu/impl/unordered_base_detail.cuh +++ b/src/stdgpu/impl/unordered_base_detail.cuh @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include #include @@ -972,9 +972,6 @@ inline STDGPU_DEVICE_ONLY { pair result(end(), operation_status::failed_collision); - // Wave-serialize to avoid livelock on AMD wave64/wave32: multiple lanes - // of the same wavefront contending for the same bucket spin forever - // because the winner cannot make progress while waiting at reconvergence. warp_convergent_execute( [&]() { @@ -1026,7 +1023,6 @@ unordered_base::erase( { operation_status result = operation_status::failed_collision; - // Wave-serialize to avoid livelock on AMD wave64/wave32 warp_convergent_execute( [&]() { diff --git a/src/stdgpu/impl/vector_detail.cuh b/src/stdgpu/impl/vector_detail.cuh index a31f2c020..0f78b8c76 100644 --- a/src/stdgpu/impl/vector_detail.cuh +++ b/src/stdgpu/impl/vector_detail.cuh @@ -20,7 +20,7 @@ #include #include -#include +#include #include #include #include @@ -196,7 +196,6 @@ vector::push_back(const T& element) // Check position if (0 <= push_position && push_position < capacity()) { - // Wave-serialize to avoid livelock on AMD wave64/wave32 detail::warp_convergent_execute( [&]() { @@ -257,7 +256,6 @@ vector::pop_back() // Check position if (0 <= pop_position && pop_position < capacity()) { - // Wave-serialize to avoid livelock on AMD wave64/wave32 detail::warp_convergent_execute( [&]() { From cf3edae93bd8a0a309d3dfee5175a5ecb053487a Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 2 Jul 2026 19:23:55 +0000 Subject: [PATCH 09/10] [ROCm] Guard execution_detail.h include for device compilers only The 72687fe commit added #include 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 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 --- src/stdgpu/execution.h | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/src/stdgpu/execution.h b/src/stdgpu/execution.h index d8e4a6142..85deb8850 100644 --- a/src/stdgpu/execution.h +++ b/src/stdgpu/execution.h @@ -23,6 +23,7 @@ #include +#include #include /** @@ -92,6 +93,13 @@ constexpr host_policy host; } // namespace stdgpu::execution -#include +// execution_detail.h declares device-qualified functions. For the CUDA backend, +// g++ compiles plain .cpp translation units without __CUDACC__ defined, so it +// cannot see __device__-qualified declarations. Include execution_detail.h only +// when a device compiler is active (__CUDACC__ for nvcc/clang-cuda, __HIP__ for +// clang++/HIP), or when using the OpenMP backend (which has no device qualifiers). +#if defined(__CUDACC__) || defined(__HIP__) || STDGPU_BACKEND == STDGPU_BACKEND_OPENMP + #include +#endif #endif // STDGPU_EXECUTION_H From 5e660960902821dbcb3b370b3e0cd2c66959bdea Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 6 Jul 2026 03:20:59 +0000 Subject: [PATCH 10/10] [ROCm] Use IS_DEVICE_COMPILED guard; rename execution_detail.cuh 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 --- src/stdgpu/CMakeLists.txt | 2 +- src/stdgpu/execution.h | 12 +++++------- .../{execution_detail.h => execution_detail.cuh} | 0 src/stdgpu/impl/unordered_base_detail.cuh | 2 +- 4 files changed, 7 insertions(+), 9 deletions(-) rename src/stdgpu/impl/{execution_detail.h => execution_detail.cuh} (100%) diff --git a/src/stdgpu/CMakeLists.txt b/src/stdgpu/CMakeLists.txt index 2808b016f..937b60d1f 100644 --- a/src/stdgpu/CMakeLists.txt +++ b/src/stdgpu/CMakeLists.txt @@ -86,7 +86,7 @@ if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.23) impl/bit_detail.h impl/bitset_detail.cuh impl/deque_detail.cuh - impl/execution_detail.h + impl/execution_detail.cuh impl/functional_detail.h impl/iterator_detail.h impl/limits_detail.h diff --git a/src/stdgpu/execution.h b/src/stdgpu/execution.h index 85deb8850..2ae40dd4d 100644 --- a/src/stdgpu/execution.h +++ b/src/stdgpu/execution.h @@ -93,13 +93,11 @@ constexpr host_policy host; } // namespace stdgpu::execution -// execution_detail.h declares device-qualified functions. For the CUDA backend, -// g++ compiles plain .cpp translation units without __CUDACC__ defined, so it -// cannot see __device__-qualified declarations. Include execution_detail.h only -// when a device compiler is active (__CUDACC__ for nvcc/clang-cuda, __HIP__ for -// clang++/HIP), or when using the OpenMP backend (which has no device qualifiers). -#if defined(__CUDACC__) || defined(__HIP__) || STDGPU_BACKEND == STDGPU_BACKEND_OPENMP - #include +// execution_detail.cuh declares device-qualified functions, so include it only +// when the device compiler is active (host-only translation units, e.g. plain +// .cpp files in the CUDA backend, must not see the __device__ declarations). +#if STDGPU_DETAIL_IS_DEVICE_COMPILED + #include #endif #endif // STDGPU_EXECUTION_H diff --git a/src/stdgpu/impl/execution_detail.h b/src/stdgpu/impl/execution_detail.cuh similarity index 100% rename from src/stdgpu/impl/execution_detail.h rename to src/stdgpu/impl/execution_detail.cuh diff --git a/src/stdgpu/impl/unordered_base_detail.cuh b/src/stdgpu/impl/unordered_base_detail.cuh index a69979a42..f7e6d5c79 100644 --- a/src/stdgpu/impl/unordered_base_detail.cuh +++ b/src/stdgpu/impl/unordered_base_detail.cuh @@ -22,8 +22,8 @@ #include #include #include -#include #include +#include #include #include #include