The CUDA Handbook ▸ Chapter 5. Memory

5.3 Constant Memory

Constant memory is optimized for read-only broadcast to multiple threads. As the name implies, the compiler uses constant memory to hold constants that couldn’t be easily computed or otherwise compiled directly into the machine code. Constant memory resides in device memory, but is accessed using different instructions that cause the GPU to access it using a special “constant cache.”

64K of memory is available to the compiler for constants, to use at its discretion.

Another 64K of memory is available to the developer, to be called out with the __constant__ keyword.

These limits are per-module (for driver API applications) or per-file (for CUDA runtime applications).

Naively, one might expect __constant__ memory to be analogous to the const keyword in C/C++, where it cannot be changed after initialization. But __constant__ memory can be changed, either by memory copies or by querying the pointer to __constant__ memory and writing to it with a kernel.

CUDA kernels must not write to __constant__ memory ranges that they may be accessing, because the constant cache is not kept coherent with respect to the rest of the memory hierarchy.

For a worked example of constant memory in a compute-intensive application, see the constant memory formulation of N-body in Section 14.4: because every thread reads the same body description at the same time, the broadcast-optimized constant cache serves the inner loop well, and holding the body descriptions in constant memory leaves shared memory free for other uses.

5.3.1 Host and Device __constant__ memory

Mark Harris describes an idiom that uses the predefined macro __CUDA_ARCH__ to maintain host and device copies of __constant__ memory that are conveniently accessed by both the CPU and GPU¹³:

__constant__ double dc_vals[2] = { 0.0, 1000.0 };
       const double hc_vals[2] = { 0.0, 1000.0 };

__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
    return dc_vals[i];
#else
    return hc_vals[i];
#endif
}

5.3.2 Accessing __constant__ Memory

Besides the accesses to constant memory implicitly caused by C/C++ operators, developers can copy to and from constant memory, and even query the pointer to a constant memory allocation.

CUDA Runtime

CUDA runtime applications can copy to and from __constant__ memory using cudaMemcpyToSymbol() and cudaMemcpyFromSymbol(), respectively.

The pointer to __constant__ memory can be queried with cudaGetSymbolAddress().

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

This pointer may be used to write to constant memory with a kernel, though developers must take care not to write to the constant memory while another kernel is reading it.

Driver API

Driver API applications can query the device pointer of constant memory using cuModuleGetGlobal(). The driver API does not include a special memory copy function, since it does not have the language integration of the CUDA runtime – applications must query the address with cuModuleGetGlobal() and then use that pointer to copy to/from the symbol.

The amount of constant memory used by a kernel may be queried with cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_CONSTANT_SIZE_BYTES).