CUDA Cheatsheet

Threads, Blocks, and Grids

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

Thread Hierarchy Overview

Grid
└── Block [0,0]  Block [1,0]  Block [2,0]  …
    └── Thread [0,0]  Thread [1,0]  …

Three levels: grid → block → thread. Each level is up to 3-dimensional. All dimensions are available in device code as built-in variables.

Built-in Variables

Variable	Type	Value
`threadIdx`	`uint3`	Thread index within its block (x, y, z)
`blockIdx`	`uint3`	Block index within the grid (x, y, z)
`blockDim`	`dim3`	Block dimensions in threads (x, y, z)
`gridDim`	`dim3`	Grid dimensions in blocks (x, y, z)
`warpSize`	`int`	Always 32 on all current hardware

__global__ void indexDemo() {
    // 1-D flat thread index
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    // 2-D: processing a matrix row/col
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    // 3-D: processing a volume
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;
}

dim3 Type

dim3 is a struct with .x, .y, .z; unspecified dimensions default to 1.

dim3 block(16, 16);       // 16×16×1 = 256 threads
dim3 grid(64, 64);        // 64×64×1 blocks

dim3 block3d(8, 8, 8);    // 512 threads/block (3-D)
dim3 grid3d(16, 16, 4);

myKernel<<<grid, block>>>(args);

1-D, 2-D, 3-D Index Patterns

1-D array

__global__ void process1D(float* data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) data[i] *= 2.0f;
}

2-D matrix (row-major, width W)

__global__ void process2D(float* mat, int H, int W) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < H && col < W)
        mat[row * W + col] += 1.0f;
}

// Launch
dim3 block(16, 16);
dim3 grid((W + 15) / 16, (H + 15) / 16);
process2D<<<grid, block>>>(dMat, H, W);

3-D volume

__global__ void process3D(float* vol, int D, int H, int W) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;
    if (x < W && y < H && z < D)
        vol[z*H*W + y*W + x] = 0.0f;
}

Warps

A warp is 32 consecutive threads within a block that execute in lockstep (SIMT). All current NVIDIA GPUs use warp size 32.

Block of 128 threads → 4 warps
  Warp 0: threads  0–31
  Warp 1: threads 32–63
  Warp 2: threads 64–95
  Warp 3: threads 96–127

Warp index helpers

__device__ int laneId()  { return threadIdx.x & 31; }          // 0–31
__device__ int warpId()  { return threadIdx.x >> 5; }          // warp within block

Warp-level intrinsics (CC ≥ 7.0 — use `__activemask()`)

__global__ void warpExample(int* data) {
    unsigned mask = __activemask();          // bitmask of active lanes

    int val = data[threadIdx.x];
    // Broadcast lane 0's value to all active lanes
    int broadcast = __shfl_sync(mask, val, 0);

    // Prefix sum within warp
    for (int offset = 1; offset < 32; offset <<= 1) {
        int n = __shfl_up_sync(mask, val, offset);
        if (laneId() >= offset) val += n;
    }
}

`__shfl_sync` family

Intrinsic	Description
`__shfl_sync(mask, var, srcLane)`	Broadcast from `srcLane`
`__shfl_up_sync(mask, var, delta)`	Shift down (higher lanes get lower lanes' data)
`__shfl_down_sync(mask, var, delta)`	Shift up (lower lanes get higher lanes' data)
`__shfl_xor_sync(mask, var, laneMask)`	Butterfly exchange

Supported types: int, unsigned, float, long long, unsigned long long, double.

Thread Blocks

All threads in a block: - Share shared memory (__shared__) - Can synchronize with __syncthreads() - Reside on the same SM (streaming multiprocessor) - Are limited to 1024 threads total

// Common block sizes (multiples of 32)
dim3 block1d(256);         // 1-D: 256 threads
dim3 block2d(16, 16);      // 2-D: 256 threads
dim3 block3d(8, 8, 4);     // 3-D: 256 threads

Gotcha: blockDim.x * blockDim.y * blockDim.z ≤ 1024. Exceeding this causes a launch error.

Grid-Stride Loop (best practice for large N)

Handles any N regardless of grid size; works well with persistent kernels.

__global__ void gridStride(float* data, int n) {
    int stride = gridDim.x * blockDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride)
        data[i] = sqrtf(data[i]);
}

// Can launch with ANY grid size; correctness is guaranteed
gridStride<<<256, 256>>>(d_data, n);

Thread Block Clusters (CC 9.0 / Hopper)

A cluster groups multiple blocks that can share distributed shared memory and synchronize cluster-wide.

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void __cluster_dims__(2, 1, 1)   // 2-block cluster
clusterKernel(int* data) {
    auto cluster = cg::this_cluster();
    auto block   = cg::this_thread_block();

    cluster.sync();   // cluster-wide barrier (cheaper than grid sync)
}

// Launch via cudaLaunchKernelEx
cudaLaunchConfig_t cfg = {};
cfg.gridDim  = 8;
cfg.blockDim = 128;
cudaLaunchAttribute attr;
attr.id = cudaLaunchAttributeClusterDimension;
attr.val.clusterDim = {2, 1, 1};
cfg.attrs    = &attr;
cfg.numAttrs = 1;
cudaLaunchKernelEx(&cfg, clusterKernel, data);

Occupancy

Occupancy = active warps / max warps per SM. Higher occupancy hides latency.

// Automatic occupancy tuning
int blockSize, minGridSize;
cudaOccupancyMaxPotentialBlockSize(
    &minGridSize, &blockSize, myKernel,
    0,    // dynamic shared mem per block
    0);   // block size limit (0 = no limit)

int gridSize = (n + blockSize - 1) / blockSize;
myKernel<<<gridSize, blockSize>>>(args);

Factors limiting occupancy

Resource	Effect
Register count per thread	Limits resident warps per SM
Shared memory per block	Limits resident blocks per SM
Block size	Too small → not enough warps; too large → register pressure
`__launch_bounds__`	Guides register allocator

// Query register / shared mem usage after compilation:
// nvcc --ptxas-options=-v myfile.cu