The Logical Map & How It Lands on Hardware

The four words you write — kernel, grid, block, thread — and exactly how each becomes silicon you already know.

Lessons 1–2 built the physical map: a GPU is SMs, each SM is sub-partitions of schedulers and lanes, executing warps of 32 in SIMT. You never touch any of that directly. What you actually write lives on the logical map — and this lesson draws it, then bolts it onto the hardware. This mapping is the whole reason you came here.

The four logical levels

You write one function and ask for it to run across a huge collection of threads. That's the entire idea. It has exactly four named levels:

Kernel — a function you mark to run on the GPU. You write its body once; it executes on every thread.
Thread — one execution of that kernel body, with its own registers and its own unique coordinates.
Block — a group of threads (≤ 1024) that can share fast memory and synchronize. The unit the hardware places.
Grid — all the blocks of one kernel launch. The whole job.

You launch a kernel by choosing two shapes — how many threads per block, and how many blocks in the grid:

addArrays<<< numBlocks, threadsPerBlock >>>(a, b, c, N);
// ╰─ grid shape ╯ ╰─ block shape ╯

Those angle brackets are the only place you specify parallelism. Pick threadsPerBlock = 256 and enough blocks to cover your data, and the runtime launches all of them. You do not — and cannot — say which SM anything runs on. (Lesson 2's Q&A: that's the GigaThread Engine's job.)

ONE GRID — e.g. 3,907 blocks of 256 threads to cover 1,000,000 elements

Block 0threads 0–255

Block 1threads 0–255

Block 2threads 0–255

· · ·

Block 3906threads 0–255

ZOOM INTO ONE BLOCK — its 256 threads are cut into 8 warps of 32

Warp 0 · threads 0–31

Warp 1 · threads 32–63

…6 more warps… The warp isn't in your code — the hardware carves each block into warps of 32 by thread index. It's where the logical model quietly becomes physical.

How a thread knows which data is "its own"

Every thread runs the identical kernel body — so how does thread 2,565 work on element 2,565 and not trample element 0? Through four built-in variables the hardware hands each thread:

threadIdx — this thread's index within its block (0–255 here).
blockIdx — this block's index within the grid (0–3906).
blockDim — threads per block (256).
gridDim — blocks per grid (3907).

From these, one line — the single most important line in all of CUDA — turns a thread's local coordinates into a unique global index:

// every thread computes a different i: int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) c[i] = a[i] + b[i]; // guard: grid rounds up past N

Worked example: the thread with threadIdx.x = 5 inside blockIdx.x = 10 computes 10 × 256 + 5 = 2565. It handles element 2,565. Its neighbour in the next block, same local index, gets 11 × 256 + 5 = 2821. No two threads collide; together they blanket the array. That guard if (i < N) matters because 3,907 × 256 = 1,000,192 threads — the last 192 have no element and must do nothing.

Blocks and grids can be 1-, 2-, or 3-dimensional (.x, .y, .z) — a convenience for mapping threads onto 2D images or 3D volumes. It's purely for your indexing sanity; the hardware still just sees a linear stream of warps. Start by thinking 1D.

The mapping — the table this whole course was built for

Now bolt the two maps together. Left is what you write; middle is what the hardware does; right is the rule connecting them — every one of which you already derived in the last two lessons:

Logical — you write	Physical — hardware runs	The rule
Kernel	The instruction stream every lane executes	Written once; run by every thread in the grid.
Grid	Work spread across all the SMs	The GigaThread Engine streams blocks onto SMs as they free up — in waves if there are more than fit.
Block (≤1024 threads)	Assigned whole to one SM, resident for its lifetime	Indivisible — never spans SMs (it needs SM-local shared memory & barriers). Several blocks may share an SM if resources allow.
Warp (32 threads of a block)	One sub-partition's scheduler + lanes, SIMT lockstep	Hardware cuts each block into warps by thread index. Resident warps interleave cycle-by-cycle (latency hiding).
Thread	One CUDA-core lane, with its own registers	thread ↔ lane. Its registers come from the sub-partition's register-file slice.

The line down the middle of your mission

You choose the logical shape — grid and block dimensions, in the launch brackets. The hardware chooses all physical placement — which SM, which sub-partition, which cycle. You never cross that line by hand; you shape the left column and the right column happens to you.

Why the model is built this way: blocks must be independent

One consequence ties it all together. Because the hardware decides block placement and order — and runs blocks in waves — blocks cannot depend on each other. They must be safe to run in any order, all at once or one at a time. Threads within a block can cooperate (shared memory, __syncthreads()); threads in different blocks essentially cannot, mid-kernel.

That constraint is the model's superpower, exactly as you spotted earlier: the same grid runs unchanged on a 20-SM laptop chip or your 84-SM RTX 5080 or a 132-SM H100. The hardware just pours blocks onto however many SMs exist. Logical stays fixed; physical scales. That decoupling — the gap between the two maps you named in your very first message — is not an accident of GPUs. It's the entire design goal.¹

Check yourself

From memory — one click locks each answer.

ApplyBlocks hold 256 threads. The thread with threadIdx.x = 3 in blockIdx.x = 4 works on which global element?

MappingA single block of 512 threads is launched. On the hardware, those 512 threads —

ConceptCUDA forbids one block from depending on another mid-kernel chiefly because —

Primary source to read next (≈20 min)

CUDA C++ Programming Guide — §"Programming Model": Kernels, Thread Hierarchy, and the built-in variables. NVIDIA's own definition of grid/block/thread and threadIdx/blockIdx, with the same indexing you just did. This is the canonical statement of today's logical map. Read it here.

You're now through the compute half of your mission. Great next threads: what actually happens at __syncthreads(), and why can't it work across blocks? · how do I pick a good block size? · or we start the memory hierarchy (Lessons 4–5), where "shared memory," which kept appearing above, finally gets defined. Ask in chat.

Notes & citations

CUDA C++ Programming Guide, "Thread Block Clusters / Scalability" — thread blocks are required to execute independently, in any order, which is what lets a compiled program scale to any number of SMs. docs.nvidia.com. See also resources.html.

← Lesson 2 ⌂ Home Lesson 4 →