binary_semaphore.md

grand_parent	parent	nav_order
Extended API	Synchronization Primitives	5

cuda::binary_semaphore

Defined in header <cuda/semaphore>:

namespace cuda {

template <cuda::thread_scope Scope>
using binary_semaphore = cuda::std::counting_semaphore<Scope, 1>;

}

The class template cuda::binary_semaphore is an extended form of cuda::std::binary_semaphore that takes an additional cuda::thread_scope argument. cuda::binary_semaphore has the same interface and semantics as cuda::std::binary_semaphore, but cuda::binary_semaphore is a class template.

Concurrency Restrictions

An object of type cuda::binary_semaphore or cuda::std::binary_semaphore, shall not be accessed concurrently by CPU and GPU threads unless:

• it is in unified memory and the concurrentManagedAccess property is 1, or
• it is in CPU memory and the hostNativeAtomicSupported property is 1.

Note, for objects of scopes other than cuda::thread_scope_system this is a data-race, and thefore also prohibited regardless of memory characteristics.

Under CUDA Compute Capability 6 (Pascal) or prior, an object of type cuda::binary_semaphore or cuda::std::binary_semaphore may not be used.

Implementation-Defined Behavior

For each cuda::thread_scope S, cuda::binary_semaphore<S>::max() is as follows:

cuda::thread_scope S	cuda::binary_semaphore<S>::max()
Any	1

Example

#include <cuda/semaphore>

__global__ void example_kernel() {
  // This semaphore is suitable for all threads in the system.
  cuda::binary_semaphore<cuda::thread_scope_system> a;

  // This semaphore has the same type as the previous one (`a`).
  cuda::std::binary_semaphore<> b;

  // This semaphore is suitable for all threads on the current processor (e.g. GPU).
  cuda::binary_semaphore<cuda::thread_scope_device> c;

  // This semaphore is suitable for all threads in the same thread block.
  cuda::binary_semaphore<cuda::thread_scope_block> d;
}

See it on Godbolt{: .btn }