From 809fbea7aab5cb97c966ba9bbb8555752dec3411 Mon Sep 17 00:00:00 2001 From: Adya <114347520+AdyaGMD@users.noreply.github.com> Date: Thu, 12 Nov 2026 00:32:00 +0000 Subject: [PATCH 1/3] Task and mesh shaders as compute on GPUs without mesh-shading hardware: Ocarina of Time's slow spot on an M1, GPU 35.9 -> 34.6 ms a frame (-1.3 ms, 3.6 %) The M1 and M2 have no mesh-shading hardware (Apple9, M3 on, has it): a render pass with any mesh draw costs 0.14-0.24 ms more, and the mesh stage runs about 20 % slower per meshlet than the same code as a kernel. With DM_NVN_MESH_COMPUTE=1 (off by default) the draw that opens a pass looks ahead through it (mk_hoist_try, as cs_hoist_try does for dispatches) and runs every task / mesh draw it meets as kernels in one compute encoder before the pass: indirect records to dispatch and draw arguments (mk_ind), the task kernels, a scan of their launch counts (mk_scan), the mesh kernels, with a barrier between the phases. Replay draws each of them as a plain indexed draw: the side library's pass-through vertex function and the game's fragment shader (mk_pso, built in the background; the native mesh draw until then). The kernels come from a side library per shader, _bk.metallib in the DM_MSL_DIR overlay, generated from the translated source by the new dm/gfx/tools/dmsh_meshcompute.py (its build command writes them all for a game). The scan stops at a draw that may write memory, a dispatch cs_hoist_try did not move, another target and other ops; it passes uniform updates and keeps native a draw whose uniform buffers or indirect record one of them touched (DM_NVN_MK_UBSTOP=1 stops there). A job is taken only in the pass its scan opened (mk_take). Indirect draws: their {count, first} is written by the GPU, so each (record, mesh shader) has a feedback slot the GPU writes: the largest count seen and, with a task shader, the largest meshlet total, which size the next frames' slots (half again and a margin). A task draw's record seen for the first time is drawn natively while a probe keeps its count. Sized at count x 32 meshlets a task, as at first, a shader writing a large vertex record went past the 64 MB limit and stayed native while its depth prepass, another shader with a 16-byte record, converted; the colour pass's equal depth test then failed on scattered pixels, where the kernel's position differs from the mesh stage's in the last bit, and grass blades flickered white. Keeping native the shaders drawn with an equal depth test (DM_NVN_MK_EQ=1) made it worse, since the prepass shader still converted. Measured (Ocarina of Time, 0400486024DEC000, Kokiri Forest looking down at the map, MacBook M1 8-core GPU, fullscreen 2880x1800): mesh compute alternated on and off every 600 frames in one run, frames 1500-4800 (three turns each, the second 300 frames of every turn): GPU p50 34.3-35.7 ms (mean 34.6) on, 35.1-36.6 (mean 35.9) off; frame p50 36.9 on, 38.2 off; on was lower in every adjacent pair. 82 % of the mesh draws convert (78,829 of 95,766 in 1,800 frames); the rest are shaders with no side library (1,376 draws: per-primitive outputs, point topology), first sights and pipelines still building. Runs compared across processes drifted 3-4 ms apart, more than the gain, so DM_NVN_MESH_COMPUTE takes ab for the DM_AB phases. Known: for about 3 s after the start, while the pipelines build and records are first seen, a prepass and its colour pass can take different paths and the grass flickers. Not run: the DM_AB form of the switch in game (the measurement used an equivalent switch of its own, folded into DM_AB here), any test suite, iPhone, M3 and later (where it stays off unless set), other games. Co-Authored-By: Claude Opus 5.5 --- dm/gfx/nvn/nvn_geometry.inc | 13 +- dm/gfx/nvn/nvn_mtl.m | 331 ++++++++++++++++++++++++++++++- dm/gfx/tools/dmsh_meshcompute.py | 183 +++++++++++++++++ docs/env_switches.tsv | 4 + 4 files changed, 528 insertions(+), 3 deletions(-) create mode 100644 dm/gfx/tools/dmsh_meshcompute.py diff --git a/dm/gfx/nvn/nvn_geometry.inc b/dm/gfx/nvn/nvn_geometry.inc index fc1abe5438..e7510ab54f 100644 --- a/dm/gfx/nvn/nvn_geometry.inc +++ b/dm/gfx/nvn/nvn_geometry.inc @@ -560,7 +560,9 @@ static void geo_mesh_encode(id en_) { // Metal holds a mesh grid to DMSH_MESH_GRID threadgroups, so it launches the draw's workgroups in blocks of that many, one object threadgroup a block. Native mesh draws have no // equivalent fallback: the demanded pipeline must finish compiling before its first draw, so each one built at a draw is recorded as an M row (pso_user.tsv) and the prewarm builds // the rows before the game reaches them (docs/reference/shaders.md, Native mesh pipelines compiled at their first draw) -static struct { int on, miss; NvShader *ts, *ms; uint32_t first, count; uint64_t ind, ind_rec; } g_mtk; // ind: an indirect draw's threadgroups (mdx_mask), ind_rec: its {count, first} record +static struct { int on, miss; NvShader *ts, *ms; uint32_t first, count; uint64_t ind, ind_rec; void *mk; } g_mtk; // ind: an indirect draw's threadgroups (mdx_mask), ind_rec: its {count, first} record +// Mesh as compute on GPUs without mesh-shading hardware (nvn_mtl.m mk_hoist_try): mk, the job this draw draws from (its vertices and indices written before the pass) +static int mk_on(void); static void mk_hoist_try(const uint8_t *from); static void *mk_take(const uint8_t *op); static void mk_note(int ind); static int mk_eq_on(void); static void mk_eq_note(NvShader *ms); static id mk_pso(NvProg *raster); static void mk_encode(id en); static void draw_(const p_draw *dw, const p_draw_idx *di); static uint64_t g_mtk_n, g_mtk_wait; static nvn_hmap g_pso_mtk = {.klen = sizeof(pso_key)}; // only the submit thread uses it: prewarm workers build pipelines (mtk_build) and the submit thread claims them (warm_claim) @@ -629,6 +631,7 @@ static void mtk_claim(pso_key *k, NvShader *ts, NvShader *ms, NvShader *fs, id mtk_pso(NvProg *raster) { if (!raster || !g_mtk.ms || !g_mtk.ms->fn) return nil; + if (g_mtk.mk) { id ps = mk_pso(raster); if (ps) return ps; g_mtk.mk = NULL; } // its pipeline is still being built: the native mesh draw (the job ran for nothing) NvShader *ts = g_mtk.ts, *ms = g_mtk.ms, *fs = raster->sh[1]; pso_key k; pso_key_fill(&k, ms, fs, g_pass.cfmt, g_pass.dfmt, g_pass.sfmt, g_pass.samples, 0); k.border |= border_need(ts, 3); @@ -656,6 +659,7 @@ static id mtk_pso(NvProg *raster) { return ps; } static void mtk_encode(id en_) { + if (g_mtk.mk) { mk_encode(en_); return; } id en = (id)en_; NvShader *ts = g_mtk.ts, *ms = g_mtk.ms; [en setArgumentTable:g_vt atStages:MTLRenderStageMesh]; uint64_t draw_gpu = g_mtk.ind_rec; @@ -693,12 +697,17 @@ static void mesh_tasks_draw(const p_draw *dw) { if (!ind || !mtl_gpu_addr(dw->ind, 8, &mb, &mo, &ind_rec)) { RED_ONCE(154, "DrawMeshTasksIndirect record at %#llx is in no memory pool: draw skipped", (unsigned long long)dw->ind); return; } } PD("m %llu mesh tasks task=%s mesh=%s fs=%s first=%d count=%d indirect=%#llx", g_stat_pass + !g_enc, ts ? ts->key : "-", ms->key, pr->sh[1] ? pr->sh[1]->key : "-", dw->first, dw->count, (unsigned long long)dw->ind); // the draw_ below is a stand-in triangle + if (mk_on()) { if (!g_enc && !tbdr_on()) mk_hoist_try(g_exec_cur); } // the draw opens the pass: its mesh jobs and the later ones of the pass run first + void *mk = mk_on() ? mk_take(g_exec_cur) : NULL; + if (mk_on() && mk_eq_on()) { mk_eq_note(ms); if (ms->mk_eq) mk = NULL; } // the draw's own state: a shader found here drawn with an equal depth test stays native from now on + if (mk_on()) mk_note(ind != 0); g_mtk.on = 1; g_mtk.miss = 0; g_mtk.ts = ts; g_mtk.ms = ms; g_mtk.first = (uint32_t)dw->first; g_mtk.count = (uint32_t)dw->count; + g_mtk.mk = mk; g_mtk.ind = ind; g_mtk.ind_rec = ind_rec; g_geo_program = raster; p_draw triangle = {.prim = 4, .count = 3, .instances = 1}; draw_(&triangle, NULL); - g_geo_program = nil; g_mtk.on = 0; args_dirty_all(); + g_geo_program = nil; g_mtk.on = 0; g_mtk.mk = NULL; args_dirty_all(); if (g_mtk.miss) { g_mtk_wait++; return; } swrite_note(pr); g_wr_epoch++; if (!g_mtk_n++) MLOG("task / mesh: first draw, task %s mesh %s fragment %s, %u workgroups from %u%s", ts ? ts->key : "-", ms->key, pr->sh[1] ? pr->sh[1]->key : "-", g_mtk.count, g_mtk.first, ind ? " (indirect)" : ""); diff --git a/dm/gfx/nvn/nvn_mtl.m b/dm/gfx/nvn/nvn_mtl.m index 1db5ed4f05..149813145c 100644 --- a/dm/gfx/nvn/nvn_mtl.m +++ b/dm/gfx/nvn/nvn_mtl.m @@ -5931,7 +5931,7 @@ static void shmap_read(void) { @interface NvShader : NSObject { @public id fn, tf_fn, discard_fn, fn_u, fn_m, fn_d, geo_fn; uint8_t mask_m; // fn_u: the integer-output entry point (dmsh_entry.py fs_uint), for pipelines whose colour attachments are all unsigned integer (pso_desc); fn_m: the mixed one, its outputs in mask_m unsigned integer; fn_d: the dual-source one (dmsh_entry.py fs_dual) uint16_t tf_stride[4]; uint32_t cb_ext[16], sb_ro; nvn_refl rf; id ae; NSUInteger arg_len; int stage, range_ro, layer_out; const char *builtin; uint64_t logged; char key[24]; id c1buf; uint64_t c1_gpu; size_t c1off; uint32_t geo_vertices, geo_invocations, geo_output, geo_stride, geo_index_stride; - uint8_t mesh_kind, mesh_notask; uint16_t mesh_threads, mesh_payload; // an NVN2 task (1) or mesh (2) shader, a mesh shader with no task, its workgroup's threads, its payload's bytes (dmsh_meshstage.py) + uint8_t mesh_kind, mesh_notask, mk_eq; uint16_t mesh_threads, mesh_payload; // mk_eq: drawn with an equal depth test somewhere, so its draws stay native everywhere (mk_eq_note) // an NVN2 task (1) or mesh (2) shader, a mesh shader with no task, its workgroup's threads, its payload's bytes (dmsh_meshstage.py) id mesh_obj; // a mesh shader with no task: the object function that launches its workgroups in blocks of DMSH_MESH_GRID (dmsh_meshstage.py main0_o) id geo_pso; int geo_pso_failed, cps_failed; // cps_failed: Metal refused this compute shader's pipeline (dispatch): its dispatches are skipped without asking again int16_t *geo_slot; uint32_t geo_words; // a lean passthrough kernel (dmsh_geometry.lean): capture word -> record slot (-1: not read), record words @@ -10688,7 +10688,11 @@ static inline void args_sync_dirty(void) { } } static NSUInteger g_args_mesh; // nonzero: geometry-stage arguments go to this slot of the vertex table, which the mesh stage shares (geo_mesh_draw) +// Mesh as compute (mk_hoist_try): a task / mesh stage's arguments are taken down as (slot, address) pairs instead, set in the compute table when the job's kernel is encoded +enum { MK_CAP = 16 }; +static int g_mk_cap; static int g_mk_cap_n; static uint32_t g_mk_cap_i[MK_CAP]; static uint64_t g_mk_cap_a[MK_CAP]; static void bind_args(id en, int nvn_stage, int metal_vertex, id ab, uint64_t gpu, size_t off, NSUInteger ai) { + if (g_mk_cap && (nvn_stage == 3 || nvn_stage == 4)) { if (g_mk_cap_n < MK_CAP) { g_mk_cap_i[g_mk_cap_n] = (uint32_t)ai; g_mk_cap_a[g_mk_cap_n++] = gpu; } else g_mk_cap = 2; return; } if (g_args_mesh && nvn_stage == 2) { enc_r_vbuf((id)en, g_args_mesh, gpu); return; } if (nvn_stage == 3) { if (rec_on()) rec_table_op(RO_T_ADDR, g_ot, gpu, ai); else [g_ot setAddress:gpu atIndex:ai]; return; } // a task shader: the object stage's table (mesh_tasks_draw) if (nvn_stage == 5 || nvn_stage == 2) enc_c_args((id)en, gpu, ai); @@ -13956,6 +13960,330 @@ static void cs_hoist_try(NvProg *pr0, const p_draw *dw0, const p_draw_idx *di0) memcpy(&S, &Ssave, sizeof S); args_dirty_all(); R = Rsave; g_pso_dirty = pso_dirty || moved; g_dss_dirty = dss_dirty || moved; prog_resolve(); { static int once; if (moved && !once++) MLOG("compute dispatches moved ahead of the render pass they would split (DM_NVN_CS_HOIST): first case pass %llu, %d dispatches", g_stat_pass + 1, moved); } } +// ---------- NVN2 mesh shaders as compute on GPUs without mesh-shading hardware (DM_NVN_MESH_COMPUTE) ---------- +// The M1 and M2 have no mesh-shading hardware (Apple added it with the M3, MTLGPUFamilyApple9), so Metal runs a mesh pipeline its own way, and a render pass with even one mesh draw +// in it costs 0.14-0.24 ms more on an M1 whatever its size (39 vertex draws and one mesh draw: 0.323 ms a pass at 1366x768 against 0.085 with vertex draws only). Ocarina of Time +// draws its scene's geometry through task and mesh shaders in about 11-13 passes a frame. Here a pass's mesh draws are compute jobs instead: when a draw opens the pass, the ops +// after it are scanned up to the first that would end the pass (as cs_hoist_try scans for dispatches), and each task / mesh draw met on the way runs, before the pass opens, as +// the side library's kernels (_bk / _bk, apps/runs/oot-meshk/genk.py for now): the task kernel writes each task workgroup's payload and launch count, a scan +// turns the counts into offsets and the indirect arguments, and the mesh kernel writes each meshlet's vertices and index triples (degenerate past its primitive count) into +// scratch memory. Replay then draws the draw as one indexed draw of those triangles with the library's pass-through vertex function (main0_kv) and the game's fragment shader: +// the pass holds plain draws only. A draw moves ahead of the draws before it only if they write no memory (draw_note_writes), and no uniform update the scan passed touches a +// uniform buffer of its stages; dispatches the scan cannot move, copies, other targets stop it. Indirect mesh draws, and shaders with no side library, stay native. +// DM_NVN_MESH_COMPUTE=1 on, 0 off; off by default while it is measured. ab alternates it with the DM_AB phases: separate runs of Ocarina of Time's slow spot drifted 3-4 ms +// apart (window state, the time into the session), more than the gain measured +static int mk_on(void) { + static int m = -1; + if (m < 0) { m = tbdr_on() ? 0 : sw_mode("DM_NVN_MESH_COMPUTE", 0); if (m) MLOG("DM_NVN_MESH_COMPUTE: task / mesh draws as compute before their pass (GPU family Apple9: %d)", (int)[g_dev supportsFamily:MTLGPUFamilyApple9]); } + return sw_on(m); +} +@interface NvMk : NSObject { @public id k[2], t[2]; id kv; uint32_t maxv, maxp, rec; int ok; } @end // k: mesh kernel (border corrections on / off), t: task kernel, kv: the vertex function +@implementation NvMk @end +static NSMutableDictionary *g_mk_lib; // NvShader pointer → NvMk, or NSNull: no side library +static NvMk *mk_lib(NvShader *sh, int task) { + if (!g_mk_lib) g_mk_lib = [NSMutableDictionary dictionary]; + NSValue *key = [NSValue valueWithPointer:(__bridge const void *)sh]; id o = g_mk_lib[key]; + if (o) return o == NSNull.null ? nil : o; + NvMk *m = nil; NSError *e = nil; + id lib = lib_from_file([NSString stringWithFormat:@"%@/%.16s_bk.metallib", g_msl_dir, sh->key], &e); + { static int n; if (!lib && n++ < 4) MLOG("mesh as compute: no side library %s/%.16s_bk.metallib (%s)", g_msl_dir.UTF8String, sh->key, e ? e.localizedDescription.UTF8String : "-"); } + if (lib) { + m = [NvMk new]; + if (task) { id f = geo_mesh_entry(lib, @"main0_tk"); for (int off = 0; f && off < 2; off++) m->t[off] = cps_new_b(f, off, &e); m->ok = m->t[0] && m->t[1]; } + else { + id f = geo_mesh_entry(lib, [sh->fn.name isEqualToString:@"main0_z"] ? @"main0_k_z" : @"main0_k"); + for (int off = 0; f && off < 2; off++) m->k[off] = cps_new_b(f, off, &e); + m->kv = geo_mesh_entry(lib, @"main0_kv"); + for (NSString *n in lib.functionNames) if ([n hasPrefix:@"dm_kinfo_"]) { NSArray *p = [[n substringFromIndex:9] componentsSeparatedByString:@"_"]; + if (p.count == 3) { m->maxv = (uint32_t)[p[0] intValue]; m->maxp = (uint32_t)[p[1] intValue]; m->rec = (uint32_t)[p[2] intValue]; } } + m->ok = m->k[0] && m->k[1] && m->kv && m->maxv && m->maxp && m->rec; + } + if (!m->ok) MLOG("[RED] mesh as compute: the side library of %s did not give its %s: %s", sh->key, task ? "task kernel" : "kernels", e ? e.localizedDescription.UTF8String : "entry missing"); + } + g_mk_lib[key] = m && m->ok ? m : NSNull.null; + return m && m->ok ? m : nil; +} +static id g_mk_ind_ps; +static int mk_fb_slot(uint64_t ind, const void *ms, uint32_t first_cap, uint32_t *cap, uint32_t *gcap); +static id mk_scan_pso(void) { + static id ps; static int failed; + if (ps || failed) return ps; + NSError *e = nil; + id l = [g_dev newLibraryWithSource:@"#include \nusing namespace metal;\n" + // groups: the meshlet slots the job has; slot: an indirect draw's feedback slot (~0u: none), whose meshlet total sizes the next frames (ov[4 + slot]) + "struct mk_scan_job { uint ntask, bound, maxp3, groups, slot, pad0, pad1, pad2; };\n" + // ov: [0] the largest launch count past the bound, [1] how many task workgroups were cut to it (diagnostic), [2] totals cut to the slots, [4 + slot] each record's largest total + "kernel void mk_scan(constant mk_scan_job& j [[buffer(0)]], device const uint* cnt [[buffer(1)]], device uint* moff [[buffer(2)]], device uint* args [[buffer(3)]], device const uint* dr [[buffer(4)]], device atomic_uint* ov [[buffer(5)]], uint lid [[thread_position_in_threadgroup]]) {\n" + " if (lid) return; uint t = 0, nt = dr[2];\n" + " for (uint i = 0; i < nt; i++) { moff[i] = t; uint c = cnt[i]; if (c > j.bound) { atomic_fetch_max_explicit(&ov[0], c, memory_order_relaxed); atomic_fetch_add_explicit(&ov[1], 1u, memory_order_relaxed); } t += min(c, j.bound); }\n" + // past the count up to the slots (j.ntask): the total, so a mesh kernel whose search bound (its constant dm_draw[2]) read the capacity, not the count mk_ind wrote, still finds its task + " for (uint i = nt; i <= j.ntask; i++) moff[i] = t;\n" + " if (j.slot != ~0u) atomic_fetch_max_explicit(&ov[4 + j.slot], t, memory_order_relaxed);\n" + " if (t > j.groups) { atomic_fetch_add_explicit(&ov[2], 1u, memory_order_relaxed); t = j.groups; }\n" // the later tasks' meshlets are dropped that frame + " args[0] = t; args[1] = 1; args[2] = 1; args[3] = 0;\n" + " args[4] = t * j.maxp3; args[5] = 1; args[6] = 0; args[7] = 0; args[8] = 0; }\n" + // An indirect draw's {count, first} as the GPU wrote it: clamped to the slots the job has, the real count kept for the next frames' size (feedback slot) + "struct mk_ind_job { uint cap, maxp3, slot, pad; };\n" + "kernel void mk_ind(constant mk_ind_job& j [[buffer(0)]], device const uint* rec [[buffer(1)]], device uint* draw [[buffer(2)]], device uint* ia [[buffer(3)]], device uint* fb [[buffer(4)]], uint lid [[thread_position_in_threadgroup]]) {\n" + " if (lid) return; uint c = rec[0], n = min(c, j.cap); fb[j.slot] = max(fb[j.slot], c);\n" + " draw[0] = n; draw[1] = rec[1]; draw[2] = n; draw[3] = 0;\n" + " ia[0] = n; ia[1] = 1; ia[2] = 1; ia[3] = 0; ia[4] = n * j.maxp3; ia[5] = 1; ia[6] = 0; ia[7] = 0; ia[8] = 0; }\n" options:nil error:&e]; + id f = [l newFunctionWithName:@"mk_scan"], fi = [l newFunctionWithName:@"mk_ind"]; + if (f) ps = [g_dev newComputePipelineStateWithFunction:f error:&e]; + if (fi) g_mk_ind_ps = [g_dev newComputePipelineStateWithFunction:fi error:&e]; + if (!ps || !g_mk_ind_ps) { failed = 1; ps = nil; MLOG("[RED] mesh as compute: the scan kernels did not build: %s", e.localizedDescription.UTF8String); } + if (ps) { uint32_t cap, gcap; mk_fb_slot(0, NULL, 0, &cap, &gcap); } // creates the feedback buffer, which the scans also write (their cut launches); key 0 is no record + return ps; +} +// Feedback for indirect draws: the most workgroups each (record address, mesh shader) has asked for, which the GPU writes (mk_ind) and the next frames size their slots by +enum { MK_FB = 4096 }; +static id g_mk_fb; static struct { uint64_t ind; const void *ms; uint32_t cap, gcap; } g_mk_fbk[MK_FB]; // gcap: meshlet slots for a task draw, from the largest total the scans kept (0: not known yet) +static uint64_t g_mk_over, g_mk_first, g_mk_eq_sh, g_mk_eq_rej; // eq: mesh shaders drawn with an equal depth test, draws they kept native (mk_eq_note) +static int mk_fb_slot(uint64_t ind, const void *ms, uint32_t first_cap, uint32_t *cap, uint32_t *gcap) { + if (!g_mk_fb) { g_mk_fb = [g_dev newBufferWithLength:(2 * MK_FB + 4) * 4 options:MTLResourceStorageModeShared]; memset(g_mk_fb.contents, 0, (2 * MK_FB + 4) * 4); res_add(g_mk_fb); } // past MK_FB: mk_scan's counters, then each slot's largest meshlet total + uint32_t h = (uint32_t)(((ind >> 4) ^ ((uintptr_t)ms >> 4)) * 2654435761u) & (MK_FB - 1); + for (int k = 0; k < 64; k++, h = (h + 1) & (MK_FB - 1)) { + if (g_mk_fbk[h].ind == ind && g_mk_fbk[h].ms == ms) { + uint32_t tot = ((volatile uint32_t *)g_mk_fb.contents)[MK_FB + 4 + h]; + if (tot > g_mk_fbk[h].gcap || (tot && !g_mk_fbk[h].gcap)) g_mk_fbk[h].gcap = tot + tot / 2 + 64; // half again and a margin: totals move with the camera + *gcap = g_mk_fbk[h].gcap; + uint32_t seen = ((volatile uint32_t *)g_mk_fb.contents)[h]; + if (!g_mk_fbk[h].cap) { if (seen) g_mk_fbk[h].cap = seen + seen / 2 + 8; } // probed (mk_prep): sized from the count the GPU kept; still 0 until the probe has run + else if (seen > g_mk_fbk[h].cap) { g_mk_over++; static int n; if (n++ < 8) MLOG("mesh as compute: an indirect draw asked for %u workgroups with %u slots (cut that frame); now %u", seen, g_mk_fbk[h].cap, seen + seen / 2 + 8); g_mk_fbk[h].cap = seen + seen / 2 + 8; } + *cap = g_mk_fbk[h].cap; return (int)h; } + if (!g_mk_fbk[h].ind) { g_mk_fbk[h].ind = ind; g_mk_fbk[h].ms = ms; g_mk_fbk[h].cap = first_cap; g_mk_fbk[h].gcap = 0; ((uint32_t *)g_mk_fb.contents)[h] = 0; ((uint32_t *)g_mk_fb.contents)[MK_FB + 4 + h] = 0; *cap = first_cap; *gcap = 0; return (int)h; } + } + return -1; +} +enum { MK_JOBS = 128, MK_KV = 27, MK_KI = 28, MK_DRAW = 29, MK_PL = 26, MK_CNT = 25, MK_MOFF = 24 }; +typedef struct { + const uint8_t *op; __unsafe_unretained NvMk *m, *t; __unsafe_unretained NvShader *ms, *ts; + int moff_b, toff_b, nm, nt; uint32_t mi[MK_CAP], ti[MK_CAP]; uint64_t ma[MK_CAP], ta[MK_CAP]; // the stages' arguments (bind_args, g_mk_cap) + int probe; // only records an indirect record's count (mk_prep): no task or mesh phase, never drawn + uint32_t count, first, groups; int slot; uint64_t kv, ki, ki_bytes, draw, pl, cnt, moff, args, scan, ind, ia, ij; // ind: an indirect draw's record (count: its slots then); ia: its dispatch {x,1,1} and, with no task, draw arguments (mk_ind) // groups: the meshlet slots (count with no task; count × the task's launch bound) +} mk_job; +static mk_job g_mkj[MK_JOBS]; static int g_mkj_n; static uint64_t g_mk_met; +static uint64_t g_mk_rej[8]; // why a task / mesh draw stayed native: 0 indirect, 1 count / feedback / discard, 2 shaders, 3 no mesh library, 4 no task library, 5 size, 6 arguments +static uint64_t g_mk_jobs, g_mk_drawn, g_mk_scans, g_mk_stop[8], g_mk_native; // jobs run, drawn; scans; why scans stopped (0 draw writes memory, 1 uniform update, 2 dispatch, 3 other op, 4 full, 5 targets); mesh draws drawn natively +// The jobs belong to the pass their scan opened (g_mk_pass): a pass opened without a scan (a geometry-emulated draw, a clear) must not take them, though replay meets the same ops +// again when a command list runs twice (another view, other uniforms): the old job's vertices were drawn with this pass's state (flickering stray triangles) +static uint64_t g_mk_pass, g_mk_stale; +static void *mk_take(const uint8_t *op) { + if (g_mkj_n && g_mk_pass != g_stat_pass + (g_enc ? 0 : 1)) { + for (int i = 0; i < g_mkj_n; i++) if (g_mkj[i].op == op) { g_mk_stale++; static int n; if (n++ < 4) MLOG("mesh as compute: a draw met a job of pass %llu in pass %llu (drawn natively)", (unsigned long long)g_mk_pass, (unsigned long long)(g_stat_pass + (g_enc ? 0 : 1))); break; } + g_mkj_n = 0; } + for (int i = 0; i < g_mkj_n; i++) if (g_mkj[i].op == op) { g_mkj[i].op = NULL; g_mk_drawn++; return &g_mkj[i]; } + g_mk_native++; + return NULL; +} +// Every task / mesh draw (mesh_tasks_draw), indirect ones too: a line every 300 frames +static uint64_t g_mk_all, g_mk_ind; +static void mk_note(int ind) { + g_mk_all++; g_mk_ind += ind != 0; + static uint64_t last; uint64_t f = atomic_load(&g_cur_frame); + if (f / 300 == last / 300) return; + last = f; + const uint32_t *ov = g_mk_fb ? (const uint32_t *)g_mk_fb.contents + MK_FB : NULL; + if (g_mk_first) { int used = 0; for (int i = 0; i < MK_FB; i++) used += g_mk_fbk[i].ind != 0; MLOG("mesh as compute: %llu indirect task draws probed (first sight, drawn natively); %d records known", (unsigned long long)g_mk_first, used); } + if (g_mk_eq_sh) MLOG("mesh as compute: %llu mesh shaders drawn with an equal depth test, %llu of their draws left native", (unsigned long long)g_mk_eq_sh, (unsigned long long)g_mk_eq_rej); + if (g_mk_stale) MLOG("mesh as compute: %llu draws met a job of another pass (drawn natively)", (unsigned long long)g_mk_stale); + if (ov && ov[2]) MLOG("mesh as compute: %u indirect task draws had more meshlets than their slots (later tasks dropped that frame)", ov[2]); + if (ov && ov[1]) MLOG("mesh as compute: %u task workgroups launched more mesh workgroups than their bound (largest %u): those were cut", ov[1], ov[0]); + MLOG("mesh as compute up to frame %llu: %llu task / mesh draws, %llu indirect, %llu indirect draws cut to their slots, %llu drawn from compute, %llu native; %llu jobs in %llu scans; scans stopped by a draw writing memory %llu, a uniform update %llu, a dispatch %llu, another op %llu, a full table %llu, other targets %llu; left native: indirect %llu, count %llu, shaders %llu, no mesh library %llu, no task library %llu, size %llu, arguments %llu; mesh draws met by scans %llu", + (unsigned long long)f, g_mk_all, g_mk_ind, g_mk_over, g_mk_drawn, g_mk_native, g_mk_jobs, g_mk_scans, g_mk_stop[0], g_mk_stop[1], g_mk_stop[2], g_mk_stop[3], g_mk_stop[4], g_mk_stop[5], g_mk_rej[0], g_mk_rej[1], g_mk_rej[2], g_mk_rej[3], g_mk_rej[4], g_mk_rej[5], g_mk_rej[6], g_mk_met); +} +// Whether this draw may write memory (draw_note_writes, without keeping its note) +static int mk_draw_writes(NvProg *pr) { + int w0 = g_pass_memw, n0 = g_mpw_n; uint64_t k0[6]; memcpy(k0, g_pw_kind, sizeof k0); + g_pass_memw = 0; draw_note_writes(pr); + int w = g_pass_memw || g_mpw_n > n0; + g_pass_memw = w0; g_mpw_n = n0; memcpy(g_pw_kind, k0, sizeof k0); + return w; +} +static int mk_ub_hit(NvShader *sh, int st, const csh_rng *up, int nup) { + if (!sh || !nup) return 0; + for (int i = 0; i < sh->rf.nb; i++) { const nvn_rbind *b = &sh->rf.b[i]; + if ((b->kind == NVN_RB_CBUF || b->kind == NVN_RB_UBOP) && b->reg < 16) { csh_rng r = {S.ubo[st][b->reg].addr, S.ubo[st][b->reg].addr + S.ubo[st][b->reg].size}; + if (r.lo && csh_hit(&r, 1, up, nup)) return 1; } } + return 0; +} +// A task / mesh draw at its place in the scan (S as it is there): its job's memory and arguments, 0 when it stays native +#define MK_REJ(i) do { if (!g_mk_rej[i]++) MLOG("mesh as compute: first draw left native for reason %d (task %s mesh %s count %d ind %llx)", i, pr && pr->sh[3] ? pr->sh[3]->key : "-", pr && pr->sh[4] ? pr->sh[4]->key : "-", dw->count, (unsigned long long)dw->ind); return 0; } while (0) +// A mesh shader drawn with an equal depth test (a depth prepass, then the colour pass on equal depth): both draws must take the same path, since the kernel's position can differ +// from the native mesh stage's in the last bit and the test then fails on scattered pixels (Ocarina of Time's grass: blades flickering white). +// DM_NVN_MK_EQ=1 (diagnostic, off by default) keeps such a shader native; it made things worse in game, since the depth prepass draws with another shader that still converted +static int mk_eq_on(void) { static int on = -1; if (on < 0) { const char *e = getenv("DM_NVN_MK_EQ"); on = e && !strcmp(e, "1"); if (on) MLOG("DM_NVN_MK_EQ=1: mesh shaders drawn with an equal depth test stay native (diagnostic: their depth prepass, another shader, still converts)"); } return on; } +static void mk_eq_note(NvShader *ms) { + if (!ms || ms->mk_eq || !R.rt.depth || !S.ds.dt || S.ds.df != NVN_DEPTH_FUNC_EQUAL) return; + ms->mk_eq = 1; g_mk_eq_sh++; + static int n; if (n++ < 8) MLOG("mesh as compute: mesh %s is drawn with an equal depth test: it stays native", ms->key); +} +static int mk_prep(NvProg *pr, const p_draw *dw, mk_job *j) { + if ((!dw->ind && dw->count <= 0) || S.tf_on || S.rast_discard || !pr) MK_REJ(1); + NvShader *ts = pr->sh[3], *ms = pr->sh[4]; + if (mk_eq_on()) { mk_eq_note(ms); if (ms && ms->mk_eq) { g_mk_eq_rej++; return 0; } } + if (ts && ts->mesh_kind != 1) ts = nil; + if (!ms || ms->mesh_kind != 2 || !ms->fn || (!ts && !ms->mesh_notask)) MK_REJ(2); + NvMk *m = mk_lib(ms, 0), *t = ts ? mk_lib(ts, 1) : nil; + if (!m) MK_REJ(3); + if (ts && !t) MK_REJ(4); + if (!mk_scan_pso()) MK_REJ(6); // also builds mk_ind + uint32_t count = (uint32_t)dw->count; uint64_t ind_rec = 0; int slot = -1; + // DM_NVN_MK_IND (diagnostic): 0 every indirect task / mesh draw stays native, task only those with a task shader convert, notask only those without + { static int ind_on = -1; if (ind_on < 0) { const char *e = getenv("DM_NVN_MK_IND"); ind_on = !e ? 1 : !strcmp(e, "0") ? 0 : !strcmp(e, "task") ? 2 : !strcmp(e, "notask") ? 3 : 1; + if (ind_on != 1) MLOG("DM_NVN_MK_IND=%s: %s", e, ind_on == 0 ? "indirect task / mesh draws stay native" : ind_on == 2 ? "only indirect draws with a task shader convert" : "only indirect draws without a task shader convert"); } + if (dw->ind && (!ind_on || (ind_on == 2 && !ts) || (ind_on == 3 && ts))) MK_REJ(0); } + uint32_t gcap = 0; + if (dw->ind) { void *mb; uint64_t mo; uint32_t cap = 0; + if (!mtl_gpu_addr(dw->ind, 8, &mb, &mo, &ind_rec) || (slot = mk_fb_slot(dw->ind, (__bridge const void *)ms, ts ? 0 : 256, &cap, &gcap)) < 0) MK_REJ(0); + // A task draw's record seen for the first time: drawn natively, with a probe (mk_ind with no slots) that keeps its count for the next frames. A guess cut it silently + // (Ocarina of Time's grass: 16 of about 300 task workgroups, other ones every frame, blades flickering white), since only a record met again reports its cut + if (!cap) { id b; size_t o; g_mk_first++; + *j = (mk_job){.probe = 1, .ind = ind_rec, .slot = slot}; + o = arena_alloc(16, &b); j->draw = b.gpuAddress + o; + o = arena_alloc(48, &b); j->ia = b.gpuAddress + o; + o = arena_alloc(16, &b); { uint32_t ij[4] = {0, 0, (uint32_t)slot, 0}; memcpy((uint8_t *)b.contents + o, ij, 16); } j->ij = b.gpuAddress + o; + return 1; } + count = cap; } + uint32_t bound = ts ? ts->mesh_threads : 1, groups = count * bound; + // An indirect task draw: slots for the meshlets its scans have counted (gcap), not count × the launch bound, which is many times what the tasks launch: sized that way a record + // whose shader writes a large vertex record went past the size limit and stayed native while its depth prepass, a small record, converted, and the colour pass's equal depth test + // failed on scattered pixels (Ocarina of Time's grass blades flickering white). Not known yet: 4 a task, which the first scan corrects + if (ts && dw->ind) groups = MIN(groups, gcap ? gcap : count * MIN(bound, 4u)); + uint64_t kvb = (uint64_t)groups * m->maxv * m->rec, kib = (uint64_t)groups * m->maxp * 12; + if (!groups || kvb + kib > (64u << 20)) MK_REJ(5); + *j = (mk_job){.m = m, .t = t, .ms = ms, .ts = ts, .count = count, .first = (uint32_t)dw->first, .groups = groups, .ki_bytes = kib, .ind = ind_rec, .slot = slot}; + j->moff_b = !border_need(ms, 4); j->toff_b = ts ? !border_need(ts, 3) : 0; + image_prepare(pr, 1u << 3 | 1u << 4); + g_mk_cap = 1; + if (ts) { g_mk_cap_n = 0; args_dirty_all(); encode_args(nil, ts, 3, 0); j->nt = g_mk_cap_n; memcpy(j->ti, g_mk_cap_i, sizeof j->ti); memcpy(j->ta, g_mk_cap_a, sizeof j->ta); } + g_mk_cap_n = 0; args_dirty_all(); encode_args(nil, ms, 4, 1); j->nm = g_mk_cap_n; memcpy(j->mi, g_mk_cap_i, sizeof j->mi); memcpy(j->ma, g_mk_cap_a, sizeof j->ma); + int over = g_mk_cap == 2; g_mk_cap = 0; g_mk_cap_n = 0; args_dirty_all(); + if (over || !j->nm) MK_REJ(6); + id b; size_t o; + o = arena_alloc(kvb, &b); j->kv = b.gpuAddress + o; + o = arena_alloc(kib, &b); j->ki = b.gpuAddress + o; + o = arena_alloc(16, &b); { uint32_t d[4] = {j->count, j->first, j->count, 0}; memcpy((uint8_t *)b.contents + o, d, 16); } j->draw = b.gpuAddress + o; + if (ts) { + o = arena_alloc((size_t)j->count * ((ts->mesh_payload + 15u) & ~15u), &b); j->pl = b.gpuAddress + o; + o = arena_alloc((size_t)j->count * 4, &b); j->cnt = b.gpuAddress + o; + o = arena_alloc((size_t)j->count * 4 + 4, &b); j->moff = b.gpuAddress + o; + o = arena_alloc(48, &b); j->args = b.gpuAddress + o; + o = arena_alloc(32, &b); { uint32_t sj[8] = {j->count, bound, m->maxp * 3, groups, dw->ind ? (uint32_t)slot : ~0u, 0, 0, 0}; memcpy((uint8_t *)b.contents + o, sj, 32); } j->scan = b.gpuAddress + o; + } + if (j->ind) { + o = arena_alloc(48, &b); j->ia = b.gpuAddress + o; + o = arena_alloc(16, &b); { uint32_t ij[4] = {j->count, m->maxp * 3, (uint32_t)slot, 0}; memcpy((uint8_t *)b.contents + o, ij, 16); } j->ij = b.gpuAddress + o; + } + return 1; +} +static void mk_set(int n, const uint32_t *ix, const uint64_t *a) { for (int i = 0; i < n; i++) enc_c_args(nil, a[i], ix[i]); } +// The jobs found, in three phases with a barrier between: task kernels, scans, mesh kernels +static void mk_run(int n) { + int tasks = 0, inds = 0; for (int i = 0; i < n; i++) { tasks += g_mkj[i].ts != nil; inds += g_mkj[i].ind != 0; } + id ce = dispatch_enc(); + if (inds) { + enc_c_pso(ce, g_mk_ind_ps); + for (int i = 0; i < n; i++) { mk_job *j = &g_mkj[i]; if (!j->ind) continue; + enc_c_args(ce, j->ij, 0); enc_c_args(ce, j->ind, 1); enc_c_args(ce, j->draw, 2); enc_c_args(ce, j->ia, 3); enc_c_args(ce, g_mk_fb.gpuAddress, 4); + enc_c_dispatch(ce, MTLSizeMake(1, 1, 1), MTLSizeMake(32, 1, 1)); } + [ce barrierAfterEncoderStages:MTLStageDispatch beforeEncoderStages:MTLStageDispatch visibilityOptions:MTL4VisibilityOptionDevice]; + } + if (tasks) { + for (int i = 0; i < n; i++) { mk_job *j = &g_mkj[i]; if (!j->ts || j->probe) continue; + mk_set(j->nt, j->ti, j->ta); enc_c_args(ce, j->draw, MK_DRAW); enc_c_args(ce, j->pl, MK_PL); enc_c_args(ce, j->cnt, MK_CNT); + enc_c_pso(ce, j->t->t[j->toff_b]); + if (j->ind) enc_c_dispatch_indirect(ce, j->ia, MTLSizeMake(j->ts->mesh_threads, 1, 1)); else enc_c_dispatch(ce, MTLSizeMake(j->count, 1, 1), MTLSizeMake(j->ts->mesh_threads, 1, 1)); } + [ce barrierAfterEncoderStages:MTLStageDispatch beforeEncoderStages:MTLStageDispatch visibilityOptions:MTL4VisibilityOptionDevice]; + enc_c_pso(ce, mk_scan_pso()); + for (int i = 0; i < n; i++) { mk_job *j = &g_mkj[i]; if (!j->ts) continue; + enc_c_args(ce, j->scan, 0); enc_c_args(ce, j->cnt, 1); enc_c_args(ce, j->moff, 2); enc_c_args(ce, j->args, 3); enc_c_args(ce, j->draw, 4); enc_c_args(ce, g_mk_fb.gpuAddress + MK_FB * 4, 5); + enc_c_dispatch(ce, MTLSizeMake(1, 1, 1), MTLSizeMake(32, 1, 1)); } + [ce barrierAfterEncoderStages:MTLStageDispatch beforeEncoderStages:MTLStageDispatch | MTLStageBlit visibilityOptions:MTL4VisibilityOptionDevice]; + } + for (int i = 0; i < n; i++) { mk_job *j = &g_mkj[i]; if (j->probe) continue; + mk_set(j->nm, j->mi, j->ma); enc_c_args(ce, j->kv, MK_KV); enc_c_args(ce, j->ki, MK_KI); enc_c_args(ce, j->draw, MK_DRAW); + if (j->ts) { enc_c_args(ce, j->pl, MK_PL); enc_c_args(ce, j->moff, MK_MOFF); } + enc_c_pso(ce, j->m->k[j->moff_b]); + if (j->ts) enc_c_dispatch_indirect(ce, j->args, MTLSizeMake(j->ms->mesh_threads, 1, 1)); + else if (j->ind) enc_c_dispatch_indirect(ce, j->ia, MTLSizeMake(j->ms->mesh_threads, 1, 1)); + else enc_c_dispatch(ce, MTLSizeMake(j->count, 1, 1), MTLSizeMake(j->ms->mesh_threads, 1, 1)); } + args_dirty_all(); + g_mk_jobs += (uint64_t)n; + { static int once; if (!once++) MLOG("mesh as compute: first jobs, %d draws (%d with a task shader) before pass %llu", n, tasks, g_stat_pass + 1); } +} +// When a draw opens a pass (from: the first op to look at; the opening task / mesh draw itself when it is one) +static void mk_hoist_try(const uint8_t *from) { + g_mkj_n = 0; // what an earlier scan left is not drawn now: replay never reached it in its pass + g_mk_pass = g_stat_pass + 1; // the pass the draw is about to open (mk_take) + if (!from || g_ret_n > 8 || g_fz_checking || g_defer_replaying) return; + fz_cursor cur = {.p = from}; for (int k = 0; k < g_ret_n; k++) cur.ret[cur.ret_n++] = g_ret[k]; + static typeof(S) Ssave; memcpy(&Ssave, &S, sizeof S); typeof(R) Rsave = R; int pso_dirty = g_pso_dirty, dss_dirty = g_dss_dirty; + static csh_rng up[CSH_RN]; int nup = 0, n = 0, landed = 0, stop = -1; + g_mk_scans++; + for (int guard = 0; guard < 4096 && stop < 0; guard++) { + const uint8_t *q; op_hdr hq; const void *aq; + if (!fz_next(&cur, &q, &hq, &aq)) break; + switch (hq.op) { + case OP_DRAW: case OP_DRAW_IDX: { + p_draw tdw; p_draw_idx tdi; const p_draw *d = hq.op == OP_DRAW ? draw_rec(aq, hq.size, &tdw) : NULL; if (hq.op == OP_DRAW_IDX) draw_idx_rec(aq, hq.size, &tdi); + prog_resolve(); __unsafe_unretained NvProg *pr = S.prog ? prog_get(S.prog) : nil; + if (d && d->prim == DM_PRIM_MESH_TASKS) g_mk_met++; + csh_rng irec = {d ? d->ind : 0, d ? d->ind + 8 : 0}; + if (d && d->prim == DM_PRIM_MESH_TASKS && pr && !mk_ub_hit(pr->sh[3], 3, up, nup) && !mk_ub_hit(pr->sh[4], 4, up, nup) && !(d->ind && nup && csh_hit(&irec, 1, up, nup))) { + if (n == MK_JOBS) { stop = 4; break; } + if (!landed && g_nubd) { ubd_land(); landed = 1; } // deferred uniform updates land first: the kernels read memory, not the draws' shadows + if (mk_prep(pr, d, &g_mkj[n])) { g_mkj[n].op = g_mkj[n].probe ? NULL : q; n++; } + } + if (mk_draw_writes(pr)) stop = 0; // later jobs could read what it writes + break; } + case OP_UPDATE_UB: { uint64_t dd, nn; memcpy(&dd, aq, 8); memcpy(&nn, (const uint8_t *)aq + 8, 8); if (nn && !csh_add(up, &nup, dd, nn)) stop = 1; + static int ubstop = -1; if (ubstop < 0) { const char *e = getenv("DM_NVN_MK_UBSTOP"); ubstop = e && !strcmp(e, "1"); if (ubstop) MLOG("DM_NVN_MK_UBSTOP=1: a uniform update stops the mesh scan (the kernels may read it through global memory)"); } + if (ubstop && nn) stop = 1; break; } + case OP_BARRIER: case OP_RENDER_COND: break; + case OP_REPORT_COUNTER: case OP_RESET_COUNTER: { uint64_t t; memcpy(&t, aq, 8); if (t != NVN_COUNTER_TYPE_SAMPLES_PASSED) stop = 3; break; } + case OP_DISPATCH: { int moved = 0; for (int i = 0; i < g_csh_n; i++) moved |= g_csh_op[i] == q; if (!moved) stop = 2; break; } // one cs_hoist_try ran ahead already ran before these jobs + case OP_SET_RT: if (!csh_rt_same(q, aq)) stop = 5; break; + default: if (!state_op(hq.op, aq)) stop = 3; break; + } + } + if (stop >= 0) g_mk_stop[stop]++; + memcpy(&S, &Ssave, sizeof S); args_dirty_all(); R = Rsave; g_pso_dirty = pso_dirty || n; g_dss_dirty = dss_dirty || n; prog_resolve(); + if (n) mk_run(n); + g_mkj_n = n; +} +// The draw's pipeline: the job's vertex function and the program's fragment shader, built in the background (nil until then: the native mesh draw) +static nvn_hmap g_pso_mk = {.klen = sizeof(pso_key)}; static os_unfair_lock g_pso_mk_lock = OS_UNFAIR_LOCK_INIT; static NSObject *g_pso_mk_building; +static id mk_pso(NvProg *raster) { + mk_job *j = g_mtk.mk; NvShader *fs = raster->sh[1]; + pso_key k; pso_key_fill(&k, j->ms, fs, g_pass.cfmt, g_pass.dfmt, g_pass.sfmt, g_pass.samples, 0); + memset(k.attr, 0, sizeof k.attr); memset(k.attr_hi, 0, sizeof k.attr_hi); memset(k.stream, 0, sizeof k.stream); + k.vs = (__bridge void *)j->ms; k.fs = (__bridge void *)fs; k.fsfn = (__bridge void *)j->m->kv; + os_unfair_lock_lock(&g_pso_mk_lock); id ps = hm_get_id(&g_pso_mk, &k); + if (!ps) { if (!g_pso_mk_building) g_pso_mk_building = [NSObject new]; hm_put_id(&g_pso_mk, &k, g_pso_mk_building); } + os_unfair_lock_unlock(&g_pso_mk_lock); + if (!ps) { + MTLRenderPipelineDescriptor *d = pso_desc(&k, nil, fs, nil, 0); + if (!pso_ds_valid((uint32_t)d.depthAttachmentPixelFormat, (uint32_t)d.stencilAttachmentPixelFormat)) { os_unfair_lock_lock(&g_pso_mk_lock); hm_put_id(&g_pso_mk, &k, NSNull.null); os_unfair_lock_unlock(&g_pso_mk_lock); return nil; } + d.vertexFunction = fn_dev_b(j->m->kv, 1); if (d.fragmentFunction) d.fragmentFunction = fn_dev_b(d.fragmentFunction, !k.border); d.vertexDescriptor = nil; + const char *mk_ = j->ms->key, *fk = fs ? fs->key : "-"; pso_key kk = k; + dispatch_async(pp_bg_queue(), ^{ + NSError *e = nil; uint64_t t0 = bp_now(); + id p = [g_dev newRenderPipelineStateWithDescriptor:d error:&e]; + static int n; if (!p || n++ < 8) MLOG("%smesh as compute: pipeline for mesh %s fragment %s, %.1f ms in the background%s%s", p ? "" : "[RED] ", mk_, fk, (bp_now() - t0) / 1e6, p ? "" : ": ", p ? "" : e.localizedDescription.UTF8String); + os_unfair_lock_lock(&g_pso_mk_lock); hm_put_id(&g_pso_mk, &kk, p ?: NSNull.null); os_unfair_lock_unlock(&g_pso_mk_lock); + }); + return nil; + } + return ps == NSNull.null || ps == g_pso_mk_building ? nil : ps; +} +static void mk_encode(id en_) { + id en = (id)en_; mk_job *j = g_mtk.mk; + enc_r_vbuf(en, MK_KV, j->kv); + if (j->ts) enc_r_draw_indexed_indirect(en, MTLPrimitiveTypeTriangle, MTLIndexTypeUInt32, j->ki, j->ki_bytes, j->args + 16); + else if (j->ind) enc_r_draw_indexed_indirect(en, MTLPrimitiveTypeTriangle, MTLIndexTypeUInt32, j->ki, j->ki_bytes, j->ia + 16); + else enc_r_draw_indexed(en, MTLPrimitiveTypeTriangle, (NSUInteger)j->groups * j->m->maxp * 3, MTLIndexTypeUInt32, j->ki, j->ki_bytes, 1, 0, 0); +} void mtl_geomesh_stats(uint64_t *n, uint64_t *wait, uint64_t *run_draws, uint64_t *runs, uint64_t ends[4]) { *n = g_geo_mesh_n; *wait = g_geo_mesh_wait; *run_draws = g_mrun_draws; *runs = g_mrun_runs; memcpy(ends, g_mrun_ends, sizeof g_mrun_ends); } void mtl_geohoist_stats(uint64_t *moved, uint64_t *runs, uint64_t *stops) { *moved = g_geo_hoisted; *runs = g_geo_hoist_runs; memset(stops, 0, 6 * sizeof *stops); stops[3] = g_geo_dependency_splits; } void mtl_cshoist_stats(uint64_t *moved, uint64_t *depth, uint64_t *stops) { *moved = g_csh_moved; *depth = g_csh_depth; for (int i = 0; i < 8; i++) stops[i] = g_csh_stop[i]; } @@ -14114,6 +14442,7 @@ static void draw_(const p_draw *dw, const p_draw_idx *di) { 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 && !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 if (!g_enc && !tbdr_on() && !g_geo_program && cshoist_on()) cs_hoist_try(pr, dw, di); // see DM_NVN_CS_HOIST + if (!g_enc && !tbdr_on() && !g_geo_program && mk_on()) mk_hoist_try(g_exec_next); // the pass's mesh draws as compute before it opens (see mk_hoist_try) 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; } diff --git a/dm/gfx/tools/dmsh_meshcompute.py b/dm/gfx/tools/dmsh_meshcompute.py new file mode 100644 index 0000000000..78bb2d5a1e --- /dev/null +++ b/dm/gfx/tools/dmsh_meshcompute.py @@ -0,0 +1,183 @@ +"""NVN2 task and mesh shaders as compute, for GPUs without mesh-shading hardware (Apple M1 and M2; Apple9, M3 on, has it). + +On those GPUs a render pass with any mesh draw costs 0.14-0.24 ms more, and the mesh stage runs about 20 % slower per meshlet than the same code as a kernel. This tool appends +to a translated _b.metal (dmsh_meshstage.py output) + + main0_k / main0_k_z the mesh wrapper as a kernel: one threadgroup a meshlet slot, vertices to + dm_kv[slot * MAXV + n], index triples to dm_ki[slot * MAXP * 3 + n] + (absolute vertex numbers; degenerate past the primitive count) + main0_kv pass-through vertex function: dm_kv[vertex_id], every user output a float4 + (Metal links a vertex output to a fragment input only when their types are + equal), the locations no output writes 0 + main0_tk (task shaders) the object wrapper as a kernel: payload to dm_pl[tg], + launch count to dm_cnt[tg] + dm_kinfo___ an empty kernel the host reads the limits from + +Buffers (compute table): 24 dm_moff (task path: exclusive offsets of the launch counts, n + 1), +25 dm_cnt, 26 dm_pl, 27 dm_kv, 28 dm_ki, 29 dm_draw {count, first, ntask}. Vertex function: 27 dm_kv. +Not taken: per-primitive outputs, non-triangle topology (those draws stay native). + +The host (nvn_mtl.m, mk_*; DM_NVN_MESH_COMPUTE=1) loads _bk.metallib from the DM_MSL_DIR overlay. Build them for a game with + + python3 dm/gfx/tools/dmsh_meshcompute.py build /Contents/Resources/dm/shader_map.tsv + +which writes /_bk.metallib for every mesh and task shader of the map; point DM_MSL_DIR at a folder holding + as b-/ (and the b-.img of the app beside it). +""" +import re +import sys + +KV, KI, DRAW, PL, CNT, MOFF = 27, 28, 29, 26, 25, 24 + + +def _func(s, head): + i = s.find(head) + if i < 0: return None, None + j = s.find('\n{', i); depth = 0; k = j + 1 + while True: + c = s[k] + if c == '{': depth += 1 + elif c == '}': + depth -= 1 + if not depth: return i, k + 1 + k += 1 + + +def mesh_kernels(s): + m = re.search(r'metal::mesh', s) + if not m: return None, 'no mesh type' + if m.group(1) != 'void': return None, 'per-primitive outputs' + if m.group(4) != 'triangle': return None, 'topology ' + m.group(4) + maxv, maxp = int(m.group(2)), int(m.group(3)) + out = [f'\n// ---- mesh as compute (genk.py): MAXV {maxv} MAXP {maxp}\n'] + task = None + for name, kname in (('main0', 'main0_k'), ('main0_z', 'main0_k_z')): + a, b = _func(s, f'[[mesh]] void {name}(') + if a is None: continue + f = s[a:b] + sig_end = f.find(')\n{') + sig, body = f[:sig_end], f[sig_end:] + task = 'dm_payload_t& dm_payload' in sig + sig = sig.replace(f'[[mesh]] void {name}(', f'kernel void {kname}(') + sig = re.sub(r',\s*dm_mesh_t dm_out', '', sig) + sig = re.sub(r',\s*const object_data dm_(?:payload|block)_t& dm_(?:payload|block) \[\[payload\]\]', '', sig) + if task: + sig = sig.replace('uint dm_tg [[threadgroup_position_in_grid]]', 'uint dm_gtg [[threadgroup_position_in_grid]]') + sig += (f', device dm_kvrec_t* dm_kv [[buffer({KV})]], device uint* dm_ki [[buffer({KI})]], constant uint* dm_draw [[buffer({DRAW})]]') + if task: sig += f', const device dm_payload_t* dm_pl [[buffer({PL})]], const device uint* dm_moff [[buffer({MOFF})]]' + if task: + pre = ('\n uint dm_lo = 0u, dm_hi = dm_draw[2];\n' + ' while (dm_hi - dm_lo > 1u) { uint md = (dm_lo + dm_hi) >> 1; if (dm_moff[md] <= dm_gtg) dm_lo = md; else dm_hi = md; }\n' + ' uint dm_tg = dm_gtg - dm_moff[dm_lo]; uint dm_slot = dm_gtg;\n' + ' const device dm_payload_t& dm_payload = dm_pl[dm_lo];\n') + else: + pre = '\n uint dm_slot = dm_tg;\n' + body = body.replace('dm_block.first + dm_tg', 'dm_draw[1] + dm_tg') + body = body.replace(')\n{', ')\n{' + pre, 1) + body = re.sub(r'\n\s*if \(dm_lid == 0u\) dm_out\.set_primitive_count\(np\);', '', body) + body = body.replace('dm_out.set_vertex(n, v);', 'dm_kv[dm_slot * %du + n] = dm_kv_pack(v);' % maxv) + # index loop: run over every index slot, degenerate past the primitive count + lm = re.search(r'for \(uint n = dm_lid; n < np \* 3u; n \+= (\d+)u\)', body) + if not lm: return None, 'index loop' + body = body.replace(lm.group(0), 'for (uint n = dm_lid; n < %du; n += %su)' % (maxp * 3, lm.group(1))) + im = re.search(r'dm_out\.set_index\(n, (.*)\);', body) + if not im: return None, 'set_index' + body = body.replace(im.group(0), 'dm_ki[dm_slot * %du + n] = dm_slot * %du + (n < np * 3u ? uint(%s) : 0u);' % (maxp * 3, maxv, im.group(1))) + if 'dm_out' in body or 'dm_block' in body: return None, 'leftover mesh output use' + out.append(sig + body + '\n') + if len(out) < 2: return None, 'no mesh entry' + # record type and packer from dm_vertex_t + vm = re.search(r'struct dm_vertex_t \{(.*?)\};', s, re.S) + fields = re.findall(r'^\s*(\w+) (\w+) \[\[[^\]]*\]\];', vm.group(1), re.M) + rec = ['struct dm_kvrec_t {'] + for t, n in fields: + pt = t if t in ('float', 'int', 'uint', 'half', 'short', 'ushort') else 'packed_' + t + rec.append(f' {pt} {n};') + rec.append('};') + pack = ['static inline dm_kvrec_t dm_kv_pack(thread const dm_vertex_t& v) { dm_kvrec_t r;'] + [f' r.{n} = v.{n};' for _, n in fields] + [' return r; }'] + # Metal links a vertex output to a fragment input only when their types are equal, where a mesh output may be narrower (Metal accepted float3 into float4 for mesh + # pipelines). The game's fragment shaders take every input as float4 (SPIR-V Cross), so each user output is a float4 here, its missing components 0 + attrs = dict(re.findall(r'^\s*\w+ (\w+) (\[\[[^\]]*\]\]);', vm.group(1), re.M)) + out_t = ['struct dm_kvout_t {'] + [f' {"float4" if n != "position" else t} {n} {attrs[n]};' for t, n in fields] + ['};'] + have = {int(x) for x in re.findall(r'user\(locn(\d+)\)', vm.group(1))} + pad = [i for i in range(16) if i not in have] # a fragment input no output feeds fails the link: those are 0 + out_t = out_t[:-1] + [f' float4 dm_z{i} [[user(locn{i})]];' for i in pad] + ['};'] + def widen(t, n): + d = int(re.fullmatch(r'\w+?([234]?)', t).group(1) or 1) + return f'float4(float{d if d > 1 else ""}(r.{n})' + ', 0.0' * (4 - d) + ')' if d < 4 else f'float4(r.{n})' + vs = out_t + [f'vertex dm_kvout_t main0_kv(uint dm_vid [[vertex_id]], const device dm_kvrec_t* dm_kv [[buffer({KV})]])', '{', + ' dm_kvrec_t r = dm_kv[dm_vid]; dm_kvout_t v;'] + [f' v.{n} = {widen(t, n)};' for t, n in fields] + [f' v.dm_z{i} = float4(0.0);' for i in pad] + [' return v;', '}'] + size = {'float': 4, 'int': 4, 'uint': 4} + rb = 0 + for t, _ in fields: + m2 = re.fullmatch(r'(float|int|uint)([234]?)', t) + if not m2: return None, 'vertex field type ' + t + rb += 4 * int(m2.group(2) or 1) + # the host reads the limits from this entry's name: dm_kinfo___ + info = f'kernel void dm_kinfo_{maxv}_{maxp}_{rb}() {{}}\n' + return '\n'.join(rec + pack) + '\n' + ''.join(out) + '\n'.join(vs) + '\n' + info, f'{"task" if task else "notask"} {maxv} {maxp} rec {rb} bytes' + + +def task_kernel(s): + a, b = _func(s, '[[object]] void main0(') + if a is None: return None, 'no object entry' + f = s[a:b] + f = f.replace('[[object]] void main0(', 'kernel void main0_tk(', 1) + f = re.sub(r',\s*object_data dm_payload_t& dm_payload \[\[payload\]\], mesh_grid_properties dm_grid', '', f) + f = f.replace('constant uint* dm_draw [[buffer(29)]])', f'constant uint* dm_draw [[buffer({DRAW})]], device dm_payload_t* dm_pl [[buffer({PL})]], device uint* dm_cnt [[buffer({CNT})]])') + f = f.replace('dm_payload.w[', 'dm_pl[dm_tg].w[') + gm = re.search(r'dm_grid\.set_threadgroups_per_grid\(uint3\((.*), 1u, 1u\)\);', f) + if not gm: return None, 'grid' + f = f.replace(gm.group(0), f'dm_cnt[dm_tg] = {gm.group(1)};') + if 'dm_grid' in f or 'dm_payload.' in f: return None, 'leftover' + return '\n// ---- task as compute (genk.py)\n' + f + '\n', 'task' + + + +def _compile(job): + import subprocess, os + k, text, out = job + air = os.path.join(out, k + '.air') + p = subprocess.run(['xcrun', '-sdk', 'macosx', 'metal', '-c', '-O1', '-fpreserve-invariance', '-fwrapv', '-fbracket-depth=1024', '-x', 'metal', '-', '-o', air], + input=text.encode(), capture_output=True) + if p.returncode == 0: + p = subprocess.run(['xcrun', '-sdk', 'macosx', 'metallib', air, '-o', os.path.join(out, k + '.metallib')], capture_output=True) + if os.path.exists(air): os.remove(air) + return k, p.returncode, p.stderr.decode()[-400:] + + +def build(shader_map, lib, out): + import csv, os, pathlib, collections + from concurrent.futures import ProcessPoolExecutor + sys.path.insert(0, str(pathlib.Path(__file__).resolve().parents[2] / 'tools')) + import shader_lib + v = pathlib.Path(lib) if os.path.isdir(lib) else shader_lib.view(lib) + rows = [r for r in csv.reader(open(shader_map), delimiter='\t') if len(r) > 3 and r[2] == 'mesh'] + keys = sorted({r[3] for r in rows}) + os.makedirs(out, exist_ok=True) + jobs, notes = [], collections.Counter() + for k in keys: + try: s = (v / f'{k}_b.metal').read_text() + except Exception: notes['no source'] += 1; continue + is_task = '[[object]] void main0(' in s and '[[mesh]] void main0(' not in s + t, note = (task_kernel if is_task else mesh_kernels)(s) + notes[note.split(' rec ')[0] if t else 'skip: ' + note] += 1 + if t: jobs.append((k + '_bk', s + t, out)) + failed = 0 + with ProcessPoolExecutor() as ex: + for k, rc, err in ex.map(_compile, jobs): + if rc: failed += 1; print('FAILED', k, err, file=sys.stderr) + for n, c in sorted(notes.items()): print(f'{c:5d} {n}') + print(f'compiled {len(jobs) - failed}, failed {failed}') + return 1 if failed else 0 + + +if __name__ == '__main__': + import sys + if len(sys.argv) == 5 and sys.argv[1] == 'build': sys.exit(build(*sys.argv[2:])) + if len(sys.argv) != 4 or sys.argv[2] not in ('mesh', 'task'): + sys.exit('usage: dmsh_meshcompute.py build \n dmsh_meshcompute.py _b.metal mesh|task ') + s = open(sys.argv[1]).read() + t, note = (mesh_kernels if sys.argv[2] == 'mesh' else task_kernel)(s) + print(note, file=sys.stderr) + if t: open(sys.argv[3], 'w').write(s + t) diff --git a/docs/env_switches.tsv b/docs/env_switches.tsv index 435f71eebd..c1e4f5e64d 100644 --- a/docs/env_switches.tsv +++ b/docs/env_switches.tsv @@ -177,6 +177,10 @@ DM_NVN_LAZY_CLEAR switch `0` makes clears with no draw after them open a pass of DM_NVN_LOADPROBE diagnostic `=1`, `2` or `3`: attachments whose load is skipped are cleared to a sentinel colour (`3` NaN; `1` NaN for float formats) and the sentinel pixels left at pass end are counted (should be 0) DM_NVN_LOG diagnostic `=1` logs the creation of NVN objects DM_NVN_MAGSCAN diagnostic `=[:every[:phase]]`: device self-check that reads each pass's colour attachments every `every` frames (default 30) and logs `[mag]` lines for solid magenta 16×16 blocks or non-finite pixels +DM_NVN_MESH_COMPUTE diagnostic `=1` runs task and mesh shaders as compute kernels before their render pass and draws their triangles with a plain indexed draw, for GPUs without mesh-shading hardware (Apple M1 and M2); needs the shaders' side libraries (`_bk.metallib`, dm/gfx/tools/dmsh_meshcompute.py) in the DM_MSL_DIR overlay; `ab` alternates it with the DM_AB phases; off by default while it is measured +DM_NVN_MK_EQ diagnostic `=1` keeps a mesh shader native once one of its draws has an equal depth test (with DM_NVN_MESH_COMPUTE); it made Ocarina of Time's grass worse, since its depth prepass draws with another shader +DM_NVN_MK_IND diagnostic which indirect task / mesh draws DM_NVN_MESH_COMPUTE converts: `0` none, `task` only those with a task shader, `notask` only those without (all by default) +DM_NVN_MK_UBSTOP diagnostic `=1` ends DM_NVN_MESH_COMPUTE's look-ahead at the first uniform-buffer update of the pass (it passes them by default and keeps native the draws that read an updated buffer) DM_NVN_PARTIAL_UPLOAD switch `0` reads a texture whole from its storage at every re-upload; by default a re-upload reads only the images (a layer's level) with a page written since the last upload and copies the others from the texture it replaces DM_NVN_PASSDUMP diagnostic `=`, `-` or `every:`: logs every render pass of those frames (attachments, clears, shader keys, bindings, draw state, presents, and the host function and line that ended each pass) for the `tbdr_*.py` tools DM_NVN_PASSDUMP_ALL diagnostic `=1` makes `DM_NVN_PASSDUMP` record every draw and binding change, not only the first per pass and stage, including copies and inline uniform updates From 2cbf2c4dfeeeb633ee4e5694670128883686b8ac Mon Sep 17 00:00:00 2001 From: Adya <114347520+AdyaGMD@users.noreply.github.com> Date: Thu, 12 Nov 2026 10:34:18 +0000 Subject: [PATCH 2/3] Mesh compute on by default on a Mac without mesh-shading hardware (M1, M2), its side libraries built into the app, and the branch compiling on main DM_NVN_MESH_COMPUTE now defaults to on where the GPU lacks MTLGPUFamilyApple9 on macOS (the M1 and M2); =0 turns it off, and it stays off elsewhere (the iPhone is not measured). The side libraries come with the app instead of a DM_MSL_DIR overlay: a mac-app build (build_ipa_from_nsp.py mesh_side_libs) runs dmsh_meshcompute.py build on the Mac library into /shader/mk-b-, reused while the library's image and the generator are unchanged, and pack_app.py --msl-extra adds them to the app's shader image (shader_lib.subset extra); a generator failure leaves them out, and those draws stay native. The first commit did not compile on main: mk_draw_writes used draw_note_writes and its counters, and the scan's target check used csh_rt_same, both from a working tree with other changes. Now mk_draw_writes is prog_swrites (any shader of the draw stores to a buffer or image), and a target change ends the scan. The measurement in the first commit was made with the working tree's versions; the scan may now stop sooner in a pass that changes targets. Checked: the host compiles. Not run: this branch's host in game. Co-Authored-By: Claude Opus 5.5 --- dm/gfx/nvn/nvn_mtl.m | 20 ++++++++++---------- dm/mac/pack_app.py | 15 ++++++++++----- dm/tools/shader_lib.py | 4 ++-- docs/env_switches.tsv | 2 +- scripts/build_ipa_from_nsp.py | 18 ++++++++++++++++++ 5 files changed, 41 insertions(+), 18 deletions(-) diff --git a/dm/gfx/nvn/nvn_mtl.m b/dm/gfx/nvn/nvn_mtl.m index 149813145c..2437948a58 100644 --- a/dm/gfx/nvn/nvn_mtl.m +++ b/dm/gfx/nvn/nvn_mtl.m @@ -13974,7 +13974,13 @@ static void cs_hoist_try(NvProg *pr0, const p_draw *dw0, const p_draw_idx *di0) // apart (window state, the time into the session), more than the gain measured static int mk_on(void) { static int m = -1; - if (m < 0) { m = tbdr_on() ? 0 : sw_mode("DM_NVN_MESH_COMPUTE", 0); if (m) MLOG("DM_NVN_MESH_COMPUTE: task / mesh draws as compute before their pass (GPU family Apple9: %d)", (int)[g_dev supportsFamily:MTLGPUFamilyApple9]); } + if (m < 0) { +#if TARGET_OS_IPHONE + int dflt = 0; // not measured there +#else + int dflt = ![g_dev supportsFamily:MTLGPUFamilyApple9]; // on by default on a Mac without mesh-shading hardware (M1, M2); =0 turns it off +#endif + m = tbdr_on() ? 0 : sw_mode("DM_NVN_MESH_COMPUTE", dflt); if (m) MLOG("DM_NVN_MESH_COMPUTE: task / mesh draws as compute before their pass (GPU family Apple9: %d)", (int)[g_dev supportsFamily:MTLGPUFamilyApple9]); } return sw_on(m); } @interface NvMk : NSObject { @public id k[2], t[2]; id kv; uint32_t maxv, maxp, rec; int ok; } @end // k: mesh kernel (border corrections on / off), t: task kernel, kv: the vertex function @@ -14092,14 +14098,8 @@ static void mk_note(int ind) { MLOG("mesh as compute up to frame %llu: %llu task / mesh draws, %llu indirect, %llu indirect draws cut to their slots, %llu drawn from compute, %llu native; %llu jobs in %llu scans; scans stopped by a draw writing memory %llu, a uniform update %llu, a dispatch %llu, another op %llu, a full table %llu, other targets %llu; left native: indirect %llu, count %llu, shaders %llu, no mesh library %llu, no task library %llu, size %llu, arguments %llu; mesh draws met by scans %llu", (unsigned long long)f, g_mk_all, g_mk_ind, g_mk_over, g_mk_drawn, g_mk_native, g_mk_jobs, g_mk_scans, g_mk_stop[0], g_mk_stop[1], g_mk_stop[2], g_mk_stop[3], g_mk_stop[4], g_mk_stop[5], g_mk_rej[0], g_mk_rej[1], g_mk_rej[2], g_mk_rej[3], g_mk_rej[4], g_mk_rej[5], g_mk_rej[6], g_mk_met); } -// Whether this draw may write memory (draw_note_writes, without keeping its note) -static int mk_draw_writes(NvProg *pr) { - int w0 = g_pass_memw, n0 = g_mpw_n; uint64_t k0[6]; memcpy(k0, g_pw_kind, sizeof k0); - g_pass_memw = 0; draw_note_writes(pr); - int w = g_pass_memw || g_mpw_n > n0; - g_pass_memw = w0; g_mpw_n = n0; memcpy(g_pw_kind, k0, sizeof k0); - return w; -} +// Whether this draw may write memory: any of its shaders stores to a buffer or image (prog_swrites) +static int mk_draw_writes(NvProg *pr) { return prog_swrites(pr); } static int mk_ub_hit(NvShader *sh, int st, const csh_rng *up, int nup) { if (!sh || !nup) return 0; for (int i = 0; i < sh->rf.nb; i++) { const nvn_rbind *b = &sh->rf.b[i]; @@ -14243,7 +14243,7 @@ static void mk_hoist_try(const uint8_t *from) { case OP_BARRIER: case OP_RENDER_COND: break; case OP_REPORT_COUNTER: case OP_RESET_COUNTER: { uint64_t t; memcpy(&t, aq, 8); if (t != NVN_COUNTER_TYPE_SAMPLES_PASSED) stop = 3; break; } case OP_DISPATCH: { int moved = 0; for (int i = 0; i < g_csh_n; i++) moved |= g_csh_op[i] == q; if (!moved) stop = 2; break; } // one cs_hoist_try ran ahead already ran before these jobs - case OP_SET_RT: if (!csh_rt_same(q, aq)) stop = 5; break; + case OP_SET_RT: stop = 5; break; // a target change ends the pass's scan default: if (!state_op(hq.op, aq)) stop = 3; break; } } diff --git a/dm/mac/pack_app.py b/dm/mac/pack_app.py index 062c2cb775..2d413f472f 100755 --- a/dm/mac/pack_app.py +++ b/dm/mac/pack_app.py @@ -195,9 +195,13 @@ def msl_inputs(msl): return files if any(p.suffix == '.metallib' for p in files) else [] -def copy_msl(msl, msls, dst, clone): - """The shader libraries to dst, msl/ in the app: an NVN game's as the one file dst.img, of what the host reads. Writing it costs the space of a copy, not a clone, as it leaves out the Mac build's sources""" - if is_chain_library(msl): shader_lib.subset(msl, f'{dst}.img'); return +def msl_extra(a): # --msl-extra: side libraries (_bk.metallib, mesh shaders as compute) that go into the chain library's image beside its own files + return sorted(Path(a.msl_extra).glob('*.metallib')) if getattr(a, 'msl_extra', None) and is_chain_library(a.msl) else [] + + +def copy_msl(msl, msls, dst, clone, extra=()): + """The shader libraries to dst, msl/ in the app: an NVN game's as the one file dst.img, of what the host reads (and the extra files). Writing it costs the space of a copy, not a clone, as it leaves out the Mac build's sources""" + if is_chain_library(msl): shader_lib.subset(msl, f'{dst}.img', extra={p.name: str(p) for p in extra}); return dst.mkdir(parents=True); copy(msls, dst, clone) @@ -220,7 +224,7 @@ def selected_inputs(a): tables = [Path(a.tables) / n for n in TABLES if (Path(a.tables) / n).is_file()] pairs = Path(a.pso_pairs) if a.pso_pairs and not any(p.name == 'pso_pairs.tsv' for p in tables) else None fonts = sorted(p for p in Path(a.font).glob('*.ttf') if not detritus(p.name)) - files = [Path(a.host)] + mods + msls + pso + tables + fonts + ([pairs] if pairs else []) + files = [Path(a.host)] + mods + msls + msl_extra(a) + pso + tables + fonts + ([pairs] if pairs else []) files += [Path(p) for p in (a.env_file, a.settings, a.icon) if p] for p in files: if not p.is_file(): raise SystemExit(f'no input file at {p}') @@ -323,7 +327,7 @@ def pack(a): if fonts: copy(fonts, dm / 'font', False) if a.env_file: shutil.copyfile(a.env_file, dm / 'env.txt') if a.settings: shutil.copyfile(a.settings, dm / 'launcher.json') - step('shader libraries'); msl_dst = res / 'msl' / (a.msl_name or msl.name); msl_dst.parent.mkdir(exist_ok=True); copy_msl(msl, msls, msl_dst, clone['msl']) + step('shader libraries'); msl_dst = res / 'msl' / (a.msl_name or msl.name); msl_dst.parent.mkdir(exist_ok=True); copy_msl(msl, msls, msl_dst, clone['msl'], msl_extra(a)) if aoc: step('add-on content'); copy([aoc], res / 'aoc', clone['aoc']) if pso: step('pipeline archive'); (res / 'pso').mkdir(); copy(pso, res / 'pso', clone['pso']) step('game files'); copy([romfs], res / ('romfs.img' if romfs.is_file() else 'romfs'), clone['romfs']) @@ -373,6 +377,7 @@ def parse(argv): ap.add_argument('--patched', required=True, help='the rewritten modules (patch_nso.py output)') ap.add_argument('--romfs', required=True, help='the RomFS of the packed version') ap.add_argument('--msl', required=True, help='the Mac shader libraries: /cache/b-, or a Vulkan or GL game\'s /vk or gl') + ap.add_argument('--msl-extra', help='a folder of side libraries (*.metallib: mesh shaders as compute, dmsh_meshcompute.py) to add to the chain library\'s image') ap.add_argument('--msl-name', help='its folder name in the app (default: its own)') ap.add_argument('--tables', required=True, help='the game profile folder (shader_map.tsv and the tables beside it)') ap.add_argument('--font', required=True, help='the folder of the shared-font stand-ins (nintendo_ext.ttf)') diff --git a/dm/tools/shader_lib.py b/dm/tools/shader_lib.py index ed6f105945..4725794a0e 100644 --- a/dm/tools/shader_lib.py +++ b/dm/tools/shader_lib.py @@ -156,8 +156,8 @@ def app_files(lib, skip_suffix=()): return {p.name: str(p) for p in view(lib).iterdir() if p.name.endswith(APP_SUFFIXES) and p.name not in APP_LEAVE_OUT and not p.name.startswith('.') and '.tmp' not in p.name and not p.name.endswith(tuple(skip_suffix) or ('\0',))} -def subset(lib, dst, skip_suffix=()): # → the number of files in the image dst - files = app_files(lib, skip_suffix) +def subset(lib, dst, skip_suffix=(), extra=None): # → the number of files in the image dst; extra: {name: path} added (side libraries the build made beside the library) + files = app_files(lib, skip_suffix) | dict(extra or {}) if not any(n.endswith('.metallib') for n in files): raise ValueError(f'no shader libraries at {lib}') write(files, os.fspath(dst)) return len(files) diff --git a/docs/env_switches.tsv b/docs/env_switches.tsv index c1e4f5e64d..51eb917861 100644 --- a/docs/env_switches.tsv +++ b/docs/env_switches.tsv @@ -177,7 +177,7 @@ DM_NVN_LAZY_CLEAR switch `0` makes clears with no draw after them open a pass of DM_NVN_LOADPROBE diagnostic `=1`, `2` or `3`: attachments whose load is skipped are cleared to a sentinel colour (`3` NaN; `1` NaN for float formats) and the sentinel pixels left at pass end are counted (should be 0) DM_NVN_LOG diagnostic `=1` logs the creation of NVN objects DM_NVN_MAGSCAN diagnostic `=[:every[:phase]]`: device self-check that reads each pass's colour attachments every `every` frames (default 30) and logs `[mag]` lines for solid magenta 16×16 blocks or non-finite pixels -DM_NVN_MESH_COMPUTE diagnostic `=1` runs task and mesh shaders as compute kernels before their render pass and draws their triangles with a plain indexed draw, for GPUs without mesh-shading hardware (Apple M1 and M2); needs the shaders' side libraries (`_bk.metallib`, dm/gfx/tools/dmsh_meshcompute.py) in the DM_MSL_DIR overlay; `ab` alternates it with the DM_AB phases; off by default while it is measured +DM_NVN_MESH_COMPUTE switch `=1` runs task and mesh shaders as compute kernels before their render pass and draws their triangles with a plain indexed draw, for GPUs without mesh-shading hardware (Apple M1 and M2); needs the shaders' side libraries (`_bk.metallib`, dm/gfx/tools/dmsh_meshcompute.py), which a Mac app build adds to its shader image (a draw without one stays native); `ab` alternates it with the DM_AB phases; on by default on a Mac without mesh-shading hardware, `=0` turns it off; off elsewhere DM_NVN_MK_EQ diagnostic `=1` keeps a mesh shader native once one of its draws has an equal depth test (with DM_NVN_MESH_COMPUTE); it made Ocarina of Time's grass worse, since its depth prepass draws with another shader DM_NVN_MK_IND diagnostic which indirect task / mesh draws DM_NVN_MESH_COMPUTE converts: `0` none, `task` only those with a task shader, `notask` only those without (all by default) DM_NVN_MK_UBSTOP diagnostic `=1` ends DM_NVN_MESH_COMPUTE's look-ahead at the first uniform-buffer update of the pass (it passes them by default and keeps native the draws that read an updated buffer) diff --git a/scripts/build_ipa_from_nsp.py b/scripts/build_ipa_from_nsp.py index 31e2e27756..189abd1c6a 100755 --- a/scripts/build_ipa_from_nsp.py +++ b/scripts/build_ipa_from_nsp.py @@ -2752,6 +2752,22 @@ def mac_app_name(name): # the app's file name: the game name with the characte return re.sub(r'\s*[:/]\s*', ' - ', name).strip(' .-') or 'nxnative' +def mesh_side_libs(py, repo, work, lib, g, env, logs): + """Task and mesh shaders as compute, for Macs without mesh-shading hardware (M1, M2; the host turns it on there, DM_NVN_MESH_COMPUTE): a side library _bk.metallib for each, + generated from the library's translated sources by dm/gfx/tools/dmsh_meshcompute.py into /shader/mk-b-, which pack_app adds to the app's shader image (--msl-extra). + Reused while the library's image and the generator are unchanged. A failure leaves them out (those draws stay native) and does not stop the build""" + out, tool = work / 'shader' / f'mk-{lib.name}', repo / 'dm/gfx/tools/dmsh_meshcompute.py' + src = Path(shader_lib.image_of(lib)); src = src if src.is_file() else Path(lib) + if not src.exists(): return None + st = src.stat(); key = f'{st.st_size} {st.st_mtime_ns} {hashlib.sha256(tool.read_bytes()).hexdigest()}\n' + kf = out / 'key' + if kf.is_file() and kf.read_text() == key and any(out.glob('*_bk.metallib')): return out + rmtree(out); PROG.set(step='mesh shaders as compute') + if run([py, tool, 'build', g.table('shader_map.tsv'), lib, out], logs / 'mesh_compute.log', env=env, fatal=False): say('mesh shaders as compute: some side libraries failed (mesh_compute.log); those draws stay native') + if not any(out.glob('*_bk.metallib')): return None + kf.write_text(key); return out + + def mac_app(py, repo, out, app, work, game, patched, host, font, msl, aoc, gfx, lang, g, name, version, bundle_id, icon, sign, env, logs, env_file='', gpus='', programs=None, control=None): """The last stage of PLATFORM 'mac-app': pack the Mac app (dm/mac/pack_app.py) from what a Mac-only build made: the Mac host and font (mac_host), the packed version's RomFS and rewritten modules, the Mac shader libraries (msl: /cache, whose b- goes in, or a Vulkan or GL game's , whose vk / gl goes in), the DLC, the profile's tables (with the shader pairs @@ -2761,11 +2777,13 @@ def mac_app(py, repo, out, app, work, game, patched, host, font, msl, aoc, gfx, lib = msl / (f'b-{chain_of(repo)}' if gfx == 'nvn' else 'vk' if gfx == 'vulkan' else 'gl') pso = mac_pso(py, repo, work, lib, g, gpus, env, logs) if gfx == 'nvn' and g.table('shader_pairs.tsv').is_file() else None # a stage of its own, before the app's PROG.stage('Mac app', step='starting') + mk = mesh_side_libs(py, repo, work, lib, g, env, logs) if gfx == 'nvn' else None # The launcher settings the IPA carries (dm/games/settings.py), without the iPhone's own switches: the app's settings window and Mac front ends show them (docs/frontends.md) settings = work / 'app/launcher.json'; settings.parent.mkdir(parents=True, exist_ok=True) run([py, repo / 'dm/games/settings.py', g.id, patched, control or game / 'control.nacp', settings], logs / 'mac_settings.log', env=dict(env, GFXAPI=gfx, SETTINGS_PLATFORM='mac', DM_LANG=lang)) cmd = [py, repo / 'dm/mac/pack_app.py', out, '--host', host, '--patched', patched, '--romfs', rfs.source(game), '--msl', lib, '--tables', g.table('shader_map.tsv').parent, '--font', font, '--title-id', f'{app:016X}', '--name', name, '--bundle-id', bundle_id, '--lang', lang, '--sign', sign or '-'] + if mk: cmd += ['--msl-extra', mk] split_pairs = work / 'shader/pso_pairs.tsv' # the split's copy of pso_pairs.tsv, for a profile without the table (as mac_run's DM_PSO_PAIRS) for opt, v in (('--version', version), ('--category', g.get('app_category')), ('--marker', g.get('romfs_marker')), ('--aoc', aoc), ('--icon', icon), ('--pso', pso), ('--pso-pairs', split_pairs if gfx == 'nvn' and not g.table('pso_pairs.tsv').is_file() and split_pairs.is_file() else None), From 9bed004ae65ecbbb3e0dfb10f01b7a48758e46bb Mon Sep 17 00:00:00 2001 From: Adya <114347520+AdyaGMD@users.noreply.github.com> Date: Thu, 12 Nov 2026 10:48:35 +0000 Subject: [PATCH 3/3] dmsh_meshcompute.py build: task shaders too build took only the shader map's rows of stage mesh, and the map names a task shader's stage task, so no task shader got its kernel and every task draw stayed native ("no task library": in Ocarina of Time 47,746 of 90,656 mesh-stage draws in 2,100 frames, a third converting instead of 82 %). The measurement in the first commit used side libraries made before this tool, which had them. Co-Authored-By: Claude Opus 5.5 --- dm/gfx/tools/dmsh_meshcompute.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dm/gfx/tools/dmsh_meshcompute.py b/dm/gfx/tools/dmsh_meshcompute.py index 78bb2d5a1e..ca4eb6fda5 100644 --- a/dm/gfx/tools/dmsh_meshcompute.py +++ b/dm/gfx/tools/dmsh_meshcompute.py @@ -152,7 +152,7 @@ def build(shader_map, lib, out): sys.path.insert(0, str(pathlib.Path(__file__).resolve().parents[2] / 'tools')) import shader_lib v = pathlib.Path(lib) if os.path.isdir(lib) else shader_lib.view(lib) - rows = [r for r in csv.reader(open(shader_map), delimiter='\t') if len(r) > 3 and r[2] == 'mesh'] + rows = [r for r in csv.reader(open(shader_map), delimiter='\t') if len(r) > 3 and r[2] in ('mesh', 'task')] # the map names task shaders' stage task keys = sorted({r[3] for r in rows}) os.makedirs(out, exist_ok=True) jobs, notes = [], collections.Counter()