The CUDA Handbook ▸ Chapter 5. Memory

5.4 Local Memory

Local memory contains the stack for every thread in a CUDA kernel. It is used as follows:

• To implement the application binary interface (ABI), i.e. calling convention;

• to spill data out of registers;

• to hold arrays whose indices cannot be resolved by the compiler.

A kernel parameter passed by value can also land in local memory. Parameters arrive in a fast, read-only bank of constant memory, but if the code takes the address of a by-value parameter – or passes it by reference to a __device__ function – the compiler must first copy it into the thread’s own local memory so that the pointer has a per-thread object to refer to. For a large parameter structure, that copy is pure overhead. The __grid_constant__ qualifier, applied to a const by-value parameter, promises the compiler that every thread in the grid observes the same unmodified value, so its address resolves directly to the constant-memory copy and the per-thread duplication is skipped.

In early implementations of CUDA hardware, any use of local memory was the “kiss of death” – it slowed things down so much that developers were encouraged to take whatever measure was needed to get rid of the local memory usage. With the advent of an L1 cache in Fermi, these performance concerns are less urgent, provided the local memory traffic is confined to L1¹⁴.

Developers can make the compiler report the amount of local memory needed by a given kernel with the nvcc options: -Xptxas –v,abi=no

At runtime, the amount of local memory used by a kernel may be queried with cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES).

Paulius Micikevicius of NVIDIA gave a good presentation on how to determine whether local memory usage was impacting performance, and what to do about it. Register spilling can incur two costs: an increased number of instructions, and an increase in the amount of memory traffic.

The L1 and L2 performance counters can be used to determine if the memory traffic is impacting performance. Strategies to improve performance in this case include:

• At compile time, specify a higher limit in –maxregcount. By increasing the number of registers available to the thread, both the instruction count and the memory traffic will decrease. The __launch_bounds__ directive may be used to tune this parameter when the kernel is being compiled online by PTXAS.

• Use non-caching loads for global memory. (nvcc –Xptxas -dlcm=cg)

• Increase the L1 size to 48K. (Call cudaFuncSetCacheConfig() or cudaDeviceSetCacheConfig()).

When launching a kernel that uses more than the default amount of memory allocated for local memory, the CUDA driver must allocate a new local memory buffer before the kernel can launch. As a result, the kernel launch may take extra time; may cause unexpected CPU/GPU synchronization; and, if the driver is unable to allocate the buffer for local memory, may fail¹⁵. By default, the CUDA driver will free these larger local memory allocations after the kernel has launched. This behavior can be inhibited by specifying the CU_CTX_RESIZE_LMEM_TO_MAX flag to cuCtxCreate(), or calling cudaDeviceSetFlags() with the cudaDeviceLmemResizeToMax flag set.

It is not difficult to build a templated function that illustrates the “performance cliff” when register spills occur. The templated GlobalCopy() kernel of Listing 5-10 implements a simple memcpy routine that uses a local array temp to stage global memory references. The template parameter n specifies the number of elements in temp, and hence the number of loads and stores to perform in the inner loop of the memory copy.

As a quick review of the SASS microcode emitted by the compiler will confirm, the compiler can keep temp in registers until n becomes too large.

template<class T, const int n> __global__ voidGlobalCopy( T *out, const T *in, size_t N ){    T temp[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;            temp[j] = in[index];        }        for ( int j = 0; j < n; j++ ) {            size_t index = i+j*blockDim.x;            out[index] = temp[j];        }    }    // 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++ ) {        for ( int j = 0; j < n; j++ ) {            size_t index = i+j*blockDim.x;            if ( index<N ) temp[j] = in[index];        }        for ( int j = 0; j < n; j++ ) {            size_t index = i+j*blockDim.x;            if ( index<N ) out[index] = temp[j];        }    }}

Listing 5-11 shows an excerpt of the output from globalCopy.cu on a GK104 GPU: the copy performance of 64-bit operands only. The degradation in performance due to register spilling becomes obvious in the row corresponding to a loop unroll of 12, where the delivered bandwidth decreases from 117 GB/s to less than 90 GB/s, and degrades further to under 30 GB/s as the loop unroll increases to 16.

Table 5-9 summarizes the register and local memory usage for the kernels corresponding to the unrolled loops. The performance degradation of the copy corresponds to the local memory usage. In this case, every thread always spills in the inner loop; presumably, the performance wouldn’t degrade so much if only some of the threads were spilling (for example, when executing a divergent code path).

Table 5-9. globalCopy() register and local memory usage.

Operand size: 8 bytesInput size: 16M operands                      Block SizeUnroll	32	64	128	256	512	maxBW	maxThreads1	75.57	102.57	116.03	124.51	126.21	126.21	5122	105.73	117.09	121.84	123.07	124.00	124.00	5123	112.49	120.88	121.56	123.09	123.44	123.44	5124	115.54	122.89	122.38	122.15	121.22	122.89	645	113.81	121.29	120.11	119.69	116.02	121.29	646	114.84	119.49	120.56	118.09	117.88	120.56	1287	117.53	122.94	118.74	116.52	110.99	122.94	648	116.89	121.68	119.00	113.49	105.69	121.68	649	116.10	120.73	115.96	109.48	99.60	120.73	6410	115.02	116.70	115.30	106.31	93.56	116.70	6411	113.67	117.36	111.48	102.84	88.31	117.36	6412	88.16	86.91	83.68	73.78	58.55	88.16	3213	85.27	85.58	80.09	68.51	52.66	85.58	6414	78.60	76.30	69.50	56.59	41.29	78.60	3215	69.00	65.78	59.82	48.41	34.65	69.00	3216	65.68	62.16	54.71	43.02	29.92	65.68	32

The measurements in Listing 5-11 and Table 5-9 were taken on a Kepler-class GK104, whose kernels were capped at 63 registers; once the staging array outgrew that budget, at an unroll of 12, the overflow spilled to local memory and copy bandwidth fell off a cliff. Ampere and later GPUs carry a much larger register file – up to 255 registers per thread – so the same globalCopy.cu kernel no longer spills across this range. On a GeForce RTX 3060, the 64-bit copy climbs from 10 registers at an unroll of 1 to 126 at an unroll of 16 with no local memory whatsoever, and its best-case bandwidth holds flat at roughly 305 GB/s throughout. The performance cliff is still real – it simply takes a larger register footprint, or an explicit --maxrregcount limit, to provoke it.