The CUDA Handbook ▸ Chapter 5. Memory

5.2 Global Memory

Global memory is the main abstraction by which CUDA kernels read or write device memory⁴. Since device memory is directly attached to the GPU and read and written using a memory controller integrated into the GPU, the peak bandwidth is extremely high – more than 600GB/s is typical for high end CUDA cards.

Device memory can be accessed by CUDA kernels using device pointers. A simple memset kernel gives an example:

template<class T>
__global__ void
GPUmemset( int *base, int value, size_t N )
{
    for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;
          i < N;
          i += gridDim.x*blockDim.x )
    {
        base[i] = value;
    }
}

The device pointer base resides in the device address space, separate from the CPU address space used by the host code in the CUDA program. As a result, host code in the CUDA program can perform pointer arithmetic on device pointers, but typically the CPU may not dereference them.

There are two exceptions to this rule, one in each direction. The first is mapped pinned host memory (see Section 5.1.3): memory that resides in system memory but can be accessed by the GPU. On non-UVA systems, the host and device pointers to such memory are different – the application must call cudaHostGetDevicePointer() or cuMemHostGetDevicePointer() to map the host pointer to the corresponding device pointer – but when UVA is in effect, the pointers are the same.

The second exception is the inverse: mapped device memory, device memory that has been mapped for direct access by the CPU. On Linux, NVIDIA enables this with GDRCopy, an API set built on the GPUDirect RDMA infrastructure that includes a kernel mode driver (gdrdrv). GDRCopy pins a range of device memory and maps it into the CPU’s address space through the GPU’s PCI Express base address register (BAR) window, after which the CPU can dereference a pointer to device memory directly⁵. The mapping is asymmetric in cost: CPU stores are write-combined and perform well for small transfers, but CPU loads are uncached reads that traverse PCI Express and are best avoided. The benefit is latency – for small payloads, storing directly through the mapping sidesteps the driver overhead of a memcpy call, which is why the technique is a staple of low-latency networking and HPC communication libraries.

This kernel writes the integer value into the address range given by base and N. The references to blockIdx, blockDim, and gridDim enable the kernel to operate correctly, using whatever block and grid parameters were specified to the kernel launch.

5.2.1 Pointers

When using the CUDA runtime, device pointers and host pointers both are typed as void *.

The driver API uses an integer-valued typedef called CUdeviceptr that is the same width as host pointers (i.e., 32 bits on 32-bit hosts and 64 bits on 64-bit hosts):

#if defined(__x86_64) || defined(AMD64) || defined(_M_AMD64)
typedef unsigned long long CUdeviceptr;
#else
typedef unsigned int CUdeviceptr;
#endif

The uintptr_t type, available in <stdint.h> and introduced in C++0x, may be used to portably convert between host pointers (void *) and device pointers (CUdeviceptr):

CUdeviceptr devicePtr;
void *p;
p = (void *) (uintptr_t) devicePtr;
devicePtr = (CUdeviceptr) (uintptr_t) p;

The host can do pointer arithmetic on device pointers to pass to a kernel or memcpy call; but the host cannot read or write device memory with these pointers.

32- and 64-bit Pointers in the Driver API

Because the original driver API definition for a pointer was 32-bit, the addition of 64-bit support to CUDA required the definition of CUdeviceptr and, in turn, all driver API functions that took CUdeviceptr as a parameter, to change⁶. cuMemAlloc(), for example, changed from:

CUresult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, unsigned int bytesize);

to:

CUresult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, size_t bytesize);

To accommodate both old applications (which linked against a cuMemAlloc() with 32-bit CUdeviceptr and size) and new ones, cuda.h includes two blocks of code that use the preprocessor to change the bindings without requiring function names to be changed as developers update to the new API.

First, a block of code surreptitiously changes function names to map to newer functions that have different semantics:

#if defined(__CUDA_API_VERSION_INTERNAL) || __CUDA_API_VERSION >= 3020

    #define cuDeviceTotalMem() cuDeviceTotalMem_v2()

…

    #define cuTexRefGetAddress() cuTexRefGetAddress_v2()

#endif /* __CUDA_API_VERSION_INTERNAL || __CUDA_API_VERSION >= 3020 */

This way, the client code uses the same old function names, but the compiled code generates references to the new function names with _v2 appended.

Later in the header, the old functions are defined as they were.

As a result, developers compiling for the latest version of CUDA get the latest function definitions and semantics.

cuda.h uses a similar strategy for functions whose semantics changed from one version to the next, such as cuStreamDestroy().

5.2.2 Pointer Queries

CUDA tracks all of its memory allocations, and provides APIs that enable applications to query CUDA about pointers that were passed in from some other party. Libraries or plugins may wish to pursue different strategies based on this information.

CUDA Runtime

The cudaPointerGetAttributes() function takes a pointer as input, and passes back a cudaPointerAttributes structure containing information about the pointer.

    struct cudaPointerAttributes {
        enum cudaMemoryType memoryType;
        int device;
        void *devicePointer;
        void *hostPointer;
    }

When UVA (unified virtual addressing) is in effect, pointers are unique process-wide, so there is no ambiguity as to the input pointer’s address space. When UVA is not in effect, the input pointer is assumed to be in the current device’s address space.

Table 5-3. cudaPointerAttributes members.

Driver API

Developers can query the address range where a given device pointer resides using the cuMemGetAddressRange() function:

CUresult CUDAAPI cuMemGetAddressRange(CUdeviceptr *pbase, size_t *psize, CUdeviceptr dptr);

This function takes a device pointer as input and passes back the base and size of the allocation containing that device pointer.

With unified virtual addressing (UVA), developers can query CUDA to get even more information about an address using cuPointerGetAttribute():

CUresult CUDAAPI cuPointerGetAttribute(void *data, CUpointer_attribute attribute, CUdeviceptr ptr);

This function takes a device pointer as input and passes back the information corresponding to the attribute parameter, as shown in Table 5-4.

Table 5-4. cuPointerAttribute usage.

Note that for unified addresses, using CU_POINTER_ATTRIBUTE_DEVICE_POINTER or CU_POINTER_ATTRIBUTE_HOST_POINTER will cause the same pointer value to be returned as the one passed in.

Kernel Queries

On SM 2.x (Fermi) hardware and later, developers can query whether a given pointer points into global space. The __isGlobal() intrinsic:

