To build a good enough mental model of CUDA, it is useful to think in terms of how memory and compute are organized.

1. Memory types in GPU

   • Global memory (HBM): ~1-3 TB/s bandwidth, ~400 cycles latency.
   • Shared memory (SRAM): ~20 TB/s bandwidth, ~30 cycles latency.
   • Registers (on the compute unit itself): ~1 cycle latency.

2. Compute

   • A thread is the smallest unit of execution; it runs one instance of the kernel code.
   • Threads are grouped into blocks.
   • Blocks are grouped into a grid.
   • Grid and block dimensions are chosen at kernel launch time, e.g. my_kernel<<<gridDim, blockDim>>>(...).

Memory: HBM is expensive, SRAM is cheap

• Reading and writing a single data point thrice takes 2400 cycles.

// Don't worry about the `__global__` term in the line below
// It will be explained later
__global__ void naive_kernel(float* A, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= N) return;
    for (int k = 0; k < N_OPS; k++) {
        // ~400 cycles from HBM
        float x = A[i];
        // ~400 cycles back to HBM
        A[i]    = f(x);
    }
}
// Cost per thread: N_OPS * ~800 cycles.

Reading from and writing to SRAM is much faster. Loading a piece of data from HBM to SRAM and then doing the three read/writes between SRAM and the register followed by writing the data back to the HBM is much faster. Visualize it step by step:

• HBM → SRAM

// declare a 256-element SRAM buffer
__shared__ float smem[256];
// ~400 cycles ONCE
smem[tid] = A[i];

• SRAM ↔ register

for (int k = 0; k < N_OPS; k++) {
    // read smem, compute in register, write smem: ~30 cycles (SRAM)
    smem[tid] = f(smem[tid]);
}

• HBM ← SRAM

// ~400 cycles ONCE
A[i] = smem[tid];

The second data movement is much faster.

The goal of every CUDA kernel: minimize HBM traffic.

Latency numbers are rounded; they vary Ampere->Hopper, cached vs uncached. The order-of-magnitude ratio (~10-15x) is what matters.

Compute: Grid → Block → Thread

A kernel is the function you write to run on the GPU (marked __global__ in CUDA). When a kernel is launched (i.e. when the function is called), the GPU runs one copy of it on every single thread, in parallel. Each thread executes the exact same code on a different piece of data.

• Show a thread: a single thread runs the kernel function on one piece of data.

• Show a block: a group of threads is arranged in a block A block can be , or dimensional. Each thread in the block has a threadIdx (its position within the thread).

• Show the grid: Blocks are grouped together in a unit called grid. A grid can also be , or dimensional. Each block in the grid has a blockIdx (its position within the grid).

• blockDim and gridDim: the shape of a block, and the shape of the grid. blockDim.x/.y count the threads along each axis of one block; gridDim.x/.y count the blocks along each axis of the grid. Both are fixed at launch and identical for every thread.

• Calculating global index: consider the block (2,3) and the thread (3,1) in that block. The global index can be calculated like so:

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;

A block only knows where a thread sits inside itself - threadIdx restarts at 0 in every block. The global index is that thread's position across the whole grid, so no two threads in the launch share one. That is what lets each thread pick out its own element of the data: A[row][col].

• Source code for defining and launching a kernel.

// Host side: choose the shapes and launch.
// blockDim: threads per block (4 wide, 2 tall = 8 threads)
// gridDim: blocks in the grid (6 wide, 4 tall = 24 blocks)
dim3 blockDim(4, 2);
dim3 gridDim(6, 4);
my_kernel<<<gridDim, blockDim>>>(A, width, height);

// Device side: every one of the 192 threads runs this, and each
// one lands on a different (row, col).
__global__ void my_kernel(float* A, int width, int height) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    // The grid may overhang the data, so drop the extra threads.
    if (col >= width || row >= height) return;

    A[row * width + col] = f(A[row * width + col]);
}

Scene C: how the grid lands on the chip

• Streaming multiprocessors (SMs): each SM is an independent processing unit on the GPU, with its own SRAM. The scheduler hands blocks from the grid to whichever SM has room to run them.

• The blocks are assigned across the SMs.

• Each SM has its own SRAM which is shared across the blocks.

A block's threads share that SM's SRAM; blocks on different SMs do not share SRAM and must go through HBM to exchange data.