Fixes from Jorge's ocean benchmarks

CHANGELOG.md

@@ -5,6 +5,17 @@ Booth — Changelog
 
 ### Frontend
 
+- 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
+  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)
+
 - #142: parse function pointer declarators, and constructors and destructors
   (Zane Hambly, 2026-07-27)
 
@@ -40,6 +51,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 +100,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,11 +140,6 @@ Booth — Changelog
   block from the kernel rather than a fixed 64 bytes
   (Zane Hambly, 2026-07-27)
 
-- 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)
 
@@ -132,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)
 

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.

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)

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;

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

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)",

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) */

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) {

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);

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) {

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:

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);
+}

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)