unsigned int __isGlobal( const void *p );

returns 1 if the input pointer refers to global memory, and 0 otherwise.

5.2.3 Dynamic Allocations

Most global memory in CUDA is obtained through dynamic allocation. Using the CUDA runtime, the functions:

cudaError_t cudaMalloc( void **, size_t );
cudaError_t cudaFree( void );

allocate and free global memory, respectively. The corresponding driver API functions are:

CUresult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, size_t bytesize);
CUresult CUDAAPI cuMemFree(CUdeviceptr dptr);

Allocating global memory is expensive. The CUDA driver implements a suballocator to satisfy small allocation requests, but if the suballocator must create a new memory block, that requires an expensive operating system call to the kernel mode driver. If that happens, the CUDA driver also must synchronize with the GPU, which may break CPU/GPU concurrency. As a result, it’s good practice to avoid allocating or freeing global memory in performance-sensitive code.

Stream-Ordered Allocation

The synchronous cost just described – the kernel-mode call and the device synchronization – is what the stream-ordered allocator was built to avoid. cudaMallocAsync() and cudaFreeAsync() take a stream and are ordered within it like any other stream operation: an allocation becomes valid for work issued into the stream after the cudaMallocAsync(), and a free takes effect once the work issued before the cudaFreeAsync() has completed. Neither call synchronizes the device. The allocator arrived in CUDA 11.2 and runs on Pascal (compute capability 6.0) and newer; because architecture alone does not guarantee it, confirm support at runtime by querying the cudaDevAttrMemoryPoolsSupported device attribute.

Freed memory is not returned to the operating system but held in a per-device memory pool and reused to satisfy later requests. A loop that allocates and frees a scratch buffer each iteration – the per-timestep pattern the synchronous allocator stalls on – touches the driver once to grow the pool and reuses that block from then on. On a GeForce RTX 3060, mallocAsyncSpeed.cu measures the gap:

Table 5-13. Per-iteration cost of allocate-then-free, synchronous versus stream-ordered (GeForce RTX 3060).

The synchronous pair costs tens of microseconds even for a small buffer, nearly all of it the device synchronization, and rises with the allocation size; the stream-ordered pair costs a fraction of a microsecond once the pool is warm and barely varies with size, because it is reusing memory rather than obtaining it.

Unlike the in-kernel malloc() available to device code, whose heap must be sized ahead of time with cudaDeviceSetLimit(cudaLimitMallocHeapSize), the pool is not fixed at a preset size: it grows on demand and is bounded only by free device memory. Its retention is governed by a release threshold (cudaMemPoolAttrReleaseThreshold) – by default the pool returns unused memory to the operating system at the next synchronization point, so raising the threshold is what keeps a large pool resident between synchronizations. The one discipline the API adds is the stream ordering itself: because an allocation is valid only after its stream reaches the cudaMallocAsync() and a buffer is reclaimed after its cudaFreeAsync(), using memory out of that order – or from another stream without ordering the two with an event – is a use-before-allocation or use-after-free hazard, not the always-safe pointer the synchronous cudaMalloc() returns.

Virtual Memory Management

cudaMalloc() reserves virtual address space and commits physical memory to back it in one call, and the two cannot be separated: to grow an allocation, the application has to allocate a larger block, copy the old contents across, and free the original, invalidating every pointer into it. The driver’s virtual memory management (VMM) API splits those operations so they can be controlled independently. It is a driver-API facility, exposed through cuMem* functions with no runtime-API equivalents.

The three steps are separate calls:

• cuMemAddressReserve() reserves a range of virtual addresses, without backing them.
• cuMemCreate() allocates a physical memory handle – a block of device memory that has no address of its own.
• cuMemMap() maps a handle into part of a reserved range, after which cuMemSetAccess() grants one or more devices permission to access it.

cuMemUnmap() and cuMemRelease() reverse the mapping and the allocation. Because reservation, allocation, and mapping are separate steps, an application can reserve a large address range once and map physical memory into it a piece at a time as a buffer grows – the base pointer never moves and no data is copied, which is how a growable device container sidesteps the reallocate-and-copy cost of cudaMalloc().

The same physical handles can be shared between processes. cuMemExportToShareableHandle() turns a handle into an operating-system shareable handle – a file descriptor on Linux, an NT handle on Windows – that another process imports with cuMemImportFromShareableHandle() and maps into its own address space, the same handle machinery the external-resource API of Section 3.11 uses for graphics interop. This layer is also what the stream-ordered pool allocator uses to grow and shrink its pools.

Pitched Allocations

The coalescing constraints, coupled with alignment restrictions for texturing and 2D memory copy, motivated the creation of pitched memory allocations. The idea is that when creating a 2D array, a pointer into the array should have the same alignment characteristics when updated to point to a different row. The pitch of the array is the number of bytes per row of the array⁷. The pitch allocations take a width (in bytes) and height, pad the width to a suitable hardware-specific pitch, and pass back the base pointer and pitch of the allocation. By using these allocation functions to delegate selection of the pitch to the driver, developers can future-proof their code against architectures that widen alignment requirements⁸.

CUDA programs often must adhere to alignment constraints enforced by the hardware, not only on base addresses but also on the widths (in bytes) of memory copies and linear memory bound to textures. Because the alignment constraints are hardware-specific, CUDA provides APIs that enable developers to delegate the selection of the appropriate alignment to the driver. Using these APIs enables CUDA applications to implement hardware-independent code and to be “future-proof” against CUDA architectures that have not yet shipped.

Figure 5-1 shows a pitch allocation being performed on an array that is 352 bytes wide. The pitch is padded to the next multiple of 64 bytes before allocating the memory.

Figure 5-1. Pitch versus width.

Given the pitch of the array in addition to the row and column, the address of an array element can be computed as follows:

inline T *
getElement( T *base, size_t Pitch, int row, int col )
{
    return (T *) ((char *) base + row*Pitch) + col;
}

The CUDA runtime function to perform a pitched allocation is as follows:

template<class T>
__inline__ __host__ cudaError_t cudaMallocPitch(
  T **devPtr,
  size_t *pitch,
  size_t widthInBytes,
  size_t height
);

The CUDA runtime also includes the function cudaMalloc3D(), which allocates 3D memory regions using the cudaPitchedPtr and cudaExtent structures:

