Shared memory is used to exchange data between CUDA threads within a block. Physically, it is implemented with a per-SM memory that can be accessed very quickly. In terms of speed, shared memory is perhaps 10x slower than register accesses, but 10x faster than accesses to global memory. As a result, shared memory is often a critical resource to reduce the external bandwidth needed by CUDA kernels.
Since developers explicitly allocate and reference shared memory, it can be thought of as a “software managed” cache or “scratchpad” memory.
Developers can request different cache configurations at both the
kernel and the device level:
cudaDeviceSetCacheConfig()/cuCtxSetCacheConfig() specify the preferred
cache configuration for a CUDA device, while cudaFuncSetCacheConfig()/
cuFuncSetCacheConfig() specify the preferred cache configuration for a
given kernel. If both are specified, the per-kernel request takes
precedence, but in any case the requirements of the kernel may override
the developer’s preference.
Kernels that use shared memory typically are written in three phases:
Load shared memory and __syncthreads(),
Process shared memory and __syncthreads(),
Write results.
Developers can make the compiler report the amount of shared memory
used by a given kernel with the nvcc options:
-Xptxas --v,abi=no
At runtime, the amount of shared memory used by a kernel may be
queried with
cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES).
Any shared memory declared in the kernel itself is automatically allocated for each block when the kernel is launched. If the kernel also includes an unsized declaration of shared memory, the amount of memory needed by that declaration must be specified when the kernel is launched.
If there is more than one extern __shared__ memory declaration, they
are aliased with respect to one another – so this declaration:
extern __shared__ char sharedChars[];
extern __shared__ int sharedInts[];
enables the same shared memory to be addressed as 8- or 32-bit integers, as needed. One motivation for using this type of aliasing is to use wider types when possible to read and write global memory, while using the narrow ones for kernel computations.
Note that if you have more than one kernel that uses unsized shared memory, they must be compiled in separate files.
Older code sometimes dropped the __syncthreads() barriers between the
final steps of a shared-memory reduction once it narrowed to a single
warp, declaring the shared-memory pointers volatile in the
belief that doing so made the idiom safe. It does not:
volatile is not a memory barrier, and the idiom also relied
on the lockstep warp execution that Independent Thread Scheduling
(Section 7.5) removed. This deprecated technique, and its correct
replacement using __syncwarp(), are discussed in Section 7.5.
It is valid – and often convenient – to use pointers to refer to
shared memory. Example kernels that use this idiom include the reduction
kernels in Chapter 12 (e.g., Listing 12-3) and the scanBlock() kernel
(Listing 13-3).