Appendix C. Glossary and Quiz Bank

The first half of this appendix defines the vocabulary of the book, as uopc uses it. Where a term has a broader meaning in the literature, the definition gives the general sense first and the uopc specifics second, with the module where it lives. The second half collects quiz questions for every week of the course, in the style of the in-class quizzes, with answers in collapsed blocks.

Where this fits: op-level details are in Appendix A; testing practice is in Appendix B; the last appendix, Appendix D, lists the full source.

C.1 Glossary

Accumulator. Scratch memory a reduction adds (or maxes) into. Rangeify creates one size-1 ALLOC per reduction, initialized to the identity, updated inside the reduce loop, and read after it. The C renderer turns a size-1 LOCAL accumulator into a plain scalar (float acc0;).

AddrSpace. The third inferred property of a UOp (after dtype and shape): ALU for values, GLOBAL for buffers passed to a kernel, LOCAL for call-scope scratch memory, SHARED for GPU shared memory. LOAD and STORE require a non-ALU source (uopc/dtype.py).

AFTER. AFTER(mem, *deps) is mem, but only once deps have executed. It is how an SSA graph, which has no inherent order, expresses "read this memory after that write" or "this accumulator lives inside that loop". It renders as nothing.

All-reduce. A collective operation in which every device contributes an array and every device receives their sum (or another reduction). Data-parallel training all-reduces gradients. A ring all-reduce moves about 2N elements per device regardless of the device count.

ALLOC. Call-scope storage, float a[n]; in C. Created by rangeify for accumulators or written by hand in loop IR. The linearizer declares it in the deepest scope shared by all its users; SHARED allocs go at the top of the kernel.

ALU op. An op that computes a value from values: CAST, RECIP, EXP2, LOG2, SQRT, SIN, ADD, MUL, MAX, CMPLT, CMPNE, FLOORDIV, FLOORMOD, WHERE (uopc.uop.ALU). Elementwise on shaped values.

Arithmetic intensity. FLOPs performed per byte (or per element) moved from memory. Low intensity means memory-bound. Register tiling raises it by reusing each loaded value many times.

Autodiff (reverse mode). Computing the gradient of a scalar with respect to many inputs by walking the graph from the output back to the inputs, applying a local rule at each node and summing contributions where paths merge. In uopc, autodiff.gradient builds a new graph for the gradient, which the compiler then compiles like any other.

AxisType. The kind of a RANGE: LOOP, REDUCE, UPCAST, THREAD, GLOBAL, LOCAL. It decides how the renderer emits the loop: a for, an OpenMP parallel for, or a GPU block/thread index.

Bank conflict. On a GPU, shared memory is split into banks; when threads of one warp access different addresses in the same bank in one instruction, the accesses serialize.

Barrier. A point every thread of a GPU block must reach before any continues (__syncthreads()). Ops.BARRIER in the loop IR; it cannot be emulated by running threads one after another.

Beta reduction. Replacing a function application by the function's body with arguments substituted for parameters. inline_call performs it on CALL nodes, scoped so nested CALL bodies are untouched.

bf16 (bfloat16). A 16-bit float with f32's 8-bit exponent and a 7-bit mantissa: f32's range at reduced precision. Not in uopc.

Bounds analysis. Computing, for every integer expression, an interval [vmin, vmax] guaranteed to contain its value (symbolic.bounds). Drives div/mod simplification, the choice of / vs floordiv_i32 in the renderer, and constant-folding of comparisons.

Broadcasting. Making shapes compatible by adding leading size-1 axes (RESHAPE) and stretching size-1 axes (EXPAND). The IR never broadcasts implicitly; the Tensor API and the .uir parser insert the ops.

BUFFER. Global, persistent, flat memory: a kernel argument. In the Tensor runtime a BUFFER UOp owns its numpy array.

CALL / PARAM. A function application and its formal parameters. CALL(body, *args); the body's PARAMs are numbered by slot. Week 1 only: they are eliminated by inlining, never rendered.

Canonical form. A chosen representative among equivalent expressions, so equal things look equal. The symbolic rules move constants to the right of commutative ops; hash consing then makes equal canonical forms the same object.

Coalescing. When the threads of a warp access consecutive addresses, the hardware combines the accesses into few memory transactions. Map the contiguous axis to threadIdx.x (NOTES 23).

Common subexpression elimination (CSE). Computing a repeated expression once. Hash consing gives it for free: building the same subexpression twice returns the same node.

Confluence. A rewrite system is confluent if the order in which rules are applied does not change the final normal form. Real rule sets, including uopc's, are designed to be close to confluent but are not proven so.

