pragma once, and a buffer that never said it was full

CHANGELOG.md

@@ -23,6 +23,11 @@ Booth — Changelog
   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)

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

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

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

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",
 };

src/fe/parser.c

@@ -270,6 +270,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;
@@ -757,14 +760,37 @@ 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;
     *cuda = 0;
 
     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
@@ -1693,6 +1719,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;