extern __host__ cudaError_t CUDARTAPI cudaMalloc3D(struct cudaPitchedPtr* pitchedDevPtr, struct cudaExtent extent);

cudaPitchedPtr, which receives the allocated memory, is defined as follows:

struct cudaPitchedPtr
{
   void *ptr;
   size_t pitch;
   size_t xsize;
   size_t ysize;
};

cudaPitchedPtr::ptr specifies the pointer; cudaPitchedPtr::pitch specifies the pitch (width in bytes) of the allocation; cudaPitchedPtr::xsize and cudaPitchedPtr::ysize are the logical width and height of the allocation, respectively.

cudaExtent is defined as follows:

struct cudaExtent
{
   size_t width;
   size_t height;
   size_t depth;
};

cudaExtent::width is treated differently for arrays and linear device memory. For arrays, it specifies the width in array elements; for linear device memory, it specifies the pitch (width in bytes).

The driver API function to allocate memory with a pitch is as follows:

CUresult CUDAAPI cuMemAllocPitch(CUdeviceptr *dptr, size_t *pPitch, size_t WidthInBytes, size_t Height, unsigned int ElementSizeBytes);

The ElementSizeBytes parameter may be 4, 8 or 16 bytes, and causes the allocation pitch to be padded to 64-, 128-, or 256-byte boundaries. Those are the alignment requirements for coalescing of 4-, 8-, and 16-byte memory transactions on SM 1.0 and SM 1.1 hardware. Applications that are not concerned with running well on that hardware can specify 4.

The pitch returned by cudaMallocPitch()/cuMemAllocPitch() is the width-in-bytes passed in by the caller, padded to an alignment that meets the alignment constraints for both coalescing of global load/store operations, and texture bind APIs. The amount of memory allocated is height*pitch.

For 3D arrays, developers can multiply the height by the depth before performing the allocation. This consideration only applies to arrays that will be accessed via global loads and stores, since 3D textures cannot be bound to global memory.

Allocations Within Kernels

Fermi-class hardware can dynamically allocate global memory using malloc(). Since this may require the GPU to interrupt the CPU, it is potentially slow. The sample program devrtMallocSpeed.cu measures the performance of malloc() and free() in kernels.

Listing 5-3 shows the key kernels and timing routine in devrtMallocSpeed.cu. As an important note, the cudaDeviceSetLimit() function must be called with cudaLimitMallocHeapSize before malloc() may be called in kernels. The invocation in devrtMallocSpeed.cu requests a full gigabyte (2³⁰ bytes):

cuda(DeviceSetLimit(cudaLimitMallocHeapSize, 1<<30) );

When cudaDeviceSetLimit() is called, the requested amount of memory is allocated and may not be used for any other purpose.

__global__ voidAllocateBuffers( void **out, size_t N ){    size_t i = blockIdx.x*blockDim.x + threadIdx.x;    out[i] = malloc( N );} __global__ voidFreeBuffers( void **in ){    size_t i = blockIdx.x*blockDim.x + threadIdx.x;    free( in[i] );} cudaError_tMallocSpeed( double *msPerAlloc, double *msPerFree,              void **devicePointers, size_t N,             cudaEvent_t evStart, cudaEvent_t evStop,             int cBlocks, int cThreads ){    float etAlloc, etFree;    cudaError_t status;     cuda(EventRecord( evStart ) );    AllocateBuffers<<<cBlocks,cThreads>>>( devicePointers, N );    cuda(EventRecord( evStop ) );    cuda(DeviceSynchronize() );    cuda(GetLastError() );     cuda(EventElapsedTime( &etAlloc, evStart, evStop ) );     cuda(EventRecord( evStart ) );    FreeBuffers<<<cBlocks,cThreads>>>( devicePointers );    cuda(EventRecord( evStop ) );    cuda(DeviceSynchronize() );    cuda(GetLastError() );     cuda(EventElapsedTime( &etFree, evStart, evStop ) );     *msPerAlloc = etAlloc / (double) (cBlocks*cThreads);    *msPerFree = etFree / (double) (cBlocks*cThreads); Error:    return status;}

Listing 5-4 shows the output from a sample run of devrtMallocSpeed.cu on a GeForce RTX 3060. It is clear that the functionality is optimized for small allocations: the 64-byte allocations take about 0.1 microseconds to perform, and 12K allocations run well under a microsecond except at the largest block sizes. The first result (about 5 microseconds per allocation) is having 1 thread per each of 500 blocks allocate a 1MB buffer.

Microseconds per alloc/free (1 thread per block):alloc	free4.76	1.92	 Microseconds per alloc/free (32-512 threads per block, 12K allocations):32		64		128		256		512		alloc	free	alloc	free	alloc	free	alloc	free	alloc	free	0.82	0.49	0.78	0.49	0.81	0.49	0.94	0.48	1.65	0.48	 Microseconds per alloc/free (32-512 threads per block, 64-byte allocations):32		64		128		256		512		alloc	free	alloc	free	alloc	free	alloc	free	alloc	free	0.15	0.22	0.14	0.21	0.12	0.20	0.11	0.20	0.10	0.15

5.2.4 Static Allocations

Applications can statically allocate global memory by annotating a memory declaration with the __device__ keyword. This memory is allocated by the CUDA driver when the module is loaded.

CUDA Runtime

Memory copies to and from statically allocated memory can be performed by cudaMemcpyToSymbol() and cudaMemcpyFromSymbol():

cudaError_t cudaMemcpyToSymbol(
        char *symbol,
  const void *src,
        size_t count,
        size_t offset = 0,
        enum cudaMemcpyKind kind = cudaMemcpyHostToDevice
);
cudaError_t cudaMemcpyFromSymbol(
  void *dst,
  char *symbol,
  size_t count,
  size_t offset = 0,
  enum cudaMemcpyKind kind = cudaMemcpyDeviceToHost
);

When calling cudaMemcpyToSymbol() or cudaMemcpyFromSymbol(), do not enclose the symbol name in quotation marks, i.e.:

    cudaMemcpyToSymbol(g_xOffset, poffsetx, Width*Height*sizeof(int));

not

    cudaMemcpyToSymbol(“g_xOffset,” poffsetx, ... );

