CUDA Cheatsheet

Optimization

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

Optimization Workflow

Profile first (nsys, ncu) — never guess.
Fix memory bandwidth bottlenecks (coalescing, shared memory tiling).
Improve occupancy (register pressure, shared memory size).
Reduce warp divergence.
Hide latency (instruction-level parallelism, async copies).
Consider algorithmic changes last.

# Quick profiling
nsys profile --stats=true ./myapp        # timeline + API summary
ncu --set full -o report ./myapp         # detailed kernel metrics
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./myapp

Memory Coalescing

The single biggest bandwidth win. A warp's 32 threads accessing consecutive 4-byte words → 1 × 128-byte cache line transaction.

// GOOD — stride-1 (coalesced)
__global__ void goodKernel(float* a, float* b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) b[i] = a[i] + 1.0f;
}

// BAD — stride-N (column access of row-major matrix)
__global__ void badKernel(float* mat, float* out, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    out[i] = mat[i * N];   // 32 separate cache lines
}

// FIX: transpose the matrix first, or change the data layout, or use shared mem transpose

Shared memory transpose (classic coalescing fix)

__global__ void transpose(float* out, const float* in, int N) {
    __shared__ float tile[32][33];   // +1 padding eliminates bank conflicts

    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;

    if (x < N && y < N) tile[threadIdx.y][threadIdx.x] = in[y * N + x];
    __syncthreads();

    x = blockIdx.y * 32 + threadIdx.x;
    y = blockIdx.x * 32 + threadIdx.y;
    if (x < N && y < N) out[y * N + x] = tile[threadIdx.x][threadIdx.y];
}

Occupancy Tuning

High occupancy hides memory latency by giving the warp scheduler more warps to switch between.

// Auto-tune block size for maximum occupancy
int blockSize, minGridSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, myKernel, 0, 0);
int gridSize = (n + blockSize - 1) / blockSize;
myKernel<<<gridSize, blockSize>>>(args);

// Check achieved occupancy
int maxActiveBlocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &maxActiveBlocks, myKernel, blockSize, 0);

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
float occupancy = (float)(maxActiveBlocks * blockSize)
                / prop.maxThreadsPerMultiProcessor;
printf("Achieved occupancy: %.1f%%\n", occupancy * 100.0f);

`__launch_bounds__` reduces register spilling

// Tells compiler: never use more than 128 threads/block, keep ≥2 blocks/SM
__launch_bounds__(128, 2)
__global__ void myKernel(float* data) { /* ... */ }

Register limit with `--maxrregcount`

nvcc --maxrregcount=32 myfile.cu   # force ≤32 regs/thread (may cause spills)

Shared Memory Tiling (Matrix Multiply)

Classic pattern for reusing global memory reads.

template <int TILE>
__global__ void matMul(const float* A, const float* B, float* C, int N) {
    __shared__ float sA[TILE][TILE];
    __shared__ float sB[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float acc = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
        int tCol = t * TILE + threadIdx.x;
        int tRow = t * TILE + threadIdx.y;

        sA[threadIdx.y][threadIdx.x] = (row < N && tCol < N) ? A[row * N + tCol] : 0.0f;
        sB[threadIdx.y][threadIdx.x] = (tRow < N && col < N) ? B[tRow * N + col] : 0.0f;
        __syncthreads();

        for (int k = 0; k < TILE; k++) acc += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        __syncthreads();
    }

    if (row < N && col < N) C[row * N + col] = acc;
}

Warp Divergence Reduction

// Warp-aligned branch (no divergence if condition splits on warp boundary)
if ((threadIdx.x / 32) % 2 == 0) doA();
else doB();

// Predication (compiler does it automatically for small branches)
float val = (threadIdx.x < 16) ? a : b;   // predicated select — no branch

Replace dynamic branches with lookup tables. Initialized __constant__ / __device__ arrays must be declared at file scope, not inside the kernel:

__constant__ float kLut[4] = {1.0f, 2.0f, 4.0f, 8.0f};   // file scope

__global__ void useLut(float* out) {
    out[threadIdx.x] = kLut[threadIdx.x % 4];   // no branch, broadcast-cached
}

