From b686d4a85d4eddf4c085fa86526bcf8c5afd3168 Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Thu, 6 Aug 2026 15:31:03 +1200 Subject: [PATCH 1/5] metal: refuse fp64 rather than narrowing to float --- src/metal/emit.c | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/src/metal/emit.c b/src/metal/emit.c index 06c110c..6edd3d0 100644 --- a/src/metal/emit.c +++ b/src/metal/emit.c @@ -132,8 +132,14 @@ static int mt_etype(metal_module_t *mm, uint32_t ti) switch (T->width) { case 16: return mt_wstr(mm, "half"); case 32: return mt_wstr(mm, "float"); - case 64: /* MSL does not have double on most Apple GPUs */ - return mt_wstr(mm, "float"); + case 64: + /* Apple GPUs have no fp64 and MSL has no double, so there is + * nothing honest to emit here. Silently narrowing to float + * would compile a kernel that quietly computes to half the + * precision the source asked for. */ + fprintf(stderr, "metal: fp64 is not available on Apple GPUs; " + "kernel uses double and cannot be lowered\n"); + return 0; default: return mt_wstr(mm, "float"); } case BIR_TYPE_BFLOAT: From 5bbb950f8e21459a04d7a1f10e63eb59f0aca082 Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Thu, 6 Aug 2026 15:31:03 +1200 Subject: [PATCH 2/5] fe: double fmax/fmin/fmod, and a call argument cap that does not lie --- lang/en.txt | 1 + src/amdgpu/isel.c | 2 +- src/barracuda.h | 2 ++ src/fe/bc_err.c | 3 ++- src/fe/bc_err.h | 1 + src/fe/sema.c | 14 +++++++++++--- src/ir/bir_lower.c | 30 +++++++++++++++++++----------- 7 files changed, 37 insertions(+), 16 deletions(-) diff --git a/lang/en.txt b/lang/en.txt index d7d37df..9ffe76e 100644 --- a/lang/en.txt +++ b/lang/en.txt @@ -50,6 +50,7 @@ E078=while-condition must be scalar type E079=do-while condition must be scalar type E080=switch expression must be integer type E081=__global__ function must return void +E082='%s' passes more than %d arguments # ---- Lowering (E100-E129) ---- E100=too many labels (max 256) diff --git a/src/amdgpu/isel.c b/src/amdgpu/isel.c index cf53055..f6d034f 100644 --- a/src/amdgpu/isel.c +++ b/src/amdgpu/isel.c @@ -1249,7 +1249,7 @@ static void isel_conversion(uint32_t idx, const bir_inst_t *I, int div) {BIR_CEIL,AMD_V_CEIL_F32},{BIR_FTRUNC,AMD_V_TRUNC_F32}, {BIR_RNDNE,AMD_V_RNDNE_F32}, }; - for (int mi = 0; mi < 11; mi++) { + for (int mi = 0; mi < (int)(sizeof m1 / sizeof m1[0]); mi++) { if (m1[mi].bo == I->op) { emit1(m1[mi].ao, mop_vreg_v((uint16_t)vr), ensure_vgpr(src)); break; diff --git a/src/barracuda.h b/src/barracuda.h index 66e7d67..9192623 100644 --- a/src/barracuda.h +++ b/src/barracuda.h @@ -16,6 +16,8 @@ #define BC_MAX_TOKENS (1 << 20) #define BC_MAX_IDENT 256 #define BC_MAX_ERRORS 64 +/* Arguments in one call. Real ocean kernels pass 23, so 16 was not enough. */ +#define BC_MAX_ARGS 64 #define BC_MAX_PATH 512 #define BC_MAX_DEPTH 256 #define CUDA_GLOBAL 0x0001 diff --git a/src/fe/bc_err.c b/src/fe/bc_err.c index 5f76224..4ac66e1 100644 --- a/src/fe/bc_err.c +++ b/src/fe/bc_err.c @@ -63,7 +63,8 @@ static const char *bc_dflt[BC_EID_MAX] = { /* E079 */ "do-while condition must be scalar type", /* E080 */ "switch expression must be integer type", /* E081 */ "__global__ function must return void", - /* E082-E099 */ NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL, + /* E082 */ "'%s' passes more than %d arguments", + /* E083-E099 */ NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL, /* ---- Lowering ---- */ /* E100 */ "too many labels (max 256)", diff --git a/src/fe/bc_err.h b/src/fe/bc_err.h index 50cb861..259bac0 100644 --- a/src/fe/bc_err.h +++ b/src/fe/bc_err.h @@ -68,6 +68,7 @@ typedef enum { BC_E079 = 79, /* do-while condition must be scalar type */ BC_E080 = 80, /* switch expression must be integer type */ BC_E081 = 81, /* __global__ function must return void */ + BC_E082 = 82, /* '%s' passes more than %d arguments */ /* ---- Lowering (E100-E129) ---- */ BC_E100 = 100, /* too many labels (max 256) */ diff --git a/src/fe/sema.c b/src/fe/sema.c index aa218ec..58f607b 100644 --- a/src/fe/sema.c +++ b/src/fe/sema.c @@ -623,7 +623,8 @@ static const cuda_builtin_t cuda_builtins[] = { {"__cosf",1,0,0},{"tanf",1,0,0},{"fabsf",1,0,0},{"fabs",1,0,0}, {"floorf",1,0,0},{"ceilf",1,0,0},{"truncf",1,0,0},{"roundf",1,0,0}, {"rintf",1,0,0},{"tanhf",1,0,0},{"copysignf",2,0,0}, - {"fmaxf",2,0,0},{"fminf",2,0,0},{"fmodf",2,0,0},{"powf",2,0,0},{"__powf",2,0,0}, + {"fmaxf",2,0,0},{"fminf",2,0,0},{"fmax",2,0,0},{"fmin",2,0,0}, + {"fmodf",2,0,0},{"powf",2,0,0},{"__powf",2,0,0}, {NULL, 0, 0, 0} }; @@ -966,13 +967,20 @@ static uint32_t check_expr(sema_ctx_t *S, uint32_t node) char cname[128]; get_text(S, callee_n, cname, sizeof(cname)); - uint32_t arg_types[16]; + uint32_t arg_types[BC_MAX_ARGS]; int nargs = 0; uint32_t arg = ND(S, callee_n)->next_sibling; - while (arg && nargs < 16) { + while (arg && nargs < BC_MAX_ARGS) { arg_types[nargs++] = check_expr(S, arg); arg = ND(S, arg)->next_sibling; } + /* Running past the cap used to leave nargs sitting at it, so the arity + * checks below reported a count we had invented rather than the one + * the call actually has. Say what happened instead of guessing. */ + if (arg) { + sema_error(S, node, BC_E082, cname, BC_MAX_ARGS); + return annotate(S, node, st_int(S)); + } /* ---- Cooperative groups: namespace::func() and handle.method() ---- */ if (ND(S, callee_n)->type == AST_SCOPE_RES) { diff --git a/src/ir/bir_lower.c b/src/ir/bir_lower.c index 8e8ba51..ea4100c 100644 --- a/src/ir/bir_lower.c +++ b/src/ir/bir_lower.c @@ -1410,7 +1410,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) {"atomicMax", BIR_ATOMIC_MAX}, {"atomicExch",BIR_ATOMIC_XCHG}, }; int matched = 0; - for (int bi = 0; bi < 8; bi++) { + for (int bi = 0; bi < (int)(sizeof atab / sizeof atab[0]); bi++) { if (strcmp(cname, atab[bi].n) != 0) continue; uint32_t an = ND(L, callee_n)->next_sibling; uint32_t a0 = lower_expr(L, an); @@ -1470,7 +1470,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) {"__shfl_down", BIR_SHFL_DOWN,0}, {"__shfl_xor", BIR_SHFL_XOR, 0}, }; - for (int bi = 0; bi < 8; bi++) { + for (int bi = 0; bi < (int)(sizeof stab / sizeof stab[0]); bi++) { if (strcmp(cname, stab[bi].n) != 0) continue; uint32_t sa[4]; int sn = 0; @@ -1502,7 +1502,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) {"__any", BIR_VOTE_ANY, 0}, {"__all", BIR_VOTE_ALL, 0}, }; - for (int bi = 0; bi < 6; bi++) { + for (int bi = 0; bi < (int)(sizeof vtab / sizeof vtab[0]); bi++) { if (strcmp(cname, vtab[bi].n) != 0) continue; uint32_t an = ND(L, callee_n)->next_sibling; uint32_t a0, a1; @@ -1672,7 +1672,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) {"floorf",BIR_FLOOR},{"ceilf",BIR_CEIL}, {"truncf",BIR_FTRUNC},{"roundf",BIR_RNDNE},{"rintf",BIR_RNDNE}, }; - for (int mi = 0; mi < 16; mi++) { + for (int mi = 0; mi < (int)(sizeof mt1 / sizeof mt1[0]); mi++) { if (strcmp(cname, mt1[mi].n) != 0) continue; uint32_t an = ND(L, callee_n)->next_sibling; uint32_t v = lower_expr(L, an); @@ -1687,8 +1687,9 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) { static const struct { const char *n; uint16_t op; } mt2[] = { {"fmaxf",BIR_FMAX},{"fminf",BIR_FMIN},{"fmodf",BIR_FREM}, + {"fmax",BIR_FMAX},{"fmin",BIR_FMIN},{"fmod",BIR_FREM}, }; - for (int mi = 0; mi < 3; mi++) { + for (int mi = 0; mi < (int)(sizeof mt2 / sizeof mt2[0]); mi++) { if (strcmp(cname, mt2[mi].n) != 0) continue; uint32_t an = ND(L, callee_n)->next_sibling; uint32_t a0 = lower_expr(L, an); @@ -1861,7 +1862,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) {"f64_16x16x4f64", 21}, }; const char *sfx = cname + 22; - for (int mi = 0; mi < 22; mi++) { + for (int mi = 0; mi < (int)(sizeof mfma_tab / sizeof mfma_tab[0]); mi++) { if (strcmp(sfx, mfma_tab[mi].sfx) != 0) continue; /* 3 args: A, B, C(accum) */ uint32_t an = ND(L, callee_n)->next_sibling; @@ -1888,7 +1889,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) {"num_groups", BIR_GRID_DIM}, }; const char *rest = cname + 11; - for (int oi = 0; oi < 4; oi++) { + for (int oi = 0; oi < (int)(sizeof ockl / sizeof ockl[0]); oi++) { if (strcmp(rest, ockl[oi].sfx) != 0) continue; /* arg is literal dim (0/1/2) */ uint32_t an = ND(L, callee_n)->next_sibling; @@ -1913,7 +1914,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) {"floor", 5, BIR_FLOOR}, {"ceil", 4, BIR_CEIL}, {"trunc", 5, BIR_FTRUNC},{"rint", 4, BIR_RNDNE}, }; - for (int oi = 0; oi < 8; oi++) { + for (int oi = 0; oi < (int)(sizeof ou / sizeof ou[0]); oi++) { if (strncmp(rest, ou[oi].n, ou[oi].len) == 0 && rest[ou[oi].len] == '_') { uint32_t an = ND(L, callee_n)->next_sibling; @@ -1928,7 +1929,7 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) static const struct { const char *n; size_t len; uint16_t op; } ob[] = { {"fmax", 4, BIR_FMAX}, {"fmin", 4, BIR_FMIN}, }; - for (int oi = 0; oi < 2; oi++) { + for (int oi = 0; oi < (int)(sizeof ob / sizeof ob[0]); oi++) { if (strncmp(rest, ob[oi].n, ob[oi].len) == 0 && rest[ob[oi].len] == '_') { uint32_t an = ND(L, callee_n)->next_sibling; @@ -1978,13 +1979,20 @@ static uint32_t lower_expr(lower_t *L, uint32_t node) uint32_t ret_t = L->M->types[ftype].inner; /* Lower arguments */ - uint32_t args[16]; + uint32_t args[BC_MAX_ARGS]; int nargs = 0; uint32_t arg = ND(L, callee_n)->next_sibling; - while (arg && nargs < 16) { + while (arg && nargs < BC_MAX_ARGS) { args[nargs++] = lower_expr(L, arg); arg = ND(L, arg)->next_sibling; } + /* Sema rejects this first, so reaching it means the two caps have + * drifted apart. Dropping the tail would emit a call with the wrong + * operands and no sign anything was lost. */ + if (arg) { + lower_error(L, node, BC_E082, "call", BC_MAX_ARGS); + return BIR_VAL_NONE; + } if (1 + nargs <= BIR_OPERANDS_INLINE) { uint32_t inst = emit(L, BIR_CALL, ret_t, (uint8_t)(1+nargs), 0); From 3efef9b2fb43eb1eba73c0381bdcbabe55122de9 Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Thu, 6 Aug 2026 15:31:03 +1200 Subject: [PATCH 3/5] driver: semantic errors fail the compile --- CHANGELOG.md | 29 +++++++++++++++++++++ src/main.c | 9 +++++++ tests/many_args.cu | 24 +++++++++++++++++ tests/terrs.c | 65 ++++++++++++++++++++++++++++++++++++++++++++++ 4 files changed, 127 insertions(+) create mode 100644 tests/many_args.cu diff --git a/CHANGELOG.md b/CHANGELOG.md index 095b46f..b55c8b1 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -5,6 +5,21 @@ Booth — Changelog ### Frontend +- double-precision `fmax`, `fmin` and `fmod`. The ocean kernels in Jorge + Galvez's do-concurrent benchmarks call them, and only the `f`-suffixed + single-precision forms were recognised + (Zane Hambly, 2026-08-06) + +- raise the cap on arguments in one call to 64, and say so when a call goes + past it. Sema stopped counting at 16 and then reported an arity mismatch + against the count it had stopped at, so a correct 23-argument call in the + same ocean benchmarks was rejected and told the wrong number + (Zane Hambly, 2026-08-06) + +- size the builtin tables from the tables themselves rather than a + hand-counted length that had to be edited in step + (Zane Hambly, 2026-08-06) + - #142: parse function pointer declarators, and constructors and destructors (Zane Hambly, 2026-07-27) @@ -40,6 +55,11 @@ Booth — Changelog where it cannot be done (Zane Hambly, 2026-07-26) +- metal: refuse a kernel that uses `double` rather than narrowing it to + `float`. Apple GPUs have no fp64, and quietly halving the precision the + source asked for is worse than saying so + (Zane Hambly, 2026-08-06) + - a divergent return masks lanes instead of ending the wave, so AMD kernels no longer lose the lanes that did not take the branch (Zane Hambly, 2026-07-25) @@ -84,6 +104,11 @@ Booth — Changelog ### Driver +- semantic errors fail the compile. Every mode but `--sema` printed them and + then carried on into codegen, wrote an output file and exited zero, so a + build system saw a clean compile of source we had already rejected + (Zane Hambly, 2026-08-06) + - collapse the C99 mode gates into one cascade (Zane Hambly, 2026-07-23) @@ -119,6 +144,10 @@ Booth — Changelog block from the kernel rather than a fixed 64 bytes (Zane Hambly, 2026-07-27) +- cover the 23-argument call, the argument cap, and semantic errors being + fatal + (Zane Hambly, 2026-08-06) + - build an example CMake consumer against a staged install, and check the target list in the package config has not drifted from the flags the compiler accepts diff --git a/src/main.c b/src/main.c index 70a5f3b..d5f5326 100644 --- a/src/main.c +++ b/src/main.c @@ -966,6 +966,15 @@ int main(int argc, char *argv[]) free(sema_ctx); return sema_rc; } + + /* Only --sema used to act on these. Every other mode printed the + * errors, carried on into codegen, wrote an output file and exited + * zero, so a build system saw a clean compile and a kernel that + * had been lowered from source we had already rejected. */ + if (sema_ctx->num_errors > 0) { + free(sema_ctx); + return 1; + } } if (want_bir && P.num_errors == 0) { diff --git a/tests/many_args.cu b/tests/many_args.cu new file mode 100644 index 0000000..39a40a2 --- /dev/null +++ b/tests/many_args.cu @@ -0,0 +1,24 @@ +/* A 23-argument device call, the shape Jorge Galvez's ocean kernels use. + * The frontend used to stop counting at 16 and then blame the call for an + * argument count it had invented. */ + +__device__ double acc(const double *a, const double *b, const double *c, + const double *d, const double *e, const double *f, + const double *g, const double *h, const double *i, + const double *j, const double *k, const double *l, + const double *m, + double p, double q, double r, double s, + int u, int v, int w, int nx, int ny, int nz) +{ + return a[u] + b[v] + c[w] + d[0] + e[0] + f[0] + g[0] + h[0] + + i[0] + j[0] + k[0] + l[0] + m[0] + + p + q + r + s + (double)(nx + ny + nz); +} + +__global__ void many(const double *x, double *out, int n) +{ + int t = blockIdx.x * blockDim.x + threadIdx.x; + if (t < n) + out[t] = acc(x, x, x, x, x, x, x, x, x, x, x, x, x, + 1.0, 2.0, 3.0, 4.0, t, t, t, n, n, n); +} diff --git a/tests/terrs.c b/tests/terrs.c index e2e5a69..11026af 100644 --- a/tests/terrs.c +++ b/tests/terrs.c @@ -70,3 +70,68 @@ static void err_diag_render(void) PASS(); } TH_REG("errors", err_diag_render) + +/* ---- calls: past sixteen arguments ---- + * Jorge Galvez's ocean kernels pass 23. Sema stopped counting at 16 and then + * reported an arity mismatch against a count it had made up. */ + +static void arg_many(void) +{ + int rc = th_run(BC_BIN " --nvidia-ptx tests/many_args.cu -o many_args.ptx", + obuf, TH_BUFSZ); + CHEQ(rc, 0); + CHECK(strstr(obuf, "error") == NULL); + remove("many_args.ptx"); + PASS(); +} +TH_REG("errors", arg_many) + +/* ---- calls: past the cap ---- + * Overflowing has to say so. Silently dropping the tail would emit a call + * with the wrong operands and nothing to show for it. */ + +static void arg_cap(void) +{ + FILE *f = fopen("argcap_test.cu", "w"); + CHECK(f != NULL); + fprintf(f, "__device__ double g(double a){return a;}\n"); + fprintf(f, "__global__ void k(double *o){ o[0] = g(0.0"); + for (int i = 1; i < 70; i++) fprintf(f, ",%d.0", i); + fprintf(f, "); }\n"); + fclose(f); + + int rc = th_run(BC_BIN " --nvidia-ptx argcap_test.cu -o argcap_test.ptx", + obuf, TH_BUFSZ); + CHNE(rc, 0); + CHECK(strstr(obuf, "E082") != NULL); + remove("argcap_test.cu"); + remove("argcap_test.ptx"); + PASS(); +} +TH_REG("errors", arg_cap) + +/* ---- errors: sema errors are fatal ---- + * They used to be printed and then ignored by every mode but --sema, so the + * backend ran on source we had already rejected and the exit status said the + * compile went fine. */ + +static void err_sema_fatal(void) +{ + FILE *f = fopen("semafail_test.cu", "w"); + CHECK(f != NULL); + fprintf(f, "__global__ void k(float *o){ o[0] = nosuchfn(1, 2); }\n"); + fclose(f); + + remove("semafail_test.ptx"); + int rc = th_run(BC_BIN " --nvidia-ptx semafail_test.cu -o semafail_test.ptx", + obuf, TH_BUFSZ); + CHNE(rc, 0); + /* and nothing written, so a build system cannot pick up a stale artefact */ + FILE *o = fopen("semafail_test.ptx", "r"); + CHECK(o == NULL); + if (o) fclose(o); + remove("semafail_test.cu"); + remove("semafail_test.ptx"); + PASS(); +} +TH_REG("errors", err_sema_fatal) From b7fc2685a643bcb4cc94ac2e2e0e1d19752a73ba Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Thu, 6 Aug 2026 15:36:26 +1200 Subject: [PATCH 4/5] docs: credit Jorge Galvez for the ocean benchmarks --- CHANGELOG.md | 6 +++--- README.md | 1 + 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index b55c8b1..c7c8835 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -5,9 +5,9 @@ Booth — Changelog ### Frontend -- double-precision `fmax`, `fmin` and `fmod`. The ocean kernels in Jorge - Galvez's do-concurrent benchmarks call them, and only the `f`-suffixed - single-precision forms were recognised +- double-precision `fmax`, `fmin` and `fmod`. The ocean kernels in + [Jorge Galvez](https://github.com/JorgeG94)'s do-concurrent benchmarks call + them, and only the `f`-suffixed single-precision forms were recognised (Zane Hambly, 2026-08-06) - raise the cap on arguments in one call to 64, and say so when a call goes diff --git a/README.md b/README.md index 4ddc808..e9a153c 100644 --- a/README.md +++ b/README.md @@ -66,6 +66,7 @@ Based in New Zealand, where it's already tomorrow and the GPUs are just as confu ## Acknowledgements - **Fernando Magno Quintão Pereira** and the **Compilers Lab at UFMG** (Universidade Federal de Minas Gerais). Fernando reached out after seeing the project, pointed me to the divergence analysis papers, and offered guidance. The SSA register allocator exists because of that conversation. +- **[Jorge Galvez](https://github.com/JorgeG94)** for sending me his `do concurrent` Fortran benchmarks and letting me run tests on them. Three real frontend bugs turned up in an afternoon, which is exactly what you want somebody else's code to do. - **The academic community**: Cooper, Harvey & Kennedy for dominators; Braun & Hack for SSA spilling; Sampaio, Souza, Collange & Pereira for divergence analysis. I'm just a hobbyist who reads papers and writes C. The actual hard work was done by the researchers. - **Steven Muchnick** for *Advanced Compiler Design and Implementation*. If this compiler does anything right, that book is why. - **Low Level** for the Zero to Hero C course and the YouTube channel. That's where I learnt C. From 81b60a7cc443fb5efaeae91d40a22e42db948c0f Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Thu, 6 Aug 2026 15:52:31 +1200 Subject: [PATCH 5/5] docs: trim the changelog to user-visible changes --- CHANGELOG.md | 17 ----------------- 1 file changed, 17 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index c7c8835..05ab037 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -16,10 +16,6 @@ Booth — Changelog same ocean benchmarks was rejected and told the wrong number (Zane Hambly, 2026-08-06) -- size the builtin tables from the tables themselves rather than a - hand-counted length that had to be edited in step - (Zane Hambly, 2026-08-06) - - #142: parse function pointer declarators, and constructors and destructors (Zane Hambly, 2026-07-27) @@ -144,15 +140,6 @@ Booth — Changelog block from the kernel rather than a fixed 64 bytes (Zane Hambly, 2026-07-27) -- cover the 23-argument call, the argument cap, and semantic errors being - fatal - (Zane Hambly, 2026-08-06) - -- build an example CMake consumer against a staged install, and check the - target list in the package config has not drifted from the flags the - compiler accepts - (Zane Hambly, 2026-08-06) - - validate Tensix ELFs against tt-metal's loader and run RV64 under QEMU (Zane Hambly, 2026-07-23) @@ -161,10 +148,6 @@ Booth — Changelog ### Documentation -- document consuming Booth from CMake in `docs/cmake.md`, and link it from the - README - (Zane Hambly, 2026-08-06) - - drop the LLVM requirement from the usage documentation (Zane Hambly, 2026-07-27)