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..2437948a58 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) { +#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 +@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: 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]; + 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: stop = 5; break; // a target change ends the pass's scan + 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..ca4eb6fda5 --- /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] 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() + 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/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 435f71eebd..51eb917861 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 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) 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 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),