Variadic packs, rank-2 matrix, multiple TUs and the i1 fixes

.gitignore

@@ -53,6 +53,12 @@ __pycache__/
 /tri_*.ttinsn
 /tri_*.elf
 
+# Line-splicing fixtures. The suite writes these itself because
+# .gitattributes checks the tree out as eol=lf, so a committed CRLF
+# fixture would arrive as an LF one and test nothing.
+/tests/spl_lf.cu
+/tests/spl_crlf.cu
+
 # Ad-hoc test kernels and runners that piled up during the Moa GPU
 # port and the parameter-count and scratch-allocation investigations.
 # Listed by name rather than pattern so future intentional tests are

CHANGELOG.md

@@ -12,6 +12,10 @@ Booth — Changelog
 
 ### 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)

Makefile

@@ -127,7 +127,7 @@ TARGET  = kath
 # strict flags without libhsa or libcuda having to be present. Neither is a
 # trunner test; both carry their own main and want real hardware.
 EXAMPLE_SRC = $(wildcard examples/*.c)
-HOSTCHK     = $(OBJDIR)/tests/tnv_rt.o               $(patsubst examples/%.c,$(OBJDIR)/examples/%.o,$(EXAMPLE_SRC))
+HOSTCHK     = $(OBJDIR)/tests/tnv_rt.o $(OBJDIR)/tests/tnv_i1.o $(OBJDIR)/tests/tnv_bstr.o $(OBJDIR)/tests/gpu_mma.o $(patsubst examples/%.c,$(OBJDIR)/examples/%.o,$(EXAMPLE_SRC))
 
 all: $(TARGET) $(ALT_RT) $(HOSTCHK)
 
@@ -157,6 +157,7 @@ TSRC    = tests/tmain.c tests/tsmoke.c tests/tcomp.c tests/tenc.c \
           tests/tra_ssa.c \
           tests/tguard.c \
           tests/ttriton.c \
+          tests/tmma.c \
           tests/ttdf.c \
           tests/ttmc.c \
           tests/trv_enc.c tests/trv_buf.c tests/trv_elf.c tests/trv_isel.c \
@@ -168,7 +169,9 @@ TSRC    = tests/tmain.c tests/tsmoke.c tests/tcomp.c tests/tenc.c \
           tests/trpi.c \
           tests/tmlir.c \
           tests/tbir.c \
-          tests/tocm.c
+          tests/tocm.c \
+          tests/tpack.c \
+          tests/tmtu.c
 
 TOBJS   = $(TSRC:%.c=$(OBJDIR)/%.o)
 COBJS   = $(OBJDIR)/src/kauri_impl.o $(OBJDIR)/src/ir/bir.o $(OBJDIR)/src/ir/bir_print.o $(OBJDIR)/src/ir/bir_lower.o $(OBJDIR)/src/ir/bir_mem2reg.o $(OBJDIR)/src/ir/bir_cfold.o $(OBJDIR)/src/ir/bir_dce.o $(OBJDIR)/src/ir/bir_struct.o $(OBJDIR)/src/ir/bir_insert.o $(OBJDIR)/src/ir/bir_sroa.o $(OBJDIR)/src/ir/bir_inline.o \

lang/en.txt

@@ -83,6 +83,10 @@ E123=expected identifier after 'import'
 E124=expected attribute name after '.'
 E125=unrecognised expression
 
+# ---- Linking translation units (E126-E127) ----
+E126='%s' is defined in more than one translation unit
+E127=too many globals to reference (max %d)
+
 # ---- ABEND Messages (hex IDs) ----
 A0C1=illegal GPU instruction
 A0C4=memory access violation (page not mapped)

mfrg.s

@@ -0,0 +1,40 @@
+    .amdgcn_target "amdgcn-amd-amdhsa--gfx942"
+    .text
+
+    .globl mf16
+    .p2align 8
+    .type mf16,@function
+mf16:
+    ; 1 SGPRs, 248 VGPRs, 0 LDS bytes, 0 scratch bytes
+.LBB0:
+    s_load_dwordx2 s[4:5], s[0:1], 0
+    s_load_dwordx2 s[6:7], s[0:1], 8
+    s_load_dwordx2 s[8:9], s[0:1], 16
+    v_mov_b32 v6, 0
+    v_mov_b32 v5, 0
+    v_mov_b32 v0, 0
+    v_add_u32 v4, 4, v5
+    v_add_u32 v3, 4, v0
+    s_waitcnt lgkmcnt(0)
+    global_load_dword v224, v6, s[8:9]
+    v_add_u32 v2, 4, v6
+    global_load_dword v240, v0, s[4:5]
+    v_add_u32 v1, 8, v6
+    global_load_dword v244, v5, s[6:7]
+    v_add_u32 v0, 12, v6
+    global_load_dword v245, v4, s[6:7]
+    global_load_dword v227, v0, s[8:9]
+    global_load_dword v241, v3, s[4:5]
+    global_load_dword v225, v2, s[8:9]
+    global_load_dword v226, v1, s[8:9]
+    s_waitcnt vmcnt(0)
+    v_mfma_f32_16x16x16f16 v[208:211], v[240:241], v[244:245], v[224:227]
+    v_add_u32 v1, 4, v6
+    global_store_dword v6, v208, s[8:9]
+    v_add_u32 v0, 8, v6
+    global_store_dword v1, v209, s[8:9]
+    v_add_u32 v1, 12, v6
+    global_store_dword v0, v210, s[8:9]
+    global_store_dword v1, v211, s[8:9]
+    s_endpgm
+

src/amdgpu/amdgpu.h

@@ -20,10 +20,6 @@
 #define AMD_WAVE_SIZE       32
 #define AMD_WAVE64          64
 
-/* Pre-loaded SGPRs — layout depends on kernel needs.
- * Default (no dispatch_ptr): s[0:1]=kernarg, s2+=TGID → reserved=3
- * With dispatch_ptr:         s[0:1]=dispatch, s[2:3]=kernarg, s4+=TGID → reserved=5
- * Actual positions computed per-kernel in isel. */
 #define AMD_KERN_MIN_RESERVED  2  /* absolute floor: kernarg pair */
 
 /* Pre-loaded VGPRs */
@@ -287,6 +283,9 @@ typedef enum {
     AMD_V_MFMA_F32_32X32X16_BF8_FP8,
     AMD_V_MFMA_F32_32X32X16_BF8_BF8,
     /* F64 matrix (gfx942 CDNA3) */
+    AMD_V_MFMA_I32_16X16X32_I8,
+    AMD_V_MFMA_I32_32X32X16_I8,
+
     AMD_V_MFMA_F64_4X4X4_F64,
     AMD_V_MFMA_F64_16X16X4_F64,
 
@@ -313,11 +312,25 @@ typedef enum {
 
 typedef struct {
     uint8_t  kind;       /* mop_kind_t */
-    uint8_t  pad;
+    uint8_t  nreg;       /* registers covered; 0 and 1 both mean one */
     uint16_t reg_num;    /* physical or virtual reg number */
     int32_t  imm;        /* immediate value or special reg ID */
 } moperand_t;            /* 8 bytes */
 
+#define MFMA_NONE 0xFFu
+
+#define RA_MFMA_LO 208u  /* v208..v247: the aligned fragment window */
+#define RA_MFMA_HI 248u
+
+typedef struct {
+    uint16_t op;              /* amd_op_t */
+    uint8_t  wa, wb, wd;      /* A, B and C/D widths in registers */
+    uint8_t  op90a, op942;
+} amd_mfma_t;
+
+const amd_mfma_t *amd_mfsh(uint16_t op);
+uint8_t amd_mfop(const amd_mfma_t *sh, int target);
+
 /* ---- Machine Instruction ---- */
 
 #define MINST_MAX_OPS 6
@@ -363,12 +376,10 @@ typedef struct {
     uint16_t bir_func;         /* BIR func index (for rplan BIR scan) */
     uint8_t  needs_dispatch;   /* 1 if kernel uses blockDim/gridDim (dispatch_ptr) */
     uint8_t  max_dim;          /* highest dim used: 0=x, 1=xy, 2=xyz */
+    uint8_t  uses_mfma;        /* reserve the aligned MFMA fragment window */
     uint32_t launch_bounds_max; /* 0 = unconstrained. >0 = programmer's optimistic thread count */
     uint32_t launch_bounds_min; /* 0 = not set */
 
-    /* Resource plan — stamped by amd_rplan(), read by isel + emit.
-     * Target decisions made once.  No is_cdna() downstream.
-     * Like pre-flight checks: argue with the checklist, not the runway. */
     uint8_t  exec_w;     /* 0=B32 (Wave32), 1=B64 (Wave64) */
     uint8_t  smem_hz;    /* 1=SMEM→SALU hazard, promote to VALU */
     uint8_t  scr_afs;    /* 1=architected flat scratch (no prologue) */
@@ -399,6 +410,7 @@ typedef struct {
 #define AMD_MAX_MINSTS    (1 << 18)   /* 256K machine instructions */
 #define AMD_MAX_MBLOCKS   (1 << 16)
 #define AMD_MAX_MFUNCS    (1 << 12)
+#define AMD_MAX_ELFK      64
 #define AMD_MAX_VREGS     (1 << 16)
 #define AMD_CODE_SIZE     (4*1024*1024)
 #define AMD_ASM_SIZE      (4*1024*1024)
@@ -445,6 +457,8 @@ typedef struct {
 
     char        asm_buf[AMD_ASM_SIZE];
     uint32_t    asm_len;
+
+    int         enc_err;   /* encoder refused something; fail the compile */
 } amd_module_t;
 
 /* ---- Encoding Table Entry ---- */

src/amdgpu/emit.c

@@ -688,7 +688,8 @@ static void ra_lin(amd_module_t *A, uint32_t mf_idx)
      * for spill relays.  GFX942's AccVGPRs are MFMA-only — the
      * 8-bit encoding fields in VOP/FLAT literally can't see them.
      * We tried.  The hardware was unimpressed. */
-    for (uint16_t r = RA_VGPR_CEIL; r-- > 0; )
+    uint16_t vceil = F->uses_mfma ? RA_MFMA_LO : RA_VGPR_CEIL;
+    for (uint16_t r = vceil; r-- > 0; )
         RA.vgpr_free[RA.num_vgpr_free++] = (uint8_t)r;
 
     ra_nspill = 0;
@@ -791,6 +792,7 @@ static void ra_lin(amd_module_t *A, uint32_t mf_idx)
     if (F->is_kernel && F->num_sgprs < F->first_alloc_sgpr)
         F->num_sgprs = F->first_alloc_sgpr;
     F->num_vgprs = RA.max_vgpr;
+    if (F->uses_mfma && F->num_vgprs < RA_MFMA_HI) F->num_vgprs = RA_MFMA_HI;
 
     /* Match old regalloc_function order exactly:
      * 1. min SGPR/VGPR + launch_bounds
@@ -1952,6 +1954,7 @@ static void ra_gc(amd_module_t *A, uint32_t mf_idx)
 
         F->num_sgprs = max_sgpr;
         F->num_vgprs = max_vgpr;
+        if (F->uses_mfma && F->num_vgprs < RA_MFMA_HI) F->num_vgprs = RA_MFMA_HI;
         gc_success = 1;
         break; /* success, no spills */
     }
@@ -1999,6 +2002,11 @@ static void asm_append(amd_module_t *A, const char *fmt, ...)
 
 static void print_operand(amd_module_t *A, const moperand_t *op)
 {
+    if (op->nreg > 1 && (op->kind == MOP_SGPR || op->kind == MOP_VGPR)) {
+        asm_append(A, "%c[%u:%u]", op->kind == MOP_SGPR ? 's' : 'v',
+                   op->reg_num, op->reg_num + op->nreg - 1);
+        return;
+    }
     switch (op->kind) {
     case MOP_SGPR:
         asm_append(A, "s%u", op->reg_num);
@@ -2464,13 +2472,17 @@ int amdgpu_emit_elf(amd_module_t *A, const char *path)
      * like a well-organised criminal enterprise. */
     static uint8_t rodata[16384];   /* up to ~256 KDs with alignment */
     uint32_t rodata_len = 0;
-    static uint32_t rodata_kd_off[64]; /* KD offset within .rodata */
-    static uint32_t code_offsets[64];  /* code offset within .text */
+    static uint32_t rodata_kd_off[AMD_MAX_ELFK]; /* KD offset within .rodata */
+    static uint32_t code_offsets[AMD_MAX_ELFK];  /* code offset within .text */
     uint32_t num_kernels = 0;
 
     for (uint32_t fi = 0; fi < A->num_mfuncs; fi++) {
         if (!A->mfuncs[fi].is_kernel) continue;
-        if (num_kernels >= 64) break;
+        if (num_kernels >= AMD_MAX_ELFK) {
+            fprintf(stderr, "amdgpu: more kernels than one object holds "
+                            "(max %u)\n", AMD_MAX_ELFK);
+            return BC_ERR_AMDGPU;
+        }
 
         mfunc_t *F = &A->mfuncs[fi];
 
@@ -2718,8 +2730,7 @@ int amdgpu_emit_elf(amd_module_t *A, const char *path)
                     uint32_t pt_idx = A->bir->type_fields[ft->count + pi];
                     const bir_type_t *pt = &A->bir->types[pt_idx];
                     is_ptr = (pt->kind == BIR_TYPE_PTR);
-                    if (!is_ptr)
-                        arg_sz = (pt->width > 0) ? (uint32_t)(pt->width / 8) : 4;
+                    if (!is_ptr) arg_sz = bir_bsz(A->bir, pt_idx, 8);
                 }
                 if (is_ptr) {
                     mp_fixmap(mp_buf, &mp_pos, 4);
@@ -2831,7 +2842,8 @@ int amdgpu_emit_elf(amd_module_t *A, const char *path)
     static uint32_t sk_name[64], sf_name[64]; /* .strtab offsets */
 
     ki = 0;
-    for (uint32_t fi = 0; fi < A->num_mfuncs && ki < num_kernels && ki < 64; fi++) {
+    for (uint32_t fi = 0; fi < A->num_mfuncs && ki < num_kernels
+                          && ki < AMD_MAX_ELFK; fi++) {
         if (!A->mfuncs[fi].is_kernel) continue;
         const char *name = A->bir->strings + A->mfuncs[fi].name;
         char kd[256];
@@ -2908,7 +2920,8 @@ int amdgpu_emit_elf(amd_module_t *A, const char *path)
     uint32_t si = 1;
 
     ki = 0;
-    for (uint32_t fi = 0; fi < A->num_mfuncs && ki < num_kernels && ki < 64; fi++) {
+    for (uint32_t fi = 0; fi < A->num_mfuncs && ki < num_kernels
+                          && ki < AMD_MAX_ELFK; fi++) {
         if (!A->mfuncs[fi].is_kernel) continue;
 
         /* .kd descriptor (STT_OBJECT) in .rodata (section 5) */

src/amdgpu/enc_tab.c

@@ -597,6 +597,8 @@ const amd_enc_entry_t amd_enc_table_gfx9[AMD_OP_COUNT] = {
     [AMD_V_MFMA_F32_32X32X16_FP8_BF8] = { AMD_FMT_VOP3P_MAI, 0x76, "v_mfma_f32_32x32x16_fp8_bf8" },
     [AMD_V_MFMA_F32_32X32X16_BF8_FP8] = { AMD_FMT_VOP3P_MAI, 0x75, "v_mfma_f32_32x32x16_bf8_fp8" },
     [AMD_V_MFMA_F32_32X32X16_BF8_BF8] = { AMD_FMT_VOP3P_MAI, 0x74, "v_mfma_f32_32x32x16_bf8_bf8" },
+    [AMD_V_MFMA_I32_16X16X32_I8]      = { AMD_FMT_VOP3P_MAI, 0x57, "v_mfma_i32_16x16x32_i8" },
+    [AMD_V_MFMA_I32_32X32X16_I8]      = { AMD_FMT_VOP3P_MAI, 0x56, "v_mfma_i32_32x32x16_i8" },
     /* F64 — for when your matrix really needs 52 bits of mantissa */
     [AMD_V_MFMA_F64_4X4X4_F64]        = { AMD_FMT_VOP3P_MAI, 0x6F, "v_mfma_f64_4x4x4f64"         },
     [AMD_V_MFMA_F64_16X16X4_F64]      = { AMD_FMT_VOP3P_MAI, 0x6E, "v_mfma_f64_16x16x4f64"       },
@@ -606,3 +608,45 @@ const amd_enc_entry_t amd_enc_table_gfx9[AMD_OP_COUNT] = {
     [AMD_PSEUDO_COPY]            = { AMD_FMT_PSEUDO, 0, "PSEUDO_COPY" },
     [AMD_PSEUDO_DEF]             = { AMD_FMT_PSEUDO, 0, "PSEUDO_DEF"  },
 };
+
+static const amd_mfma_t mfma_shapes[] = {
+    { AMD_V_MFMA_F32_4X4X4_F16,           2, 2,  4, 0x4A, 0x4A },
+    { AMD_V_MFMA_F32_16X16X16_F16,        2, 2,  4, 0x4D, 0x4D },
+    { AMD_V_MFMA_F32_32X32X8_F16,         2, 2, 16, 0x4C, 0x4C },
+    { AMD_V_MFMA_F32_4X4X4_BF16_1K,       2, 2,  4, 0x65, 0x5F },
+    { AMD_V_MFMA_F32_16X16X16_BF16_1K,    2, 2,  4, 0x67, 0x61 },
+    { AMD_V_MFMA_F32_32X32X8_BF16_1K,     2, 2, 16, 0x66, 0x60 },
+    { AMD_V_MFMA_F32_4X4X1_F32,           1, 1,  4, 0x42, 0x42 },
+    { AMD_V_MFMA_F32_16X16X4_F32,         1, 1,  4, 0x45, 0x45 },
+    { AMD_V_MFMA_F32_32X32X2_F32,         1, 1, 16, 0x44, 0x44 },
+    { AMD_V_MFMA_I32_4X4X4_I8,            1, 1,  4, 0x52, 0x52 },
+    { AMD_V_MFMA_I32_16X16X16_I8,         1, 1,  4, 0x55, MFMA_NONE },
+    { AMD_V_MFMA_I32_32X32X8_I8,          1, 1, 16, 0x54, MFMA_NONE },
+    { AMD_V_MFMA_I32_16X16X32_I8,         2, 2,  4, MFMA_NONE, 0x57 },
+    { AMD_V_MFMA_I32_32X32X16_I8,         2, 2, 16, MFMA_NONE, 0x56 },
+    { AMD_V_MFMA_F32_16X16X32_FP8_FP8,    2, 2,  4, MFMA_NONE, 0x73 },
+    { AMD_V_MFMA_F32_16X16X32_FP8_BF8,    2, 2,  4, MFMA_NONE, 0x72 },
+    { AMD_V_MFMA_F32_16X16X32_BF8_FP8,    2, 2,  4, MFMA_NONE, 0x71 },
+    { AMD_V_MFMA_F32_16X16X32_BF8_BF8,    2, 2,  4, MFMA_NONE, 0x70 },
+    { AMD_V_MFMA_F32_32X32X16_FP8_FP8,    2, 2, 16, MFMA_NONE, 0x77 },
+    { AMD_V_MFMA_F32_32X32X16_FP8_BF8,    2, 2, 16, MFMA_NONE, 0x76 },
+    { AMD_V_MFMA_F32_32X32X16_BF8_FP8,    2, 2, 16, MFMA_NONE, 0x75 },
+    { AMD_V_MFMA_F32_32X32X16_BF8_BF8,    2, 2, 16, MFMA_NONE, 0x74 },
+    { AMD_V_MFMA_F64_4X4X4_F64,           2, 2,  2, 0x6F, 0x6F },
+    { AMD_V_MFMA_F64_16X16X4_F64,         2, 2,  8, 0x6E, 0x6E },
+};
+
+const amd_mfma_t *amd_mfsh(uint16_t op)
+{
+    for (unsigned i = 0; i < sizeof mfma_shapes / sizeof mfma_shapes[0]; i++)
+        if (mfma_shapes[i].op == op) return &mfma_shapes[i];
+    return NULL;
+}
+
+uint8_t amd_mfop(const amd_mfma_t *sh, int target)
+{
+    if (!sh) return MFMA_NONE;
+    if (target == AMD_TARGET_GFX90A) return sh->op90a;
+    if (target == AMD_TARGET_GFX942) return sh->op942;
+    return MFMA_NONE;
+}

src/amdgpu/encode.c

@@ -456,6 +456,12 @@ static void encode_flat_global(amd_module_t *A, const minst_t *mi, uint16_t hw_o
 
 static void encode_vop3p_mai(amd_module_t *A, const minst_t *mi, uint16_t hw_op)
 {
+    const amd_mfma_t *sh = amd_mfsh(mi->op);
+    if (sh) {
+        uint8_t t = amd_mfop(sh, A->target);
+        if (t == MFMA_NONE) { A->enc_err = 1; return; }
+        hw_op = t;
+    }
     /* VOP3P-MAI: 64-bit encoding for MFMA matrix instructions.
        DW0: [31:23]=0x1A7(prefix) [22:16]=OP(7b) [15]=ACC_CD [14:11]=ABID
             [10:8]=CBSZ [7:0]=VDST