The film
About 5:29, hosted by Pylon. It asks you two questions and waits for your answer before it shows you. The guide's voice is not recorded yet, so for now the film plays with captions; the full words are also below.
Nothing loads until you press play.
The lab: the cannery floor
Choose a crew size, a stride and a list length; count every thread and every trip with the emulator; then, if your browser has WebGPU, run the real kernel on your own chip, check it and time it. In the film, reads eight apart took 5.67 times the trips of neighbouring reads.
Counting runs in this page. Timing runs the kernel on your own graphics chip, only when you press Time. Nothing is sent anywhere.
Everything the film says
Threads, workgroups and memory
Threads, workgroups and memory.
A graphics chip does not do one sum very fast. It does thousands of small sums at the same time. Tonight Pylon runs the cannery where that happens, and counts every trip to the warehouse.
One worker, one can
The job: add two lists of numbers, element by element, into a third. A list of 1,048,576 numbers.
On a graphics chip, each element gets its own worker, called a thread. Thread 40 reads number 40 from each list, adds them, and writes the sum to place 40. Every thread runs the same few lines of code.
Crews
Threads come in crews called workgroups. You choose the crew size; in WebGPU a crew may hold up to 256 threads by default.
A thread finds its element from its crew's number and its place in the crew: crew number times crew size, plus its place. With crews of 16, crew 2, place 8, is element 40.
Crews come in whole numbers. Add 1,000 numbers with crews of 64 and you need 16 crews, 1,024 threads. The last 24 have nothing to do, so the code must check that its element exists.
One past the end
Get that check wrong by one, and one idle thread writes one place past the end of the list. In a list of 40 numbers, numbered 0 to 39, element 40 has no place of its own.
If the check is off by one, does the program crash? (the film pauses here)
On a graphics chip that need not crash. The WGSL rules call it an error, but let it end quietly: the stray write may land anywhere in the same buffer, be dropped, or stop that one thread. So the output can look right while something else in the same buffer was spoiled.
Our checker compares every sum with a slow, careful one on the main processor, and plants guard values just past the end. Of 12 test cases, nine came from three kernels we broke on purpose, each run three ways. Together the two checks catch seven of the nine. The two they miss are stray writes the chip dropped, which leave every sum right. Only our emulator, which counts every write, sees those.
Trips to the warehouse
On NVIDIA chips, each crew is split into teams of 32 threads, called warps; other chips use other team sizes, from 4 to 128, and our counts use NVIDIA's 32. When a team reads from memory, the chip fetches whole blocks at a time; the current CUDA guide describes blocks of 32 bytes.
If the 32 threads read 32 neighbouring numbers, 4 bytes each, that is 128 bytes: 4 blocks. One team, a few trips.
Now spread the reads out, every eighth number. The same 32 reads now touch 32 blocks, one per thread.
Counting the trips
Count every trip for the whole job, reads and writes. With neighbouring reads: 393,216 blocks. With reads spread eight apart: 2,228,224. That is 5.67 times the trips, not 8, because the writes still go to neighbouring places. Most of each block fetched is thrown away.
Reads spread eight apart. How many times the trips? (the film pauses here)
Spreading them further changes nothing more in this model, because each thread already has a block of its own.
Your own chip
These are counts from a model of the chip, not times. Real chips have caches that hide some of the cost, and they differ from one another.
So the lab runs the real kernel on the chip in your own computer, if your browser allows it, and times it. If the time for spread-out reads is far from what the count predicts, then caches, a different team or block size, or the cost of starting the job are at work.
Goodnight
Thousands of small workers, in crews, fetching in blocks. Fast code keeps neighbouring workers on neighbouring numbers, and checks the edges.
Next in the cannery: multiplying whole tables, and carrying trays instead of single cans. Goodnight.
Sources
Every claim in the film, with where it comes from and what that source does not prove. The numbers about the model are computed from its weights, which you can load and check: facts.json.
Night School's GPU emulator and kernels, built for this film. media/school/threads/facts.json, written by tools/school/fill_threads.mjs from system/school/threads.mjs: WGSL kernels that add two arrays (contiguous and strided), an exact JS emulator of thread grouping, idle threads and memory transactions in 32-byte and 128-byte models, and a checker with canaries tested on nine broken kernels. https://github.com/HarperZ9/HarperZ9.github.io/tree/main/media/school/threads
Used for: n = 1,000 with workgroups of 64: 16 workgroups, 1,024 threads, 24 idle. 32-byte model: contiguous 393,216 transactions for the whole job; stride 8 and above 2,228,224 (efficiency 0.1765). Of nine broken cases (three kernels, three allowed outcomes each), the output check with canaries misses two (stray writes dropped); the emulator's out-of-bounds store count catches those, and no case escapes both.
Does not prove: Any timing on a real GPU; caches are not modelled.
W3C WebGPU specification. W3C (2026). WebGPU. Candidate Recommendation Draft; supported limits table. https://www.w3.org/TR/webgpu/
Used for: Default maxComputeInvocationsPerWorkgroup 256 and maxComputeWorkgroupSizeX 256.
Does not prove: Anything about a particular device.
W3C WGSL specification. W3C (2026). WebGPU Shading Language. Candidate Recommendation Draft; built-in values and out-of-bounds access. https://www.w3.org/TR/WGSL/
Used for: global_invocation_id = workgroup_id * workgroup_size + local_invocation_id; an out-of-bounds store may write to any location of the bound buffer or not be executed.
Does not prove: Anything about a particular device.
NVIDIA CUDA C++ Programming Guide. NVIDIA (2026). CUDA Programming Guide, version 13.4.2: section 1.2.2.2 (warps) and section 2.3.4.1 (global memory, 32-byte transactions). https://docs.nvidia.com/cuda/cuda-programming-guide/
Used for: Threads execute in warps of 32; global memory is accessed in 32-byte transactions.
Does not prove: Anything about WebGPU devices from other vendors.
What this lesson leaves out
The film's numbers are counts from a model of the chip, not times; caches are not modelled.
Teams of 32 and 32-byte blocks are NVIDIA's; other chips use other sizes, so a time on your chip can differ from the count.
The kernels were checked against the emulator and the main processor; they were not timed on a graphics chip for this film.
Before this: Matrices move space (F2). Earlier at the prism: Evaluating a language model (L16).