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 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