The CUDA Handbook ▸ Chapter 8. Streaming Multiprocessors

8.1 Memory

8.1.1 Registers

Each SM contains thousands of 32-bit registers, which are allocated to threads as specified when the kernel is launched. Registers are both the fastest and most plentiful memory in the SM. As an example, the Kepler-class (SM 3.0) SMX contains 65,536 registers or 256K, while the texture cache is only 48K.

CUDA registers can contain integer or floating-point data; for hardware capable of performing double-precision arithmetic (SM 1.3 and higher), the operands are contained in even-valued register pairs. On SM 2.0 and higher hardware, register pairs also can hold 64-bit addresses.

CUDA hardware also supports wider memory transactions: the built-in int2/float2 and int4/float4 data types, residing in aligned register pairs or quads respectively, may be read or written in single 64- or 128-bit-wide loads or stores. Once in registers, the individual data elements can be referenced as .x/.y (for int2/float2) or .x/.y/.z/.w (for int4/float4).

Developers can cause nvcc to report the number of registers used by a kernel by specifying the command-line option --ptxas-options —verbose. The number of registers used by a kernel affects the number of threads that can fit in an SM, and often must be tuned carefully for optimal performance.

The maximum number of registers used for a compilation may be specified with:

--ptxas-options --maxregcount N

Register Aliasing

Because registers can hold floating-point or integer data, some intrinsics serve only to coerce the compiler into changing its view of a variable. The __int_as_float() and __float_as_int() intrinsics cause a variable to “change personalities” between 32-bit integer and single-precision floating point:

float __int_as_float( int i );
int __float_as_int( float f );

The __double2loint(), __double2hiint(), and __hiloint2double() intrinsics similarly cause registers to change personality (usually in-place). __double_as_longlong() and __longlong_as_double() coerce register pairs in-place; __double2loint() and __double2hiint() return the least- and most-significant 32 bits of the input operand, respectively; and __hiloint2double() constructs a double out of the high and low halves.

int double2loint( double d );
int double2hiint( double d );
int hiloint2double( int hi, int lo );
double long_as_double(long long int i );
long long int __double_as_longlong( double d );

8.1.2 Local Memory

Local memory is used to spill registers, and also to hold local variables that are indexed and whose indices cannot be computed at compile time. Local memory is backed by the same pool of device memory as global memory, so exhibits the same latency characteristics and benefits from the L1 and L2 cache hierarchy on Fermi and later hardware. Local memory is addressed in such a way that the memory transactions are automatically coalesced.

The hardware includes special instructions to load and store local memory: the SASS variants are LLD/LST for Tesla and LDL/STL for Fermi and Kepler.

8.1.3 Global Memory

The SMs can read or write global memory using GLD/GST instructions (on Tesla) and LD/ST instructions (on Fermi and Kepler). Developers can use standard C operators to compute and dereference addresses, including pointer arithmetic and the dereferencing operators *, [], and ->. Operating on 64- or 128-bit built-in data types (int2/float2/int4/float4) automatically causes the compiler to issue 64- or 128-bit load and store instructions.

Maximum memory performance is achieved through coalescing of memory transactions, described in Section 5.2.9.

Tesla-class hardware (SM 1.x) uses special address registers to hold pointers; later hardware implements a load/store architecture that uses the same register file for pointers, integer and floating point values and the same address space for constant memory, shared memory and global memory¹.

Fermi-class hardware includes several features not available on older hardware:

• 64-bit addressing is supported via “wide” load/store instructions in which addresses are held in even-numbered register pairs. 64-bit addressing is not supported on 32-bit host platforms; on 64-bit host platforms, 64-bit addressing is enabled automatically. As a result, code generated for the same kernels compiled for 32- and 64-bit host platforms may have different register counts and performance.

• The L1 cache may be configured to be 16K or 48K in size² (Kepler added the ability to split the cache as 32K L1/32K shared). Load instructions can include cacheability hints (to tell the hardware to pull the read into L1, or bypass the L1 and keep the data only in L2). These may be accessed via inline PTX, or through the command line option –X ptxas –dlcm=ca (cache in L1 and L2, the default setting) or –X ptxas –dlcm=cg (cache only in L2).

Atomic operations (or just “atomics”) update a memory location in a way that works correctly even when multiple GPU threads are operating on the same memory location. The hardware enforces mutual exclusion on the memory location for the duration of the operation. Since the order of operations is not guaranteed, the operators supported generally are associative³.

Atomics first became available for global memory for SM 1.1 and greater, and for shared memory for SM 1.2 and greater. Until the Kepler generation of hardware, however, global memory atomics were too slow to be useful

The global atomic intrinsics, summarized in Table 8-2, become automatically available when the appropriate architecture is specified to nvcc via --gpu-architecture.

Table 8-2. Atomic operations

All of these intrinsics can operate on 32-bit integers. 64-bit support for atomicAdd(), atomicExch(), and atomicCAS() was added in SM 1.2. atomicAdd() of 32-bit floating point values (float) was added in SM 2.0. 64-bit support for atomicMin(), atomicMax(), atomicAnd(), atomicOr(), and atomicXor() was added in SM 3.5.

