v0.5.3
Zaneham/Boothv0.5.3Sep 3, 2026by Zaneham
AI Summary
Booth introduces major new capabilities including full OCaml support, a vendored MLIR frontend, and enhanced CUDA compilation features, along with significant bug fixes and the addition of new C++ standard library features.
Key Highlights
- Full OCaml support with kernel examples and MLIR frontend vendoring
- Multiple translation units for CUDA compilation
- Added C++ features: variadic templates, fold expressions, and sizeof...
- Added intrinsics (__popc, __clz, __ffs, __brev) to all backends
- New reproducibility and mutation testing commands
Breaking Changes
- Parsing fix for parenthesized identifiers restores addition behavior for expressions like `(a) + (b)`
- Runtime split by execution context rather than vendor
New Features
- OCaml support and kernel examples
- MLIR frontend vendoring
- Multiple translation units for CUDA compilation
- Variadic template parameter packs, fold expressions, and sizeof...
- mma.sync and MFMA lowering
- Runtime split by execution context
- make repro and make mutate commands
- Logo added
Full Release Notes
Kia ora, G'day and Hello, Here is Booth 0.5.3! There's been some big changes recently which I am happy to show ya'll! OCaml makes an appearance and you can begin using OCaml and have been using it to make a few kernels in my own spare time. Quite a bit of the code has been ported over from my other OCaml projects and luckily I've been working within the OCaml compiler. I even did a write up on my website at https://zanehambly.com/ocaml. If you'd like to see an exemplar Kernel I put some of my finance education to the test which you can see at `src/ocaml/Asian.ml`. An MLIR frontend has also been vendored in with huge thanks to @certik. The code is largely his with some edits here and there to fit in with the makefile and is released under his license. CUDA has also seen a tonne of love with the addition of multiple translation units allowing the compilation of multiple `.cu` files with one invocation. Elsewhere: - `(a) + (b)` adds again. Any parenthesised identifier was being read as a cast without checking whether it named a type, so the left operand vanished with no diagnostic. `((a) + (b))` is what every defensive macro expands to, and one of Booth's own test fixtures had been quietly returning the wrong answer. - Variadic template parameter packs, fold expressions and `sizeof...`. - `mma.sync` and MFMA lower, checked against llvm-mc byte for byte. - `--bir-in` reads BIR text, so a compiler outside this tree can target Booth without linking against it. - `__popc`, `__clz`, `__ffs` and `__brev` on every backend. - The runtime is split by where it runs rather than by vendor. - Kauri is vendored as `src/kauri.h` under MIT, so `KA_GUARD`, `KA_CHK` and `KA_PNEW` are available tree-wide. - Tests are named for their family and position, `rvi01` and `tdf39` rather than `rv_isel_max_frame_slots_in_range`, after z390's `TESTDCB1`. - `make repro` compiles every test file twice under `--amdgpu`, `--nvidia-ptx` and `--ir` and compares the bytes, so the deterministic layout `bir.h` claims is checked rather than assumed. - `make mutate` bends one line of Booth at a time and checks the suite notices. - Bug fixes: Triton was answering a full block pool with index 0, which is a live block, so an arena exhaustion was invisible. Triton blocks are named, where all four blocks of a loop kernel had been carrying the kernel's own name. DCE and mem2reg move instructions without moving `inst_lines[]`, so every line number past the first deleted instruction pointed at the wrong source. And `--help` no longer calls `--metal` a stub; it is 838 lines of lowering and `--intel-spirv` is the 52-line one. - WE HAVE A LOGO!! woohoo The current target is getting llama.cpp compiling. All 67 of ggml-cuda's files now make it through the preprocessor, where none did before, which took `#pragma once`, variadic macros and a buffer that had been quietly truncating rather than saying so. What stops them now is C++ class templates, and that is the next block of work for me. I've been contributing to LLVM and reading a bit of Clang to see how they do it and I am lowkey dreading it. The docs have also been updated so please check them out if you'd like to see how to use these new features. Thanks to everyone who's pointed me in the right direction Kind regards, Zane