Add compute-backed geometry shader replay

Capture vertex outputs, run geometry shaders as compute, and rasterize compacted primitives through a general NVN pipeline. Rebuild the pinned translator and AI shader chain, add geometry regression coverage, and document supported behavior and remaining limits.

dm/games/010089B00D09C000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-fc6c0368c7b9 2026-10-23T14:27:31+08:00
+# chain b-3b3046f76f22 2026-10-23T16:15:22+08:00
 0ed22b53c3db7a71
 623b686baab08b30
 738bc739393f913b

dm/games/010089B00D09C000/cb_extent.tsv

@@ -1,4 +1,4 @@
-# chain b-fc6c0368c7b9 2026-10-23T14:27:34+08:00
+# chain b-3b3046f76f22 2026-10-23T16:15:23+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 00005528b1b0e2cd	0:512
 0003d30f7a6c71b2	0:320
@@ -1391,6 +1391,7 @@
 1b1799926c0773c2	0:656
 1b1d17dd5f090ad7	0:192
 1b2ce675610985a3	0:112
+1b378c5c014e3a32	0:32
 1b4020b5e1aca7b2	0:528
 1b422c152ca6a320	0:80
 1b4d0cf3e4e1aba3	0:144
@@ -1778,6 +1779,7 @@
 23242d9bb538257a	0:144
 233afc6c0d721d35	0:96
 233d40c1c1d1ff28	0:576
+2341c1be4fcbf690	0:16
 234203fd460a0ad1	0:208
 2346eac64d1c80ec	0:256
 2349946407a153c4	0:144
@@ -10795,6 +10797,7 @@ d57139953a05361b	0:160
 d571931632ab0f18	0:224
 d57468710502ce13	0:176
 d5748fd1206e9ca8	0:144
+d57cdb168fdf6927	0:32
 d57ff4397b96f525	0:256
 d5836c4e977eb06b	0:208
 d5864edbaaeb0ef3	0:80

dm/games/010089B00D09C000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010089B00D09C000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "fc6c0368c7b9"
+#define DMSH_CHAIN "3b3046f76f22"

dm/gfx/nvn/hle_nvn.c

@@ -2222,7 +2222,7 @@ NVN(nvnCommandBufferSetViewportSwizzles, void, NVNcommandBuffer *c, int first, i
 NVN(nvnCommandBufferSetAlphaRef, void, NVNcommandBuffer *c, float r) { (void)c; ONCE_LOG("SetAlphaRef %.3f: the host has no alpha test (ColorStateSetAlphaTest is only recorded)", r); }
 NVN(nvnCommandBufferSetLineWidth, void, NVNcommandBuffer *c, float w) { (void)c; if (w != 1.0f) ONCE_LOG("SetLineWidth %.2f: Metal has only 1-pixel lines", w); }
 NVN(nvnCommandBufferSetPointSize, void, NVNcommandBuffer *c, float s) { (void)c; (void)s; ONCE_LOG("SetPointSize: the point size comes from the shader"); }
-NVN(nvnCommandBufferSetPrimitiveRestart, void, NVNcommandBuffer *c, uint32_t en, int idx) { (void)c; if ((en & 0xff) && idx != -1 && idx != 0xffff) ONCE_LOG("[RED] SetPrimitiveRestart index %#x: Metal only accepts all ones", idx); }
+NVN(nvnCommandBufferSetPrimitiveRestart, void, NVNcommandBuffer *c, uint32_t en, int idx) { *REC(c, OP_PRIM_RESTART, p_prim_restart) = (p_prim_restart){(en & 0xff) != 0, (uint32_t)idx}; }
 // Rasterizer discard: draws run their vertex stage only (with transform feedback, the capture); recorded, and at replay such a draw rasterizes nothing (nvn_mtl.m)
 NVN(nvnCommandBufferSetRasterizerDiscard, void, NVNcommandBuffer *c, uint32_t en) { REC(c, OP_RAST_DISCARD, p_rast_discard)->on = (en & 0xff) != 0; }
 NVN(nvnCommandBufferSetSampleMask, void, NVNcommandBuffer *c, int m) { (void)c; if ((m & 0xffff) != 0xffff) ONCE_LOG("SetSampleMask %#x: ignored by the host for now", m); }

dm/gfx/nvn/nvn_geometry.h

@@ -0,0 +1,88 @@
+// Primitive assembly for geometry compute. Values are capture-record indices;
+// restart separates runs before strips/fans are assembled. No game-specific state.
+#pragma once
+#include <stdint.h>
+#include <stddef.h>
+
+static int nvn_geo_topology(int prim) {
+    switch (prim) {
+        case 0: return 0;
+        case 1: case 2: case 3: return 1;
+        case 10: case 11: return 2;
+        case 4: case 5: case 6: case 7: return 3;
+        case 12: return 4;
+        default: return -1;
+    }
+}
+static int nvn_geo_width(int topology) { static const int width[] = {1, 2, 4, 3, 6}; return topology >= 0 && topology < 5 ? width[topology] : 0; }
+
+// Returns words required/written. SIZE_MAX means unsupported topology or overflow.
+// dst == NULL counts only. A short destination is never partially written.
+static size_t nvn_geo_assemble(int prim, const uint32_t *src, size_t count, int restart, uint32_t restart_index, uint32_t *dst, size_t capacity) {
+    int width = nvn_geo_width(nvn_geo_topology(prim));
+    if (!width) return SIZE_MAX;
+    if (dst) {
+        size_t needed = nvn_geo_assemble(prim, src, count, restart, restart_index, NULL, 0);
+        if (needed == SIZE_MAX || needed > capacity) return SIZE_MAX;
+    }
+    size_t words = 0;
+    for (size_t begin = 0; begin < count;) {
+        size_t end = begin;
+        while (end < count && !(restart && src[end] == restart_index)) end++;
+        size_t n = end - begin, np = 0;
+        switch (prim) {
+            case 0: np = n; break;
+            case 1: np = n / 2; break;
+            case 2: np = n > 1 ? n : 0; break;
+            case 3: np = n > 1 ? n - 1 : 0; break;
+            case 4: np = n / 3; break;
+            case 5: case 6: np = n > 2 ? n - 2 : 0; break;
+            case 7: np = n / 4 * 2; break;
+            case 10: np = n / 4; break;
+            case 11: np = n > 3 ? n - 3 : 0; break;
+            case 12: np = n / 6; break;
+        }
+        if (np > (SIZE_MAX - words) / (size_t)width) return SIZE_MAX;
+        if (dst && words + np * width > capacity) return SIZE_MAX;
+        for (size_t i = 0; dst && i < np; i++) {
+            size_t k[6] = {0};
+            switch (prim) {
+                case 0: k[0] = i; break;
+                case 1: k[0] = 2*i; k[1] = 2*i+1; break;
+                case 2: k[0] = i; k[1] = (i+1)%n; break;
+                case 3: k[0] = i; k[1] = i+1; break;
+                case 4: for (int c=0;c<3;c++) k[c]=3*i+c; break;
+                case 5: k[0]=i+(i&1); k[1]=i+1-(i&1); k[2]=i+2; break;
+                case 6: k[0]=0; k[1]=i+1; k[2]=i+2; break;
+                case 7: k[0]=4*(i/2); k[1]=k[0]+1+(i&1); k[2]=k[1]+1; break;
+                case 10: for (int c=0;c<4;c++) k[c]=4*i+c; break;
+                case 11: for (int c=0;c<4;c++) k[c]=i+c; break;
+                case 12: for (int c=0;c<6;c++) k[c]=6*i+c; break;
+            }
+            for (int c=0;c<width;c++) dst[words+i*width+c]=src[begin+k[c]];
+        }
+        words += np * width;
+        begin = end < count ? end + 1 : end;
+    }
+    return words;
+}
+
+// One ordered compaction per draw. Output is a point/line/triangle list, so it
+// does not depend on Metal primitive-restart behaviour. Preserve triangle-strip
+// winding and the last vertex. Unwritten GS index slots are UINT_MAX.
+static const char k_nvn_geo_pack_msl[] =
+"#include <metal_stdlib>\nusing namespace metal;\n"
+"struct geo_pack { device const uint *src; device uint *dst; device uint *args; uint groups, stride, topology, capacity; };\n"
+"kernel void geo_pack_indices(constant geo_pack& j [[buffer(0)]]) {\n"
+" uint outn=0;\n"
+" for(uint g=0;g<j.groups;g++) { uint a=0,b=0,n=0;\n"
+"  for(uint i=0;i<j.stride;i++) { uint v=j.src[g*j.stride+i];\n"
+"   if(v==0xffffffffu) {n=0;continue;}\n"
+"   if(j.topology==1) {if(outn<j.capacity) j.dst[outn++]=v;}\n"
+"   else if(j.topology==6) {if(n && outn+2<=j.capacity) {j.dst[outn++]=b;j.dst[outn++]=v;}}\n"
+"   else if(n>=2 && outn+3<=j.capacity) {j.dst[outn++]=(n&1)?b:a;j.dst[outn++]=(n&1)?a:b;j.dst[outn++]=v;}\n"
+"   a=b;b=v;n++;\n"
+"  }\n"
+" }\n"
+" j.args[0]=outn;j.args[1]=1;j.args[2]=0;j.args[3]=0;j.args[4]=0;\n"
+"}\n";

dm/gfx/nvn/nvn_geometry.inc

@@ -0,0 +1,158 @@
+// General geometry replay: native VS capture -> GS compute -> ordered index
+// compaction -> ordinary rendering. Included after the resource/pipeline helpers.
+static NvProg *g_geo_program;
+static uint64_t g_geo_indices, g_geo_index_bytes, g_geo_args;
+static void draw(const p_draw *dw, const p_draw_idx *di);
+static void lc_land(void);
+
+static int geo_cmp_index(const void *a, const void *b) {
+    uint32_t x = *(const uint32_t *)a, y = *(const uint32_t *)b;
+    return (x > y) - (x < y);
+}
+static id<MTLRenderPipelineState> geo_capture_pso(NvShader *vs) {
+    pso_key k; static const uint32_t cf[8]; pso_key_fill(&k, vs, nil, cf, 0, 0, 1, 0); k.fsfn = GEO_MARK;
+    id<MTLRenderPipelineState> ps = hm_get_id(&g_pso, &k);
+    if (ps || hm_get_id(&g_pso_fail, &k)) return ps;
+    NSError *e = nil;
+    ps = [g_dev newRenderPipelineStateWithDescriptor:pso_desc(&k, vs, nil, nil, 0) error:&e];
+    if (ps) hm_put_id(&g_pso, &k, ps);
+    else { hm_put_id(&g_pso_fail, &k, NSNull.null); MLOG("[RED] geometry vertex capture %s: %s", vs->key, e.localizedDescription.UTF8String); }
+    return ps;
+}
+static void geo_draw(NvProg *pr, const p_draw *dw0, const p_draw_idx *di0) {
+    NvShader *vs = pr->sh[0], *base = pr->sh[2];
+    if (!vs || !vs->geo_fn || !base || !base->geo_vertices || pr->sh[3] || pr->sh[4] || S.tf_on) {
+        MLOG("[RED] geometry draw cannot capture this pipeline (vertex %s, geometry %s, tessellation %d, feedback %d)", vs ? vs->key : "-", base ? base->key : "-", !!(pr->sh[3] || pr->sh[4]), S.tf_on); return;
+    }
+    p_draw dw = dw0 ? *dw0 : (p_draw){0}; p_draw_idx di = di0 ? *di0 : (p_draw_idx){0};
+    if (di0 && (di.itype < 0 || di.itype > 2)) { MLOG("[RED] geometry index type %d is invalid", di.itype); return; }
+    uint64_t ind = dw0 ? dw.ind : di.ind;
+    // Indirect records and GPU-written index data need a coherent CPU snapshot
+    // for primitive assembly. This is deliberately a correctness path; no such
+    // wait is added to ordinary VS/FS draws.
+    int wait = ind != 0 || di0 != NULL;  // CPU assembly must also see writes from earlier submissions.
+    for (int i = 0; i < 16; i++) if (S.vbuf[i].size && pend_hits(S.vbuf[i].address, S.vbuf[i].size)) wait = 1;
+    for (int i = 0; i < vs->rf.nv; i++) if (vconv_wanted(&vs->rf, i, NULL, 0)) wait = 1;
+    if (defer_has_active()) defer_fallback();
+    lc_land(); rp_flush(); rp_end(); ubd_land();
+    if (wait) { mdi_wait(); if (g_mdi_ev.signaledValue < g_mdi_v) return; memset(g_vconv, 0, sizeof g_vconv); }
+    if (ind) {
+        uint32_t a[5] = {0}; size_t n = dw0 ? 16 : 20;
+        if (!nvn_host_readable(ind, n)) { MLOG("[RED] geometry indirect arguments are not readable"); return; }
+        memcpy(a, (void *)(uintptr_t)ind, n);
+        if (!a[0] || !a[1]) return;
+        if (a[0] > INT32_MAX || a[1] > INT32_MAX) { MLOG("[RED] geometry indirect draw exceeds the host count range"); return; }
+        if (dw0) { dw.count = a[0]; dw.instances = a[1]; dw.first = a[2]; dw.base_instance = a[3]; dw.ind = dw.mdi_cnt = 0; }
+        else { di.count = a[0]; di.instances = a[1]; di.ibuf += (uint64_t)a[2] << di.itype; di.base_vertex = (int32_t)a[3]; di.base_instance = a[4]; di.ind = di.mdi_cnt = 0; }
+    }
+    int prim = dw0 ? dw.prim : di.prim, topology = nvn_geo_topology(prim);
+    int count = dw0 ? dw.count : di.count, instances = dw0 ? dw.instances : di.instances;
+    uint32_t first_instance = (uint32_t)(dw0 ? dw.base_instance : di.base_instance);
+    g_draw_binst = first_instance;
+    if (count <= 0 || instances < 0) return;
+    if (!instances) instances = 1;
+    if (topology < 0 || count > (1 << 22)) { MLOG("[RED] geometry primitive %d or count %d exceeds the assembly path", prim, count); return; }
+    NSString *key = [[NSString stringWithUTF8String:base->key] substringToIndex:16];
+    NvShader *gs = shader_get([key stringByAppendingFormat:@"_bg%d", topology], 2);
+    NvShader *reader = shader_get([key stringByAppendingString:@"_bgv"], 0);
+    if (!gs || !reader) return;
+    id<MTLRenderPipelineState> capture = geo_capture_pso(vs); if (!capture) return;
+    if (!gs->geo_pso) { NSError *e = nil; gs->geo_pso = [g_dev newComputePipelineStateWithFunction:gs->fn error:&e];
+        if (!gs->geo_pso) { MLOG("[RED] geometry compute %s: %s", gs->key, e.localizedDescription.UTF8String); return; } }
+    if (gs->geo_pso.maxTotalThreadsPerThreadgroup < 32) { MLOG("[RED] geometry kernel cannot run a 32-thread group"); return; }
+    static id<MTLComputePipelineState> pack;
+    if (!pack) { NSError *e = nil; id<MTLLibrary> lib = [g_dev newLibraryWithSource:@(k_nvn_geo_pack_msl) options:nil error:&e];
+        if (lib) pack = [g_dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"geo_pack_indices"] error:&e];
+        if (!pack) { MLOG("[RED] geometry index compaction: %s", e.localizedDescription.UTF8String); return; } }
+
+    uint32_t *order = malloc((size_t)count * 4), *unique = malloc((size_t)count * 4);
+    if (!order || !unique) { free(order); free(unique); MLOG("[RED] geometry index allocation failed"); return; }
+    uint32_t lo = UINT32_MAX, hi = 0; size_t nu = 0;
+    int restart = di0 && S.restart.on;
+    int isz = di0 ? 1 << di.itype : 0;
+    uint32_t restart_index = S.restart.index & (isz == 1 ? UINT8_MAX : isz == 2 ? UINT16_MAX : UINT32_MAX);
+    if (di0 && (di.itype < 0 || di.itype > 2 || !nvn_host_readable(di.ibuf, (uint64_t)count * isz))) {
+        free(order); free(unique); MLOG("[RED] geometry index buffer is not readable"); return;
+    }
+    for (int i = 0; i < count; i++) {
+        uint32_t v = (uint32_t)i;
+        if (di0) { v = 0; memcpy(&v, (const uint8_t *)(uintptr_t)di.ibuf + (size_t)i * isz, isz); }
+        order[i] = v;
+        if (restart && v == restart_index) continue;
+        lo = MIN(lo, v); hi = MAX(hi, v); unique[nu++] = v;
+    }
+    if (!nu) { free(order); free(unique); return; }
+    int64_t first = di0 ? (int64_t)lo + di.base_vertex : dw.first;
+    uint64_t span = (uint64_t)hi - lo + 1;
+    if (first < 0 || first + span > (uint64_t)UINT32_MAX + 1) { free(order); free(unique); MLOG("[RED] geometry vertex range is outside 32-bit vertex IDs"); return; }
+    // Normalize only after assembly, so a restart value cannot collide with a
+    // normalized vertex number. Indexed captures use a sorted unique index list.
+    size_t words = nvn_geo_assemble(prim, order, count, restart, restart_index, NULL, 0);
+    int width = nvn_geo_width(topology);
+    uint64_t np = words == SIZE_MAX ? 0 : words / width;
+    if (!np) { free(order); free(unique); return; }
+    uint64_t groups = np * (uint64_t)instances * gs->geo_invocations;
+    if (groups > UINT32_MAX || span * (uint64_t)instances > (256ull << 20) / (144 * 4) ||
+        (uint64_t)first_instance + instances > (uint64_t)UINT32_MAX + 1) {
+        free(order); free(unique); MLOG("[RED] geometry draw exceeds the capture or invocation bound"); return;
+    }
+    uint64_t input_size = span * instances * 144 * 4;
+    uint64_t output_size = groups * gs->geo_vertices * gs->geo_stride * 4;
+    uint64_t raw_size = groups * gs->geo_index_stride * 4;
+    uint64_t packed_size = groups * gs->geo_vertices * (gs->geo_output == 7 ? 3 : gs->geo_output == 6 ? 2 : 1) * 4;
+    uint64_t total = input_size + output_size + raw_size + packed_size + words * 4 + nu * 4 + 4096;
+    if (groups > UINT32_MAX || total > (256ull << 20) || np > UINT32_MAX) {
+        free(order); free(unique); MLOG("[RED] geometry draw needs %llu MiB of scratch, above the 256 MiB bound", (unsigned long long)(total >> 20)); return;
+    }
+    id<MTLBuffer> scratch; size_t off = arena_alloc((size_t)total, &scratch);
+    uint8_t *cpu = (uint8_t *)scratch.contents + off; uint64_t gpu = scratch.gpuAddress + off;
+    size_t used = 0;
+    #define GEO_ALLOC(N, G, P) do { used = (used + 255) & ~(size_t)255; (G) = gpu + used; (P) = cpu + used; used += (size_t)(N); } while (0)
+    uint64_t input, output, raw, packed, map, capture_ix, info, capture_info, args, job;
+    void *pinput, *poutput, *praw, *ppacked, *pmap, *pcapture_ix, *pinfo, *pcapture_info, *pargs, *pjob;
+    GEO_ALLOC(input_size, input, pinput); GEO_ALLOC(output_size, output, poutput); GEO_ALLOC(raw_size, raw, praw); GEO_ALLOC(packed_size, packed, ppacked);
+    GEO_ALLOC(words * 4, map, pmap); GEO_ALLOC(nu * 4, capture_ix, pcapture_ix); GEO_ALLOC(32, info, pinfo); GEO_ALLOC(32, capture_info, pcapture_info);
+    GEO_ALLOC(20, args, pargs); GEO_ALLOC(40, job, pjob);
+    #undef GEO_ALLOC
+    (void)pinput; (void)poutput; (void)praw; (void)ppacked;
+    if (used > total) MDIE("geometry scratch size accounting");
+    nvn_geo_assemble(prim, order, count, restart, restart_index, pmap, words);
+    for (size_t i = 0; i < words; i++) ((uint32_t *)pmap)[i] -= lo;
+    qsort(unique, nu, 4, geo_cmp_index);
+    size_t unique_count = 0;
+    for (size_t i = 0; i < nu; i++) if (!i || unique[i] != unique[i - 1]) ((uint32_t *)pcapture_ix)[unique_count++] = unique[i];
+    free(order); free(unique);
+    uint32_t vi[8] = {(uint32_t)span, (uint32_t)instances, 0, 0, (uint32_t)np, 0, 0, 0}; memcpy(pinfo, vi, 32);
+    uint64_t ci[2] = {input, 0}; memcpy(pcapture_info, ci, 16);
+    uint32_t cc[4] = {(uint32_t)first, (uint32_t)span, first_instance, (uint32_t)instances}; memcpy((uint8_t *)pcapture_info + 16, cc, 16);
+    memset(pargs, 0, 20);
+    struct { uint64_t src, dst, args; uint32_t groups, stride, topology, capacity; } pj = {raw, packed, args, (uint32_t)groups, gs->geo_index_stride, gs->geo_output, (uint32_t)(packed_size / 4)};
+    memcpy(pjob, &pj, sizeof pj);
+    if (g_blit) { enc_b_end(g_blit); g_blit = nil; }
+    MTL4RenderPassDescriptor *rd = [MTL4RenderPassDescriptor new]; rd.renderTargetWidth = rd.renderTargetHeight = rd.defaultRasterSampleCount = 1;
+    id<MTL4RenderCommandEncoder> re = enc_r_begin(rd); enc_r_pso(re, capture); bind_vbufs(re);
+    p_draw vertex_range = {.first = (int32_t)first, .count = (int32_t)span, .instances = instances, .base_instance = (int32_t)first_instance};
+    vconv_bind(re, vs, 1, &vertex_range, NULL);
+    args_dirty_all(); encode_args(re, vs, 0, 1); enc_r_vbuf(re, 13, capture_info);
+    if (di0) enc_r_draw_indexed(re, MTLPrimitiveTypePoint, unique_count, MTLIndexTypeUInt32, capture_ix, nu * 4, instances, di.base_vertex, first_instance);
+    else enc_r_draw_instanced(re, MTLPrimitiveTypePoint, dw.first, dw.count, instances, first_instance);
+    enc_r_end(re);
+    g_geo_bind[0] = input; g_geo_bind[1] = output; g_geo_bind[2] = raw; g_geo_bind[3] = info; g_geo_bind[4] = map;
+    id<MTL4ComputeCommandEncoder> ce = enc_c_begin(NULL);
+    args_dirty_all(); encode_args(ce, gs, 2, 0); enc_c_pso(ce, gs->geo_pso);
+    enc_c_dispatch(ce, MTLSizeMake((np + 31) / 32, instances, gs->geo_invocations), MTLSizeMake(32, 1, 1)); enc_c_end(ce);
+    ce = enc_c_begin(NULL); enc_c_pso(ce, pack); enc_c_args(ce, job, 0);
+    enc_c_dispatch(ce, MTLSizeMake(1, 1, 1), MTLSizeMake(1, 1, 1)); enc_c_end(ce);
+    args_dirty_all(); g_enc_pso = nil; g_enc_dss = nil; g_ed.ok = 0;
+    swrite_note(pr); udc_note_draw(pr); g_wr_epoch++; g_mdx_epoch++;
+    if (S.rast_discard) return;
+    NvProg *raster = pr->geo_raster;
+    if (!raster || raster->sh[0] != reader || raster->sh[1] != pr->sh[1]) {
+        raster = [NvProg new]; raster->sh[0] = reader; raster->sh[1] = pr->sh[1]; pr->geo_raster = raster;
+    }
+    g_geo_program = raster; g_geo_indices = packed; g_geo_index_bytes = packed_size; g_geo_args = args;
+    p_draw synthetic = {.prim = gs->geo_output == 1 ? 0 : gs->geo_output == 6 ? 1 : 4, .count = 3, .instances = 1, .base_instance = (int32_t)first_instance};
+    draw(&synthetic, NULL);
+    g_geo_program = nil; args_dirty_all();
+    static int once; if (!once++) MLOG("geometry compute: vertex %s, geometry %s, %llu input primitives, %d instances", vs->key, gs->key, (unsigned long long)np, instances);
+}

