feat(semaphore): support <cuda/std/semaphore> and <cuda/semaphore> on HIP/AMD

.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

.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
 

.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
 

.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
 

.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
 

README.rst

@@ -166,17 +166,16 @@ 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.
+- `<cuda/std/semaphore>` and `<cuda/semaphore>` 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:
 
 =========================================  ======================  ========================================================================================
 Group                                      API                     Header  Description
 =========================================  ======================  ========================================================================================
 Synchronization Library                    `<cuda/std/latch>`      Single-phase asynchronous thread-coordination mechanism
 Synchronization Library                    `<cuda/std/barrier>`    Multi-phase asynchronous thread-coordination mechanism
-Synchronization Library                    `<cuda/std/semaphore>`  Primitives for constraining concurrent access
 Extended Synchronization Library           `<cuda/latch>`          System-wide `cuda::std::latch` single-phase asynchronous thread coordination mechanism.
 Extended Synchronization Library           `<cuda/barrier>`        System-wide `cuda::std::barrier` multi-phase asynchronous thread coordination mechanism.
-Extended Synchronization Library           `<cuda/semaphore>`      System-wide primitives for constraining concurrent access.
 Extended Synchronization Library           `<cuda/pipeline>`       Coordination mechanisms to sequence asynchronous operations.
 Extended Memory Access Properties Library  `<cuda/annotated_ptr>`  Memory access properties for pointers.
 PTX API                                    `<cuda/ptx>`            The `cuda::ptx` namespace contains functions that map to Nvidia PTX instructions. 

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

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