I read Programming Massively Parallel Processors (Kirk & Hwu) cover to cover. These are the notes I kept, first as a private revision set, then as Abi’s Concise Notes on gpuengineering.com. The PDF is still there if you want the original.
This is the same material as a post: the four truths that make GPUs strange, then enough CUDA to write a kernel, then the hierarchy, the memory model, and the patterns you actually need.
For decades, software had a free lunch. Every 18 months a new CPU arrived and sequential code got faster. Around 2003 that stopped. Energy and heat killed single-thread frequency scaling. The new frontier was parallel programming, and the most powerful parallel processors, GPUs, run on rules that are counter-intuitive if you were trained on sequential machines.
Four truths
1. GPUs do not reduce slowness. They hide it.
An individual GPU operation is often slower than the CPU equivalent. The secret is not latency. It is throughput.
CPUs are latency-oriented. Huge caches, branch prediction, a lot of silicon spent making one thread finish sooner.
GPUs are throughput-oriented. They tolerate latency. When one warp waits on memory, the scheduler swaps in another warp that is ready. Arithmetic units stay busy. Smaller caches exist mostly to save bandwidth (fetch once, reuse across threads) not to make a single access fast.
A post office with one clerk: the latency clerk waits while you fill a form. The throughput clerk sends you aside and serves the next person. More customers per hour. Your personal wait might be longer.
Reducing latency is much more expensive, in power and area, than increasing throughput. That is why the hardware looks the way it does.
2. Your “parallel” threads are not truly independent.
Threads run in warps of 32. Same instruction, same time. SIMD. Efficient when everyone does the same thing. Disastrous when they don’t.
That failure is control divergence. An if / else where half the warp takes each path: the hardware serializes. First the if (the other half idle), then the else (the first half idle). Ten instructions plus eight is eighteen for the warp, even though each thread only needed one path.
You must know which threads share a warp and keep them on the same path.
3. To go faster, sometimes you must do more work.
On a CPU, O(N) beats O(N log N). On a GPU, a work-inefficient algorithm can still win.
Sequential scan of N elements: ~N additions, N steps, not parallelizable. Kogge–Stone parallel scan: O(N log N) work, but log N parallel steps. For N = 1024 that is ~1,024 sequential steps versus 10 parallel steps, and nearly 10,000 additions. More work. Far less time. The metric that matters is the longest chain of dependent calculations, not total FLOPs.
4. Scalability is designing for independence.
Write the kernel once. It should run on a laptop GPU or a datacenter GPU without changing a line. That is transparent scalability.
The trick is independent thread blocks. Threads inside a block may share memory and synchronize. Blocks must not wait on other blocks. The runtime can then run many blocks at once on a fat GPU, or one after another on a thin one. Order does not matter. The result is still correct.
The chain: hide latency with massive thread counts → lockstep warps (so they are not independent) → work-inefficient algorithms become fastest → independence between blocks makes the system scalable.
Parallelism is not “do more things at once.” It is a different relationship between software and hardware.
Why use a GPU at all?
Think of the CPU as a skilled manager, sequential, branchy, low latency. The GPU is an army that does one simple thing on millions of data points.
| CPU | GPU | |
|---|---|---|
| Goal | Minimize latency for one task | Maximize throughput for many |
| Cores | Few, powerful | Many, smaller, cheaper |
| Analogy | Specialists | An army doing the same job |
Heterogeneous computing assigns each processor the work it is built for. CUDA is how you orchestrate that.
The CUDA model
Host (CPU) starts main, owns the story, launches work. Device (GPU) runs the kernel: a function executed by thousands or millions of threads.
Host and device have separate memory. The usual five steps:
cudaMalloc: allocate on the devicecudaMemcpyhost → device- Launch the kernel with
<<<grid, block>>> cudaMemcpydevice → hostcudaFree
Thread hierarchy
- Thread: one execution of the kernel. GPU threads are cheap (a few cycles to create). CPU threads are not.
- Warp: 32 threads. The real scheduling unit. Lockstep.
- Block: threads that may share memory and
__syncthreads(). Must be reachable by every thread in the block or you deadlock. Blocks do not synchronize with other blocks, that is what makes the code scale. - Grid: all blocks for one kernel launch.
Built-ins: threadIdx (in the block), blockIdx (in the grid), blockDim, gridDim.
Vector add
__global__ void vecAddKernel(float* A, float* B, float* C, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
C[i] = A[i] + B[i];
}
}Launch: vecAddKernel<<<ceil(n / 256.0), 256>>>(A_d, B_d, C_d, n).
The index line is the whole trick. Block 2, thread 5, 256 threads per block: 2 * 256 + 5 = 517. Every thread gets one element. No collisions.
Host side is the five-step dance: malloc three buffers, copy A and B up, launch, copy C down, free.
Memory
| Global | Shared | |
|---|---|---|
| Where | Off-chip DRAM | On-chip |
| Speed | Slow, high latency | Fast |
| Size | Gigabytes | Kilobytes per block |
| Who | Any thread in the grid | Only that block |
| Life | The application | The block |
Global is where the host can write. Shared is a user-managed cache: load a tile from global, reuse it, write back. That is how you stop paying DRAM for every multiply.
Registers are fastest and per-thread. Constant memory is small, cached, read-only, grid-wide.
Two performance laws
Control divergence. Threads in a warp take different branches → both paths run, half idle each time. Structure data so a warp decides the same way. A few divergent warps at a boundary are fine if thousands are not.
Memory coalescing. A warp’s 32 loads become one transaction if they hit a contiguous, aligned chunk of global memory. Consecutive threadIdx.x must hit consecutive addresses. Scatter is death.
The other two you will hear constantly: occupancy (enough warps resident to hide latency) and tile + reuse (shared memory so you do not refetch).
Parallel patterns
Reduction: many → one (sum, max). A tree. Watch divergence as threads go idle; rearrange work so warps stay full.
Scan (prefix sum): running totals. Inclusive or exclusive. Kogge–Stone: more work, fewer steps. Brent–Kung: work-efficient, O(N), different parallelism trade-off.
Convolution: sliding window, high reuse. Naive: every output reloads the neighborhood from global. Fix: tiling into shared memory.
Histogram: output interference. Atomics or privatization.
Sparse / graph: irregular access, load imbalance. Layout is the algorithm.
Deep learning: mostly GEMM. That is why Tensor Cores and CUTLASS exist.
How to look at a slow kernel
USE method, for every resource: utilization, saturation, errors.
Then ask:
- Coalesced global access?
- Tiling / caching before you invent a new kernel?
- Occupancy high enough, or registers spilling?
- Divergence?
- Work balanced, or a few threads doing everything?
- Atomics / privatization for write conflicts?
- Synchronization only where you must, barriers idle everyone else?
Profile. Then intervene. Then profile again.
What to take
CPU vs GPU is latency vs throughput. CUDA is host managing device memory and launching a grid of blocks. The index formula is how threads don’t step on each other. Warps, coalescing, shared-memory tiles, and independent blocks are why the same kernel can be 10× or 100× apart.
This is not a niche skill anymore. It is how you reason about the hardware under an LLM.
Practical guidance
- Learn warps, the memory hierarchy, and occupancy before writing a “clever” kernel.
- Check coalescing and divergence before you invent a new algorithm.
- Profile, intervene, profile again. USE: utilization, saturation, errors.
- Then return to the inference path, this chapter buys milliseconds you have already named.
Limitations
These notes follow Kirk & Hwu. They are a reading record, not original architecture research, and they predate a specific serving-engine bake-off. NVIDIA warp width and memory names are the vocabulary here; AMD and other stacks rhyme but do not copy the names.
Original notes: PDF. Book: Kirk & Hwu, Programming Massively Parallel Processors.