dm/gfx/nvn/nvn_mtl.m

@@ -2796,8 +2796,10 @@ static void shmap_load(void) {
     warm_start();
 }
 
-@interface NvShader : NSObject { @public id<MTLFunction> fn, tf_fn, fn_u;   // fn_u: the integer-output entry point (dmsh_build.py fs_uint), for pipelines whose colour attachments are all unsigned integer (pso_desc)
+@interface NvShader : NSObject { @public id<MTLFunction> fn, tf_fn, fn_u, geo_fn;   // fn_u: the integer-output entry point (dmsh_build.py fs_uint), for pipelines whose colour attachments are all unsigned integer (pso_desc)
     uint16_t tf_stride[4]; uint32_t cb_ext[16], sb_ro; nvn_refl rf; id<MTLArgumentEncoder> ae; NSUInteger arg_len; int stage, range_ro, afr, layer_out; const char *builtin; uint64_t logged; char key[24]; id<MTLBuffer> c1buf; uint64_t c1_gpu; size_t c1off;
+    uint32_t geo_vertices, geo_invocations, geo_output, geo_stride, geo_index_stride;
+    id<MTLComputePipelineState> geo_pso;
     int gs_loc, lay_loc; char gs_comp, lay_comp;   // a slice-writing geometry shader (no function: the layer comes from input gs_loc.gs_comp) / a vertex shader with layer entry points writing output lay_loc.lay_comp (dmsh_build.py gs_slice, vs_layer_entry); -1 = none
     // The argument buffer encoded last time
     id<MTLBuffer> ab_buf; uint64_t ab_gpu; size_t ab_off; uint64_t ab_pass; uint32_t ab_binst;   // ab_pass: serial of the render pass that encoded this argument block - reused only within the same pass (the arena holding a block from another pass / another submission may already have been reclaimed and rewritten)
@@ -2808,7 +2810,8 @@ static void shmap_load(void) {
 - (void)dealloc { free(rf.c1); }
 @end
 #include "nvn_pso_skey.h"
-@interface NvProg : NSObject { @public NvShader *sh[6]; char gs_miss[17]; } @end   // gs_miss: the key of a geometry shader with no library (no Metal counterpart), "" = none (draw_gs_check)
+#include "nvn_geometry.h"
+@interface NvProg : NSObject { @public NvShader *sh[6]; char gs_miss[17]; NvProg *geo_raster; } @end   // geo_raster: GS output reader paired with this program's fragment stage
 @implementation NvProg @end
 // Warm-up entry: built in the background, waiting to be taken by the submit thread. vs / fs are strong references (in pso they are only raw pointers, and the background dictionary is released when the task ends)
 @interface WarmItem : NSObject { @public int kind; NSString *key, *vs_key, *fs_key; NvShader *shader, *vs, *fs; pso_key pso; id<MTLFunction> fsfn; id object; } @end
@@ -2820,7 +2823,7 @@ static void c1_store(NvShader *o) {
     NSString *guardKey = [NSString stringWithUTF8String:o->key];
     if ([guardKey hasSuffix:@"_b"]) guardKey = [guardKey substringToIndex:guardKey.length - 2];
     os_unfair_lock_lock(&g_c1_lock);   // Both the warm-up thread and the submit thread load shaders
-    if (!g_c1_full_valid || [g_c1_full containsObject:guardKey]) {   // Runtime index not proven bounded: exclusive 64 KiB of all zeros (same as the original implementation)
+    if (o->geo_vertices || !g_c1_full_valid || [g_c1_full containsObject:guardKey]) {   // Geometry topology variants have independent control flow; only the base shader is in the guard table.
         o->c1buf = [g_dev newBufferWithLength:C1_DECL_SIZE options:MTLResourceStorageModeShared]; o->c1_gpu = o->c1buf.gpuAddress;
         if (!o->c1buf) { os_unfair_lock_unlock(&g_c1_lock); MDIE("c1 dedicated buffer allocation failed"); }
         res_add(o->c1buf);
@@ -2937,7 +2940,7 @@ static NvShader *shader_load(NSString *pc, int stage, int layer) {
         o->gs_loc = [j[@"gs_slice"][@"loc"] intValue]; o->gs_comp = [[j[@"gs_slice"][@"comp"] description] UTF8String][0];
         return o;
     }
-    if (stage == 2 && !j) return nil;   // any other geometry shader has no library (no Metal counterpart): the caller's red line says so
+    if (stage == 2 && !j) return nil;   // missing geometry translation: the caller's red line says so
     NSString *lp = [NSString stringWithFormat:@"%@/%@.metallib", g_msl_dir, pc];
     id<MTLLibrary> lib = lib_from_file(lp, &e);
     if (!lib || !j) { MLOG("[RED] failed to load shader %s: %s", pc.UTF8String, e ? e.localizedDescription.UTF8String : "no reflection json"); return nil; }
@@ -2957,6 +2960,15 @@ static NvShader *shader_load(NSString *pc, int stage, int layer) {
         if (atomic_load(&g_depth_zero)) { if (j[@"entry_z"]) entry = j[@"entry_z"]; else RED_ONCE(85, "vertex shader %s has no near=0 entry (entry_z): depth stays in the GL convention", pc.UTF8String); }
         if (layer) entry = atomic_load(&g_depth_zero) ? j[@"entry_layer_z"] : j[@"entry_layer"];
     }
+    if (j[@"entry_geometry_capture"]) {
+        NSError *ge = nil;
+        o->geo_fn = [lib newFunctionWithName:j[@"entry_geometry_capture"] constantValues:[MTLFunctionConstantValues new] error:&ge];
+        if (!o->geo_fn) MLOG("[RED] vertex shader %s: geometry capture entry failed: %s", pc.UTF8String, ge.localizedDescription.UTF8String);
+    }
+    NSDictionary *geo = j[@"geometry"];
+    if (j[@"geometry_reader"]) o->geo_output = [j[@"geometry_output"] unsignedIntValue];
+    if (geo) { o->geo_vertices = [geo[@"vertices"] unsignedIntValue]; o->geo_invocations = [geo[@"invocations"] unsignedIntValue];
+        o->geo_output = [geo[@"output"] unsignedIntValue]; o->geo_stride = [geo[@"stride"] unsignedIntValue]; o->geo_index_stride = [geo[@"index_stride"] unsignedIntValue]; }
     o->fn = [lib newFunctionWithName:entry]; o->stage = stage; strlcpy(o->key, pc.UTF8String, sizeof o->key); o->ventry = entry;
     // writes the layer (render_target_array_index, named in the library's metadata): its passes are layered and its pipelines get an input class (pso_desc). A library with layer entry points
     // names it for those alone: its main0 / main0_z write no layer
@@ -3536,17 +3548,19 @@ static inline __attribute__((always_inline)) void *samp_fast(int sid) {
 @interface NvArena : NSObject { @public id<MTLBuffer> buf; uint64_t gpu, fence; size_t used; } @end
 @implementation NvArena @end
 static NSMutableArray<NvArena *> *g_arenas; static NvArena *g_cur_arena; static NSMutableArray<NvArena *> *g_sub_arenas;
-static NvArena *arena_new(void) {
+static NvArena *arena_new(size_t minimum) {
     // The for-in loop variable does not retain the object: take a strong reference first and then remove it from the array, otherwise what is returned after the array releases it is a dangling pointer
-    for (NvArena *it in g_arenas) if (it->fence && it->fence <= g_done) { NvArena *r = it; [g_arenas removeObject:r]; r->fence = 0; r->used = 0; return r; }
+    for (NvArena *it in g_arenas) if (it->fence && it->fence <= g_done && it->buf.length >= minimum) { NvArena *r = it; [g_arenas removeObject:r]; r->fence = 0; r->used = 0; return r; }
-    NvArena *a = [NvArena new]; a->buf = [g_dev newBufferWithLength:4 << 20 options:MTLResourceStorageModeShared]; a->gpu = a->buf.gpuAddress; res_add(a->buf); g_arena_made++; return a;
+    NvArena *a = [NvArena new]; a->buf = [g_dev newBufferWithLength:MAX(minimum, 4u << 20) options:MTLResourceStorageModeShared];
+    if (!a->buf) MDIE("cannot allocate %zu bytes of frame scratch", minimum);
+    a->gpu = a->buf.gpuAddress; res_add(a->buf); g_arena_made++; return a;
 }
 void arena_begin(void) { if (!g_arenas) { g_arenas = [NSMutableArray array]; g_sub_arenas = [NSMutableArray array]; } g_cur_arena = nil; }
 void arena_commit(uint64_t fence) { for (NvArena *a in g_sub_arenas) { a->fence = fence; [g_arenas addObject:a]; } [g_sub_arenas removeAllObjects]; g_cur_arena = nil; }
 static size_t arena_alloc(size_t n, id<MTLBuffer> *buf) {
     n = (n + 255) & ~(size_t)255;
     if (!g_cur_arena || g_cur_arena->used + n > g_cur_arena->buf.length) {
-        pthread_mutex_lock(&g_mu); g_cur_arena = arena_new(); pthread_mutex_unlock(&g_mu);
+        pthread_mutex_lock(&g_mu); g_cur_arena = arena_new(n); pthread_mutex_unlock(&g_mu);
         [g_sub_arenas addObject:g_cur_arena];
     }
     size_t off = g_cur_arena->used; g_cur_arena->used += n; *buf = g_cur_arena->buf; return off;
@@ -3556,6 +3570,7 @@ static size_t arena_alloc(size_t n, id<MTLBuffer> *buf) {
 static struct {
     const prog_t *prog;
     const prog_t *cprog;       // compute program: NVN binds it apart from the graphics stages (see OP_BIND_PROGRAM in state_op)
+    p_prim_restart restart;
     const prog_t *bst[5];      // program bound per stage (NVN stage bits 1 vertex 2 fragment 4 geometry 8 / 16 tessellation, indices as in NvProg.sh), see bind_stages
     blend_st blend[8]; color_st color; chmask_st chmask; ms_st ms; poly_st poly; ds_st ds;
     va_st attr[16]; int nattr; vs_st stream[16]; int nstream;
@@ -4007,6 +4022,8 @@ static NSMutableDictionary<NSString *, id<MTLLibrary>> *g_bf_lib_cache;
 static nvn_hmap g_pso = {.klen = sizeof(pso_key)}, g_dss;
 static nvn_hmap g_pso_fail = {.klen = sizeof(pso_key)};   // pipelines that failed to build: remembered and never retried (previously every draw recompiled it, over 10,000 times in one iPhone session)
 static const char g_tf_mark;
+static char g_geo_mark;
+#define GEO_MARK ((const void *)&g_geo_mark)
 #define TF_MARK ((const void *)&g_tf_mark)   // pso_key.fsfn of a transform feedback capture pipeline (tf_pso)
 static int mtl_fmt_uint(MTLPixelFormat f) {
     switch (f) { case MTLPixelFormatR8Uint: case MTLPixelFormatR16Uint: case MTLPixelFormatR32Uint: case MTLPixelFormatRG8Uint: case MTLPixelFormatRG16Uint: case MTLPixelFormatRG32Uint:
@@ -4043,8 +4060,8 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS
             if (hl) d.fragmentFunction = [hl newFunctionWithName:[NSString stringWithFormat:@"hl%d", i]];
             break; } }
     if (fs && fs->fn_u && !fsfn && d.fragmentFunction == fs->fn && pso_all_uint(k, fs, merged)) d.fragmentFunction = fs->fn_u;   // integer colour targets (dmsh_build.py fs_uint)
-    if (vs->layer_out) d.inputPrimitiveTopology = MTLPrimitiveTopologyClassTriangle;   // a vertex shader that writes the layer (pso_get_)
+    if (vs->layer_out && k->fsfn != GEO_MARK) d.inputPrimitiveTopology = vs->geo_output == 1 ? MTLPrimitiveTopologyClassPoint : vs->geo_output == 6 ? MTLPrimitiveTopologyClassLine : MTLPrimitiveTopologyClassTriangle;
-    if (k->fsfn == TF_MARK) { d.vertexFunction = vs->tf_fn; d.fragmentFunction = nil; d.rasterizationEnabled = NO; }   // transform feedback capture (tf_draw)
+    if (k->fsfn == TF_MARK || k->fsfn == GEO_MARK) { d.vertexFunction = k->fsfn == GEO_MARK ? vs->geo_fn : vs->tf_fn; d.fragmentFunction = nil; d.rasterizationEnabled = NO; }   // transform feedback capture (tf_draw)
     d.rasterSampleCount = k->samples ? k->samples : 1;
     d.alphaToCoverageEnabled = k->a2c;
     for (int i = 0; i < 8; i++) {
@@ -4068,7 +4085,7 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS
     }
     d.depthAttachmentPixelFormat = (MTLPixelFormat)k->dfmt; d.stencilAttachmentPixelFormat = (MTLPixelFormat)k->sfmt;
     MTLVertexDescriptor *vd = [MTLVertexDescriptor vertexDescriptor];
-    const nvn_refl *rf = &vs->rf; int nconv = 0, on = vin_on(vs, k, k->fsfn == TF_MARK);
+    const nvn_refl *rf = &vs->rf; int nconv = 0, on = vin_on(vs, k, k->fsfn == TF_MARK || k->fsfn == GEO_MARK);
     if (on) { d.vertexFunction = vs->vfn; g_vin_psos++; }   // the shadow-input entry point (see vin_on)
     #define ZERO_ATTR(i, ct_) do { vd.attributes[i].format = (ct_) == NVN_CT_UINT ? MTLVertexFormatUInt4 : (ct_) == NVN_CT_INT ? MTLVertexFormatInt4 : MTLVertexFormatFloat4; vd.attributes[i].offset = 0; vd.attributes[i].bufferIndex = 30; \
                                   vd.layouts[30].stride = 16; vd.layouts[30].stepFunction = MTLVertexStepFunctionConstant; vd.layouts[30].stepRate = 0; } while (0)   // reads the zero buffer
@@ -4375,7 +4392,7 @@ static inline void args_sync_dirty(void) {
     }
 }
 static void bind_args(id<MTLCommandEncoder> en, int nvn_stage, int metal_vertex, id<MTLBuffer> ab, uint64_t gpu, size_t off, NSUInteger ai) {
-    if (nvn_stage == 5) enc_c_args((id<MTLComputeCommandEncoder>)en, gpu, ai);
+    if (nvn_stage == 5 || nvn_stage == 2) enc_c_args((id<MTLComputeCommandEncoder>)en, gpu, ai);
     else enc_r_args((id<MTLRenderCommandEncoder>)en, metal_vertex ? MTLRenderStageVertex : MTLRenderStageFragment, gpu, ai);
 }
 // Static constant buffers of route B: the Ryujinx support_buffer (layout in Ryujinx.Graphics.Shader/SupportBuffer.cs: alpha_test u32, is_bgra u32[8],
@@ -4415,6 +4432,7 @@ static int ubo_samp_id(uint32_t h) { int s = (int)(h >> 20); return s ? s + S.sa
 static uint32_t g_dispatch_dim[3];   // the current dispatch's workgroup counts (driver segment c0 + 0x200 for compute)
 static uint32_t g_draw_binst;   // the current draw's base instance (draw), for c0 + 0x644
 static __unsafe_unretained id<MTLBuffer> g_c0_dim_buf; static size_t g_c0_dim_off;   // where encode_args put them for this dispatch (nil: the shader reads no c0); an indirect dispatch copies its GPU-written counts there
+static uint64_t g_geo_bind[5];
 static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, int metal_vertex);
 static int encode_args(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, int metal_vertex) { AP_B; int r = encode_args_(en, sh, nvn_stage, metal_vertex); AP_E(3); return r; }
 static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, int metal_vertex) {
@@ -4436,6 +4454,7 @@ static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, i
     for (int i = 0; i < sh->rf.nb; i++) {
         const nvn_rbind *b = &sh->rf.b[i];
         int reg = b->reg, aid = b->aid;
+        if (b->kind == NVN_RB_GEOMETRY) { slots[aid] = g_geo_bind[reg]; continue; }
         if (b->kind == NVN_RB_CBUF) {
             const p_bind_buf *u = reg < 16 ? &S.ubo[nvn_stage][reg] : NULL;
             void *mb; uint64_t mo, gpu;
@@ -4928,6 +4947,7 @@ static int state_op(uint32_t op, const void *a) {
         case OP_TF_BIND: { const p_tf_bind *b = a; if (b->index >= 0 && b->index < 4) S.tf[b->index] = *b; break; }
         case OP_TF_STATE: { const p_tf_state *b = a;   // Begin writes from the start of each buffer; Pause / Resume keep the offsets (the host keeps them itself, the game's control buffer is not used)
             if (b->what == 0) { S.tf_on = 1; memset(S.tf_off, 0, sizeof S.tf_off); } else S.tf_on = b->what == 3; break; }
+        case OP_PRIM_RESTART: S.restart = *(const p_prim_restart *)a; break;
         case OP_RAST_DISCARD: S.rast_discard = ((const p_rast_discard *)a)->on; break;
         default: return 0;
     }
@@ -6529,6 +6549,8 @@ static void tf_draw(NvProg *pr, const p_draw *dw, const p_draw_idx *di) {
        (unsigned long long)S.tf_off[0], (unsigned long long)S.tf_off[1], (unsigned long long)S.tf_off[2], (unsigned long long)S.tf_off[3]);
 }
 
+#include "nvn_geometry.inc"
+
 // Lazy clears. A set of targets that gets only clears before the game sets other targets used to open a pass just to land them (rp_flush), storing the cleared attachments, and the next
 // pass that attached them loaded them back. Breath of the Wild clears its G-buffer one target at a time (depth with the lighting target, then albedo, normals, the fourth), then attaches all
 // five: four passes stored 27 MB of cleared memory at 1600x900 and the G-buffer pass loaded it again. Such clears now wait (g_lc): the next pass that attaches the same texture level and
@@ -6608,7 +6630,7 @@ static int lc_keeps(uint32_t op) {
         case OP_SCISSOR: case OP_BIND_PROGRAM: case OP_BIND_STATE: case OP_BIND_VATTR: case OP_BIND_VSTREAM: case OP_BIND_VBUF:
         case OP_BIND_UBO: case OP_BIND_SSBO: case OP_BIND_TEX: case OP_BIND_TEX1: case OP_BIND_SAMP: case OP_BIND_IMG: case OP_STENCIL:
         case OP_POLY_OFFSET: case OP_VIEWPORTS: case OP_DEPTH_RANGE: case OP_DEPTH_CLAMP: case OP_VP_SWIZZLE: case OP_BLEND_COLOR: case OP_COPY_STRIDE:
-        case OP_TF_BIND: case OP_TF_STATE: case OP_RAST_DISCARD: case OP_BARRIER: case OP_SET_TEX_POOL: case OP_SET_SAMP_POOL: case OP_RENDER_COND:
+        case OP_TF_BIND: case OP_TF_STATE: case OP_RAST_DISCARD: case OP_PRIM_RESTART: case OP_BARRIER: case OP_SET_TEX_POOL: case OP_SET_SAMP_POOL: case OP_RENDER_COND:
         case OP_UPDATE_UB: case OP_COPY_B2B: case OP_CLEAR_BUF: return 1;
         default: return 0;
     }
@@ -6616,12 +6638,19 @@ static int lc_keeps(uint32_t op) {
 
 static void draw(const p_draw *dw, const p_draw_idx *di) {
     g_draw_binst = (uint32_t)(dw ? dw->base_instance : di ? di->base_instance : 0);   // c0 + 0x644 of the stages that read the driver segment (encode_args)
+    if (!g_geo_program && S.prog) {
+        NvProg *geometry_program = prog_get(S.prog);
+        if (geometry_program && geometry_program->sh[2] && geometry_program->sh[2]->geo_vertices) {
+            if ((dw && dw->mdi_cnt && !mdi_live(dw->mdi_cnt, dw->mdi_i)) || (di && di->mdi_cnt && !mdi_live(di->mdi_cnt, di->mdi_i))) return;
+            geo_draw(geometry_program, dw, di); return;
+        }
+    }
     int mdx_on = (dw && dw->mdi_cnt) || (di && di->mdi_cnt) ? mdx_wanted(dw, di) : 0;   // an indirect-count multi-draw's record masked on the GPU (mdx_run), else read after a wait
     if (!mdx_on && ((di && di->mdi_cnt && !mdi_live(di->mdi_cnt, di->mdi_i)) || (dw && dw->mdi_cnt && !mdi_live(dw->mdi_cnt, dw->mdi_i)))) return;   // past the multi-draw's count
     if (mdx_on && mdx_next(dw, di)) return;   // the next draw of a masked multi-draw: only the draw is encoded
     g_mdxc.next = NULL;
     if (!S.prog) { RED_ONCE(18, "no program bound before the draw"); PD("Z %llu draw skipped: no program bound", g_stat_pass + !g_enc); return; }
-    __unsafe_unretained NvProg *pr = prog_get(S.prog);
+    __unsafe_unretained NvProg *pr = g_geo_program ? g_geo_program : prog_get(S.prog);
     draw_gs_check(pr);   // a geometry stage left with no function (see prog_fold_gs)
     if (S.rast_discard) { tf_draw(pr, dw, di); return; }   // the vertex stage only: transform feedback capture, or nothing (tf_draw)
     if (g_nubd && (g_enc || ubd_lazy()) && ubd_conflict(pr, dw, di)) { g_ubd_conflicts++; rp_end(); ubd_land(); }   // land the deferred inline updates before this draw (see ubd_conflict)
@@ -6635,7 +6664,7 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
         if (!mdx_find(ind, cnt, &mdx)) { mdx_run(dw, di); if (!mdx_find(ind, cnt, &mdx)) mdx = 0; }
         if (!mdx && !mdi_live(cnt, i)) return; }
     if (di && wide_wanted(di) && !wide_find(di)) wide_run(di);   // 8-bit indices written earlier in this submission: widened on the GPU before the render pass (see wide_run)
-    if (!g_defer_replaying) {
+    if (!g_defer_replaying && !g_geo_program) {
         if (g_defer.active && g_defer.in_p && !g_defer.at_q) { if (defer_take(dw, di)) return; }
     }
     // Draws that cannot draw anything (both faces culled, empty scissor) return before the render pass is opened: otherwise the render pass is already open and on TBDR the whole attachment is loaded and stored for nothing
@@ -6674,10 +6703,10 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
     if (!g_enc) rr_begin();
 #endif
     int defer_q_draw = !g_enc && !g_defer_replaying && g_defer_natq > 0;   // The draw that opens Q's render pass: after the render pass is opened as usual through hoist_try / fuse_try, replay P first and then draw this one
-    if (!g_enc && tbdr_on()) hoist_try();   // First hoist and execute the pure-clear render passes sandwiched between this render pass and its continuation (before rp_begin_ex, not occupying this render pass)
+    if (!g_enc && tbdr_on() && !g_geo_program) hoist_try();   // First hoist and execute the pure-clear render passes sandwiched between this render pass and its continuation (before rp_begin_ex, not occupying this render pass)
     // The deferral decision comes after hoist_try: hoisted pure-clear render passes are skipped in the lookahead, and only then can P's same-target continuation segment be recognized
-    if (!g_enc && tbdr_on() && !defer_has_in_p() && !g_defer_replaying && !g_fz.active && !g_fz_checking && g_exec_next) { defer_try(dw, di); if (g_defer_cur->active && g_defer_cur->in_p) return; }
+    if (!g_enc && tbdr_on() && !g_geo_program && !defer_has_in_p() && !g_defer_replaying && !g_fz.active && !g_fz_checking && g_exec_next) { defer_try(dw, di); if (g_defer_cur->active && g_defer_cur->in_p) return; }
-    if (!g_enc && tbdr_on() && !g_fz_checking) fuse_try();   // Readable tile render pass merge is on by default; when the structure / same-pixel proof does not hold it still falls back to a normal render pass, and it is not re-entered during the CHECK redraw
+    if (!g_enc && tbdr_on() && !g_fz_checking && !g_geo_program) fuse_try();   // Readable tile render pass merge is on by default; when the structure / same-pixel proof does not hold it still falls back to a normal render pass, and it is not re-entered during the CHECK redraw
     int need = g_enc || !tbdr_on() ? 3 : pass_ds_need(), mode = defer_q_draw ? 0 : need == 0 ? 2 : need == 2 ? 3 : need == 1 ? 4 : 0;   // 2 attach neither, 3 attach stencil only, 4 attach depth only
     int skip_d = mode != 0;
     if (!g_enc) { g_fd_dw = dw; g_fd_di = di; }
@@ -6818,6 +6847,10 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
 #endif
     udc_note_draw(pr);
     swrite_note(pr);
+    if (g_geo_program) {
+        enc_r_draw_indexed_indirect(en, mp, MTLIndexTypeUInt32, g_geo_indices, g_geo_index_bytes, g_geo_args);
+        return;
+    }
     // A draw with 0 primitives (the game does send count=0) draws no pixels, so skip encoding; Metal requires the count to be non-zero, otherwise the validation layer aborts on the spot (measured on the camp / wilderness / combat pages)
     if ((dw && dw->ind) || (di && di->ind)) {   // DrawArraysIndirect / DrawElementsIndirect: NVN's argument records {count, instances, first, base instance} and {count, instances, first index,
         // base vertex, base instance} have the layouts of MTLDrawPrimitivesIndirectArguments and MTLDrawIndexedPrimitivesIndirectArguments; quads, fans and 8-bit indices use the converted ones (ixc_run)
@@ -7592,6 +7625,7 @@ void mtl_exec(const uint8_t *p) {
             case OP_TF_BIND: { const p_tf_bind *b = a; PD("Y tf bind %d addr=%#llx size=%llu", b->index, (unsigned long long)b->addr, (unsigned long long)b->size); state_op(h.op, a); break; }
             case OP_TF_STATE: { const p_tf_state *b = a; static const char *const w[4] = {"begin", "end", "pause", "resume"}; PD("Y tf %s ctl=%#llx written %llu/%llu/%llu/%llu", w[b->what & 3], (unsigned long long)b->ctl,
                 (unsigned long long)S.tf_off[0], (unsigned long long)S.tf_off[1], (unsigned long long)S.tf_off[2], (unsigned long long)S.tf_off[3]); state_op(h.op, a); break; }
+            case OP_PRIM_RESTART: state_op(h.op, a); break;
             case OP_RAST_DISCARD: PD("Y rasterizer discard %d", ((const p_rast_discard *)a)->on); state_op(h.op, a); break;
             case OP_COPY_B2T: copy_b2t(a); break;
             case OP_COPY_T2B: copy_t2b(a); break;

dm/gfx/nvn/nvn_obj.h

@@ -102,7 +102,7 @@ enum {
     OP_BIND_PROGRAM, OP_BIND_STATE, OP_BIND_VATTR, OP_BIND_VSTREAM, OP_BIND_VBUF,
     OP_BIND_UBO, OP_BIND_SSBO, OP_BIND_TEX, OP_BIND_SAMP,
     OP_STENCIL, OP_POLY_OFFSET, OP_DRAW, OP_DRAW_IDX, OP_DISPATCH, OP_BARRIER,
-    OP_SET_TEX_POOL, OP_SET_SAMP_POOL, OP_COPY_B2T, OP_COPY_STRIDE, OP_BLEND_COLOR, OP_COPY_T2B, OP_BIND_IMG, OP_COPY_T2T, OP_RESET_COUNTER, OP_CLEAR_TEX, OP_CLEAR_BUF, OP_DEPTH_CLAMP, OP_RENDER_COND, OP_FENCE_SYNC, OP_TF_BIND, OP_TF_STATE, OP_RAST_DISCARD, OP_VP_SWIZZLE, OP_BIND_TEX1, OP_COUNT_
+    OP_SET_TEX_POOL, OP_SET_SAMP_POOL, OP_COPY_B2T, OP_COPY_STRIDE, OP_BLEND_COLOR, OP_COPY_T2B, OP_BIND_IMG, OP_COPY_T2T, OP_RESET_COUNTER, OP_CLEAR_TEX, OP_CLEAR_BUF, OP_DEPTH_CLAMP, OP_RENDER_COND, OP_FENCE_SYNC, OP_TF_BIND, OP_TF_STATE, OP_RAST_DISCARD, OP_VP_SWIZZLE, OP_BIND_TEX1, OP_PRIM_RESTART, OP_COUNT_
 };
 // Payload (immediately follows op_hdr)
 typedef struct { int32_t n; const tex_t *color[8]; view_t cview[8]; uint8_t has_cview[8]; const tex_t *depth; view_t dview; uint8_t has_dview; } p_set_rt;
@@ -112,6 +112,7 @@ typedef struct { float depth; int32_t depth_mask, stencil, stencil_mask; } p_cle
 typedef struct { int32_t first, n; float v[16][6]; } p_viewports;          // x, y, w, h, near, far (near and far are overridden by DEPTH_RANGE)
 typedef struct { int32_t x, y, w, h; } p_scissor;
 typedef struct { float n, f; } p_depth_range;
+typedef struct { uint32_t on, index; } p_prim_restart;
 typedef struct { int32_t on; } p_depth_clamp;   // SetDepthClamp: depth outside [near, far] is clamped instead of clipped (Metal's depth clip mode)
 typedef struct { int32_t first, n; uint8_t sw[16][4]; } p_vp_swizzle;   // SetViewportSwizzles: per viewport the source of clip x, y, z, w (NVNviewportSwizzle) XOR the identity (0, 2, 4, 6), so 0 = unchanged
 typedef struct { const prog_t *prog; int32_t stages; } p_bind_program;

dm/gfx/nvn/nvn_refl.h

@@ -7,7 +7,7 @@
 #include <stdint.h>
 
 enum { NVN_RB_CBUF = 1, NVN_RB_SAMP, NVN_RB_TEX, NVN_RB_IMG, NVN_RB_TBUF,
-       NVN_RB_SBUFP, NVN_RB_SUPPORT, NVN_RB_C1, NVN_RB_DRV0, NVN_RB_DRV2, NVN_RB_SAMPT, NVN_RB_UBOP, NVN_RB_TEXU, NVN_RB_SAMPU, NVN_RB_SAMPI, NVN_RB_SBUFU, NVN_RB_TEXD, NVN_RB_SAMPD, NVN_RB_IMGA, NVN_RB_TEXA, NVN_RB_SAMPA };
+       NVN_RB_SBUFP, NVN_RB_SUPPORT, NVN_RB_C1, NVN_RB_DRV0, NVN_RB_DRV2, NVN_RB_SAMPT, NVN_RB_UBOP, NVN_RB_TEXU, NVN_RB_SAMPU, NVN_RB_SAMPI, NVN_RB_SBUFU, NVN_RB_TEXD, NVN_RB_SAMPD, NVN_RB_IMGA, NVN_RB_TEXA, NVN_RB_SAMPA, NVN_RB_GEOMETRY };
 // TEXA / SAMPA: an array of n handles in the driver table c2 that the shader indexes at run time (dmsh_build.py c2_array): element k = the texture / sampler of the slot whose c2 word lies
 // at byte off + 8k (combined, else separate); words outside both tables give the dummy. Luigi's Mansion 3's shadow cascades
 // IMGA: SPIRV-Cross's MSL 2.3 image atomics: a device atomic_uint / atomic_int pointer (<image name>_atomic) beside the image of NVN image slot r; the replay binds the image's storage address

dm/gfx/nvn/nvn_refl.m

@@ -81,6 +81,7 @@ int nvn_refl_from_json(NSDictionary *j, nvn_refl *o, char *err, int errlen) {
             else if ([k isEqualToString:@"sbuf"]) { r->kind = NVN_RB_SBUFP; r->rw = sbuf_rw(b); }
             else if ([k isEqualToString:@"ubo_ptr"]) r->kind = NVN_RB_UBOP;
             else if ([k isEqualToString:@"sbuf_ubo"]) { r->kind = NVN_RB_SBUFU; r->rw = sbuf_rw(b); r->off = (uint16_t)[b[@"off"] intValue]; }
+            else if ([k isEqualToString:@"geometry"]) { if (reg < 0 || reg > 4) FAIL("invalid geometry helper slot"); r->kind = NVN_RB_GEOMETRY; }
             else if ([k isEqualToString:@"support"]) r->kind = NVN_RB_SUPPORT;
             else if ([k isEqualToString:@"c1"]) r->kind = NVN_RB_C1;
             else if ([k isEqualToString:@"drvcb"]) r->kind = [b[@"bank"] intValue] == 2 ? NVN_RB_DRV2 : NVN_RB_DRV0;

dm/gfx/tests/.gitignore

@@ -44,6 +44,8 @@ t_ub
 t_ub.dSYM/
 t_vbuf
 t_vbuf.dSYM/
+t_geometry
+t_geometry.dSYM/
 t_dirty_race
 t_dirty_race.dSYM/
 t_vin

dm/gfx/tests/Makefile

@@ -59,6 +59,8 @@ t_ub: t_ub.m $(MTL_T_DEP)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_ub.m $(MTL_T_SRC) $(FW)
 t_vbuf: t_vbuf.m $(MTL_T_DEP)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_vbuf.m $(MTL_T_SRC) $(FW)
+t_geometry: t_geometry.m $(MTL_T_DEP)
+	clang $(CFLAGS) -fobjc-arc -o $@ t_geometry.m $(MTL_T_SRC) $(FW)
 t_mip_extent: t_mip_extent.m ../nvn/hle_nvn.c ../../host/crash.c ../../host/hle_bigstack.c t_host_stubs.c $(GFX) $(GFX_H)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_mip_extent.m ../../host/crash.c ../../host/hle_bigstack.c t_host_stubs.c $(filter-out ../nvn/nvn_mtl.m,$(GFX)) $(FW)
 t_rt_range_shape: t_rt_range_shape.m $(MTL_T_DEP)
@@ -80,7 +82,7 @@ t_pfv: t_pfv.m $(MTL_T_DEP)
 # The Vulkan layer (#included) over the whole NVN layer, whose device and present path it uses
 t_vk_desc: t_vk_desc.m ../vk/hle_vk.m $(MTL_T_DEP)
 	clang $(CFLAGS) -I../../third_party/vulkan-headers/include -fobjc-arc -o $@ t_vk_desc.m $(MTL_T_SRC) ../nvn/nvn_mtl.m $(FW)
-run: t_nvn t_dirty_slots t_dirty_race t_builtin t_vin t_tbuf2d t_deswz_gpu t_pace t_fstep t_refl t_hmap t_cover t_fuse_check t_present_crop t_reorder_chain t_defer_chain t_zskip t_xv t_ub t_vbuf t_mip_extent t_rt_range_shape t_rt_range_ro t_rt_range t_rt_depth_ro t_pso_skey t_pfv t_view_levels t_view_swizzle t_vk_desc
+run: t_nvn t_dirty_slots t_dirty_race t_builtin t_vin t_tbuf2d t_deswz_gpu t_pace t_fstep t_refl t_hmap t_cover t_fuse_check t_present_crop t_reorder_chain t_defer_chain t_zskip t_xv t_ub t_vbuf t_geometry t_mip_extent t_rt_range_shape t_rt_range_ro t_rt_range t_rt_depth_ro t_pso_skey t_pfv t_view_levels t_view_swizzle t_vk_desc
 	./t_nvn
 	./t_dirty_slots
 	./t_dirty_race
@@ -101,6 +103,7 @@ run: t_nvn t_dirty_slots t_dirty_race t_builtin t_vin t_tbuf2d t_deswz_gpu t_pac
 	./t_xv
 	./t_ub
 	MTL_DEBUG_LAYER=1 ./t_vbuf
+	MTL_DEBUG_LAYER=1 ./t_geometry
 	DM_NVN_VBUF_CACHE=0 MTL_DEBUG_LAYER=1 ./t_vbuf
 	MTL_DEBUG_LAYER=1 ./t_mip_extent
 	MTL_DEBUG_LAYER=1 ./t_mip_extent --stencil-only

dm/gfx/tests/t_geometry.m

@@ -0,0 +1,180 @@
+// Geometry assembly, GPU strip compaction, and production capture/compute/raster replay.
+#include "../nvn/nvn_mtl.m"
+void *front_metal_layer(void) { return NULL; }
+module_t *module_of(const void *p) { (void)p; return NULL; }
+module_t g_mod[MAX_MOD]; int g_nmod;
+void svc_thread_dump(void) {}
+uint64_t now_ns(void) { return bp_now(); }
+static int checks, failed;
+#define EXPECT(name, want, got) do { uint64_t check_want=(want), check_got=(got); checks++; if(check_want!=check_got) { fprintf(stderr,"%s: expected %llu, got %llu\n",name,(unsigned long long)check_want,(unsigned long long)check_got); failed++; } } while(0)
+static void finish(void) {
+    EXPECT("submission completed",1,mtl_wait(mtl_sub_commit(),5000000000ull));
+    EXPECT("GPU error",0,mtl_last_commit_error());
+}
+static id<MTLLibrary> library(NSString *source) {
+    NSError *e=nil; id<MTLLibrary> l=[g_dev newLibraryWithSource:source options:nil error:&e];
+    if(!l) { fprintf(stderr,"%s\n",e.localizedDescription.UTF8String); abort(); } return l;
+}
+static void assembly(void) {
+    uint32_t a[]={10,11,12,13,UINT32_MAX,20,21,22}, out[32];
+    uint32_t want[]={10,11,12,12,11,13,20,21,22};
+    EXPECT("strip count",9,nvn_geo_assemble(5,a,8,1,UINT32_MAX,NULL,0));
+    EXPECT("strip write",9,nvn_geo_assemble(5,a,8,1,UINT32_MAX,out,32));
+    EXPECT("strip winding and restart",0,memcmp(out,want,sizeof want));
+    for(int i=0;i<32;i++) out[i]=123;
+    EXPECT("short output rejected",SIZE_MAX,nvn_geo_assemble(5,a,8,1,UINT32_MAX,out,8));
+    for(int i=0;i<32;i++) EXPECT("short output untouched",123,out[i]);
+    uint32_t fan[]={10,11,12,10,12,13};
+    EXPECT("fan",6,nvn_geo_assemble(6,a,4,0,0,out,32));
+    EXPECT("fan order",0,memcmp(out,fan,sizeof fan));
+    uint32_t loop[]={10,11,11,12,12,10};
+    EXPECT("loop",6,nvn_geo_assemble(2,a,3,0,0,out,32));
+    EXPECT("loop closure",0,memcmp(out,loop,sizeof loop));
+    EXPECT("line adjacency",8,nvn_geo_assemble(11,a,5,0,0,out,32));
+    EXPECT("adjacent second edge",11,out[4]);
+    EXPECT("triangle adjacency",6,nvn_geo_assemble(12,a,8,0,0,out,32));
+    EXPECT("unsupported",SIZE_MAX,nvn_geo_assemble(99,a,8,0,0,NULL,0));
+}
+static void pack_test(void) {
+    id<MTLLibrary> l=library(@(k_nvn_geo_pack_msl)); NSError *e=nil;
+    id<MTLComputePipelineState> p=[g_dev newComputePipelineStateWithFunction:[l newFunctionWithName:@"geo_pack_indices"] error:&e];
+    if(!p) abort();
+    uint32_t data[]={0,1,2,3,UINT32_MAX,4,5,6, UINT32_MAX,UINT32_MAX,UINT32_MAX,UINT32_MAX,UINT32_MAX,UINT32_MAX,UINT32_MAX,UINT32_MAX, 8,9,10,UINT32_MAX,11,12,13,UINT32_MAX};
+    uint32_t want[]={0,1,2,2,1,3,4,5,6,8,9,10,11,12,13};
+    id<MTLBuffer> b=[g_dev newBufferWithLength:4096 options:MTLResourceStorageModeShared]; res_add(b);
+    for(int t=0;t<3;t++) {
+        memset(b.contents,0xcc,4096); memcpy(b.contents,data,sizeof data);
+        struct {uint64_t src,dst,args;uint32_t groups,stride,topology,capacity;} j={b.gpuAddress,b.gpuAddress+512,b.gpuAddress+1024,3,8,t==0?7:t==1?6:1,100};
+        memcpy((char*)b.contents+2048,&j,sizeof j);
+        mtl_sub_begin(); id<MTL4ComputeCommandEncoder> en=enc_c_begin(NULL); enc_c_pso(en,p);enc_c_args(en,b.gpuAddress+2048,0);
+        enc_c_dispatch(en,MTLSizeMake(1,1,1),MTLSizeMake(1,1,1));enc_c_end(en);finish();
+        uint32_t *args=(uint32_t*)((char*)b.contents+1024);
+        EXPECT("compacted index count",t==0?15:t==1?18:13,args[0]);
+        EXPECT("indirect instance count",1,args[1]);
+        if(!t) EXPECT("strip order across empty invocations",0,memcmp((char*)b.contents+512,want,sizeof want));
+        EXPECT("output bound sentinel",0xcccccccc,((uint32_t*)((char*)b.contents+512))[args[0]]);
+    }
+}
+static NvShader *shader(id<MTLLibrary> l, NSString *name, int stage, const int *helpers, int n) {
+    NvShader *s=[NvShader new];s->stage=stage;s->fn=[l newFunctionWithName:name];s->rf.nb=n;
+    if(n) {s->ae=[s->fn newArgumentEncoderWithBufferIndex:0];s->arg_len=s->ae.encodedLength;}
+    for(int i=0;i<n;i++)s->rf.b[i]=(nvn_rbind){.kind=NVN_RB_GEOMETRY,.reg=helpers[i],.aid=i};
+    return s;
+}
+static void replay(int mode) {
+    id<MTLLibrary> l=library(@"#include <metal_stdlib>\nusing namespace metal;\n"
+        "struct C {device float *dst;uint4 info;};\n"
+        "vertex void capture(uint v[[vertex_id]],uint i[[instance_id]],constant C&c[[buffer(13)]]) {uint n=v-c.info.x,inst=i-c.info.z;device float*p=c.dst+(inst*c.info.y+n)*144; p[0]=float(n);p[1]=float(inst);p[2]=0;p[3]=1;}\n"
+        "vertex float4 unused(uint v[[vertex_id]]) {return float4(0,0,0,1);}\n"
+        "struct G{device const float*in[[id(0)]];device float4*out[[id(1)]];device uint*ix[[id(2)]];constant uint4*info[[id(3)]];device const uint*map[[id(4)]];};\n"
+        "kernel void geometry(constant G&g[[buffer(0)]],uint3 t[[thread_position_in_grid]]) {if(t.x>=g.info[1].x||t.y>=g.info[0].y)return;uint k=(t.y*g.info[1].x+t.x)*2+t.z;uint v=g.map[t.x];float2 q=float2(g.in[(t.y*g.info[0].x+v)*144],g.in[(t.y*g.info[0].x+v)*144+1]);float2 a=q*0.5-0.9;float2 d[4]={float2(0),float2(.4,0),float2(0,.4),float2(.4)};for(uint n=0;n<4;n++){g.out[k*4+n]=float4(a+d[n],0,1);g.ix[k*5+n]=k*4+n;}g.ix[k*5+4]=0xffffffffu;}\n"
+        "struct V{device const float4*p[[id(0)]];};vertex float4 reader(uint v[[vertex_id]],constant V&a[[buffer(0)]]){return a.p[v];}\n"
+        "fragment float4 red(){return float4(1,0,0,1);}\n");
+    int helpers[]={0,1,2,3,4}, rh[]={1};
+    NvProg *pr=[NvProg new];pr->sh[0]=shader(l,@"unused",0,NULL,0);pr->sh[0]->geo_fn=[l newFunctionWithName:@"capture"];
+    pr->sh[1]=shader(l,@"red",1,NULL,0);pr->sh[2]=shader(l,@"geometry",2,helpers,5);
+    NvShader *gs=pr->sh[2];strcpy(gs->key,"0123456789abcdef_b");gs->geo_vertices=4;gs->geo_invocations=2;gs->geo_output=7;gs->geo_stride=4;gs->geo_index_stride=5;
+    g_shcache=[NSMutableDictionary dictionary];g_shcache[@"0123456789abcdef_bg0"]=gs;g_shcache[@"0123456789abcdef_bgv"]=shader(l,@"reader",0,rh,1);
+    static NSMutableArray *keep;if(!keep)keep=[NSMutableArray array];[keep addObject:pr];[keep addObject:g_shcache[@"0123456789abcdef_bgv"]];
+    static prog_t cp[7];cp[mode].serial=8700+mode;serial_put(cp[mode].serial,pr);
+    tex_t target={.magic=MAGIC('T','E','X',0),.serial=8800+mode};target.d.width=target.d.height=32;target.d.depth=target.d.samples=target.d.levels=1;target.d.target=NVN_TEXTURE_TARGET_2D;target.d.format=NVN_FORMAT_RGBA8;
+    uint8_t *commands; if(posix_memalign((void**)&commands,getpagesize(),65536))abort();memset(commands,0,65536);mtl_pool_add(commands,65536);
+    size_t used=0;
+    #define EMIT(OP,X) do {size_t z=(8+sizeof(X)+7)&~7ull;op_hdr h={OP,(uint32_t)z};memcpy(commands+used,&h,8);memcpy(commands+used+8,&(X),sizeof(X));used+=z;}while(0)
+    p_set_rt rt={.n=1};rt.color[0]=&target;EMIT(OP_SET_RT,rt);
+    p_viewports vp={.n=1};vp.v[0][2]=vp.v[0][3]=32;EMIT(OP_VIEWPORTS,vp);p_scissor sc={0,0,32,32};EMIT(OP_SCISSOR,sc);
+    p_clear_color cl={.index=0,.mask=15,.rgba={0,0,0,1}};EMIT(OP_CLEAR_COLOR,cl);
+    p_bind_program bp={&cp[mode],7};EMIT(OP_BIND_PROGRAM,bp);
+    p_draw d={.prim=0,.first=7,.count=mode==6?9000:3,.instances=2,.base_instance=5};
+    if((mode>=1&&mode<=3)||mode==5) {
+        int type=mode==5?1:mode-1;uint32_t ix[]={12,10,UINT32_MAX,11};
+        for(int i=0;i<4;i++)memcpy(commands+8192+(i<<type),ix+i,1<<type);
+        p_prim_restart restart={1,UINT32_MAX};EMIT(OP_PRIM_RESTART,restart);
+        p_draw_idx di={.prim=0,.itype=type,.count=4,.instances=2,.base_vertex=-3,.base_instance=5,.ibuf=(uint64_t)(commands+8192)};
+        if(mode==5){uint32_t args[]={4,2,0,(uint32_t)-3,5};memcpy(commands+16384,args,sizeof args);di.ind=(uint64_t)(commands+16384);}
+        EMIT(OP_DRAW_IDX,di);
+    } else {
+        if(mode==4){uint32_t args[]={3,2,7,5};memcpy(commands+16384,args,sizeof args);d.ind=(uint64_t)(commands+16384);}
+        EMIT(OP_DRAW,d);
+    }
+    op_hdr end={OP_END,8};memcpy(commands+used,&end,8);
+    #undef EMIT
+    state_reset();memset(&R,0,sizeof R);mtl_sub_begin();uint64_t h=(uint64_t)commands;mtl_sub_index(&h,1);mtl_exec(commands);
+    id<MTLBuffer> b=[g_dev newBufferWithLength:4096 options:MTLResourceStorageModeShared];
+    uint32_t *args=NULL;id<MTLBuffer> argsBuffer;
+    for(NvArena *a in g_sub_arenas)if(g_geo_args>=a->gpu&&g_geo_args+20<=a->gpu+a->buf.length){argsBuffer=a->buf;args=(uint32_t*)((char*)argsBuffer.contents+g_geo_args-a->gpu);}
+    [blit() copyFromTexture:tex_mtl(&target) sourceSlice:0 sourceLevel:0 sourceOrigin:MTLOriginMake(0,0,0) sourceSize:MTLSizeMake(32,32,1) toBuffer:b destinationOffset:0 destinationBytesPerRow:128 destinationBytesPerImage:4096];finish();
+    EXPECT("two GS invocations per input primitive",mode==6?216000:72,args?args[0]:0);
+    int red=0;uint8_t *px=b.contents;for(int i=0;i<1024;i++)if(px[i*4]>200)red++;
+    EXPECT("expanded quads rasterized",mode==6?288:216,red);
+    EXPECT("each instance has independent capture and output",1,px[(24*32+4)*4]>200);
+    mtl_pool_del(commands);mtl_drain();free(commands);
+}
+// Optional real-game fixtures stay outside the repository. The supplied directory
+// is a dmsh_build result for the current chain, containing these three GS keys.
+static void translated(const char *directory) {
+    g_msl_dir=@(directory);g_shcache=[NSMutableDictionary dictionary];
+    const char *keys[]={"1b378c5c014e3a32","d57cdb168fdf6927","2341c1be4fcbf690"};
+    float *uniforms;if(posix_memalign((void**)&uniforms,getpagesize(),65536))abort();memset(uniforms,0,65536);mtl_pool_add(uniforms,65536);
+    NvShader *vs=shader_get(@"ba2512f7f6973541_b",0);
+    EXPECT("real capture vertex loaded",1,vs&&vs->geo_fn);
+    if(vs&&vs->geo_fn) {
+        for(int v=0;v<3;v++)for(int c=0;c<7;c++)uniforms[v*7+c]=(v*7+c)*.03125f;
+        id<MTLBuffer> b=[g_dev newBufferWithLength:4096 options:MTLResourceStorageModeShared];res_add(b);memset(b.contents,0xcc,4096);
+        uint64_t head[2]={b.gpuAddress,0};uint32_t info[4]={1,2,5,2};memcpy((char*)b.contents+3072,head,16);memcpy((char*)b.contents+3088,info,16);
+        state_reset();S.ssbo[0][0]=(p_bind_buf){.stage=0,.index=0,.addr=(uint64_t)uniforms,.size=65536};
+        mtl_sub_begin();MTL4RenderPassDescriptor *d=[MTL4RenderPassDescriptor new];d.renderTargetWidth=d.renderTargetHeight=d.defaultRasterSampleCount=1;
+        id<MTL4RenderCommandEncoder> en=enc_r_begin(d);enc_r_pso(en,geo_capture_pso(vs));args_dirty_all();encode_args(en,vs,0,1);enc_r_vbuf(en,13,b.gpuAddress+3072);
+        enc_r_draw_instanced(en,MTLPrimitiveTypePoint,1,2,2,5);enc_r_end(en);finish();
+        for(int inst=0;inst<2;inst++)for(int v=0;v<2;v++) {
+            float *p=(float*)b.contents+(inst*2+v)*144,*q=uniforms+(v+1)*7;
+            EXPECT("capture vertex ID",1,p[0]==q[0]*2-1);EXPECT("capture y",1,p[1]==1-q[1]*2);
+            EXPECT("capture retains guest clip depth",1,p[2]==0);EXPECT("capture w",1,p[3]==1);
+            EXPECT("capture loc0",1,p[16]==q[2]);for(int c=0;c<4;c++)EXPECT("capture loc1",1,p[20+c]==q[c+3]);
+            EXPECT("missing varying default w",1,p[27]==1);
+        }
+        EXPECT("capture range guard",0xcccccccc,((uint32_t*)b.contents)[4*144]);
+    }
+    for(int test=0;test<3;test++) {
+        NvShader *s=shader_get([NSString stringWithFormat:@"%s_bg%d",keys[test],test==2?3:0],2);
+        EXPECT("translated geometry loaded",1,s!=nil);if(!s)continue;
+        NSError *e=nil;id<MTLComputePipelineState> ps=[g_dev newComputePipelineStateWithFunction:s->fn error:&e];if(!ps){fprintf(stderr,"%s\n",e.localizedDescription.UTF8String);abort();}
+        id<MTLBuffer> b=[g_dev newBufferWithLength:32768 options:MTLResourceStorageModeShared];res_add(b);memset(b.contents,0xcc,32768);
+        // Two primitives and two instances: vertex count and primitive count
+        // differ for triangles, exposing either stride being used in its place.
+        int width=test==2?3:1, vertices=width*2;float *in=b.contents;
+        memset(in,0,vertices*2*144*4);
+        for(int inst=0;inst<2;inst++)for(int v=0;v<vertices;v++) {
+            float *p=in+(inst*vertices+v)*144;p[0]=(v%3==1?1:0)+v/3*3+inst*8;p[1]=v%3==2?1:0;p[2]=.25f;p[3]=1;
+            p[16]=p[20]=.4f;p[17]=p[21]=.5f;p[18]=p[22]=.6f;p[19]=p[23]=1;
+        }
+        uniforms[0]=uniforms[1]=1;uniforms[2]=0;uniforms[3]=0;uniforms[4]=.1f;uniforms[5]=.2f;
+        uint32_t vi[8]={(uint32_t)vertices,2,0,0,2,0,0,0};memcpy((char*)b.contents+8192,vi,sizeof vi);
+        uint32_t *map=(uint32_t*)((char*)b.contents+8448);for(int v=0;v<vertices;v++)map[v]=v;
+        state_reset();S.ubo[2][0]=(p_bind_buf){.stage=2,.index=0,.addr=(uint64_t)uniforms,.size=65536};
+        g_geo_bind[0]=b.gpuAddress;g_geo_bind[1]=b.gpuAddress+12288;g_geo_bind[2]=b.gpuAddress+16384;g_geo_bind[3]=b.gpuAddress+8192;g_geo_bind[4]=b.gpuAddress+8448;
+        mtl_sub_begin();id<MTL4ComputeCommandEncoder> en=enc_c_begin(NULL);args_dirty_all();encode_args(en,s,2,0);enc_c_pso(en,ps);enc_c_dispatch(en,MTLSizeMake(1,2,1),MTLSizeMake(32,1,1));enc_c_end(en);finish();
+        float *out=(float*)((char*)b.contents+12288);uint32_t *ix=(uint32_t*)((char*)b.contents+16384);
+        for(int group=0;group<4;group++) {
+            for(unsigned v=0;v<s->geo_vertices;v++) {
+                EXPECT("emission index and instance stride",group*s->geo_vertices+v,ix[group*s->geo_index_stride+v]);
+                float *p=out+(group*s->geo_vertices+v)*s->geo_stride;
+                int iv=(group/2)*vertices+(group%2)*width+(test==2?v:0);float *q=in+iv*144;
+                float x=q[0]+(test==2?0:(v&1)?.3f:-.3f),y=q[1]+(test==2?0:v<2?.6f:-.6f);
+                EXPECT("emitted x",1,fabsf(p[0]-x)<1e-5f);EXPECT("emitted y",1,fabsf(p[1]-y)<1e-5f);
+                EXPECT("emitted z",1,fabsf(p[2]-.25f)<1e-5f);EXPECT("emitted w",1,p[3]==1);
+                EXPECT("default point size",1,p[4]==1);
+                if(test==2) {for(int c=0;c<3;c++) {float want=c==v?(v==0?sqrtf(.5f):1):0;EXPECT("wireframe edge distance",1,fabsf(p[9+c]-want)<1e-5f);}}
+            }
+            for(unsigned v=s->geo_vertices;v<s->geo_index_stride;v++)EXPECT("restart and unused index padding",UINT32_MAX,ix[group*s->geo_index_stride+v]);
+        }
+        EXPECT("output bound sentinel",0xcccccccc,((uint32_t*)out)[4*s->geo_vertices*s->geo_stride]);
+        EXPECT("index bound sentinel",0xcccccccc,ix[4*s->geo_index_stride]);
+    }
+    mtl_pool_del(uniforms);mtl_drain();free(uniforms);
+}
+int main(int argc,char **argv) {@autoreleasepool {
+    setvbuf(stdout,NULL,_IONBF,0);assembly();mtl_init();pack_test();for(int mode=0;mode<7;mode++)replay(mode);
+    if(argc>1)translated(argv[1]);
+    printf("t_geometry: %d checks, %d failed\n",checks,failed);return failed!=0;
+}}

dm/gfx/tools/dmsh_build.py

@@ -7,6 +7,7 @@
 # Environment IOSOUT=dir: the same MSL also goes through dm/ios/tools/ios_metal.compile to produce an iphoneos metallib + reflection (chain_metal records the iOS compile convention)
 # Final summary: success / failure counts and the distribution of reasons; distribution of driver segment read offsets (those outside the range the replay layer can fill are listed separately)
 import ast, inspect, os, re, sys, json, struct, hashlib, shutil, subprocess, tempfile, collections, threading, concurrent.futures as cf
+import dmsh_geometry
 DM_ROOT = os.environ.get('DM_ROOT') or os.path.expanduser('~/dm-work')   # Workspace root (game files, build outputs and run artifacts are all under it)
 HOME = os.path.expanduser('~')
 dm_dir, out = sys.argv[1:3]; want = set(sys.argv[3:]); cache_root = out
@@ -106,6 +107,7 @@ def cache_outputs(key, digest, cached=False):
             ('ios', IOSOUT, [key + '_b.metallib', key + '_b.json'])):
         if not directory: continue
         base = os.path.join(root, label) if cached else directory
+        names += [key + s + ext for s in ('_bg0', '_bg1', '_bg2', '_bg3', '_bg4', '_bgv') for ext in ('.metallib', '.json', '.metal') if label == 'mac' or ext != '.metal']
         found = [n for n in names if os.path.exists(os.path.join(base, n))]
         if key + '_b.metallib' not in found or key + '_b.json' not in found: return None
         files[label] = found
@@ -447,7 +449,8 @@ def reflect(key, stage, msl, res):
     for m in re.finditer(r'^\s*(.+?)\s+(\w+)\s*\[\[id\((\d+)\)\]\];', body.group(1) if body else '', re.M):
         ty, nm, aid = m.group(1).strip(), m.group(2), int(m.group(3))
         e = {'msl_name': nm, 'msl_type': ty, 'arg_id': aid}
-        if nm == 'support_buffer': e.update(kind='support')
+        if nm in dmsh_geometry.HELPERS: e.update(kind='geometry', reg=dmsh_geometry.HELPERS[nm])
+        elif nm == 'support_buffer': e.update(kind='support')
         elif re.fullmatch(r'[a-z]+_c(\d+)', nm):
             n = int(re.fullmatch(r'[a-z]+_c(\d+)', nm).group(1))
             e.update(kind='cbuf', reg=n - 3) if 3 <= n <= 16 else e.update(kind='c1' if n == 1 else 'drvcb', bank=n)
@@ -512,7 +515,7 @@ def unfillable(dr):
 
 RULES = []   # MSL rewrite rules attached by shader name (function, applicability test); the former dark-range remap toe has been removed (back to the original Switch brightness)
 # The first-stage code fingerprint does not include dmsh_fetch.py; readable-tile variants get a separate second-stage cache keyed on the generated MSL.
-CODE_FP = hashlib.sha256(source_without_rule_bodies().encode() + b''.join(b'\0' + open(m.__file__, 'rb').read() for m in (dmsh_chain,) + ((ios_metal,) if IOSOUT else ()))).hexdigest()
+CODE_FP = hashlib.sha256(source_without_rule_bodies().encode() + b''.join(b'\0' + open(m.__file__, 'rb').read() for m in (dmsh_chain, dmsh_geometry) + ((ios_metal,) if IOSOUT else ()))).hexdigest()
 TOOL_FP = tool_fingerprint().encode()
 
 _CAST = ('as_type<int>(', 'as_type<uint>(')
@@ -723,16 +726,71 @@ def _ep_drop(x, ids):   # OpEntryPoint without the given interface IDs: model, f
     y[0] = len(y) << 16 | (x[0] & 0xffff)
     return y
 
+def geometry_finish(key, w, j, c1):
+    """Build all five input-topology variants; the draw selects one at replay.
+    Resources retain logical NVN stage 2 even though Metal executes a kernel.
+    """
+    code = open(j['code'], 'rb').read()
+    w2, w3, w4 = struct.unpack_from('<3I', code, 8)
+    nv, inv, topology = w4 & 4095, w2 >> 24, w3 >> 24 & 15
+    if topology not in (1, 6, 7) or not nv or not inv:
+        return 'geometry failed', 'invalid output declaration', None
+    nidx = nv * 2  # one restart after every emitted vertex, including incomplete strips
+    jobs = [{**j, 'key': key + '_g' + str(t), 'geometry_topology': t} for t in range(5)]
+    jp = os.path.join(w, key + '_geometry.json'); atomic_json(jp, jobs)
+    rx = subprocess.run([XLATE, jp, w], capture_output=True, text=True)
+    if rx.returncode: return 'geometry failed', rx.stderr.strip()[:160], None
+    base = None
+    for t, job in enumerate(jobs):
+        k = job['key']; spv = os.path.join(w, k + '.spv')
+        dmsh_geometry.flatten_sets(spv)
+        r = subprocess.run(['spirv-cross', *XC, spv], capture_output=True, text=True)
+        if r.returncode: return 'geometry MSL failed', r.stderr.strip()[:160], None
+        src = dmsh_geometry.topology_pointer(r.stdout)
+        res = json.load(open(os.path.join(w, k + '.res.json')))
+        refl = {'key': key + '_bg' + str(t), 'route': 'B', 'chain': CHAIN, 'dm_ctl': key,
+                'source': f'{loc[key][0]}#{loc[key][1]}', 'pc_paired': pc_of[key], 'stage': 'gs',
+                'entry': 'main0', 'msl_entry_kind': 'kernel', 'threadgroup_size': [32, 1, 1],
+                'geometry': {'vertices': nv, 'invocations': inv, 'output': topology,
+                             'stride': res['geometry_stride'], 'index_stride': nidx},
+                'c1': c1.hex(), 'notes': res['notes'], 'drv_reads': drv_reads(src), 'drv_unfillable': [],
+                'bindings': reflect(key, 'gs', src, res), 'vertex_inputs': [], 'raster': {}}
+        def save(suffix, text, meta):
+            path = os.path.join(out, key + suffix + '.metal'); open(path, 'w').write(text)
+            MTL.compile('mac', path, os.path.join(out, key + suffix + '.metallib'))
+            atomic_json(os.path.join(out, key + suffix + '.json'), dict(meta, chain_metal=MAC_RECORD))
+            if IOSOUT:
+                MTL.compile('ios', path, os.path.join(IOSOUT, key + suffix + '.metallib'))
+                atomic_json(os.path.join(IOSOUT, key + suffix + '.json'), dict(meta, chain_metal=ios_metal.RECORD))
+        save('_bg' + str(t), src, refl)
+        if t == 0:
+            base = {**refl, 'key': key + '_b'}
+            save('_b', src, base)
+            spv = os.path.join(w, k + '.reader.spv'); dmsh_geometry.flatten_sets(spv)
+            r = subprocess.run(['spirv-cross', *XC, spv], capture_output=True, text=True)
+            if r.returncode: return 'geometry reader failed', r.stderr.strip()[:160], None
+            reader = vs_layer(vs_init(r.stdout))
+            analysis, full = vs_depth(reader)
+            meta = {'key': key + '_bgv', 'route': 'B', 'chain': CHAIN, 'dm_ctl': key,
+                    'stage': 'vs', 'entry': 'main0', 'entry_z': 'main0_z', 'msl_entry_kind': 'vertex',
+                    'geometry_reader': True, 'geometry_output': topology, 'c1': '', 'notes': [], 'drv_reads': {}, 'drv_unfillable': [],
+                    'bindings': reflect(key, 'vs', analysis, {'sbufs': []}), 'vertex_inputs': [], 'raster': {}}
+            save('_bgv', full, meta)
+    return 'ok', '', base
+
+
 def finish(key, w, j, c1):
     name, k = loc[key]
     msl = os.path.join(w, key + '.metal'); lib = os.path.join(out, key + '_b.metallib')
-    if j['stage'] == 'gs':   # a geometry shader that only picks the layer is folded into the vertex shader by the host (see gs_slice); any other has no Metal counterpart and fails below
+    if j['stage'] == 'gs':   # keep the layer-only fold; other geometry runs as compute
         s = gs_slice(os.path.join(w, key + '.spv'))
         if s:
             refl = {'key': key + '_b', 'route': 'B', 'chain': CHAIN, 'dm_ctl': key, 'source': f'{name}#{k}', 'pc_paired': pc_of[key], 'stage': 'gs', 'gs_slice': s,
                     'notes': [], 'drv_reads': '', 'drv_unfillable': [], 'bindings': []}   # no library: the host draws with the vertex shader's layer entry points
             write_refl(key, refl)
             return 'ok', '', refl
+        try: return geometry_finish(key, w, j, c1)
+        except (ValueError, subprocess.CalledProcessError) as e: return 'geometry failed', str(e)[:160], None
     if j['stage'] == 'vs': spv_single_lane(os.path.join(w, key + '.spv'))   # one-lane subgroups (see spv_single_lane)
     r = subprocess.run(['spirv-cross', *XC, os.path.join(w, key + '.spv'), '--output', msl], capture_output=True, text=True)
     if r.returncode: return 'MSL failed', r.stderr.strip()[:160], None
@@ -755,7 +813,9 @@ def finish(key, w, j, c1):
         if pad: open(msl, 'w').write(src)
     lay = None
     if j['stage'] == 'vs':   # the two depth-convention entry points (see vs_depth); the analysis below uses the source with only main0
+        capture = dmsh_geometry.capture(src)
         src, full = vs_depth(src)
+        full += capture
         if j.get('tf'): full += vs_tf(src, j['tf'])   # transform feedback capture (see vs_tf)
         plain = full
         full, lay = vs_layer_entry(full)   # layer entry points for a slice-writing geometry shader (see vs_layer_entry)
@@ -776,7 +836,7 @@ def finish(key, w, j, c1):
     entry = re.search(r'(vertex|fragment|kernel)\s+\S+\s+(\w+)\s*\(', src)
     dr = drv_reads(src)
     refl = {'key': key + '_b', 'route': 'B', 'chain': CHAIN, 'dm_ctl': key, 'source': f'{name}#{k}', 'pc_paired': pc_of[key], 'stage': j['stage'],
-            'entry': entry.group(2) if entry else 'main0', 'msl_entry_kind': entry.group(1) if entry else None, **({'entry_z': 'main0_z'} if j['stage'] == 'vs' else {}),
+            'entry': entry.group(2) if entry else 'main0', 'msl_entry_kind': entry.group(1) if entry else None, **({'entry_z': 'main0_z', 'entry_geometry_capture': 'main0_gs'} if j['stage'] == 'vs' else {}),
             **({'entry_tf': 'main0_tf', 'tf_strides': [e[2] if e else 0 for e in j['tf']]} if j.get('tf') else {}),
             **({'entry_layer': 'main0_l', 'entry_layer_z': 'main0_lz', 'layer_from': lay} if lay else {}), **({'entry_u': 'main0_u'} if fu else {}),
             'threadgroup_size': j.get('local'), 'c1': c1.hex(), 'notes': res['notes'] + (['not in the answer table, queries use defaults'] if j.get('no_facts') else []), 'drv_reads': dr, 'drv_unfillable': unfillable(dr),

dm/gfx/tools/dmsh_chain.lock

@@ -1,8 +1,8 @@
 # Translation chain tool fingerprints (written by dmsh_chain.py --update-lock; regenerate after changing tools and commit it)
-dmsh_xlate.dll	34dc61131479cb40
+dmsh_xlate.dll	08d31903d666ea0c
-Ryujinx.Graphics.Shader.dll	d0b54941cc2237d9
+Ryujinx.Graphics.Shader.dll	eb315f0da558ca5a
-Ryujinx.Common.dll	78c38eb78e9e8b4c
+Ryujinx.Common.dll	b99db57e72c1fa16
-xlate_source	7202f939c87f92ea
+xlate_source	16bbef6c762573ee
 spirv-cross	d9dc197f061b42a5
 metal_macosx	Apple metal version 32023.921 (metalfe-32023.921.6)
 metal_iphoneos	Apple metal version 32023.921 (metalfe-32023.921.6)

dm/gfx/tools/dmsh_chain.py

@@ -14,7 +14,7 @@ DM_ROOT = os.environ.get('DM_ROOT') or os.path.expanduser('~/dm-work')   # Works
 HERE = os.path.dirname(os.path.abspath(__file__)); GFX = os.path.dirname(HERE); HOME = os.path.expanduser('~')
 sys.path.insert(0, os.path.join(GFX, '../games')); import game as _game   # game profiles
 GAME = _game.load()
-FILES = ['tools/dmsh_chain.py', 'tools/dmsh_build.py', 'tools/dmsh_fetch.py', 'tools/dmsh_vin.py', 'tools/dmsh_facts.py', 'tools/dmsh_xlate/Program.cs', 'tools/dmsh_xlate/dmsh_xlate.csproj', '../ios/tools/ios_metal.py', 'tools/dmsh_chain.lock']
+FILES = ['tools/dmsh_chain.py', 'tools/dmsh_build.py', 'tools/dmsh_fetch.py', 'tools/dmsh_vin.py', 'tools/dmsh_geometry.py', 'tools/dmsh_facts.py', 'tools/dmsh_xlate/Program.cs', 'tools/dmsh_xlate/dmsh_xlate.csproj', '../ios/tools/ios_metal.py', 'tools/dmsh_chain.lock']
 TABLES = ['shader_facts.tsv', 'shader_map.tsv', 'shader_map.tsv.conflicts.tsv', 'shader_pairs.tsv']   # in the game profile folder
 LOCK = os.path.join(HERE, 'dmsh_chain.lock'); HEADER = str(GAME.table('dmsh_chain.h'))
 # the translator defaults to the in-repo build that matches the lock file (a rebuilt dll has different fingerprints)

dm/gfx/tools/dmsh_geometry.py

@@ -0,0 +1,119 @@
+"""Offline geometry helpers. No game IDs or shader keys participate in this ABI.
+
+VS capture records are 144 words: position 0, point size 4, clip distances
+5..12, layer 13, viewport 14, reserved 15, generic attributes 16..143.
+This matches ResourceReservations.FillGeometryInputMap. Capture uses the native
+vertex fetch, including its instance divisors and format conversions.
+"""
+import re
+import struct
+
+WORDS = 144
+CAPTURE_SLOT = 13
+HELPERS = {'vertex_output': 0, 'geometry_vb_output': 1,
+           'geometry_ib_output': 2, 'vb_info': 3, 'trb_data': 4,
+           'vb_input': 1}
+
+
+def capture(src):
+    """Append a rasterizer-discard vertex entry to the unmodified-depth source."""
+    m = re.search(r'\bvertex main0_out main0\((.*?)\)\s*\{', src, re.S)
+    if not m:
+        raise ValueError('geometry capture: missing main0 vertex entry')
+    params = m.group(1)
+    body = src[m.end():]
+    if not body.rstrip().endswith('}'):
+        raise ValueError('geometry capture: unexpected function ending')
+    pro = ''
+    for semantic, name in [('vertex_id', 'dm_vid'), ('instance_id', 'dm_iid')]:
+        found = re.search(r'uint (\w+) \[\[' + semantic + r'\]\]', params)
+        if found:
+            pro += f'    uint {name} = {found.group(1)};\n'
+        else:
+            params += (', ' if params else '') + f'uint {name} [[{semantic}]]'
+    st = re.search(r'struct main0_out\s*\{(.*?)\};', src, re.S)
+    if not st:
+        raise ValueError('geometry capture: missing output struct')
+    stores = []
+    for line in st.group(1).splitlines():
+        line = line.strip()
+        if not line:
+            continue
+        v = re.fullmatch(r'(float4|uint4|int4) (\w+) \[\[user\(locn(\d+)\)\]\];', line)
+        if v:
+            if int(v[3]) >= 32:
+                raise ValueError('geometry capture: attribute outside Maxwell map')
+            stores += [f'dm_p[{16 + int(v[3]) * 4 + c}] = as_type<uint>(out.{v[2]}.{s});' for c, s in enumerate('xyzw')]
+            continue
+        v = re.fullmatch(r'float4 (\w+) \[\[position(?:, invariant)?\]\];', line)
+        if v:
+            stores += [f'dm_p[{c}] = as_type<uint>(out.{v[1]}.{s});' for c, s in enumerate('xyzw')]
+            continue
+        v = re.fullmatch(r'(float|uint|int) (\w+) \[\[(point_size|render_target_array_index|viewport_array_index)\]\];', line)
+        if v:
+            offset = {'point_size': 4, 'render_target_array_index': 13, 'viewport_array_index': 14}[v[3]]
+            stores.append(f'dm_p[{offset}] = as_type<uint>(out.{v[2]});')
+            continue
+        v = re.fullmatch(r'float (\w+)\[(\d+)\] \[\[clip_distance\]\];', line)
+        if v and int(v[2]) <= 8:
+            stores += [f'dm_p[{5 + c}] = as_type<uint>(out.{v[1]}[{c}]);' for c in range(int(v[2]))]
+            continue
+        raise ValueError('geometry capture: unknown output ' + line)
+    params += f', constant dm_geometry_capture& dm_gs [[buffer({CAPTURE_SLOT})]]'
+    pro += (f'    uint4 dm_ci = dm_gs.info; device uint* dm_capture = dm_gs.dst;\n'
+            f'    if (dm_vid - dm_ci.x >= dm_ci.y || dm_iid - dm_ci.z >= dm_ci.w) return;\n'
+            f'    device uint* dm_p = dm_capture + ((ulong(dm_iid - dm_ci.z) * dm_ci.y + dm_vid - dm_ci.x) * {WORDS});\n')
+    # Initialize missing attributes exactly as the normal vertex builder does.
+    epilogue = ('{ for (uint dm_k = 0; dm_k < 144; ++dm_k) dm_p[dm_k] = 0; '
+                'dm_p[3] = dm_p[4] = 0x3f800000u; '
+                'for (uint dm_k = 19; dm_k < 144; dm_k += 4) dm_p[dm_k] = 0x3f800000u; '
+                + ' '.join(stores) + ' return; }')
+    if 'return out;' not in body:
+        raise ValueError('geometry capture: missing return')
+    return '\nstruct dm_geometry_capture { device uint* dst; uint4 info; };\nvertex void main0_gs(' + params + ')\n{\n' + pro + body.replace('return out;', epilogue)
+
+
+def flatten_sets(path):
+    """Give auxiliary sets their own binding ranges in the ordinary set-0 ABI."""
+    data = open(path, 'rb').read()
+    w = list(struct.unpack('<%dI' % (len(data) // 4), data))
+    sets, bindings = {}, {}
+    i = 5
+    while i < len(w):
+        n, op = w[i] >> 16, w[i] & 65535
+        if not n or i + n > len(w):
+            raise ValueError('invalid SPIR-V instruction')
+        if op == 71 and n == 4:
+            if w[i + 2] == 34: sets[w[i + 1]] = (w[i + 3], i + 3)
+            if w[i + 2] == 33: bindings[w[i + 1]] = i + 3
+        i += n
+    for ident, (s, at) in sets.items():
+        if s:
+            if s not in (1, 2) or ident not in bindings:
+                raise ValueError('unknown geometry auxiliary descriptor set')
+            w[bindings[ident]] += 10000 if s == 1 else 11000
+            w[at] = 0
+    open(path, 'wb').write(struct.pack('<%dI' % len(w), *w))
+
+
+def topology_pointer(src):
+    """Topology remap uses uint words; avoid allocating a texture view per draw."""
+    src, n = re.subn(r'texture_buffer<float> trb_data', 'device const uint* trb_data', src)
+    if n != 1:
+        raise ValueError('missing geometry topology buffer')
+    # Match the balanced read argument rather than relying on expression shape.
+    prefix = 'spvDescriptorSet0.trb_data.read('
+    while prefix in src:
+        start = src.index(prefix)
+        at, depth = start + len(prefix), 1
+        end = at
+        while depth:
+            if src[end] == '(': depth += 1
+            if src[end] == ')': depth -= 1
+            end += 1
+        if src[end:end + 2] != '.x':
+            raise ValueError('unexpected topology read')
+        src = src[:start] + 'as_type<float>(spvDescriptorSet0.trb_data[' + src[at:end - 1] + '])' + src[end + 2:]
+    if 'trb_data.read' in src:
+        raise ValueError('unconverted geometry topology read')
+    return src

dm/gfx/tools/dmsh_xlate/Program.cs

@@ -30,7 +30,7 @@ static class SwshXlate
     {
         readonly byte[] _code; readonly JsonObject _job; int _binding = 1; public readonly List<string> Notes = new();
         readonly byte[] _c1;
-        public Accessor(byte[] code, JsonObject job, Dictionary<int, SamplerType> imageTypes = null) { _code = code; _job = job; _c1 = job["c1"] is JsonNode h ? Convert.FromHexString((string)h) : Array.Empty<byte>(); _imageTypes = imageTypes; }
+        public Accessor(byte[] code, JsonObject job, Dictionary<int, SamplerType> imageTypes = null) { if (job["geometry_topology"] != null) _binding = 40; _code = code; _job = job; _c1 = job["c1"] is JsonNode h ? Convert.FromHexString((string)h) : Array.Empty<byte>(); _imageTypes = imageTypes; }
         // A texture query on an image handle (TXQ on c2 0x120 + 8 r: Unreal's GetDimensions on the RWTexture3D a compute shader writes) asks the type of that handle, which a table built
         // without the game's PC shaders does not answer, and the 2D default took a 3D image's depth from its width. The shader's own surface instructions give the image's type: the
         // first translation records that a query went unanswered (ImageQueryGuessed), and Main translates again with the types of the images it found, by c2 handle (ImageTypesOf)
@@ -106,12 +106,13 @@ static class SwshXlate
         public int QueryComputeLocalSizeZ() => Local(2);
         public int QueryComputeSharedMemorySize() => _job["shared"] is JsonNode n ? (int)n : 0;   // the shared memory size in the compute shader's control section (shader_split.py)
         static AttributeType Attr(string s) => s switch { "sint" => AttributeType.Sint, "uint" => AttributeType.Uint, _ => AttributeType.Float };
+        public bool QueryGeometryInputFixed() => _job["geometry_topology"] != null;
         public GpuGraphicsState QueryGraphicsState()
         {
             Array32<AttributeType> attr = default; Array8<AttributeType> fout = default;
             if (_job["attr"] is JsonObject a) foreach (var kv in a) attr[int.Parse(kv.Key)] = Attr((string)kv.Value);
             if (_job["fout"] is JsonObject f) foreach (var kv in f) fout[int.Parse(kv.Key)] = Attr((string)kv.Value);
-            return new GpuGraphicsState(false, InputTopology.Triangles, false, TessPatchType.Triangles, TessSpacing.EqualSpacing,
+            return new GpuGraphicsState(false, _job["geometry_topology"] is JsonNode gt ? (InputTopology)(int)gt : InputTopology.Triangles, false, TessPatchType.Triangles, TessSpacing.EqualSpacing,
                 false, false, false, false /* depth not rewritten: dmsh_build.py adds the GL entry point (see the file header) */, false, 1f, AlphaTestOp.Always, 0f, attr, false, fout, false,
                 false /* no Y flip */, true /* origin upper-left */, false);
         }
@@ -157,25 +158,33 @@ static class SwshXlate
                 // Job debug = false: turn off DebugMode (DebugMode skips the AST optimizer and the output is full of temporaries; semantics unchanged)
                 var flags = (job["debug"] is JsonNode dbg && !(bool)dbg ? 0 : TranslationFlags.DebugMode) | ((string)job["stage"] == "cs" ? TranslationFlags.Compute : 0);
                 var ctx = Translator.CreateContext(0, acc, new TranslationOptions(TargetLanguage.Spirv, TargetApi.Vulkan, flags));
-                ShaderProgram p = ctx.Translate();
+                bool geometry = job["geometry_topology"] != null;
+                ShaderProgram p = ctx.Translate(geometry);
                 Dictionary<int, SamplerType> imageTypes = null;
                 if (acc.ImageQueryGuessed && (imageTypes = Accessor.ImageTypesOf(p)).Count > 0)   // a query on an image handle went unanswered: again with the images' types (see Accessor)
                 {
                     acc = new Accessor(code, job, imageTypes);
                     ctx = Translator.CreateContext(0, acc, new TranslationOptions(TargetLanguage.Spirv, TargetApi.Vulkan, flags));
-                    p = ctx.Translate();
+                    p = ctx.Translate(geometry);
                 }
                 if (UsesFloat64(p.BinaryCode))   // Tears of the Kingdom 1.4.3: 30 shaders used double and failed in the Metal compiler
                 {
                     job["softfloat64"] = true; acc = new Accessor(code, job, imageTypes);
                     ctx = Translator.CreateContext(0, acc, new TranslationOptions(TargetLanguage.Spirv, TargetApi.Vulkan, flags));
-                    p = ctx.Translate(); acc.Notes.Add("FP64 lowered to FP32 (DoubleToFloat)");
+                    p = ctx.Translate(geometry); acc.Notes.Add("FP64 lowered to FP32 (DoubleToFloat)");
+                }
+                if (geometry)
+                {
+                    // A fresh context: translation mutates the decoded program.
+                    var reader = Translator.CreateContext(0, new Accessor(code, job, imageTypes), new TranslationOptions(TargetLanguage.Spirv, TargetApi.Vulkan, flags));
+                    File.WriteAllBytes(Path.Combine(outDir, key + ".reader.spv"), reader.GenerateVertexPassthroughForCompute().BinaryCode);
                 }
                 File.WriteAllBytes(Path.Combine(outDir, key + ".spv"), VertexCounterLowering.Apply(code, (string)job["stage"], p.BinaryCode, acc.Notes));
                 var res = new JsonObject
                 {
                     ["key"] = key, ["stage"] = (string)job["stage"], ["notes"] = new JsonArray(acc.Notes.ConvertAll(n => (JsonNode)n).ToArray()),
                     ["cbufs"] = new JsonArray(), ["sbufs"] = new JsonArray(), ["textures"] = new JsonArray(), ["images"] = new JsonArray(),
+                    ["geometry_stride"] = geometry ? ctx.GetResourceReservations().OutputSizePerInvocation : 0,
                     ["uses_frag_coord"] = p.Info.UsesFragCoord, ["frag_out_map"] = p.Info.FragmentOutputMap,
                 };
                 foreach (var b in p.Info.CBuffers) res["cbufs"].AsArray().Add(new JsonObject { ["binding"] = b.Binding, ["slot"] = b.Slot });

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/IGpuAccessor.cs

@@ -8,6 +8,9 @@ namespace Ryujinx.Graphics.Shader
     /// </summary>
     public interface IGpuAccessor : ILogger
     {
+        // Fixed VS capture ABI used by the offline Metal geometry path.
+        bool QueryGeometryInputFixed() => false;
+
         /// <summary>
         /// Reads data from the constant buffer 1.
         /// </summary>

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/EmitterContext.cs

@@ -156,6 +156,20 @@ namespace Ryujinx.Graphics.Shader.Translation
 
                     this.Store(StorageKind.LocalMemory, ResourceManager.LocalGeometryOutputVertexCountMemoryId, Const(0));
                     this.Store(StorageKind.LocalMemory, ResourceManager.LocalGeometryOutputIndexCountMemoryId, Const(0));
+                    if (ResourceManager.Reservations.GeometryInputFixed)
+                    {
+                        this.Store(StorageKind.LocalMemory, ResourceManager.LocalGeometryPrimitiveOpenMemoryId, Const(0));
+                        for (int i = 0; i < ResourceManager.Reservations.OutputSizePerInvocation; i++)
+                            this.Store(StorageKind.LocalMemory, ResourceManager.LocalVertexDataMemoryId, Const(i), ConstF(0));
+                        void DefaultOne(IoVariable variable, int location, int component)
+                        {
+                            if (ResourceManager.Reservations.TryGetOffset(StorageKind.Output, variable, location, component, out int offset))
+                                this.Store(StorageKind.LocalMemory, ResourceManager.LocalVertexDataMemoryId, Const(offset), ConstF(1));
+                        }
+                        DefaultOne(IoVariable.Position, 0, 3);
+                        DefaultOne(IoVariable.PointSize, 0, 0);
+                        for (int i = 0; i < 32; i++) DefaultOne(IoVariable.UserDefined, i, 3);
+                    }
                 }
             }
         }
@@ -560,7 +574,7 @@ namespace Ryujinx.Graphics.Shader.Translation
 
                     Operand writtenIndices = this.Load(StorageKind.LocalMemory, ResourceManager.LocalGeometryOutputIndexCountMemoryId);
 
-                    int maxIndicesPerPrimitiveInvocation = TranslatorContext.Definitions.GetGeometryOutputIndexBufferStridePerInstance();
+                    int maxIndicesPerPrimitiveInvocation = TranslatorContext.Definitions.GetGeometryOutputIndexBufferStridePerInstance(ResourceManager.Reservations.GeometryInputFixed);
                     int maxIndicesPerPrimitive = maxIndicesPerPrimitiveInvocation * TranslatorContext.Definitions.ThreadsPerInputPrimitive;
 
                     this.BranchIfTrue(lblExit, this.ICompareGreaterOrEqualUnsigned(writtenIndices, Const(maxIndicesPerPrimitiveInvocation)));
@@ -570,7 +584,7 @@ namespace Ryujinx.Graphics.Shader.Translation
                     Operand primitiveIndex = this.Load(StorageKind.Input, IoVariable.GlobalId, Const(0));
                     Operand instanceIndex = this.Load(StorageKind.Input, IoVariable.GlobalId, Const(1));
                     Operand invocationId = this.Load(StorageKind.Input, IoVariable.GlobalId, Const(2));
-                    Operand vertexCount = this.Load(StorageKind.ConstantBuffer, vertexInfoCbBinding, Const((int)VertexInfoBufferField.VertexCounts), Const(0));
+                    Operand vertexCount = this.Load(StorageKind.ConstantBuffer, vertexInfoCbBinding, Const((int)(ResourceManager.Reservations.GeometryInputFixed ? VertexInfoBufferField.GeometryCounts : VertexInfoBufferField.VertexCounts)), Const(0));
                     Operand primitiveId = this.IAdd(this.IMultiply(instanceIndex, vertexCount), primitiveIndex);
                     Operand ibOffset = this.IMultiply(primitiveId, Const(maxIndicesPerPrimitive));
                     ibOffset = this.IAdd(ibOffset, this.IMultiply(invocationId, Const(maxIndicesPerPrimitiveInvocation)));

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/EmitterContextInsts.cs

@@ -233,11 +233,31 @@ namespace Ryujinx.Graphics.Shader.Translation
 
         public static Operand EmitVertex(this EmitterContext context)
         {
+            if (context.VertexAsCompute && context.ResourceManager.Reservations.GeometryInputFixed)
+            {
+                Operand end = Label();
+                Operand count = context.Load(StorageKind.LocalMemory, context.ResourceManager.LocalGeometryOutputVertexCountMemoryId);
+                context.BranchIfTrue(end, context.ICompareGreaterOrEqualUnsigned(count, Const(context.TranslatorContext.Definitions.MaxOutputVertices)));
+                context.Store(StorageKind.LocalMemory, context.ResourceManager.LocalGeometryPrimitiveOpenMemoryId, Const(1));
+                context.Add(Instruction.EmitVertex);
+                context.MarkLabel(end);
+                return null;
+            }
             return context.Add(Instruction.EmitVertex);
         }
 
         public static Operand EndPrimitive(this EmitterContext context)
         {
+            if (context.VertexAsCompute && context.ResourceManager.Reservations.GeometryInputFixed)
+            {
+                Operand end = Label();
+                Operand open = context.Load(StorageKind.LocalMemory, context.ResourceManager.LocalGeometryPrimitiveOpenMemoryId);
+                context.BranchIfTrue(end, context.ICompareEqual(open, Const(0)));
+                context.Store(StorageKind.LocalMemory, context.ResourceManager.LocalGeometryPrimitiveOpenMemoryId, Const(0));
+                context.Add(Instruction.EndPrimitive);
+                context.MarkLabel(end);
+                return null;
+            }
             return context.Add(Instruction.EndPrimitive);
         }
 

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ResourceManager.cs

@@ -52,6 +52,7 @@ namespace Ryujinx.Graphics.Shader.Translation
         public int LocalVertexIndexInstanceRateMemoryId { get; private set; }
         public int LocalGeometryOutputVertexCountMemoryId { get; private set; }
         public int LocalGeometryOutputIndexCountMemoryId { get; private set; }
+        public int LocalGeometryPrimitiveOpenMemoryId { get; private set; }
 
         public ShaderProperties Properties { get; }
 
@@ -137,6 +138,8 @@ namespace Ryujinx.Graphics.Shader.Translation
 
                 LocalGeometryOutputVertexCountMemoryId = AddMemoryDefinition("local_geometry_output_vertex", AggregateType.U32);
                 LocalGeometryOutputIndexCountMemoryId = AddMemoryDefinition("local_geometry_output_index", AggregateType.U32);
+                if (Reservations.GeometryInputFixed)
+                    LocalGeometryPrimitiveOpenMemoryId = AddMemoryDefinition("local_geometry_primitive_open", AggregateType.U32);
             }
         }
 

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ResourceReservations.cs

@@ -24,6 +24,7 @@ namespace Ryujinx.Graphics.Shader.Translation
         public int ReservedStorageBuffers { get; }
         public int ReservedTextures { get; }
         public int ReservedImages { get; }
+        public bool GeometryInputFixed { get; }
         public int InputSizePerInvocation { get; }
         public int OutputSizePerInvocation { get; }
         public int OutputSizeInBytesPerInvocation => OutputSizePerInvocation * sizeof(uint);
@@ -81,16 +82,30 @@ namespace Ryujinx.Graphics.Shader.Translation
             if (vertexAsCompute)
             {
                 _offsets = new();
+                GeometryInputFixed = gpuAccessor.QueryGeometryInputFixed();
 
                 if (vacInput.HasValue)
                 {
-                    InputSizePerInvocation = FillIoOffsetMap(gpuAccessor, StorageKind.Input, vacInput.Value);
+                    InputSizePerInvocation = GeometryInputFixed ? FillGeometryInputMap() : FillIoOffsetMap(gpuAccessor, StorageKind.Input, vacInput.Value);
                 }
 
                 OutputSizePerInvocation = FillIoOffsetMap(gpuAccessor, StorageKind.Output, vacOutput);
             }
         }
 
+        // Matches dmsh_geometry.py: independent stages meet through the Maxwell attribute map.
+        private int FillGeometryInputMap()
+        {
+            for (int c = 0; c < 4; c++) _offsets.Add(new(StorageKind.Input, IoVariable.Position, 0, c), c);
+            _offsets.Add(new(StorageKind.Input, IoVariable.PointSize), 4);
+            for (int c = 0; c < 8; c++) _offsets.Add(new(StorageKind.Input, IoVariable.ClipDistance, 0, c), 5 + c);
+            _offsets.Add(new(StorageKind.Input, IoVariable.Layer), 13);
+            _offsets.Add(new(StorageKind.Input, IoVariable.ViewportIndex), 14);
+            for (int a = 0; a < 32; a++) for (int c = 0; c < 4; c++)
+                _offsets.Add(new(StorageKind.Input, IoVariable.UserDefined, a, c), 16 + 4 * a + c);
+            return 144;
+        }
+
         private int FillIoOffsetMap(IGpuAccessor gpuAccessor, StorageKind storageKind, IoUsage vacUsage)
         {
             int offset = 0;

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ShaderDefinitions.cs

@@ -337,8 +337,12 @@ namespace Ryujinx.Graphics.Shader.Translation
             return _graphicsState.AttributeTypes[location].HasFlag(AttributeType.PackedRgb10A2Signed);
         }
 
-        public int GetGeometryOutputIndexBufferStridePerInstance()
+        public int GetGeometryOutputIndexBufferStridePerInstance(bool fixedInput = false)
         {
+            // Even a strip with fewer than three vertices can be ended. Empty
+            // ends are ignored by the fixed-input emitter, so at most one
+            // restart per emitted vertex is needed.
+            if (fixedInput) return MaxOutputVertices * 2;
             return MaxOutputVertices + OutputTopology switch
             {
                 OutputTopology.LineStrip => MaxOutputVertices / 2,
@@ -347,9 +351,9 @@ namespace Ryujinx.Graphics.Shader.Translation
             };
         }
 
-        public int GetGeometryOutputIndexBufferStride()
+        public int GetGeometryOutputIndexBufferStride(bool fixedInput = false)
         {
-            return GetGeometryOutputIndexBufferStridePerInstance() * ThreadsPerInputPrimitive;
+            return GetGeometryOutputIndexBufferStridePerInstance(fixedInput) * ThreadsPerInputPrimitive;
         }
     }
 }

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/Transforms/GeometryToCompute.cs

@@ -142,7 +142,7 @@ namespace Ryujinx.Graphics.Shader.Translation.Transforms
             Operand baseIndexOffset = GenerateBaseOffset(
                 resourceManager,
                 node,
-                definitions.GetGeometryOutputIndexBufferStride(),
+                definitions.GetGeometryOutputIndexBufferStride(resourceManager.Reservations.GeometryInputFixed),
                 definitions.ThreadsPerInputPrimitive);
             Operand outputBaseIndex = Local();
             node.List.AddBefore(node, new Operation(Instruction.Add, outputBaseIndex, new[] { baseIndexOffset, outputPrimIndex }));
@@ -197,7 +197,7 @@ namespace Ryujinx.Graphics.Shader.Translation.Transforms
             Operand baseIndexOffset = GenerateBaseOffset(
                 resourceManager,
                 node,
-                definitions.GetGeometryOutputIndexBufferStride(),
+                definitions.GetGeometryOutputIndexBufferStride(resourceManager.Reservations.GeometryInputFixed),
                 definitions.ThreadsPerInputPrimitive);
             Operand outputBaseIndex = Local();
             node.List.AddBefore(node, new Operation(Instruction.Add, outputBaseIndex, new[] { baseIndexOffset, outputPrimIndex }));
@@ -212,7 +212,7 @@ namespace Ryujinx.Graphics.Shader.Translation.Transforms
         private static Operand GenerateBaseOffset(ResourceManager resourceManager, LinkedListNode<INode> node, int stride, int threadsPerInputPrimitive)
         {
             Operand primitiveId = Local();
-            GeneratePrimitiveId(resourceManager, node, primitiveId);
+            GeneratePrimitiveId(resourceManager, node, primitiveId, true);
 
             Operand baseOffset = Local();
             node.List.AddBefore(node, new Operation(Instruction.Multiply, baseOffset, new[] { primitiveId, Const(stride) }));
@@ -303,8 +303,10 @@ namespace Ryujinx.Graphics.Shader.Translation.Transforms
             return vertexElemOffset;
         }
 
-        private static LinkedListNode<INode> GeneratePrimitiveId(ResourceManager resourceManager, LinkedListNode<INode> node, Operand dest)
+        private static LinkedListNode<INode> GeneratePrimitiveId(ResourceManager resourceManager, LinkedListNode<INode> node, Operand dest, bool outputIndex = false)
         {
+            if (resourceManager.Reservations.GeometryInputFixed && !outputIndex)
+                return node.List.AddBefore(node, new Operation(Instruction.Load, StorageKind.Input, dest, new[] { Const((int)IoVariable.GlobalId), Const(0) }));
             int vertexInfoCbBinding = resourceManager.Reservations.VertexInfoConstantBufferBinding;
 
             Operand vertexCount = Local();
@@ -312,7 +314,7 @@ namespace Ryujinx.Graphics.Shader.Translation.Transforms
                 Instruction.Load,
                 StorageKind.ConstantBuffer,
                 vertexCount,
-                new[] { Const(vertexInfoCbBinding), Const((int)VertexInfoBufferField.VertexCounts), Const(0) }));
+                new[] { Const(vertexInfoCbBinding), Const((int)(resourceManager.Reservations.GeometryInputFixed ? VertexInfoBufferField.GeometryCounts : VertexInfoBufferField.VertexCounts)), Const(0) }));
 
             Operand vertexIndex = Local();
             node.List.AddBefore(node, new Operation(

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/TranslatorContext.cs

@@ -460,7 +460,7 @@ namespace Ryujinx.Graphics.Shader.Translation
         {
             if (vertexAsCompute)
             {
-                return new ShaderDefinitions(ShaderStage.Compute, 32, 32, 1);
+                return new ShaderDefinitions(ShaderStage.Compute, 32, GpuAccessor.QueryGeometryInputFixed() ? 1 : 32, 1);
             }
             else
             {

docs/development.md

@@ -592,7 +592,7 @@ The game's own work and its frame rate stay as they are; the host avoids work of
 **AI: THE SOMNIUM FILES (`010089B00D09C000`)**
 - **Status:** brought up on the Mac: the Spike Chunsoft logo movie, the title (its background movies), a new game (language, auto save notice), the opening at the merry-go-round and the first investigation with its conversation tutorial, both sticks; loading the autosave from the title. 60 fps (present interval 1), about 2.27 GB resident, flat over 5 minutes; audio through FMOD's AudioOut. Playable on an iPhone.
 - **Engine:** Unity 2017.4.17f1 IL2CPP with FMOD on SDK 6.4.2, update 1.0.1 (v65536). `subsdk0` is the movie module; the 168 movies (`StreamingAssets/Movie`, H.264 and AAC in MP4) play through Unity's video player, NV12 into CPU buffers copied to an R8 + RG8 texture pair each frame.
-- **Shaders:** 13,057 (6,572 vertex, 6,482 pixel, 3 geometry) in the AssetBundles of `StreamingAssets` (Cuphead's Unity 2017.4 layout) and `main`. The 3 geometry shaders (Unity's DX11 depth of field, a wireframe effect) do not build; none was drawn. The ComputeShader objects of `Data/resources.assets` have no type tree (red line in the split, as in Cuphead); no dispatch has missed one so far.
+- **Shaders:** 13,057 (6,572 vertex, 6,482 pixel, 3 geometry) in the AssetBundles of `StreamingAssets` (Cuphead's Unity 2017.4 layout) and `main`. All build for Mac and iOS. The 3 geometry shaders (Unity's DX11 depth of field, a wireframe effect) use the general native-VS-capture / GS-compute path; GPU readback tests check their output, but none was drawn in the opening observed so far. See [geometry support and limits](geometry-metal-research.md). The ComputeShader objects of `Data/resources.assets` have no type tree (red line in the split, as in Cuphead); no dispatch has missed one so far.
 - **Host changes:** an H.264 drain held up by a full set of output slots (the logo movie never ended, [03](agents/03-bringup.md)), and the per-stage composite programs restored after the precompile prediction ([04](agents/04-graphics.md)); both general.
 - **Memory:** no system resource; Unity sets a 3.11 GB heap (pool − used) in the 4 GB layout, `mem_used` 104 MB for the rewritten modules.
 - **Input:** controller only, docked (1280x720 in the cropped 1920x1080 window, by the game's choice), with the 3D pad layout: the investigation moves its cursor with the left stick and the view with the right one.

docs/geometry-metal-research.md

@@ -0,0 +1,94 @@
+# Geometry shader support on Metal 4
+
+Research and implementation checked on 2026-10-23 against Apple documentation, the installed Xcode SDK, this worktree, and AI: THE SOMNIUM FILES. The implementation now has GPU readback tests; the initial compilation probe described below predates it.
+
+## Implemented compute path
+
+The offline builder keeps the existing slice-writer fold and builds other geometry shaders as compute kernels, with five input-topology variants and an output-reader vertex function. Selection uses shader metadata and the NVN draw topology, with no game IDs or shader-key rules. Each vertex library also contains a rasterizer-discard capture entry. This runs native vertex fetching, including instance divisors and the host's existing format conversions, and writes a fixed 144-word record: position, point size, clip distances, layer, viewport and 32 generic attributes. Separately bound VS and GS programs therefore need no precomputed pair table.
+
+Replay captures the VS once per unique vertex and instance, dispatches the GS using logical NVN stage 2 resources, compacts emitted strips into ordered lists, and draws the output through the original fragment shader. Depth conversion occurs in the final vertex reader. Each input primitive, instance and GS invocation has its own bounded output region. Empty `EndPrimitive` operations are ignored, short strips have room for their restart markers, and output beyond the declared vertex maximum is bounded. The compactor preserves strip winding and the last vertex of each primitive.
+
+`dm/gfx/tests/t_geometry.m` checks primitive assembly, strip compaction, native capture, indexed and indirect replay, restart, base vertex/instance, multiple GS invocations, and a draw exceeding the previous 4 MiB arena allocation. With the external AI shader cache passed as its argument, it additionally checks the generated capture entry and all three translated GS kernels against numeric expectations. On Apple M4 Pro, **498 checks passed with Metal API validation enabled**. The vertex-binding regression passed 43 checks and the draw-skip regression passed 38. All **13,057 AI shaders** built for Mac and iOS, including the previously failing three geometry shaders; the profile's chain header, constant-buffer guards and extents were regenerated and the Mac host rebuilt. The changed Metal replay translation unit also compiled against the iOS SDK.
+
+A 90-second Mac run reached the title screen at 60 fps with no `[RED]` diagnostics. Its final `CRASH sig=14` is the test harness's time limit, not a game failure. The run used a separate test save folder. Evidence: `run/geometry-boot.log`, `geometry-build/boot-3000.png`, `geometry-build/full-build.log` and `geometry-build/test-results.log` under the game's work folder.
+
+Current limits are explicit: GS transform feedback and tessellation combinations are rejected; triangle-strip adjacency, quad strips and polygons are not assembled. Geometry storage-buffer addresses in the NVN driver table remain uncharacterized (addresses carried in guest uniform buffers use the existing binding path). Provoking-vertex mode, GS-written primitive IDs, viewport-index outputs and multiple output streams are not validated. Indexed and indirect draws wait for a coherent CPU snapshot for assembly; compaction is one ordered GPU thread per draw. Scratch is capped at 256 MiB and input counts at 4,194,304, with a red diagnostic instead of truncation. Other games and iOS device execution are unvalidated. AI's opening did not previously issue these GS draws, so a successful opening alone cannot establish their in-game appearance or performance.
+
+The external build and test artifacts are under `010089B00D09C000-work/geometry-build/` beside the game's NSPs. To repeat the GPU tests, build `t_geometry` with `GAME=010089B00D09C000`, then run it with the current `cache/b-<chain>` directory as its optional first argument. No extracted shader code is stored in the repository.
+
+## Conclusion
+
+General geometry shader emulation is feasible. A mesh pipeline can keep expansion inside a render pass; compute expansion followed by an ordinary draw is a more explicit implementation with temporary buffers. Neither is a switch that makes the existing geometry-stage SPIR-V compile unchanged. The mesh route replaces the vertex stage too, and therefore needs the vertex and geometry programs connected through a new translation ABI. Keep the existing proven slice-writer fold as an optimization. These are engineering conclusions from the interfaces and current implementation below.
+
+**The implementation uses the compute route.** The vendored translator already lowers GS to compute and generates a vertex function to read its output. Native VS capture supplies an independent-stage ABI. The original isolated probe established compilation feasibility; the GPU validation above now establishes the tested output semantics. Mesh shaders remain a potential optimization.
+
+## AI: THE SOMNIUM FILES evidence
+
+The registered profile contains exactly three geometry keys. Existing build logs fail these with an `unknown` MSL stage declaration. The profile maps the following effects; candidate vertex keys come from adjacent entries of the same Unity shader, not a captured runtime binding. [Shader inventory](../dm/games/010089B00D09C000/shader_map.tsv), [existing bring-up observations](development.md).
+
+| GS key | Effect | Body and declared output | Candidate VS |
+| --- | --- | --- | --- |
+| `1b378c5c014e3a32` | `Hidden/DepthOfField/BokehSplatting` | Reads input vertex 0, expands its screen position to a four-vertex triangle strip; maximum 4, one invocation | `ba2512f7f6973541` |
+| `d57cdb168fdf6927` | `Hidden/Dof/DX11Dof` | Same expansion math as the first, with input varying locations 0 and 1 exchanged; maximum 4, one invocation | `771d597a6eb545aa` |
+| `2341c1be4fcbf690` | `UCLA Game Lab/Wireframe/Single-Sided` | Reads all three positions, computes screen-space edge distances, emits three vertices in input order; maximum 3, one invocation | `ce9cbd7dca9fda74` |
+
+The output maxima and invocation counts were read from shader headers. Input topology is supplied by the translator accessor: the probe selected points for the bokeh shaders and triangles for wireframe, consistent with their bodies. These are not independently captured NVN draw topologies. Wireframe cannot be implemented by simply dropping the GS: its fragment inputs depend on the other two vertices. The bokeh shaders actually expand geometry and also fail the existing layer-only fold.
+
+Probe artifacts remain beside the game, outside the repository:
+
+`/Volumes/T7/nsp/AI THE SOMNIUM FILES/010089B00D09C000-work/geometry-probe/`
+
+This contains `Program.cs`, `Probe.csproj`, `translation.log`, `compile-results.json`, native GS GLSL, compute/passthrough SPIR-V and MSL, and Mac/iOS libraries. The small C# probe references the pinned binaries in this worktree's `dmsh_xlate_bin`; it does not rebuild or replace the production translator. It reads package records directly and uses float attribute defaults and deliberately separate helper/resource bindings. It is not the production reflection or binding ABI.
+
+Validation: nine generated programs (three VS kernels, three GS kernels, three output-reader vertex functions) compiled and linked for each platform: **18 successful metallibs, zero failed final compile/link steps**. The final probe uses fresh translator contexts for each translation. Reusing a context first for native GS and then compute produced a translator warning and different output; that preliminary result was replaced. SPIRV-Cross uses the repository's MSL 2.3/native texture-buffer flags; Metal compilation uses its Mac/iOS flags. No pipeline creation, dispatch, pixel comparison, or iPhone execution was performed. Existing bring-up documentation says none of these GS effects was drawn in the tested opening; compiling them does not establish a visible improvement there.
+
+## Existing implementation to reuse
+
+[`TranslatorContext.Translate(asCompute: true)`](../dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/TranslatorContext.cs) activates [`VertexToCompute`](../dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/Transforms/VertexToCompute.cs) and [`GeometryToCompute`](../dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/Transforms/GeometryToCompute.cs). Connect the stages with `SetNextStage` and `SetVertexOutputMapForGeometryAsCompute`; `GenerateVertexPassthroughForCompute` supplies the final vertex reader. The current [offline wrapper](../dm/gfx/tools/dmsh_xlate/Program.cs) calls only ordinary `Translate()`.
+
+The compute lowering already writes output to fixed regions per input primitive and GS invocation. `EmitVertex` stores attributes and indices; `EndPrimitive` stores restart markers, and the emitter fills unused index positions with restart markers. This preserves an ordered layout without a global atomic append cursor. It still needs host allocation and output-topology handling. [`ResourceReservations`](../dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/ResourceReservations.cs), [`VertexInfoBuffer`](../dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/VertexInfoBuffer.cs), and [`EmitterContext`](../dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Translation/EmitterContext.cs) define the auxiliary binding/layout contract. The default compute local size is 32×32×1; validate pipeline limits and occupancy before retaining it.
+
+## Original implementation plan and remaining coverage
+
+1. Add explicit offline variants and reflection for VS-as-compute, GS-as-compute, and the output reader. Include input topology, invocation count, maximum output, interface strides, helper bindings, and logical NVN stage in metadata. Solve VS/GS interface selection for separately bound programs: the current build translates stages independently, while the existing GS lowering consumes its VS output map. A canonical interface or offline compatible-pair variants are design choices; runtime compilation is not available.
+2. Extend host resource encoding for logical geometry stage 2 while executing a Metal compute function. `encode_args` already accepts an NVN stage, but `c0_ssbo_base` currently supports only VS, FS, and CS. Do not accidentally source GS uniforms from compute stage 5. Supply helper buffers and vertex-format decoding, index/topology remapping, base vertex/instance and draw counts.
+3. Add a draw path: finish preceding conflicting render work, execute VS into an intermediate stream, execute GS into bounded output/index buffers, then draw through the output reader with the original fragment shader. `ixc_run` provides an existing compute-before-draw integration example. Preserve original order and restart/winding semantics, apply clip-depth conversion only at the final rasterizing stage, and account for indirect counts without a CPU readback. Bound multiplication/allocation and split large work without silently truncating geometry.
+4. Integrate resource lifetimes, barriers, dirty-buffer tracking, deferred uniform updates, layered attachments, program/pipeline caches, and prewarming. Existing `tf_draw` captures VS outputs only; a GS path must capture the final pre-rasterization stage when transform feedback is enabled. Keep the layer-only fast path.
+5. Validate GPU output against synthetic point expansion and triangle edge-distance cases before testing an actual effect in game. Cover zero/variable emission, multiple strips, indexed/instanced/indirect draws, adjacency, order-sensitive blending, layer writes and resource writes as support expands. Use Big Brain Academy's documented, actively drawn line-renderer GS as an additional end-to-end target; AI's opening is not a sufficient visual test. Run the full required shader-chain rebuild procedure for every affected profile after production changes.
+
+The implemented subset and measured validation are recorded above. GPU-only indirect assembly, parallel compaction, GS transform feedback, additional topology coverage, and validation on further games remain follow-up work.
+
+## Verified platform capabilities
+
+Apple lists mesh shading for Metal 3 and 4 starting at Apple7 (A14 and M1), but indirect mesh arguments and mesh draws in indirect command buffers start at Apple9 (A17 Pro and M3). Ordinary indirect draw/dispatch arguments start at Apple3. Consequently an indirect mesh implementation must not assume that all Metal 4 devices support it. The table gives a 16,384-byte payload maximum, reduced by up to 32 bytes for specified built-ins/debugging, and a 4 GB per-draw payload/geometry limit on Apple7/8. [Apple feature tables, pages 2, 4–5, 9–10](https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf).
+
+Mesh shaders replace the vertex stage with object and mesh stages, and feed rasterization directly. Traditional and mesh draws can share a render encoder. Apple's explanation describes point, line, or triangle output; 256 vertices, 512 primitives, and 16 KB total mesh output per mesh are the documented baseline limits. These constrain one-to-one mapping of guest GS invocations with large output declarations. [Apple WWDC22: Transform your geometry with Metal mesh shaders](https://developer.apple.com/videos/play/wwdc2022/10162/).
+
+The installed SDK confirms `MTL4MeshRenderPipelineDescriptor`, available on macOS/iOS 26, with object, mesh, and fragment function descriptors, payload length, and threadgroup limits. `MTL4RenderCommandEncoder` has direct and indirect mesh dispatch methods; its object stage is optional. Its indirect method consumes `MTLDispatchThreadgroupsIndirectArguments` at a 4-byte-aligned GPU address, and the buffer must be resident. Sources: installed `Metal.framework/Headers/MTL4MeshRenderPipeline.h` and `MTL4RenderCommandEncoder.h` under `xcrun --show-sdk-path`; corresponding [mesh descriptor](https://developer.apple.com/documentation/metal/mtl4meshrenderpipelinedescriptor) and [render encoder](https://developer.apple.com/documentation/metal/mtl4rendercommandencoder) documentation.
+
+Apple also ships a sample converting HLSL geometry/tessellation pipelines through mesh shading. This establishes a supported implementation approach, not compatibility with this repository's translator. Shader Converter's normal input is DXIL, whereas the current build consumes the existing translator's SPIR-V. Adding that tool would require another translation and binding path. [Apple sample catalogue](https://developer.apple.com/metal/sample-code/), [Shader Converter input and requirements](https://developer.apple.com/metal/shader-converter/).
+
+## Candidate implementations
+
+| Route | Proposed mapping | Main integration cost |
+| --- | --- | --- |
+| Mesh | One mesh workgroup per input primitive/GS invocation initially; run the required vertex work, then GS logic, then publish vertices and primitive indices. An object stage is optional. | Vertex fetching, indexing, instance IDs, and the VS-to-GS interface must be supplied explicitly. Output limits may require splitting. GPU-driven dispatch needs a fallback on Apple7/8. |
+| Compute plus ordinary draw | Run VS/GS work into an output stream; write indirect draw arguments; a simple vertex shader fetches the stream and feeds the existing fragment shader. | Output allocation, variable emitted counts, primitive conversion, explicit barriers, and render-pass interruption. |
+
+These are proposed designs, not measured implementations. Apple's mesh presentation identifies the extra device-memory buffer and compute-to-render synchronization that the second route needs and the first can avoid. Neither route's speed can be inferred for this game's workloads without a capture. [Apple WWDC22 mesh pipeline discussion](https://developer.apple.com/videos/play/wwdc2022/10162/).
+
+A correctness-first compute prototype can reserve a fixed output region per guest input primitive/invocation, recording emitted counts separately. Compact in original order with a prefix sum if needed. A single atomic append cursor is insufficient as an ordering design: scheduling decides append positions, potentially changing blending order. For mesh lowering, explicitly verify Metal primitive order against guest input/invocation/emit order before treating it as equivalent; the sources examined here do not establish that full correspondence.
+
+Both routes need deliberate handling of indexed/nonindexed draws, primitive restart, strips and adjacency, base vertex/base instance, `PrimitiveId` and invocation ID, `EmitVertex`/`EndPrimitive`, layer/viewport writes, flat varyings and provoking vertices, clipping, and transform feedback. Treat these as coverage requirements to check against the actual shader inventory rather than requirements proven to occur in the example game. Re-running VS independently for every GS primitive can also duplicate work or shader side effects; an intermediate VS stream may be needed for the general case.
+
+## Synchronization and repository integration
+
+Metal 4 requires explicit synchronization of conflicting resource accesses. A compute expansion writes vertices and draw arguments that the following draw reads: the dependency needs both execution order and device-memory visibility. [Apple resource synchronization](https://developer.apple.com/documentation/metal/resource-synchronization).
+
+The installed `MTL4CommandEncoder.h` distinguishes queue producer/consumer barriers from barriers within one encoder. `MTL4VisibilityOptionNone` is execution-only; `MTL4VisibilityOptionDevice` flushes to device coherence; `ResourceAlias` additionally covers aliased virtual addresses. Queue barriers do not synchronize commands within the same encoder. Thus count clearing, expansion, and argument generation inside a compute encoder need encoder barriers where their memory accesses depend on one another. Cross-queue dependencies need events rather than assuming a fence applies across queues. [Apple fence synchronization](https://developer.apple.com/documentation/metal/synchronizing-passes-with-a-fence).
+
+Mesh draws execute object/mesh/fragment stages, rather than vertex/fragment stages. [Apple render-encoder stage mapping](https://developer.apple.com/documentation/metal/mtl4rendercommandencoder?language=objc). In this worktree, `enc_r_begin_pass` binds only vertex and fragment argument tables, while `enc_r_end_pass` has a narrowed producer mask containing vertex/fragment. Mesh integration therefore requires auditing stage bindings, barrier masks, resource lifetimes, and timing instrumentation; changing the PSO and draw call alone is incomplete. [Host implementation](../dm/gfx/nvn/nvn_mtl.m).
+
+The build retains the narrow GS pattern in `gs_slice`, whose metadata lets `prog_fold_gs` select a suitable vertex variant. Other supported GS stages use the compute path, with reflection for output topology, maximum emitted vertices, invocation count, stride and helper resources. [Shader builder](../dm/gfx/tools/dmsh_build.py), [host program composition](../dm/gfx/nvn/nvn_mtl.m).
+
+The compilation probe established that the example shader bodies can pass through the existing compute lowering; the subsequent GPU tests exercise their output. In-game use of an active effect and iOS device execution remain validation gaps. Production shader-chain changes require the header, Mac/iOS shader, guard/extent, and host rebuild procedure for affected profiles.