Constant folding. Evaluating an op whose inputs are all constants at compile time. exec_alu is the reference semantics; it refuses to fold undefined behavior (casting inf to int).

Data parallelism. Splitting a batch across devices that each hold a full copy of the model, then all-reducing gradients.

DEBUG. Environment variable controlling diagnostics in kernelize.realize: 1 schedule summary, 2 per-kernel timing, 3 warm re-run timing, 4 generated C.

DType. A data type record (name, itemsize, C type, numpy format): void, bool, i32, f32, f16.

Emulator. gpu.emulate: rendering a GPU kernel with the C renderer, where GLOBAL and LOCAL ranges are ordinary loops, so its index mapping can be tested on a CPU.

END. END(x, *ranges) closes loops: nodes that depend on it run after those loops finish.

EXPAND. A movement op that stretches size-1 axes; reading an expanded axis reads index 0. Its gradient is a sum over the expanded axes.

f16. IEEE half precision: 5-bit exponent, 10-bit mantissa, maximum 65504. _Float16 in the generated C, half in generated CUDA.

Finite differences. Approximating a derivative by (f(x+ε) − f(x−ε)) / 2ε. The oracle for gradient tests.

Fixed point. A state that a transformation no longer changes. graph_rewrite rewrites until no rule fires.

Floor division. Division rounding toward negative infinity (Python's //), as opposed to C's truncation toward zero. FLOORDIV/FLOORMOD use floor semantics, which makes (a*c + b) // c = a + b // c unconditional.

Fusion. Computing several ops in one kernel without writing intermediates to memory. Its cost is recompute when a fused value is consumed through an EXPAND.

fuse_reduce. A kernelize mode in which REDUCEs are not materialization points, so reductions nest inside each other (used for fused attention); it also skips stage_broadcast_compute).

Guard. A run-time check that the assumptions a compiled program was specialized on still hold (a shape, a bound on a symbolic size). Failing a guard means recompiling or rejecting the input.

Gradient check. Comparing an autodiff gradient against finite differences.

graph_rewrite. The rewrite engine: bottom-up, greedy, to a fixed point, iterative, memoized by node (uopc/rewrite.py).

Hash consing. Constructing structurally identical objects only once, so identity is structural equality. UOp.__new__ looks up (op, src, _argkey(arg)) in a weak-value cache.

Identity element. The value a reduction starts from: 0 for ADD, -inf (or INT_MIN) for MAX.

im2col. Rewriting a convolution as a matrix multiply by gathering each receptive field into a row. Tensor.conv2d builds the columns from SHRINKs.

INDEX. Indexing: on a value it selects a sub-array (popping leading axes); on memory it computes an address that LOAD and STORE use.

Interval arithmetic. Computing bounds of an expression from bounds of its parts ([a0,a1] + [b0,b1] = [a0+b0, a1+b1]). Sound but unable to see relations between variables (it cannot prove j < n inside a loop to n).

JIT cache (structural cache). kernelize._plans, keyed by _structure_key: the graph's structure with buffers anonymized. A training step with the same structure as the previous one reuses the compiled plan.

Kernel. One compiled function that reads buffers and writes one output buffer. A plan is a list of kernels plus the temp buffers between them.

Kernelize. Cutting a shape graph into kernels at materialization points, and building the plan (uopc/kernelize.py).

KV cache. In autoregressive decoding, the stored key and value vectors of all past tokens, per layer, so each new token only computes its own.

Layout pass. opt.layout_pass: copying an operand whose reduce axis is strided into a contiguous temp (a STAGE) before a reduction reads it, after pushing movement ops down to the data (push_views).

Linearizer. The pass that turns a graph into an ordered list of UOps respecting loop scopes (uopc/linearize.py), by list scheduling over the loop tree.

LOAD / STORE. Read a value from memory / write a value to memory, through an INDEX address.

Loop-invariant hoisting. Computing an expression outside the loops it does not depend on. The linearizer does it automatically: a node is emitted in the scope of its sources, not of its users.

lower. rangeify.lower(root, out, ctx): turn one shaped value into a SINK of loops that store it into a buffer.

Materialization point. A node whose value is written to memory, ending a kernel: every root, every STAGE, and every REDUCE unless fuse_reduce is set.

Mixed precision. Storing values in a low-precision type (f16, bf16) and accumulating in a higher one (f32).

Movement op. RESHAPE, EXPAND, PERMUTE, SHRINK: ops that change how indices map to data without computing anything.

Mutation testing. Measuring a test suite by injecting small defects (mutants) and checking that some test fails for each.

