You have seen the machine: streaming multiprocessors, warps, a scheduler that hides latency by switching threads. This lesson looks at the same territory from the opposite side — as a contract. Before we can design GPU hardware in the coming modules, we need to state precisely what the software is promised: what a kernel launch means, which guarantees the silicon must honour, and — just as important — which guarantees are deliberately withheld. The withheld ones are where all the hardware freedom lives.

The contract is called SIMT — Single Instruction, Multiple Threads — and its genius is a sleight of hand: the programmer writes one ordinary scalar program, and the silicon runs vector operations. Understanding exactly where the illusion sits is the whole lesson.

The kernel: one program, a grid of threads

The unit of GPU work is a kernel: a scalar function describing what one thread does. You launch it over a grid of thread blocks, each block holding up to ~1024 threads. Every thread runs the same code and discovers who it is from its coordinates — here in CUDA-style pseudocode:

The index arithmetic is the idiom you will type ten thousand times: i = \text{blockIdx} \times \text{blockDim} + \text{threadIdx}. Thread 3 of block 5, with 256 threads per block, owns element 5 \times 256 + 3 = 1283. And note the guard: the grid is built from whole blocks, so launching \lceil n / 256 \rceil blocks over-provisions threads at the ragged edge — the if (i < n) keeps the spares from scribbling out of bounds.

The hardware's secret: warps

Now the sleight of hand. The hardware does not run those threads independently. It bundles them, 32 at a time, into warps that execute in lock-step: one instruction issues, and all 32 lanes perform it on their own registers. Each thread keeps the illusion of a private program counter and private registers — but per warp there is really one instruction stream driving a 32-wide vector unit.

Compare the two honest alternatives you already know. With explicit SIMD intrinsics, you manage the vector width: your source code says "add these 8 floats to those 8 floats", and widening the machine means rewriting the code. SIMT hides the width entirely: you write scalar code per thread, and the compiler and hardware pick how to pack threads into lanes. That is why the same kernel source has survived warp schedulers, dual-issue SMs and four hardware generations unchanged — the width was never in the program. SIMT is SIMD with a programmer-friendly face: vector efficiency, scalar ergonomics.

The four levels, drawn

Step through the decomposition. Two of these levels (grid, block) belong to the programming model; the warp level belongs to the microarchitecture — keep that boundary in sight, because the "Watch out!" below turns on it:

Be the machine: a SIMT executor in TypeScript

The best way to internalise the model is to build it. Here is a toy SIMT machine: it takes a scalar kernel and runs it over a grid warp by warp, exactly as an SM's scheduler would (warp width 4 instead of 32, so the printout fits). Watch the index arithmetic produce each lane's global i, and watch the final ragged warp run with its out-of-range lanes masked:

The inner lane loop is the part that is hardware on a real GPU — all four (or 32) calls happen in the same cycle, driven by one instruction. Only the outer loops are genuinely sequential decisions of the scheduler, and even those interleave freely across SMs.

The block: the unit of togetherness — and the scalability contract

Why does the model bother with blocks at all, rather than one flat sea of threads? Because the contract needs a boundary between "threads that may cooperate" and "threads that may not":

• Within a block: threads are guaranteed co-resident on one SM, so they can share the SM's fast shared memory and synchronise at a barrier (__syncthreads()-style: everyone waits until all arrive).
• Between blocks: nothing. Blocks are independent — the hardware may run them in any order, on any SM, in any degree of parallelism.

That second, deliberately empty promise is the scalability contract, and it is why the GPU ecosystem works at all: a kernel that launches 10,000 independent blocks runs correctly on a laptop GPU with 8 SMs (blocks queue up, ~1250 each) and on a datacentre GPU with 140 SMs (blocks spread wide) — the same binary, no recompilation, near-linear scaling. The hardware designer gets to choose the SM count per market segment precisely because the software was never allowed to care.