Both formulations work, but the latter formulation will compile for any symbol name (even undefined symbols). If you want compile errors for invalid symbols, avoid the quotation marks.

CUDA runtime applications can query the pointer corresponding to a static allocation by calling cudaGetSymbolAddress():

cudaError_t cudaGetSymbolAddress( void **devPtr, char *symbol );

Beware: it is all too easy to pass the symbol for a statically-declared device memory allocation to a CUDA kernel; but this does not work. You must call cudaGetSymbolAddress() and use the resulting pointer.

Driver API

Developers using the driver API can obtain pointers to statically allocated memory by calling cuModuleGetGlobal():

CUresult CUDAAPI cuModuleGetGlobal(CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name);

Note that cuModuleGetGlobal() passes back both the base pointer and the size of the object. If the size is not needed, developers can pass NULL for the bytes parameter.

Once this pointer has been obtained, the memory can be accessed by passing the CUdeviceptr to memory copy calls or CUDA kernel invocations.

5.2.5 Querying The Amount of Global Memory

The amount of global memory in a system may be queried even before CUDA has been initialized.

CUDA Runtime

Call cudaGetDeviceProperties() and examine cudaDeviceProp.totalGlobalMem:

size_t totalGlobalMem; /**< Global memory on device in bytes */

Driver API

Call this driver API function:

CUresult CUDAAPI cuDeviceTotalMem(size_t *bytes, CUdevice dev);
WDDM (Windows Display Driver Model) and Available Memory

The WDDM driver model introduced with Windows Vista changed the model for memory management by display drivers, so that chunks of video memory could be swapped in and out of host memory as needed to perform rendering. As a result, the amount of memory reported by cuDeviceTotalMem() / cudaDeviceProp::totalGlobalMem will not exactly reflect the amount of physical memory on the card.

5.2.6 Memset APIs

For developer convenience, CUDA provides 1D and 2D memset functions. Since they are implemented using kernels, they are asynchronous even when no stream parameter is specified; but for applications that must serialize the execution of a memset within a stream, there are *Async() variants that take a stream parameter.

CUDA Runtime

The CUDA runtime supports byte-sized memset only:

cudaError_t cudaMemset(void *devPtr, int value, size_t count);
cudaError_t cudaMemset2D(void *devPtr, size_t pitch, int value, size_t width, size_t height);

The pitch parameter specifies the bytes per row of the memset operation.

Driver API

The driver API supports 1D and 2D memset of a variety of sizes:

Table 5-2. Memset variations

These memset functions take the destination pointer, value to set, and number of values to write starting at the base address. The pitch parameter is the bytes per row (not elements per row!).

CUresult CUDAAPI cuMemsetD8(CUdeviceptr dstDevice, unsigned char uc, size_t N);
CUresult CUDAAPI cuMemsetD16(CUdeviceptr dstDevice, unsigned short us, size_t N);
CUresult CUDAAPI cuMemsetD32(CUdeviceptr dstDevice, unsigned int ui, size_t N);
CUresult CUDAAPI cuMemsetD2D8(CUdeviceptr dstDevice, size_t dstPitch, unsigned char uc, size_t Width, size_t Height);
CUresult CUDAAPI cuMemsetD2D16(CUdeviceptr dstDevice, size_t dstPitch, unsigned short us, size_t Width, size_t Height);
CUresult CUDAAPI cuMemsetD2D32(CUdeviceptr dstDevice, size_t dstPitch, unsigned int ui, size_t Width, size_t Height);

Now that CUDA runtime and driver API functions can peacefully coexist in the same application, CUDA runtime developers can use these functions as needed. The unsigned char, unsigned short and unsigned int parameters just specify a bit pattern; to fill a global memory range with some other type, such as float, use a volatile union to coerce the float to unsigned int.

5.2.7 Reading and Writing Global Memory

CUDA kernels can read or write global memory using standard C semantics such as pointer indirection (operator*, operator->) or array subscripting (operator[]).

A simple templatized kernel to write a constant into a memory range is as follows:

template<class T>
__global__ void
GlobalWrites( T *out, T value, size_t N )
{
    for ( size_t i = blockIdx.x*blockDim.x+threadIdx.x;
          i < N;
          i += blockDim.x*gridDim.x ) {
        out[i] = value;
    }
}

This kernel works correctly for any inputs: any component size, any block size, any grid size. Its code is intended more for illustrative purposes than maximum performance – CUDA kernels that use more registers and operate on multiple values in the inner loop go faster – but for some block and grid configurations, its performance is perfectly acceptable. In particular, provided the base address and block size are specified correctly, it performs coalesced memory transactions that maximize memory bandwidth.

5.2.8 Coalescing Constraints

For best performance when reading and writing data, CUDA kernels must perform coalesced memory transactions. Any memory transaction that does not meet the full set of criteria needed for coalescing is “uncoalesced.” The penalty for uncoalesced memory transactions varies from 2x to 8x, depending on the chip implementation. Coalesced memory transactions have a much less dramatic impact on performance on more recent hardware, as shown in Table 5-5.

Transactions are coalesced on a per-warp basis. A simplified set of criteria must be met in order for the memory read or write being performed by the warp to be coalesced:

• The words must be at least 32 bits in size. Reading or writing bytes or 16-bit words is always uncoalesced.

• The addresses being accessed by the threads of the warp must be contiguous and increasing (i.e., offset by the thread ID).

• The base address of the warp (the address being accessed by the first thread in the warp) must be aligned as shown in Table 5-6.

The ElementSizeBytes parameter to cuMemAllocPitch() is intended to accommodate the size restriction: it specifies the size in bytes of the memory accesses intended by the application, so the pitch guarantees that a set of coalesced memory transactions for a given row of the allocation also will be coalesced for other rows.

Most kernels in this book perform coalesced memory transactions, provided the input addresses are properly aligned. NVIDIA has provided more detailed, architecture-specific information on how global memory transactions are handled, as detailed below.

Table 5-5. Bandwidth Penalties for Uncoalesced Memory Access

* 8- and 16-bit memory accesses are always uncoalesced.

Table 5-6. Alignment Criteria for Coalescing

SM 1.x (Tesla)

SM 1.0 and SM 1.1 hardware require that each thread in a warp access adjacent memory locations in sequence, as described above.

