Monster Hunter Rise in Kamura Village: indirect-count multi-draws masked on the GPU instead of a wait for the count (36–54 fps on an iPhone), and the sparse texture layout in 64 KB tiles (magenta clothes)
- Indirect-count multi-draws: a compute pass copies each draw's record into the frame arena, zeros past the GPU-written count, and the draw reads the copy (nvn_mtl.m mdx_run); copies are keyed by record and count address, and back-to-back draws of one multi-draw are encoded alone (mdx_next). The replay used to commit and wait for the GPU before each GPU-driven pass: 14 of the 20 ms a frame the iPhone's submitting thread took. On the Mac the village submit went from 20 ms of wall time (6.1 ms of work) to 5.1 ms, all of it work. - Sparse textures (flag SPARSE): tiled levels are laid out with their width padded to whole 64 KB tiles, the mip tail starts at the first level smaller than a tile, storage in whole tiles, 64 KB aligned (nvn_tex.c level_walk, nvn_tex_level_gobs; the deswizzlers address rows by it). The game's streamed textures come pre-swizzled that way: unpadded, Hinoa's and Minoto's 1024×1024 ASTC 6×6 clothes read thousands of all-zero blocks, decoded as ASTC's magenta error colour. - DM_NVN_ZBLOCK: counts all-zero ASTC blocks in uploaded textures' bound levels, and dumps their raw storage.
This commit is contained in:
seedds committed
1 parent
e4a88c614a
commit
f0ff511b4a
9 files changed
+209
-32
No files matched your search
@@ -2011,7 +2011,7 @@ NVN(nvnCommandBufferDrawArraysInstanced, void, NVNcommandBuffer *c, NVNdrawPrimi
|
||||
// Indirect, as DrawArraysIndirect: the record at ind is read by the GPU (nvn_mtl.m draw)
|
||||
// Indirect-count multi-draws (Monster Hunter Rise's GPU-driven scene: compute culling writes the draw records and a u32 count, then up to max draws are made from them; 1–49 per call,
|
||||
// stride 20, the records and counts in its sparse virtual pool). Metal has no draw count read by the GPU, so each multi-draw becomes max ordinary indirect draws that carry the count's address
|
||||
// and their index; the replay reads the count once the GPU has written it and skips draws past it (nvn_mtl.m mdi_live). Both used to be empty, and the whole scene stayed undrawn
|
||||
// and their index; the replay masks their records on the GPU, zeros past the count (nvn_mtl.m mdx_run). Both used to be empty, and the whole scene stayed undrawn
|
||||
enum { MDI_MAX = 1024 };
|
||||
NVN(nvnCommandBufferMultiDrawElementsIndirectCount, void, NVNcommandBuffer *c, NVNdrawPrimitive pr, NVNindexType it, NVNbufferAddress ib, NVNbufferAddress ind, NVNbufferAddress cnt, int max, ptrdiff_t stride) {
|
||||
if (max > MDI_MAX) { ONCE_LOG("[RED] indirect-count multi-draw of up to %d draws: only the first %d are made", max, MDI_MAX); max = MDI_MAX; }
|
||||
@@ -2155,10 +2155,11 @@ static int write_region_bl(const tex_t *t, unsigned level, unsigned layer, const
|
||||
uint32_t x0 = (uint32_t)(r->xoffset / bw) * (uint32_t)bpb, y0 = (uint32_t)(r->yoffset / bh);
|
||||
size_t wb = (size_t)((r->width + bw - 1) / bw) * (size_t)bpb, hb = (size_t)((r->height + bh - 1) / bh);
|
||||
if (x0 + wb > row_bytes || y0 + hb > rows) return -1;
|
||||
uint32_t lb = nvn_tex_level_gobs(&t->d, level) * 64; if (lb < row_bytes) lb = row_bytes; // the row as laid out (padded in a sparse texture's tiled levels)
|
||||
for (size_t y = 0; y < hb; y++)
|
||||
for (size_t x = 0; x < wb;) { // 16-byte runs are contiguous within a GOB
|
||||
uint32_t bx = x0 + (uint32_t)x; size_t n = 16 - bx % 16; if (n > wb - x) n = wb - x;
|
||||
memcpy(t->mem + base + nvn_bl_offset(bx, y0 + (uint32_t)y, 0, row_bytes, rows, th, td), src + y * row + x, n); x += n;
|
||||
memcpy(t->mem + base + nvn_bl_offset(bx, y0 + (uint32_t)y, 0, lb, rows, th, td), src + y * row + x, n); x += n;
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
+141
-13
@@ -751,6 +751,7 @@ static int bw_hit(uint64_t lo, uint64_t hi) { for (int i = 0; i < g_bw_n; i++) i
|
||||
static int br_hit(uint64_t lo, uint64_t hi) { for (int i = 0; i < g_br_n; i++) if (lo < g_br_hi[i] && g_br_lo[i] < hi) return 1; return 0; }
|
||||
// buffer → buffer / fill: reads [rlo, rhi) (0, 0 when nothing is read), writes [wlo, whi)
|
||||
static uint64_t g_wr_epoch; // Bumped by every buffer write outside a render pass (copy / update / fill, texture → buffer, dispatch): indirect-draw conversions made before it are stale (ixc_run)
|
||||
static uint64_t g_mdx_epoch; // Bumped by GPU writes that can change an indirect-count multi-draw's records or count (copy, fill, texture → buffer, dispatch; see mdx_run)
|
||||
static id<MTL4ComputeCommandEncoder> blit_buf(uint64_t rlo, uint64_t rhi, uint64_t wlo, uint64_t whi) {
|
||||
g_wr_epoch++;
|
||||
id<MTL4ComputeCommandEncoder> e = blit_raw();
|
||||
@@ -768,7 +769,7 @@ static id<MTL4ComputeCommandEncoder> blit(void) {
|
||||
}
|
||||
|
||||
void mtl_sub_copy(uint64_t src, uint64_t dst, uint64_t n) {
|
||||
pend_add(dst, n, 2);
|
||||
pend_add(dst, n, 2); g_mdx_epoch++;
|
||||
void *sb, *db; uint64_t so, dof;
|
||||
if (!mtl_addr(src, n, &sb, &so) || !mtl_addr(dst, n, &db, &dof))
|
||||
MDIE("buffer copy address is in no memory pool src=%#llx dst=%#llx n=%#llx", (unsigned long long)src, (unsigned long long)dst, (unsigned long long)n);
|
||||
@@ -940,7 +941,7 @@ void mtl_sub_update(const void *data, uint64_t dst, uint64_t n) {
|
||||
// In-order buffer fill within the submission: all four bytes equal → fillBuffer; otherwise copied in chunks from a 64 KiB pattern buffer
|
||||
void mtl_sub_fill(uint64_t dst, uint64_t n, uint32_t value) {
|
||||
if (!n) return;
|
||||
pend_add(dst, n, 3);
|
||||
pend_add(dst, n, 3); g_mdx_epoch++;
|
||||
void *db; uint64_t dof;
|
||||
if (!mtl_addr(dst, n, &db, &dof)) MDIE("buffer fill address is in no memory pool dst=%#llx n=%#llx", (unsigned long long)dst, (unsigned long long)n);
|
||||
uint8_t b0 = (uint8_t)value;
|
||||
@@ -1516,6 +1517,30 @@ uint64_t mtl_upload_gpu_ns_take(void) { return atomic_exchange(&g_upload_gpu_ns,
|
||||
void mtl_create_stats(uint64_t *lib, uint64_t *pso) { *lib = g_bp_lib_n; *pso = g_bp_pso_n; }
|
||||
void mtl_create_ns(uint64_t *lib_ns, uint64_t *pso_ns) { *lib_ns = g_bp_lib_ns; *pso_ns = g_bp_pso_ns; } // Cumulative time spent synchronously creating shader libraries / pipelines on the spot (the long-frame log takes the per-frame difference)
|
||||
uint64_t mtl_upload_bytes(void) { return g_bp_up_bytes; } // Cumulative texture upload bytes
|
||||
// Diagnostic DM_NVN_ZBLOCK=1: all-zero 16-byte blocks of each ASTC texture uploaded on the CPU, per level (an all-zero ASTC block decodes as the error colour, magenta). tex_mtl_ logs
|
||||
// a texture whose bound levels hold any (zb_report), with each level's zero blocks, its total and whether all its pages are mapped (sparse pools)
|
||||
static int zblock_on(void) { static int v = -1; if (v < 0) { v = getenv("DM_NVN_ZBLOCK") != NULL; if (v) MLOG("DM_NVN_ZBLOCK: counting all-zero blocks of ASTC textures at upload (diagnostic)"); } return v; }
|
||||
static uint64_t g_zb_ser; static uint32_t g_zb[16], g_zb_tot[16]; static int g_zb_n;
|
||||
static void zb_report(const tex_t *t, unsigned base, const char *why) {
|
||||
if (!zblock_on() || g_zb_ser != t->serial) return;
|
||||
int bad = 0; for (int l = (int)base; l < g_zb_n; l++) bad |= g_zb[l] != 0;
|
||||
g_zb_ser = 0; if (!bad) return;
|
||||
static int nl; if (nl++ >= 400) return;
|
||||
char b[512]; size_t o = 0; int sp = t->pool && nvn_sparse_pool(t->pool);
|
||||
for (int l = 0; l < g_zb_n && o < sizeof b - 48; l++) { uint64_t lo = nvn_tex_level_offset(&t->d, (unsigned)l), hi = l + 1 < (int)t->d.levels ? nvn_tex_level_offset(&t->d, (unsigned)l + 1) : t->L.size;
|
||||
o += (size_t)snprintf(b + o, sizeof b - o, " L%d %u/%u%s", l, g_zb[l], g_zb_tot[l], sp ? (nvn_sparse_mapped(t->pool, (uint64_t)(uintptr_t)t->mem + lo, hi - lo) ? "" : " unmapped") : ""); }
|
||||
MLOG("[zb] %s ser=%llu format %#x %ux%u levels %u flags %#x mem=%p size %#llx sparse %d bound from level %u frame %llu:%s", why, (unsigned long long)t->serial, t->d.format, t->d.width, t->d.height, t->d.levels,
|
||||
t->d.flags, t->mem, (unsigned long long)t->L.size, sp, base, (unsigned long long)atomic_load(&g_cur_frame), b);
|
||||
const char *dir = getenv("DM_NVN_ZBLOCK"); static int nd; static uint64_t seen[16]; // DM_NVN_ZBLOCK=<folder>: also the raw storage of the first such texture of each format and size (up to 16), zb_<serial>.bin (unmapped pages as 0)
|
||||
uint64_t kind = (uint64_t)t->d.format << 48 | (uint64_t)t->d.width << 24 | t->d.height; int dup = 0; for (int i = 0; i < nd; i++) dup |= seen[i] == kind;
|
||||
if (dir && dir[0] == '/' && nd < 16 && !dup) { seen[nd] = kind; uint64_t n = sp ? MAX(t->L.size, 2u << 20) : t->L.size; uint8_t *m = malloc(n); char fp[1024]; FILE *f; // a sparse texture: 2 MB, and zbm_<serial>.bin a byte per 4 KB page (1 mapped)
|
||||
if (m) { if (sp) nvn_sparse_copy(t->pool, m, (uint64_t)(uintptr_t)t->mem, n); else memcpy(m, t->mem, n);
|
||||
snprintf(fp, sizeof fp, "%s/zb_%llu.bin", dir, (unsigned long long)t->serial); if ((f = fopen(fp, "wb"))) { fwrite(m, 1, n, f); fclose(f); nd++; }
|
||||
if (sp) { for (uint64_t o = 0; o < n; o += 0x1000) m[o >> 12] = (uint8_t)nvn_sparse_mapped(t->pool, (uint64_t)(uintptr_t)t->mem + o, 0x1000);
|
||||
snprintf(fp, sizeof fp, "%s/zbm_%llu.bin", dir, (unsigned long long)t->serial); if ((f = fopen(fp, "wb"))) { fwrite(m, 1, n >> 12, f); fclose(f); } }
|
||||
free(m); }
|
||||
MLOG("[zb] ser=%llu: block height log2 %d (packaged %d), stride %u, depth %u, target %d, samples %d", (unsigned long long)t->serial, (int)t->L.tileH, t->d.pkg_block_h, t->d.stride, t->d.depth, t->d.target, t->d.samples); }
|
||||
}
|
||||
static id<MTLTexture> pkg_upload_(const tex_t *t, uint64_t *stamp, id<MTLTexture> reuse) {
|
||||
MTLTextureDescriptor *d = tex_desc(t, 1);
|
||||
if (!d) return nil;
|
||||
@@ -1576,11 +1601,15 @@ static id<MTLTexture> pkg_upload_(const tex_t *t, uint64_t *stamp, id<MTLTexture
|
||||
id<MTLBuffer> stg = [g_dev newBufferWithLength:total ? total : 16 options:MTLResourceStorageModeShared];
|
||||
res_add(stg); [g_up_keep addObject:stg];
|
||||
size_t off = 0;
|
||||
int zb = zblock_on() && t->d.format >= 0x79 && t->d.format <= 0x94; if (zb) { g_zb_ser = t->serial; g_zb_n = (int)MIN(d.mipmapLevelCount, 16u); memset(g_zb, 0, sizeof g_zb); memset(g_zb_tot, 0, sizeof g_zb_tot); }
|
||||
for (unsigned l = 0; l < d.mipmapLevelCount; l++) {
|
||||
uint32_t w = MAX(t->d.width >> l, 1u), h = MAX(t->d.height >> l, 1u), z = t->d.target == 2 ? MAX(t->d.depth >> l, 1u) : 1;
|
||||
size_t pitch = (size_t)((w + bw - 1) / bw) * bpb, rows = (h + bh - 1) / bh;
|
||||
for (unsigned y = 0; y < layers; y++) {
|
||||
if (nvn_deswizzle(&t->d, l, y, src, (uint8_t *)stg.contents + off, pitch)) { RED_ONCE(9, "deswizzle failed format %#x target %d", t->d.format, t->d.target); break; }
|
||||
if (zb && l < 16) { const uint64_t *q = (const uint64_t *)((const uint8_t *)stg.contents + off); size_t nb = pitch * rows * z / 16;
|
||||
for (size_t i = 0; i < nb; i++) g_zb[l] += !(q[2 * i] | q[2 * i + 1]);
|
||||
g_zb_tot[l] += (uint32_t)nb; }
|
||||
up_blit_add(stg, off, pitch, pitch * rows, MTLSizeMake(w, h, z), mt, y, l);
|
||||
off += pitch * rows * z;
|
||||
}
|
||||
@@ -1774,6 +1803,7 @@ static id<MTLTexture> tex_mtl_(const tex_t *ct) {
|
||||
if (nt) { id<MTLTexture> old = pk->tex; pk->tex = nt;
|
||||
if (pk->spare && pk->spare != nt) defer_obj_release(pk->spare); // still in use, or of another shape: retired as before
|
||||
if (pk->sparse) sparse_view(pk, t);
|
||||
zb_report(t, pk->sparse ? (unsigned)pk->vbase : 0, "re-upload");
|
||||
pk->spare = old; pk->spare_fence = g_fence + 1; g_ser_gen++; pk->stamp = st; pk->vgen = pk->vseen = vg; pk->span = nvn_dirty_span(t->mem, t->L.size); pk->seen = nvn_dirty_now(); g_pkg_reuploads++; if (dirtycheck() && !pk->sparse) pk->hash = full_hash(t->mem, t->L.size); }
|
||||
return pk->view ? pk->view : pk->tex;
|
||||
}
|
||||
@@ -1784,6 +1814,7 @@ static id<MTLTexture> tex_mtl_(const tex_t *ct) {
|
||||
pk->span = nvn_dirty_span(t->mem, t->L.size);
|
||||
pk->seen = nvn_dirty_now(); g_pkg_uploads++; if (dirtycheck() && !pk->sparse) pk->hash = full_hash(t->mem, t->L.size);
|
||||
if (pk->sparse) sparse_view(pk, t);
|
||||
zb_report(t, pk->sparse ? (unsigned)pk->vbase : 0, "upload");
|
||||
tex_put(t->serial, pk);
|
||||
return pk->view ? pk->view : pk->tex;
|
||||
}
|
||||
@@ -4188,6 +4219,7 @@ static id<MTL4RenderCommandEncoder> rp_begin(void);
|
||||
// not one byte of the depth content is touched, and on TBDR the load and store of the whole depth texture are saved (measured on a real device: a single-draw fullscreen pass takes half the GPU, with an extra D32S8 load and store attached).
|
||||
// The scan follows real replay order across sub-list returns and across the remaining handles of this submission; end of submission = the pass must end. Undecidable cases such as a mid-pass clear or too-deep nesting always count as 'used'.
|
||||
static const uint8_t *g_exec_next; // The op following the current draw in the replay
|
||||
static const uint8_t *g_exec_cur; // The current draw's own op (set with g_exec_next)
|
||||
static const uint8_t *g_ret[16]; static int g_ret_n; // Call return stack of the replay (at OP_CALL the caller's next op is pushed)
|
||||
static const uint64_t *g_rest; static int g_nrest; // Handles in this submission still to be replayed after the current handle
|
||||
void mtl_exec_rest(const uint64_t *h, int n) { g_rest = h; g_nrest = n; }
|
||||
@@ -5849,9 +5881,9 @@ static int ubd_conflict(NvProg *pr, const p_draw *dw, const p_draw_idx *di) {
|
||||
return 0;
|
||||
}
|
||||
// Draws of an indirect-count multi-draw (hle_nvn.c nvnCommandBufferMultiDrawElementsIndirectCount): draw i is made only while i < the u32 count the game's GPU culling wrote earlier in the
|
||||
// submission. Metal cannot skip draws by a count in GPU memory, so before the first such draw after GPU writes (a dispatch, copy, update or fill since the last wait, g_wr_epoch; or a new
|
||||
// submission) the replay ends the render pass, commits what it has encoded and waits for it, then reads the count on the CPU; later draws of the group use the same value. One wait per
|
||||
// GPU-driven pass instead of a compute pass that masks the records (not done yet). The records themselves are still read by the GPU (ordinary indirect draws)
|
||||
// submission. Metal cannot skip draws by a count in GPU memory. Most such draws are masked on the GPU (mdx_run below); the rest (quads, fans and 8-bit indices, whose records ixc_run
|
||||
// converts, and records or counts in no Metal buffer) wait here: before the first such draw after GPU writes (a dispatch, copy, update or fill since the last wait, g_wr_epoch; or a new
|
||||
// submission) the replay ends the render pass, commits what it has encoded and waits for it, then reads the count on the CPU; later draws of the group use the same value
|
||||
static id<MTLSharedEvent> g_mdi_ev; static uint64_t g_mdi_v, g_mdi_waits, g_mdi_epoch = ~0ull, g_mdi_sub, g_mdi_addr; static uint32_t g_mdi_count;
|
||||
static void mdi_wait(void) {
|
||||
rp_end();
|
||||
@@ -5877,6 +5909,93 @@ static int mdi_live(uint64_t cnt, int i) {
|
||||
}
|
||||
return (uint32_t)i < g_mdi_count;
|
||||
}
|
||||
// ---------- Indirect-count multi-draws masked on the GPU ----------
|
||||
// The wait above held the submitting thread for the GPU at every GPU-driven pass: in Monster Hunter Rise's Kamura Village, 14 of the 20 ms a submission took on an iPhone were that wait
|
||||
// (6 ms of work), and the game ran at 36–54 fps. Instead a compute pass copies each such draw's record into the frame arena, or zeros for a draw past the count, and the draw reads the copy:
|
||||
// past the count it draws nothing, and the CPU never waits. As in ixc_run, a compute pass cannot run inside a render pass: at the first such draw with no copy yet, mdx_run copies it and every
|
||||
// such draw after it in replay order up to the next write that could change a record or count (a dispatch, copy, fill or texture → buffer copy, which bump g_mdx_epoch; a uniform update only
|
||||
// when it covers one: g_mdx_stop), so only a draw met inside an open render pass splits it. Copies are found by record and count address in an open-addressing table, valid for one g_mdx_tag
|
||||
// (Kamura Village draws the same records with two counts between the same writes: keyed by the record alone, the second multi-draw found no copy and waited for the GPU each frame)
|
||||
#define MDX_SLOTS 16384 // a power of two; a scan fills at most a quarter
|
||||
typedef struct { uint64_t ind, cnt, out, tag; } mdx_ent; // out: GPU address of the copy, 0 when the record or count is in no Metal buffer (those draws wait, mdi_live)
|
||||
static mdx_ent g_mdx[MDX_SLOTS]; static int g_mdx_n; static uint64_t g_mdx_tag = 1, g_mdx_ep = ~0ull, g_mdx_sub, g_mdx_runs, g_mdx_splits, g_mdx_draws, g_mdx_fast; static const uint8_t *g_mdx_stop;
|
||||
static const char k_mdx_msl[] =
|
||||
"#include <metal_stdlib>\nusing namespace metal;\n"
|
||||
"struct mjob { device const uint *rec; device const uint *cnt; device uint *out; uint i, words; };\n" // draw i's record (words long), its multi-draw's count, the copy (null: padding)
|
||||
"kernel void mdx_mask(device const mjob *j [[buffer(0)]], uint t [[thread_position_in_grid]]) {\n"
|
||||
" mjob m = j[t]; if (!m.out) return; bool live = m.i < *m.cnt;\n"
|
||||
" for (uint k = 0; k < m.words; k++) m.out[k] = live ? m.rec[k] : 0u; }\n";
|
||||
static int mdx_wanted(const p_draw *dw, const p_draw_idx *di) { return dw ? dw->mdi_cnt && !ixc_wanted(dw, NULL) : di->mdi_cnt && !ixc_wanted(NULL, di); }
|
||||
static mdx_ent *mdx_slot(uint64_t ind, uint64_t cnt) { // its entry in this tag, or the free slot for it
|
||||
for (uint64_t h = ((ind ^ cnt * 0xD6E8FEB86659FD93ull) * 0x9E3779B97F4A7C15ull) >> 50;; h = (h + 1) & (MDX_SLOTS - 1)) { mdx_ent *e = &g_mdx[h]; if (e->tag != g_mdx_tag || (e->ind == ind && e->cnt == cnt)) return e; }
|
||||
}
|
||||
static int mdx_find(uint64_t ind, uint64_t cnt, uint64_t *out) { // → 1 when mdx_run has taken this draw (*out: its copy, or 0)
|
||||
if (g_mdx_ep != g_mdx_epoch || g_mdx_sub != g_sub_serial) return 0;
|
||||
const mdx_ent *e = mdx_slot(ind, cnt); if (e->tag != g_mdx_tag) return 0;
|
||||
*out = e->out; return 1;
|
||||
}
|
||||
static void mdx_run(const p_draw *dw, const p_draw_idx *di) {
|
||||
static id<MTLComputePipelineState> ps; static id<MTL4ArgumentTable> at; static int failed;
|
||||
if (failed) return;
|
||||
if (!at) { NSError *e = nil; id<MTLLibrary> lib = [g_dev newLibraryWithSource:@(k_mdx_msl) options:nil error:&e]; id<MTLFunction> f = [lib newFunctionWithName:@"mdx_mask"];
|
||||
ps = f ? [g_dev newComputePipelineStateWithFunction:f error:&e] : nil;
|
||||
MTL4ArgumentTableDescriptor *ad = [MTL4ArgumentTableDescriptor new]; ad.maxBufferBindCount = 1; at = ps ? [g_dev newArgumentTableWithDescriptor:ad error:&e] : nil;
|
||||
if (!at) { failed = 1; MLOG("[RED] indirect-count multi-draw masking: cannot create the compute pipeline %s (those draws wait for the GPU)", e.localizedDescription.UTF8String); return; } }
|
||||
if (g_mdx_ep != g_mdx_epoch || g_mdx_sub != g_sub_serial || g_mdx_n > MDX_SLOTS / 4) { g_mdx_tag++; g_mdx_n = 0; g_mdx_ep = g_mdx_epoch; g_mdx_sub = g_sub_serial; }
|
||||
enum { QMAX = MDX_SLOTS / 4 };
|
||||
static struct { mdx_ent *e; uint32_t i, words; } q[QMAX]; int nq = 0;
|
||||
uint64_t lo = ~0ull, hi = 0; // the hull of the records and counts taken, for the uniform update check
|
||||
#define MDX_TAKE(IND, CNT, I, W) do { mdx_ent *e_ = mdx_slot((IND), (CNT)); if (e_->tag != g_mdx_tag && nq < QMAX) { *e_ = (mdx_ent){(IND), (CNT), 0, g_mdx_tag}; g_mdx_n++; \
|
||||
q[nq].e = e_; q[nq].i = (uint32_t)(I); q[nq++].words = (W); lo = MIN(lo, MIN((IND), (CNT))); hi = MAX(hi, MAX((IND) + 4 * (W), (CNT) + 4)); } } while (0)
|
||||
if (dw) MDX_TAKE(dw->ind, dw->mdi_cnt, dw->mdi_i, 4); else MDX_TAKE(di->ind, di->mdi_cnt, di->mdi_i, 5);
|
||||
if (g_exec_next && g_ret_n <= 8) { // later draws in replay order (current list → callers → the rest of this submission)
|
||||
fz_cursor cur = {.p = g_exec_next, .ret_n = g_ret_n}; memcpy(cur.ret, g_ret, sizeof(*g_ret) * (size_t)g_ret_n);
|
||||
const uint8_t *p; op_hdr h; const void *a;
|
||||
for (int guard = 0; guard < 200000 && nq < QMAX && fz_next(&cur, &p, &h, &a); guard++) {
|
||||
if (h.op == OP_DISPATCH || h.op == OP_COPY_B2B || h.op == OP_CLEAR_BUF || h.op == OP_COPY_T2B) break;
|
||||
if (h.op == OP_UPDATE_UB) { uint64_t d, n; memcpy(&d, a, 8); memcpy(&n, (const uint8_t *)a + 8, 8);
|
||||
int hit = 0; for (int k = 0; d < hi && lo < d + n && k < nq && !hit; k++) { const mdx_ent *e = q[k].e; hit = (d < e->ind + 4 * q[k].words && e->ind < d + n) || (d < e->cnt + 4 && e->cnt < d + n); }
|
||||
if (hit) { g_mdx_stop = p; break; } continue; }
|
||||
p_draw dt; p_draw_idx it; const p_draw *d = h.op == OP_DRAW ? draw_rec(a, h.size, &dt) : NULL; const p_draw_idx *x = h.op == OP_DRAW_IDX ? draw_idx_rec(a, h.size, &it) : NULL;
|
||||
if (d && mdx_wanted(d, NULL)) MDX_TAKE(d->ind, d->mdi_cnt, d->mdi_i, 4);
|
||||
else if (x && mdx_wanted(NULL, x)) MDX_TAKE(x->ind, x->mdi_cnt, x->mdi_i, 5);
|
||||
}
|
||||
}
|
||||
#undef MDX_TAKE
|
||||
if (!nq) return;
|
||||
if (g_enc) { g_mdx_splits++; PD("M %llu indirect-count multi-draw masking: render pass split here", g_stat_pass); rp_end(); }
|
||||
if (g_blit) { enc_b_end(g_blit); g_blit = nil; }
|
||||
typedef struct { uint64_t rec, cnt, out; uint32_t i, words; } mjob;
|
||||
int nj = (nq + 63) & ~63;
|
||||
id<MTLBuffer> ab; size_t ao = arena_alloc((size_t)nj * sizeof(mjob) + (size_t)nq * 32, &ab); // the jobs (padded to whole threadgroups), then the copies 32 bytes apart
|
||||
mjob *j = (mjob *)((uint8_t *)ab.contents + ao); uint64_t og = ab.gpuAddress + ao + (uint64_t)nj * sizeof(mjob);
|
||||
memset(j, 0, (size_t)nj * sizeof(mjob));
|
||||
for (int k = 0; k < nq; k++) { mdx_ent *e = q[k].e; void *mb; uint64_t mo, rg, cg;
|
||||
if (!mtl_gpu_addr(e->ind, 4 * q[k].words, &mb, &mo, &rg) || !mtl_gpu_addr(e->cnt, 4, &mb, &mo, &cg)) continue;
|
||||
e->out = og + 32 * (uint64_t)k; j[k] = (mjob){rg, cg, e->out, q[k].i, q[k].words}; }
|
||||
id<MTL4ComputeCommandEncoder> ce = enc_c_begin(NULL); [ce setArgumentTable:at]; [at setAddress:ab.gpuAddress + ao atIndex:0]; [ce setComputePipelineState:ps];
|
||||
[ce dispatchThreadgroups:MTLSizeMake((NSUInteger)nj / 64, 1, 1) threadsPerThreadgroup:MTLSizeMake(64, 1, 1)];
|
||||
enc_c_end(ce);
|
||||
args_dirty_all();
|
||||
g_mdx_draws += (uint64_t)nq;
|
||||
if (!g_mdx_runs++) MLOG("indirect-count multi-draw masking: first time, %d draws (their records are copied on the GPU, zeros past the count, before the render pass)", nq);
|
||||
else if (g_mdx_runs % 3600 == 0) MLOG("indirect-count multi-draw masking: %llu times, %llu draws (%llu encoded alone, mdx_next), render pass split %llu times; waits for the GPU-written count %llu",
|
||||
(unsigned long long)g_mdx_runs, (unsigned long long)g_mdx_draws, (unsigned long long)g_mdx_fast, (unsigned long long)g_mdx_splits, (unsigned long long)g_mdi_waits);
|
||||
}
|
||||
// The draws of one multi-draw are recorded back to back with the same state, and only the first few are live: each one replayed right after the one before it (no op between them,
|
||||
// the same open pass) is encoded alone, without the rest of draw's work, which cost a whole draw for each masked one (about 1.5 µs on an M4 Pro, 1,100 a frame in Kamura Village)
|
||||
static struct { const uint8_t *next; const void *enc; uint64_t cnt, ig, il, pass; int i; MTLPrimitiveType mp; MTLIndexType it; } g_mdxc;
|
||||
static void mdx_chain(uint64_t cnt, int i, MTLPrimitiveType mp, MTLIndexType it, uint64_t ig, uint64_t il) {
|
||||
g_mdxc = g_defer_replaying || !g_exec_next ? (typeof(g_mdxc)){0} : (typeof(g_mdxc)){g_exec_next, (__bridge const void *)g_enc, cnt, ig, il, g_stat_pass, i, mp, it};
|
||||
}
|
||||
static int mdx_next(const p_draw *dw, const p_draw_idx *di) {
|
||||
uint64_t cnt = dw ? dw->mdi_cnt : di->mdi_cnt, out; int i = dw ? dw->mdi_i : di->mdi_i;
|
||||
if (!g_mdxc.next || g_mdxc.next != g_exec_cur || !g_enc || g_mdxc.enc != (__bridge const void *)g_enc || g_mdxc.pass != g_stat_pass || g_mdxc.cnt != cnt || g_mdxc.i + 1 != i
|
||||
|| g_defer_replaying || passdump() || !mdx_find(dw ? dw->ind : di->ind, cnt, &out) || !out) return 0;
|
||||
if (dw) enc_r_draw_indirect(g_enc, g_mdxc.mp, out); else enc_r_draw_indexed_indirect(g_enc, g_mdxc.mp, g_mdxc.it, g_mdxc.ig, g_mdxc.il, out);
|
||||
g_stat_draw++; g_mdx_fast++; g_mdxc.next = g_exec_next; g_mdxc.i = i;
|
||||
return 1;
|
||||
}
|
||||
// ---------- Transform feedback ----------
|
||||
// A draw under rasterizer discard rasterizes nothing; with feedback active it captures its vertex shader's outputs into the bound feedback buffers. Pokémon Legends: Z-A skins its characters
|
||||
// so: a vertex-only skinning program writes positions (12 bytes a vertex) and packed normal / tangent bits (4 bytes each), and the character draws read those buffers as vertex streams.
|
||||
@@ -6016,7 +6135,10 @@ static int lc_keeps(uint32_t op) {
|
||||
|
||||
static void draw(const p_draw *dw, const p_draw_idx *di) {
|
||||
g_draw_binst = (uint32_t)(dw ? dw->base_instance : di ? di->base_instance : 0); // c0 + 0x644 of the stages that read the driver segment (encode_args)
|
||||
if ((di && di->mdi_cnt && !mdi_live(di->mdi_cnt, di->mdi_i)) || (dw && dw->mdi_cnt && !mdi_live(dw->mdi_cnt, dw->mdi_i))) return; // past the multi-draw's count
|
||||
int mdx_on = (dw && dw->mdi_cnt) || (di && di->mdi_cnt) ? mdx_wanted(dw, di) : 0; // an indirect-count multi-draw's record masked on the GPU (mdx_run), else read after a wait
|
||||
if (!mdx_on && ((di && di->mdi_cnt && !mdi_live(di->mdi_cnt, di->mdi_i)) || (dw && dw->mdi_cnt && !mdi_live(dw->mdi_cnt, dw->mdi_i)))) return; // past the multi-draw's count
|
||||
if (mdx_on && mdx_next(dw, di)) return; // the next draw of a masked multi-draw: only the draw is encoded
|
||||
g_mdxc.next = NULL;
|
||||
if (!S.prog) { RED_ONCE(18, "no program bound before the draw"); PD("Z %llu draw skipped: no program bound", g_stat_pass + !g_enc); return; }
|
||||
__unsafe_unretained NvProg *pr = prog_get(S.prog);
|
||||
if (S.rast_discard) { tf_draw(pr, dw, di); return; } // the vertex stage only: transform feedback capture, or nothing (tf_draw)
|
||||
@@ -6026,6 +6148,10 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
|
||||
PD("N %llu zero-source draw skipped", g_stat_pass + 1); return; }
|
||||
if (g_nlc && !g_enc) lc_take(pr); // waiting clears: this pass's load actions, or landed before it (see lc_take)
|
||||
if (ixc_wanted(dw, di) && !ixc_find(dw, di)) ixc_run(dw, di); // indirect quads / fans / 8-bit indices: converted on the GPU before the render pass, at their own place in the replay (see ixc_run)
|
||||
uint64_t mdx = 0; // the masked copy of the record (mdx_run); a draw it could not copy waits for the count instead
|
||||
if (mdx_on) { uint64_t ind = dw ? dw->ind : di->ind, cnt = dw ? dw->mdi_cnt : di->mdi_cnt; int i = dw ? dw->mdi_i : di->mdi_i;
|
||||
if (!mdx_find(ind, cnt, &mdx)) { mdx_run(dw, di); if (!mdx_find(ind, cnt, &mdx)) mdx = 0; }
|
||||
if (!mdx && !mdi_live(cnt, i)) return; }
|
||||
if (di && wide_wanted(di) && !wide_find(di)) wide_run(di); // 8-bit indices written earlier in this submission: widened on the GPU before the render pass (see wide_run)
|
||||
if (!g_defer_replaying) {
|
||||
if (g_defer.active && g_defer.in_p && !g_defer.at_q) { if (defer_take(dw, di)) return; }
|
||||
@@ -6214,10 +6340,11 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
|
||||
const ixc_ent *e = ixc_wanted(dw, di) ? ixc_find(dw, di) : NULL; void *mb, *ab; uint64_t mo, ao, ig;
|
||||
uint64_t ind = dw ? dw->ind : di->ind;
|
||||
if (ixc_wanted(dw, di)) { if (e && e->args) enc_r_draw_indexed_indirect(en, mp, MTLIndexTypeUInt32, e->ixb, e->ixlen, e->args); else RED_ONCE(88, "indirect draw skipped: the arguments for primitive %d were not converted (argument address %#llx is in no memory pool, or the conversion pipeline was not created)", prim, (unsigned long long)ind); }
|
||||
else if (!mtl_addr(ind, dw ? 16 : 20, &ab, &ao)) RED_ONCE(88, "indirect draw skipped: argument address %#llx is in no memory pool", (unsigned long long)ind);
|
||||
else if (dw) enc_r_draw_indirect(en, mp, ((__bridge id<MTLBuffer>)ab).gpuAddress + ao);
|
||||
else if (!mdx && !mtl_addr(ind, dw ? 16 : 20, &ab, &ao)) RED_ONCE(88, "indirect draw skipped: argument address %#llx is in no memory pool", (unsigned long long)ind);
|
||||
else if (dw) { enc_r_draw_indirect(en, mp, mdx ? mdx : ((__bridge id<MTLBuffer>)ab).gpuAddress + ao); if (mdx) mdx_chain(dw->mdi_cnt, dw->mdi_i, mp, MTLIndexTypeUInt32, 0, 0); }
|
||||
else if (!mtl_gpu_addr(di->ibuf, 1, &mb, &mo, &ig)) RED_ONCE(22, "index buffer is in no memory pool");
|
||||
else enc_r_draw_indexed_indirect(en, mp, di->itype == 1 ? MTLIndexTypeUInt16 : MTLIndexTypeUInt32, ig, ((__bridge id<MTLBuffer>)mb).length - mo, ((__bridge id<MTLBuffer>)ab).gpuAddress + ao);
|
||||
else { MTLIndexType ty = di->itype == 1 ? MTLIndexTypeUInt16 : MTLIndexTypeUInt32; uint64_t il = ((__bridge id<MTLBuffer>)mb).length - mo;
|
||||
enc_r_draw_indexed_indirect(en, mp, ty, ig, il, mdx ? mdx : ((__bridge id<MTLBuffer>)ab).gpuAddress + ao); if (mdx) mdx_chain(di->mdi_cnt, di->mdi_i, mp, ty, ig, il); }
|
||||
} else if (dw && prim == 7) { // Quads: shared indices 0,1,2,0,2,3
|
||||
quad_ib_make();
|
||||
if (dw->count >= 4) enc_r_draw_indexed(en, mp, (dw->count / 4) * 6, MTLIndexTypeUInt32, g_quad_gpu, g_quad_ib.length, dw->instances > 0 ? dw->instances : 1, dw->first, dw->base_instance);
|
||||
@@ -6321,7 +6448,7 @@ static id<MTL4ComputeCommandEncoder> dispatch_enc(void) {
|
||||
}
|
||||
static void dispatch(const p_dispatch *dp) {
|
||||
zskip_clear();
|
||||
g_wr_epoch++;
|
||||
g_wr_epoch++; g_mdx_epoch++;
|
||||
if (!S.cprog) { RED_ONCE(27, "no program bound before the dispatch"); return; }
|
||||
// DispatchComputeIndirect: the group counts {x, y, z} at ind (MTLDispatchThreadgroupsIndirectArguments) are usually written by the GPU earlier in the frame; Metal reads them when the dispatch runs
|
||||
void *ib = NULL; uint64_t io = 0, ig = 0;
|
||||
@@ -6375,7 +6502,7 @@ static void dispatch(const p_dispatch *dp) {
|
||||
}
|
||||
// GPU: texture -> buffer (row stride taken from SetCopyRowStride, layer stride from SetCopyImageStride, 0 = tightly packed)
|
||||
static void copy_t2b(const p_copy_b2t *c) {
|
||||
g_wr_epoch++;
|
||||
g_wr_epoch++; g_mdx_epoch++;
|
||||
if (c->tex) xv_check(c->tex->serial, c->has_view && c->view.set_levels ? c->view.base_level : 0, c->has_view && c->view.set_layers ? c->view.min_layer : 0, "texture-to-buffer copy");
|
||||
id<MTLTexture> t = tex_mtl(c->tex);
|
||||
if (!t) return;
|
||||
@@ -6900,6 +7027,7 @@ void mtl_exec(const uint8_t *p) {
|
||||
case OP_COPY_B2B: { uint64_t s_, d, n; memcpy(&s_, a, 8); memcpy(&d, (const uint8_t *)a + 8, 8); memcpy(&n, (const uint8_t *)a + 16, 8);
|
||||
PD("Y copy_b2b cmd=%p src=%#llx dst=%#llx bytes=%llu", p, (unsigned long long)s_, (unsigned long long)d, (unsigned long long)n); mtl_sub_copy(s_, d, n); break; }
|
||||
case OP_UPDATE_UB: { uint64_t d, n; memcpy(&d, a, 8); memcpy(&n, (const uint8_t *)a + 8, 8);
|
||||
if (p == g_mdx_stop) { g_mdx_epoch++; g_mdx_stop = NULL; } // it writes a record or count that masked copies were made of (mdx_run)
|
||||
PD("Y update_ub cmd=%p src=%#llx dst=%#llx bytes=%llu f=%g,%g,%g,%g", p, (unsigned long long)((const uint8_t *)a + 16), (unsigned long long)d, (unsigned long long)n,
|
||||
n >= 4 ? ((const float *)((const uint8_t *)a + 16))[0] : 0.f, n >= 8 ? ((const float *)((const uint8_t *)a + 16))[1] : 0.f, n >= 12 ? ((const float *)((const uint8_t *)a + 16))[2] : 0.f, n >= 16 ? ((const float *)((const uint8_t *)a + 16))[3] : 0.f); mtl_sub_update((const uint8_t *)a + 16, d, n); break; }
|
||||
case OP_SET_RT: {
|
||||
@@ -6952,8 +7080,8 @@ void mtl_exec(const uint8_t *p) {
|
||||
PD("c %llu depth-stencil clear recorded ser=%llu depth=%g depth_mask=%d stencil=%d stencil_mask=%#x", g_stat_pass, R.rt.depth ? (unsigned long long)R.rt.depth->serial : 0ull, c->depth, c->depth_mask, c->stencil, c->stencil_mask);
|
||||
break;
|
||||
}
|
||||
case OP_DRAW: { if (g_rc_skip) break; p_draw t; g_exec_next = p + h.size; draw(draw_rec(a, h.size, &t), NULL); g_exec_next = NULL; break; } // Lookahead start point
|
||||
case OP_DRAW_IDX: { if (g_rc_skip) break; p_draw_idx t; g_exec_next = p + h.size; draw(NULL, draw_idx_rec(a, h.size, &t)); g_exec_next = NULL; break; }
|
||||
case OP_DRAW: { if (g_rc_skip) break; p_draw t; g_exec_cur = p; g_exec_next = p + h.size; draw(draw_rec(a, h.size, &t), NULL); g_exec_next = NULL; break; } // Lookahead start point
|
||||
case OP_DRAW_IDX: { if (g_rc_skip) break; p_draw_idx t; g_exec_cur = p; g_exec_next = p + h.size; draw(NULL, draw_idx_rec(a, h.size, &t)); g_exec_next = NULL; break; }
|
||||
// Do not mark dirty if the state did not change (the game often re-sends the same program / state block / vertex format before each draw; the dirty flag only decides whether to re-fetch the pipeline / depth-stencil object, and equal values mean the same object)
|
||||
case OP_SCISSOR: case OP_BIND_PROGRAM: case OP_BIND_STATE: case OP_BIND_VATTR: case OP_BIND_VSTREAM: case OP_BIND_VBUF:
|
||||
case OP_BIND_UBO: case OP_BIND_SSBO: case OP_BIND_TEX: case OP_BIND_SAMP: case OP_BIND_IMG: case OP_STENCIL:
|
||||
|
||||
+33
-7
@@ -62,21 +62,36 @@ static void samples_xy(int s, uint32_t *sx, uint32_t *sy) {
|
||||
if (s == 2) *sx = 2; else if (s == 4) *sx = *sy = 2; else if (s == 8) { *sx = 4; *sy = 2; }
|
||||
}
|
||||
|
||||
static uint64_t level_offset(const nvn_texdesc *d, unsigned level, unsigned dpl, uint32_t w, uint32_t h, uint32_t z,
|
||||
unsigned tileH, unsigned tileD, fmt_t t) {
|
||||
// Sparse textures (flag SPARSE, block-linear) are mapped in 64 KB tiles: one row of blocks, 128 GOBs / (block height × block depth) wide. A level at least one tile wide, high and deep
|
||||
// is laid out with its width padded to whole tiles, so it fills whole tiles; the levels from the first smaller one on form the mip tail, packed as usual from there, and the
|
||||
// storage is whole tiles, 64 KB aligned. Monster Hunter Rise's streamed textures come pre-swizzled that way in its files and are placed at the level offsets the host reports: laid out unpadded,
|
||||
// its 1024×1024 ASTC 6×6 textures (48 GOBs wide instead of 43, level 1 at 0x90000 instead of 0x81000) read 2,183 all-zero blocks in level 0 and 3,043 in level 1, which ASTC
|
||||
// decodes as its error colour (magenta clothes)
|
||||
enum { SPARSE_TILE = 0x10000 };
|
||||
static int sparse_of(const nvn_texdesc *d) { return (d->flags & 0x20) && !(d->flags & 0x10) && d->target != 10; }
|
||||
// Offset of level `level` from the start of its layer, and the GOB columns it is laid out with (*gobs, may be NULL)
|
||||
static uint64_t level_walk(const nvn_texdesc *d, unsigned level, unsigned dpl, uint32_t w, uint32_t h, uint32_t z,
|
||||
unsigned tileH, unsigned tileD, fmt_t t, uint32_t *gobs) {
|
||||
uint32_t sx, sy; samples_xy(d->samples, &sx, &sy);
|
||||
uint32_t tw = 128u >> (tileH + tileD); if (!tw) tw = 1; // GOB columns of a 64 KB tile
|
||||
int tiled = sparse_of(d); // still in a sparse texture's tiled levels
|
||||
uint64_t off = 0;
|
||||
for (unsigned i = 0; i < level; i++) {
|
||||
for (unsigned i = 0;; i++) {
|
||||
uint32_t lw = blk(mip(w * sx, i), t.bw);
|
||||
uint32_t lh = dpl >= 2 ? blk(mip(h * sy, i), t.bh) : 1;
|
||||
uint32_t ld = dpl >= 3 ? mip(z, i) : 1;
|
||||
uint32_t lwb = lw * t.bpb;
|
||||
unsigned th = adj_tile(tileH, 8, lh), td = adj_tile(tileD, 1, ld);
|
||||
uint32_t wg = (lwb + 63) / 64, hg = (lh + 7) / 8;
|
||||
uint32_t wt = wg, ht = (hg + (1u << th) - 1) >> th, dt = (ld + (1u << td) - 1) >> td;
|
||||
off += (uint64_t)wt * ht * dt << (9 + th + td);
|
||||
tiled = tiled && lwb >= tw * 64 && lh >= 8u << tileH && ld >= 1u << tileD;
|
||||
if (tiled) wg = (wg + tw - 1) / tw * tw;
|
||||
if (i == level) { if (gobs) *gobs = wg; return off; }
|
||||
uint32_t ht = (hg + (1u << th) - 1) >> th, dt = (ld + (1u << td) - 1) >> td;
|
||||
off += (uint64_t)wg * ht * dt << (9 + th + td);
|
||||
}
|
||||
return off;
|
||||
}
|
||||
static uint64_t level_offset(const nvn_texdesc *d, unsigned level, unsigned dpl, uint32_t w, uint32_t h, uint32_t z, unsigned tileH, unsigned tileD, fmt_t t) {
|
||||
return level_walk(d, level, dpl, w, h, z, tileH, tileD, t, 0);
|
||||
}
|
||||
|
||||
void nvn_tex_layout(const nvn_texdesc *d, nvn_texlayout *o) {
|
||||
@@ -102,6 +117,7 @@ void nvn_tex_layout(const nvn_texdesc *d, nvn_texlayout *o) {
|
||||
o->tileH = tileH; o->tileD = tileD;
|
||||
unsigned levels = d->levels ? d->levels : 1;
|
||||
uint64_t ls = level_offset(d, levels, dpl, w, h, z, tileH, tileD, t);
|
||||
if (sparse_of(d)) ls = (ls + SPARSE_TILE - 1) & ~(uint64_t)(SPARSE_TILE - 1);
|
||||
if (layered) { // Layer size is aligned to the block
|
||||
uint32_t hh = blk(h, t.bh); unsigned th = tileH, td = tileD;
|
||||
while (th && hh <= (8u << (th - 1))) th--;
|
||||
@@ -111,7 +127,7 @@ void nvn_tex_layout(const nvn_texdesc *d, nvn_texlayout *o) {
|
||||
}
|
||||
o->layer_size = ls;
|
||||
o->size = layered ? ls * z : ls;
|
||||
o->align = 512;
|
||||
o->align = sparse_of(d) ? SPARSE_TILE : 512;
|
||||
}
|
||||
|
||||
uint64_t nvn_tex_level_offset(const nvn_texdesc *d, unsigned level) {
|
||||
@@ -123,6 +139,16 @@ uint64_t nvn_tex_level_offset(const nvn_texdesc *d, unsigned level) {
|
||||
return level_offset(d, level, dpl, w, h, z, L.tileH, L.tileD, t);
|
||||
}
|
||||
|
||||
uint32_t nvn_tex_level_gobs(const nvn_texdesc *d, unsigned level) {
|
||||
fmt_t t = fmt_of(d->format);
|
||||
unsigned dpl; int layered; uint32_t w, h, z, g = 0;
|
||||
shape(d, &dpl, &layered, &w, &h, &z);
|
||||
if (!t.bpb || (d->flags & 0x10) || d->target == 10) return 0;
|
||||
nvn_texlayout L; nvn_tex_layout(d, &L);
|
||||
level_walk(d, level, dpl, w, h, z, L.tileH, L.tileD, t, &g);
|
||||
return g;
|
||||
}
|
||||
|
||||
void nvn_tex_level_tiles(const nvn_texdesc *d, unsigned level, unsigned *tileH, unsigned *tileD) {
|
||||
fmt_t t = fmt_of(d->format);
|
||||
unsigned dpl; int layered; uint32_t w, h, z;
|
||||
|
||||
@@ -16,6 +16,7 @@ typedef struct {
|
||||
|
||||
void nvn_tex_layout(const nvn_texdesc *d, nvn_texlayout *o);
|
||||
uint64_t nvn_tex_level_offset(const nvn_texdesc *d, unsigned level);
|
||||
uint32_t nvn_tex_level_gobs(const nvn_texdesc *d, unsigned level); // GOB columns level `level` is laid out with (its width in GOBs; padded to whole 64 KB tiles in a sparse texture's tiled levels)
|
||||
int nvn_fmt_bpb(int format, int *bw, int *bh); // Bytes per block (0 = unknown format), block width and height
|
||||
// Block height / block depth actually used by level `level` (log2 of the GOB count): the per-level block height shrinks with that level's row count (block rows), and deswizzling uses it level by level
|
||||
void nvn_tex_level_tiles(const nvn_texdesc *d, unsigned level, unsigned *tileH, unsigned *tileD);
|
||||
@@ -45,13 +45,14 @@ static int move(const nvn_texdesc *d, unsigned level, unsigned layer, uint8_t *b
|
||||
uint32_t row_bytes, rows, slices, bpb; uint64_t base; unsigned th, td;
|
||||
if (nvn_deswizzle_geom(d, level, layer, &row_bytes, &rows, &slices, &base, &th, &td, &bpb)) return -1;
|
||||
if (pitch < row_bytes) return -1;
|
||||
bl += base;
|
||||
bl += base; uint32_t lb = nvn_tex_level_gobs(d, level) * 64; // the row as laid out (padded in a sparse texture's tiled levels)
|
||||
if (lb < row_bytes) lb = row_bytes;
|
||||
for (uint32_t z = 0; z < slices; z++)
|
||||
for (uint32_t y = 0; y < rows; y++) {
|
||||
uint8_t *row = lin + ((size_t)z * rows + y) * pitch;
|
||||
for (uint32_t x = 0; x < row_bytes; x += 16) { // 16-byte segments: 16 bytes are contiguous within a GOB, and 16 divides all block widths
|
||||
uint32_t n = row_bytes - x < 16 ? row_bytes - x : 16;
|
||||
uint8_t *p = bl + nvn_bl_offset(x, y, z, row_bytes, rows, th, td);
|
||||
uint8_t *p = bl + nvn_bl_offset(x, y, z, lb, rows, th, td);
|
||||
if (to_linear) memcpy(row + x, p, n); else memcpy(p, row + x, n);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -11,7 +11,8 @@
|
||||
int nvn_deswizzle(const nvn_texdesc *d, unsigned level, unsigned layer, const void *src_storage, void *dst_linear, size_t dst_pitch);
|
||||
// Reverse (for unit test round trips): linear src_linear → written into the corresponding level / layer of the whole-texture storage dst_storage
|
||||
int nvn_swizzle(const nvn_texdesc *d, unsigned level, unsigned layer, const void *src_linear, size_t src_pitch, void *dst_storage);
|
||||
// Single-point address: within this level and layer, linear (x byte, y block row, z slice) → block-linear byte offset relative to the start of this level and layer (unit tests verify against hand calculation)
|
||||
// Single-point address: within this level and layer, linear (x byte, y block row, z slice) → block-linear byte offset relative to the start of this level and layer (unit tests verify against hand calculation).
|
||||
// row_bytes: the row as laid out, which a sparse texture's tiled levels pad to whole tiles (nvn_tex_level_gobs × 64)
|
||||
uint64_t nvn_bl_offset(uint32_t x, uint32_t y, uint32_t z, uint32_t row_bytes, uint32_t rows, unsigned tileH, unsigned tileD);
|
||||
|
||||
int nvn_deswizzle_geom(const nvn_texdesc *d, unsigned level, unsigned layer,
|
||||
|
||||
@@ -5,9 +5,9 @@ static id<MTLComputePipelineState> g_nvn_deswz_pipeline;
|
||||
static const char *const k_nvn_deswz_msl =
|
||||
"#include <metal_stdlib>\n"
|
||||
"using namespace metal;\n"
|
||||
"struct P { ulong src_off, dst_off, base, pitch; uint row_bytes, rows, slices, th, td; };\n"
|
||||
"struct P { ulong src_off, dst_off, base, pitch; uint row_bytes, rows, slices, th, td, gobs_x; };\n"
|
||||
"kernel void deswz(device const uchar *src [[buffer(0)]], device uchar *dst [[buffer(1)]], constant P& p [[buffer(2)]], uint3 id [[thread_position_in_grid]]) {\n"
|
||||
" uint bh = 1u << p.th, bd = 1u << p.td, gobs_x = (p.row_bytes + 63u) >> 6, blocks_y = (p.rows + (8u * bh) - 1u) >> (3u + p.th);\n"
|
||||
" uint bh = 1u << p.th, bd = 1u << p.td, gobs_x = p.gobs_x, blocks_y = (p.rows + (8u * bh) - 1u) >> (3u + p.th);\n"
|
||||
" uint tid = id.x & 31u, gob_x = id.x >> 5, gob_y = id.y, z = id.z; uint xseg = ((tid >> 4) << 1) | ((tid >> 1) & 1u), y = (gob_y << 3) | (((tid >> 2) & 3u) << 1) | (tid & 1u), x = (gob_x << 6) | (xseg << 4);\n"
|
||||
" if (x >= p.row_bytes || y >= p.rows || z >= p.slices) return;\n"
|
||||
" uint n = min(16u, p.row_bytes - x); ulong block = 512ul * bh * bd; ulong gob_in_block = ((ulong)(z & (bd - 1u)) << p.th) + (gob_y & (bh - 1u)); ulong so = p.src_off + p.base + ((ulong)(z >> p.td) * gobs_x * blocks_y + (gob_y >> p.th) * gobs_x + gob_x) * block + gob_in_block * 512ul + (ulong)tid * 16ul; ulong doff = p.dst_off + ((ulong)z * p.rows + y) * p.pitch + x;\n"
|
||||
@@ -35,8 +35,8 @@ int nvn_deswz_gpu_encode(id<MTLComputeCommandEncoder> enc, id<MTLBuffer> src, ui
|
||||
uint64_t base;
|
||||
unsigned th, td;
|
||||
if (nvn_deswizzle_geom(d, level, layer, &row_bytes, &rows, &slices, &base, &th, &td, &bpb) || dst_pitch < row_bytes) return -1;
|
||||
struct { uint64_t src_off, dst_off, base, pitch; uint32_t row_bytes, rows, slices, th, td; } p = {
|
||||
src_off, dst_off, base, dst_pitch, row_bytes, rows, slices, th, td
|
||||
struct { uint64_t src_off, dst_off, base, pitch; uint32_t row_bytes, rows, slices, th, td, gobs_x; } p = {
|
||||
src_off, dst_off, base, dst_pitch, row_bytes, rows, slices, th, td, MAX(nvn_tex_level_gobs(d, level), (row_bytes + 63u) >> 6) // GOB columns as laid out (nvn_tex_level_gobs)
|
||||
};
|
||||
[enc setComputePipelineState:g_nvn_deswz_pipeline];
|
||||
[enc setBuffer:src offset:0 atIndex:0];
|
||||
@@ -57,8 +57,8 @@ int nvn_deswz_gpu_encode_m4(id<MTL4ComputeCommandEncoder> enc, id<MTL4ArgumentTa
|
||||
uint64_t base;
|
||||
unsigned th, td;
|
||||
if (nvn_deswizzle_geom(d, level, layer, &row_bytes, &rows, &slices, &base, &th, &td, &bpb) || dst_pitch < row_bytes) return -1;
|
||||
struct { uint64_t src_off, dst_off, base, pitch; uint32_t row_bytes, rows, slices, th, td; } p = {
|
||||
src_off, dst_off, base, dst_pitch, row_bytes, rows, slices, th, td
|
||||
struct { uint64_t src_off, dst_off, base, pitch; uint32_t row_bytes, rows, slices, th, td, gobs_x; } p = {
|
||||
src_off, dst_off, base, dst_pitch, row_bytes, rows, slices, th, td, MAX(nvn_tex_level_gobs(d, level), (row_bytes + 63u) >> 6) // GOB columns as laid out (nvn_tex_level_gobs)
|
||||
};
|
||||
memcpy((uint8_t *)params.contents + params_off, &p, sizeof p);
|
||||
[enc setArgumentTable:table];
|
||||
|
||||
@@ -111,6 +111,23 @@ int main(void) {
|
||||
printf(" E 32³ RGBA16F (the shape of a 32³ color lookup table) block height/depth %u/%u, storage %llu (expected 0/4, 262144)\n", L.tileH, L.tileD, (unsigned long long)L.size);
|
||||
}
|
||||
|
||||
// ---------- S sparse layout: Monster Hunter Rise's streamed 1024×1024 ASTC 6×6 sRGB textures (flag SPARSE, packaged block height 8 GOBs), level 0 checked against the game's data ----------
|
||||
// A 64 KB tile is 16 GOBs × 8 GOBs here. Level 0 (171 block columns = 43 GOBs) is padded to 48 GOBs × 3 block rows = 9 tiles: its dumped data decodes with no all-zero block that way
|
||||
// (2,183 unpadded). Level 1 (86 = 22 GOBs → 32, 2 block rows) is 4 tiles; level 2 (43 block columns, less than a tile's 64) starts the mip tail, packed as before; storage in whole tiles
|
||||
printf("S sparse layout (64 KB tiles: tiled levels padded to whole tiles, then the mip tail)\n");
|
||||
{
|
||||
nvn_texdesc d = {1, 0x8B, 0x20, 0, 1024, 1024, 1, 8, 0, 4}; nvn_texlayout L; nvn_tex_layout(&d, &L);
|
||||
static const uint64_t off[8] = {0, 0x90000, 0xD0000, 0xDB000, 0xDE000, 0xDEC00, 0xDF000, 0xDF200};
|
||||
static const uint32_t gobs[3] = {48, 32, 11};
|
||||
for (unsigned l = 0; l < 8; l++) CHECK(nvn_tex_level_offset(&d, l) == off[l], "S level %u offset %#llx expected %#llx", l, (unsigned long long)nvn_tex_level_offset(&d, l), (unsigned long long)off[l]);
|
||||
for (unsigned l = 0; l < 3; l++) CHECK(nvn_tex_level_gobs(&d, l) == gobs[l], "S level %u laid out %u GOBs wide, expected %u", l, nvn_tex_level_gobs(&d, l), gobs[l]);
|
||||
CHECK(L.size == 0xE0000 && L.align == 0x10000, "S storage %#llx alignment %#x, expected 0xe0000 / 0x10000", (unsigned long long)L.size, L.align);
|
||||
printf(" S 1024×1024 ASTC 6×6 sparse: levels at 0, %#llx, %#llx …, storage %#llx (expected 0x90000, 0xd0000, 0xe0000)\n", (unsigned long long)nvn_tex_level_offset(&d, 1), (unsigned long long)nvn_tex_level_offset(&d, 2), (unsigned long long)L.size);
|
||||
d.flags = 0; nvn_tex_layout(&d, &L); // the same texture not sparse keeps the packed layout
|
||||
CHECK(nvn_tex_level_offset(&d, 1) == 0x81000 && L.size == 0xBC400 && nvn_tex_level_gobs(&d, 0) == 43, "S not sparse: level 1 %#llx storage %#llx, expected 0x81000 / 0xbc400",
|
||||
(unsigned long long)nvn_tex_level_offset(&d, 1), (unsigned long long)L.size);
|
||||
}
|
||||
|
||||
printf("%s: %d/%d checks passed\n", fails ? "RED" : "all passed", checks - fails, checks);
|
||||
return fails ? 1 : 0;
|
||||
}
|
||||
@@ -227,6 +227,7 @@ Game Freak's Trinity engine on SDK 19 (NVN 55); every shader comes from the Oodl
|
||||
| `DM_NVN_ALLOCPROF=1` | Metal allocation growth per replayed op type and per function, command allocator sizes, new textures |
|
||||
| `DM_NVN_CBDUMP_KEY=vskey` (with `CBDUMP` and `PASSDUMP_ALL`) | Only that vertex shader's draws: uniform buffers plus vertex streams and index buffer (`F<frame>_p<pass>_vb<stream>_<addr>.bin`, `_ib<type>_<addr>.bin`) |
|
||||
| `DM_NVN_TEXINFO=WxH` | Storage and description (pool, offset, flags, packaged or not, storage size and how many of its bytes in guest memory are non-zero, i.e. written by the CPU) of every texture of that size on first use, and each pool registration of one (`TEXINFO register`: builder channel mapping, view and view mapping) |
|
||||
| `DM_NVN_ZBLOCK=1` or `=<folder>` | Each ASTC texture uploaded on the CPU whose bound levels hold all-zero blocks (ASTC's error colour, magenta): `[zb]` lines with each level's zero blocks, its total and whether its pages are mapped (sparse pools). With a folder, also the raw storage of the first one of each format and size (`zb_<serial>.bin`; 2 MB and a byte-per-4 KB-page map `zbm_<serial>.bin` for sparse textures) |
|
||||
| `DM_NVN_TEXRAW=info` (with `TEXDUMP`) | Raw bytes of the non-render-target textures `TEXINFO` picked out, at the `TEXDUMP` frame |
|
||||
| `DM_NVN_TEXRAW=ser,…` (with `TEXDUMP`) | Raw bytes of those textures (by serial, as `PASSDUMP` prints them) at the `TEXDUMP` frame; a Depth32Float_Stencil8 texture gives its 32-bit float depth plane (a shadow map) |
|
||||
| `DM_NVN_T2T_TRACE=WxH:file` (+ `DM_NVN_T2T_FROM=f`) | After each texture copy into an RGBA16F texture of that size, a GPU-timeline copy of the result into a trace buffer; 240 entries written to `file` (header line, 240 frame numbers, then the texels) |
|
||||
@@ -260,8 +261,9 @@ A first-frame bring-up that worked for Mario: screenshot is black → `TEXDUMP`
|
||||
|
||||
- **Subgroup operations in vertex shaders** (about 4 % of them): `SPIRV-Cross threw an exception: Subgroup builtins are not available in this type of function`. Metal has no SIMD-group functions in vertex functions, and SPIRV-Cross's own subgroup emulation covers only `Elect`. RE Engine scalarizes uniform loads: ballot the active lanes, shuffle the first lane's value to all and take a fast path when all agree. With a one-lane subgroup that is exact, so `dmsh_build.py` rewrites the vertex shaders' SPIR-V that way before SPIRV-Cross (`spv_single_lane`: the ballot is the lane's own bit, shuffles and broadcasts return its own value, all / any are the predicate, all-equal and elect are true, lane ID 0, size 1).
|
||||
- **Geometry shaders** (53) do not translate: SPIRV-Cross writes `unknown main0_out main0()` for the geometry stage, and Metal has none.
|
||||
- **Indirect-count multi-draws** (`nvnCommandBufferMultiDrawElementsIndirectCount` / `MultiDrawArraysIndirectCount`): the scene is GPU-driven. Compute culling writes the draw records (stride 20) and a u32 count, then up to `max` draws (1–49 per call) are made from them. Both were empty, and the G-buffer pass had no draws: the title screen showed only the UI. Metal has no draw count read by the GPU, so each call is recorded as `max` ordinary indirect draws that carry the count's address and their index (capped at 1,024). The replay reads the count once the GPU has written it: at the first such draw after a GPU write it commits the work so far and waits for it (`mdi_wait` in `nvn_mtl.m`), then skips draws past the count. That costs one wait per pass that uses it; hoisting the count into a GPU mask (as `ixc_run` does for index ranges) would remove it.
|
||||
- **Indirect-count multi-draws** (`nvnCommandBufferMultiDrawElementsIndirectCount` / `MultiDrawArraysIndirectCount`): the scene is GPU-driven. Compute culling writes the draw records (stride 20) and a u32 count, then up to `max` draws (1–49 per call) are made from them. Both were empty, and the G-buffer pass had no draws: the title screen showed only the UI. Metal has no draw count read by the GPU, so each call is recorded as `max` ordinary indirect draws that carry the count's address and their index (capped at 1,024). A compute pass copies each such draw's record into the frame arena, or zeros past the count, and the draw reads the copy, so a draw past the count draws nothing (`mdx_run` in `nvn_mtl.m`). Like `ixc_run`, it runs at the first such draw with no copy and takes every later one up to the next write that could change a record or count; copies are keyed by record and count address, since the game draws the same records with two counts. Draws of one multi-draw replayed back to back are encoded without the rest of `draw` (`mdx_next`). The replay used to commit its work and wait for the GPU before the first such draw after each GPU write, then read the count on the CPU (`mdi_wait`, now only for quads, fans and 8-bit indices). In Kamura Village that wait held the submitting thread 14 of its 20 ms a frame on an iPhone (36–54 fps). On the Mac the same scene's submit took 20 ms of wall time for 6.1 ms of work, and now 5.1 ms, all of it work.
|
||||
- **Textures in the sparse virtual pool hold what the CPU wrote into physical pages before mapping them**: the streamed textures (flag 0x20, packaged layout) and plain sampled ones with no packaged layout. The replay took the plain ones for GPU textures, so their Metal copies stayed 0: the hunter's innerwear (ASTC 6x6) came out magenta, the ASTC error colour, and the face black. A sampled texture in a sparse pool (not a render target, image, linear or buffer texture) is now uploaded like a packaged one, through the pool's page table (03, sparse pools).
|
||||
- **Sparse textures have a layout of their own** (`nvn_tex.c` `level_walk`): they are mapped in 64 KB tiles, one row of blocks 128 GOBs / (block height × block depth) wide. A level at least one tile wide, high and deep is laid out with its width padded to whole tiles, the levels from the first smaller one on form the mip tail (packed as usual), and the storage is whole tiles, 64 KB aligned. The game's files hold its streamed textures pre-swizzled that way, and it places each level at the offset the host reports. Laid out unpadded, the 1024×1024 ASTC 6×6 clothes of Hinoa and Minoto (43 GOBs wide instead of 48, level 1 at 0x81000 instead of 0x90000) read 2,183 all-zero blocks in level 0 and 3,043 in level 1, which ASTC decodes as its error colour: their trousers, sleeves and gloves were magenta, worst in Kamura Village. `DM_NVN_ZBLOCK` finds such textures.
|
||||
- **Sparse textures are sampled only where the game has mapped them**: the translator drops TEX's LOD clamp and its residency predicate, so streamed textures were sampled at levels the game had not mapped yet, which read 0 (magenta on the title screen's clothes). A sparse texture is bound as a view starting at its first level whose pages, and those of every smaller level, are all mapped (`sparse_view` in `nvn_mtl.m`). That samples exactly as a minimum LOD would; the view follows each upload.
|
||||
- **Bindless handles in uniform buffers written earlier in the same submission**: the replay reads handles on the CPU, and a write queued earlier in the submission has not landed yet. Inline uniform updates keep where their data sits, so the handle is read from the newest one covering it (`pend_read`). The game queues more than the old limit of 256 writes per submission, after which every address counted as written; pending writes are now indexed by 4 KB page (8,192 per submission).
|
||||
- **Clear of a texture made without render-target use** (a 32x80 packaged one): a whole-level clear upgrades the texture to render-target use first, as format views are upgraded (`tex_upgrade`).
|
||||
|
||||
Reference in new issue
Block a user