Week 8 study guide

Week 8 writes kernels the way a GPU runs them: a grid of blocks, each a group of threads with its own shared memory and a barrier. Two loop axes describe that, a stronger rule keeps threads and blocks from colliding, and render-cuda writes the CUDA. No GPU is needed: the grader's emulator checks every bit. This page recaps the week, collects the rules worth memorizing, and walks the lab one step at a time.

The week in one page

8.1 How a GPU runs code

A kernel is one function every thread runs. Threads come in blocks of up to 1024, a launch starts a grid of blocks, and each thread finds its work from blockIdx.x and threadIdx.x. Blocks run in any order and never wait for each other. Memory is global (everyone), shared (one block) or registers (one thread). -fmad=false and -prec-div=true keep the bits identical to the CPU's.

8.2 Global and local loops

axis=global is the grid, axis=local a block's threads, with a literal bound from 1 to 1024. Neither changes a result. Offsets that are stored, or loaded where the same code stores, must be mixed-radix numbers: digits are loop indices (or their div or mod by a literal) times positive literals, and sorted by coefficient each must beat the most the smaller digits add up to. The offset contains the parallel index, and every access to a buffer is the same expression.

8.3 Shared memory and barriers

Consecutive local loops are phases, with a barrier between them. The rule is checked per phase, so a later phase may read any shared slot an earlier one wrote. An alloc used only in one global loop is shared memory, zeroed for each block. A missing barrier can pass by luck in a block of one or two warps, so the suite also uses blocks of 1024.

8.4 Rendering CUDA

Each top-level global loop is a __global__ function. Shared allocs are __shared__ arrays, cleared by the threads. Each phase is a block on threadIdx.x followed by __syncthreads(). A host kernel launches the kernels in order with UIR_LAUNCH, which the grader defines for real CUDA or for the emulator. For the CPU, a global loop renders like a thread loop.

8.5 Getting a GPU

The emulator grades everything but speed. With an NVIDIA GPU: nvcc on Linux, or --cuda-docker on Windows. Otherwise Colab, Kaggle, or a rented machine on Vast.ai or RunPod for well under a dollar an hour. Bring no secrets to a rented machine, and stop it when you're done.

8.6 Lab

The suite has 39 tests: run tests for the rules, r_ tests for the CPU render, c_ tests for CUDA (four of them must be refused), and one GPU speed test that needs half the reference's speed. Weeks 1 to 7 must stay green.

Cheat sheet

Mistakes people make

Each of these is a bug that was planted in the reference compiler to prove the suite catches it. The week 8 test guide lists every test with the bugs known to fail it.

1. No barrier between phases. Caught by c_reverse_shared, c_reverse_shared_1024 and the speed test.
2. Not clearing shared memory. On a GPU it holds an earlier block's leftovers. Caught by c_shared_dirty_then_fresh.
3. Running a phase on one thread. The bits stay right, but the GPU sits idle. Caught by s_cuda_gemm_512 on a GPU.
4. Letting digits touch. A coefficient equal to the most the smaller digits add is a collision. Caught by err_digits_exactly_touch.
5. Forgetting the parallel index. Every block writing o[t]. Caught by err_blocks_write_one_place and err_threads_write_one_cell.
6. Checking only stores in a phase. Caught by err_same_phase_neighbor.
7. Missing a shape rule. A store outside a local loop, a local loop outside a global one, two block sizes, or two layouts for one buffer. Caught by err_store_outside_local, err_local_outside_global, err_local_sizes_differ and err_block_layouts_differ.

Stuck on the lab?

Build in this order, and don't move on until the check passes.

1. Axes as plain loops. Parse global and local, check local bounds, then ignore them. The run tests that expect a result pass, except block_shared_starts_at_zero.
2. Shared memory in run. Done when this program's buffer is 0 1 2 1 2 3 (and 0 1 2 1 3 5 without the axes):

   %o = buffer 6 dtype=i32 slot=0
   %sh = alloc 3 dtype=i32
   %b = range 2 axis=global
   %t = range 3 axis=local
   %sv = shrink %sh %t shape=[]
   %cur = load %sv
   %bt = add %b %t
   %n = add %cur %bt
   %s = store %sv %n
   %e1 = end %s %t
   %sh1 = after %sh %e1
   %u = range 3 axis=local
   %rv = shrink %sh1 %u shape=[]
   %r = load %rv
   %m = mul %b 3
   %j = add %m %u
   %ov = shrink %o %j shape=[]
   %s2 = store %ov %r
   %e2 = end %s2 %u
   %eb = end %e1 %e2 %b
   %o1 = after %o %eb

3. The rules. Shape rules first, then the mixed-radix rule per phase and per kernel. Done when every run test passes.
4. The CPU render. Reuse your thread-loop outlining. Done when the r_ tests pass.
5. render-cuda. Done when every c_ test passes in the emulator, then on a GPU with the speed test.

Check yourself