Normal form. An expression no rule can rewrite further.

OpenMP. A C extension for shared-memory parallelism; #pragma omp parallel for splits a loop's iterations across threads. Emitted for THREAD ranges.

OPT. Environment variable selecting the kernel optimizer in tensor.py: 0 none, 1 HandOpt (default), 2 SearchOpt.

Opts. The per-kernel schedule: upcast factors, a thread axis, unroll factor, tile (cache block) sizes, and the GPU flag with local sizes (rangeify.Opts).

PatternMatcher / UPat. A list of (pattern, function) rewrite rules, indexed by root op; UPat matches UOps structurally, binds names, and with a list src tries all orders of a commutative op's sources.

PERMUTE. A movement op that reorders axes. Its gradient is the inverse permutation.

Placeholder buffer. A BUFFER with a negative slot that kernelize substitutes for a materialized value while building each kernel, later bound to a temp, input or output.

Property test. A test that checks a general claim on many generated inputs (for example, that simplify preserves values) rather than on a few chosen examples.

Rangeify. Lowering a shape graph to scalar loop code by giving every output axis a RANGE and pushing the index tuple down through every op (uopc/rangeify.py).

RANGE. A loop: its value is the loop counter, its source is the trip count, its arg is RangeArg(idx, kind). Loops nest in increasing idx.

ravel / unravel. Converting a multi-dimensional index to a flat offset in row-major order, and back.

Recompute. Computing the same value more than once because it was fused rather than materialized. The trade-off of fusion; stage_broadcast_compute bounds it with a cost rule.

REDUCE. REDUCE(x, (op, n)) reduces the last n axes with ADD or MAX.

Register tiling (upcasting). Computing a small tile of outputs per loop iteration, held in registers, so each loaded input is reused across the tile.

RESHAPE. A movement op that changes the shape without changing the element count or order.

Roofline model. Attainable performance = min(peak FLOP/s, arithmetic intensity × memory bandwidth).

Scope. The set of loops a node must sit inside: every RANGE it depends on, minus those closed by an END on the path. The linearizer closes scopes under "encloses" so each is a path in the loop tree.

Shape graph. A graph of shaped UOps (with movement ops and reductions), before rangeify.

Shared memory. Fast, per-block, programmer-managed GPU memory (__shared__). AddrSpace.SHARED.

SHRINK. A movement op that keeps a sub-box [begin, end) per axis. Its gradient is a zero-pad.

SIMT. Single instruction, multiple threads: the GPU execution model in which a warp of threads executes one instruction stream on different data.

SINK. The root of a program: groups side effects.

SSA. Static single assignment: every value is defined once. A UOp graph is SSA by construction.

STACK. Creates a new leading axis from same-shaped sources. Rangeify turns it into a WHERE chain.

STAGE. "Materialize this into a temp buffer": a user- or pass-inserted materialization point. With arg="detach" it is instead a stop-gradient (Tensor.detach()), removed by kernelize before scheduling.

Symbolic shape. A shape entry known only at run time, represented as an expression (usually a variable with declared bounds) rather than an integer. Not in uopc; see Chapter 15.

Tensor core. GPU hardware that performs a small matrix multiply-accumulate (for example 16×16×16 with f16 inputs and f32 accumulation) per instruction, issued cooperatively by a warp.

Term rewriting system. A set of rules lhs → rhs applied to expression trees until none applies. The basis of uopc's optimizer and its week-1 interpreter.

Termination. A rewrite system terminates if no infinite sequence of rewrites exists. graph_rewrite has a step limit as a safety net, not a proof.

THREAD. The AxisType of a loop run in parallel across CPU threads with OpenMP.

Toposort. An ordering of graph nodes with every node after its sources. UOp.toposort is iterative.

UOp. The IR node: UOp(op, src, arg), immutable and hash-consed, with inferred dtype, shape and addrspace.

.uir. The course's text interchange format: one SSA statement per line, parsed by uir.parse and printed by uir.to_uir.

Unroll. Replicating a loop body so fewer iterations run; in uopc, Opts.unroll replicates the innermost reduce body and combines the copies as a balanced tree.

UPCAST. The AxisType for a fully unrolled axis. In uopc, upcasting is done during rangeify by instantiating the body per offset, so no UPCAST range is ever created.

vmin / vmax. The lower and upper bounds computed by bounds analysis.

Warp. A group of 32 GPU threads that execute in lockstep.

WHERE. Select: WHERE(c, a, b) is a where c is true, else b.

