Synchronization Primitives

Synchronization Primitives#

2026-10-01

3 min read time

Applies to Linux and Windows

This page covers the synchronization primitives extended API, including thread-scoped atomics, latches, barriers, semaphores, and pipelines for coordinating host and device threads.

Atomics

cuda::atomic	System-wide std::atomic objects and operations	libhipcxx 1.0.0 / CCCL 2.0.0 / CUDA 10.2
cuda::atomic_ref	System-wide std::atomic_ref objects and operations	libhipcxx 1.7.0 / CCCL 2.0.0 / CUDA 11.6

Latches

cuda::latch

System-wide std::latch single-phase asynchronous thread coordination mechanism

libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0

Barriers

cuda::barrier

System wide std::barrier multi-phase asynchronous thread coordination mechanism

libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0

Semaphores

cuda::counting_semaphore	System wide std::counting_semaphore primitive for constraining concurrent access	libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0
cuda::binary_semaphore	System wide std::binary_semaphore primitive for mutual exclusion	libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0

Pipelines

The pipeline library is included in the CUDA Toolkit, but is not part of the open source libhipcxx distribution.

cuda::pipeline	Coordination mechanism for sequencing asynchronous operations	libhipcxx 1.2.0 / CCCL 2.0.0 / CUDA 11.1
cuda::pipeline_shared_state	cuda::pipeline shared state object	libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0
cuda::pipeline_role	Defines producer/consumer role for a thread participating in a pipeline	libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0
cuda::make_pipeline	Creates a cuda::pipeline	libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0
cuda::pipeline_consumer_wait_prior	Blocks the current thread until all operations committed up to a prior pipeline stage complete	libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0
cuda::pipeline_producer_commit	Binds operations previously issued by the current thread to a cuda::barrier	libhipcxx 1.1.0 / CCCL 2.0.0 / CUDA 11.0