SM 1.2 and 1.3 hardware relaxed the coalescing constraints somewhat. To issue a coalesced memory request, divide each 32-thread warp into two “half warps,” lanes 0-15 and lanes 16-31. To service the memory request from each half-warp, the hardware performs the following algorithm:

1. Find the active thread with the lowest thread ID, and locate the memory segment that contains that thread’s requested address. The segment size depends on the word size: 1-byte requests result in 32-byte segments; 2-byte requests result in 64-byte segments; all other requests result in 128-byte segments.

2. Find all other active threads whose requested address lies in the same segment;

3. If possible, reduce the segment transaction size to 64 or 32 bytes;

4. Carry out the transaction and mark the serviced threads as inactive.

5. Repeat steps 1-4 until all threads in the half-warp have been serviced.

Although these requirements are somewhat relaxed as compared to the SM 1.0-1.1 constraints, a great deal of locality is still required for effective coalescing. In practice, the relaxed coalescing means the threads within a warp can permute the inputs within small segments of memory, if desired.

SM 2.x (Fermi)

SM 2.x and later hardware includes L1 and L2 caches. The L2 cache services the entire chip; the L1 caches are per-SM and may be configured to be 16K or 48K in size.

The cache lines are 128 bytes and map to 128-byte aligned segments in device memory. Memory accesses that are cached in both L1 and L2 are serviced with 128-byte memory transactions whereas memory accesses that are cached in L2 only are serviced with 32-byte memory transactions. Caching in L2 only can therefore reduce over-fetch, for example, in the case of scattered memory accesses.

The hardware can specify the cacheability of global memory accesses on a per-instruction basis. By default, the compiler emits instructions that cache memory accesses in both L1 and L2 (-Xptxas -dlcm=ca); this can be changed to cache only in L2 by specifying -Xptxas -dlcm=cg. Memory accesses that are not present in L1, but cached only in L2, are serviced with 32-byte memory transactions, which may improve cache utilization for applications that are performing scattered memory accesses.

Reading via pointers that are declared volatile causes any cached results to be discarded and for the data to be re-fetched. This idiom is mainly useful for polling host memory locations.

Table 5-7. SM 2.x Cache Line Requests

Table 5-7 summarizes how memory requests by a warp are broken down into 128-byte cache line requests.

On SM 2.x and higher architectures, threads within a warp can access any words in any order, including the same words.

SM 3.x (Kepler)

The L2 cache architecture is the same as SM 2.x; SM 3.x does not cache global memory accesses in L1.

In SM 3.5, global memory may be accessed via the texture cache (which is 48K per SM in size), by accessing memory via const restrict pointers, or by using the __ldg() intrinsics in sm_35_intrinsics.h. As when texturing directly from device memory, it is important not to access memory that might be accessed concurrently by other means, since this cache is not kept coherent with respect to the L2.

5.2.9 Microbenchmarks: Peak Memory Bandwidth

The source code accompanying this book includes microbenchmarks that determine which combination of operand size, loop unroll factor, and block size maximizes bandwidth for a given GPU. Rewriting the earlier GlobalWrites() code as a template that takes an additional parameter n (the number of writes to perform in the inner loop) yields the kernel of Listing 5-5.

