Foundations — CUDA programming model basics

This is the D1 Foundations deep-dive for CUDA programming model basics. Before you touch a kernel, you must be fluent in every symbol the parent note throws at you. We build each one from nothing: plain words → a picture → why the topic can't live without it.

1. Host and Device — two separate worlds

The picture: imagine two islands. The Host island has a few clever workers. The Device island has thousands of simple workers. A single narrow bridge — the PCIe bus — connects them. Nothing on one island can reach across and grab a book from the other island's library directly; you must ship it across the bridge.

Why the topic needs it: every CUDA program is a story of moving data across that bridge, running work on the Device island, then shipping results back. The parent note's steps "Copy Host → Device" and "Copy Device → Host" only make sense once you see there are genuinely two memories that cannot see each other. See Memory Hierarchy for the full map of the Device's own internal libraries.

2. Thread — the smallest worker

The picture: a single dot on the warehouse floor, holding a slip of paper that says "I handle element 7." It does not know about element 6 or 8 — only its own.

Why the topic needs it: the whole point of a GPU is to stop writing loops. Instead of "do this N times, one after another," you say "here is the work for ONE element — now spawn N threads and let them all run at once." The thread is the atom of that idea.

3. Block — a team of threads

The picture: cluster the dots into square huddles of, say, 256 workers. Each huddle is a block. Workers inside one huddle can pass notes to each other cheaply; workers in different huddles cannot.

Why the topic needs it: cooperation is expensive at scale. Letting all 10,000 workers talk to all others is chaos. Grouping them into blocks means only the ~256 workers who need to collaborate share a table. A whole block is scheduled onto one physical Streaming Multiprocessor (SM).

4. Grid — all the teams together

The picture: zoom out. The entire warehouse floor, tiled with huddles, is the grid. One kernel launch = one grid.

Why the topic needs it: the grid is the scalability layer. You describe your problem as "this many blocks, this many threads each," and the GPU figures out how to pour those blocks onto however many SMs it happens to have — 20 SMs on a small card, 100+ on a big one. You never hardcode the hardware. This is the logical→physical mapping the parent note draws: Grid → GPU, Blocks → SMs, Threads → CUDA cores.

5. The dot-notation symbols: blockIdx, blockDim, threadIdx

These three built-in variables are how a thread discovers who it is. Each has an .x (and for 2D, .y). Read .x as "the horizontal component."

The picture: you are a worker. blockIdx.x = 2 means "I'm in the 3rd huddle (huddles counted from 0)." blockDim.x = 256 means "each huddle holds 256 people." threadIdx.x = 5 means "I'm the 6th person in my huddle."

6. The global index formula — earning every piece

Now we assemble the parent note's central equation. We build it, not just quote it.

WHAT it does: turns a thread's local seat number into a unique global number from $0$ to $N-1$ across the whole grid.

WHY, step by step (look at figure s03):

• Each block holds blockDim.x threads. So before my block, there were blockIdx.x full blocks.
• The number of threads that came before my block = blockIdx.x × blockDim.x. That is my huddle's starting offset — the red bracket in the figure.
• Add my seat inside the huddle, threadIdx.x, and I land on my own unique slot.

A concrete check. Block $2$, block size $256$, seat $5$: $\text{globalIdx}=2\times 256+5=517.$ Thread $517$ handles C[517]. No two threads ever collide, because each (blockIdx, threadIdx) pair gives a different sum — exactly like (page × 256) + line uniquely locates a line in a book.

Why multiplication and not something fancier? We need each block to claim a contiguous, non-overlapping range of indices. Multiplying the block number by the block size is precisely "skip over all the earlier blocks' slots." No trigonometry, no exotic tool — just counting.

7. ceil, and the ceiling-division trick

To cover $N$ elements with blocks of size $T$, how many blocks do we launch?

The picture: you have $N$ eggs and cartons holding $T$ eggs each. Even if the last carton is only partly full, you still need that whole carton. Rounding down would leave eggs homeless.

WHY the + T - 1 trick? Computers do integer division by chopping off the remainder (rounding down). To force a round-up, you nudge the numerator up by $T-1$ first. If $N$ divides evenly, that nudge isn't enough to spill into the next block; if there's any remainder, it is. Example: $N=1000$, $T=256$: $\frac{1000+255}{256}=\frac{1255}{256}=4\text{(integer division)}.$ Four blocks. But $4\times 256=1024>1000$ — which is why we need the next symbol.

8. The boundary check if (idx < N)

Four blocks give 1024 threads, but we only have 1000 elements. Threads $1000$ through $1023$ are extra — they have no valid egg to touch.

The picture: the last carton has 24 empty egg slots. If an extra worker reaches into slot $1000$ of an array that stops at $999$, it grabs memory that isn't ours — a crash or garbage.

This is the parent note's Mistake 1. It is not optional — because block sizes are fixed, numBlocks × T almost always overshoots $N$.

9. Symbols in the memory table: cycle, KB/MB/GB

The parent's memory table uses these units — define them so the table is readable.

The picture — why closer = faster: registers live inside the compute core (millimetres of wire → 1 cycle). Global VRAM sits off-chip (centimetres of wire + a memory controller → hundreds of cycles). Host RAM is across the PCIe bridge (metres of signalling → thousands of cycles). Distance literally costs time because electrical signals take time to travel. The full tiered picture lives in Memory Hierarchy.

Prerequisite map

Each foundation on the left feeds the CUDA execution model. If any node is unclear, revisit its section above before reading the parent note. Related paths onward: Thread Warps and SIMT, Streaming Multiprocessors, GPU Architecture Overview.

Equipment checklist

Cover the right side and answer aloud. If you stumble, that section is your next stop.