WMMA. Warp matrix multiply-accumulate: CUDA's API for tensor cores (nvcuda::wmma). Ops.WMMA is reserved in uopc; tensor-core kernels are produced by pattern (gpu.wmma_gemm).

C.2 Quiz bank

Questions are grouped by week. Answers refer to uopc and to measured numbers from book/NOTES.md where relevant.

Week 1: UOps, rewriting, symbolic

1. Why does uopc's hash-consing cache key use _argkey(arg) rather than arg?
2. What is -7 // 2 under FLOORDIV, and what does C's -7 / 2 give?
3. Why is the rule (x + c1) + c2 → x + (c1 + c2) restricted to integers?
4. In graph_rewrite, why is the result of a rule that fired rewritten again?
5. x is an i32 PARAM. What does simplify((x*6 + 4) % 2) return, and why is it valid for every x?
6. Why are toposort and graph_rewrite written with explicit stacks?
7. What are the shape and value of INDEX(STACK(STACK(1,2,3), STACK(4,5,6)), CONST 1)?
8. When inlining an outer CALL, why must the substitution not descend into a nested CALL's body?
9. What does folding 16777216.0 + 1.0 in f32 give, and why must the folder round through f32?

Week 2: memory, rendering, runtime

1. What does AFTER(mem, dep) express, and why does an SSA graph need it?
2. Exactly when does the C renderer emit / for a FLOORDIV?
3. Which UOps does the renderer inline as expressions, and which always get a named variable?
4. Name one error that address-space inference catches and dtype/shape inference cannot.
5. Why does compile_c compile in a temporary directory inside the cache directory?

Week 3: loops and the linearizer

1. Define the scope of a node.
2. Why did the first linearizer design ("emit a node when its scope equals the open loops") deadlock on GEMM?
3. Where is an accumulator ALLOC declared, and what goes wrong with OpenMP if it is declared at the top of the kernel?
4. In x - max(x) computed per row, which two loops are siblings, and which must run first?
5. How does the loop-IR GEMM in tests/test_w03_loops.py keep its accumulator inside the i and j loops but outside k?

Week 4: movement ops and rangeify

1. What does the syllabus's gemm_256x256.uir compute? Give the correct version.
2. State the index map rangeify uses for EXPAND, SHRINK and PERMUTE.
3. Why does rangeify simplify the index at every RESHAPE? Quote the measured effect.
4. REDUCE pops which axes? How do you sum axis 0 of a (3, 4) tensor?
5. What does rangeify turn a STACK into?
6. How is PAD built without a pad primitive?

Week 5: kernelize, autodiff, the Tensor library

1. What are the materialization points of kernelize?
2. Why can fusing an elementwise result into a reduce through an EXPAND be a bad idea? What rule does uopc use, and what did it save on GPT?
3. Why does make_plan renumber buffers 0..n-1 within each kernel?
4. Why does AdamW pass its bias corrections as one-element tensors rather than Python floats?
5. Give the gradient rules for EXPAND, SHRINK and PERMUTE.
6. How does uopc split the gradient of REDUCE(MAX) among tied maxima?
7. Why does autodiff refuse STORE, CALL, AFTER, RANGE and END?
8. Why is nn.Embedding a one-hot matmul?

Week 6: fast CPU kernels

1. Why does register tiling raise arithmetic intensity? Use the 4×8 tile from opt.py.
2. Why did adding threads make the strided 1024³ GEMM slower?
3. What does -fassociative-math allow, and what is the cost?
4. With Opts(upcast={0: 2, 1: 2}), how many for loops and accumulators does a GEMM kernel have?
5. What does the layout pass do, and when does it stage an operand?
6. How do hand_opt and search_opt differ?

Weeks 7–8: GPUs, shared memory, tensor cores

1. Which kernel axis gets threadIdx.x, and why?
2. Why does smem_gemm_ir need two barriers per tile step?
3. What does gpu.emulate let you test without a GPU, and what can it not test?
4. By what factor does T×T tiling in shared memory cut global-memory traffic of a GEMM?
5. What are the operand and accumulator types of a typical tensor-core MMA, and what tile does wmma_gemm use?

Weeks 9–10: real models

1. Why does examples/gpt.py pad the sampling context to a fixed length?
2. How does examples/gpt.py make attention causal?
3. What state does online softmax keep per query row?
4. What were the measured effects, on GPT, of stage_broadcast_compute and of identity-based hashing plus memoized simplify?

Beyond (Chapter 15)

1. Give the KV-cache size formula and evaluate it for 2 layers, dim 64, 64 tokens, f32.
2. What do declared bounds on a symbolic size buy the compiler?
3. In data-parallel training, what is communicated each step?