Instruction-Level Parallelism (ILP)

Unroll loops to issue multiple independent instructions per thread, hiding latency.

// Manual ILP=4
__global__ void saxpy4(float* y, const float* x, float a, int n) {
    int i = (blockIdx.x * blockDim.x + threadIdx.x) * 4;
    if (i + 3 < n) {
        float x0 = x[i], x1 = x[i+1], x2 = x[i+2], x3 = x[i+3];
        y[i]   = a * x0 + y[i];
        y[i+1] = a * x1 + y[i+1];
        y[i+2] = a * x2 + y[i+2];
        y[i+3] = a * x3 + y[i+3];
    }
}
// Or use #pragma unroll
__global__ void saxpyU(float* y, const float* x, float a, int n) {
    int base = (blockIdx.x * blockDim.x + threadIdx.x) * 4;
    #pragma unroll 4
    for (int k = 0; k < 4; k++) {
        int i = base + k;
        if (i < n) y[i] = a * x[i] + y[i];
    }
}

Fast Math

# Enable fast (approximate) math for the whole file
nvcc --use_fast_math myfile.cu
# Equivalent to: -ftz=true -prec-div=false -prec-sqrt=false -fmad=true

Fast intrinsics vs. standard functions

Standard	Fast intrinsic	Notes
`sinf(x)`	`__sinf(x)`	~22-bit, ~4x faster
`cosf(x)`	`__cosf(x)`	~22-bit
`expf(x)`	`__expf(x)`	~22-bit
`logf(x)`	`__logf(x)`	~22-bit
`powf(x,y)`	`__powf(x,y)`	large ULP error
`fdividef(a,b)`	`__fdividef(a,b)`	24-bit
`sqrtf(x)`	`__fsqrt_rn(x)`	or use `rsqrtf` + multiply
`1.0f/sqrtf(x)`	`rsqrtf(x)`	reciprocal sqrt

// FMA (fused multiply-add) — enabled automatically with -fmad
float r = __fmaf_rn(a, b, c);   // r = a*b + c, single rounding

Asynchronous Memory Copy (Ampere, CC ≥ 8.0)

cuda::memcpy_async copies from global to shared memory without staging through registers. For a block-scope collective, the pipeline must be created with cuda::make_pipeline(block, &state) over a __shared__ state, and every argument to memcpy_async must be uniform across the block. Prefetch stage k+1 before waiting on stage k to actually overlap copy with compute.

#include <cooperative_groups.h>
#include <cuda/pipeline>
namespace cg = cooperative_groups;

// Grid-stride, double-buffered load → compute.
// Assumes n is a multiple of blockDim.x * gridDim.x (blockDim.x = 256).
__global__ void pipelinedKernel(const float* __restrict__ g,
                                float* __restrict__ out, int n) {
    __shared__ float smem[2][256];                 // 2 stages
    auto block = cg::this_thread_block();

    __shared__ cuda::pipeline_shared_state<cuda::thread_scope_block, 2> state;
    auto pipe = cuda::make_pipeline(block, &state);   // block-scope pipeline

    const int tileSize = block.size();
    const int stride   = gridDim.x * tileSize;
    const int first    = blockIdx.x * tileSize;    // uniform per block

    // Prime stage 0
    pipe.producer_acquire();
    cuda::memcpy_async(block, smem[0], g + first, sizeof(float) * tileSize, pipe);
    pipe.producer_commit();

    int stage = 0;
    for (int base = first; base < n; base += stride, stage++) {
        int next = base + stride;
        if (next < n) {                            // prefetch next tile first
            pipe.producer_acquire();
            cuda::memcpy_async(block, smem[(stage + 1) & 1], g + next,
                               sizeof(float) * tileSize, pipe);
            pipe.producer_commit();
        }
        pipe.consumer_wait();                      // wait only for current tile
        out[base + block.thread_rank()] = smem[stage & 1][block.thread_rank()] * 2.0f;
        pipe.consumer_release();
    }
}

Warp Reduction Patterns

// Fast warp reduction using shuffle (no shared memory needed)
__device__ float warpReduceSum(float val) {
    unsigned mask = __activemask();
    for (int offset = 16; offset >= 1; offset >>= 1)
        val += __shfl_down_sync(mask, val, offset);
    return val;   // lane 0 holds the sum
}

