From 700370af4f88af9ebee94ab4ca72d36328c5160f Mon Sep 17 00:00:00 2001
From: ZaneHam <zanehambly@gmail.com>
Date: Thu, 3 Sep 2026 20:22:54 +1200
Subject: [PATCH] pragma once, and a buffer that never said it was full

---
 CHANGELOG.md      |   5 +
 lang/en.txt       |   1 +
 src/fe/bc_err.c   |   3 +-
 src/fe/bc_err.h   |   1 +
 src/fe/lexer.c    |   2 +
 src/fe/parser.c   |  30 +++++-
 src/fe/preproc.c  | 237 +++++++++++++++++++++++++++++++++++++++++-----
 src/fe/preproc.h  |  15 ++-
 src/fe/token.h    |   2 +
 src/main.c        |   7 ++
 tests/noretn.cu   |  16 ++++
 tests/ppcyc_a.cuh |   5 +
 tests/ppcyc_b.cuh |   5 +
 tests/ppcycle.cu  |  10 ++
 tests/ppovfl.cu   |  48 ++++++++++
 tests/ppovfl.opt  |   2 +
 tests/ppvarg.cu   |  27 ++++++
 tests/tphase.c    |  62 ++++++++++++
 18 files changed, 451 insertions(+), 27 deletions(-)
 create mode 100644 tests/noretn.cu
 create mode 100644 tests/ppcyc_a.cuh
 create mode 100644 tests/ppcyc_b.cuh
 create mode 100644 tests/ppcycle.cu
 create mode 100644 tests/ppovfl.cu
 create mode 100644 tests/ppovfl.opt
 create mode 100644 tests/ppvarg.cu

diff --git a/CHANGELOG.md b/CHANGELOG.md
index 04a9c7f..40e5e87 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -38,6 +38,11 @@ Booth — Changelog
 
 ### Frontend
 
+- 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)
diff --git a/lang/en.txt b/lang/en.txt
index 9ffe76e..a31f278 100644
--- a/lang/en.txt
+++ b/lang/en.txt
@@ -36,6 +36,7 @@ E049=#include: nesting too deep (max %d)
 E050=#include: cannot read '%s'
 E051=unknown directive: #%s
 E052=unterminated #if/#ifdef (missing %d #endif)
+E053=preprocessor output overflow (max %d bytes)
 
 # ---- Sema (E070-E099) ----
 E070=arrow on non-pointer
