CUDA Cheatsheet

Memory Model

Use this CUDA reference while you build software engineering projects, review code for technical interview prep, or polish examples for a software engineer resume.

Memory Space Overview

Memory	Qualifier	Scope	Lifetime	Cached	Latency
Global	`__device__`	All threads + host	Application	L2 (L1 opt-in)	~200–800 cycles
Shared	`__shared__`	Block	Kernel	On-chip (SRAM)	~1–5 cycles
Local	(automatic, spilled)	Thread	Kernel	L1/L2	Same as global
Registers	(automatic)	Thread	Kernel	On-chip	~1 cycle
Constant	`__constant__`	All threads	Application	Const cache	~1 cycle (hit)
Texture	texture objects	All threads	Application	Texture cache	~10–50 cycles
Unified/Managed	`__managed__`	Host + device	Application	Migrates	Variable

Registers

Each thread gets its own private registers — fastest storage. Spills go to local memory (in global memory, slow).

__global__ void regExample(float* out, int n) {
    float acc = 0.0f;    // in a register
    int   idx = blockIdx.x * blockDim.x + threadIdx.x;  // register
    if (idx < n) out[idx] = acc + idx;
}
// Check register usage: nvcc --ptxas-options=-v

Gotcha: Arrays whose index is not a compile-time constant are likely placed in local memory (global memory), not registers.

Global Memory

Largest pool, accessible by all threads and the host. Coalesced accesses (contiguous, aligned, same warp) get the best bandwidth.

// COALESCED: stride-1, each thread reads next float
__global__ void goodAccess(float* a, float* b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) b[i] = a[i] * 2.0f;    // all 32 threads read 128 consecutive bytes
}

// UNCOALESCED: stride-32 (column of a matrix stored row-major)
__global__ void badAccess(float* mat, int stride, float* out) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    out[i] = mat[i * stride];    // 32 cache lines fetched for 32 floats — bad
}

Coalescing rules (CC ≥ 2.0)

A warp's memory access is coalesced into the minimum number of 128-byte cache-line transactions when: 1. All threads access within the same aligned 128-byte segment. 2. Addresses are consecutive (stride 1 = ideal).

`restrict` hint

__global__ void noAlias(float* __restrict__ a,
                        float* __restrict__ b,
                        float* __restrict__ c, int n) {
    // Compiler assumes a, b, c do not alias → enables more aggressive loads
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

Shared Memory

On-chip SRAM shared within a block. Declared with __shared__. Used for tiling, reuse, and inter-thread communication.

Static allocation

__global__ void staticShared(float* in, float* out, int n) {
    __shared__ float tile[256];   // size fixed at compile time

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    tile[threadIdx.x] = (i < n) ? in[i] : 0.0f;
    __syncthreads();

    // Reverse within block
    out[i] = tile[blockDim.x - 1 - threadIdx.x];
}

Dynamic allocation

// Kernel signature: extern __shared__ (unsized array)
__global__ void dynShared(float* in, float* out) {
    extern __shared__ float smem[];   // size set at launch
    smem[threadIdx.x] = in[blockIdx.x * blockDim.x + threadIdx.x];
    __syncthreads();
    out[blockIdx.x * blockDim.x + threadIdx.x] = smem[blockDim.x - 1 - threadIdx.x];
}

// Launch: pass shared bytes as 3rd chevron arg
dynShared<<<grid, 256, 256 * sizeof(float)>>>(d_in, d_out);

Bank conflicts

Shared memory is divided into 32 banks (4 bytes wide). Simultaneous accesses to the same bank by multiple threads in a warp serialize.

// NO conflict: each thread accesses a different bank
smem[threadIdx.x]        // thread t → bank (t % 32) — stride 1, all different

// CONFLICT (32-way): all threads access smem[0]
float v = smem[0];       // broadcast (no conflict on CC ≥ 2.0 for read)

// CONFLICT: stride 32 → all threads hit bank 0
float v = smem[threadIdx.x * 32];  // 32-way conflict!

// FIX: pad the array
__shared__ float smem[32][33];    // +1 column padding eliminates bank conflicts
float v = smem[threadIdx.y][threadIdx.x];

Constant Memory

64 KB total for __constant__ variables, broadcast-cached: if all threads in a warp read the same address, it costs 1 transaction.

__constant__ float cCoeffs[64];    // device-side declaration (file scope)

// Host sets values before kernel launch:
float h[64] = { /* ... */ };
cudaMemcpyToSymbol(cCoeffs, h, sizeof(h));

__global__ void applyFilter(float* data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) data[i] *= cCoeffs[i % 64];   // all threads reading same index = free
}

