diff --git a/AGENTS.md b/AGENTS.md index 52fde40..5fbf8ef 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -4,7 +4,7 @@ - Each game's constants and shader tables form a game profile in `dm/games//`, where `` is the base game's title ID: `game.tsv` plus the tables and `dmsh_chain.h`. Mario's profile is `010015100B514000`. Tools pick the profile with `DM_GAME`; `make -C dm/host` and `build_ipa.sh` pick it with `GAME`. There is no default game: set it every time (the tools stop and list the profiles otherwise). Put per-game values in the profile, not in code ([dm/games/README.md](dm/games/README.md)). - Per-build knowledge is keyed by the NSO build ID: rewriter cuts in `dm/tools/patch_nso.py`, the profile's `build` rows (the builds the NSP script accepts, and their display versions), and host behaviour that uses fixed module addresses checks `module_t.build_id` (from `fixups.tsv`; `dm_main_build()`, `DM_D2R_BUILD` for code inherited from D2R). When a new game build appears, carry addresses over with `dm/tools/nso_fnmap.py` and check them by disassembly; never let fixed addresses run on a build they were not found in. - To analyse or port **another** Switch title, first run `scripts/build_ipa_from_nsp.py` on its NSP folder. With no matching profile, it writes a survey to `/_survey/survey.md`, then a provisional profile guessed from it (`<NSP folder>/<title ID>_build/profile/<title ID>/`, found by the tools through `DM_GAMES_DIR`) and a best-effort IPA that may not boot. Then follow [docs/agents](docs/agents/README.md) (extraction, binary survey, host bring-up failure catalogue, graphics specifics, debugging recipes) and add a profile. -- Several games are often ported at once, each in its own session. When you find a bug or a fix that another game may hit (host services, the NVN or Vulkan layer, the rewriter, the shader chain, the scripts, the machine), add it at the top of `_agent/findings.md` in the folder that holds the NSP folders, then message the other busy sessions (`ListAgents`, then `SendMessage`). Read that file before debugging a new symptom. Don't edit or commit another session's uncommitted files; ask that session first. +- Several games are often ported at once, each in its own session. When you find a bug or a fix that another game may hit (host services, the NVN or Vulkan layer, the rewriter, the shader chain, the scripts, the machine), add it at the top of `_agent/findings.md` in the folder that holds the NSP folders, then message the other busy sessions (`ListAgents`, then `SendMessage`). Read that file before debugging a new symptom. Don't edit or commit another session's uncommitted files; ask that session first. Only the coordinating session runs the multi-game smoke test (`_agent/smoke9.sh`, or copies of it) and pushes: check your own game, then tell it when your change is in the tree and ready. - Scripts the user runs live in `scripts/` at the repository root (now `build_ipa_from_nsp.py`). Each one runs both in Spyder (F5, configured by the settings block at its top) and from a terminal (command-line options override those settings; non-zero exit on error). Tools that only agents or other scripts call stay under `dm/`. - Game files, keys and anything extracted from them live in the game's NSP folder, never in this repository and never in the home folder (it is nearly full): the work folder `<NSP folder>/<id>-work` for extraction, runs and analysis, beside the NSP script's `<id>_build/`. Tool venvs and other large scratch data go on that drive too. Do not print key values. - Language: everything is in English: `README.md`, docs, code comments and other files. Log messages stay as they are, and red log lines start with `【红】`. Exception until a shader-chain rebuild: the files the chain version hashes (`dm/gfx/tools/dmsh_chain.py`, `dmsh_build.py`, `dmsh_fetch.py`, `dmsh_facts.py`, `dmsh_xlate/Program.cs`, `dmsh_xlate.csproj`, `dm/ios/tools/ios_metal.py`) and the `dmsh_chain.h` headers they generate still have bilingual comments; editing them changes every game's chain version. diff --git a/README.md b/README.md index 9573596..7a98d18 100644 --- a/README.md +++ b/README.md @@ -37,7 +37,7 @@ Being ported, playable on the Mac, with iPhone builds not tested yet: - Hades (base game + update 1.0.38233): the script cannot pack it yet. - Hyrule Warriors: Age of Calamity (base game + update 1.3.0, + DLC): up to its first battle. -- Fire Emblem: Three Houses (base game + update 1.2.0, + DLC): through its prologue. Packing it needs a one-time capture of its run-time shaders on the Mac ([docs/development.md](docs/development.md#game-notes)). +- Fire Emblem: Three Houses (base game + update 1.2.0, + DLC): through its prologue. While packing, the script starts the game on the Mac for a few seconds to collect the shaders it builds as it runs. | | | |:---:|:---:| diff --git a/dm/games/01007EF00011E000/game.tsv b/dm/games/01007EF00011E000/game.tsv index 3535d41..857d307 100644 --- a/dm/games/01007EF00011E000/game.tsv +++ b/dm/games/01007EF00011E000/game.tsv @@ -23,6 +23,8 @@ pad_layout full # Layout captures (CaptureOn panes, about 140 layouts) keep their blurred alpha in the red channel of an R8 render target mapped R 0 0 1; every pixel-shader variant of the layout library # (main's BNSH at 0x1d83000) reads such textures with alpha = red: soft prompt shadows and glows instead of hard-edged boxes and cloud-filled rectangles alpha_from_red d898eb245efdd495 e6a88e12dbae9743 be1d4af059205375 3a0122d004ea74ca 57984f1f48115868 f29dde3d00fb2972 cc37329bf8b7d773 87d900168952ea49 aa209f57b133e862 af07227eeb90304b adcbe3846d1d1ed7 43593beb906b7c6f 11bb9006c60d49d6 e7004fcb0aa6ce8c 896bfbc0413048c9 e1d0faffe4d027a2 8064cd6f18b1b5ed 71c2b3751a9e65f7 3dfe1aa3d996e91e e8ec1269c0f3f2db 302afdb62561cc5c 31e27998ede8f19f 0d532fb24cf738e5 6599d9842b242912 9a09e7e1e7c4cf84 5c490d6406951546 +# The frame-interval controller (60 fps cuts in patch_nso.py) and dynamic resolution read the GPU time between the frame's two timestamp reports: each report takes its segment's GPU start / end +gpu_timestamps segment # ---- 版本 | Versions ---- # 只解了合并 1.9.0 的那份；基础版单独的构建号用到时再补 | only the merged 1.9.0 copy is extracted; the base game's own build IDs are added when needed diff --git a/dm/games/0100F2C0115B6000/game.tsv b/dm/games/0100F2C0115B6000/game.tsv index 68808f0..729e62a 100644 --- a/dm/games/0100F2C0115B6000/game.tsv +++ b/dm/games/0100F2C0115B6000/game.tsv @@ -38,6 +38,8 @@ app_pool 0xCD500000 # ---- Graphics ---- # Streamed textures are written by the CPU straight into pool memory (already in GPU layout) and published with nvnMemoryPoolFlushMappedRange (hle_nvn.c ft_add) tex_flush_upload on +render_cond on +tex_dirty_upload on # ---- NVN driver constant segments (Breath of the Wild's values, the same NVN build) ---- nvn_tex_base 0x20 diff --git a/dm/games/README.md b/dm/games/README.md index 689579c..2a70056 100644 --- a/dm/games/README.md +++ b/dm/games/README.md @@ -47,11 +47,14 @@ None of them has a default game: without one they stop and list the profiles. To | `feedback_snap` | Fragment shader keys (prefixes) whose draws sample a colour attachment of their own render pass and must read a snapshot instead (`fb_prepare` in `nvn_mtl.m`). Writing and sampling one attachment in a render pass is undefined in Metal: the Mac happens to read the old content, the iPhone sometimes reads whole 16×16 blocks as magenta. Each such draw costs a pass split and a copy, so only the shaders seen to break are listed (Super Mario 3D World + Bowser's Fury: its final composite). `DM_NVN_FB_SNAP=0` / `all` / `<keys>` overrides, also from `Documents/env.txt` | | `alpha_from_red` | Fragment shader keys (prefixes, up to 48) that read a single-channel render target whose channel mapping is exactly (R, 0, 0, 1) as (R, 0, 0, R) (`tex_alpha_red` in `nvn_mtl.m`). Breath of the Wild's layout captures (panes tagged `A8`) keep a blurred alpha mask in the red channel of such an R8 texture, and every layout pane that samples one treats it as alpha; read literally, HUD prompts get hard-edged shadow boxes and cloud-filled glow rectangles. BotW lists all 26 pixel-shader variants of its layout library. `DM_NVN_ALPHA_RED=0` / `<keys>` overrides, also from `Documents/env.txt` | | `tex_flush_upload` | `on` / `off` (the default): a texture the game fills from the CPU in place (no packaged layout, no `WriteTexels`) and publishes with `nvnMemoryPoolFlushMappedRange` is uploaded from guest memory like a packaged one, and re-uploaded when its pages change (`ft_add` in `hle_nvn.c`). Tears of the Kingdom streams all its textures this way; without it the 3D scene is black. Off by default because a flushed range in another game could cover an image a shader writes, whose GPU content the upload would replace | +| `render_cond` | `on` / `off` (the default): `SetRenderEnableConditional` / `SetRenderEnable` are recorded and evaluated at replay, and a timestamp report's counter field is 0 as on the GPU (`OP_RENDER_COND` in `nvn_mtl.m`). Tears of the Kingdom wraps its clears in conditions that skip clearing a target nothing has drawn into since its last clear; without the key every clear runs and its 3D scene is black. Off by default: other games draw everything, as before | +| `tex_dirty_upload` | `on` / `off` (the default), needs `tex_flush_upload`: a registered texture whose pages the CPU writes after `nvnTextureInitialize` without any NVN flush is uploaded as well (its pages are write-protected at registration, checked on first use; `nvn_ft_written` in `hle_nvn.c`). Tears of the Kingdom fills its ASTC terrain material arrays this way; without it the cave is magenta | +| `gpu_timestamps` | What a timestamp report (`nvnCommandBufferReportCounter`) records: `submission` (the default) = the time the whole submission completed, the same for every report in it; `segment` = the GPU start (`TIMESTAMP_TOP`) or end (the other types) of the segment the report was encoded in (the host commits a submission in segments of 64 draws, `sub_cut` in `nvn_mtl.m`), so two reports of one submission bracket the GPU time between them, as on hardware. Breath of the Wild needs `segment`: its frame-interval controller and its dynamic resolution read the GPU time between the frame's first and last report, and with one time for both they read 0 (the game kept 60 fps in slow motion when the GPU could not keep up). `DM_NVN_GPU_TS=segment` / `submission` overrides it, also from `Documents/env.txt` | | `nvn_tex_base`, `nvn_tex_slots`, `nvn_samp_base`, `nvn_img_base`, `nvn_img_slots`, `nvn_ssbo_bases` | NVN driver constant segments, measured over all shaders | | `shader_groups`, `shader_note` | `shader_split.py` source groups and the map header note (Engage: `engine` for `data.unity3d` and `Resources/`, `assets` for the Addressables bundles) | | `shader_main_glslc` | `on` / `off` (the default): `shader_split.py` also reads GLSLC compiler outputs (magic `0x19866891`) in `main`. 13 Sentinels keeps all its shaders there; Engage's `main` has 11 that it does not use, so the key stays off for other games | | `shader_separable` | `on` / `off` (the default): the game builds single-stage programs and combines them per draw, so the real vertex / pixel pairs exist only at run time. `shader_split.py` then pairs every vertex shader with a cover of pixel shaders that together read every input any pixel shader reads, and `dmsh_build.py`'s `vs_pad` makes every vertex shader link with any pixel shader (13 Sentinels) | -| `shader_glslc_capture` | `on` / `off` (the default): the game compiles some shaders at run time with its own GLSLC (Fire Emblem: Three Houses: 9 built-in KTGL shaders at boot), so no game file holds them. They are captured once on the Mac with `DM_GLSLC_CAPTURE=<dir>` (`dm/host/hle_glslc.c`), and `shader_split.py` reads `<version dir>/glslc_capture/*.glslc` (beside that version's `RomFS`) or `$DM_GLSLC_CAPTURE`; without the folder those shaders get no translation and the tables differ from the profile's | +| `shader_glslc_capture` | `on` / `off` (the default): the game compiles some shaders at run time with its own GLSLC (Fire Emblem: Three Houses: 9 built-in KTGL shaders at boot), so no game file holds them. They are captured on the Mac with `DM_GLSLC_CAPTURE=<dir>` (`dm/host/hle_glslc.c`; the NSP script does this itself, booting the game in the Mac host until every capture file the profile's `shader_map.tsv` names is there), and `shader_split.py` reads `<version dir>/glslc_capture/*.glslc` (beside that version's `RomFS`) or `$DM_GLSLC_CAPTURE`; without the folder those shaders get no translation and the tables differ from the profile's | The translator (`dmsh_xlate/Program.cs`) still has Mario's texture / sampler bases and the depth convention built in. It uses the bases only to look up per-slot texture / image types in `shader_facts.tsv`; `dmsh_build.py` decodes the bindings from the profile. So a profile with other bases (Engage: `0x168` / `0x568`) builds as long as its facts table has no per-slot texture / image entries; otherwise `dmsh_build.py` stops, and the translator must take the bases from the profile first. diff --git a/dm/games/game.py b/dm/games/game.py index 176bcdd..8cbaa28 100644 --- a/dm/games/game.py +++ b/dm/games/game.py @@ -99,8 +99,14 @@ def header(g): if any(not re.fullmatch(r'[0-9a-f]{8,16}', k) for k in fbs): sys.exit(f'{g.id}: feedback_snap 要写片元着色器键（8–16 位小写十六进制）| feedback_snap takes fragment shader keys (8-16 lowercase hex digits): {fbs}') afr = g.words('alpha_from_red', '') # fragment shader key prefixes that read single-channel textures mapped (R, 0, 0, 1) with alpha = red (nvn_mtl.m tex_alpha_red); none by default if any(not re.fullmatch(r'[0-9a-f]{8,16}', k) for k in afr): sys.exit(f'{g.id}: alpha_from_red 要写片元着色器键（8–16 位小写十六进制）| alpha_from_red takes fragment shader keys (8-16 lowercase hex digits): {afr}') + rcd = g.get('render_cond', 'off') # NVN conditional rendering evaluated at replay, timestamp reports with counter 0 (nvn_mtl.m OP_RENDER_COND): on / off + if rcd not in ('on', 'off'): sys.exit(f'{g.id}: render_cond 只能是 on / off | render_cond must be on or off: {rcd}') tfu = g.get('tex_flush_upload', 'off') # textures written by the CPU into a pool and published with nvnMemoryPoolFlushMappedRange are uploaded like packaged ones (hle_nvn.c): on / off if tfu not in ('on', 'off'): sys.exit(f'{g.id}: tex_flush_upload 只能是 on / off | tex_flush_upload must be on or off: {tfu}') + tdu = g.get('tex_dirty_upload', 'off') # with tex_flush_upload: a registered texture the CPU writes after init without any flush is uploaded too (dirty pages, hle_nvn.c nvn_ft_written): on / off + if tdu not in ('on', 'off') or (tdu == 'on' and tfu != 'on'): sys.exit(f'{g.id}: tex_dirty_upload 只能是 on / off，且要 tex_flush_upload on | tex_dirty_upload must be on or off, and needs tex_flush_upload on: {tdu}') + gts = g.get('gpu_timestamps', 'submission') # counter timestamps (nvn_mtl.m rep_ticks): submission = one time per submission, segment = the GPU start / end of the segment holding the report + if gts not in ('submission', 'segment'): sys.exit(f'{g.id}: gpu_timestamps 只能是 submission / segment | gpu_timestamps must be submission or segment: {gts}') api = g.get('device_api', '55.15') # the NVN API major.minor version nvnDeviceGetInteger reports (hle_nvn.c) if not re.fullmatch(r'\d+\.\d+', api): sys.exit(f'{g.id}: device_api 要写成 主.次 | device_api must be major.minor: {api}') src = f'dm/games/{g.id}' if g.dir.parent == HERE else str(g.dir) @@ -131,6 +137,9 @@ def header(g): f'#define DM_GAME_FEEDBACK_SNAP {q(" ".join(fbs))}', f'#define DM_GAME_ALPHA_FROM_RED {q(" ".join(afr))}', f'#define DM_GAME_TEX_FLUSH_UPLOAD {int(tfu == "on")}', + f'#define DM_GAME_RENDER_COND {int(rcd == "on")}', + f'#define DM_GAME_TEX_DIRTY_UPLOAD {int(tdu == "on")}', + f'#define DM_GAME_GPU_TS_SEGMENT {int(gts == "segment")}', f'#define DM_GAME_NVN_API_MAJOR {api.split(".")[0]}', f'#define DM_GAME_NVN_API_MINOR {api.split(".")[1]}', ''] diff --git a/dm/gfx/nvn/hle_nvn.c b/dm/gfx/nvn/hle_nvn.c index 5b58df5..fae9642 100644 --- a/dm/gfx/nvn/hle_nvn.c +++ b/dm/gfx/nvn/hle_nvn.c @@ -277,7 +277,7 @@ NVN(nvnTextureBuilderGetStorageAlignment, size_t, const NVNtextureBuilder *b) { // textures with content 0, and the G-buffer's base colour came out black. Textures that can hold such data (not render targets, not linear, no packaged layout, in a pool // the CPU can write) are registered by address range; a flush that covers part of one marks it CPU-written (mtl_tex_cpu_written: uploaded from guest memory, re-uploaded // when its pages change). Off by default: in another game a flushed range could cover an image a shader writes, whose GPU content the upload would replace -typedef struct { const uint8_t *mem; uint64_t size; tex_t t; } flush_tex; +typedef struct { const uint8_t *mem; uint64_t size; tex_t t; uint64_t stamp; } flush_tex; // stamp: dirty-page generation at registration (tex_dirty_upload) static flush_tex *g_ft; static size_t g_nft, g_cft; static pthread_mutex_t g_ft_mu = PTHREAD_MUTEX_INITIALIZER; static void ft_add(const tex_t *t) { // Buffer textures are views onto the pool buffer the GPU reads directly (tex_mtl in nvn_mtl.m); an upload would push Age of Calamity's 1,048,576-texel R32F one through tex_desc, past Metal's 2D limit @@ -285,7 +285,8 @@ static void ft_add(const tex_t *t) { (t->d.flags & (NVN_TEXTURE_FLAGS_LINEAR | NVN_TEXTURE_FLAGS_LINEAR_RENDER_TARGET | NVN_TEXTURE_FLAGS_COMPRESSIBLE))) return; pthread_mutex_lock(&g_ft_mu); if (g_nft == g_cft && !(g_ft = realloc(g_ft, (g_cft = g_cft ? 2 * g_cft : 1024) * sizeof *g_ft))) die("flush_tex 分配失败"); - g_ft[g_nft++] = (flush_tex){t->mem, t->L.size, *t}; + uint64_t st = nvn_dirty_now(); if (DM_GAME_TEX_DIRTY_UPLOAD) nvn_dirty_protect(t->mem, t->L.size); // see nvn_ft_written + g_ft[g_nft++] = (flush_tex){t->mem, t->L.size, *t, st}; pthread_mutex_unlock(&g_ft_mu); } static void ft_del(uint64_t serial) { @@ -294,6 +295,19 @@ static void ft_del(uint64_t serial) { for (size_t i = 0; i < g_nft; i++) if (g_ft[i].t.serial == serial) { g_ft[i] = g_ft[--g_nft]; break; } pthread_mutex_unlock(&g_ft_mu); } +// Profile tex_dirty_upload: Tears of the Kingdom fills some textures from the CPU after creating them and publishes them with no NVN flush at all (os::FlushDataCache runs +// natively; its 121-layer ASTC terrain material arrays, 42 MB each). They stayed GPU-only textures with content 0, which ASTC decodes as the error colour, and the cave came out magenta. +// ft_add write-protects a registered texture's pages, and on its first resolution (tex_mtl, submit thread) a write since then makes it CPU-written like a flushed one +int nvn_ft_written(const tex_t *t) { + if (!DM_GAME_TEX_DIRTY_UPLOAD || !t->serial) return 0; + int r = 0; pthread_mutex_lock(&g_ft_mu); + for (size_t i = 0; i < g_nft; i++) if (g_ft[i].t.serial == t->serial) { + if (nvn_dirty_since(g_ft[i].mem, g_ft[i].size, g_ft[i].stamp)) { r = 1; g_ft[i] = g_ft[--g_nft]; } + break; } + pthread_mutex_unlock(&g_ft_mu); + if (r) mtl_tex_cpu_written(t); + return r; +} static void ft_flush(const uint8_t *a, uint64_t n) { if (!DM_GAME_TEX_FLUSH_UPLOAD || !a || !n) return; pthread_mutex_lock(&g_ft_mu); @@ -311,7 +325,7 @@ NVN(nvnTextureInitialize, NVNboolean, NVNtexture *t, const NVNtextureBuilder *b) o->has_layout = s->has_layout; memcpy(o->layout, s->layout, 8); o->set_swz = s->set_swz; o->swz = s->swz; nvn_tex_layout(&o->d, &o->L); o->mem = s->pool && s->pool->mem ? s->pool->mem + s->off : NULL; - o->serial = nvn_new_serial(); + o->serial = nvn_new_serial(); mtl_tex_init(o); // See SetPackagedTextureData. Linear textures are not re-uploaded as packaged: in a CPU-visible pool they are created on the pool buffer (tex_linear) and read what was copied in if (s->pkg_data && o->mem) { memcpy(o->mem, s->pkg_data, o->L.size); @@ -1581,8 +1595,15 @@ NVN(nvnCommandBufferReportValue, void, NVNcommandBuffer *c, uint32_t v, NVNbuffe NVN(nvnCommandBufferClearBuffer, void, NVNcommandBuffer *c, NVNbufferAddress a, size_t n, uint32_t v) { *REC(c, OP_CLEAR_BUF, p_clear_buf) = (p_clear_buf){a, n, v, 0}; } // State the replay layer does not have yet: logged once (conditional rendering is treated as "always draw") -NVN(nvnCommandBufferSetRenderEnable, void, NVNcommandBuffer *c, uint32_t en) { (void)c; if (!(en & 0xff)) ONCE_LOG("【红】SetRenderEnable(false)：宿主暂按恒画"); } -NVN(nvnCommandBufferSetRenderEnableConditional, void, NVNcommandBuffer *c, NVNconditionalRenderMode m, NVNbufferAddress a) { (void)c; (void)m; (void)a; ONCE_LOG("【红】SetRenderEnableConditional：宿主暂按恒画"); } +// Render enable (profile render_cond, see OP_RENDER_COND in nvn_mtl.m): recorded and evaluated at replay; without the key everything draws, as before +NVN(nvnCommandBufferSetRenderEnable, void, NVNcommandBuffer *c, uint32_t en) { + if (DM_GAME_RENDER_COND) *REC(c, OP_RENDER_COND, p_render_cond) = (p_render_cond){(en & 0xff) ? 2 : 3, 0, 0}; + else if (!(en & 0xff)) ONCE_LOG("【红】SetRenderEnable(false)：宿主暂按恒画"); +} +NVN(nvnCommandBufferSetRenderEnableConditional, void, NVNcommandBuffer *c, NVNconditionalRenderMode m, NVNbufferAddress a) { + if (DM_GAME_RENDER_COND) *REC(c, OP_RENDER_COND, p_render_cond) = (p_render_cond){(int)m & 1, 0, a}; + else ONCE_LOG("【红】SetRenderEnableConditional：宿主暂按恒画"); +} // Fire Emblem: Three Houses turns depth clamp on for its scenes; ignored, the ground right under its close-up battle cameras was clipped away and the lighting pass painted the sky's fog colour there NVN(nvnCommandBufferSetDepthClamp, void, NVNcommandBuffer *c, uint32_t en) { REC(c, OP_DEPTH_CLAMP, p_depth_clamp)->on = (en & 0xff) != 0; } NVN(nvnCommandBufferSetLineWidth, void, NVNcommandBuffer *c, float w) { (void)c; if (w != 1.0f) ONCE_LOG("SetLineWidth %.2f：Metal 只有 1 像素线", w); } diff --git a/dm/gfx/nvn/nvn_dirty.c b/dm/gfx/nvn/nvn_dirty.c index 4ccb5e6..7f2a0fb 100644 --- a/dm/gfx/nvn/nvn_dirty.c +++ b/dm/gfx/nvn/nvn_dirty.c @@ -20,12 +20,16 @@ enum { PG_NONE = 0, PG_PROT = 1, PG_BUSY = 2 }; enum { BUSY_WAIT_MS = 2000 }; // Tracked ranges = the memory pools (registered by MemoryPoolInitialize). After a pool is destroyed its live=0 slot is reused; the signal handler scans without locks. -#define MAX_RANGE 4096 // was 256: with more memory pools than that the excess pools are untracked and their textures are treated as dirty and re-uploaded on every use (for Diablo's measured peak see "memory pools" in the memory readout line; Zelda has about 300 and hit this). The scan only covers slots in use, so a larger capacity costs nothing +// The table grows in blocks of 1,024 slots, up to 262,144; a block is never moved or freed, so the signal handler reads it without locks. A pool past the end goes untracked: its textures count +// as dirty and are re-uploaded at every use. The table was a fixed 256 slots (Zelda has about 300 pools at the title), then a fixed 4,096, and Breath of the Wild reaches 4,006 pools in its first +// shrine. The scans only cover slots in use, so the capacity itself costs nothing +enum { BLK_SHIFT = 10, BLK = 1 << BLK_SHIFT, NBLK = 256, MAX_RANGE = BLK * NBLK }; // Ranges are rounded to host pages (16 KB) but guest memory pools are only 4 KB aligned: adjacent pools share their first / last pages, so one address can fall in several registered ranges. // So: "is this texture tracked" looks for a range that holds its whole span; a fault / host pre-write unprotects and stamps every range covering that page; pool destruction identifies its range by the pool's exact start mem. // (Previously the first range containing the address was always taken: two SMBW ASTC 6x5 textures whose pool's first page is shared with a small preceding pool were judged untracked and re-uploaded on every use, about 6 MB per frame) typedef struct { uintptr_t base, end, mem; _Atomic uint8_t *state; _Atomic uint64_t *gen; size_t cap; _Atomic int live; } range_t; -static range_t g_r[MAX_RANGE]; +static _Atomic(range_t *) g_blk[NBLK]; +static inline range_t *rg(int i) { return &atomic_load_explicit(&g_blk[i >> BLK_SHIFT], memory_order_acquire)[i & (BLK - 1)]; } static _Atomic int g_nr; static uintptr_t g_pg, g_shift; static _Atomic uint64_t g_now = 1; @@ -56,11 +60,11 @@ static int busy_wait(range_t *r, uintptr_t p) { } return 1; } -static int live_at(int i) { return atomic_load_explicit(&g_r[i].live, memory_order_acquire); } +static int live_at(int i) { return atomic_load_explicit(&rg(i)->live, memory_order_acquire); } // The registered range that holds all of [a, a+len) (used for a texture's tracking / protection / dirty check) static range_t *find_span(uintptr_t a, size_t len) { int n = atomic_load(&g_nr); - for (int i = 0; i < n; i++) if (live_at(i) && a >= g_r[i].base && a + len <= g_r[i].end) return &g_r[i]; + for (int i = 0; i < n; i++) { range_t *r = rg(i); if (live_at(i) && a >= r->base && a + len <= r->end) return r; } return NULL; } static void unprotect_page(range_t *r, uintptr_t p) { // The caller has already changed the state from PROT to BUSY @@ -83,14 +87,14 @@ static int on_fault(void *addr, int is_write, void *uc) { if (!is_write) return 0; uintptr_t a = (uintptr_t)addr; int ours = 0, n = atomic_load(&g_nr); for (int i = 0; i < n; i++) // Every range sharing the page must be released: releasing only one leaves the other range's state at "protected", and later writes no longer fault and are missed - if (live_at(i) && a >= g_r[i].base && a < g_r[i].end) ours |= release_page(&g_r[i], (a - g_r[i].base) >> g_shift, &g_faults); + { range_t *r = rg(i); if (live_at(i) && a >= r->base && a < r->end) ours |= release_page(r, (a - r->base) >> g_shift, &g_faults); } return ours; } static void prep_write(void *addr, size_t len) { uintptr_t a = (uintptr_t)addr, e = a + len; int n = atomic_load(&g_nr); if (!len) return; for (int i = 0; i < n; i++) { - range_t *r = &g_r[i]; + range_t *r = rg(i); if (!live_at(i) || e <= r->base || a >= r->end) continue; uintptr_t lo = a > r->base ? a : r->base, hi = e < r->end ? e : r->end; for (uintptr_t p = (lo - r->base) >> g_shift, q = (hi - 1 - r->base) >> g_shift; p <= q; p++) release_page(r, p, &g_preps); @@ -104,7 +108,7 @@ void nvn_dirty_add_range(const void *mem, size_t size) { pthread_mutex_lock(&g_slot_mu); int n = atomic_load(&g_nr); int slot = -1; - for (int i = 0; i < n; i++) if (!atomic_load_explicit(&g_r[i].live, memory_order_acquire)) { slot = i; break; } + for (int i = 0; i < n; i++) if (!live_at(i)) { slot = i; break; } if (slot < 0) { if (n >= MAX_RANGE) { int expected = 0; @@ -112,9 +116,13 @@ void nvn_dirty_add_range(const void *mem, size_t size) { pthread_mutex_unlock(&g_slot_mu); return; } slot = n; + if (!atomic_load_explicit(&g_blk[n >> BLK_SHIFT], memory_order_acquire)) { // published before g_nr counts the slot (the scans read g_nr first) + range_t *b = calloc(BLK, sizeof *b); if (!b) die("脏页跟踪区间表分配失败（第 %d 块）", n >> BLK_SHIFT); + atomic_store_explicit(&g_blk[n >> BLK_SHIFT], b, memory_order_release); + } } { int lv = atomic_fetch_add(&g_live, 1) + 1, pk = atomic_load(&g_live_peak); while (lv > pk && !atomic_compare_exchange_weak(&g_live_peak, &pk, lv)) {} } - range_t *r = &g_r[slot]; + range_t *r = rg(slot); uintptr_t base = (uintptr_t)mem & ~(g_pg - 1), end = ((uintptr_t)mem + size + g_pg - 1) & ~(g_pg - 1); size_t np = (end - base) >> g_shift; _Atomic uint8_t *state = r->state; @@ -136,7 +144,7 @@ void nvn_dirty_add_range(const void *mem, size_t size) { void nvn_dirty_remove_range(const void *mem) { pthread_mutex_lock(&g_slot_mu); range_t *r = NULL; int n = atomic_load(&g_nr); - for (int i = 0; i < n && !r; i++) if (live_at(i) && g_r[i].mem == (uintptr_t)mem) r = &g_r[i]; // Match by the pool's exact start: an adjacent pool sharing a page may have the same rounded base + for (int i = 0; i < n && !r; i++) if (live_at(i) && rg(i)->mem == (uintptr_t)mem) r = rg(i); // Match by the pool's exact start: an adjacent pool sharing a page may have the same rounded base if (r) { prep_write((void *)r->base, r->end - r->base); atomic_store_explicit(&r->live, 0, memory_order_release); atomic_fetch_sub(&g_live, 1); } pthread_mutex_unlock(&g_slot_mu); } @@ -169,7 +177,7 @@ int nvn_dirty_since(const void *addr, size_t len, uint64_t stamp) { } void nvn_dirty_slot_stats(int *live, int *used_slots, uint64_t *adds_total) { int n = atomic_load(&g_nr), active = 0; - for (int i = 0; i < n; i++) active += atomic_load_explicit(&g_r[i].live, memory_order_acquire) != 0; + for (int i = 0; i < n; i++) active += live_at(i) != 0; *live = active; *used_slots = n; *adds_total = atomic_load(&g_adds_total); } void nvn_dirty_stats(uint64_t *faults, uint64_t *preps) { *faults = atomic_load(&g_faults); *preps = atomic_load(&g_preps); } diff --git a/dm/gfx/nvn/nvn_mtl.m b/dm/gfx/nvn/nvn_mtl.m index 396648c..04e6745 100644 --- a/dm/gfx/nvn/nvn_mtl.m +++ b/dm/gfx/nvn/nvn_mtl.m @@ -4,6 +4,7 @@ #import <QuartzCore/CAMetalDisplayLink.h> #import <ImageIO/ImageIO.h> #import <UniformTypeIdentifiers/UniformTypeIdentifiers.h> +#include <mach/mach_time.h> #include <dlfcn.h> #include <sys/stat.h> #include <sys/mman.h> @@ -407,7 +408,35 @@ int mtl_gpu_addr(uint64_t addr, uint64_t len, void **buf, uint64_t *off, uint64_ // ---------- Submission ---------- static void defer_reset(void); static void defer_fallback(void); -typedef struct { int kind; volatile uint32_t *val; uint32_t mode, v; uint64_t addr; uint32_t type; void *vis; } post_op; // vis: the visibility slot table for kind 3 +typedef struct { int kind; volatile uint32_t *val; uint32_t mode, v; uint64_t addr; uint32_t type; void *vis; uint32_t seg; } post_op; // vis: the visibility slot table for kind 3; seg: the segment it was encoded in (sub_cut) +static uint32_t g_seg; // Index of the segment being encoded in this submission (0 until the first sub_cut) +typedef struct seg_clk seg_clk; static seg_clk *g_sclk; static void seg_clk_put(seg_clk *c); // Counter timestamps by segment (see gpu_ticks) +// ---------- Conditional rendering (profile render_cond) ---------- +// SetRenderEnableConditional(mode, rec) enables rendering if the 64-bit counter of the report at rec equals (RENDER_IF_EQUAL) or differs from (RENDER_IF_NOT_EQUAL) the one at +// rec + 16, until the next SetRenderEnable(true). Tears of the Kingdom wraps nearly every ClearColor / ClearDepthStencil this way: rec is set to 1 from the CPU when the target +// needs a clear, each clear is followed by ReportCounter(TIMESTAMP_TOP, rec), and rec + 16 is never written. A timestamp report puts 0 in the counter field (the time goes in the +// second field), so after one clear the two compare equal and later conditional clears of that target are skipped. The host used to ignore the condition and write the time +// into both fields, so every clear ran: the first tile-lighting round, the half-resolution scene depth its light volumes test against and the light accumulation were cleared +// before the composite read them, and the 3D scene came out black. +// Metal has no predication, so the condition is evaluated here at replay, on the CPU, from the records in memory. A report replayed earlier in this submission is not written yet +// (that happens when the GPU completes): a timestamp one counts with its known counter 0; a samples-passed result (an occlusion query, written from the visibility buffers) is +// unknown, and a condition on it keeps drawing, as before (occlusion culling only saves work). Inside a block evaluated false, draws and clears are dropped +static int g_rc_skip; // a render-enable block evaluated false is active +#define RC_PEND 256 +static struct { uint64_t addr; int kind; } g_rc_pend[RC_PEND]; static int g_rc_npend; // reports replayed in this submission: kind 1 = timestamp (counter 0), 2 = another counter (unknown) +static uint64_t g_rc_occ[4096]; // addresses ever written by a non-timestamp report (open addressing): conditions on them keep drawing +static void rc_occ_add(uint64_t a) { for (uint32_t i = (uint32_t)(a >> 4) & 4095, k = 0; k < 4096; k++, i = (i + 1) & 4095) { if (g_rc_occ[i] == a) return; if (!g_rc_occ[i]) { g_rc_occ[i] = a; return; } } } +static int rc_occ_has(uint64_t a) { for (uint32_t i = (uint32_t)(a >> 4) & 4095, k = 0; k < 4096; k++, i = (i + 1) & 4095) { if (g_rc_occ[i] == a) return 1; if (!g_rc_occ[i]) return 0; } return 0; } +static void rc_note(uint64_t a, int kind) { + if (kind == 2) rc_occ_add(a); + for (int i = 0; i < g_rc_npend; i++) if (g_rc_pend[i].addr == a) { g_rc_pend[i].kind = kind; return; } + if (g_rc_npend < RC_PEND) { g_rc_pend[g_rc_npend].addr = a; g_rc_pend[g_rc_npend++].kind = kind; } +} +static int rc_value(uint64_t a, uint64_t *v) { // 0 = unknown at encode time + if (rc_occ_has(a)) return 0; + for (int i = g_rc_npend - 1; i >= 0; i--) if (g_rc_pend[i].addr == a) { if (g_rc_pend[i].kind != 1) return 0; *v = 0; return 1; } + *v = *(const volatile uint64_t *)a; return 1; +} // ---------- Occlusion queries (NVN_COUNTER_TYPE_SAMPLES_PASSED) ---------- // NVN semantics: samples passed is a cumulative counter on the queue - ResetCounter zeroes it, every draw afterwards adds to it, and ReportCounter writes out the cumulative value at that moment (+ timestamp) when the GPU completes. // Metal: every render pass attaches a visibility result buffer; draws during counting use counting mode, with a new slot per pass (no reliance on accumulation across passes); a reset empties the 'slot table since the reset', @@ -470,7 +499,7 @@ static int tbdr_on(void); static uint64_t g_sub_serial; // Submission index (used by the parallel precompile's 'predict only once per submission') void mtl_sub_begin(void) { - mtl_init(); g_sub_serial++; + mtl_init(); g_sub_serial++; g_seg = 0; seg_clk_put(g_sclk); g_sclk = NULL; g_rc_npend = 0; g_rc_skip = 0; if (tbdr_on() == 0 && g_ruler_pending) { g_ruler_truth = 0; g_ruler_pending = 0; MLOG("DM_NVN_RULER：切到被测模式"); } defer_drain(); defer_reset(); @@ -636,13 +665,33 @@ static void post(post_op o) { g_post[g_npost++] = o; } void mtl_sub_event(volatile uint32_t *val, uint32_t mode, uint32_t v) { post((post_op){1, val, mode, v, 0, 0, NULL}); } -void mtl_sub_counter(uint64_t addr, uint32_t type) { post((post_op){2, NULL, 0, 0, addr, type, NULL}); } +void mtl_sub_counter(uint64_t addr, uint32_t type) { post((post_op){2, NULL, 0, 0, addr, type, NULL, g_seg}); } void mtl_counter_stats(uint64_t rep[32], uint64_t rst[32]) { memcpy(rep, g_cnt_report, sizeof g_cnt_report); memcpy(rst, g_cnt_reset, sizeof g_cnt_reset); } // Counter timestamps (19.2 MHz) are taken from the guest clock (hle_os.c guest_ns, the same source as svc GetSystemTick / cntpct): on hardware GPU timestamps and the CPU counter share one timebase, and games compare them. // They used to come from CLOCK_MONOTONIC, which keeps running while the Mac sleeps while the guest clock (mach_absolute_time) does not, so the two differed by the sleep time; once timestamps were converted to ns, Mario halved its frame rate. In graphics unit tests guest_ns is a weak definition returning 0; fall back to the monotonic clock uint64_t guest_ns(void); static uint64_t gpu_ticks(void) { uint64_t g = guest_ns(); if (!g) { struct timespec t; clock_gettime(CLOCK_MONOTONIC, &t); g = (uint64_t)t.tv_sec * 1000000000ull + t.tv_nsec; } return g * 12 / 625; } +// Counter timestamps by segment (profile gpu_timestamps segment; DM_NVN_GPU_TS=segment / submission overrides it). On hardware a timestamp report is written when the GPU reaches it, so two reports of one submission bracket the GPU work between them. +// The host writes all reports of a submission in its completion handler; with one time for all of them a game that measures its GPU time that way reads 0 (Breath of the Wild: frame-interval controller and dynamic resolution). +// In segment mode a report takes the GPU time of the segment it was encoded in (sub_cut): TIMESTAMP_TOP (0x10) that segment's GPU start, other types its GPU end. The writes stay in the last segment's handler, before the fence, +// so a game that waits for the fence reads one submission's reports together; a segment whose feedback has not arrived by then takes the last segment's times +#define SEG_CLK_N 1024 +struct seg_clk { _Atomic int refs; _Atomic uint64_t t[SEG_CLK_N][2]; }; // per submission (g_sclk, allocated at its first cut): guest ns of each cut segment's GPU start / end (0 = no feedback yet) +static int seg_ts(void) { static int v = -1; if (v < 0) { const char *e = getenv("DM_NVN_GPU_TS"); v = e && *e ? !strcmp(e, "segment") : DM_GAME_GPU_TS_SEGMENT; MLOG("计数器时间戳：%s（游戏档案 gpu_timestamps；DM_NVN_GPU_TS=segment / submission）", v ? "按所在段的 GPU 起止时间" : "整次提交同一时刻"); } return v; } +static void seg_clk_put(seg_clk *c) { if (c && atomic_fetch_sub(&c->refs, 1) == 1) free(c); } +static uint64_t gpu_s_guest_ns(double s) { // a feedback GPU time (seconds on the mach absolute time base) on the guest clock; 0 when unknown + static mach_timebase_info_data_t tb; if (!tb.denom) mach_timebase_info(&tb); + uint64_t g = guest_ns(), now = tb.denom ? mach_absolute_time() * tb.numer / tb.denom : 0, at = (uint64_t)(s * 1e9); + if (s <= 0 || !g || at > now || now - at > g) return 0; + return g - (now - at); +} +static uint64_t rep_ticks(const post_op *o, seg_clk *sc, uint32_t last, uint64_t c0, uint64_t c1) { // the time a counter report records (19.2 MHz); c0 / c1: the last segment's GPU start / end on the guest clock + int top = o->type == 0x10; uint64_t ns = 0; + if (sc && o->seg < last && o->seg < SEG_CLK_N) ns = atomic_load(&sc->t[o->seg][!top]); + if (!ns) ns = top ? c0 : c1; + return ns ? ns * 12 / 625 : gpu_ticks(); +} // ---------- Boot instrumentation (DM_BOOTPROF=1): metallib loading, pipeline creation (both happen synchronously during replay on the submit thread), texture upload, first-frame present ---------- static uint64_t bp_now(void) { struct timespec t; clock_gettime(CLOCK_MONOTONIC, &t); return (uint64_t)t.tv_sec * 1000000000ull + t.tv_nsec; } @@ -756,8 +805,10 @@ uint64_t mtl_sub_commit(void) { [o addFeedbackHandler:^(id<MTL4CommitFeedback> f) { (void)f; for (dispatch_block_t block in commit_blocks) block(); }]; uint64_t gpu_f = atomic_load(&g_cur_frame); gpu_tag_begin(gpu_f); NSString *cctx = @(cctx_snap()); + seg_clk *sc = g_sclk; uint32_t seg_last = g_seg; g_sclk = NULL; [o addFeedbackHandler:^(id<MTL4CommitFeedback> cb) { gpu_tag_done(gpu_f, cb.GPUStartTime, cb.GPUEndTime); + uint64_t c0 = 0, c1 = 0; if (seg_ts()) { c0 = gpu_s_guest_ns(cb.GPUStartTime); c1 = gpu_s_guest_ns(cb.GPUEndTime); } if (cb.error) { atomic_store(&g_last_commit_error, 1); static int n, nv; int victim = [cb.error.description containsString:@"InnocentVictim"]; if (victim ? nv++ < 1 : n++ < 8) MLOG("【红】命令缓冲错误详情 fence=%llu：%s | GPU %.3f→%.3f", (unsigned long long)fence, cb.error.description.UTF8String, cb.GPUStartTime, cb.GPUEndTime); if (!victim) MLOG("【红】出错那次提交的现场 fence=%llu：%s", (unsigned long long)fence, cctx.UTF8String); MLOG("【红】命令缓冲执行出错：%s", cb.error.localizedDescription.UTF8String); } for (int i = 0; i < n; i++) { if (ops[i].kind == 1) { if (ops[i].mode == 1) *ops[i].val += ops[i].v; else *ops[i].val = ops[i].v; } @@ -766,9 +817,9 @@ uint64_t mtl_sub_commit(void) { for (int k = 0; k < l->n; k++) { sum += ((const uint64_t *)((__bridge id<MTLBuffer>)l->s[k].buf).contents)[l->s[k].slot]; CFRelease(l->s[k].buf); } free(l); uint64_t rec[2] = {sum, gpu_ticks()}; memcpy((void *)ops[i].addr, rec, sizeof rec); } - else { uint64_t t = gpu_ticks(), rec[2] = {ops[i].type == 0 || ops[i].type == 0x10 ? t : 0, t}; memcpy((void *)ops[i].addr, rec, sizeof rec); } + else { uint64_t t = rep_ticks(&ops[i], sc, seg_last, c0, c1), rec[2] = {(ops[i].type == 0 || ops[i].type == 0x10) && !DM_GAME_RENDER_COND ? t : 0, t}; memcpy((void *)ops[i].addr, rec, sizeof rec); } // render_cond: a timestamp report's counter field is 0, as on the GPU (see rc_value) } - free(ops); + free(ops); seg_clk_put(sc); pthread_mutex_lock(&g_mu); if (fence > g_done) g_done = fence; pthread_cond_broadcast(&g_cv); pthread_mutex_unlock(&g_mu); defer_drain(); }]; @@ -881,6 +932,32 @@ void mtl_alias_drop_pool(const void *pool) { static __unsafe_unretained id serial_get(uint64_t k) { pthread_mutex_lock(&g_smu); __unsafe_unretained id o = hm_get_id(&g_by_serial, &k); pthread_mutex_unlock(&g_smu); return o; } static void serial_put(uint64_t k, id o) { pthread_mutex_lock(&g_smu); id old = hm_get_id(&g_by_serial, &k); if (old) { if (alias_has(old)) defer_obj_unref(old); else defer_obj_release(old); } hm_put_id(&g_by_serial, &k, o); g_ser_gen++; pthread_mutex_unlock(&g_smu); } static void serial_del(uint64_t k) { pthread_mutex_lock(&g_smu); id old = hm_get_id(&g_by_serial, &k); if (old) { if (alias_has(old)) defer_obj_unref(old); else defer_obj_release(old); } hm_del_id(&g_by_serial, &k); g_ser_gen++; pthread_mutex_unlock(&g_smu); } +// Texture serials owned by a live NVN texture object (TextureInitialize until Finalize or re-initialization, hle_nvn.c). Command records and texture pool registrations snapshot the +// texture object (tex_snap in hle_nvn.c, mtl_reg_tex), so replay can resolve a texture the game has already finalized: Breath of the Wild records each frame's transient render targets +// (the HUD captures, the 640x360 half-resolution buffers) and finalizes their objects before it submits. Those serials went back into g_by_serial and stayed there for good, as no second +// Finalize comes: about 10 a frame, 200,000 after an hour in the first shrine. A dead serial resolves through g_ghost instead, a small table whose oldest entry goes when a new one comes, +// so a record replayed again soon still finds its texture +enum { NGHOST = 512 }; +static nvn_hmap g_tex_live = {.klen = 8}; // serial → (void *)1, under g_smu +static nvn_hmap g_ghost = {.klen = 8}; static uint64_t g_ghost_ring[NGHOST]; static unsigned g_ghost_next; static uint64_t g_ghost_puts; +void mtl_tex_init(const tex_t *t) { if (!t->serial) return; pthread_mutex_lock(&g_smu); nvn_hm_put(&g_tex_live, &t->serial, (void *)1); pthread_mutex_unlock(&g_smu); } +static void obj_drop(id o) { if (alias_has(o)) defer_obj_unref(o); else defer_obj_release(o); } // caller holds g_smu +static __unsafe_unretained id tex_get(uint64_t k) { + pthread_mutex_lock(&g_smu); __unsafe_unretained id o = hm_get_id(&g_by_serial, &k); if (!o && g_ghost.n) o = hm_get_id(&g_ghost, &k); pthread_mutex_unlock(&g_smu); return o; +} +static void tex_put(uint64_t k, id o) { // serials from 2^62 up are external textures (mtl_ext_tex), never initialized by the game + pthread_mutex_lock(&g_smu); + nvn_hmap *m = k >= (1ull << 62) || nvn_hm_get(&g_tex_live, &k) ? &g_by_serial : &g_ghost; + id old = hm_get_id(m, &k); + if (old) obj_drop(old); + else if (m == &g_ghost) { + uint64_t *slot = &g_ghost_ring[g_ghost_next++ % NGHOST]; + if (*slot) { id ev = hm_get_id(&g_ghost, slot); if (ev) obj_drop(ev); hm_del_id(&g_ghost, slot); } + *slot = k; g_ghost_puts++; + } + hm_put_id(m, &k, o); g_ser_gen++; + pthread_mutex_unlock(&g_smu); +} // Host record of a packaged texture: the current MTLTexture, the upload stamp, and the global write generation at the last check @interface NvPkg : NSObject { @public id<MTLTexture> tex; uint64_t stamp, seen, hash; } @end // hash: full hash of the uploaded content under DM_NVN_DIRTYCHECK @implementation NvPkg @@ -941,10 +1018,10 @@ static id<MTLTexture> tex_upgrade_pfv(const tex_t *t, id<MTLTexture> base, MTLPi else for (NSUInteger s = 0; s < slices; s++) enc_b_copy_texture(b, base, s, l, MTLOriginMake(0, 0, 0), MTLSizeMake(w, h, 1), upgraded, s, l, MTLOriginMake(0, 0, 0)); } - id holder = serial_get(t->serial); + id holder = tex_get(t->serial); if ([holder isKindOfClass:[NvPkg class]]) { id<MTLTexture> old = ((NvPkg *)holder)->tex; ((NvPkg *)holder)->tex = upgraded; defer_obj_release(old); g_ser_gen++; } else { - serial_put(t->serial, upgraded); + tex_put(t->serial, upgraded); alias_key ak; memset(&ak, 0, sizeof ak); ak.pool = (uint64_t)(uintptr_t)t->pool; ak.off = t->off; ak.d = t->d; // Switch the sharing table to the upgraded one too, so later objects on the same storage get it directly pthread_mutex_lock(&g_smu); if (hm_get_id(&g_alias, &ak) == base) { hm_put_id(&g_alias, &ak, upgraded); CFSetAddValue(g_alias_set, (__bridge const void *)upgraded); } pthread_mutex_unlock(&g_smu); } @@ -1282,7 +1359,8 @@ static id<MTLTexture> tex_mtl_(const tex_t *ct) { t->d.width, t->d.height, t->d.depth, t->d.levels, t->d.stride, t->d.pkg_block_h, (unsigned long long)t->L.size, (unsigned long long)nz); } } } - id have = serial_get(t->serial); + id have = tex_get(t->serial); + if (!packaged && !have && DM_GAME_TEX_DIRTY_UPLOAD && t->mem && on_sub_thread()) { int nvn_ft_written(const tex_t *); packaged = nvn_ft_written(t); } // first use of a texture the CPU wrote since init without a flush (hle_nvn.c) if (packaged) { NvPkg *pk = [have isKindOfClass:[NvPkg class]] ? have : nil; // resolved as GPU-only before the CPU wrote it: dropped and rebuilt as packaged if (pk) { // Already uploaded: check the covered pages only if there has been a guest write since the last check @@ -1304,7 +1382,7 @@ static id<MTLTexture> tex_mtl_(const tex_t *ct) { pk->tex = pkg_upload(t, &pk->stamp); if (!pk->tex) return nil; pk->seen = nvn_dirty_now(); g_pkg_uploads++; if (dirtycheck()) pk->hash = full_hash(t->mem, t->L.size); - serial_put(t->serial, pk); + tex_put(t->serial, pk); return pk->tex; } if (have) return have; @@ -1322,12 +1400,12 @@ static id<MTLTexture> tex_mtl_(const tex_t *ct) { id<MTLTexture> mt = [b newTextureWithDescriptor:d offset:mo bytesPerRow:row]; if (!mt) { RED_ONCE(39, "缓冲贴图建不出：%u 元素 格式 %#x", t->d.width, t->d.format); return nil; } res_add(mt); - serial_put(t->serial, mt); + tex_put(t->serial, mt); return mt; } alias_key ak; memset(&ak, 0, sizeof ak); ak.pool = (uint64_t)(uintptr_t)t->pool; ak.off = t->off; ak.d = t->d; pthread_mutex_lock(&g_smu); __unsafe_unretained id shared = hm_get_id(&g_alias, &ak); pthread_mutex_unlock(&g_smu); - if (shared) { serial_put(t->serial, shared); return shared; } + if (shared) { tex_put(t->serial, shared); return shared; } id<MTLTexture> mt = tex_linear(t); if (!mt) { MTLTextureDescriptor *d = tex_desc(t, 0); @@ -1343,7 +1421,7 @@ static id<MTLTexture> tex_mtl_(const tex_t *ct) { if (!g_alias_set) g_alias_set = CFSetCreateMutable(NULL, 0, NULL); hm_put_id(&g_alias, &ak, mt); CFSetAddValue(g_alias_set, (__bridge const void *)mt); pthread_mutex_unlock(&g_smu); - serial_put(t->serial, mt); + tex_put(t->serial, mt); return mt; } // Memory readout line (always on, low frequency, capped): resident memory, GPU allocated, deferred-release queue, residency set object count, registered memory pools / peak / tracking table capacity, cumulative texture re-uploads. @@ -1356,7 +1434,7 @@ void mtl_mem_stat(unsigned long long frame, unsigned long long rss_mb) { pthread_mutex_lock(&g_deferred_mu); unsigned long nd = (unsigned long)g_deferred.count; pthread_mutex_unlock(&g_deferred_mu); pthread_mutex_lock(&g_res_mu); unsigned long nr = (unsigned long)g_res_objects.count; pthread_mutex_unlock(&g_res_mu); int live = 0, peak = 0, cap = 0; nvn_dirty_stat(&live, &peak, &cap); - pthread_mutex_lock(&g_smu); unsigned ns = g_by_serial.n, ashared = g_alias.n; pthread_mutex_unlock(&g_smu); + pthread_mutex_lock(&g_smu); unsigned ns = g_by_serial.n, ashared = g_alias.n, nghost = g_ghost.n; uint64_t ghost_puts = g_ghost_puts; pthread_mutex_unlock(&g_smu); unsigned long dead_n = 0; unsigned long long dead_mb = 0; // Objects read from a weak table come back retained + autoreleased: the frame hook runs on a guest thread with no autorelease pool that drains, so without this wrapper every buffer counted would be pinned for good (the readout creating the leak itself) int pres = 0; @autoreleasepool { pthread_mutex_lock(&g_mu); pres = g_pool_res_n; for (id<MTLBuffer> b in g_dead_pools) { dead_n++; dead_mb += b.length; } pthread_mutex_unlock(&g_mu); } dead_mb >>= 20; // host objects held by serial (textures / packaged textures / programs): if it climbs along with GPU allocated, textures are not being released @@ -1365,7 +1443,7 @@ void mtl_mem_stat(unsigned long long frame, unsigned long long rss_mb) { #if TARGET_OS_IPHONE snprintf(avail, sizeof avail, "%llu MB", (unsigned long long)(os_proc_available_memory() >> 20)); #endif - MLOG("内存 至帧 %llu：常驻 %llu MB，显卡已分配 %llu MB，延迟释放队列 %lu，驻留集 %lu 个，序号对象 %u 个，内存池 %d 个（其中已驻留 %d；峰值 %d / 跟踪表容量 %d；已释放仍活 %lu 个 %llu MB），共用贴图 %u 张（随池淘汰累计 %llu），贴图重传累计 %llu 次，每帧竞技场累计新建 %llu 块；离系统内存上限还有 %s", frame, rss_mb, gpu, nd, nr, ns, live, pres, peak, cap, dead_n, dead_mb, ashared, (unsigned long long)g_alias_evicted, (unsigned long long)g_pkg_reuploads, (unsigned long long)g_arena_made, avail); + MLOG("内存 至帧 %llu：常驻 %llu MB，显卡已分配 %llu MB，延迟释放队列 %lu，驻留集 %lu 个，序号对象 %u 个，内存池 %d 个（其中已驻留 %d；峰值 %d / 跟踪表容量 %d；已释放仍活 %lu 个 %llu MB），共用贴图 %u 张（随池淘汰累计 %llu），贴图重传累计 %llu 次，每帧竞技场累计新建 %llu 块，已 Finalize 仍被引用的贴图 %u 张（小表，累计放入 %llu）；离系统内存上限还有 %s", frame, rss_mb, gpu, nd, nr, ns, live, pres, peak, cap, dead_n, dead_mb, ashared, (unsigned long long)g_alias_evicted, (unsigned long long)g_pkg_reuploads, (unsigned long long)g_arena_made, nghost, (unsigned long long)ghost_puts, avail); if (g_allocprof > 0) { char ab[256]; int o = 0; for (int i = 0; i < g_alloc_n && o < 200; i++) o += snprintf(ab + o, sizeof ab - o, " #%d %llu MB", i, (unsigned long long)(g_alloc4[i].allocatedSize >> 20)); for (int i = 0; i < 3 && o < 240; i++) o += snprintf(ab + o, sizeof ab - o, " 传%d %llu 出%d %llu", i, (unsigned long long)(g_up_alloc4[i].allocatedSize >> 20), i, (unsigned long long)(g_present_alloc4[i].allocatedSize >> 20)); MLOG("命令分配器占用：%s", ab); @@ -1375,7 +1453,11 @@ void mtl_mem_stat(unsigned long long frame, unsigned long long rss_mb) { } void mtl_pkg_stats(uint64_t *uploads, uint64_t *reuploads) { *uploads = g_pkg_uploads; *reuploads = g_pkg_reuploads; } uint64_t mtl_pkg_miss(void) { return g_pkg_miss; } -void mtl_tex_release(tex_t *t) { if (t->serial) serial_del(t->serial); } +void mtl_tex_release(tex_t *t) { + if (!t->serial) return; + pthread_mutex_lock(&g_smu); nvn_hm_del(&g_tex_live, &t->serial); pthread_mutex_unlock(&g_smu); + serial_del(t->serial); +} // ---------- Render passes: SetRenderTargets opens a new pass; clears before the first command in the pass are folded into the load action ---------- static struct { @@ -2586,7 +2668,8 @@ static void state_reset(void) { } // Correct the vertex format by the shader input's component type: Metal requires integer inputs to get integer formats of the same signedness, and float inputs float / normalized formats. -// Integer input (inferred at translation time as "only used as an integer", dmsh_build.py int_inputs) with an integer format: on a signedness mismatch switch to the same width with the other sign (bit pattern unchanged); with a 32-bit float format read it as an integer of the same width (bit pattern unchanged). +// Integer input (inferred at translation time as "only used as an integer", dmsh_build.py int_inputs) with an integer format: on a signedness mismatch switch to the same width with the other sign (bit pattern unchanged for 32-bit components; +// 8 / 16-bit ones would be widened the wrong way, so vconv_wanted converts them first and this switch is only the fallback when the conversion slots run out); with a 32-bit float format read it as an integer of the same width (bit pattern unchanged). // Other mismatches (integer input with a normalized / half format, float input with an integer format) cannot keep the bit pattern: report once and keep the format. static MTLVertexFormat vfmt_for_ct(MTLVertexFormat f, int ct, int nvn_fmt) { static const MTLVertexFormat U[] = { MTLVertexFormatUChar, MTLVertexFormatUChar2, MTLVertexFormatUChar3, MTLVertexFormatUChar4, MTLVertexFormatUShort, MTLVertexFormatUShort2, MTLVertexFormatUShort3, MTLVertexFormatUShort4, MTLVertexFormatUInt, MTLVertexFormatUInt2, MTLVertexFormatUInt3, MTLVertexFormatUInt4 }; @@ -2633,7 +2716,17 @@ static int vconv_ncomp(int f) { // components of an NVN float-type vertex form } return 0; } +// 8- and 16-bit integer formats: Maxwell's fetch widens each component into the 32-bit register by the format's own signedness (sign for SINT, zero for UINT). Metal reads them only into +// an input of the same signedness, and vfmt_for_ct's switch to the other sign keeps the 8 / 16 bits but widens them the other way. So an input of the other signedness takes the 32-bit +// values converted here too. Fire Emblem: Three Houses: the terrain's heights are RGBA16I into a uint input; negative heights read as 32768-65535 and the ground under close-up +// cameras was never drawn (the lighting then saw depth 0 and painted fog there) +static int vconv_int(int f, int *bits, int *sgn) { // components of an 8 / 16-bit integer format (with its component size and signedness); 0 = not one + static const uint8_t T[16] = { 0x03, 0x04, 0x08, 0x09, 0x0F, 0x10, 0x14, 0x15, 0x1B, 0x1C, 0x20, 0x21, 0x27, 0x28, 0x2C, 0x2D }; // per component count: UI8, I8, UI16, I16 + for (int i = 0; i < 16; i++) if (T[i] == f) { *bits = i & 2 ? 16 : 8; *sgn = i & 1; return i / 4 + 1; } + return 0; +} static int vconv_size(int f) { // bytes of one element + int ib, is, in = vconv_int(f, &ib, &is); if (in) return in * ib / 8; switch (f) { case 0x01: case 0x02: return 1; case 0x05: case 0x06: case 0x07: case 0x0D: case 0x0E: return 2; case 0x19: case 0x1A: return 3; case 0x1D: case 0x1E: case 0x1F: return 6; case 0x29: case 0x2A: case 0x2B: return 8; } return 4; @@ -2645,6 +2738,13 @@ static float vconv_uf(uint32_t v, int mbits) { // unsigned small float of R11G return ldexpf(1.0f + (float)m / (float)(1u << mbits), (int)e - 15); } static void vconv_elem(int f, const uint8_t *p, float o[4]) { + int ib, is, in = vconv_int(f, &ib, &is); + if (in) { // integer format: the register's 32-bit value, the missing components (0, 0, 0, 1) as integers + uint32_t u[4] = { 0, 0, 0, 1 }; + for (int i = 0; i < in; i++) { + if (ib == 8) u[i] = is ? (uint32_t)(int32_t)(int8_t)p[i] : p[i]; + else { uint16_t h; memcpy(&h, p + 2 * i, 2); u[i] = is ? (uint32_t)(int32_t)(int16_t)h : h; } } + memcpy(o, u, sizeof u); return; } o[0] = o[1] = o[2] = 0; o[3] = 1; int n = vconv_ncomp(f); for (int i = 0; i < n; i++) switch (f) { @@ -2664,8 +2764,10 @@ static void vconv_elem(int f, const uint8_t *p, float o[4]) { } } static int vconv_wanted(const nvn_refl *rf, int v, const pso_key *k) { // does Metal input v of this pipeline take converted data - int a = rf->v[v].nvn; - return (rf->v[v].ct == NVN_CT_UINT || rf->v[v].ct == NVN_CT_INT) && a < 16 && vconv_ncomp(k ? k->attr[a].fmt : S.attr[a].fmt); + int a = rf->v[v].nvn, ct = rf->v[v].ct, ib, is; + if ((ct != NVN_CT_UINT && ct != NVN_CT_INT) || a >= 16) return 0; + int f = k ? k->attr[a].fmt : S.attr[a].fmt; + return vconv_ncomp(f) || (vconv_int(f, &ib, &is) && is != (ct == NVN_CT_INT)); } // Per submission, one conversion per (range, format, layout): KTGL draws the same mesh several times a frame (depth, shadow, G-buffer) static struct { uint64_t addr, size, serial; uint32_t fmt, off, stride; uint64_t gpu; } g_vconv[64]; static int g_vconv_next; @@ -2683,7 +2785,8 @@ static void vconv_bind(id<MTL4RenderCommandEncoder> en, const nvn_refl *rf) { for (uint64_t i = 0; i < nv; i++) vconv_elem(f, (const uint8_t *)(uintptr_t)(addr + off + i * stride), d + 4 * i); gpu = b.gpuAddress + o; int c = g_vconv_next++ % 64; g_vconv[c].addr = addr; g_vconv[c].size = size; g_vconv[c].serial = g_sub_serial; g_vconv[c].fmt = (uint32_t)f; g_vconv[c].off = off; g_vconv[c].stride = stride; g_vconv[c].gpu = gpu; - RED_ONCE(89, "顶点属性格式 %#x 喂给整数输入：回放时在 CPU 上换成 float32 位模式另绑（Maxwell 取顶点时的转换）", f); + if (vconv_ncomp(f)) RED_ONCE(89, "顶点属性格式 %#x 喂给整数输入：回放时在 CPU 上换成 float32 位模式另绑（Maxwell 取顶点时的转换）", f); + else RED_ONCE(96, "顶点属性格式 %#x（8 / 16 位整数）喂给符号相反的整数输入：回放时在 CPU 上按格式本身的符号扩成 32 位另绑", f); } else if (nv) RED_ONCE(90, "顶点属性换算：范围 %llu 个顶点超出帧内暂存（格式 %#x）", (unsigned long long)nv, f); } if (gpu) enc_r_vbuf(en, (NSUInteger)slot, gpu); @@ -2777,7 +2880,7 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS continue; } int st = k->attr[a].stream; - if (vconv_wanted(rf, v, k) && nconv < VCONV_MAX) { // converted on the CPU at draw time (vconv_bind): float32 bits, 4 components, its own slot + if (vconv_wanted(rf, v, k) && nconv < VCONV_MAX) { // converted on the CPU at draw time (vconv_bind): the register's 32-bit values, 4 components, its own slot int slot = VCONV_SLOT0 + nconv++; vd.attributes[loc].format = rf->v[v].ct == NVN_CT_UINT ? MTLVertexFormatUInt4 : MTLVertexFormatInt4; vd.attributes[loc].offset = 0; vd.attributes[loc].bufferIndex = slot; vd.layouts[slot].stride = 16; @@ -3833,7 +3936,7 @@ static void hoist_try(void) { for (int guard = 0; guard < 200000; guard++) { op_hdr h; memcpy(&h, p, sizeof h); const void *a = p + 8; switch (h.op) { - case OP_END: case OP_CALL: case OP_COPY_T2T: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: return; + case OP_END: case OP_CALL: case OP_COPY_T2T: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: case OP_RENDER_COND: return; case OP_JUMP: { const uint8_t *t; memcpy(&t, a, 8); p = t; continue; } case OP_DRAW: case OP_DRAW_IDX: if (in_grp) return; break; case OP_SCISSOR: sc = *(const p_scissor *)a; has_sc = 1; break; @@ -4141,7 +4244,7 @@ static void defer_try(const p_draw *dw, const p_draw_idx *di) { for (int k = 0; k < g_defer.nreads && !bad; k++) if (h_ser(g_defer.p_reads[k]) == ds_) bad = 1; if (g_defer.nreads > 512) bad = 1; if (bad) { g_defer.reason = 3; g_defer.reason_op = h.op; goto fail; } break; } - case OP_REPORT_COUNTER: case OP_UPDATE_UB: case OP_COPY_B2B: case OP_CLEAR_BUF: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: g_defer.reason = 3; g_defer.reason_op = h.op; goto fail; + case OP_REPORT_COUNTER: case OP_RENDER_COND: case OP_UPDATE_UB: case OP_COPY_B2B: case OP_CLEAR_BUF: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: g_defer.reason = 3; g_defer.reason_op = h.op; goto fail; case OP_CLEAR_COLOR: { const p_clear_color *c = a; if (!p_end) { g_defer.reason = 16; goto fail; } // A clear inside P's continuation segment: do not defer if (c->index >= 0 && c->index < scan_rt.n && (defer_rt_hits_reads(&scan_rt, c->index) || (scan_rt.color[c->index] && defer_rt_hits_reads(&g_defer.p_rt, c->index)))) { g_defer.reason = 5; goto fail; } break; } case OP_CLEAR_DS: if (!p_end) { g_defer.reason = 16; goto fail; } break; @@ -4224,7 +4327,7 @@ static void fuse_try(void) { for (int guard = 0; guard < 200000 && !r2; guard++) { if (!fz_next(&cur, &p, &h, &a)) return; switch (h.op) { - case OP_COPY_T2T: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: + case OP_COPY_T2T: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: case OP_RENDER_COND: case OP_CLEAR_COLOR: case OP_CLEAR_DS: case OP_UPDATE_UB: case OP_COPY_B2B: case OP_CLEAR_BUF: return; case OP_SCISSOR: sc = *(const p_scissor *)a; has_sc = 1; break; case OP_SET_RT: { @@ -4245,7 +4348,7 @@ static void fuse_try(void) { for (int guard = 0; guard < 200000; guard++) { if (h.op == OP_DRAW || h.op == OP_DRAW_IDX) break; switch (h.op) { - case OP_COPY_T2T: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: case OP_SET_RT: case OP_UPDATE_UB: case OP_COPY_B2B: case OP_CLEAR_BUF: return; + case OP_COPY_T2T: case OP_COPY_B2T: case OP_COPY_T2B: case OP_CLEAR_TEX: case OP_DISPATCH: case OP_SET_RT: case OP_UPDATE_UB: case OP_COPY_B2B: case OP_CLEAR_BUF: case OP_RENDER_COND: return; case OP_SCISSOR: sc = *(const p_scissor *)a; has_sc = 1; break; case OP_CLEAR_COLOR: { const p_clear_color *c = a; if (ncl >= 8 || c->index < 0 || c->index >= r2->n || !r2->color[c->index] || c->mask != 0xf) return; const tex_t *t = r2->color[c->index]; @@ -4921,6 +5024,65 @@ static void ixc_run(const p_draw *dw, const p_draw_idx *di) { int prim = e->prim, it = e->itype; unsigned long long fr = (unsigned long long)atomic_load(&g_cur_frame), pass = g_stat_pass, splits = g_ixc_splits; id<MTLBuffer> keep = ab; [g_commit_blocks addObject:^{ (void)keep; MLOG("间接绘制换算 帧 %llu 通道 %llu：图元 %d 索引类型 %d 记录 %u,%u,%u,%u → 参数 %u,%u,%u,%d,%u（累计切开 %llu）", fr, pass, prim, it, src[0], src[1], src[2], src[3], arg[0], arg[1], arg[2], (int)arg[3], arg[4], splits); }]; } } +// ---------- 8-bit indices written by the GPU earlier in the submission ---------- +// Metal has no 8-bit index type, so draw widens 8-bit indices to 16-bit on the CPU at replay. When a buffer write encoded earlier in this submission (copy / update / fill, +// see pend_hits) covers them, the CPU reads the bytes from before that write (Tears of the Kingdom). Those draws are widened on the GPU instead: a compute pass reads the +// 8-bit indices from a 16-byte-aligned address, with the byte shift as a constant (index data may start anywhere), and writes 16-bit ones into the frame arena, which lives +// until the submission completes. As in ixc_run, a compute pass cannot run inside a render pass: the first such draw widens itself and every such draw after it up to the next +// buffer write outside a render pass, so only a draw met inside an open render pass with no widening yet splits it. Draws whose indices no write covers keep the CPU widening. +// Adapted from MacNativeGames' widen_u8_gpu (github.com/ly05010419/MacNativeGames, MIT). +#define WIDE_MAX 256 +typedef struct { uint64_t ibuf; int32_t count; uint64_t gpu; } wide_ent; // gpu: the 16-bit indices in the frame arena (0: not widened) +static wide_ent g_wide[WIDE_MAX]; static int g_wide_n; static uint64_t g_wide_epoch, g_wide_sub, g_wide_runs, g_wide_draws, g_wide_splits; +static const char k_wide_msl[] = +"#include <metal_stdlib>\nusing namespace metal;\n" +"struct wjob { device const uchar *src; device ushort *dst; uint n, shift; };\n" // src: 16-byte aligned, the indices start shift bytes after it +"kernel void wide_u8(device const wjob &j [[buffer(0)]], uint i [[thread_position_in_grid]]) { if (i < j.n) j.dst[i] = ushort(j.src[j.shift + i]); }\n"; // as the CPU widening (0xff stays 0x00ff) +static int wide_wanted(const p_draw_idx *di) { return !di->ind && di->itype == 0 && di->prim != 6 && di->prim != 7 && di->count > 0 && pend_hits(di->ibuf, (uint64_t)di->count); } +static const wide_ent *wide_find(const p_draw_idx *di) { + if (g_wide_epoch != g_wr_epoch || g_wide_sub != g_sub_serial) return NULL; + for (int i = 0; i < g_wide_n; i++) if (g_wide[i].ibuf == di->ibuf && g_wide[i].count == di->count) return &g_wide[i]; + return NULL; +} +static void wide_take(const p_draw_idx *di) { if (g_wide_n < WIDE_MAX && !wide_find(di)) g_wide[g_wide_n++] = (wide_ent){di->ibuf, di->count, 0}; } +static void wide_run(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_wide_msl) options:nil error:&e]; id<MTLFunction> f = lib ? [lib newFunctionWithName:@"wide_u8"] : nil; + 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("【红】8 位索引 GPU 展开：计算管线建不出 %s（照旧由 CPU 展开）", e.localizedDescription.UTF8String); return; } } + if (g_wide_epoch != g_wr_epoch || g_wide_sub != g_sub_serial || g_wide_n > WIDE_MAX - 8) { g_wide_n = 0; g_wide_epoch = g_wr_epoch; g_wide_sub = g_sub_serial; } + int n0 = g_wide_n; wide_take(di); + if (g_exec_next && g_ret_n <= 8) { // Later draws in replay order (current list → callers → the rest of this submission), up to the next buffer write + 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 *q; op_hdr h; const void *a; + for (int guard = 0; guard < 200000 && g_wide_n < WIDE_MAX && fz_next(&cur, &q, &h, &a); guard++) { + if (h.op == OP_DISPATCH || h.op == OP_COPY_B2B || h.op == OP_UPDATE_UB || h.op == OP_CLEAR_BUF || h.op == OP_COPY_T2B) break; + if (h.op == OP_DRAW_IDX && wide_wanted(a)) wide_take(a); + } + } + int n = g_wide_n - n0; if (n <= 0) return; + if (g_enc) { g_wide_splits++; PD("M %llu 8 位索引 GPU 展开：在此切开渲染通道", g_stat_pass); rp_end(); } + if (g_blit) { enc_b_end(g_blit); g_blit = nil; } + typedef struct { uint64_t src, dst; uint32_t n, shift; } wjob; + id<MTLBuffer> jb; size_t jo = arena_alloc((size_t)n * sizeof(wjob), &jb); wjob *jp = (wjob *)((uint8_t *)jb.contents + jo); + id<MTL4ComputeCommandEncoder> ce = enc_c_begin(NULL); [ce setArgumentTable:at]; [ce setComputePipelineState:ps]; + for (int i = 0; i < n; i++) { + wide_ent *e = &g_wide[n0 + i]; void *mb; uint64_t mo, src; e->gpu = 0; + if (!mtl_gpu_addr(e->ibuf, (uint64_t)e->count, &mb, &mo, &src)) continue; // the draw reports it (RED_ONCE 22) + id<MTLBuffer> ob; size_t oo = arena_alloc((size_t)e->count * 2, &ob); + jp[i] = (wjob){src & ~15ull, ob.gpuAddress + oo, (uint32_t)e->count, (uint32_t)(src & 15)}; + [at setAddress:jb.gpuAddress + jo + (uint64_t)i * sizeof(wjob) atIndex:0]; + [ce dispatchThreadgroups:MTLSizeMake(((NSUInteger)e->count + 63) / 64, 1, 1) threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; + e->gpu = ob.gpuAddress + oo; + } + enc_c_end(ce); + args_dirty_all(); + g_wide_draws += (uint64_t)n; + if (!g_wide_runs++) MLOG("8 位索引 GPU 展开：首次 %d 个绘制（索引本提交里有 GPU 侧写入，在渲染通道前由计算展开成 16 位）", n); + else if (g_wide_runs % 3600 == 0) MLOG("8 位索引 GPU 展开：%llu 次 %llu 个绘制，其中切开渲染通道 %llu 次", (unsigned long long)g_wide_runs, (unsigned long long)g_wide_draws, (unsigned long long)g_wide_splits); +} static void draw(const p_draw *dw, const p_draw_idx *di) { if (!S.prog) { RED_ONCE(18, "绘制前没有绑定程序"); PD("Z %llu 绘制跳过：没有绑定程序", g_stat_pass + !g_enc); return; } __unsafe_unretained NvProg *pr = prog_get(S.prog); @@ -4932,6 +5094,7 @@ static void draw(const p_draw *dw, const p_draw_idx *di) { g_zskip_count++; { static int once; if (!once++) MLOG("零源跳过：首例 通道 %llu", g_stat_pass + 1); } PD("N %llu 零源跳过", g_stat_pass + 1); return; } 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) + 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; } } @@ -5087,7 +5250,7 @@ frz_plain:; static id<MTLBuffer> zero; if (!zero) { zero = [g_dev newBufferWithLength:64 options:MTLResourceStorageModeShared]; res_add(zero); } static uint64_t zero_gpu; if (!zero_gpu) zero_gpu = zero.gpuAddress; enc_r_vbuf(en, 30, zero_gpu); - if (pr->sh[0]) vconv_bind(en, &pr->sh[0]->rf); // integer inputs fed float-type formats (see vconv_ncomp) + if (pr->sh[0]) vconv_bind(en, &pr->sh[0]->rf); // integer inputs fed float-type formats or 8 / 16-bit integers of the other sign (see vconv_ncomp, vconv_int) g_kd_force = kd_draw(pr, dw, di); if (pr->sh[0]) encode_args(en, pr->sh[0], 0, 1); if (pr->sh[1]) encode_args(en, pr->sh[1], 1, 0); @@ -5119,11 +5282,12 @@ frz_plain:; if (enc_label_on() && !g_el_draws++ && pr->sh[1]) { memcpy(g_el_key, pr->sh[1]->key, 8); g_el_key[8] = 0; } if (passdump()) { // Characteristics of a render pass's first draw (to judge whether the load can be omitted: full screen, no blending, writes all channels, does not depend on depth); with DM_NVN_PASSDUMP_ALL every draw is recorded static uint64_t last_pass = ~0ull; static int all = -1; if (all < 0) all = getenv("DM_NVN_PASSDUMP_ALL") != NULL; - if (all || g_kd_force) PD("d %llu bv=%d bi=%d first=%d vs=%s fs=%s vc=%d inst=%d blend=%#x b0=%#llx dt=%d dw=%d df=%d cull=%d st=%d sf=%d,%d sop=%d,%d,%d/%d,%d,%d smask=%#x,%#x/%#x,%#x sref=%#x,%#x ib=%#llx it=%d cmask=%#x b1=%#llx b2=%#llx b3=%#llx", g_stat_pass, di ? di->base_vertex : 0, di ? di->base_instance : 0, dw ? dw->first : 0, pr->sh[0] ? pr->sh[0]->key : "-", pr->sh[1] ? pr->sh[1]->key : "-", + if (all || g_kd_force) PD("d %llu bv=%d bi=%d first=%d vs=%s fs=%s vc=%d inst=%d blend=%#x b0=%#llx dt=%d dw=%d df=%d cull=%d st=%d sf=%d,%d sop=%d,%d,%d/%d,%d,%d smask=%#x,%#x/%#x,%#x sref=%#x,%#x ib=%#llx it=%d cmask=%#x b1=%#llx b2=%#llx b3=%#llx po=%d/%g,%g,%g dc=%d", g_stat_pass, di ? di->base_vertex : 0, di ? di->base_instance : 0, dw ? dw->first : 0, pr->sh[0] ? pr->sh[0]->key : "-", pr->sh[1] ? pr->sh[1]->key : "-", dw ? dw->count : di->count, dw ? dw->instances : di->instances, S.color.blend_en, (unsigned long long)*(uint64_t *)&S.blend[0], (int)S.ds.dt, (int)S.ds.dw, (int)S.ds.df, (int)S.poly.cull, (int)S.ds.st, (int)S.ds.sf_f, (int)S.ds.sf_b, (int)S.ds.fail_f, (int)S.ds.zfail_f, (int)S.ds.zpass_f, (int)S.ds.fail_b, (int)S.ds.zfail_b, (int)S.ds.zpass_b, S.st_vmask[0], S.st_vmask[1], S.st_wmask[0], S.st_wmask[1], S.st_ref[0], S.st_ref[1], - di ? (unsigned long long)di->ibuf : 0ull, di ? di->itype : 0, S.chmask.mask, (unsigned long long)*(uint64_t *)&S.blend[1], (unsigned long long)*(uint64_t *)&S.blend[2], (unsigned long long)*(uint64_t *)&S.blend[3]); + di ? (unsigned long long)di->ibuf : 0ull, di ? di->itype : 0, S.chmask.mask, (unsigned long long)*(uint64_t *)&S.blend[1], (unsigned long long)*(uint64_t *)&S.blend[2], (unsigned long long)*(uint64_t *)&S.blend[3], + (int)S.poly.offset_en, S.po.factor, S.po.units, S.po.clamp, S.depth_clamp); if (all && pr->sh[0]) { // Content fingerprint of each constant buffer of the vertex stage (to judge whether a draw is completely identical across frames: used by the shadow cache etc.); when unreadable (a write pending execution in this submit) record P char ub[400]; size_t o = 0; NvShader *v = pr->sh[0]; for (int i = 0; i < v->rf.nb && o < sizeof ub - 40; i++) if (v->rf.b[i].kind == NVN_RB_CBUF && v->rf.b[i].reg < 16) { @@ -5211,9 +5375,16 @@ frz_plain:; if (di->itype == 0) { // 8-bit indices: Metal only has 16 / 32-bit. At replay they are read from guest memory and widened to 16-bit in this frame's arena (SMBW's full-screen passes and the final copy into the window all use 8-bit indices; without this the whole screen is black) if (!di->count) return; + const wide_ent *w = wide_wanted(di) ? wide_find(di) : NULL; + if (w && w->gpu) { // written earlier in this submission: widened on the GPU before this pass (wide_run) + use_rec("buf", di->ibuf, (uint64_t)di->count, 0, NULL); + enc_r_draw_indexed(en, mp, di->count, MTLIndexTypeUInt16, w->gpu, (NSUInteger)di->count * 2, di->instances ? di->instances : 1, di->base_vertex, di->base_instance); + if (g_rt_range_pending) { PD("j rt_level_range pass=%llu vc=%u", g_stat_pass, (unsigned)di->count); g_rt_range_pending = 0; } + return; + } void *mb; uint64_t mo; const uint8_t *src; if (!mtl_addr(di->ibuf, di->count, &mb, &mo) || !(src = [(__bridge id<MTLBuffer>)mb contents])) { RED_ONCE(22, "索引缓冲不在内存池里"); return; } - if (pend_hits(di->ibuf, di->count)) RED_ONCE(48, "8 位索引缓冲本提交里有 GPU 侧写入：按当前内容展开"); + if (pend_hits(di->ibuf, di->count)) RED_ONCE(48, "8 位索引缓冲本提交里有 GPU 侧写入，GPU 展开没成：按当前内容展开"); id<MTLBuffer> ab; size_t off = arena_alloc(((size_t)di->count * 2 + 15) & ~(size_t)15, &ab); uint16_t *d = (uint16_t *)((uint8_t *)ab.contents + off); src += mo; for (int32_t i = 0; i < di->count; i++) d[i] = src[i]; @@ -5844,9 +6015,12 @@ static void sub_cut(void) { if (g_blit) { enc_b_end(g_blit); g_blit = nil; } up_flush(); uint64_t f = atomic_load(&g_cur_frame); gpu_tag_begin(f); + uint32_t k = g_seg++; seg_clk *sc = NULL; // reports encoded so far keep this segment's GPU times (seg_ts) + if (seg_ts()) { if (!g_sclk) { g_sclk = calloc(1, sizeof *g_sclk); atomic_store(&g_sclk->refs, 1); } sc = g_sclk; atomic_fetch_add(&sc->refs, 1); } MTL4CommitOptions *o = [MTL4CommitOptions new]; [o addFeedbackHandler:^(id<MTL4CommitFeedback> cb) { gpu_tag_done(f, cb.GPUStartTime, cb.GPUEndTime); + if (sc) { if (k < SEG_CLK_N) { atomic_store(&sc->t[k][0], gpu_s_guest_ns(cb.GPUStartTime)); atomic_store(&sc->t[k][1], gpu_s_guest_ns(cb.GPUEndTime)); } seg_clk_put(sc); } if (cb.error) { atomic_store(&g_last_commit_error, 1); static int k; if (k++ < 8) MLOG("【红】分段命令缓冲执行出错：%s", cb.error.description.UTF8String); } }]; res_commit(); @@ -5859,6 +6033,130 @@ static void sub_cut(void) { if (!g_cuts++) MLOG("分段提交：首段在帧 %llu 通道 %llu 交出", (unsigned long long)f, (unsigned long long)g_stat_pass); } +// ---------- Clears a load action cannot express ---------- +// A colour clear issued inside an open render pass, with a partial channel mask or with a scissor smaller than the target, and a depth / stencil clear with a scissor inside an +// open pass (Tears of the Kingdom issues all of them) are drawn as one triangle over the whole pass in the render encoder. It writes the clear value into the masked channels of +// that attachment, or the depth / the stencil bits of the clear's write mask, inside the scissor, with the depth test ALWAYS and no visibility counting. The value is read from +// the frame arena through the triangle's own argument tables (the draws' g_vt / g_ft are put back after it), so the pipelines depend only on the pass formats, the slot and the +// mask. Clears the load action expresses keep it. Adapted from MacNativeGames' clear_quad / clear_ds_quad (github.com/ly05010419/MacNativeGames, MIT). +static int pf_int_kind(MTLPixelFormat f) { // the fragment output type a colour attachment takes: 0 float (also normalized), 1 uint, 2 int + switch (f) { + case MTLPixelFormatR8Uint: case MTLPixelFormatR16Uint: case MTLPixelFormatR32Uint: case MTLPixelFormatRG8Uint: case MTLPixelFormatRG16Uint: case MTLPixelFormatRG32Uint: + case MTLPixelFormatRGBA8Uint: case MTLPixelFormatRGB10A2Uint: case MTLPixelFormatRGBA16Uint: case MTLPixelFormatRGBA32Uint: return 1; + case MTLPixelFormatR8Sint: case MTLPixelFormatR16Sint: case MTLPixelFormatR32Sint: case MTLPixelFormatRG8Sint: case MTLPixelFormatRG16Sint: case MTLPixelFormatRG32Sint: + case MTLPixelFormatRGBA8Sint: case MTLPixelFormatRGBA16Sint: case MTLPixelFormatRGBA32Sint: return 2; + default: return 0; + } +} +static id<MTLLibrary> clr_lib(int slot, int kind) { // vertex clr_vs (the triangle, depth from buffer 0) and fragment clr_fs (colour slot `slot` of type `kind` from buffer 0) + static id<MTLLibrary> libs[8][3]; static uint8_t failed[8][3]; + if (libs[slot][kind] || failed[slot][kind]) return libs[slot][kind]; + static const char *ty[3] = {"float", "uint", "int"}; char src[720]; + snprintf(src, sizeof src, "#include <metal_stdlib>\nusing namespace metal;\nstruct V { float4 p [[position]]; };\n" + "vertex V clr_vs(uint i [[vertex_id]], device const float *z [[buffer(0)]]) { V o; o.p = float4(float2(float((i << 1) & 2), float(i & 2)) * 2.0 - 1.0, z[0], 1.0); return o; }\n" + "struct O { %s4 c [[color(%d)]]; };\nfragment O clr_fs(device const %s4 *v [[buffer(0)]]) { O o; o.c = v[0]; return o; }\n", ty[kind], slot, ty[kind]); + NSError *e = nil; libs[slot][kind] = [g_dev newLibraryWithSource:@(src) options:nil error:&e]; + if (!libs[slot][kind]) { failed[slot][kind] = 1; MLOG("【红】清除三角形着色器编译失败（槽 %d 类型 %d）：%s", slot, kind, e.localizedDescription.UTF8String); } + return libs[slot][kind]; +} +typedef struct { int32_t slot, kind, mask; uint32_t cfmt[8], dfmt, sfmt, samples; } clr_key; // slot -1: a depth / stencil clear (no fragment function) +static id<MTLRenderPipelineState> clr_pso(const clr_key *k) { + static NSMutableDictionary<NSData *, id<MTLRenderPipelineState>> *cache; if (!cache) cache = [NSMutableDictionary dictionary]; + NSData *kd = [NSData dataWithBytes:k length:sizeof *k]; id<MTLRenderPipelineState> ps = cache[kd]; if (ps) return ps; + id<MTLLibrary> lib = clr_lib(k->slot < 0 ? 0 : k->slot, k->slot < 0 ? 0 : k->kind); if (!lib) return nil; + MTLRenderPipelineDescriptor *d = [MTLRenderPipelineDescriptor new]; + d.vertexFunction = [lib newFunctionWithName:@"clr_vs"]; if (k->slot >= 0) d.fragmentFunction = [lib newFunctionWithName:@"clr_fs"]; + for (int i = 0; i < 8; i++) if (k->cfmt[i]) { d.colorAttachments[i].pixelFormat = (MTLPixelFormat)k->cfmt[i]; d.colorAttachments[i].writeMask = MTLColorWriteMaskNone; } // the other attachments stay as they are + if (k->slot >= 0) d.colorAttachments[k->slot].writeMask = (MTLColorWriteMask)((k->mask & 1 ? MTLColorWriteMaskRed : 0) | (k->mask & 2 ? MTLColorWriteMaskGreen : 0) | (k->mask & 4 ? MTLColorWriteMaskBlue : 0) | (k->mask & 8 ? MTLColorWriteMaskAlpha : 0)); + d.depthAttachmentPixelFormat = (MTLPixelFormat)k->dfmt; d.stencilAttachmentPixelFormat = (MTLPixelFormat)k->sfmt; d.rasterSampleCount = k->samples ? k->samples : 1; + NSError *e = nil; ps = [g_dev newRenderPipelineStateWithDescriptor:d error:&e]; + if (!ps) { static int n; if (n++ < 4) MLOG("【红】清除三角形管线建不出（槽 %d 颜色格式 %u 深度 %u 模板 %u 采样 %u）：%s，这次清除跳过", k->slot, k->slot >= 0 ? k->cfmt[k->slot] : 0, k->dfmt, k->sfmt, k->samples, e.localizedDescription.UTF8String); return nil; } + cache[kd] = ps; return ps; +} +static id<MTLDepthStencilState> clr_dss(int dw, unsigned sm) { // depth ALWAYS, written when dw; stencil ALWAYS → REPLACE through write mask sm (no stencil test when 0) + static id<MTLDepthStencilState> c[2][256]; + if (c[dw][sm]) return c[dw][sm]; + MTLDepthStencilDescriptor *d = [MTLDepthStencilDescriptor new]; d.depthCompareFunction = MTLCompareFunctionAlways; d.depthWriteEnabled = dw != 0; + if (sm) { MTLStencilDescriptor *st = [MTLStencilDescriptor new]; st.stencilCompareFunction = MTLCompareFunctionAlways; st.stencilFailureOperation = st.depthFailureOperation = MTLStencilOperationKeep; + st.depthStencilPassOperation = MTLStencilOperationReplace; st.readMask = 0xff; st.writeMask = sm; d.frontFaceStencil = d.backFaceStencil = st; } + return c[dw][sm] = [g_dev newDepthStencilStateWithDescriptor:d]; +} +static int clr_scissor(MTLScissorRect *sc) { // the clear's scissor in the open pass, clamped as for draws; 0 = empty + NSUInteger tw = g_pass.tw, th = g_pass.th; *sc = (MTLScissorRect){0, 0, tw, th}; + if (!R.has_sc) return tw && th; + long x0 = R.sc.x < 0 ? 0 : R.sc.x, y0 = R.sc.y < 0 ? 0 : R.sc.y, x1 = (long)R.sc.x + R.sc.w, y1 = (long)R.sc.y + R.sc.h; + if (x1 > (long)tw) x1 = (long)tw; if (y1 > (long)th) y1 = (long)th; + if (x1 <= x0 || y1 <= y0) return 0; + *sc = (MTLScissorRect){(NSUInteger)x0, (NSUInteger)y0, (NSUInteger)(x1 - x0), (NSUInteger)(y1 - y0)}; return 1; +} +static void clr_tri(id<MTL4RenderCommandEncoder> en, id<MTLRenderPipelineState> ps, id<MTLDepthStencilState> ds, MTLScissorRect sc, const uint32_t val[4], float z, uint32_t sref) { + static id<MTL4ArgumentTable> vt, ft; + if (!ft) { NSError *e = nil; MTL4ArgumentTableDescriptor *ad = [MTL4ArgumentTableDescriptor new]; ad.maxBufferBindCount = 1; + vt = [g_dev newArgumentTableWithDescriptor:ad error:&e]; ft = vt ? [g_dev newArgumentTableWithDescriptor:ad error:&e] : nil; + if (!ft) { vt = nil; static int n; if (!n++) MLOG("【红】清除三角形参数表建不出：%s，这次清除跳过", e.localizedDescription.UTF8String); return; } } + id<MTLBuffer> b; size_t o = arena_alloc(32, &b); uint8_t *p = (uint8_t *)b.contents + o; memcpy(p, val, 16); memcpy(p + 16, &z, 4); // colour at 0, depth at 16 + [ft setAddress:b.gpuAddress + o atIndex:0]; [vt setAddress:b.gpuAddress + o + 16 atIndex:0]; + [en setArgumentTable:vt atStages:MTLRenderStageVertex]; [en setArgumentTable:ft atStages:MTLRenderStageFragment]; + enc_r_pso(en, ps); g_enc_pso = nil; enc_r_dss(en, ds); g_enc_dss = nil; // the next draw sets its own again + if (g_enc_vis) { enc_r_visibility(en, MTLVisibilityResultModeDisabled, 0); g_enc_vis = 0; } // not counted by an occlusion query; the next counted draw turns counting back on + enc_r_viewport(en, (MTLViewport){0, 0, (double)g_pass.tw, (double)g_pass.th, 0, 1}); enc_r_scissor(en, sc); + enc_r_cull(en, MTLCullModeNone); enc_r_fill(en, MTLTriangleFillModeFill); enc_r_depth_bias(en, 0, 0, 0); enc_r_depth_clip(en, YES); enc_r_stencil_ref(en, sref, sref); + enc_r_draw(en, MTLPrimitiveTypeTriangle, 0, 3); + [en setArgumentTable:g_vt atStages:MTLRenderStageVertex]; [en setArgumentTable:g_ft atStages:MTLRenderStageFragment]; +} +static void clr_note(int slot, int depth, int stencil) { // as udc_note_draw for a draw: the attachment was written (zero-source skip, unwritten-attachment store) + if (slot >= 0 && R.rt.color[slot]) { zskip_del(R.rt.color[slot]->serial); g_zskip_cand_ok[slot] = 0; } + for (int i = 0; i < g_nsd; i++) { sd_att *t = &g_sd[i]; + if ((t->depth == 0 && slot >= 0 && t->idx == slot) || (t->depth == 1 && depth) || (t->depth == 2 && stencil)) { t->wrote = 1; t->wwhy = 2; } } +} +static uint64_t g_clr_tri[2]; // clears drawn as a triangle: colour, depth / stencil +static void clear_color_tri(const p_clear_color *c) { + int slot = c->index, mask = c->mask & 0xf; + if (!mask || slot < 0 || slot >= 8 || slot >= R.rt.n || !R.rt.color[slot]) return; + id<MTL4RenderCommandEncoder> en = rp_begin(); // a clear outside a pass opens it as a draw would (pending clears land first) + if (en && !g_pass.cfmt[slot] && g_enc_mode) { rp_end(); en = rp_begin(); } // opened without this attachment (clear-only / TBDR modes): reopen with all of them + if (!en) return; + int at = slot; + if (!g_pass.cfmt[at] && rt_dup_slot(&R.rt, slot)) // a slot attached twice is attached only once (rt_dup_slot): clear the first one + for (int j = 0; j < slot; j++) { const tex_t *a = R.rt.color[j], *b = R.rt.color[slot]; if (!a || a->serial != b->serial || !g_pass.cfmt[j]) continue; + const view_t *va = &R.rt.cview[j], *vb = &R.rt.cview[slot]; int ha = R.rt.has_cview[j], hb = R.rt.has_cview[slot]; + if ((ha && va->set_levels ? va->base_level : 0) == (hb && vb->set_levels ? vb->base_level : 0) && (ha && va->set_layers ? va->min_layer : 0) == (hb && vb->set_layers ? vb->min_layer : 0)) { at = j; break; } } + if (!g_pass.cfmt[at]) { PD("c %llu 清色三角形跳过：槽 %d 不在本通道", g_stat_pass, slot); return; } + MTLScissorRect sc; if (!clr_scissor(&sc)) { PD("c %llu 清色三角形跳过：剪裁为空 槽 %d", g_stat_pass, slot); return; } + int kind = pf_int_kind((MTLPixelFormat)g_pass.cfmt[at]); + clr_key k; memset(&k, 0, sizeof k); k.slot = at; k.kind = kind; k.mask = mask; memcpy(k.cfmt, g_pass.cfmt, sizeof k.cfmt); k.dfmt = g_pass.dfmt; k.sfmt = g_pass.sfmt; k.samples = g_pass.samples; + id<MTLRenderPipelineState> ps = clr_pso(&k); id<MTLDepthStencilState> ds = clr_dss(0, 0); if (!ps || !ds) return; + uint32_t v[4]; + for (int i = 0; i < 4; i++) { float f = c->rgba[i]; int32_t si; // integer attachments take the value converted, as a load-action clear does + if (kind == 1) v[i] = f > 0 ? (f < 4294967040.f ? (uint32_t)f : 0xffffffffu) : 0; + else if (kind == 2) { si = !(f == f) ? 0 : f >= 2147483520.f ? INT32_MAX : f <= -2147483648.f ? INT32_MIN : (int32_t)f; memcpy(&v[i], &si, 4); } + else memcpy(&v[i], &f, 4); } + clr_tri(en, ps, ds, sc, v, 0, 0); clr_note(at, 0, 0); + if (!g_clr_tri[0]++) MLOG("通道中途 / 部分掩码 / 部分剪裁的清色：画成整通道三角形（首例 通道 %llu 槽 %d 掩码 %#x 剪裁 %lu,%lu %lux%lu）", g_stat_pass, at, mask, + (unsigned long)sc.x, (unsigned long)sc.y, (unsigned long)sc.width, (unsigned long)sc.height); + PD("c %llu 清色画成三角形 槽%d ser=%llu 掩码=%#x 剪裁=%lu,%lu,%lu,%lu 色=%g,%g,%g,%g", g_stat_pass, at, (unsigned long long)R.rt.color[at]->serial, mask, + (unsigned long)sc.x, (unsigned long)sc.y, (unsigned long)sc.width, (unsigned long)sc.height, c->rgba[0], c->rgba[1], c->rgba[2], c->rgba[3]); +} +static void clear_ds_tri(const p_clear_ds *c) { + unsigned want_s = (unsigned)c->stencil_mask & 0xff; + if (!c->depth_mask && !want_s) return; + id<MTL4RenderCommandEncoder> en = rp_begin(); + if (en && g_enc_mode && ((c->depth_mask && !g_pass.hd) || (want_s && !g_pass.hs))) { rp_end(); en = rp_begin(); } // opened without the depth or stencil plane (TBDR lookahead): reopen with all attachments + if (!en) return; + int dw = c->depth_mask && g_pass.hd; unsigned sm = g_pass.hs ? want_s : 0; + if (!dw && !sm) return; + MTLScissorRect sc; if (!clr_scissor(&sc)) { PD("c %llu 深度模板清除三角形跳过：剪裁为空", g_stat_pass); return; } + clr_key k; memset(&k, 0, sizeof k); k.slot = -1; memcpy(k.cfmt, g_pass.cfmt, sizeof k.cfmt); k.dfmt = g_pass.dfmt; k.sfmt = g_pass.sfmt; k.samples = g_pass.samples; + id<MTLRenderPipelineState> ps = clr_pso(&k); id<MTLDepthStencilState> ds = clr_dss(dw, sm); if (!ps || !ds) return; + float z = c->depth > 0 ? (c->depth < 1 ? c->depth : 1) : 0; // as clearDepth: within [0, 1] (NaN → 0) + const uint32_t v[4] = {0}; + clr_tri(en, ps, ds, sc, v, z, (uint32_t)c->stencil & 0xff); clr_note(-1, dw, sm != 0); + if (!g_clr_tri[1]++) MLOG("通道中途的剪裁深度模板清除：画成整通道三角形（首例 通道 %llu 深度 %d 模板掩码 %#x 剪裁 %lu,%lu %lux%lu）", g_stat_pass, dw, sm, + (unsigned long)sc.x, (unsigned long)sc.y, (unsigned long)sc.width, (unsigned long)sc.height); + PD("c %llu 深度模板清除画成三角形 ser=%llu 深度=%g(%d) 模板=%d 掩码=%#x 剪裁=%lu,%lu,%lu,%lu", g_stat_pass, R.rt.depth ? (unsigned long long)R.rt.depth->serial : 0ull, z, dw, c->stencil, sm, + (unsigned long)sc.x, (unsigned long)sc.y, (unsigned long)sc.width, (unsigned long)sc.height); +} + void mtl_exec(const uint8_t *p) { static int inited; if (!inited) { inited = 1; state_reset(); } static int dump_all = -1; if (dump_all < 0) dump_all = getenv("DM_NVN_PASSDUMP_ALL") != NULL; @@ -5892,9 +6190,16 @@ void mtl_exec(const uint8_t *p) { PD("Y signal_event cmd=%p dst=%p mode=%u value=%u", p, val, m, v); mtl_sub_event(val, m, v); break; } case OP_REPORT_COUNTER: { uint64_t t, ad; memcpy(&t, a, 8); memcpy(&ad, (const uint8_t *)a + 8, 8); g_cnt_report[t & 31]++; PD("Q report type=%llu cmd=%p dst=%#llx", (unsigned long long)t, p, (unsigned long long)ad); + if (DM_GAME_RENDER_COND) rc_note(ad, t == 0 || t == 0x10 ? 1 : 2); if (t == NVN_COUNTER_TYPE_SAMPLES_PASSED) { post((post_op){3, NULL, 0, 0, ad, (uint32_t)t, vis_snapshot()}); g_vis_cur = -1; } else mtl_sub_counter(ad, (uint32_t)t); break; } + case OP_RENDER_COND: { const p_render_cond *rc = a; // recorded only with the profile key render_cond (hle_nvn.c) + if (rc->mode >= 2) { g_rc_skip = rc->mode == 3; PD("E %llu 渲染开关 %d", g_stat_pass, rc->mode == 2); } + else { uint64_t x = 0, y = 0; int known = rc_value(rc->addr, &x) & rc_value(rc->addr + 16, &y); + g_rc_skip = known && (rc->mode ? x == y : x != y); + PD("E %llu 条件渲染 %s addr=%#llx 值 %llu / %llu%s%s", g_stat_pass, rc->mode ? "不等才画" : "相等才画", (unsigned long long)rc->addr, (unsigned long long)x, (unsigned long long)y, known ? "" : "（未知，照画）", g_rc_skip ? " → 不画" : ""); } + break; } case OP_RESET_COUNTER: { uint64_t t; memcpy(&t, a, 8); g_cnt_reset[t & 31]++; PD("Q reset type=%llu cmd=%p", (unsigned long long)t, p); if (t == NVN_COUNTER_TYPE_SAMPLES_PASSED) { #ifdef NVN_RT_RANGE_TEST @@ -5922,37 +6227,43 @@ void mtl_exec(const uint8_t *p) { } case OP_CLEAR_COLOR: { const p_clear_color *c = a; + if (g_rc_skip) { PD("c %llu 清色因条件渲染跳过 槽%d", g_stat_pass, c->index); break; } // A scissored clear to 0 outside a pass with all channels (SMBW clears pyramid levels as sub-rectangles of 480×272 and other targets): copy from the zero buffer into the scissor rectangle in the replay stream if (!g_enc && c->mask == 0xf && !c->rgba[0] && !c->rgba[1] && !c->rgba[2] && !c->rgba[3] && R.has_sc && c->index >= 0 && c->index < R.rt.n && R.rt.color[c->index] && !scissor_full(R.rt.color[c->index], &R.rt.cview[c->index], R.rt.has_cview[c->index]) && clear_rect_zero(c->index)) break; if (g_rt_merged) rp_end(); // A clear arrived in a continued render pass: end it here, and the clear is attached to the start of the next render pass as usual if (c->index < 0 || c->index >= R.rt.n || !R.rt.color[c->index]) { RED_ONCE(5, "清色目标 %d 无效", c->index); break; } if (g_enc || c->mask != 0xf || !scissor_full(R.rt.color[c->index], &R.rt.cview[c->index], R.rt.has_cview[c->index])) { - RED_ONCE(6, "通道中途 / 部分掩码 / 部分剪裁的清色未实现（首例 掩码 %#x）", c->mask); - static uint64_t n; if (getenv("DM_NVN_CLEARLOG") && n++ < 40) MLOG("清色跳过 帧%llu 通道%llu 目标%d ser=%llu %ux%u 中途=%d 掩码=%#x 剪裁=%d,%d,%d,%d 色=%g,%g,%g,%g", (unsigned long long)atomic_load(&g_cur_frame), g_stat_pass, c->index, + // Inside a pass, with a partial mask or a partial scissor: drawn as a triangle (clear_color_tri) + static uint64_t n; if (getenv("DM_NVN_CLEARLOG") && n++ < 40) MLOG("清色画成三角形 帧%llu 通道%llu 目标%d ser=%llu %ux%u 中途=%d 掩码=%#x 剪裁=%d,%d,%d,%d 色=%g,%g,%g,%g", (unsigned long long)atomic_load(&g_cur_frame), g_stat_pass, c->index, (unsigned long long)R.rt.color[c->index]->serial, R.rt.color[c->index]->d.width, R.rt.color[c->index]->d.height, g_enc != nil, c->mask, R.sc.x, R.sc.y, R.sc.w, R.sc.h, c->rgba[0], c->rgba[1], c->rgba[2], c->rgba[3]); - break; } + clear_color_tri(c); break; } R.clear_c[c->index] = 1; R.cc[c->index] = MTLClearColorMake(c->rgba[0], c->rgba[1], c->rgba[2], c->rgba[3]); PD("c %llu 清色记下 槽%d ser=%llu 掩码=%#x 剪裁=%d:%d,%d,%d,%d", g_stat_pass, c->index, (unsigned long long)R.rt.color[c->index]->serial, c->mask, R.has_sc, R.sc.x, R.sc.y, R.sc.w, R.sc.h); break; } case OP_CLEAR_DS: { const p_clear_ds *c = a; + if (g_rc_skip) { PD("c %llu 深度模板清除因条件渲染跳过", g_stat_pass); break; } if (g_rt_merged) rp_end(); // Depth / stencil clear in the middle of a pass: end the current pass and attach the clear to the start of the next one (color attachments load as usual, as a continued pass does above when a clear arrives). It used to be skipped, and later draws depth-tested against stale depth. - // When the scissor does not cover the depth target, Metal's load-action clear cannot express a partial clear; still skipped + // When the scissor does not cover the depth target, Metal's load-action clear cannot express a partial clear: drawn as a triangle inside the pass (clear_ds_tri) + int mid = g_enc != nil; if (g_enc) { - if (R.rt.depth && !scissor_full(R.rt.depth, &R.rt.dview, R.rt.has_dview)) { RED_ONCE(7, "通道中途的剪裁深度模板清除未实现"); break; } + if (R.rt.depth && !scissor_full(R.rt.depth, &R.rt.dview, R.rt.has_dview)) { clear_ds_tri(c); break; } { static int once; if (!once++) MLOG("通道中途的深度模板清除：结束本通道，清除挂到下一个通道开头（首例 通道 %llu）", g_stat_pass); } rp_end(); } + if (R.rt.depth && (!scissor_full(R.rt.depth, &R.rt.dview, R.rt.has_dview) || ((c->stencil_mask & 0xff) && (c->stencil_mask & 0xff) != 0xff))) { + static int once; if (!once++) MLOG("深度模板清除带剪裁或部分模板掩码，按整张清（首例 通道 %llu 中途 %d 深度掩码 %d 模板掩码 %#x 剪裁 %d,%d %dx%d 目标 %ux%u）", g_stat_pass, mid, c->depth_mask, c->stencil_mask, + R.sc.x, R.sc.y, R.sc.w, R.sc.h, R.rt.depth->d.width, R.rt.depth->d.height); } if (c->depth_mask) { R.clear_d = 1; R.cd = c->depth; } if (c->stencil_mask) { R.clear_s = 1; R.cs = (uint32_t)c->stencil; } PD("c %llu 深度模板清除记下 ser=%llu 深度=%g 深度掩码=%d 模板=%d 模板掩码=%#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: g_exec_next = p + h.size; draw(a, NULL); g_exec_next = NULL; break; // Lookahead start point - case OP_DRAW_IDX: g_exec_next = p + h.size; draw(NULL, a); g_exec_next = NULL; break; + case OP_DRAW: if (g_rc_skip) break; g_exec_next = p + h.size; draw(a, NULL); g_exec_next = NULL; break; // Lookahead start point + case OP_DRAW_IDX: if (g_rc_skip) break; g_exec_next = p + h.size; draw(NULL, a); 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: diff --git a/dm/gfx/nvn/nvn_obj.h b/dm/gfx/nvn/nvn_obj.h index 0017b7b..070b36f 100644 --- a/dm/gfx/nvn/nvn_obj.h +++ b/dm/gfx/nvn/nvn_obj.h @@ -100,11 +100,12 @@ enum { OP_BIND_PROGRAM, OP_BIND_STATE, OP_BIND_VATTR, OP_BIND_VSTREAM, OP_BIND_VBUF, OP_BIND_UBO, OP_BIND_SSBO, OP_BIND_TEX, OP_BIND_SAMP, OP_STENCIL, OP_POLY_OFFSET, OP_DRAW, OP_DRAW_IDX, OP_DISPATCH, OP_BARRIER, - OP_SET_TEX_POOL, OP_SET_SAMP_POOL, OP_COPY_B2T, OP_COPY_STRIDE, OP_BLEND_COLOR, OP_COPY_T2B, OP_BIND_IMG, OP_COPY_T2T, OP_RESET_COUNTER, OP_CLEAR_TEX, OP_CLEAR_BUF, OP_DEPTH_CLAMP, OP_COUNT_ + OP_SET_TEX_POOL, OP_SET_SAMP_POOL, OP_COPY_B2T, OP_COPY_STRIDE, OP_BLEND_COLOR, OP_COPY_T2B, OP_BIND_IMG, OP_COPY_T2T, OP_RESET_COUNTER, OP_CLEAR_TEX, OP_CLEAR_BUF, OP_DEPTH_CLAMP, OP_RENDER_COND, OP_COUNT_ }; // Payload (immediately follows op_hdr) typedef struct { int32_t n; const tex_t *color[8]; view_t cview[8]; uint8_t has_cview[8]; const tex_t *depth; view_t dview; uint8_t has_dview; } p_set_rt; typedef struct { int32_t index, mask; float rgba[4]; } p_clear_color; +typedef struct { int32_t mode, pad; uint64_t addr; } p_render_cond; // mode 0 / 1: SetRenderEnableConditional (render if the 64-bit counters at addr and addr + 16 are equal / not equal), 2 / 3: SetRenderEnable(true / false) typedef struct { float depth; int32_t depth_mask, stencil, stencil_mask; } p_clear_ds; typedef struct { int32_t first, n; float v[16][6]; } p_viewports; // x, y, w, h, near, far (near and far are overridden by DEPTH_RANGE) typedef struct { int32_t x, y, w, h; } p_scissor; @@ -166,6 +167,7 @@ int mtl_link_wait_tick(uint64_t timeout_ns); double mtl_presented_time(uint64_t n); // Actual presentation time (seconds) of the n-th present (counted from 1); 0 = not ready, −1 = dropped // Texture release (at Finalize) +void mtl_tex_init(const tex_t *t); // TextureInitialize: the serial belongs to a live texture object until mtl_tex_release void mtl_tex_release(tex_t *t); void mtl_pkg_stats(uint64_t *uploads, uint64_t *reuploads); uint64_t mtl_pkg_miss(void); diff --git a/dm/gfx/tests/t_dirty_slots.c b/dm/gfx/tests/t_dirty_slots.c index 836fbe8..8461302 100644 --- a/dm/gfx/tests/t_dirty_slots.c +++ b/dm/gfx/tests/t_dirty_slots.c @@ -58,6 +58,20 @@ int main(void) { munmap(live[i], 4096); } + // More pools alive at once than the table's fixed 4,096 slots used to hold (Breath of the Wild reaches 4,006 in its first shrine): the table grows, and the last pool is still tracked + enum { NMANY = 5000 }; + void **many = malloc(NMANY * sizeof *many); + for (int i = 0; i < NMANY; i++) { many[i] = page_map(); nvn_dirty_add_range(many[i], 4096); } + void *last = many[NMANY - 1]; + EXPECT("同时存活 5000 个池：最后一个在跟踪", 1, nvn_dirty_tracked(last, 4096)); + nvn_dirty_protect(last, 4096); + stamp = nvn_dirty_now(); + EXPECT("第 5000 个池保护后没人写：不脏", 0, nvn_dirty_since(last, 4096, stamp)); + ((volatile uint8_t *)last)[0] = 5; + EXPECT("第 5000 个池被游戏写：脏", 1, nvn_dirty_since(last, 4096, stamp)); + for (int i = 0; i < NMANY; i++) { nvn_dirty_remove_range(many[i]); munmap(many[i], 4096); } + free(many); + // Adjacent pools sharing a host page (pools are 4 KB aligned, host pages 16 KB): the first page of SMBW's ASTC texture pool is shared with a small preceding pool, and was judged untracked and re-uploaded on every use size_t pg = (size_t)getpagesize(); uint8_t *blk = mmap(NULL, 4 * pg, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANON, -1, 0); diff --git a/dm/third_party/README.md b/dm/third_party/README.md index 1ba5cf6..44e7dee 100644 --- a/dm/third_party/README.md +++ b/dm/third_party/README.md @@ -15,4 +15,4 @@ A few more notes: - `flip/tool/` also carries three single-header libraries used by the FLIP command-line tool: `stb_image.h` (2.23), `stb_image_write.h` (1.13) and `tinyexr.h`. Their license notices are in the files themselves. - `libdeflate/VERSION.md`, `libvpx/README.md`, `stb_vorbis/README.md`, `ryujinx/README.md` and `vulkan-headers/README.md` record where each came from and what it is used for. -- Two pieces of third-party code live outside this directory: `dm/gfx/nvn/include/nvn.h` (MIT, author named in the file header) and the layout size algorithm in `dm/gfx/nvn/nvn_tex.c`, ported from deko3d (zlib license). +- Three pieces of third-party code live outside this directory: `dm/gfx/nvn/include/nvn.h` (MIT, author named in the file header), the layout size algorithm in `dm/gfx/nvn/nvn_tex.c`, ported from deko3d (zlib license), and in `dm/gfx/nvn/nvn_mtl.m` the GPU widening of 8-bit indices (`wide_*`) and the in-pass clear triangles (`clear_color_tri` / `clear_ds_tri`), adapted from MacNativeGames' `widen_u8_gpu` and `clear_quad` / `clear_ds_quad` (MIT, license text in `macnativegames/LICENSE`). diff --git a/dm/third_party/macnativegames/LICENSE b/dm/third_party/macnativegames/LICENSE new file mode 100644 index 0000000..a94417a --- /dev/null +++ b/dm/third_party/macnativegames/LICENSE @@ -0,0 +1,21 @@ +MIT License + +Copyright (c) 2026 MacNativeGames contributors + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. diff --git a/dm/tools/patch_nso.py b/dm/tools/patch_nso.py index 149cd96..8acfa2c 100644 --- a/dm/tools/patch_nso.py +++ b/dm/tools/patch_nso.py @@ -146,13 +146,18 @@ CUTS = { # The game logic time step follows it: ksys::VFR is initialized with a base interval of 2 (constant 2 at main+0x127e740 → VFR::init main+0x151d2d0), and every frame VFR::updateInterval (main+0x151d3d0, # called at main+0xcf9ad4 / main+0x1281044) takes the framework's current interval (vtable +0x138); the frame ratio = current interval / 2 -- at interval 1 each frame advances half a 30 fps tick and the game speed stays the same; # when it struggles and steps back to 2 / 3 the ratio is 1 / 1.5, again at the same speed. So only the floor changes from 2 to 1: the lower bound of the candidate interval (two places), the step-down target recorded when stepping up, max(candidate, 3) - 1 → max(candidate, 2) - 1, - # and stepping down stops at < 3 → < 2. On hardware the same logic starting from 1 would also be 60 fps (falling back to 30 by itself when the GPU cannot keep up) + # and stepping down stops at < 3 → < 2. On hardware the same logic starting from 1 would also be 60 fps; it falls back to 30 by itself only when its input shows a slow frame (the last two cuts make that input include the GPU time) 'cd57b23fa4bbad65803d9788c01821ee': [ (0x112670c, '6041201e', '4041201e', '动态间隔候选下限：fmov s0,s11（2.0）→ fmov s0,s10（1.0）', '仅宿主'), (0x1126710, '7f0b0071', '7f070071', '动态间隔候选下限：cmp w27,#2 → cmp w27,#1', '仅宿主'), (0x112665c, '78008052', '58008052', '升档时记的降档目标 max(候选, w24) − 1：mov w24,#3 → mov w24,#2', '仅宿主'), (0x1126748, '1f0d0071', '1f090071', '升档时记的降档目标：cmp w8,#3 → cmp w8,#2', '仅宿主'), (0x1126794, '3f0d0071', '3f090071', '降档到 < 3 停 → 到 < 2 停：cmp w9,#3 → cmp w9,#2', '仅宿主'), + # The controller's input (framework +0x250, controller main+0x1126cec..) is the main thread's CPU span up to the draw step (about 2.7 ms in the first shrine), so GPU-bound frames stayed at interval 1 and the game ran slow + # (iPhone 16 Pro Max: 25 ms frames, two thirds of the speed). The game has a GPU-time path that uses max(CPU span, GPU time between the frame's two timestamp reports, framework +0x220, read after the fence wait) + # but takes it only while [main+0x1e3a638]+0x1d2b8 is set (64 of 6,046 shrine frames); both of its checks become branches into it. The GPU time needs the profile's gpu_timestamps segment + (0x112691c, 'c00000b4', '03000014', '间隔控制器恒走 GPU 时间路径（CPU 段计时）：cbz x0,#0x1126934 → b #0x1126928', '仅宿主'), + (0x1126974, 'c00000b4', '03000014', '间隔控制器恒走 GPU 时间路径（取 max(CPU 段, GPU 时间)）：cbz x0,#0x112698c → b #0x1126980', '仅宿主'), ], '56ec24747ec9e7f19b3efb062483a19d': [ (0x346430, '08fca752', '087ca052', 'AddressRegionManager 构造：节点区上限 mov w8,#0x3fe00000 → mov w8,#0x3e00000', '仅宿主'), diff --git a/docs/agents/03-bringup.md b/docs/agents/03-bringup.md index 1130d61..a1b58b5 100644 --- a/docs/agents/03-bringup.md +++ b/docs/agents/03-bringup.md @@ -151,7 +151,7 @@ Each entry: symptom → root cause → fix (where the fix lives). Check these fi ### Frame rate -- **Running a 30 fps game at 60** (Breath of the Wild): find where the present interval comes from and how the logic time step follows it before patching anything. BotW's `sead::GameFrameworkNx` main loop (`main+0x112660c`) runs a dynamic interval: from the last frame's time it computes how many vertical syncs were needed, clamps that to [2, 3] (the maximum from the interval controller built at `main+0x1126c04`), steps up after three overrunning frames and down one level every six frames, and sets the result through `setWaitVsyncInterval` (vtable +0x130, `main+0x11269b8` → `nvnWindowSetPresentInterval`). The game logic follows on its own: `ksys::VFR` is initialized with a constant base interval of 2 (`main+0x127e740`) and every frame `VFR::updateInterval` (`main+0x151d3d0`) takes the framework's current interval, giving a frame ratio of interval / 2. Five rewriter cuts lower the floor from 2 to 1 (`patch_nso.py`, build `cd57b23f…`): the candidate floor (two places) and the hysteresis targets. Probes confirmed the result: the framework presents at interval 1, VFR reports interval 1 with ratio 0.5, and when a frame runs long the game steps back to 2 with ratio 1, so the game speed does not change. The zeldaret/botw decompilation (`VFR.cpp`) gave the semantics; 1.9.0's addresses were found by the constants (`60 / interval`, the `fmov s11, #2.0` floor). +- **Running a 30 fps game at 60** (Breath of the Wild): find where the present interval comes from and how the logic time step follows it before patching anything. BotW's `sead::GameFrameworkNx` main loop (`main+0x112660c`) runs a dynamic interval: from the last frame's time it computes how many vertical syncs were needed, clamps that to [2, 3] (the maximum from the interval controller built at `main+0x1126c04`), steps up after three overrunning frames and down one level every six frames, and sets the result through `setWaitVsyncInterval` (vtable +0x130, `main+0x11269b8` → `nvnWindowSetPresentInterval`). The game logic follows on its own: `ksys::VFR` is initialized with a constant base interval of 2 (`main+0x127e740`) and every frame `VFR::updateInterval` (`main+0x151d3d0`) takes the framework's current interval, giving a frame ratio of interval / 2. Five rewriter cuts lower the floor from 2 to 1 (`patch_nso.py`, build `cd57b23f…`): the candidate floor (two places) and the hysteresis targets. Probes confirmed the result: the framework presents at interval 1, VFR reports interval 1 with ratio 0.5, and when a frame runs long the game steps back to 2 with ratio 1, so the game speed does not change. Then check what the controller measures: BotW's measured only the main thread's CPU time (about 3 ms a frame), so a GPU-bound device stayed at interval 1 and played in slow motion (GitHub issue #9). Two more cuts make it always take its GPU-time path, and the host's timestamp reports had to carry real GPU times for that path to see anything (profile key `gpu_timestamps`, [04](04-graphics.md)). The zeldaret/botw decompilation (`VFR.cpp`) gave the semantics; 1.9.0's addresses were found by the constants (`60 / interval`, the `fmov s11, #2.0` floor). ### Audio diff --git a/docs/agents/04-graphics.md b/docs/agents/04-graphics.md index a651d22..c3482e9 100644 --- a/docs/agents/04-graphics.md +++ b/docs/agents/04-graphics.md @@ -70,7 +70,7 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog - **Virtual memory pools** (flag 0x200): aliased onto the physical pool named by `MapVirtual` when all requests share one base (Mario's case); otherwise unimplemented. - **SPIRV-Cross failures**: 2 agl post-effect compute shaders still produce SPIR-V that SPIRV-Cross can't parse (`nullptr` exception). Before the compute start-offset fix it was 7: misdecoded code produced the bad SPIR-V. - **CPU-written linear textures**: a texture with `NVN_TEXTURE_FLAGS_LINEAR` or `LINEAR_RENDER_TARGET` in a CPU-visible pool (Mario's glyph cache; its four 64×1 RGBA16F sky colour ramps, without which the sky is black) is created on the pool's `MTLBuffer` (`tex_linear`), so CPU writes (Mario's glyph cache) are visible to the GPU. Block-linear CPU-written textures still need the packaged-texture re-upload path. -- **Integer vertex inputs**: without PC reflection the translator declares every vertex input `float4`. Skinned models read their bone index (`R8UI` / `RG8UI` / `RGBA8UI`) as raw integer bits (`as_type<int>(in.m_N.x) * 48`); Metal converted the integers to floats, every bone index was garbage, and all characters were stretched into streaks. `dmsh_build.py` (`int_inputs`) now re-translates a vertex shader with `uint` inputs when every use of an input is an integer bit cast (529 of 4,854 in Mario), and the host makes the Metal vertex format agree with the declared type (`vfmt_for_ct`; signedness swaps and 32-bit float→int reinterpretation keep the bits). A remaining mismatch is reported as `顶点属性格式 … 与着色器输入类型 … 对不上`. +- **Integer vertex inputs**: without PC reflection the translator declares every vertex input `float4`. Skinned models read their bone index (`R8UI` / `RG8UI` / `RGBA8UI`) as raw integer bits (`as_type<int>(in.m_N.x) * 48`); Metal converted the integers to floats, every bone index was garbage, and all characters were stretched into streaks. `dmsh_build.py` (`int_inputs`) now re-translates a vertex shader with `uint` inputs when every use of an input is an integer bit cast (529 of 4,854 in Mario), and the host makes the Metal vertex format agree with the declared type (`vfmt_for_ct`; signedness swaps of 32-bit components and 32-bit float→int reinterpretation keep the bits, while 8- and 16-bit integers of the other sign are widened on the CPU, see Three Houses below). A remaining mismatch is reported as `顶点属性格式 … 与着色器输入类型 … 对不上`. - **Vertex formats**: Mario uses `RGB10A2SN` (0x69) → `MTLVertexFormatInt1010102Normalized`. Unmapped formats fall back to a 4-component read and print `顶点属性格式 … 未映射` once. - **Instanced array draws** (`DrawArraysInstanced`) carry instance count and base instance in the `OP_DRAW` record. - **Metal 4 does not order operations inside one encoder.** The host batches consecutive copies, fills and inline uniform updates (`UpdateUniformBuffer`) into one compute encoder. Mario writes the same uniform slot twice in a row each frame (4 bytes of −1, then 16 bytes of wave coefficients); without an intra-encoder barrier the writes occasionally landed in reverse order, the fountain's GPU wave simulation (a 512×1 RGBA16F texture updated by render → copy back → splash pass every frame) ran with −1 as its self coefficient, diverged to ±65504, and its vertex shader pushed the water surface to the top of the screen (tall blue / textured panels that flickered frame to frame). `blit_buf` now inserts `barrierAfterEncoderStages` when an operation reads or writes a guest range an earlier operation in the same encoder wrote; texture copies always get one. `DM_NVN_BLIT_ORDER=0` turns the barriers off to reproduce. Render passes and compute dispatches each have their own encoder with full barriers, so they were not affected. @@ -128,7 +128,8 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog - **Cube textures sampled as cube arrays.** 1,512 shaders sample the environment probes as `texturecube_array` while a single cube texture is bound. Metal can make a cube-array view only of a 2D-array base texture, so NVN cube textures are now created as 6-layer 2D arrays and bound through cube / cube-array views by `tex_fit` (which also maps a cube array read as a cube to its first cube). - **Open: a compressed texture as a render target.** A 32×32 RGTC1 (BC4) texture is written through an `RGBA16UI` view (one texel per 8-byte block: BC4 encoded on the GPU). Metal can neither attach a compressed format nor make an uncompressed view of one, and the pipeline-descriptor validation aborts the process, so the attachment is left out and its draws do nothing (`渲染目标是压缩格式贴图 …`). Doing it properly needs an uncompressed proxy target copied into the BC4 texture through a buffer after the pass. - **Big-stack NVN entry points**: the render thread's stack is small, and Metal's driver compiles internal programs (a fast-clear blit) on the calling thread while the host ends a render pass inside `nvnQueueSubmitCommands`. That overflowed the guest stack (SIGBUS in `AGCDeserializedReply`, address in the guest stack region). All implemented NVN functions except the `nvnCommandBuffer*` recording calls now run on the per-thread 512 KB host stack ([03](03-bringup.md), big-stack entry; `nvn_big_entry` in `nvn_entry.S`). -- **60 fps** is a rewriter cut on the game's frame-interval logic ([03](03-bringup.md)); the engine's variable frame rate (`ksys::VFR`) scales its time step by the present interval itself, so the frame rate the player gets is whatever host and GPU can sustain, with the game speed unchanged. +- **60 fps** is a rewriter cut on the game's frame-interval logic ([03](03-bringup.md)); the engine's variable frame rate (`ksys::VFR`) scales its time step by the present interval, so the speed is right as long as the interval matches the frames. **Fixed: slow motion when the GPU is the bottleneck (GitHub issue #9).** The interval controller (`main+0x112660c`) measured a frame as framework +0x250, the main thread's CPU span from the start of `procFrame` to before its draw step (`main+0x1126cec` → `main+0x1126d20`): about 2.7 ms in the shrine (probe at `main+0x11266a4`, `@x19+0x250`). Only while the object at `[main+0x1e3a638] + 0x1d2b8` exists did it take max(CPU span, GPU time), in 64 of 6,046 shrine frames. So the interval stepped up only on CPU spikes: an iPhone 16 Pro Max at a 25 ms frame p50 never left interval 1 in 4,800 frames, and the game ran at about two thirds of its speed. The GPU time (framework +0x220) is the difference between the frame's two timestamp reports: `TIMESTAMP_TOP` at the start of the frame's command buffer (`main+0xe94584`) and `TIMESTAMP` at its end (`main+0xe94664`), read after the draw step has waited for the frame's fence (`main+0xe949bc`). It read 0–6 ticks, because the host wrote all 5 reports of the frame's one submission with one time in its completion handler. Two more cuts (`main+0x112691c`, `main+0x1126974`: cbz → b) make the controller always take the GPU path. The profile key `gpu_timestamps segment` makes each report take the GPU start (`TIMESTAMP_TOP`) or end (other types) of the `sub_cut` segment it was encoded in (`rep_ticks` in `nvn_mtl.m`: MTL4CommitFeedback times converted to the guest clock, written in the last segment's handler as before). The GPU time now reads the frame's GPU span, gaps where the GPU waits for the host's translation included. Measured on the Mac with two BotW runs sharing the GPU under a heavy CPU load (load average about 75, a shader rebuild running): GPU time p50 16–23 ms; the fixed run stayed at interval 2 (3 while the GPU time passed 1.9 vsyncs, 1 for under 2% of the frames) at 0.93–0.98 of real time (Σ VFR ratio / 30 per wall second, probe `main+0x151d3e0` `@x0+0x30`: +0x30 ratio, +0x34 interval). At the same time a run with the old timestamps (`DM_NVN_GPU_TS=submission`) spent 10–70% of its frames at interval 1 and ran at 0.78–0.95. The same GPU time drives the game's dynamic resolution (`main+0x1267070`): a running average as a percentage of a 60 Hz frame scales the resolution down above 175% (29 ms) and back up below 150% (25 ms). With 0 it always sat at its maximum, and it still does whenever the game holds 30. +- **Fixed: the iPhone app was killed for memory in the first shrine (GitHub issue #9, iPhone 16 Pro Max, handheld).** An 8 GB iPhone gives this app about 3.4 GB (resident plus `离系统内存上限还有` stayed at 3,375 MB). Resident memory grew about 0.2 MB a frame from 2.6 GB at the shrine load until the kill, while the residency set and the texture counts stayed flat. The cause was the deferred-release drain running on the guest submit thread with no autorelease pool, the Metroid Dread leak (`defer_drain` now has its own `@autoreleasepool`): on the Mac a host without that pool grew 1 MB a frame in the shrine (GPU 3.0 → 4.4 GB in 2,400 frames); with it the shrine stays at 2.7–2.9 GB resident and 1.5–1.7 GB GPU, idle or walking (vmmap at 2.8 GB: about 1.6 GB guest memory, 0.75 GB private textures, 0.17 GB IOAccelerator). Found on the way: (1) about 10 serial-table entries a frame were never deleted: the game records its transient render targets (the HUD captures, the 640×360 buffers) and finalizes them before it submits, and replay put the dead serials back into `g_by_serial` (276,000 after 25,000 frames). Dead serials now go to a 512-entry table (`tex_put`, `g_ghost`; the memory line's `已 Finalize 仍被引用的贴图`). (2) The memory-pool tracking table had a fixed 4,096 slots and the phone reached 4,006 pools; a pool past the end would have its textures re-uploaded at every use. It now grows in blocks of 1,024 (`nvn_dirty.c`, up to 262,144). - **Address → memory pool lookups**: BotW keeps about 1,760 memory pools alive, and `mtl_addr` / `mtl_gpu_addr` scanned them all for every buffer binding of every draw: a CPU sample of the submit thread had about 1,000 of its samples under those scans (vertex buffers in `draw`, uniform buffers in `encode_args`). They now try the last hit, then a binary search over pools sorted by start address (`pool_find`, rebuilt when pools are added or removed). Host submit time in the shrine went from 21 ms to 7 ms per frame (p50), the median frame from 38 ms to 21 ms. - **Submissions handed to the GPU in segments.** The game does not submit frame N+1 until frame N's GPU work is done: a stack sample in the shrine had the main thread in `nvnSyncWait` (the fence of the frame just submitted) for 59% of the time and the render thread in `nvnQueueSubmitCommands` for 40%. On the Switch that submit costs nothing; here it is the host's translation (about 3,600 draws and 347 passes, 7 ms), so per frame the GPU ran 13–14 ms and then sat idle 7.7 ms (`DM_NVN_TSV` rows `@gpu_beg_ms` / `@gpu_end_ms` against `@call_ms`). The replay now commits the command buffer encoded so far every 64 draws, between two passes (`sub_cut` in `nvn_mtl.m`, called after `rp_flush` in `OP_SET_RT`; `DM_NVN_CUT=<draws>`, 0 = off). It cuts only with no open encoder and nothing deferred, merged or hoisted, commits the pending upload batch first, and continues on the same command allocator. The fence signal, the completion work (`post` ops, `g_done`) and the allocator's fence stay on the last segment: the queue runs command buffers in commit order, and every encoder starts with a queue barrier. GPU idle per frame went from 7.7 to 3.8 ms; most of the rest is the game's own work between its fence wait and the next submit (about 2.5 ms). The other five games' frame times were unchanged (Dread's p95 dropped from 19.6 to 16.9 ms). - **Frame rate in the Shrine of Resurrection on the Mac (M4 Pro)**, after the above: docked (1600×900) GPU about 13 ms, frame p50 16.7 ms (60 fps; p90 17.6 ms), was 20.8 ms before segmenting; handheld (`DM_OPMODE=0`, 1280×720) GPU p50 9.3 ms and a steady 60. Host translation is 8 ms per frame (22,000 instructions per draw), spread over Metal's encoder creation and draw calls, argument encoding, Objective-C retain / release and pool lookups. The per-pass GPU timer (`DM_NVN_PASSGPU`) is switched off under Metal 4 (`pg_render` returns at once), so GPU time cannot yet be attributed to passes. The frame loads 376 MB and stores 432 MB of render-target data (`@rt_load_mb` / `@rt_store_mb`; 325 of 347 passes load and store everything), which is the next thing to look at. @@ -154,7 +155,7 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog - **Two texture pools and two sampler pools with overlapping IDs**: the scene binds a 12,544-entry texture pool and a 3,856-entry sampler pool, the 2D UI's command buffer binds 4,096-entry ones, and both fill IDs from 256. The replay had one ID table, so the scene's default 8×8 white texture overwrote the UI's textures: the title menu's labels drew as plain bars, other menus were empty except for the cursor. Each pool now has a base in the table, in creation order; `RegisterTexture` / `RegisterSampler` add it, and binds recorded on a command buffer add the base of the pool that command buffer last set (`bind_handles`, `cb_pool_set` in `hle_nvn.c`; log line `第 2 个贴图池…从 N 起`). Deferred binds and handles read from uniform buffers are not remapped (red line). The table's capacity is the sum of the pools: the scene registers each draw's textures into fresh IDs of a ring (8,729 and up a few seconds in), which used to read as empty slots and left the screen black. - **An integer vertex input fed a float-type format**: `int_inputs` declares an input `uint` when its first translation only bit-casts it, and KTGL's character and soldier shaders then take it back with `as_type<float>`; the game binds a Half2 there (UVs). Maxwell's fetch converts by the format and the register holds the float32 bits. Metal refuses such a pipeline (`Vertex attribute m_36(5) of type uint4 cannot be read using MTLAttributeFormatHalf2`: the characters' shading draws were dropped, black silhouettes), and reading the raw bytes as integers gives denormals (UVs collapse onto one texel: flat tan characters, also Age of Calamity's soldiers). The replay now converts such attributes (half, unorm / snorm 8 and 16, 10:10:10:2, 11:11:10, BGRA8 into an int / uint input) on the CPU into float32 bits, four components with the fetch's (0, 0, 0, 1.0) defaults, in the frame arena, and binds them as UInt4 / Int4 on buffer slots 1–12 (`vconv_bind`, once per range and layout per submission; red line 89). - **Depth clamp**: the game turns `SetDepthClamp` on; it is now recorded (`OP_DEPTH_CLAMP`) and replayed as Metal's depth clip mode (it used to be ignored). -- **Open: the river's surface is flat blue in close-up battle cameras.** The water is drawn without depth, and the deferred lighting treats depth 0 (reversed-Z, cleared to 0) as sky, so it paints the fog colour there; from the map camera the river looks right. +- **Fixed: the ground under close-up battle cameras was flat steel blue.** The terrain's vertex shader (`56778feb1cfdee7e`) reads its grid position from an `RGBA16I` attribute: x and z the grid cell, y the height, which spans the whole int16 range. So does the side pass that draws the river surface into targets of its own (`db725c85ca060403`). `int_inputs` declared that input `uint4`, and Metal refuses a signed format for an unsigned input, so `vfmt_for_ct` swapped `Short4` for `UShort4`. That keeps the 16 bits but widens them with zeros: every negative height read as 32,768–65,535, and those vertices went far up out of view. Maxwell widens by the format's signedness, whatever type the shader gives the register. Where the ground was missing nothing wrote the depth (reversed-Z, cleared to 0), and the lighting takes depth 0 for sky and paints the fog colour. The G-buffer colour targets are never cleared, so they kept an old image there: cobblestones of a mesh under the terrain, which from the map camera showed through the gaps like a road. 8- and 16-bit integer formats whose signedness differs from the input's now go through the CPU conversion like the float-type formats (`vconv_int` in `nvn_mtl.m`, red line 96 once): 32-bit values extended by the format's own sign, missing components (0, 0, 0, 1) as integers. How it was traced, since two steps misled: a `HILITE` colour over the gap was stale too (the never-cleared albedo), and the end-of-frame depth showed 0 there. Recomputing every draw's clip positions from `CBDUMP` data showed that no draw of the shader reached the gap. Then the side pass's layer, which showed only scattered tiles of the river, led to the vertex format. Ruled out on the way: depth clamp (Metal's clip and clamp modes gave the same depth), the pass without a depth attachment (142595 is a screen-space decal pass: it samples the depth and writes albedo and normals), polygon offset (none), dropped draws. **Still open:** 64-pixel rectangles in the albedo where decals land. The decal pass (`417a523053457352`) and the lighting (`094a2d11253f126c`, `f54830a75e07e14a`) read their decal and light lists from texture buffers that the translation declares `texture2d`, because those handles have no PC type (`贴图类型不符：着色器声明 2，贴图 9`); skipping the decal pass removes the rectangles. - **Run-time GLSLC**: the 2D primitive, text and loading-feather shaders come from the game's own compiler at boot and are captured once on the Mac ([03](03-bringup.md), SDK 6). ### Seen in Super Mario Odyssey @@ -162,16 +163,25 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog - **Fixed: the Cap Kingdom fog never parted around Mario (reported on an iPhone; fix checked on the Mac).** On the Switch the ground fog opens around Mario as he walks through it. The fog's shading reads velocity fields that a small fluid simulation updates every frame: two 256×256 and two 128×128 RG16F targets, with a two-pass blur and a blended composite. Its only input is a point-sprite splat per mover, drawn with `alPointSpriteShader`: one vertex per mover, its size in the vertex data (Mario's main splat asks for about 18 px, others 3–5 px) and a soft disc in the fragment shader (`clamp(1 - 4r², 0, 1)` × velocity). The vertex shader writes `gl_PointSize = min(size, c0[0x610])`, and the host had nothing at that offset of the driver segment, so every splat had size 0. The simulation fields stayed exactly zero (raw dumps with `DM_NVN_TEXDUMP` + `DM_NVN_SNAPRAW`; the 8-bit dump shows velocities as black either way) and the fog stayed whole. `c0 + 0x610` holds the largest point size, and the host now writes 2047 there for the graphics stages. Writing the `SetPointSize` value instead (always 1 in this game) left one-pixel splats whose disc weight came out 0 (the fields held `-0.0`). After the fix the fields hold velocities up to ±44, and the fog clears around Mario. The build's `drv_reads` scan in `dmsh_build.py` only matches fragment-shader names (`fp_c0`), so vertex and compute reads of the driver segment are not listed in the JSON; the host treats a declared driver-segment binding as read (`nvn_refl.m`). - **Fixed: the Cap Kingdom ground showed 32×32 blocks of a stale image on every fourth frame (Mac).** Presents 4k were brighter and blocky, 4k+1 and 4k+2 clean, 4k+3 partly blocked, repeating to the decimal across runs. The blocks were already in the HDR target when the G-buffer pass ended, and only with the ground's material (`b6384a85afd3ef94`) drawn. The cause: in some frames the game first draws a depth pre-pass (`alRenderDepthOnly`) for objects its occlusion queries report visible (the queries rotate through three slots, hence the cycle), then draws the G-buffer with other vertex shaders and depth test less-or-equal, without clearing the colour targets. Metal compiled the two vertex shaders' identical position math to slightly different depths, the G-buffer fragments failed the test, and those pixels kept the previous content. Every Metal compile now uses `-fpreserve-invariance` (see "Positions must be invariant" above). Ruled out on the way: the TBDR optimizations, same-pass feedback, the present path; evaluating the game's conditional rendering on the CPU from earlier frames' results made it worse (the condition belongs to the same frame on the GPU), so the host still ignores the condition and draws. +### Seen in The Legend of Zelda: Tears of the Kingdom + +- **Fixed: the 3D scene was black (conditional rendering).** TOTK wraps nearly every `ClearColor` / `ClearDepthStencil` in `SetRenderEnableConditional(RENDER_IF_NOT_EQUAL, rec)` … `SetRenderEnable(true)`. Each clear is followed by `ReportCounter(TIMESTAMP_TOP, rec)`, the CPU sets `rec` to 1 when a target needs a clear, and `rec + 16` is never written. The condition compares the 64-bit counters at `rec` and `rec + 16`. A timestamp report puts 0 in the counter field (the time goes in the second field), so after one clear the two are equal and later conditional clears of that target are skipped. The host ignored the condition and wrote the time into both fields, so every clear ran. The first tile-lighting round, the half-resolution scene depth the light volumes test against (back faces `GREATER` against a buffer just cleared to 1.0 can never pass) and the light accumulation were all cleared before the composite read them. Profile key `render_cond`: `OP_RENDER_COND` is evaluated at replay from the records in memory. A timestamp report replayed earlier in the same submission counts as 0, and a condition on a samples-passed result (an occlusion query, known only when the GPU completes) keeps drawing. A pass dump shows the switches as `E` lines (`条件渲染 … → 不画` = skipped) and a skipped clear as `c … 跳过`. +- **Fixed: the cave was magenta (textures filled with no NVN flush).** The terrain draws sample two 121-layer ASTC 8×8 arrays (42 MB each) that TOTK creates on zeroed memory and then fills from the CPU, with no `nvnMemoryPoolFlushMappedRange` or `FlushTexels` over them (its cache maintenance runs natively). `tex_flush_upload` never marked them, so they stayed GPU-only with content 0, and Metal decodes an all-zero ASTC block as the error colour (1, 0, 1). A zero BC block decodes black, which looked like missing light. Profile key `tex_dirty_upload`: `ft_add` write-protects a registered texture's pages, and on first use (submit thread) a write since then makes it CPU-written. `DM_NVN_TEXINFO=WxH` tells the cases apart: it prints how many bytes of the texture's guest memory are non-zero at first use. +- **How the tile lighting was traced**: the G-buffer class byte selects per-class half-resolution passes through a depth buffer of class values (`gl_FragDepth = mix(0.5, 1, nibble-swapped class / 255)`, then full-screen or per-tile quads with depth `EQUAL`), so `DM_NVN_DEPTHOVR` (depth test forced to `ALWAYS`) with `HILITE` tells "rejected by depth" from "never rasterized". `TEXSNAP` copies used to run before the pass they were meant to show (classic-queue command buffer committed ahead of the Metal 4 one); they now go through the replay's copy encoder. +- **Clears inside a pass, with a partial channel mask or with a partial scissor** were dropped (red 「通道中途 / 部分掩码 / 部分剪裁的清色未实现」, 「通道中途的剪裁深度模板清除未实现」). NVN clears obey the scissor and the clear's own channel or stencil write mask, and can come between draws, while a Metal load action clears a whole attachment at the start of a pass. Such clears are now drawn as one triangle over the pass in the open render encoder (`clear_color_tri` / `clear_ds_tri` in `nvn_mtl.m`, adapted from MacNativeGames, MIT). The pipeline has the pass's formats, the clear's mask as the write mask of its slot and no writes to the other slots. Depth is `ALWAYS` (and written for a depth clear), stencil `REPLACE` through the clear's write mask, the scissor is clamped as for draws, and visibility counting is off for the triangle. The value comes from the frame arena through the triangle's own argument tables, and the draws' tables are put back after it. A clear outside a pass opens the pass as a draw would. Clears a load action can express keep it, and two cases are still cleared whole, logged once as 「深度模板清除带剪裁或部分模板掩码，按整张清」: a depth / stencil clear with a scissor outside a pass, and a partial stencil write mask. TOTK issues the first (a depth clear scissored to 8×8 of a 64×64 target), and so does Wonder. With `render_cond` TOTK's scissored colour clears are all skipped; without it they are 50×50 tiles of a 500×400 target cleared to opaque black, about five a frame. Metroid Dread logged the red line once while loading on its smoke route. +- **8-bit indices written earlier in the same submission** were read stale (red 「8 位索引缓冲本提交里有 GPU 侧写入：按当前内容展开」). Metal has no 8-bit index type, so replay widens 8-bit indices to 16-bit on the CPU. When a copy, update or fill encoded earlier in the submission covers them (`pend_hits`), the CPU reads the bytes from before that write. Those draws are now widened on the GPU (`wide_run`): a compute pass reads the indices from a 16-byte-aligned address, with the byte shift as a constant, and writes 16-bit ones into the frame arena. As in `ixc_run`, the first such draw widens every such draw after it up to the next buffer write outside a render pass, so only a draw met inside an open render pass with no widening yet would split it. In TOTK about one draw per widening, up to a dozen widenings a frame on the title and load screens, no pass split (log line `8 位索引 GPU 展开`). Draws whose indices no write covers keep the CPU widening. + ## 7. Diagnostic switches (host) | Switch | Output | |---|---| | `DM_NVN_SHOT=f1,f2:path_%d.png` | Window screenshots at those presented frames | -| `DM_NVN_PASSDUMP=f` or `f1-f2` (+ `DM_NVN_PASSDUMP_ALL=1`) | Every render pass of those frames: `P` attachments (with `mtl=` the Metal texture behind them), `C` clear colours, `K` shader keys, `B` each binding's actual source (also with `mtl=`), `d`/`D` draw state (`b0`–`b3` the blend states of targets 0–3), `Y` inline updates, `X` texture clears, `T` texture copies, `N` passes skipped because no attachment could be attached. Without `_ALL` the `K` / `B` / `R` lines cover only the first draw of each pass | +| `DM_NVN_PASSDUMP=f` or `f1-f2` (+ `DM_NVN_PASSDUMP_ALL=1`) | Every render pass of those frames: `P` attachments (with `mtl=` the Metal texture behind them), `C` clear colours, `K` shader keys, `B` each binding's actual source (also with `mtl=`), `d`/`D` draw state (`b0`–`b3` the blend states of targets 0–3; on `d`, `po=` the polygon offset enables / factor, units, clamp and `dc=` depth clamp), `Y` inline updates, `X` texture clears, `T` texture copies, `N` passes skipped because no attachment could be attached. Without `_ALL` the `K` / `B` / `R` lines cover only the first draw of each pass | | `DM_NVN_TEXDUMP=f:path_%d.png` (+ `DM_NVN_TEXSNAP=i,…`) | Every colour render target of frame f, as PNG (`%d` = serial × 100 + layer × 10 + level; 2D arrays by the attached layer). `TEXSNAP` copies the attachments of in-frame pass i (2D, or slice 0 of a 2D array) in a separate command buffer, which does not run at the pass's position in the GPU timeline: a snapshot taken right after a clearing pass showed a complete image, so read it as "what this target holds around the frame", not "right after pass i" | | `DM_NVN_TEXDUMP=size=WxH:path_%d.png` | The same dump for the first frame that opens a render pass whose first colour attachment level is W×H: for one-off work (a mask built when a prompt appears) whose frame differs between runs. With `DM_NVN_SNAPRAW=1` each image also gets a `.raw` file with the alpha channel | | `DM_NVN_CBDUMP=dir` (with PASSDUMP) | Raw uniform buffer contents per binding | | `DM_NVN_HILITE=fskey,…` | Replace those fragment shaders with solid colours: tells "geometry is wrong" from "shading is wrong" | +| `DM_NVN_DEPTHOVR=fskey,…` | Those fragment shaders' draws compare depth with `ALWAYS`: with `HILITE`, tells "the depth test rejects every fragment" from "nothing rasterizes". With `PASSDUMP`, the `D` line shows the depth range (`dr=`) and logic op (`lo=`), and `c` lines each recorded clear (slot, serial, mask, scissor; depth / stencil values) | | `DM_NVN_SHDUMP=dir` | Control + code of programs missing from the shader map | | `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`) | @@ -186,6 +196,7 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog | `DM_NVN_FB_SNAP=0` / `all` / `<keys>` | Which fragment shaders get the snapshot for same-pass feedback sampling (a draw sampling a colour attachment of its own render pass reads a copy taken just before it): none, every shader (slow, see BotW above), or a comma-separated key list instead of the profile's `feedback_snap` | | `DM_NVN_ALPHA_RED=0` / `<keys>` | Which fragment shaders read single-channel render targets mapped (R, 0, 0, 1) with alpha = red: none, or a comma-separated key list instead of the profile's `alpha_from_red` (BotW's layout captures, see above) | | `DM_NVN_CUT=<draws>` | Commit the replay's command buffer every that many draws, between passes (default 64; 0 = one command buffer per submission) | +| `DM_NVN_GPU_TS=segment` / `submission` | What timestamp reports record, instead of the profile's `gpu_timestamps`: the GPU start / end of the segment holding each report, or one time for the whole submission | | `MTL_DEBUG_LAYER=1 MTL_DEBUG_LAYER_ERROR_MODE=nslog` | Metal validation without aborting | A first-frame bring-up that worked for Mario: screenshot is black → `TEXDUMP` shows which targets have content → `PASSDUMP` the pass that writes the window → `HILITE` its fragment shader → `CBDUMP` its uniforms → read the translated `.metal` of its vertex shader. diff --git a/docs/agents/05-debugging.md b/docs/agents/05-debugging.md index 8d289a2..257f176 100644 --- a/docs/agents/05-debugging.md +++ b/docs/agents/05-debugging.md @@ -93,6 +93,8 @@ To see the order of events, put temporary observation gates (`OBS` in `patch_nso 1. The host's memory report line (`内存 至帧 …`) every few hundred frames; `footprint -p <pid>` from outside. On a phone, a crash with no log usually means `离系统内存上限还有` reached 0. 2. `vmmap -summary <pid>` twice a few seconds apart: which region grows. "owned unmapped (graphics)" = private Metal resources. 3. `DM_NVN_ALLOCPROF=1`: growth per replayed op type and per function (`tex_mtl` / pass begin / pipelines / argument buffers), command allocator sizes, and a log of new textures with serial, size, format and storage. A stream of new serials with the same sizes = objects re-created every frame. +4. GPU allocated (`显卡已分配`) climbing while `驻留集` (the residency set), `共用贴图` and the upload counts stay flat: retired objects kept alive by an autorelease pool that never drains. Any objc call on the guest submit thread outside `g_sub_pool` leaks what it autoreleases (Metroid Dread; Breath of the Wild, GitHub issue #9). Build the host with the suspect wrapped in `@autoreleasepool` and compare. +5. `序号对象` climbing by the same amount every frame with GPU memory flat: entries the serial table never deletes. A census of `g_by_serial` by class and size in `mtl_mem_stat` shows which (BotW: finalized textures replay resolved again; they now go to the small table counted as `已 Finalize 仍被引用的贴图`). Instruments' Metal Resource Events didn't record this process from the command line (`xctrace record --launch`, with `--env` for the variables); the in-host probes were faster. diff --git a/docs/development.md b/docs/development.md index b66c2c9..599c40d 100644 --- a/docs/development.md +++ b/docs/development.md @@ -38,11 +38,12 @@ What the script does: 1. It finds the game profile from the title ID (the `[title ID]` in the file names, or the NSP's ticket; the update by its `[vN]` tag). 2. It extracts the NSPs. 3. It rewrites the code modules, checking each build ID against the profile. -4. It splits out the shaders and checks them against the profile's tables. -5. It builds the iOS shader libraries. -6. It calls `build_ipa.sh`. +4. For a profile with `shader_glslc_capture on` (Three Houses), it captures the shaders the game compiles at run time, which no game file holds. It builds the Mac host into `<id>_build/tools/dmhost` and boots the game in it, headless and silent, with `DM_GLSLC_CAPTURE` pointing at `<id>_build/extract/<version>/glslc_capture`. The save, temporary storage and log all stay in `<id>_build/`. It stops the game as soon as every capture file named in the profile's shader map is there (Three Houses: 9, about a second after launch), and gives up after 30 s. A folder that already holds them all is reused. +5. It splits out the shaders and checks them against the profile's tables. +6. It builds the iOS shader libraries. +7. It calls `build_ipa.sh`. -While it runs, the console shows one line per stage: building nx_extract, extracting each version, the DLC, the rewrite, the shader split, the iOS shaders, the app build and packing the IPA. Each line shows `[k/n]`, the stage and its current step. Where the work can be counted, it adds a bar with done/total, percent, elapsed time and ETA. Counted work is bytes extracted, modules rewritten, shaders (cache check, translation, compile), files compiled by Xcode, and bytes copied and zipped into the IPA. Elsewhere a spinner shows the elapsed time. A total marked `~` is an estimate. A version's first extraction is estimated from the NSPs' size, and later extractions use the previous run's exact size. The zip is estimated from the previous IPA, or from the app's size before compression. In a terminal or in Spyder the line is redrawn in place. When the stage ends, it becomes a `✓` line with the elapsed time and a summary: +While it runs, the console shows one line per stage: building nx_extract, extracting each version, the DLC, the rewrite, the run-time shader capture (only for such profiles), the shader split, the iOS shaders, the app build and packing the IPA. Each line shows `[k/n]`, the stage and its current step. Where the work can be counted, it adds a bar with done/total, percent, elapsed time and ETA. Counted work is bytes extracted, modules rewritten, shaders (cache check, translation, compile), files compiled by Xcode, and bytes copied and zipped into the IPA. Elsewhere a spinner shows the elapsed time. A total marked `~` is an estimate. A version's first extraction is estimated from the NSPs' size, and later extractions use the previous run's exact size. The zip is estimated from the previous IPA, or from the app's size before compression. In a terminal or in Spyder the line is redrawn in place. When the stage ends, it becomes a `✓` line with the elapsed time and a summary: ``` [15:24:28] ✓ [2/7] extract v327680 (0:20): 9,779 files, 5.0 GB @@ -137,7 +138,8 @@ Record on the game version you ship. At launch the log names each list it read a - **60 fps:** the Switch version runs at 30; one patch to its code switches it to the 60 fps timing the game already has, so nothing runs too fast. **The Legend of Zelda: Breath of the Wild (`01007EF00011E000`)** -- A patch to the game's frame-interval logic lifts its 30 fps limit to 60 without changing the game speed (it falls back to 30 by itself when a frame takes too long). +- Patches to the game's frame-interval logic lift its 30 fps limit to 60. The game's time step follows its frame interval, so the speed is right at a steady 60, 30 or 20. The game picks the interval from the frame's CPU time and its GPU time, which it reads from two GPU timestamps around the frame; the host gives each timestamp the GPU time of the work around it (profile key `gpu_timestamps`). So when the device can't hold 60 the game drops to 30 at the right speed. Before, it saw only the main thread's CPU time (about 3 ms a frame in the first shrine) and stayed at 60 in slow motion: at 40 fps, about two thirds of the normal speed (an iPhone 16 Pro Max in handheld mode, GitHub issue #9; [04](agents/04-graphics.md)). +- **Memory:** about 2.8 GB in the first shrine in handheld mode, idle or walking (Mac measurement), against the 3.4 GB an 8 GB iPhone gives the app. - The game waits for each frame's GPU work before submitting the next, so the host hands its translated commands to the GPU in segments ([04](agents/04-graphics.md)). - **HUD prompt shadows and glows:** each prompt bar captures itself into a blurred mask the game keeps in the red channel of an R8 texture whose channel mapping says alpha = 1; the profile's `alpha_from_red` key makes every layout shader read alpha from red for such render targets, otherwise prompts sit on hard-edged dark boxes or cream cloud rectangles ([04](agents/04-graphics.md)). @@ -184,7 +186,7 @@ Record on the game version you ship. At launch the log names each list it read a **Fire Emblem: Three Houses (`010055D009F78000`), being ported** - **Engine:** Koei Tecmo's older KTGL (the game is co-developed with Intelligent Systems), SDK 6.5.0 (the oldest so far), update 1.2.0. Modules main + subsdk0 (GLSLC, NVN's runtime shader compiler) + subsdk1 (movie) + sdk. RomFS: the LINKDATA archive (`DATA0.bin` is the index, `DATA1.bin` holds 6.9 GB), loose overrides in `patch1`–`patch4`, WebM movies and KTSL sound banks; compressed entries and `*.gz` files are Koei Tecmo's chunked zlib ([01](agents/01-extract.md)). - **Shaders:** 1,364. `shader_split.py` reads the SPKG packages (G2S containers, as in Age of Calamity) in `DATA1.bin` and `patch*/nx/shader`, a `_SPK4100` post-effect package, and the BNSH and GLSLC outputs in `main`. SDK 6's GLSLC writes 0x848 / 0x850-byte control sections and compute code with no header, which the reader used to drop ([04](agents/04-graphics.md)). -- **Shaders compiled at run time:** KTGL compiles 9 built-in GLSL shaders at boot with the game's own GLSLC (2D primitives, text, the loading quill). No game file holds their output, so they are captured once on the Mac: run the game for a few seconds with `DM_GLSLC_CAPTURE=<version dir>/glslc_capture` (`hle_glslc.c`); the profile key `shader_glslc_capture on` makes `shader_split.py` read that folder. The NSP script needs it beside the extracted RomFS (`<id>_build/extract/v327680/glslc_capture`): without it the split tables differ from the profile and the script stops. +- **Shaders compiled at run time:** KTGL compiles 9 built-in GLSL shaders at boot with the game's own GLSLC (2D primitives, text, the loading quill). No game file holds their output, so they are captured on the Mac. The game runs with `DM_GLSLC_CAPTURE=<version dir>/glslc_capture` (`hle_glslc.c`), and the profile key `shader_glslc_capture on` makes `shader_split.py` read that folder. The outputs are identical on every run. The NSP script does this by itself in its `GLSLC capture` stage, so the NSPs and prod.keys are enough. The stage builds the Mac host into `<id>_build/tools/` and boots the game until the 9 files are in `<id>_build/extract/v327680/glslc_capture`, which takes about 10 s in all. Without the folder the split tables would differ from the profile, and the script would stop at its table check. By hand, run the game for a few seconds with that variable (step 5 of the build by hand). - **Host changes it needed** ([03](agents/03-bringup.md), [04](agents/04-graphics.md)): - it submits on two NVN queues from two threads (submissions are serialized now); - its command buffer records into one 16 MB ring chunk and frees the ring by `GetCommandMemoryUsed`, which now counts from the last `AddCommandMemory` as NVN does (frames stopped after about 15 s); @@ -193,7 +195,8 @@ Record on the game version you ship. At launch the log names each list it read a - SDK 6's audio command header has its own layout (no sound); - its character shaders read Half2 UVs through integer-declared inputs, now converted to float32 bits at replay (black silhouettes, then flat tan characters); - depth clamp is replayed. -- **Open:** the river surface shows as flat blue in close-up battle cameras (the water writes no depth and the deferred lighting takes depth 0 for sky). One red line per run says a texture buffer is bound where a shader samples 2D. +- **Fixed:** the ground under close-up battle cameras showed as flat blue. The terrain's signed 16-bit heights were read zero-extended, so the vertices with negative heights went out of view ([04](agents/04-graphics.md), Three Houses). +- **Open:** 64-pixel rectangles in the ground where decals land. One red line per run says a texture buffer is bound where a shader samples 2D: the decal and light lists. - **Memory:** about 3.5 GB resident and 2.8 GB of Metal allocations on the Mac; the game itself makes a 1.5 GB GPU-only memory pool. - **Frame rate:** 30 fps, as on the Switch. - **Mac input:** a fresh save needs `DM_KBD=<name>` for the name entry, and its confirm dialogs default to No. From a new game into the prologue battle (frame-based, about 6 minutes): `DM_KEYS="f360:A*185/45,f1200:Plus*35/219,f8700:A*120/250,f8760:Left~0.3*120/250,f8800:A*120/250,f9010:Plus*100/281"`; the game offers a save right before the battle. diff --git a/scripts/build_ipa_from_nsp.py b/scripts/build_ipa_from_nsp.py index 21be8a5..38c5961 100755 --- a/scripts/build_ipa_from_nsp.py +++ b/scripts/build_ipa_from_nsp.py @@ -1,16 +1,17 @@ #!/usr/bin/env python3 -# Unsigned IPA straight from the NSPs: identify the game (title ID → game profile dm/games/<id>/) → extract (nx_extract) → rewrite the code modules (patch_nso.py) → split out the shaders and check them against the profile's tables (shader_split.py) +# Unsigned IPA straight from the NSPs: identify the game (title ID → game profile dm/games/<id>/) → extract (nx_extract) → rewrite the code modules (patch_nso.py) +# → (profiles with shader_glslc_capture: capture the shaders the game compiles at run time, by booting it in the Mac host for a few seconds) → split out the shaders and check them against the profile's tables (shader_split.py) # → iOS shader libraries (dmsh_build.py) → build_ipa.sh. A game with no matching profile is surveyed (<NSP folder>/<title ID>_survey/survey.md), then gets a provisional profile guessed from # the survey (<NSP folder>/<title ID>_build/profile/<title ID>/, outside the repository, found through DM_GAMES_DIR) and a best-effort IPA that may not boot; SURVEY_ONLY / --survey-only stops after the survey. # Two ways to run it: open this file in Spyder and press F5 (it uses the values under "Settings" below; an empty NSP folder or prod.keys is asked for at run time); # or in a terminal: python3 scripts/build_ipa_from_nsp.py <NSP folder> -k <prod.keys> [options] (values given on the command line override "Settings", -h lists every option; # without keyboard input a missing value is not asked for but reported as an error; exit code 0 on success, 1 on error, 130 when interrupted). Command-line options entered in Spyder's run configuration work the same way. # Uses only three things: this repository, a folder holding one game's NSPs (the base game plus the updates listed in the profile's shader_versions are required; DLC NSPs, when present, are packed too and unlock the add-on content already in the game), and prod.keys. -# Every intermediate file is under <id>_build/ in the NSP folder: extract/ (the extracted game files, deleted after the IPA is built by default), aoc/ (each DLC's RomFS by index), tools/, patched_*/, shader/ (shader package and caches, kept so the next build is fast), logs/ (each step's full output), tmp/, ipa/. The finished IPA also goes to the NSP folder by default. +# Every intermediate file is under <id>_build/ in the NSP folder: extract/ (the extracted game files, deleted after the IPA is built by default), aoc/ (each DLC's RomFS by index), tools/ (nx_extract, and the Mac host for a run-time shader capture), patched_*/, shader/ (shader package and caches, kept so the next build is fast), logs/ (each step's full output), tmp/, ipa/. The finished IPA also goes to the NSP folder by default. # The shader translator is the in-repo dm/gfx/tools/dmsh_xlate_bin (the build whose fingerprints match dmsh_chain.lock; a rebuilt dll has different fingerprints and dmsh_build refuses it). # The Python running this file needs lz4, capstone and fonttools; also needed: Xcode, xcodegen, .NET SDK 10, zstd, spirv-cross. For Mario the first run takes about 35 minutes (mostly the shader libraries); later runs take about 10 (the shader libraries come from the cache). Key values are never printed. # The console shows one line per stage ([k/n] stage, with a bar, done/total, percent, elapsed time and ETA where the work can be counted); the tools' own output goes only to logs/ (PROGRESS / --progress). -import collections, ctypes, getpass, os, re, shutil, struct, subprocess, sys, tempfile, threading, time, types, zlib +import collections, ctypes, getpass, os, re, shlex, shutil, struct, subprocess, sys, tempfile, threading, time, types, zlib from pathlib import Path # ---- Settings ---- @@ -602,6 +603,83 @@ def rewrite(py, repo, exe_dir, out, env, logs, g): PROG.end(', '.join(mods)) +GLSLC_SECS = 30 # the longest capture run: Three Houses compiles its 9 shaders within 1.2 s of boot, and the run stops as soon as every file the profile names is there + + +def glslc_wanted(g): # the capture files the profile's shader map was built from (column 5: /glslc_capture/<name>#GLSLC@…); None without any (a provisional profile) + try: t = g.table('shader_map.tsv').read_text(errors='replace') + except OSError: return None + return set(re.findall(r'/glslc_capture/([0-9a-f]{16}\.glslc)#', t)) or None + + +def glslc_done(d): # capture files written whole: a GLSLC output (magic 0x19866891) holds its own length at +0x44 + out = set() + try: names = [n for n in os.listdir(d) if n.endswith('.glslc')] + except OSError: return out + for n in names: + try: b = (d / n).read_bytes() + except OSError: continue + if len(b) >= 0x48 and b[:4] == b'\x91\x68\x86\x19' and struct.unpack_from('<I', b, 0x44)[0] == len(b): out.add(n) + return out + + +def build_mac_host(repo, out, env, logs, g): + """The Mac host for the profile, built into out: dm/host's own compile line (make -n) with the output moved, since make takes no target path with spaces and dm/host/dmhost may be another build's""" + host = repo / 'dm/host' + run(['make', '-s', f'GAME={g.id}', f'gen/{g.id}/game_profile.h', 'gen/vpx/libvpx.a'], logs / 'mac_host_make.log', env=env, cwd=host) + dry = subprocess.run(['make', '-n', '-B', f'GAME={g.id}', 'dmhost'], env=env, cwd=host, capture_output=True, text=True) + line = next((l for l in dry.stdout.splitlines() if ' -o dmhost ' in l), None) + if not line: raise BuildError(f'make -n 没给出 Mac 宿主的编译行 | make -n gave no compile line for the Mac host: {dry.stderr.strip()[-500:]}') + cmd = shlex.split(line); cmd[cmd.index('-o') + 1] = out + out.parent.mkdir(parents=True, exist_ok=True) + run(cmd, logs / 'mac_host_build.log', env=env, cwd=host) + return out + + +def capture_glslc(repo, dirs, pkg_ver, patched, work, env, logs, g): + """Profiles with shader_glslc_capture on (Fire Emblem: Three Houses): the game compiles some shaders at run time with its own GLSLC, so no game file holds them. The game boots in the Mac host + (headless and silent) with DM_GLSLC_CAPTURE (dm/host/hle_glslc.c), which writes each output to <version dir>/glslc_capture, where shader_split.py reads them; the outputs are the same on every run. + dirs: {shader version: its extraction}; a folder that already holds every file the profile's shader map names is reused""" + want = glslc_wanted(g) + todo = [v for v, d in dirs.items() if not (glslc_done(d / 'glslc_capture') >= want if want else glslc_done(d / 'glslc_capture'))] + PROG.stage('GLSLC capture') + if not todo: + say(f'沿用捕获的运行时着色器 | reusing the captured run-time shaders: {", ".join(str(d / "glslc_capture") for d in dirs.values())}') + PROG.end('reused'); return + if todo != [pkg_ver]: raise BuildError(f'运行时着色器的捕获跑的是打包版本 {pkg_ver} 的代码，{", ".join(v for v in todo if v != pkg_ver)} 要另外捕获 | the run-time shader capture runs the packed version {pkg_ver}\'s code; {", ".join(v for v in todo if v != pkg_ver)} needs a capture of its own') + PROG.set(step='building the Mac host') + host = build_mac_host(repo, work / 'tools/dmhost', env, logs, g) + d = dirs[pkg_ver]; cap, rd, log = d / 'glslc_capture.partial', work / 'tmp/glslc_run', logs / 'glslc_capture.log' + rmtree(cap); rmtree(rd) + for s in ('save', 'temp'): (rd / s).mkdir(parents=True) + # everything the run writes stays in <id>_build (DM_ROOT is already there): a fresh save and temporary storage, no shader libraries yet, the profile's shader map (the host looks beside itself by default) + henv = dict(env, DM_ROMFS=str(d / 'RomFS'), DM_SAVE=str(rd / 'save'), DM_TEMP=str(rd / 'temp'), DM_MSL_DIR=str(rd / 'msl'), DM_NVN_SHMAP=str(g.table('shader_map.tsv')), DM_GLSLC_CAPTURE=str(cap)) + PROG.set(step='booting the game', done=0, total=len(want) if want else None, unit='shaders') + say(f'在 Mac 宿主里启动游戏，捕获它运行时编译的着色器 | booting the game in the Mac host to capture the shaders it compiles at run time → {d / "glslc_capture"}') + t0 = time.time() + with open(log, 'w') as f: + p = subprocess.Popen([str(host), str(patched)], env=henv, cwd=rd, stdout=f, stderr=subprocess.STDOUT) + try: # until every file the profile names is there (Three Houses: within its first frames), the host exits, or GLSLC_SECS pass + while True: + got = glslc_done(cap); PROG.set(done=len(got & want if want else got)) + if (want and got >= want) or p.poll() is not None or time.time() - t0 > GLSLC_SECS: break + time.sleep(0.25) + finally: + if p.poll() is None: + p.terminate() # SIGTERM: the host's graceful exit + try: p.wait(5) + except subprocess.TimeoutExpired: p.kill(); p.wait() + got = glslc_done(cap) + if not got or (want and not got >= want): + PROG.end(ok=False) + PROG.say(f'---- last lines of {log} ----\n' + ''.join(collections.deque(open(log, errors='replace'), maxlen=40)).rstrip('\n')) + raise BuildError(f'只捕获到 {len(got)} 个运行时着色器' + (f'，游戏档案要 {len(want)} 个' if want else '') + f'（宿主输出 | host output: {log}）| captured only {len(got)} run-time shaders' + + (f', the profile needs {len(want)}' if want else '') + f' (host output: {log})') + if want and got - want: say(f'【红】捕获到游戏档案的着色器表里没有的 | captured files the profile\'s shader map does not name: {", ".join(sorted(got - want))}') + swap_in(cap, d / 'glslc_capture'); rmtree(rd) + PROG.end(f'{len(got)} shaders in {hms(time.time() - t0)}') + + def data_rows(b): # a table's data rows (without # comment lines: the header changes with the shader_split.py version, not with the game files) return [l for l in b.splitlines() if not l.startswith(b'#')] @@ -824,7 +902,8 @@ def main(args=None): ex = work / 'extract'; ex.mkdir(parents=True, exist_ok=True); (work / 'tmp').mkdir(exist_ok=True) gb = sum({p.stat().st_size for b, u in picks.values() for p in (b, u) if p}) / 2**30 + 6 # extraction + shader cache + slack if free_gb(work) < gb and not all((ex / v).exists() for v in versions): raise BuildError(f'{work} 所在的卷剩余不足 {gb:.0f} GB | less than {gb:.0f} GB free on the volume of {work}') - PROG.plan(1 + len(versions) + bool(fam['dlc']) + 5) # nx_extract, each version's extraction, the DLC, rewrite, shader split, iOS shaders, app build, IPA pack + glslc = g.get('shader_glslc_capture', 'off') == 'on' # shaders the game compiles at run time, captured in the Mac host (capture_glslc) + PROG.plan(1 + len(versions) + bool(fam['dlc']) + glslc + 5) # nx_extract, each version's extraction, the DLC, rewrite, the GLSLC capture, shader split, iOS shaders, app build, IPA pack PROG.stage('build nx_extract') nx = build_nx_extract(repo, work, logs, env) PROG.end() @@ -846,6 +925,7 @@ def main(args=None): patched = work / f'patched_{pkg_ver}' rewrite(py, repo, game / 'ExeFS', patched, env, logs, g) + if glslc: capture_glslc(repo, {v: dirs[v] for v in need}, pkg_ver, patched, work, env, logs, g) sh = work / 'shader' split_shaders(py, repo, [dirs[v] / 'RomFS' for v in need], [dirs[v] / 'ExeFS/main' for v in need], sh / 'pkg', env, logs, g, provisional) ios = build_shaders(py, repo, sh, xlate, env, logs, g, provisional)