CUDA Cheatsheet

Kernels

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

Kernel Declaration

A kernel is a __global__ function. It must return void. Every parameter is copied by value into constant memory at launch time (up to 32,764 bytes total on CC ≥ 7.0 with CUDA ≥ 12.1; 4,096 bytes on older toolkits/hardware).

__global__ void addVectors(const float* __restrict__ a,
                           const float* __restrict__ b,
                           float* __restrict__ c,
                           int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

Kernel Launch Syntax

// Basic
addVectors<<<numBlocks, threadsPerBlock>>>(a, b, c, n);

// With dynamic shared memory and stream
addVectors<<<numBlocks, threadsPerBlock, sharedBytes, stream>>>(a, b, c, n);

Computing grid size from problem size:

int n = 1 << 20;
int blockSize = 256;
int gridSize  = (n + blockSize - 1) / blockSize;   // ceiling division
addVectors<<<gridSize, blockSize>>>(a, b, c, n);

Kernel Templates

template <int BLOCK_SIZE>
__global__ void tiledMatMul(const float* A, const float* B, float* C, int N) {
    __shared__ float tA[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float tB[BLOCK_SIZE][BLOCK_SIZE];
    // ...
}

// Instantiate at compile time
tiledMatMul<32><<<grid, dim3(32,32)>>>(dA, dB, dC, N);

Dynamic Parallelism (CDP2)

Kernels can launch child kernels from device code. Device-side cudaDeviceSynchronize() was deprecated in CUDA 11.6 and removed in CUDA 12 — a parent can no longer wait on its children. CDP2 orders child work with named device streams instead:

__global__ void parent(float* data, int n) {
    if (threadIdx.x == 0) {
        // Fire-and-forget: child starts immediately, no ordering with parent
        child<<<(n+31)/32, 32, 0, cudaStreamFireAndForget>>>(data, n);

        // Tail launch: child runs AFTER the entire parent grid completes
        finalize<<<1, 32, 0, cudaStreamTailLaunch>>>(data, n);
    }
}

Results a child produces are only guaranteed visible to grids launched after it (e.g. a tail launch) — chain dependent phases with cudaStreamTailLaunch.
Legacy CDP1 code that calls cudaDeviceSynchronize() in device code does not compile on CUDA 12+ toolkits.

Build flag required: nvcc -rdc=true -lcudadevrt

Cooperative Kernels (CC ≥ 6.0)

Allow grid-wide synchronization. All blocks must be resident simultaneously — use cudaOccupancyMaxActiveBlocksPerMultiprocessor to compute gridSize.

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

__global__ void coopKernel(int* data, int n) {
    auto grid = cg::this_grid();
    // ... phase 1 work ...
    grid.sync();   // grid-wide barrier
    // ... phase 2 work ...
}

// Launch via cudaLaunchCooperativeKernel
void* args[] = { &data, &n };
cudaLaunchCooperativeKernel((void*)coopKernel, gridSize, blockSize, args);

Variadic / Indirect Launch

// cudaLaunchKernel — type-erased launch (useful in libraries)
void* kernelArgs[] = { (void*)&d_in, (void*)&d_out, (void*)&n };
cudaLaunchKernel((const void*)myKernel,
                 gridDim, blockDim,
                 kernelArgs,
                 sharedMem, stream);

Kernel Attributes and Hints

// Query occupancy
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize,
                                   myKernel, 0, 0);

// Max active blocks per SM
int maxBlocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&maxBlocks,
    myKernel, blockSize, sharedBytesPerBlock);

// Set cache preference per kernel
cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferShared);
// Options: cudaFuncCachePreferNone | L1 | Shared | Equal

`__launch_bounds__`

Tells the compiler the max threads per block (and optionally min blocks per SM), enabling better register allocation.

__launch_bounds__(256, 2)          // maxThreadsPerBlock=256, minBlocksPerSM=2
__global__ void myKernel(...) { }

Predication and Warp Divergence

All threads in a warp execute the same instruction. Branches whose condition differs across a warp cause predication (serialized paths). Minimize divergence:

// Bad: each thread follows a different branch path
if (threadIdx.x % 2 == 0) doA();
else                        doB();

// Better: restructure to warp-aligned branches or use data to avoid divergence
// (threads 0-31 same branch, 32-63 same branch, etc.)

Printf in Kernels

#include <cstdio>

__global__ void debugKernel() {
    printf("Block %d Thread %d\n", blockIdx.x, threadIdx.x);
    // Buffer is flushed on cudaDeviceSynchronize / cudaStreamSynchronize
}

Output is buffered (default 1 MB). Increase with cudaDeviceSetLimit(cudaLimitPrintfFifoSize, 4<<20).

Assertion in Kernels

#include <cassert>

__global__ void safeKernel(float* data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    assert(i < n);   // triggers if condition false; compiled out in release
    data[i] *= 2.0f;
}

Compile with -DNDEBUG to strip all assert calls.

Kernel Execution Limits

Limit	Typical value
Max threads per block	1024
Max block dimension (x)	1024
Max block dimension (y)	1024
Max block dimension (z)	64
Max grid dimension (x)	2^31 − 1
Max grid dimension (y, z)	65535
Max kernel parameter size	32,764 bytes (CC ≥ 7.0, CUDA ≥ 12.1); 4,096 bytes otherwise
Warp size	32

// Check limit at runtime
cudaDeviceProp p; cudaGetDeviceProperties(&p, 0);
printf("Max threads/block: %d\n", p.maxThreadsPerBlock);
printf("Max grid x:        %d\n", p.maxGridSize[0]);