Open in tt-awesome →

Booth

community★ featured
by Zaneham · C · Apache-2.0 · 1738⭐ ·

An open-source CUDA, HIP, and Triton compiler with no LLVM anywhere in the path. Takes the same sources you would hand to nvcc, ROCm, or Triton's JIT and emits AMD RDNA 2/3/4 binaries, NVIDIA PTX, Tenstorrent Metalium C++, native RV32IM, or plain x86-64 — so a Triton matmul can run on a laptop that has never seen a GPU. Also reads a deliberate subset of MLIR (`func.func` plus the `arith` dialect) and Fortran `do concurrent` kernels via LFortran. Formerly BarraCUDA; renamed to honour Kathleen Booth.

📦 Repo
LATEST v0.5.3 2026-09-03T09:44:42Z Release notes ↗
3 previous releases
v0.5.2 2026-08-08T03:05:39Z
v5.01 2026-07-14T08:02:45Z
v0.5.0 2026-05-29T04:30:28Z
See all releases on GitHub ↗
Booth — Changelog
=================

## Booth 0.5.3

### Runtime

- #169: the runtime is split by where it runs, the `BC_ERR_*` codes no
  longer collide, and the examples and NVIDIA harness are built
  (Zane Hambly, 2026-08-23)


### Frontend

- variadic template parameter packs, several `.cu` files as separate
  translation units, `mma.sync` and `mfma` lowering, and an i1 that no
  longer strides by zero (Zane Hambly, 2026-09-03)

- `(a) + (b)` adds again; the parser treated any parenthesised identifier as a
  type name without asking whether it named one, so the left operand vanished
  into a cast with no diagnostic (Zane Hambly, 2026-09-03)

- the cast test is now the type name registry, so the registry has to be
  complete. Template type parameters, `using X = T` aliases and the type names
  sema resolves without a typedef (`size_t`, `uint32_t`, `float4` and the rest)
  all reach it. A compound literal through a typedef, `(pair){1, 2}`, parses
  for the first time, and `sizeof(name)` where the name is a type reads as a
  type rather than an expression (Zane Hambly, 2026-09-03)

- llama.cpp's ggml-cuda preprocesses, all 67 files; `#pragma once` is
  honoured, variadic and multi-line macro invocations expand, and an
  expansion too big for the output buffer is E053 rather than an
  unterminated buffer the lexer reads past (Zane Hambly, 2026-09-03)

- `kath --mlir` reads MLIR text, no LLVM in the path. Čertík's pure-C
  reader vendored under `src/mlir/vendor` (mlir 826b69c9, corec a160199d),
  reached only through `src/mlir/mlir_fe.c` (Zane Hambly, 2026-08-11)

- `src/mlir/lower.c` walks the parsed module into BIR: `func.func`, `return`,
  `arith.constant` and every arith binop, compare and conversion the reader
  classifies. From there it is the pipeline CUDA and Triton already use, and
  MLIR reaches all four backends. `--mlir --pp` reprints instead
  (Zane Hambly, 2026-08-11)

- an op outside the subset stops the lowering and names itself. Skipping it
  would leave a function that compiles and computes something else
  (Zane Hambly, 2026-08-11)

- five fixes to the vendored reader, all worth upstreaming, and four of them
  are `func.func` being unfinished where `tt.func` is not: `parser_init`
  renamed off Booth's own, `parser_error`'s `exit(1)` replaced by a
  `mlir_parse_fail()` the linker supplies, `func.func` binding its arguments
  before parsing the body rather than after, `func.func` accepting the
  `attributes` clause where MLIR actually writes it, and `arith.xori`,
  `shli` and `shrsi` added to `op_string_to_type`, which the printer could
  already write but the parser could not read back
  (Zane Hambly, 2026-08-11)

- `ml_parse` resets the reader's process-wide type interning, which upstream
  assumes one context per process. Without it a closed context left the next
  parse in freed memory (Zane Hambly, 2026-08-11)

- the Triton lowering records pool overflow through `bir_pfull`, which the C99
  one already did and it never has. I
cuda hip triton fortran mlir compiler cross-platform metalium rv32im no-llvm
blackhole