Skip to main content Link Menu Expand (external link) Document Search Copy Copied

cuda::atomic::atomic_thread_fence

Defined in header <cuda/atomic>:

__host__ __device__
void cuda::atomic_thread_fence(cuda::std::memory_order order,
                               cuda::thread_scope scope = cuda::thread_scope_system);

Establishes memory synchronization ordering of non-atomic and relaxed atomic accesses, as instructed by order, for all threads within scope without an associated atomic operation. It has the same semantics as cuda::std::atomic_thread_fence.

Example

#include <cuda/atomic>

__global__ void example_kernel(int* data) {
  *data = 42;
  cuda::atomic_thread_fence(cuda::std::memory_order_release,
                            cuda::thread_scope_device);
}

See it on Godbolt