Booth 0.5.3

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