Week 8 · Lesson 6

The week 8 lab

Teach your compiler blocks and threads. It checks the global and local rules, gives each block fresh shared memory, renders global loops for the CPU, and writes CUDA with render-cuda. You're done when every week 8 test passes, in the emulator or on a GPU, weeks 1 to 7 still pass, and you've written 10 tests of your own.

The command-line contract

commandweek 8 changes
your.exe run FILEaxis=global and axis=local run as plain loops, except that a block's shared memory starts at zero for every value of the global index. New static errors (exit 2): a local bound that isn't a literal from 1 to 1024, a local loop outside a global loop or inside another local loop, local loops of different sizes in one global loop, another parallel loop inside a global loop, a store outside a local loop, and an offset that breaks the mixed-radix rule.
your.exe render FILEA global loop renders like a week 7 thread loop, with its shared memory declared in the outlined function. Local loops are plain loops.
your.exe render-cuda FILENew. Prints CUDA: a __global__ function per global loop, __shared__ arrays zeroed by the threads, each phase as a block on threadIdx.x followed by __syncthreads(), and a host kernel that calls UIR_LAUNCH. Exits 2 if the program isn't in kernel form.

How the tests are graded

Running the suite

python tools/test.py --compiler "C:/path/to/your.exe" --week 8                     # nvcc if it's on PATH, else the emulator
python tools/test.py --compiler "C:/path/to/your.exe" --week 8 --cuda-docker nvidia/cuda:12.6.3-devel-ubuntu22.04
python tools/test.py --compiler "C:/path/to/your.exe" --week 8 -k c_reverse -v     # one group, with details
python _reference/uirc.py render-cuda tests/week8/render/c_reverse_shared.uir     # the reference's CUDA, to compare

Then run weeks 1 to 7 again, one --week at a time. The speed test reads data/a512.f32 and data/b512.f32 next to it, 1 MB each. Lesson 8.5 covers where to get a GPU, and you don't need one until the last step.

Write your 10 tests

Ten tests in path/to/your-repo/tests/week8/, checked with python tools/test.py --check-tests path/to/your-repo/tests/week8. Good targets: an offset with div and mod digits that just fits, and one where a coefficient is one too small. A vector store whose size pushes it into a neighbor's slot. A later phase reading another thread's slot, and the same read in the same phase. Shared memory read before any phase writes it. A CUDA test with two kernels where the second reads what the first wrote.

A build order that works

  1. Parse the new axes

    Accept global and local, check local bounds, and run both as plain loops.

    Checkpoint: weeks 1 to 7 pass, and the week 8 run tests that expect a result pass, except block_shared_starts_at_zero.

  2. Shared memory

    An alloc whose every load and store is inside one global loop is shared. Re-zero it for each value of the global index in run, exactly like week 7's private allocs.

    Checkpoint: block_shared_starts_at_zero and plain_loop_alloc_carries_over both pass.

  3. Shape rules

    Where local loops may be, one block size per kernel, no other parallel loops inside, no stores outside local loops.

    Checkpoint: err_local_*, err_thread_in_global and err_store_outside_local exit 2.

  4. The mixed-radix rule

    Split each offset into terms, find the digits and their largest values, sort, and compare. Do it per phase with the local index, then per kernel with the global index, skipping shared memory. Check the parallel index is there, and that every access to a buffer is the same expression.

    Checkpoint: every run test passes.

  5. Render for the CPU

    Treat a global loop as a thread loop and its shared allocs as private ones. Your week 7 outlining does the rest.

    Checkpoint: the r_ tests pass.

  6. render-cuda

    Start from your C emitter. Emit one __global__ function per top-level global loop with its free buffers as parameters, then the shared arrays and their zeroing, then each phase. Add __device__ to your helpers and emit the host kernel.

    Checkpoint: every c_ test passes in the emulator.

  7. On a GPU

    Run the suite with nvcc or --cuda-docker. If a test passes in the emulator and fails on the GPU, or the other way around, look for a missing barrier or unzeroed shared memory first.

    Checkpoint: everything passes on the GPU, including s_cuda_gemm_512.

Checklist

The AI policy, still

An AI assistant may explain concepts, draw diagrams, and help you debug by asking questions, but must never write or rewrite compiler code. "Why does my reverse pass in the emulator but fail on the GPU?" is a fine question. "Write my render-cuda" is not.

Check yourself

Your render-cuda launches every kernel with 1 thread per block, and each thread loops over the whole local range. Which tests fail?

Do the bits change if one thread does every iteration in order?

One thread running the phases in order is exactly the plain loops, so every bit matches. Only the speed test sees it, and that's why there is one. This is one of the bugs planted in the reference to prove the suite catches it.

c_reverse_shared passes on your GPU but c_reverse_shared_1024 fails. What's the likely bug?

Lesson 8.3, "A missing barrier".

The two tests differ only in block size, 64 threads against 1024. A block of 1024 has 32 warps that run at different times, so a phase can start reading before another warp has finished writing.

Without a GPU, can you finish the week?

Lesson 8.5.

The emulator checks correctness completely. The speed test builds your CUDA, and timing it is the one thing worth renting a GPU for, which lesson 8.5 shows how to do for about a dollar.