diff --git a/src/fe/bc_err.c b/src/fe/bc_err.c
index 4ac66e1..ce8e120 100644
--- a/src/fe/bc_err.c
+++ b/src/fe/bc_err.c
@@ -48,7 +48,8 @@ static const char *bc_dflt[BC_EID_MAX] = {
     /* E050 */ "#include: cannot read '%s'",
     /* E051 */ "unknown directive: #%s",
     /* E052 */ "unterminated #if/#ifdef (missing %d #endif)",
-    /* E053-E069 */ NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,
+    /* E053 */ "preprocessor output overflow (max %d bytes)",
+    /* E054-E069 */ NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,NULL,
 
     /* ---- Sema ---- */
     /* E070 */ "arrow on non-pointer",
diff --git a/src/fe/bc_err.h b/src/fe/bc_err.h
index 259bac0..89bcac5 100644
--- a/src/fe/bc_err.h
+++ b/src/fe/bc_err.h
@@ -54,6 +54,7 @@ typedef enum {
     BC_E050 = 50,  /* #include: cannot read '%s' */
     BC_E051 = 51,  /* unknown directive: #%s */
     BC_E052 = 52,  /* unterminated #if/#ifdef (missing %d #endif) */
+    BC_E053 = 53,  /* preprocessor output overflow (max %d bytes) */
 
     /* ---- Sema (E070-E099) ---- */
     BC_E070 = 70,  /* arrow on non-pointer */
diff --git a/src/fe/lexer.c b/src/fe/lexer.c
index c932845..0114402 100644
--- a/src/fe/lexer.c
+++ b/src/fe/lexer.c
@@ -10,6 +10,7 @@ typedef struct {
 } kw_entry_t;
 
 static const kw_entry_t keywords[] = {
+    {"_Noreturn",          TOK_NORETURN},
     {"__constant__",       TOK_CU_CONSTANT},
     {"__device__",         TOK_CU_DEVICE},
     {"__forceinline__",    TOK_CU_FORCEINLINE},
@@ -262,6 +263,7 @@ static const char *tok_names[] = {
     [TOK_CU_RESTRICT]   = "__restrict__",
     [TOK_CU_FORCEINLINE] = "__forceinline__",
     [TOK_CU_NOINLINE]   = "__noinline__",
+    [TOK_NORETURN]      = "_Noreturn",
     [TOK_EOF]           = "EOF",
     [TOK_ERROR]         = "ERROR",
 };
diff --git a/src/fe/parser.c b/src/fe/parser.c
index 346a732..3160649 100644
--- a/src/fe/parser.c
+++ b/src/fe/parser.c
@@ -241,6 +241,9 @@ static int is_storage_class(int type)
     switch (type) {
     case TOK_STATIC: case TOK_EXTERN: case TOK_REGISTER:
     case TOK_INLINE: case TOK_TYPEDEF:
+    /* _Noreturn carries no flag. Nothing in BIR or any backend asks whether
+       a callee returns, so accepting and dropping it is the whole semantics. */
+    case TOK_NORETURN:
         return 1;
     default:
         return 0;
@@ -728,6 +731,27 @@ static uint32_t parse_expr(parser_t *P, int min_prec)
 
 /* Where we divine the programmer's intent from a soup of keywords */
 
+/* [[noreturn]], [[maybe_unused]] and the rest. Booth models none of them and
+   they are all hints, so the run is consumed whole. Returns 1 if it ate one. */
+static int skatt(parser_t *P)
+{
+    int got = 0;
+    KA_GUARD(g, 64);
+    while (g-- && cur_type(P) == TOK_LBRACKET && peek_type(P, 1) == TOK_LBRACKET) {
+        int depth = 0;
+        KA_GUARD(h, 4096);
+        while (h-- && cur_type(P) != TOK_EOF) {
+            int t = cur_type(P);
+            if (t == TOK_LBRACKET) depth++;
+            else if (t == TOK_RBRACKET) depth--;
+            advance(P);
+            if (depth == 0) break;
+        }
+        got = 1;
+    }
+    return got;
+}
+
 static uint32_t parse_type_spec(parser_t *P, uint16_t *quals, uint16_t *cuda)
 {
     *quals = 0;
@@ -735,7 +759,9 @@ static uint32_t parse_type_spec(parser_t *P, uint16_t *quals, uint16_t *cuda)
 
     for (;;) {
         int ct = cur_type(P);
-        if (is_cuda_qualifier(ct)) {
+        if (ct == TOK_LBRACKET && peek_type(P, 1) == TOK_LBRACKET) {
+            if (!skatt(P)) break;
+        } else if (is_cuda_qualifier(ct)) {
             *cuda |= cuda_flag_for(ct);
             advance(P);
             /* __launch_bounds__(maxThreads[, minBlocks]) — actually read the numbers
@@ -1655,6 +1681,8 @@ static uint32_t parse_stmt(parser_t *P)
 
 static uint32_t parse_decl_or_stmt(parser_t *P)
 {
+    skatt(P);
+
     if (cur_type(P) == TOK_PP_LINE) {
         uint32_t pp = alloc_node(P, AST_PP_DIRECTIVE);
         P->nodes[pp].d.text.offset = cur(P)->offset;
diff --git a/src/fe/preproc.c b/src/fe/preproc.c
index 61fcaed..223a4b2 100644
--- a/src/fe/preproc.c
+++ b/src/fe/preproc.c
@@ -82,18 +82,64 @@ static int pp_is_ident_char(char c)
     return pp_is_ident_start(c) || (c >= '0' && c <= '9');
 }
 
+/* Comments are gone before macros are seen, everywhere except here: a body was
+ * kept whole, so '#define S 5 // why' commented out every line that used S, and
+ * a joined invocation would lose everything after a // on its first line.
+ * A block comment left open is not this function's business. */
+static uint32_t stripc(char *s, uint32_t n)
+{
+    uint32_t r = 0, w = 0;
+    while (r < n) {
+        if (s[r] == '"' || s[r] == '\'') {
+            char q = s[r];
+            s[w++] = s[r++];
+            while (r < n && s[r] != q) {
+                if (s[r] == '\\' && r + 1 < n) s[w++] = s[r++];
+                s[w++] = s[r++];
+            }
+            if (r < n) s[w++] = s[r++];
+            continue;
+        }
+        if (s[r] == '/' && r + 1 < n && s[r+1] == '/') break;
+        if (s[r] == '/' && r + 1 < n && s[r+1] == '*') {
+            uint32_t e = r + 2;
+            while (e + 1 < n && !(s[e] == '*' && s[e+1] == '/')) e++;
+            if (e + 1 >= n) break;
+            s[w++] = ' ';
+            r = e + 2;
+            continue;
+        }
+        s[w++] = s[r++];
+    }
+    return w;
+}
+
 /* ---- Output helpers ---- */
 
+/* One byte of out_max is held back so the buffer can always be terminated;
+ * the lexer reads until NUL and a full buffer used to run it into whatever
+ * static followed. */
+static int pp_room(preproc_t *pp, uint32_t need)
+{
+    if (pp->out_len + need < pp->out_max) return 1;
+    if (!pp->ovflw) {
+        pp->ovflw = 1;
+        pp_error(pp, BC_E053, pp->out_max);
+    }
+    return 0;
+}
+
 static void pp_emit_char(preproc_t *pp, char c)
 {
-    if (pp->out_len < pp->out_max)
+    if (pp_room(pp, 1))
         pp->out[pp->out_len++] = c;
 }
 
 static void pp_emit_str(preproc_t *pp, const char *s, uint32_t len)
 {
-    for (uint32_t i = 0; i < len && pp->out_len < pp->out_max; i++)
-        pp->out[pp->out_len++] = s[i];
+    if (!pp_room(pp, len)) return;
+    memcpy(pp->out + pp->out_len, s, len);
+    pp->out_len += len;
 }
 
 /* ---- String pool ---- */
@@ -127,16 +173,25 @@ static int pp_define_macro(preproc_t *pp, const char *name, uint32_t name_len,
                            const char *body, uint32_t body_len,
                            int num_params,
                            const char params[][BC_MAX_IDENT],
-                           const uint32_t param_lens[])
+                           const uint32_t param_lens[],
+                           int varia)
 {
     /* Check for redefinition — replace body if exists */
     pp_macro_t *existing = pp_find_macro(pp, name, name_len);
     if (existing) {
+        /* A header seen twice redefines every macro in it identically. The
+         * pool has no free list, so storing that again is pure waste. */
+        if (existing->body_len == body_len &&
+            existing->num_params == (int16_t)num_params &&
+            existing->varia == (uint8_t)(varia != 0) &&
+            memcmp(pp->pool + existing->body_off, body, body_len) == 0)
+            return BC_OK;
         existing->body_off = pool_add(pp, body, body_len);
         existing->body_len = (uint16_t)body_len;
         existing->num_params = (int16_t)num_params;
+        existing->varia = (uint8_t)(varia != 0);
         for (int i = 0; i < num_params && i < PP_MAX_PARAMS; i++) {
-            existing->param_off[i] = (uint16_t)pool_add(pp, params[i], param_lens[i]);
+            existing->param_off[i] = pool_add(pp, params[i], param_lens[i]);
             existing->param_len[i] = (uint8_t)param_lens[i];
         }
         return BC_OK;
@@ -153,9 +208,10 @@ static int pp_define_macro(preproc_t *pp, const char *name, uint32_t name_len,
     m->body_off = pool_add(pp, body, body_len);
     m->body_len = (uint16_t)body_len;
     m->num_params = (int16_t)num_params;
+    m->varia = (uint8_t)(varia != 0);
 
     for (int i = 0; i < num_params && i < PP_MAX_PARAMS; i++) {
-        m->param_off[i] = (uint16_t)pool_add(pp, params[i], param_lens[i]);
+        m->param_off[i] = pool_add(pp, params[i], param_lens[i]);
         m->param_len[i] = (uint8_t)param_lens[i];
     }
     return BC_OK;
@@ -626,6 +682,20 @@ static int64_t pp_eval_expr(preproc_t *pp, const char *expr, uint32_t len)
 #define PP_ARG_BUF  1024
 #define PP_TMP_BUF  8192
 
+/* __VA_ARGS__ is the last parameter and it swallows every argument from its
+ * own position onwards, so a parameter reference covers a range, not one arg. */
+static int vlast(const pp_macro_t *m, int pi)
+{
+    return m->varia && pi >= 0 && pi == m->num_params - 1;
+}
+
+static int varg(const pp_macro_t *m, int pi, int nargs)
+{
+    if (pi < 0) return pi;
+    int end = vlast(m, pi) ? nargs : pi + 1;
+    return end > nargs ? nargs : end;
+}
+
 static uint32_t pp_expand_text(preproc_t *pp,
                                const char *in, uint32_t in_len,
                                char *out, uint32_t out_max,
@@ -843,12 +913,16 @@ static uint32_t pp_expand_text(preproc_t *pp,
                             break;
                         }
                     }
-                    if (pi >= 0 && pi < nargs) {
+                    int a0 = pi, a1 = varg(m, pi, nargs);
+                    if (pi >= 0 && a0 < a1) {
                         if (slen < PP_TMP_BUF) subst[slen++] = '"';
-                        for (uint32_t k = 0; k < arg_lens[pi] && slen < PP_TMP_BUF; k++) {
-                            if (args[pi][k] == '"' || args[pi][k] == '\\')
-                                if (slen < PP_TMP_BUF) subst[slen++] = '\\';
-                            if (slen < PP_TMP_BUF) subst[slen++] = args[pi][k];
+                        for (int p = a0; p < a1; p++) {
+                            if (p > a0 && slen < PP_TMP_BUF) subst[slen++] = ',';
+                            for (uint32_t k = 0; k < arg_lens[p] && slen < PP_TMP_BUF; k++) {
+                                if (args[p][k] == '"' || args[p][k] == '\\')
+                                    if (slen < PP_TMP_BUF) subst[slen++] = '\\';
+                                if (slen < PP_TMP_BUF) subst[slen++] = args[p][k];
+                            }
                         }
                         if (slen < PP_TMP_BUF) subst[slen++] = '"';
                     }
@@ -868,9 +942,13 @@ static uint32_t pp_expand_text(preproc_t *pp,
                             break;
                         }
                     }
-                    if (pi >= 0 && pi < nargs) {
-                        for (uint32_t k = 0; k < arg_lens[pi] && slen < PP_TMP_BUF; k++)
-                            subst[slen++] = args[pi][k];
+                    int a0 = pi, a1 = varg(m, pi, nargs);
+                    if (pi >= 0 && (pi < nargs || vlast(m, pi))) {
+                        for (int p = a0; p < a1; p++) {
+                            if (p > a0 && slen < PP_TMP_BUF) subst[slen++] = ',';
+                            for (uint32_t k = 0; k < arg_lens[p] && slen < PP_TMP_BUF; k++)
+                                subst[slen++] = args[p][k];
+                        }
                     } else {
                         for (uint32_t k = ps; k < ps + plen && slen < PP_TMP_BUF; k++)
                             subst[slen++] = body[k];
@@ -940,6 +1018,7 @@ static void pp_dir_define(preproc_t *pp)
     char params[PP_MAX_PARAMS][BC_MAX_IDENT];
     uint32_t param_lens[PP_MAX_PARAMS];
     int num_params = -1; /* -1 = object-like */
+    int varia = 0;
 
     if (!pp_at_end(pp) && pp_cur(pp) == '(') {
         num_params = 0;
@@ -949,6 +1028,15 @@ static void pp_dir_define(preproc_t *pp)
         if (!pp_at_end(pp) && pp_cur(pp) != ')') {
             while (!pp_at_end(pp) && num_params < PP_MAX_PARAMS) {
                 pp_skip_hspace(pp);
+                if (pp_cur(pp) == '.' && pp_peek(pp, 1) == '.' &&
+                    pp_peek(pp, 2) == '.') {
+                    pp_advance(pp); pp_advance(pp); pp_advance(pp);
+                    memcpy(params[num_params], "__VA_ARGS__", 11);
+                    param_lens[num_params] = 11;
+                    num_params++;
+                    varia = 1;
+                    break;
+                }
                 if (!pp_is_ident_start(pp_cur(pp))) break;
                 uint32_t plen = pp_read_ident(pp, params[num_params], BC_MAX_IDENT);
                 param_lens[num_params] = plen;
@@ -961,7 +1049,12 @@ static void pp_dir_define(preproc_t *pp)
                     break;
             }
         }
-        pp_skip_hspace(pp);
+        /* Anything left before the ')' is a parameter list this preprocessor
+         * does not understand. Skipping it beats leaving it in the body,
+         * which is how '#define F(...)' used to expand to '...)'. */
+        KA_GUARD(g, PP_LINE_BUF);
+        while (!pp_at_end(pp) && pp_cur(pp) != ')' && pp_cur(pp) != '\n' && g--)
+            pp_advance(pp);
         if (!pp_at_end(pp) && pp_cur(pp) == ')')
             pp_advance(pp);
     }
@@ -971,13 +1064,15 @@ static void pp_dir_define(preproc_t *pp)
     /* Collect body (rest of logical line) */
     char body[PP_EXPAND_BUF];
     uint32_t body_len = pp_collect_line(pp, body, PP_EXPAND_BUF);
+    body_len = stripc(body, body_len);
 
     /* Trim trailing whitespace from body */
     while (body_len > 0 && (body[body_len-1] == ' ' || body[body_len-1] == '\t'))
         body_len--;
 
     pp_define_macro(pp, name, name_len, body, body_len,
-                    num_params, (const char (*)[BC_MAX_IDENT])params, param_lens);
+                    num_params, (const char (*)[BC_MAX_IDENT])params, param_lens,
+                    varia);
 }
 
 static void pp_dir_undef(preproc_t *pp)
@@ -1041,9 +1136,31 @@ static void pp_dir_error(preproc_t *pp)
     pp_error(pp, BC_E047, msg);
 }
 
+/* ---- #pragma once ---- */
+
+static int once_has(const preproc_t *pp, const char *path)
+{
+    for (uint32_t i = 0; i < pp->n_once; i++)
+        if (strcmp(pp->once[i], path) == 0) return 1;
+    return 0;
+}
+
+/* A full table only means the header gets read again, which is what every
+ * header did before this existed. Slower, never wrong. */
+static void once_add(preproc_t *pp, const char *path)
+{
+    if (once_has(pp, path)) return;
+    if (pp->n_once >= PP_MAX_ONCE) return;
+    snprintf(pp->once[pp->n_once++], BC_MAX_PATH, "%s", path);
+}
+
 static void pp_dir_pragma(preproc_t *pp)
 {
-    /* Just skip the line — could emit as a comment later */
+    pp_skip_hspace(pp);
+    if (pp->pos + 4 <= pp->src_len &&
+        memcmp(pp->src + pp->pos, "once", 4) == 0 &&
+        !pp_is_ident_char(pp_peek(pp, 4)))
+        once_add(pp, pp->filename);
     pp_skip_to_eol(pp);
 }
 
@@ -1161,6 +1278,9 @@ static void pp_dir_include(preproc_t *pp)
         return;
     }
 
+    if (once_has(pp, fullpath))
+        return;
+
     /* Guard against include depth overflow */
     if (pp->file_depth >= PP_MAX_FILE_DEPTH) {
         pp_error(pp, BC_E049, PP_MAX_FILE_DEPTH);
@@ -1208,6 +1328,57 @@ static int pp_pop_file(preproc_t *pp)
     return 1;
 }
 
+/* ---- Multi-line macro invocations ---- */
+
+/* Expansion runs a line at a time, so an argument list that opens on one line
+ * and closes on a later one has to be joined first. True when the line ends
+ * inside the argument list of a function-like macro. */
+static int ocall(preproc_t *pp, const char *s, uint32_t n)
+{
+    int depth = 0;
+    int bcmt = pp->in_bcmt;
+    uint32_t i = 0;
+
+    while (i < n) {
+        if (bcmt) {
+            if (s[i] == '*' && i + 1 < n && s[i+1] == '/') { bcmt = 0; i += 2; }
+            else i++;
+            continue;
+        }
+        if (s[i] == '/' && i + 1 < n && s[i+1] == '*') { bcmt = 1; i += 2; continue; }
+        if (s[i] == '/' && i + 1 < n && s[i+1] == '/') break;
+        if (s[i] == '"' || s[i] == '\'') {
+            char q = s[i++];
+            while (i < n && s[i] != q) {
+                if (s[i] == '\\' && i + 1 < n) i++;
+                i++;
+            }
+            i++;
+            continue;
+        }
+        if (depth > 0) {
+            if (s[i] == '(') depth++;
+            else if (s[i] == ')') depth--;
+            i++;
+            continue;
+        }
+        if (pp_is_ident_start(s[i])) {
+            uint32_t st = i;
+            while (i < n && pp_is_ident_char(s[i])) i++;
+            uint32_t j = i;
+            while (j < n && (s[j] == ' ' || s[j] == '\t')) j++;
+            if (j < n && s[j] == '(') {
+                pp_macro_t *m = pp_find_macro(pp, s + st, i - st);
+                if (m && m->num_params >= 0) { depth = 1; i = j + 1; }
+            }
+            continue;
+        }
+        i++;
+    }
+    return depth > 0;
+}
+
+
 /* ---- Main processing loop ---- */
 
 static void pp_process_directive(preproc_t *pp)
@@ -1290,7 +1461,7 @@ static void pp_process_directive(preproc_t *pp)
 
 int pp_process(preproc_t *pp)
 {
-    while (!pp_at_end(pp) || pp->file_depth > 0) {
+    while ((!pp_at_end(pp) || pp->file_depth > 0) && !pp->ovflw) {
         /* If we've reached the end of an included file, pop the file stack */
         if (pp_at_end(pp)) {
             if (!pp_pop_file(pp))
@@ -1330,7 +1501,24 @@ int pp_process(preproc_t *pp)
         pp->pos = line_start;
         char line[PP_LINE_BUF];
         uint32_t llen = pp_collect_line(pp, line, PP_LINE_BUF);
+
+        uint32_t joins = 0;
+        KA_GUARD(g, PP_MAX_JOIN);
+        while (g-- && llen + 2 < PP_LINE_BUF && ocall(pp, line, llen)) {
+            uint32_t nl = pp_cur(pp) == '\n' ? 1u
+                        : (pp_cur(pp) == '\r' && pp_peek(pp, 1) == '\n') ? 2u : 0u;
+            if (nl == 0) break;
+            llen = stripc(line, llen);
+            for (uint32_t k = 0; k < nl; k++)
+                pp_advance(pp);
+            line[llen++] = ' ';
+            llen += pp_collect_line(pp, line + llen, PP_LINE_BUF - llen);
+            joins++;
+        }
+        if (joins) llen = stripc(line, llen);
+
         pp_expand_and_emit(pp, line, llen);
+        while (joins--) pp_emit_char(pp, '\n');
 
         if (!pp_at_end(pp) && pp_cur(pp) == '\n') {
             pp_emit_char(pp, '\n');
@@ -1338,14 +1526,17 @@ int pp_process(preproc_t *pp)
         }
     }
 
+    /* Abandoning the run leaves the include stack loaded, and every entry
+     * owns a malloc'd buffer. */
+    KA_GUARD(g, PP_MAX_FILE_DEPTH + 1);
+    while (g-- && pp_pop_file(pp)) { }
+
     /* Check for unterminated conditionals */
-    if (pp->cond_depth > 0) {
+    if (pp->cond_depth > 0 && !pp->ovflw) {
         pp_error(pp, BC_E052, pp->cond_depth);
     }
 
-    /* Null-terminate output */
-    if (pp->out_len < pp->out_max)
-        pp->out[pp->out_len] = '\0';
+    pp->out[pp->out_len < pp->out_max ? pp->out_len : pp->out_max - 1] = '\0';
 
     return pp->num_errors > 0 ? BC_ERR_PREPROC : BC_OK;
 }
@@ -1385,5 +1576,5 @@ int pp_define(preproc_t *pp, const char *name, const char *value)
     uint32_t nlen = (uint32_t)strlen(name);
     uint32_t vlen = value ? (uint32_t)strlen(value) : 0;
     return pp_define_macro(pp, name, nlen, value ? value : "", vlen,
-                           -1, NULL, NULL);
+                           -1, NULL, NULL, 0);
 }
diff --git a/src/fe/preproc.h b/src/fe/preproc.h
index f1dada5..d66ed85 100644
--- a/src/fe/preproc.h
+++ b/src/fe/preproc.h
@@ -18,6 +18,8 @@
 #define PP_MAX_COND_DEPTH   64
 #define PP_MAX_FILE_DEPTH   16
 #define PP_MAX_EXPAND_DEPTH 32
+#define PP_MAX_ONCE         256
+#define PP_MAX_JOIN         256
 #define PP_POOL_SIZE        (256 * 1024)   /* 256KB for macro names + bodies */
 #define PP_LINE_BUF         (64 * 1024)    /* 64KB max logical line */
 #define PP_EXPAND_BUF       (128 * 1024)   /* 128KB expansion workspace */
@@ -28,9 +30,10 @@ typedef struct {
     uint32_t body_off;
     uint16_t body_len;
     int16_t  num_params;   /* -1 = object-like, 0+ = function-like */
-    uint16_t param_off[PP_MAX_PARAMS];
+    uint8_t  varia;        /* last param is __VA_ARGS__ */
+    uint32_t param_off[PP_MAX_PARAMS];
     uint8_t  param_len[PP_MAX_PARAMS];
-} pp_macro_t;   /* ~56 bytes */
+} pp_macro_t;
 
 typedef struct {
     int  active;        /* currently emitting code? */
@@ -80,6 +83,10 @@ typedef struct {
     pp_file_entry_t file_stack[PP_MAX_FILE_DEPTH];
     int         file_depth;
 
+    /* Resolved paths of files that asked for #pragma once */
+    char        once[PP_MAX_ONCE][BC_MAX_PATH];
+    uint32_t    n_once;
+
     /* Expansion workspace */
     char        line_buf[PP_LINE_BUF];
     char        exp_buf[PP_EXPAND_BUF];
@@ -88,6 +95,10 @@ typedef struct {
      * without this a macro named in comment prose gets substituted. */
     int         in_bcmt;
 
+    /* Output buffer filled. Reported once, then the run is abandoned; a
+     * truncated buffer is not a smaller program, it is a broken one. */
+    int         ovflw;
+
     /* Errors */
     bc_error_t  errors[BC_MAX_ERRORS];
     int         num_errors;
diff --git a/src/fe/token.h b/src/fe/token.h
index 4ffe907..f5c0a9f 100644
--- a/src/fe/token.h
+++ b/src/fe/token.h
@@ -153,6 +153,8 @@ typedef enum {
     TOK_CU_FORCEINLINE,     /* __forceinline__ */
     TOK_CU_NOINLINE,        /* __noinline__ */
 
+    TOK_NORETURN,           /* _Noreturn */
+
     TOK_EOF,
     TOK_ERROR,
     TOK_COUNT
diff --git a/src/main.c b/src/main.c
index 75b9044..aba15cb 100644
--- a/src/main.c
+++ b/src/main.c
@@ -727,6 +727,13 @@ int main(int argc, char *argv[])
             return prc != BC_OK ? 1 : 0;
         }
 
+        /* A truncated expansion is not shorter source, it is different
+         * source, and every phase after this would be reading a lie. */
+        if (pp->ovflw) {
+            free(pp);
+            return 1;
+        }
+
         lex_src = pp_out_buf;
         lex_len = pp->out_len;
         free(pp);
diff --git a/tests/noretn.cu b/tests/noretn.cu
new file mode 100644
index 0000000..94b2cc4
--- /dev/null
+++ b/tests/noretn.cu
@@ -0,0 +1,16 @@
+_Noreturn void nrt_die(const char *file, int line, const char *fmt, ...);
+
+[[noreturn]] void nrt_stop(int code);
+
+__device__ int nrt_bump(int a, [[maybe_unused]] int slack)
+{
+    return a + 1;
+}
+
+__global__ void nrt_kern(int *data, int n)
+{
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) {
+        data[i] = nrt_bump(data[i], 0);
+    }
+}
diff --git a/tests/ppcyc_a.cuh b/tests/ppcyc_a.cuh
new file mode 100644
index 0000000..a168913
--- /dev/null
+++ b/tests/ppcyc_a.cuh
@@ -0,0 +1,5 @@
+#pragma once
+
+#include "ppcyc_b.cuh"
+
+#define PPCYC_A 3
diff --git a/tests/ppcyc_b.cuh b/tests/ppcyc_b.cuh
new file mode 100644
index 0000000..39e456d
--- /dev/null
+++ b/tests/ppcyc_b.cuh
@@ -0,0 +1,5 @@
+#pragma once
+
+#include "ppcyc_a.cuh"
+
+#define PPCYC_B 4
diff --git a/tests/ppcycle.cu b/tests/ppcycle.cu
new file mode 100644
index 0000000..3343c9d
--- /dev/null
+++ b/tests/ppcycle.cu
@@ -0,0 +1,10 @@
+#include "ppcyc_a.cuh"
+#include "ppcyc_b.cuh"
+
+__global__ void ppcyc_kern(int *data, int n)
+{
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) {
+        data[i] = data[i] + PPCYC_A + PPCYC_B;
+    }
+}
diff --git a/tests/ppovfl.cu b/tests/ppovfl.cu
new file mode 100644
index 0000000..ba1b2ba
--- /dev/null
+++ b/tests/ppovfl.cu
@@ -0,0 +1,48 @@
+/* Expansion that outgrows the output buffer. The point is the
+ * diagnostic, not the program. */
+
+#define OVF1 zzzzzzzzzzzzzzzz,zzzzzzzzzzzzzzzz,zzzzzzzzzzzzzzzz,zzzzzzzzzzzzzzzz
+#define OVF2 OVF1 OVF1 OVF1 OVF1 OVF1 OVF1 OVF1 OVF1
+#define OVF3 OVF2 OVF2 OVF2 OVF2 OVF2 OVF2 OVF2 OVF2
+#define OVF4 OVF3 OVF3 OVF3 OVF3 OVF3 OVF3 OVF3 OVF3
+
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
+OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4 OVF4
diff --git a/tests/ppovfl.opt b/tests/ppovfl.opt
new file mode 100644
index 0000000..d95ace9
--- /dev/null
+++ b/tests/ppovfl.opt
@@ -0,0 +1,2 @@
+# Options for ppovfl.cu. See test_diag.opt for the format.
+xfail all  expansion outgrows the output buffer, refusal is the point
diff --git a/tests/ppvarg.cu b/tests/ppvarg.cu
new file mode 100644
index 0000000..4bd0fa4
--- /dev/null
+++ b/tests/ppvarg.cu
@@ -0,0 +1,27 @@
+#define PPV_DROP(...)
+#define PPV_KEEP(first, ...) first + __VA_ARGS__
+#define PPV_NAME(...) #__VA_ARGS__
+#define PPV_SHIFT 2 // shifting by a byte and a bit
+
+#define PPV_WRAP(decl, hint) decl
+
+PPV_DROP(9, 9)
+
+PPV_WRAP(
+    __device__ int ppv_add(int a, int b), // the hint is dropped
+    "use ppv_add2 instead");
+
+__device__ int ppv_add(int a, int b)
+{
+    return PPV_KEEP(a, b) >> PPV_SHIFT;
+}
+
+__global__ void ppv_kern(int *data, int n)
+{
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) {
+        data[i] = ppv_add(data[i], PPV_SHIFT);
+    }
+}
+
+const char *ppv_text = PPV_NAME(a, b);
diff --git a/tests/tphase.c b/tests/tphase.c
index 80cb4c9..73a670f 100644
--- a/tests/tphase.c
+++ b/tests/tphase.c
@@ -93,3 +93,65 @@ static void pha08(void)
     PASS();
 }
 TH_REG("pha", 8, "semantic analysis runs", pha08)
+
+/* ---- phase: declaration specifiers Booth does not model ---- */
+
+/* _Noreturn parsed as a type name, so the void after it was a syntax error,
+   and 61 of llama.cpp's 67 CUDA files start with one via GGML_NORETURN. */
+static void pha09(void)
+{
+    int rc = th_run(BC_BIN " --parse tests/noretn.cu", obuf, TH_BUFSZ);
+    CHEQ(rc, 0);
+    CHECK(strstr(obuf, "0 parse error(s)") != NULL);
+    CHECK(strstr(obuf, "(ident nrt_kern)") != NULL);
+    CHECK(strstr(obuf, "(ident nrt_bump)") != NULL);
+    PASS();
+}
+TH_REG("pha", 9, "_Noreturn and [[attributes]] are accepted", pha09)
+
+/* ---- phase: preprocessor, again ---- */
+
+/* ppcyc_a and ppcyc_b include each other. Without #pragma once the depth
+   guard is all that stops them, and it stopped ggml-cuda too. */
+static void pha10(void)
+{
+    int rc = th_run(BC_BIN " --pp tests/ppcycle.cu", obuf, TH_BUFSZ);
+    CHEQ(rc, 0);
+    CHECK(strstr(obuf, "error") == NULL);
+    CHECK(strstr(obuf, "3 + 4") != NULL);
+    PASS();
+}
+TH_REG("pha", 10, "#pragma once breaks an include cycle", pha10)
+
+/* Three that travel together: a body kept its trailing // comment, an
+   argument list that closed on a later line was never joined, and ... in a
+   parameter list ended up in the body rather than naming __VA_ARGS__. */
+static void pha11(void)
+{
+    int rc = th_run(BC_BIN " --pp tests/ppvarg.cu", obuf, TH_BUFSZ);
+    CHEQ(rc, 0);
+    CHECK(strstr(obuf, "...") == NULL);
+    CHECK(strstr(obuf, "decl") == NULL);
+    CHECK(strstr(obuf, "the hint is dropped") == NULL);
+    CHECK(strstr(obuf, "int ppv_add(int a, int b);") != NULL);
+    CHECK(strstr(obuf, ">> 2") != NULL);
+    CHECK(strstr(obuf, "\"a, b\"") != NULL);
+    PASS();
+}
+TH_REG("pha", 11, "variadic macros and multi-line calls", pha11)
+
+/* The output buffer used to fill and keep going, unterminated, so the lexer
+   read on into the source buffer that follows it and reported whatever it
+   found there. The symptom was an unterminated block comment. */
+static void pha12(void)
+{
+    int rc = th_run(BC_BIN " --nvidia-ptx tests/ppovfl.cu -o ppovfl_test.ptx",
+                    obuf, TH_BUFSZ);
+    CHNE(rc, 0);
+    CHECK(strstr(obuf, "E053") != NULL);
+    CHECK(strstr(obuf, "E002") == NULL);
+    CHECK(strstr(obuf, "E020") == NULL);
+    remove("ppovfl_test.ptx");
+    PASS();
+}
+TH_REG("pha", 12, "a truncated expansion is diagnosed", pha12)