Blog

Notes on building Vx — what the compiler does, what it refuses, and what it got wrong along the way.

• GLM on an A100, in Vx 20 October 2026

  GLM-5's decoder, with its sparse attention and 256-expert layers, runs on an A100 through 28 GPU kernels written in Vx. On random weights it matches the official model to within 1e-5 up to 78 layers, generates the same 16 tokens, and runs INT4, FP8 and BF16 weights as stored. The GPU also found a race that no CPU test could.

• Optimizing for the common case, in language design 20 October 2026

  The most common parallel loop starts one task per index. Go, C# and JavaScript first made it a trap, and Rust needs move. In Vx, a task takes each variable the way it uses it, with no move keyword, and a copy can never go stale.

• Breaking the borrow checker on purpose 15 October 2026

  We broke Vx's borrow checker 108 ways, one at a time, to see whether the tests would notice. They caught 93. Of the 15 they missed, 11 were rules no test checked, and one led to a program that read freed memory. Then 78 of Rust's borrow checker bug reports, run in Vx, found 14 more holes.

• What the borrow checker let through 14 October 2026

  A review of Vx's borrow checker found eleven programs that compiled and should not have. One printed a different number with optimization on, and one, sent by a reader, changed a variable not declared mut. All are fixed. Rust's checker has open bugs of the same two kinds, one of them reported this week.

• When the build order changes the program 9 October 2026

  C++ generates a copy assignment or a template instance only in the files that use it, and the linker keeps one copy. When the copies differ, the link order picks which one runs. Swift and Haskell share the problem, Rust rules it out, and Vx is mostly there, with one gap found while writing this.

• Programming a GPU fleet as one program 8 October 2026

  In Vx the machine is part of the program, so a serving plan that does not fit a GPU is refused when it is compiled, and every move between GPUs and machines has a price before it runs. With a 70B serving plan on an 8-GPU node, a 148-node fleet as a graph, and the gaps.

• Can a library describe the memory hierarchy? 8 October 2026

  A Rust library can put the memory space in a tensor's type, and the Rust compiler then rejects a host read of GPU memory. It cannot add up how much memory a call path uses, compute shapes on stable Rust, or check the code that runs on the device. Both versions, side by side.

• Why void is a value in Vx 8 October 2026

  In Vx, void is an ordinary type with exactly one value, as () is in Rust and unlike void in C, so generic code needs no special case for it. A bug treated that value as something that could be moved; the fix, and four gaps found while writing this.

• How Vx frees memory 7 October 2026

  Vx frees memory in four ways, and the type checker decides almost all of it. A tensor is freed right after its last use, a value with a Drop impl at the end of its block as in Rust, and GPU memory through the plugin that allocated it. With the drop points the compiler prints, and the gaps that are left.

• Proving tensor indices at compile time 5 October 2026

  The Vx compiler hands index arithmetic to the z3 theorem prover. It proves that a KV-cache offset and a copy of attention heads into shared memory stay inside their tensors, for every input, and refuses a wrong stride. Also a proof that held while the program printed a negative offset, and why.

• Vx fully embraces {ML, IR, MLIR} 3 October 2026

  Tensor programs, Vx's own intermediate representation, and a Vx MLIR dialect, shown stage by stage down to a PTX kernel with warp shuffles. Then mlir! blocks that state what they write, so the race proof can check a GPU loop that calls them. Rewritten 24 October 2026.

• Can the borrow checker algorithm be auto-vectorized? The answer is yes! 1 October 2026

  Vx's borrow checker packs four lifetimes into one 64-bit word. Rewritten without early exits, the check that one lifetime may stand in for another gives the same answers on every input tried, and LLVM turns it into SIMD code with no branches. What it costs, and when the old loop is still faster.

• What comptime can do in Vx 1 October 2026

  A comptime block runs ordinary Vx code inside the compiler and puts the answer into the program as a constant. A prime count, a lookup table, a register machine and a sort checked on every build, a z3 proof that an index stays in bounds, and the errors the compiler gives when a block cannot run.

• Premature de-optimization is the root of all evil 28 September 2026

  Popular languages hide where memory is, hide the machine's shape, and lower to a low-level IR almost at once. Programs are slow before anyone writes them, and specialist tools win part of the speed back. Vx keeps memory and topology in the program until code generation.

• Nonlinear Turing tape 23 September 2026

  Mainstream languages assume one memory that never runs out: Turing's tape with a pointer size on it. An H100 node has host DRAM, 80 GiB of HBM, 50 MiB of L2 and 132 copies of 228 KiB of shared memory. Vx reads that list as a machine file and checks programs against it.

• float * of C/C++/Rust Fails to Distinguish Host DRAM from GPU HBM 23 September 2026

  malloc and cudaMalloc return different memory, and C, C++ and Rust give both the same type. The libraries each built the missing type, and lose it again at every raw pointer. In Vx the memory space is part of the type.

• Reverse engineering trait implementation to fix the Vx's 22 September 2026

  Two impls of one trait collided on one mangled name, and picking between them looked like a language design decision. Compiling the same program in Rust showed that Rust makes no such decision: it infers the argument's type, and defaults a bare integer literal to i32.

• The library that only worked on the machine that built it 22 September 2026

  Vx 0.0.1 could not run hello-world on any machine except the one that built it. Both macOS libraries recorded the release runner's own directory, and the smoke test passed throughout because it ran on that runner. Three more bugs were hiding behind it.

• Vx can model peak memory pressure at compile time 21 September 2026

  Add up every tensor a Llama-2-7B prefill puts on the GPU and you get 58.93 GiB, which does not fit a 40 GiB A100. The tensors resident at the same time peak at 15.94 GiB. The compiler computes the peak from the program's block structure, before the card is booked.

• Tiles, without a tile type 21 September 2026

  Can Vx be used for tile-based GPU programming, the way NVIDIA's cuTile is? A block loop, a tensor in Memory::SMEM and two barriers already compile to PTX with shared storage and bar.sync. The tile algebra does not exist.

• The compiler bugs we found by writing Vx's core library 21 September 2026

  A week ago Vx had no core library. Now it has six modules of it, a random-number library, and the compiler fixes that writing them forced — a method that dispatched differently between runs among them. Every line is Vx, with no extern in it.

• The NUMA that wasn't there 12 September 2026

  A CPU's NUMA domains are the same relation as two GPUs on a node, so the compiler already knew how to price them. It predicted a remote access costs 2.26× a local one and the hardware said 2.35× — after we threw away a day of measurements taken on a machine that reports two nodes it does not have.

• What a pointer forgets 10 September 2026

  A function stages a tensor to the GPU and returns it. In C++ the return type is float * and the fact is gone at the boundary. In Vx it is part of the type.

• Vx 0.0.1: placement in the type system 10 September 2026

  The first release. What the compiler checks, what it runs on, and the parts that are honestly not built yet.

Atom feed