// Block reduction (warp → shared → warp)
__device__ float blockReduceSum(float val) {
    __shared__ float partial[32];   // one per warp
    int lane   = threadIdx.x & 31;
    int warpId = threadIdx.x >> 5;

    val = warpReduceSum(val);
    if (lane == 0) partial[warpId] = val;
    __syncthreads();

    val = (threadIdx.x < (blockDim.x >> 5)) ? partial[lane] : 0.0f;
    if (warpId == 0) val = warpReduceSum(val);
    return val;
}

Compiler Pragmas and Hints

// Unroll a loop by factor N (or fully if no arg)
#pragma unroll 4
for (int i = 0; i < 16; i++) acc += data[i];

// Disable unrolling
#pragma unroll 1
for (int i = 0; i < n; i++) acc += data[i];

// Hint for loop carried dependencies
#pragma nounroll
for (int i = 0; i < n; i++) { data[i] = prev; prev = data[i]; }

Vectorized Loads

Load 128-bit words in one instruction (4× float or 2× double).

__global__ void vectorLoad(const float4* __restrict__ in, float4* out, int n4) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) {
        float4 v = in[i];     // single 128-bit load
        v.x *= 2.0f; v.y *= 2.0f; v.z *= 2.0f; v.w *= 2.0f;
        out[i] = v;           // single 128-bit store
    }
}
// Call with n4 = n / 4; requires n divisible by 4 and aligned pointer

Tensor Cores — WMMA (`<mma.h>`, CC ≥ 7.0)

Warp-level matrix multiply-accumulate: one warp cooperatively computes a D = A×B + C tile. Compile with -arch=sm_70 or newer. half supports the shapes 16×16×16, 32×8×16, 8×32×16; __nv_bfloat16 and tf32 need CC ≥ 8.0.

#include <mma.h>
#include <cuda_fp16.h>
using namespace nvcuda;

// One warp computes one 16x16 tile of C = A*B.
// A (MxK), B (KxN), C (MxN), all column-major; M, N, K multiples of 16.
__global__ void wmmaGemm(const half* A, const half* B, float* C,
                         int M, int N, int K) {
    int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / warpSize;  // tile row
    int warpN =  blockIdx.y * blockDim.y + threadIdx.y;              // tile col

    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::col_major> aFrag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> bFrag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> accFrag;
    wmma::fill_fragment(accFrag, 0.0f);

    for (int k = 0; k < K; k += 16) {
        wmma::load_matrix_sync(aFrag, A + warpM * 16 + k * M, M);   // lda = M
        wmma::load_matrix_sync(bFrag, B + k + warpN * 16 * K, K);   // ldb = K
        wmma::mma_sync(accFrag, aFrag, bFrag, accFrag);
    }
    wmma::store_matrix_sync(C + warpM * 16 + warpN * 16 * M, accFrag, M,
                            wmma::mem_col_major);
}

All 32 threads of the warp must call each *_sync collective with the same arguments.
For production GEMMs use cuBLAS or CUTLASS — they also reach the newer Hopper (wgmma) and Blackwell (tcgen05) Tensor Core paths that WMMA does not expose.

Profiling Metrics Quick Reference

Metric (ncu)	What it tells you
`sm__throughput.avg.pct_of_peak`	Overall SM utilization
`l1tex__t_bytes.sum`	L1/Tex cache traffic
`lts__t_bytes.sum`	L2 cache traffic
`dram__bytes.sum`	DRAM (HBM) bandwidth used
`smsp__warps_issue_stalled_long_scoreboard.avg`	Memory latency stalls
`smsp__warps_issue_stalled_short_scoreboard.avg`	Instruction latency stalls
`smsp__warps_issue_stalled_wait.avg`	Synchronization stalls
`sm__warps_active.avg.pct_of_peak`	Achieved occupancy
`smsp__sass_thread_inst_executed_op_ffma.sum`	FP32 FMA count

ncu --metrics l1tex__t_bytes.sum,dram__bytes.sum -o prof ./myapp
ncu -i prof.ncu-rep --page raw    # inspect later