Gotcha: If different threads read different constant memory addresses in the same warp, the accesses serialize. Use global memory with L1 caching instead.

Texture Memory

Optimized for 2-D spatial locality, with built-in interpolation and boundary handling. Prefer texture objects (CC ≥ 3.0) over deprecated texture references.

// Create texture object
cudaArray_t cuArray;
cudaChannelFormatDesc chDesc = cudaCreateChannelDesc<float>();
cudaMallocArray(&cuArray, &chDesc, W, H);
cudaMemcpy2DToArray(cuArray, 0, 0, h_data, W*sizeof(float),
                    W*sizeof(float), H, cudaMemcpyHostToDevice);

cudaResourceDesc resDesc = {};
resDesc.resType         = cudaResourceTypeArray;
resDesc.res.array.array = cuArray;

cudaTextureDesc texDesc = {};
texDesc.addressMode[0]  = cudaAddressModeClamp;
texDesc.addressMode[1]  = cudaAddressModeClamp;
texDesc.filterMode      = cudaFilterModeLinear;   // bilinear
texDesc.readMode        = cudaReadModeElementType;
texDesc.normalizedCoords = 1;

cudaTextureObject_t texObj;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);

// In kernel
__global__ void sample(cudaTextureObject_t tex, float* out, int W, int H) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    float u = (x + 0.5f) / W;
    float v = (y + 0.5f) / H;
    out[y * W + x] = tex2D<float>(tex, u, v);
}

// Cleanup
cudaDestroyTextureObject(texObj);
cudaFreeArray(cuArray);

L1 / L2 Cache Control

// Per-access cache hints (CC ≥ 7.0, PTX level via intrinsics)
// Load with L1 cache bypass (streaming load):
float val = __ldcs(&ptr[i]);      // cache streaming (L2 only)
float val2 = __ldcg(&ptr[i]);     // cache global (L2, bypass L1)
float val3 = __ldca(&ptr[i]);     // cache at all levels (L1 + L2)
float val4 = __ldlu(&ptr[i]);     // last use — evict after load

// Store hints
__stcs(&ptr[i], val);   // streaming store (write-combine, bypass L1)
__stcg(&ptr[i], val);   // cache global
__stwb(&ptr[i], val);   // write-back all coherent levels

Global L1 cache config (pre-Volta)

// Kernel-level (deprecated on CC ≥ 7.5, ignored on Ampere+)
cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferShared);  // 48 KB shared, 16 KB L1
cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferL1);      // 48 KB L1, 16 KB shared
cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferEqual);   // 32/32

Unified / Managed Memory

Accessible from both host and device; the CUDA driver migrates pages automatically.

float* data;
cudaMallocManaged(&data, n * sizeof(float));  // allocate in unified memory

// Optionally prefetch before kernel to avoid page faults
cudaMemPrefetchAsync(data, n * sizeof(float), devId);

myKernel<<<grid, block>>>(data, n);
cudaDeviceSynchronize();

// Access from host freely after sync
printf("data[0] = %f\n", data[0]);

cudaFree(data);

Gotcha on Pascal+: Concurrent host+device access to the same managed allocation is allowed only on CC 6.x+ systems with hardware page migration (HMM/ATS). On older hardware, access from host while a kernel is running causes a segfault.

Memory Hierarchy Summary Table

Memory	Size (typical)	BW (typical A100)	Access
Register file	256 KB/SM	~20 TB/s	Per-thread
L1 / Shared	192 KB/SM (configurable)	~19 TB/s	Per-block
L2 cache	40 MB	~4 TB/s	All SMs
Global (HBM2e)	40–80 GB	2 TB/s	All threads + host
Constant cache	64 KB	~2 TB/s (all same addr)	All threads (read-only)
Texture cache	part of L1	~2 TB/s	All threads (read-only)