The CUDA Handbook ▸ Chapter 5. Memory

5.10 The CUDA C++ Standard Library

Section 5.9 introduced part of libcu++, the CUDA C++ Standard Library, when it wrapped cp.async in cuda::memcpy_async(), cuda::pipeline, and cuda::barrier. The library has two layers. The first, in the cuda::std:: namespace, is a port of the ISO C++ standard library that compiles for both the host and the device: <cuda/std/atomic>, <cuda/std/type_traits>, <cuda/std/tuple>, <cuda/std/array>, <cuda/std/mdspan>, <cuda/std/complex>, and more, each the standard facility of the same name with its host-only dependencies removed. The second, in the cuda:: namespace, holds extensions with no standard equivalent, present to expose GPU hardware: the asynchronous primitives of Section 5.9 and the scoped atomics below.

The cuda::std:: layer lets one definition serve both sides of the launch. A function template that needs std::tuple, a fixed-width integer, or a compile-time type trait can include the corresponding cuda/std header and compile unchanged for the device, rather than reimplementing the utility for device code. The facilities that require an operating system underneath – iostreams, std::thread, the filesystem – are absent, because a kernel has none to call; what remains is the portable, computational core of the standard library.

Scoped atomics are the extension that matters most for performance, and the reason is the memory model. A std::atomic<T>, and its device-capable form cuda::std::atomic<T>, orders its operations with respect to every thread in the system, as the ISO model requires. On a GPU, that guarantee is expensive: making a value visible beyond the block, to the L2 cache and the rest of the device, or beyond the device entirely, across the link to the CPU, costs coherence traffic that a counter confined to one block never needs.

The scoped types name the threads that must agree. cuda::atomic<T, Scope> and cuda::atomic_ref<T, Scope> take a thread scope – thread_scope_block, thread_scope_device, or thread_scope_system – declaring which threads must observe the operation as atomic and in order, and the hardware supplies only the coherence that scope demands: a block-scoped atomic on a shared-memory location is ordered within the block and no further, skipping the traffic a device- or system-scoped operation would incur. atomic_ref adds the atomicity without owning the storage; it binds to an ordinary object, so a plain array element or a struct field can be operated on atomically in place. The memory orders are the ISO ones – memory_order_relaxed, memory_order_acquire, and memory_order_release among them – passed to .fetch_add(), .load(), and .store(). Scope and order together are the C++ surface of the CUDA memory consistency model: the scoped acquire/release model that the PTX fence and atomic instructions implement, and that the atomicAdd_block() and atomicAdd_system() intrinsic suffixes of Section 5.2.11 select by hand.

#include <cuda/atomic>

// Each block counts matches in shared memory; that counter need only be
// ordered among the block's own threads, so it is block-scoped. One
// device-scoped add then folds each block's subtotal into the grid total.
__global__ void
countMatches( const int32_t *data, int32_t n, int32_t key, int32_t *total )
{
    __shared__ int32_t hits;
    if ( threadIdx.x == 0 ) {
        hits = 0;
    }
    __syncthreads();

    cuda::atomic_ref<int32_t, cuda::thread_scope_block> blockHits( hits );
    for ( int32_t i = blockIdx.x*blockDim.x + threadIdx.x;
                  i < n;
                  i += blockDim.x*gridDim.x ) {
        if ( data[i] == key ) {
            blockHits.fetch_add( 1, cuda::memory_order_relaxed );
        }
    }
    __syncthreads();

    if ( threadIdx.x == 0 ) {
        cuda::atomic_ref<int32_t, cuda::thread_scope_device> grandTotal( *total );
        grandTotal.fetch_add( hits, cuda::memory_order_relaxed );
    }
}

The scope’s effect on throughput is easy to measure. The scopedAtomicSpeed.cu microbench has every block increment one counter with relaxed fetch_add under full contention – block scope on a counter in shared memory, device and system scope on one in global memory, with a separate counter per block so the number of contending threads is the same in each run. On an RTX 3060 (Table 5-14), the block-scoped atomic runs about nine times faster than the global one, because it resolves on-chip instead of at the L2 cache. Device and system scope run at the same speed: on a discrete GPU whose device memory no other agent is touching, widening the scope broadens the ordering guarantee, not the cost.

Table 5-14. Relaxed fetch_add throughput by thread scope on a GeForce RTX 3060 (896 blocks of 256 threads, one counter per block, best of five).

The scope also reaches outward. A thread_scope_system atomic is ordered against the CPU and other devices, so a kernel and host code – or two peer GPUs – can share a flag or a lock in the coherent memory of a platform like Grace Hopper (Section 5.7.1). The plain intrinsics of Section 5.2.11 cannot serve that case, because they order only within the device.