Done means merged, tested, and green on both example
gates. The design records live in the repo:
docs/gpu_tiles.md (which numbers the GPU work
internally as stages 0 to 4), docs/simdgroup_plan.md,
and the
Metaxu Book, whose examples run in CI.
done
The compiler runs programs. Parser, modules,
type inference with hard conflict detection, and a strict interpreter
that defines the semantics. 21 example programs run with pinned
output, one of them a project with locked dependencies; two
negative fixtures must be rejected.
done
Programs compile to native code. An LLVM
backend whose output must match the interpreter byte for byte,
including error messages and float formatting. It is reaching for
C-like performance: small paired microbenchmarks exist to pin
regressions and keep performance in view, not to score points
against C.
done
Modes replace lifetimes. A borrow checker
over locality, mutability, and linearity that propagates properties
instead of asking for annotations. Thread-safety is inferred from
what spawned closures capture.
done
Effects run threads. Delimited one-shot
continuations; try/catch; real OS threads
and mutexes behind Thread/Mutex effects;
a handler crossing costs about 280 ns.
done
Tiles compute on the CPU.
Tile[T, R, C] values with shapes checked at compile
time, deterministic interpreter semantics, and a native lowering
tested against them.
done
Kernels launch through an effect.
Gpu.launch(n, f) defaults to a sequential CPU loop;
handlers virtualize it. Buffer reads see the launch snapshot, writes
merge through a written mask.
done
Kernels compile to Metal source. An emitter
that produces C++-compatible MSL, so a clang++ shim tests it
bit-exactly in a container with no GPU, plus a generated
self-checking harness for Apple silicon.
done
Tiles in 32 and 16 bits. f32 and f16 as tile
element kinds, stored as representable doubles with one rounding per
operation, bit-identical across all three engines.
done
The language dispatches to Metal.
run_metal installed as the Gpu handler
sends the same kernel through MLX on a Mac, or through the bit-exact
shim elsewhere. Anything outside the contract is an error, never a
silent CPU fallback.
done
Matrix units. A kernel with an 8x8 f32 or
f16 dot gets a second lowering: one launch instance is one 32-lane
simdgroup, tiles live in threadgroup memory, and the dot is
simdgroup_multiply_accumulate on
simdgroup_float8x8. Kernels and the reference
interpreter did not change; the shim emulates the collectives and
stays bit-exact.