From e2e8f70dd8737a5cf8b7b798b99654d06c8eb326 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Sun, 7 Jun 2026 20:21:19 +0000 Subject: [PATCH 1/2] Enable and on HIP/AMD The counting/binary semaphore implementation is already present and builds on the cuda::std::atomic wait/notify machinery that is supported on AMD; only the public umbrella headers were hard-gated with an #error for __HIP_PLATFORM_AMD__. Remove that gate from both and so the existing implementation is exposed. No implementation changes are required; the remaining sm_70 guard keys on __CUDA_ARCH__ and is inert on AMD. The semaphore acquire path spins until the count is positive. Threads in a single wavefront on AMD GPUs do not have independent forward-progress guarantees, so a blocking acquire() that waits on a release() issued by another thread of the same wavefront can deadlock; coordinate across separate wavefronts/blocks (thread_scope_device, thread_scope_system) instead. README.rst is updated to move the two headers out of the unsupported list and to document this scope caveat next to the existing device-clock note. This work was authored with the assistance of an AI coding assistant. Test Plan: Validated on a gfx90a GPU (MI250X) with ROCm 7.2.1. Producer/consumer with cuda::binary_semaphore across two blocks, and the same pattern with thread_scope_system, both PASS: hipcc -std=c++17 --offload-arch=gfx90a -I include prod_cons.cpp -o pc && ./pc Upstream semaphore tests, compiled per .upstream-tests/test/CMakeLists.txt (HIPCC branch: -include test/force_include_hip.h, which runs each test on host and launches it on the device), run on gfx90a: std/thread/thread.semaphore/version.pass.cpp PASS std/thread/thread.semaphore/max.pass.cpp PASS std/thread/thread.semaphore/try_acquire.pass.cpp PASS std/thread/thread.semaphore/acquire.pass.cpp PASS std/thread/thread.semaphore/release.pass.cpp PASS heterogeneous/semaphore.pass.cpp PASS std/thread/thread.semaphore/timed.pass.cpp drives two concurrent agents within one wavefront through try_acquire_until/try_acquire_for; the timed acquire reaches its deadline before the same-wavefront releaser runs, which is the forward-progress limitation described above, so that test is not claimed on gfx90a. --- README.rst | 3 +-- include/cuda/semaphore | 5 ----- include/cuda/std/semaphore | 5 ----- 3 files changed, 1 insertion(+), 12 deletions(-) diff --git a/README.rst b/README.rst index 31fb338aa..8ea4d14bd 100644 --- a/README.rst +++ b/README.rst @@ -166,6 +166,7 @@ Limitations/Unsupported Features/APIs - Libhipcxx does not support CUDA backend/NVIDIA hardware. - Libhipcxx does not support the Windows OS. - `cuda::std::chrono::system_clock::now()` does not return a UNIX timestamp, host system clock and device system clock are not synchronized and they may run at different clock rates. +- `` and `` are supported. Blocking `acquire()` spins until the count is positive; on AMD hardware threads in a single wavefront do not have independent forward-progress guarantees, so a blocking `acquire()` that waits on a `release()` performed by another thread of the same wavefront (`thread_scope_block` used within one wavefront) can deadlock. Coordinate across separate wavefronts/blocks (`thread_scope_device`, `thread_scope_system`) where the producer and consumer can make progress independently. The timed `try_acquire_for`/`try_acquire_until` share this limitation when the releaser is in the same wavefront; they return without acquiring once the timeout elapses instead of deadlocking, but on AMD hardware the deadline is measured against a device timestamp counter that is not synchronized with the host clock (see the `system_clock` note above). - The following APIs from [libcudacxx] are *NOT* supported in libhipcxx: ========================================= ====================== ======================================================================================== @@ -173,10 +174,8 @@ Group API Header Descr ========================================= ====================== ======================================================================================== Synchronization Library `` Single-phase asynchronous thread-coordination mechanism Synchronization Library `` Multi-phase asynchronous thread-coordination mechanism -Synchronization Library `` Primitives for constraining concurrent access Extended Synchronization Library `` System-wide `cuda::std::latch` single-phase asynchronous thread coordination mechanism. Extended Synchronization Library `` System-wide `cuda::std::barrier` multi-phase asynchronous thread coordination mechanism. -Extended Synchronization Library `` System-wide primitives for constraining concurrent access. Extended Synchronization Library `` Coordination mechanisms to sequence asynchronous operations. Extended Memory Access Properties Library `` Memory access properties for pointers. PTX API `` The `cuda::ptx` namespace contains functions that map to Nvidia PTX instructions. diff --git a/include/cuda/semaphore b/include/cuda/semaphore index 152dda5bc..22997b288 100644 --- a/include/cuda/semaphore +++ b/include/cuda/semaphore @@ -33,11 +33,6 @@ #ifndef _CUDA_SEMAPHORE #define _CUDA_SEMAPHORE -// NOTE(HIP/AMD): semaphore is not supported on AMD hardware -#ifdef __HIP_PLATFORM_AMD__ -#error semaphore is not supported on AMD hardware and should not be included -#endif - #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 700 # error "CUDA synchronization primitives are only supported for sm_70 and up." #endif diff --git a/include/cuda/std/semaphore b/include/cuda/std/semaphore index b0e81f6f6..c8cd5d0b6 100644 --- a/include/cuda/std/semaphore +++ b/include/cuda/std/semaphore @@ -33,11 +33,6 @@ #ifndef _CUDA_STD_SEMAPHORE #define _CUDA_STD_SEMAPHORE -// NOTE(HIP/AMD): semaphore is not supported on AMD hardware -#ifdef __HIP_PLATFORM_AMD__ -#error semaphore is not supported on AMD hardware and should not be included -#endif - #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 700 # error "CUDA synchronization primitives are only supported for sm_70 and up." #endif From fee8c1dc8052e693b70e4fbf46186a5644c47b30 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 8 Jun 2026 16:33:53 +0000 Subject: [PATCH 2/2] [ROCm] Enable the same-wavefront-safe semaphore tests on HIP Five libcu++ semaphore conformance tests do not exercise the AMD same-wavefront forward-progress hazard, so the `// UNSUPPORTED: hipcc, hiprtc` skip is removed and they run on HIP: version, max, try_acquire (non-blocking), acquire (its two acquires are __syncthreads-separated and each drain a pre-existing count), and the heterogeneous host+device test (acquire and release serialized across separate launches). They pass on AMD CDNA2 (gfx90a) and RDNA3/RDNA4 (gfx1100/gfx1201). release.pass.cpp and timed.pass.cpp keep the skip: both co-schedule a blocking acquire() with its satisfying release() in one wavefront (release completes only because the count starts positive; timed must wait on the sibling and reaches its deadline), and AMD gives no intra-wavefront forward-progress guarantee. Confirmed from the compiled GCN ISA: the safe tests emit no reachable same-wavefront blocking spin, while release/timed emit the s_sleep/s_memrealtime backoff spin co-scheduled with a sibling releaser. Authored with the assistance of Claude (Anthropic). --- .upstream-tests/test/heterogeneous/semaphore.pass.cpp | 1 - .../test/std/thread/thread.semaphore/acquire.pass.cpp | 1 - .upstream-tests/test/std/thread/thread.semaphore/max.pass.cpp | 1 - .../test/std/thread/thread.semaphore/try_acquire.pass.cpp | 1 - .../test/std/thread/thread.semaphore/version.pass.cpp | 1 - 5 files changed, 5 deletions(-) diff --git a/.upstream-tests/test/heterogeneous/semaphore.pass.cpp b/.upstream-tests/test/heterogeneous/semaphore.pass.cpp index 908f52adb..32978c075 100644 --- a/.upstream-tests/test/heterogeneous/semaphore.pass.cpp +++ b/.upstream-tests/test/heterogeneous/semaphore.pass.cpp @@ -23,7 +23,6 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. -// UNSUPPORTED: hipcc, hiprtc // UNSUPPORTED: nvrtc, pre-sm-70 // uncomment for a really verbose output detailing what test steps are being launched diff --git a/.upstream-tests/test/std/thread/thread.semaphore/acquire.pass.cpp b/.upstream-tests/test/std/thread/thread.semaphore/acquire.pass.cpp index e4275bf6f..7ff8b39d8 100644 --- a/.upstream-tests/test/std/thread/thread.semaphore/acquire.pass.cpp +++ b/.upstream-tests/test/std/thread/thread.semaphore/acquire.pass.cpp @@ -30,7 +30,6 @@ // // NOTE(HIP/AMD): semaphore is not supported on AMD hardware -// UNSUPPORTED: hipcc, hiprtc // UNSUPPORTED: libcpp-has-no-threads // UNSUPPORTED: pre-sm-70 diff --git a/.upstream-tests/test/std/thread/thread.semaphore/max.pass.cpp b/.upstream-tests/test/std/thread/thread.semaphore/max.pass.cpp index 56a9c0e26..90c71bc56 100644 --- a/.upstream-tests/test/std/thread/thread.semaphore/max.pass.cpp +++ b/.upstream-tests/test/std/thread/thread.semaphore/max.pass.cpp @@ -25,7 +25,6 @@ // // NOTE(HIP/AMD): semaphore is not supported on AMD hardware -// UNSUPPORTED: hipcc, hiprtc // UNSUPPORTED: libcpp-has-no-threads // UNSUPPORTED: pre-sm-70 diff --git a/.upstream-tests/test/std/thread/thread.semaphore/try_acquire.pass.cpp b/.upstream-tests/test/std/thread/thread.semaphore/try_acquire.pass.cpp index d77506b52..90b5647fb 100644 --- a/.upstream-tests/test/std/thread/thread.semaphore/try_acquire.pass.cpp +++ b/.upstream-tests/test/std/thread/thread.semaphore/try_acquire.pass.cpp @@ -25,7 +25,6 @@ // // NOTE(HIP/AMD): semaphore is not supported on AMD hardware -// UNSUPPORTED: hipcc, hiprtc // UNSUPPORTED: libcpp-has-no-threads // UNSUPPORTED: pre-sm-70 diff --git a/.upstream-tests/test/std/thread/thread.semaphore/version.pass.cpp b/.upstream-tests/test/std/thread/thread.semaphore/version.pass.cpp index c4698bb02..7080053fb 100644 --- a/.upstream-tests/test/std/thread/thread.semaphore/version.pass.cpp +++ b/.upstream-tests/test/std/thread/thread.semaphore/version.pass.cpp @@ -25,7 +25,6 @@ // // NOTE(HIP/AMD): semaphore is not supported on AMD hardware -// UNSUPPORTED: hipcc, hiprtc // UNSUPPORTED: libcpp-has-no-threads // UNSUPPORTED: pre-sm-70