Because atomic operations are implemented using hardware in the GPU’s integrated memory controller, they do not work across the PCI Express bus; so they do not work correctly on device memory pointers that correspond to host memory or peer memory.

At the hardware level, atomics come in two forms: atomic operations that return the value that was at the specified memory location before the operator was performed, and reduction operations that the developer can “fire and forget” at the memory location, ignoring the return value. Since the hardware can perform the operation more efficiently if there is no need to return the old value, the compiler detects whether the return value is used and, if it is not, emits different instructions. In SM 2.0, for example, the instructions are called ATOM and RED, respectively.

8.1.4 Constant Memory

Constant memory resides in device memory, but is backed by a different, read-only cache that is optimized to broadcast the results of read requests to threads that all reference the same memory location. Each SM contains a small, latency-optimized cache for purposes of servicing these read requests; making the memory (and the cache) read-only simplifies cache management, since the hardware has no need to implement write-back policies to deal with memory that has been updated.

SM 2.x and subsequent hardware includes a special optimization for memory that is not denoted as constant, but which the compiler has identified as 1) read-only, and 2) whose address is not dependent on the block or thread ID. The “load uniform” (LDU) instruction reads memory using the constant cache hierarchy, and broadcasts the data to the threads.

8.1.5 Shared Memory

Shared memory is very fast, on-chip memory in the SM that threads can use for data interchange within a thread block. Since it is a per-SM resource, shared memory usage can affect occupancy, the number of blocks and threads that the SM can keep resident.

SMs load and store shared memory with special instructions: G2R/R2G on SM 1.x, and LDS/STS on SM 2.x and later.

Shared memory is arranged as interleaved banks, and generally is optimized for 32-bit access. If more than one thread in a warp references the same bank, a bank conflict occurs, and the hardware must handle memory requests consecutively until all requests have been serviced. Typically, to avoid bank conflicts, applications access shared memory with an interleaved pattern based on the thread ID, e.g.:

extern __shared__ float shared[];
float data = shared[BaseIndex + threadIdx.x];

Having all threads in a warp read from the same 32-bit shared memory location also is fast: the hardware includes a broadcast mechanism to optimize for this case.

Writes to the same bank are serialized by the hardware, reducing performance. Writes to the same address cause race conditions, and should be avoided.

For 2D access patterns (such as tiles of pixels in an image processing kernel), it’s good practice to pad the shared memory allocation so the kernel can reference adjacent rows without causing bank conflicts. SM 2.x and subsequent hardware has 32 banks⁴; so for 2D tiles where threads in the same warp may access the data by row, it is a good strategy to pad the tile size to a multiple of 33 32-bit words.

On SM 1.x hardware, shared memory is about 16K in size⁵; on later hardware, there is a total of 64K of L1 cache that may be configured as 16K or 48K of shared memory, of which the remainder is used as L1 cache⁶.

Over the last few generations of hardware, NVIDIA has improved the hardware’s handling of operand sizes other than 32 bits. On SM 1.x hardware, 8- and 16-bit reads from the same bank caused bank conflicts, while SM 2.x and later hardware can broadcast reads of any size out of the same bank. Similarly, 64-bit operands (such as double) in shared memory were so much slower than 32-bit operands on SM 1.x that developers sometimes had to resort to storing the data as separate high and low halves.

SM 3.x hardware adds a new feature for kernels that predominantly use 64-bit operands in shared memory: a mode that increases the bank size to 64 bits.

Atomics in Shared Memory

SM 1.2 added the ability to perform atomic operations in shared memory. The earliest hardware had no dedicated shared-memory atomic unit: the compiler synthesized each operation from an explicit lock/unlock loop, built around an LDSLK (load shared and lock) instruction that returned a predicate telling whether the lock had been acquired. The update was predicated on that lock, and the warp looped until every thread had succeeded – as many as 32 iterations under full contention, since the lock was held per 32-bit word.

That implementation is long obsolete. Beginning with Maxwell (SM 5.0), shared memory has a native atomic unit, and a shared-memory atomic compiles to a single ATOMS instruction – the shared-memory counterpart of the global-memory ATOM/RED pair, with the same split between a form that returns the prior value and a cheaper form for when it is not used. No special code is required: call the atomicAdd() family (Table 8-2) on __shared__ addresses, target a modern architecture with nvcc, and the compiler emits the native instruction.

When you must hand-roll a spin-wait anyway – polling a flag written by another warp, or implementing synchronization the built-in primitives do not cover – have the loop back off with __nanosleep(ns) (Volta and later; SASS NANOSLEEP), which suspends the issuing warp for approximately ns nanoseconds so a contended loop stops hammering the memory system. Doubling the interval on each failed attempt – a capped exponential backoff – is a good default.

8.1.6 Barriers and Coherency

The familiar __syncthreads() intrinsic waits until all the threads in the thread block have arrived before proceeding. It is needed to maintain coherency of shared memory within a thread block⁷. Other, similar memory barrier instructions can be used to enforce some ordering on broader scopes of memory, as described in Table 8-3.

Table 8-3. Memory Barrier Intrinsics