__shared__ cuda::pipeline_shared_state< cuda::thread_scope::thread_scope_block, stages_count > shared_state;

producer_acquire

consumer_release

consumer_wait

producer_commit

7.27.6.6. Keep Commit and Arrive-On Operations Converged

7.27.6.5. Warp Entanglement - Arrive-On

7.27.6.4. Warp Entanglement - Wait

7.27.6.3. Warp Entanglement - Commit

Trivially copyable

On devices with compute capability 8.0, the cp.async family of instructions allows copying data from global to shared memory asynchronously. These instructions support copying 4, 8, and 16 bytes at a time. If the size provided to memcpy_async is a multiple of 4, 8, or 16, and both pointers passed to memcpy_async are aligned to a 4, 8, or 16 alignment boundary, then memcpy_async can be implemented using exclusively asynchronous memory operations.

memcpy_async

Additionally for achieving best performance when using memcpy_async API, an alignment of 128 Bytes for both shared memory and global memory is required.

the pipeline mechanism is shared among CUDA threads in the same CUDA warp. This sharing causes batches of memcpy_async to be entangled within a warp

barrier.arrive_and_wait(); // Waits for all copies to complete

__shared__ cuda::barrier<cuda::thread_scope::thread_scope_block> barrier; if (block.thread_rank() == 0) { init(&barrier, block.size()); // Friend function initializes barrier } block.sync();

incrementing the expected count of the current phase on creation, and decrementing it on completion of the copy operation, such that the phase of the barrier will only advance when all threads participating in the barrier have arrived, and all memcpy_async bound to the current phase of the barrier have completed.

cooperative_groups::memcpy_async

On devices with compute capability 8.0 or higher, memcpy_async transfers from global to shared memory can benefit from hardware acceleration, which avoids transfering the data through an intermediate register.

Until the copy operation completes, modifying the global data or reading or writing the shared data introduces a data race.

block.sync(); // Wait for compute using shared memory to finish

block.sync(); // Wait for all copies to complete

each thread block needs to synchronize after the shared[local_idx] = global[global_idx] assignment, to ensure all writes to shared memory have completed before the compute phase can begin. The thread block also needs to synchronize again after the compute phase, to prevent overwriting shared memory before all threads have completed their computations. This pattern is illustrated in the following code snippet.

With memcpy_async improves the previous example by introducing the memcpy_async and the cuda::memcpy_async APIs to directly copy data from global to shared memory without using intermediate registers.

Without memcpy_async introduces an example that does not overlap computation with data movement and uses an intermediate register to copy data.

memcpy_async API

Asynchronous