template<class T, const int n> __global__ voidGlobalWrites( T *out, T value, size_t N ){    size_t i;    for ( i = n*blockIdx.x*blockDim.x+threadIdx.x;           i < N-n*blockDim.x*gridDim.x;           i += n*blockDim.x*gridDim.x ) {        for ( int j = 0; j < n; j++ ) {            size_t index = i+j*blockDim.x;            out[index] = value;        }    }    // to avoid the (index<N) conditional in the inner loop,     // we left off some work at the end    for ( int j = 0; j < n; j++ ) {        size_t index = i+j*blockDim.x;        if ( index<N ) out[index] = value;    }}

ReportRow(), the function given in Listing 5-6 that writes one row of output by calling a template function BandwidthWrites (not shown) that reports the bandwidth for a given type, grid, and block size:

template<class T, const int n, bool bOffset>doubleReportRow( size_t N,            size_t threadStart,            size_t threadStop,            size_t cBlocks ){    int maxThreads = 0;    double maxBW = 0.0;    printf( "%d\t", n );    for ( int cThreads = threadStart;               cThreads <= threadStop;               cThreads *= 2 ) {        double bw;        bw = BandwidthWrites<T,n,bOffset>( N, cBlocks, cThreads );        if ( bw > maxBW ) {            maxBW = bw;            maxThreads = cThreads;        }        printf( "%.2f\t", bw );    }    printf( "%.2f\t%d\n", maxBW, maxThreads );    return maxBW;}

The threadStart and threadStop parameters typically are 32 and 512. 32 is the warp size and the minimum number of threads per block that can occupy the machine.

The bOffset template parameter specifies whether BandwidthWrites should offset the base pointer, causing all memory transactions to become uncoalesced. If the program is invoked with the --uncoalesced command line option, it will perform the bandwidth measurements with the offset pointer.

Note that depending on sizeof(T), kernels with n above a certain level will fall off a performance cliff as the number of temporary variables in the inner loop grows too high to hold in registers.

The 5 applications summarized in Table 5-8 implement this strategy: they measure the memory bandwidth delivered for different operand sizes (8-, 16-, 32-, 64-, and 128-bit), threadblock sizes (32, 64, 128, 256, and 512), and loop unroll factors (1-16). CUDA hardware isn’t necessarily sensitive to all of these parameters – for example, many parameter settings enable a GK104 to deliver 140GB/s of bandwidth via texturing, but only if the operand size is at least 32-bit – but for a given workload and hardware, the microbenchmarks highlight which parameters matter. Also, for small operand sizes, they highlight how loop unrolling can help increase performance (not all applications can be refactored to read larger operands).

Table 5-8. Memory Bandwidth Microbenchmarks.

Listing 5-7 gives example output from globalRead.cu, run on a GeForce GTX 680 GPU. The output is grouped by operand size, from bytes to 16-byte quads; the leftmost column of each group gives the loop unroll factor. The bandwidth delivered for blocks of sizes 32-512 is given in each column, then the maxBW and maxThreads columns give the highest bandwidth and the block size that delivered the highest bandwidth, respectively.

The GeForce GTX 680 can deliver up to 140 GB/s, so Listing 5-7 makes it clear that when reading 8- and 16-bit words on SM 3.0, global loads are not the way to go: bytes deliver at most 60 GB/s and 16-bit words deliver at most 101 GB/s⁹. For 32-bit operands, a 2x loop unroll and at least 256 threads per block are needed to get maximum bandwidth.

That picture is now a historical one, however: successive generations of CUDA hardware have gotten much better at saturating global memory bandwidth, no matter the operand size. Larger caches and smarter coalescing hardware have absorbed most of the penalty described above. On an Ampere-class GeForce RTX 3060 (2021), the same globalRead.cu microbenchmark delivers 300 GB/s for byte operands and 318 GB/s for 16-byte quads – a spread of about 6 percent, on a part with a theoretical peak of 360 GB/s – and globalWrite.cu and globalCopy.cu are similarly flat. Loop unrolling and block size still move the numbers at the margins, but the days when byte-sized loads forfeited more than half of the machine’s bandwidth are over. Reader-submitted results across GPU generations are collected at cudahandbook.com/benchmarks.

These microbenchmarks can help developers optimize their bandwidth-bound applications – choose the one whose memory access pattern most closely resembles your application, and either run the microbenchmark on the target GPU or, if possible, modify the microbenchmark to resemble the actual workload more closely and run it to determine the optimal parameters.

Running globalRead.cu microbenchmark on GeForce GTX 680Using coalesced memory transactionsOperand size: 1 byte Input size: 16M operands                      Block SizeUnroll	32	64	128	256	512	maxBW	maxThreads1	9.12	17.39	30.78	30.78	28.78	30.78	1282	18.37	34.54	56.36	53.53	49.33	56.36	1283	23.55	42.32	61.56	60.15	52.91	61.56	1284	21.25	38.26	58.99	58.09	51.26	58.99	1285	25.29	42.17	60.13	58.49	52.57	60.13	1286	25.68	42.15	59.93	55.42	47.46	59.93	1287	28.84	47.03	56.20	51.41	41.41	56.20	1288	29.88	48.55	55.75	50.68	39.96	55.75	1289	28.65	47.75	56.84	51.17	37.56	56.84	12810	27.35	45.16	52.99	46.30	32.94	52.99	12811	22.27	38.51	48.17	42.74	32.81	48.17	12812	23.39	40.51	49.78	42.42	31.89	49.78	12813	21.62	37.49	40.89	34.98	21.43	40.89	12814	18.55	32.12	36.04	31.41	19.96	36.04	12815	21.47	36.87	39.94	33.36	19.98	39.94	12816	21.59	36.79	39.49	32.71	19.42	39.49	128Operand size: 2 bytesInput size: 16M operands                      Block SizeUnroll	32	64	128	256	512	maxBW	maxThreads1	18.29	35.07	60.30	59.16	56.06	60.30	1282	34.94	64.39	94.28	92.65	85.99	94.28	1283	45.02	72.90	101.38	99.02	90.07	101.38	1284	38.54	68.35	100.30	98.29	90.28	100.30	1285	45.49	75.73	98.68	98.11	90.05	98.68	1286	47.58	77.50	100.35	97.15	86.17	100.35	1287	53.64	81.04	92.89	87.39	74.14	92.89	1288	44.79	74.02	89.19	83.96	69.65	89.19	1289	47.63	76.63	91.60	83.52	68.06	91.60	12810	51.02	79.82	93.85	84.69	66.62	93.85	12811	42.00	72.11	88.23	79.24	62.27	88.23	12812	40.53	69.27	85.75	76.32	59.73	85.75	12813	44.90	73.44	78.08	66.96	41.27	78.08	12814	39.18	68.43	74.46	63.27	39.27	74.46	12815	37.60	64.11	69.93	60.22	37.09	69.93	12816	40.36	67.90	73.07	60.79	36.66	73.07	128Operand size: 4 bytesInput size: 16M operands                      Block SizeUnroll	32	64	128	256	512	maxBW	maxThreads1	36.37	67.89	108.04	105.99	104.09	108.04	1282	73.85	120.90	139.91	139.93	136.04	139.93	2563	62.62	109.24	140.07	139.66	138.38	140.07	1284	56.02	101.73	138.70	137.42	135.10	138.70	1285	87.34	133.65	140.64	140.33	139.00	140.64	1286	100.64	137.47	140.61	139.53	127.18	140.61	1287	89.08	133.99	139.60	138.23	124.28	139.60	1288	58.46	103.09	129.24	122.28	110.58	129.24	1289	68.99	116.59	134.17	128.64	114.80	134.17	12810	54.64	97.90	123.91	118.84	106.96	123.91	12811	64.35	110.30	131.43	123.90	109.31	131.43	12812	68.03	113.89	130.95	125.40	108.02	130.95	12813	71.34	117.88	123.85	113.08	76.98	123.85	12814	54.72	97.31	109.41	101.28	71.13	109.41	12815	67.28	111.24	118.88	108.35	72.30	118.88	12816	63.32	108.56	117.77	103.24	69.76	117.77	128Operand size: 8 bytesInput size: 16M operands                      Block SizeUnroll	32	64	128	256	512	maxBW	maxThreads1	74.64	127.73	140.91	142.08	142.16	142.16	5122	123.70	140.35	141.31	141.99	142.42	142.42	5123	137.28	141.15	140.86	141.94	142.63	142.63	5124	128.38	141.39	141.85	142.56	142.00	142.56	2565	117.57	140.95	141.17	142.08	141.78	142.08	2566	112.10	140.62	141.48	141.86	141.95	141.95	5127	85.02	134.82	141.59	141.50	141.09	141.59	1288	94.44	138.71	140.86	140.25	128.91	140.86	1289	100.69	139.83	141.09	141.45	127.82	141.45	25610	92.51	137.76	140.74	140.93	126.50	140.93	25611	104.87	140.38	140.67	136.70	128.48	140.67	12812	97.71	138.62	140.12	135.74	125.37	140.12	12813	95.87	138.28	139.90	134.18	123.41	139.90	12814	85.69	134.18	133.84	131.16	120.95	134.18	6415	94.43	135.43	135.30	133.47	120.52	135.43	6416	91.62	136.69	133.59	129.95	117.99	136.69	64Operand size: 16 bytesInput size: 16M operands                      Block SizeUnroll	32	64	128	256	512	maxBW	maxThreads1	125.37	140.67	141.15	142.06	142.59	142.59	5122	131.26	141.95	141.72	142.32	142.49	142.49	5123	141.03	141.65	141.63	142.43	138.44	142.43	2564	139.90	142.70	142.62	142.20	142.84	142.84	5125	138.24	142.08	142.18	142.79	140.94	142.79	2566	131.41	142.45	142.32	142.51	142.08	142.51	2567	131.98	142.26	142.27	142.11	142.26	142.27	1288	132.70	142.47	142.10	142.67	142.19	142.67	2569	136.58	142.28	141.89	142.42	142.09	142.42	25610	135.61	142.67	141.85	142.86	142.36	142.86	25611	136.27	142.48	142.45	142.14	142.41	142.48	6412	130.62	141.79	142.06	142.39	142.16	142.39	25613	107.98	103.07	105.54	106.51	107.35	107.98	3214	103.53	95.38	96.38	98.34	102.92	103.53	3215	89.47	84.86	85.31	87.01	90.26	90.26	51216	81.53	75.49	75.82	74.36	76.91	81.53	32

5.2.10 Persisting Accesses in the L2 Cache

Every global-memory access passes through the L2 cache, which is shared by all the SMs and is the last stop before device memory. Its replacement policy is ordinarily out of the program’s hands: a heavily reused array and a once-touched streaming array compete for the same lines, and a kernel that sweeps a large working set can evict the very data it is about to reuse. Beginning with the Ampere architecture, an application can bias that competition by reserving part of the L2 for a chosen range of memory and marking accesses to it persisting, so its lines are held in preference to ordinary streaming traffic.

The control has two parts. First, a set-aside: cudaDeviceSetLimit() with cudaLimitPersistingL2CacheSize reserves a portion of the L2 – up to the device’s persistingL2CacheMaxSize – for persisting accesses. Second, an access-policy window: a cudaAccessPolicyWindow names a contiguous range of global memory (base_ptr and num_bytes, the latter no larger than accessPolicyMaxWindowSize) and the properties to apply – cudaAccessPropertyPersisting for accesses that land in the window and cudaAccessPropertyStreaming for accesses beyond its bounds. The window is attached to a stream with cudaStreamSetAttribute() – or to a graph node, or to one launch through cudaLaunchKernelEx() – and every kernel launched on that stream applies it. When the range is larger than the set-aside, the window’s hitRatio field asks the hardware to mark that fraction of its lines persisting, chosen at random, so a region too large to pin whole does not thrash the cache.

Persisting lines stay reserved until they are released, either by clearing the window with a num_bytes of zero or by calling cudaCtxResetPersistingL2Cache(), which returns the whole set-aside to normal use. A region left marked persisting after its kernel finishes wastes L2 on data no kernel is reading, so the reservation is best scoped to the kernels that benefit.

The benefit is confined to a specific case: data reused often enough to matter, under enough streaming pressure that the default policy would evict it, and small enough to fit the set-aside. A working set that already stays resident does not benefit, because the default policy was keeping it anyway, and one larger than the set-aside cannot be pinned whole. The l2Persistence.cu microbench arranges the favorable case: a gather kernel indexes a 32 MB array, far larger than the L2, with 90% of its accesses falling in a 1.5 MB hot region and the rest scattered across the whole array to keep the cache under load. Marking the hot region persisting holds it in L2 while the scattered accesses stream past (Table 5-15).

Table 5-15. Gather over a 32 MB array on a GeForce RTX 3060 (2.25 MB L2, 1.5 MB persisting set-aside), 90% of accesses in the hot region, best of six.

The set-aside is small on this consumer part – 1.55 MB of a 2.25 MB L2 – so the win is modest; on the data-center GPUs whose L2 runs to tens of megabytes, the same technique pins a proportionally larger working set.

5.2.11 Atomic Operations

Support for atomic operations was added in SM 1.x, but was prohibitively slow; atomics on global memory were improved on SM 2.x (Fermi-class) hardware and vastly improved on SM 3.x (Kepler-class) hardware.

Atomic operations are implemented in the GPU memory controller, and do not work on memory not attached to the GPU. In particular, atomic ops will not work on mapped pinned memory (obtained via cu(da)HostGetDevicePointer()) or peer memory (enabled by cudaDeviceEnablePeerAccess()/cuCtxEnablePeerAccess()).

Most atomic operations, such as atomicAdd(), enable code to be simplified by replacing reductions (which often require shared memory and synchronization) with “fire and forget” semantics. Until SM 3.x hardware arrived, however, that type of programming idiom incurred huge performance degradations because pre-Kepler hardware was not efficient at dealing with “contended” memory locations (i.e. when many GPU threads are performing atomics on the same memory location).

Atomics and Synchronization

Besides “fire and forget” semantics, atomics also may be used for synchronization between blocks. CUDA hardware supports the workhorse base abstraction for synchronization, “compare and swap” (or CAS). On CUDA, compare-and-swap (also known as compare-and-exchange, e.g. the CMPXCHG instruction in x86) is defined as follows:

int atomicCAS( int *address, int expected, int value);¹⁰

This function reads the word old at address, computes (old == expected ? value : old), stores the result back to address, and returns old. In other words, the memory location is left alone unless it was equal to the expected value specified by the caller, in which case it is updated with the new value.

A simple critical section called a “spin lock” can be built out of CAS, as follows:

void enter_spinlock( int *address )
{
    while atomicCAS( address, 0, 1 );
}

Assuming the spin lock’s value is initialized to 0, the while loop iterates until the spin lock value is 0 when the atomicCAS() is executed. When that happens, *address atomically becomes 1 (the third parameter to atomicCAS()) and any other threads trying to acquire the critical section spin waiting for the critical section value to become 0 again.

The thread owning the spin lock can give it up by atomically swapping the 0 back in:

void leave_spinlock( int *address )
{
    atomicExch( m_p, 0 );
}

On CPUs, compare-and-swap instructions are used to implement all manner of synchronization. Operating systems use them (sometimes in conjunction with the kernel-level thread context switching code) to implement higher-level synchronization primitives; CAS also may be used directly to implement “lock-free” queues and other data structures.

The CUDA execution model, however, imposes restrictions on the use of global memory atomics for synchronization. Unlike CPU threads, some CUDA threads within a kernel launch may not begin execution until other threads in the same kernel have exited. On CUDA hardware, each SM can context switch a limited number of thread blocks, so any kernel launch with more than MaxThreadBlocksPerSM*NumSMs requires the first thread blocks to exit before more thread blocks can begin execution. As a result, it is important that developers not assume all of the threads in a given kernel launch are active.

Additionally, the enter_spinlock() routine above is prone to deadlock if used for intra-block synchronization¹¹ – for which it is unsuitable in any case, since the hardware supports so many better ways for threads within the same block to communicate and synchronize with one another (shared memory and __syncthreads(), respectively).

Listing 5-8 shows the implementation of the cudaSpinlock class, which uses the algorithm listed above and is subject to the just-described limitations.

class cudaSpinlock {public:    cudaSpinlock( int *p );    void acquire();    void release();private:    int *m_p;}; inline __device__cudaSpinlock::cudaSpinlock( int *p ){    m_p = p;} inline __device__ voidcudaSpinlock::acquire( ){    while ( atomicCAS( m_p, 0, 1 ) );} inline __device__ voidcudaSpinlock::release( ){    atomicExch( m_p, 0 );}

Use of cudaSpinlock is illustrated in the spinlockReduction.cu sample, which computes the sum of an array of double values by having each block perform a reduction in shared memory, then using the spin lock to synchronize for the summation. Listing 5-9 gives the SumDoubles() function from this sample; note how adding the partial sum is performed only by thread 0 of each block.

__global__ voidSumDoubles(     double *pSum,     int *spinlock,     const double *in,     size_t N,     int *acquireCount ){    SharedMemory<double> shared;    cudaSpinlock globalSpinlock( spinlock );     for ( size_t i = blockIdx.x*blockDim.x+threadIdx.x;                  i < N;                  i += blockDim.x*gridDim.x ) {        shared[threadIdx.x] = in[i];        __syncthreads();        double blockSum = Reduce_block<double,double>( );        __syncthreads();         if ( threadIdx.x == 0 ) {            globalSpinlock.acquire( );            *pSum += blockSum;            __threadfence();            globalSpinlock.release( );        }    }}

5.2.12 Peer-to-Peer Access

Under certain circumstances, SM 2.0-class and later hardware can map memory belonging to other, similarly-capable GPUs. The following conditions apply:

• Unified virtual addressing (UVA) must be in effect.

• Both GPUs must be Fermi-class, and be based on the same chip.

• The GPUs must be on the same I/O hub.

Since peer-to-peer mapping is intrinsically a multi-GPU feature, it is described in detail in the multi-GPU chapter (Section 9.2).

5.2.13 Inter-Process Sharing (IPC)

Device memory allocations belong to the process that made them: a pointer returned by cudaMalloc() is meaningful only within its own address space, and handing the raw value to another process does not transfer access. The inter-process communication (IPC) API bridges that gap. cudaIpcGetMemHandle() turns a device pointer into a cudaIpcMemHandle_t, an opaque token that the owning process can pass to another by any ordinary means – a pipe, a socket, a file. The receiving process hands the token to cudaIpcOpenMemHandle(), which maps the same physical device memory into its own address space and returns a pointer it can dereference directly, with no copy. Events share the same way, through cudaIpcGetEventHandle() and cudaIpcOpenEventHandle(), so one process can wait on work another submitted. Both processes must be using the same device, and the importer releases its mapping with cudaIpcCloseMemHandle().

IPC is the mechanism behind multi-process pipelines and multi-process multi-GPU frameworks, where several processes cooperate on data resident on one GPU without routing it through the host. Two newer facilities meet the same need for memory the legacy handle cannot describe: the stream-ordered memory pools of Section 5.2.3 export a shareable handle with cudaMemPoolExportToShareableHandle(), and the virtual memory management API exports its physical allocations with cuMemExportToShareableHandle(). Both yield an operating-system handle – a file descriptor on Linux, an NT handle on Windows – that another process imports, the same handle machinery Section 3.11 uses to share memory with graphics and Vulkan.

5.2.14 ECC (Error Correcting Codes)

SM 2.x and later GPUs in the Tesla (i.e. server GPU) product line come with the ability to run with error correction. In exchange for a smaller amount of memory (since some memory is used to record some redundancy) and lower bandwidth, GPUs with ECC enabled can silently correct single-bit errors, and report double-bit errors.

ECC has the following characteristics.

• It reduces the amount of available memory by 12.5%. On a cg1.4xlarge instance in Amazon EC2, for example, it reduces the amount of memory from 3071MB to 2687MB.

• It makes context synchronization more expensive.

• Uncoalesced memory transactions are more expensive when ECC is enabled than otherwise.

ECC can be enabled and disabled using the nvidia-smi command-line tool (described in Section 4.6), or by using the NVML (NVIDIA Management Library).

When an uncorrectable ECC error is detected, synchronous error-reporting mechanisms will return cudaErrorECCUncorrectable (for the CUDA runtime) and CUDA_ERROR_ECC_UNCORRECTABLE (for the driver API).

The tradeoffs described above are a relic of the GDDR era, when error correction was carved out of memory that otherwise would have been available to the application. On GPUs with attached HBM (High Bandwidth Memory), ECC is not optional: at HBM’s stacked densities and data rates, error correction became a necessary feature for the memory to work at all. The HBM standards incorporate ECC directly – recent generations perform on-die error correction – so on HBM-equipped GPUs, ECC is always enabled and costs neither the capacity nor the bandwidth penalties enumerated above.

5.2.15 Texturing From Global Memory

For applications that cannot conveniently adhere to the coalescing constraints, the texture mapping hardware presents a satisfactory alternative. The hardware supports texturing from global memory (via cudaBindTexture()/cuTexRefSetAddress()), which has lower peak performance than coalesced global reads, but higher performance for less-regular access. The texture cache resources are also separate from other cache resources on the chip; a software coherency scheme is enforced by the driver invalidating the texture cache before kernel invocations that contain TEX instructions¹².

See Chapter 10 (Texturing) for details.

SM 3.x hardware added the ability to read global memory through the texture cache hierarchy without setting up and binding a texture reference. This functionality may be accessed with standard C++ language constructs: the const restrict keywords. Alternatively, you can use the __ldg() intrinsics defined in sm_35_intrinsics.h.