shaders.md

# Shaders

Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_build.py`, `dmsh_xlate`, `Program.cs`) and their pipelines. Each entry: symptom, cause, fix, the games it was found on. The [symptom index](../agents/symptoms.md) lists every symptom.

## Containers by engine

### Sumo Ninty_NX shaders and MuxFDMR streams

- **Symptom:** the split finds no shaders in `main` or RomFS, while reporting control sections in `main` with no code area. The game has `.sp4` resources and loose `.nvnbin` files rather than BNSH archives or whole GLSLC outputs.
- **Cause:** Sumo's `Ninty_NX` descriptors store control/code offsets relative to their container, whose start is recovered from the descriptor's own relative offset. The code is not necessarily adjacent to its control section. `MuxFDMR` version 2 resources wrap these containers in one or two zlib streams; their header holds both the stored and decoded lengths. Loose `NVNsh` version 1 records instead name one stage and its control/code ranges directly.
- **Fix:** `shader_split.py` reads the bounded descriptor ranges and loose records, checks each stage through the existing control/code validation, and scans decoded Mux members recursively. It validates container versions, non-overlapping ranges, complete zlib frames and exact decoded lengths before accepting their shaders. As with Needle, a magic whose container does not fit the data is a literal in executable code and is passed over; the split reads `main` whole, where a stop would end it. Detection follows the format bytes, not a filename or title ID. Original control/code bytes and shader keys are preserved; single-stage vertex programs use the existing pixel-input cover. `test_sumo.py` covers embedded containers, loose records, compressed members, malformed input and file-scanner integration.
- **Found on:** [Team Sonic Racing](../../dm/games/010092B0091D0000/README.md) 1.0.3: 337 `.sp4` resources contain 28,258 descriptors, `main` holds 586, and two `.nvnbin` files hold four stages. The split finds 4,516 unique shaders (2,116 vertex, 2,399 pixel, one compute), against zero before.

### Phyre NVN effects in Virtuos VBF archives

- **Symptom:** the shader split finds no game programs despite loose `.fx.phyre` effects and multi-GB `.vbf` archives.
- **Cause:** Phyre stores GPU code immediately before its control section, often at unaligned offsets inside `RYHP` effect containers. Virtuos `SRYK` archives keep the effects in 64 KB zlib or raw blocks. These are neither BNSH nor whole GLSLC outputs.
- **Fix:** `vbf.py` reads the Switch archive's two hash tables, file records, names and block lengths through the RomFS adapter. The split reads compiled effects and shader members without extracting textures or audio. `phyre_shaders` verifies the adjacent code header, the control's code length and stage, and preserves the original bytes and keys. Single-stage vertex programs use a pixel-input cover for padding.
- **Found on:** [FINAL FANTASY X/X-2 HD Remaster](../../dm/games/0100BC300CB48000/README.md); base and both supplied DLC packs yield 5,752 unique shaders.
- **Check:** `test_vbf.py` covers raw and zlib blocks, shared/missing members, truncation, unaligned effects, loose shader keys and inconsistent lengths. All 5,752 programs translate on the Mac.

### A file no reader reads

- **Symptom:** the shader split stops on one archive with a Python traceback (an `IndexError` or `struct.error` naming no file), or a reader's `[RED]` line about a format it does not know, and nothing is written; the NSP script loses the build at the split.
- **Cause:** a worker's exception (and `red()`, a `SystemExit`) ended the whole split through the process pool, so one archive of a format revision a reader does not know stopped every other file's shaders too. `nlg_dict` read its index without checking the layout, so an index of another revision gave an `IndexError` instead of a line naming the archive.
- **Fix:** `scan_file` leaves out a file whose reader fails, with `[RED] <file>: <why> (its shaders are left out)`, and the split goes on; at the end it prints `[RED] <n> file(s) could not be read, their shaders are left out (first: …)` and writes its tables. The NSP script shows that line and packs a provisional profile's IPA without those shaders; for a registered profile it stops there, since the profile's tables were split from every file. `nlg_dict` checks that the records and chunks fit the index and every chunk index is in range before reading it. The format itself still needs a reader: the line names the archive and, for NLG, its header bytes.
- **Found on:** Mario Strikers: Battle League (no profile; [seedds/nxnative#282](https://github.com/seedds/nxnative/pull/282)): one of its 879 Next Level Games archives, a revision `nlg_dict` does not read.

### A Yaz0 member that expands is not overlap

- **Symptom:** the split leaves out archives whose members do not overlap: `[RED] <file>/<member>.szs/#7: the archive members add up to more than 16 times the file (members that overlap)`, then `[RED] <n> file(s) could not be read`, and a registered profile's build stops there.
- **Cause:** `scan_blob` bounds archives in archives (SARC, PlatinumGames' DAT) so that a damaged or crafted file cannot make its work grow with each level: they nest at most `NEST` (16) deep, and the members a file's archives open add up to at most 16 times the file. It counted a Yaz0 member by what it expands to, then counted that member's members again, so the same decoded bytes were paid twice. A 204 KB SARC holding one layout archive that expands 9 times added up to 18 times the file.
- **Fix:** two budgets, each 16 times the file. The first counts each member by what it takes up in its archive, so it still stops members that overlap. The second counts what Yaz0 members expand to, charged from the Yaz0 header before the member is decoded. Each bound stays what it was, so the work is still within a small multiple of 16 times the file. Charging only the stored size, with no second budget, would have let the work grow by the Yaz0 ratio (about 90 at most) as well. `test_split_nesting.py` covers the layout archive's shape, the fan-out corpus input, a compressed fan-out (stopped after 19 scans; 8^6 without the bounds) and a Yaz0 member that expands 87 times, stopped before it is decoded.
- **Found on:** [Mario Kart 8 Deluxe](../../dm/games/0100152000022000/README.md) (GitHub issue #640): its race and result layout archives (`UI/<lang>/race.sarc`, `tv.sarc`, `UI/cmn/tv_result.sarc`), 14 files in 4.0.0 and 28 with 4.0.1.

### BNSH in NintendoWare archives and in main

- BNSH (`BNSH\0\0\0\0` + `grsc` block at +0x60) is found inside `bfsha`, particle effects (`esetb`, VFXB), UI layouts (`blarc`), and `main` itself.
- Variation table at +0x80 (count at +0x7c); binary program at variation +0x10; six stage pointers from program +0x08 (vs, hs, ds, gs, ps, cs).
- Stage code descriptor: control ptr +0x08, code ptr +0x10, code size +0x18, control size +0x1c. All pointers are relative to the BNSH start.

### Nintendo EPD (sead / agl / gsys): BNSH, agl sharcb and resource packs

- [Super Mario 3D World + Bowser's Fury](../../dm/games/010028600EBDA000/README.md): bfsha (BNSH) and agl sharcb (BAHS version 9, [next entry](#bahs-in-agl-sharcb)) in `ShaderData/*.szs`, particle effects (BNSH, including 25 compute shaders) in `EffectData`, six layouts and `SystemData/AglResource.szs` (BAHS), and two BNSH in `main`: 11,080 shaders. All bind with Mario's texture handle layout (`0x20 + 8 × slot`); the one failure is a geometry shader (fur fins, `RenderFinFur`), which Metal cannot express ([Geometry shaders with no Metal stage](#geometry-shaders-with-no-metal-stage)).
- [The Legend of Zelda: Breath of the Wild](../../dm/games/01007EF00011E000/README.md): 19,592 shaders: `Shader/*.sbfsha` (BNSH) and the gsys / ksys / terrain resource packs inside `Pack/Bootup_Graphics.pack` and `Pack/TitleBG.pack`, the particle effects in `Effect/*.sesetlist` (11,373 shaders, 101 compute), layouts (`.sblarc`) and 37 in `main`. All bind with Mario's texture handle layout; the full build had no failure and no unresolved binding.

### BAHS in agl sharcb

- **How it is stored:** agl `sharcb`, the engine utility shaders: code region at header +0x1c (size +0x18). Binary entries in the table area are 16 bytes `{code offset, 0, control size (0x880/0x88c; 0x848 in BAHS version 9, Super Mario 3D World), signed offset from this entry to the control section}`.
- **Pairing:** control sections are shared between programs, so pair by entry, never by searching for code near a control.
- **Found on:** [Super Mario 3D World + Bowser's Fury](../../dm/games/010028600EBDA000/README.md) (BAHS version 9).

### Zstd files recognised by their magic

- **Symptom:** `shader_split.py` finds 0 shaders in a RomFS of zstd files named `.zst`.
- **Cause:** the split decompressed only `.zs` files.
- **Fix:** it now goes by the zstd magic (`28 b5 2f fd`) whatever the name.
- **Found on:** [Paper Mario: The Thousand-Year Door](../../dm/games/0100ECD018EBE000/README.md) (8,726 of the RomFS's 9,959 files, the shader archives among them). Of the other 21 games' RomFS on hand only Super Mario Galaxy's `GDDR3.zst` (a Vulkan game, no NVN split) has such a file.

### HAL's `.cmp` containers: a size-prefixed frame

- **Symptom:** the game asks for shaders the map does not have (`[RED] shader ... is not in the shader map`, hundreds of lines), and the split finds far fewer shaders than the game uses (121 against 27,288 on Kirby and the Forgotten Land).
- **Cause:** HAL keeps its assets in `.cmp` containers: a 4-byte little-endian uncompressed size, then a frame (zstd, or LZ4 in the `ubin/` map data). The split decompressed a zstd file by its magic at offset 0, so a frame that starts at offset 4 was read raw, and the shaders inside the BFRES archives of those containers were never seen.
- **Fix:** the split opens a size-prefixed frame as well (`unsized`: zstd, with a dictionary as for a plain zstd file, or LZ4) when the file's bytes 4 to 8 are a frame magic and bytes 0 to 4 are not the zstd magic. A frame whose header records a size other than the prefix is not such a container, and the file is read as it is: Tears of the Kingdom's 1,202 `Cave/…/C.crbin.<hash>/*.quad` files are a 4-byte hash (the one in their folder name) and a zstd frame, the only files with a frame at +4 in the RomFS of 39 other games on hand. A frame that records no size must decompress to exactly the prefix's size; otherwise (it does not decompress, it holds more or less, or it names a dictionary `Pack/ZsDic.pack.zs` does not have) the split prints a `[RED]` line and reads the file as it is.
- **Found on:** [Kirby and the Forgotten Land](../../dm/games/01004D300C5AE000/README.md): 7,201 `.cmp` files, 5,121 of them zstd frames with 19,364 BNSH inside (11.2 GB decompressed) and 2,080 LZ4 frames, all of the latter in `ubin/` (0.3 GB, and no BNSH); the same layout is in 430 `.cmpbin` files (76 zstd in `lyt/`, 354 LZ4 in `snd/`). Every zstd frame records its size, the same as its prefix; no LZ4 frame does, and each decompresses to its prefix's size. The split now finds 27,288 shaders in the base game (18,476 in `gfx/`, 8,313 in `ptcl/`, 378 in `lyt/`, 67 in the loose `shader` files, 54 in `main`), against 121 before, and 27,302 with update 1.1.0 (v65536: 14 more in `lyt/`).

### BNSH in Bezel Engine archives

- **How it is stored:** the shaders are BNSH in the `.bfsha` entries of the `.bea` archives (zstd entries, [01](../agents/01-extract.md)), most of them the engine's own in `System.nx.bea` (`_BezelSystemResources/<system>/shader/*.bfsha`: PBR shadow maps 846, post effects 770, …), the rest in the layouts' `__ArchiveShader.bnsh` and the per-archive shader packs (`.bnbshpk`), plus one BNSH in `main`.
- **Reader:** `shader_split.py` scans each archive entry as a file of its own (`bea_members`).
- **Found on:** [Big Brain Academy: Brain vs. Brain](../../dm/games/0100620012D6E000/README.md) (NintendoWare Bezel Engine 1.56): 2,125 shaders in about 90 s.

### Bezel Engine 2.9 archives keep their entries inline

- **Symptom:** the split finds almost no shaders and reports `[RED] /Archive/<name>.bea: entry 0 @0x4854535341 is not an ASST entry (its shaders are left out)` for every archive (`… file(s) could not be read`); the "offset" is the bytes `ASSTH`.
- **Cause:** the 1.56 form (version 0x10100 at +8) has at +0x28 the offset of an array of entry offsets; the 2.9 form (version 0x60000) has the entries themselves there, back to back, each 0x48 bytes (its size the u32 at `ASST` +4), with an 8-byte type tag (`_G3DSHA`, `_TEXDATA`, …) at +0x1c and a 16-byte hash before the data and name offsets, which are the last two u64 of an entry in both forms.
- **Fix:** `bea_members` reads the entries in place when the array starts with `ASST`.
- **Found on:** [Super Mario Party Jamboree](../../dm/games/0100965017338000/README.md) (Bezel Engine 2.9, 949 archives, 127,203 entries): 25,693 shaders, against 246 (those in `main`) before, most of them in the 13,839 `.fshb` (`_G3DSHA`) entries; the split takes about 10 minutes.

### BNSH entries of GFPAK archives

- **Symptom:** a provisional split finds only a fraction of the shaders (3,126 of Pokémon Sword's 41,516).
- **Fix:** `shader_split.py` `gfpak_members` reads the BNSH entries of GFPAK archives (LZ4 entries, [01](../agents/01-extract.md)).
- **Found on:** [Pokémon Sword](../../dm/games/0100ABF008968000/README.md): 41,516 shaders, nearly all BNSH entries of the 7,699 GFPAK archives under `bin/archive` (the models' and effects' shader archives), beside loose BNSH files (`bin/battle/waza/model/**/*.bnsh_vsh` / `_fsh`, `bin/graphics`, `system_resource/shader`) and two in `main`. The YEBIS shaders it also draws are in `main` ([Shaders in a zlib stream in main](#shaders-in-a-zlib-stream-in-main-yebis-tfxo)).

### Oodle entries of GFPAK archives

- **Symptom:** the split counts hundreds of files as unreadable (678 on Pokémon Legends: Arceus), each with `[RED] <file>.gfpak: entry 0 has compression 3 (Oodle), which this reader does not decode`, and leaves their shaders out.
- **Cause:** every entry of Arceus's GFPAKs is compression 3, Oodle Kraken, where every entry of Sword's is LZ4 ([BNSH entries of GFPAK archives](#bnsh-entries-of-gfpak-archives)).
- **Fix:** `shader_split.py` `gfpak_members` decodes compression 3 with the game's own decoder, as it decodes [Trinity packs](#bnsh-in-trinity-packs) (`oodle`, `DM_OODLE_HOST` / `DM_OODLE_MODULES`). `profile_guess.py` `build_sources` counts the GFPAKs that hold an Oodle entry, so the NSP script sets that decoder up for the split (`oodle_packs`). `dm/gfx/tests/test_gfpak.py` covers each compression. The Mac path was checked without Arceus's files: a GFPAK of 300 Kraken streams from Pokémon Scarlet's packs (28 MB raw), read by `gfpak_members` through the Mac host, decodes the same with Scarlet's and Z-A's `OodleLZ_Decompress`, and `build_sources` sets `oodle_packs` for it but not for an LZ4-only GFPAK.
- **Found on:** Pokémon Legends: Arceus 1.0.2 (`01001F5010DFA000`, no profile yet; `main` has `OodleLZ_Decompress`): all 678 GFPAKs under `bin/archive`, Oodle in every entry. The provisional split finds 12,851 shaders with them, 69 more than without. 73 shader map rows first appear in a GFPAK: `field/resident_release.gfpak` 40, `appli/battle_resident.gfpak` 26, `appli/system_resident.gfpak` 5, `appli/poke_dex.gfpak` 2.

### BNSH in Trinity packs

- **How it is stored:** every shader is a BNSH inside the Oodle-compressed files of `arc/data.trpfs` ([01](../agents/01-extract.md)).
- **Reader:** for any game with Trinity packs (`ONEPACK` files), `trpfs_jobs` / `trinity_shaders` in `shader_split.py` name the packs from `data.trpfd`, decode every file with the game's own Oodle in the Mac host (`DM_OODLE_HOST` / `DM_OODLE_MODULES`, set by the NSP script) and scan it like a loose file.
- **Copies that differ:** the layout archives' control sections are 104 KB long, and their stage descriptors run 456 bytes past the control's own length (control +0x20), with bytes that differ between copies of one shader; `same_shader` compares copies only up to the control's own length.
- **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md) (Game Freak's engine): 28,603 shaders (11,633 vertex, 16,845 pixel, 125 compute) and 26,051 variation pairs: the material shaders in `system_resource/gfx2/shader/NX64/env_ikkaku` (`material/standard.bnsh`, `material/transparent.bnsh`, `system/light.bnsh`, and the `special/imgui.trsha` pack with the engine's system shaders), particle effects (`.ptcl` / `.trtml` in `ik_effectdata`, 19,842 of the shaders), UI layouts (`__ArchiveShader.bnsh` in the `.truiv` packs of `uidata`) and the 40 BNSH in `main`.

### GLSLC outputs in main and in loose files

- **How it is stored:** 13 Sentinels (Vanillaware's engine) has no shader file in RomFS, whose six CRI CPK archives hold textures, scripts, CRI audio and USM movies: 838 GLSLC compiler outputs sit in `main`'s read-only segment, each referenced from a `{pointer, size, 0}` table in `.data`. Other engines keep whole outputs in loose RomFS files of their own formats (below).
- **Layout:** an output starts with the magic `0x19866891`, its total size is at +0x44, and it holds exactly one control section with the code area right after it (control +0x20 = control length); a debug section with the NV assembly and the GLSL source follows. The stage comes from SPH word 0, as for BAHS. Outputs that hold several stages, and SDK 4's outputs, are read through their section headers ([GLSLC outputs holding several stages](#glslc-outputs-holding-several-stages), [GLSLC outputs read through their section headers](#glslc-outputs-read-through-their-section-headers)).
- **Reader:** `shader_split.py` reads them in every game's `main`, accepting only outputs that pass the control / code / length check. Loose RomFS files are searched the same way: it reads GLSLC outputs in any RomFS file no other reader takes. When the reader was written, of the other profiles' files only Metroid Dread's MSHD files held any, which `mshd_shaders` reads ([MSHD in MercurySteam's bshdat files](#mshd-in-mercurysteams-bshdat-files)).
- **Found on:**
  - [13 Sentinels: Aegis Rim](../../dm/games/01003FC01670C000/README.md): 838 in `main`.
  - [Fire Emblem Engage](../../dm/games/0100A6301214E000/README.md): 11 in `main`, Unity's own.
  - [FINAL FANTASY TACTICS - The Ivalice Chronicles](../../dm/games/010038B015560000/README.md): 257 whole GLSLC outputs in the engine's technique files (`*.tec`: `resources_nx/release/master/shader/*.tec`, magic `TEC\0`, each output whole after a table of parameter names) beside 49 in `main`.
  - [Absolum](../../dm/games/0100A83022A22000/README.md): 163 GLSLC outputs in `shaders.nx.bin` (`rom:/shaders.nx.bin`: a u32 count, then per entry a u32 hash, a u32 size and one whole output): 162 unique shaders, 83 pixel and 79 vertex, every vertex shader paired with the single-stage cover ([Single-stage programs paired with a cover](#single-stage-programs-paired-with-a-cover)). Its `.xnb` effects are XNA's D3D9 effects, MOJOSHADER linked in, which the GPU does not run.
  - [Catherine: Full Body](../../dm/games/0100BF00112C0000/README.md): 3,050 loose files in `data/shaders_binary` (`*.vs` / `*.ps`, one stage each, one output per file), read like FINAL FANTASY TACTICS' technique files; its vertex shaders are in two parts ([VertexA and VertexB parts merged into one vertex shader](#vertexa-and-vertexb-parts-merged-into-one-vertex-shader)).

### GLSLC outputs holding several stages

- **Symptom:** the split found no shader in the game (0 shaders).
- **Cause:** an output may hold several stages compiled together as one program (Baba Is You: the 120 outputs in `Assets.dat`, each a pixel and a vertex shader, then a reflection section), which the single-stage reader does not take.
- **Fix:** such an output is read by its section headers (+0x4c section count, section *i*'s header at +0x90 + 0x90 *i*, type 0 = GPU code; `glslc_program`; the same headers SDK 4's outputs are read through, [GLSLC outputs read through their section headers](#glslc-outputs-read-through-their-section-headers)), and its stages are paired with each other, not as single-stage programs. The NSP script now stops at the split when it finds none.
- **Found on:** [Baba Is You](../../dm/games/01002CD00A51C000/README.md#shaders).

### GLSLC outputs in a zip archive, and those NROs compile at run time

- **Symptom:** an NVN game's split finds none of its shipped shaders, and the ones it compiles at run time were not captured: the game's renderer is in NROs that call `glslcCompile`, and its whole RomFS is one zip.
- **Cause:** Portal keeps all its files in `game.zip` (41,416 members, Deflate or stored). The Source engine's D3D9 layer (`materialsystem.nro`) links each vertex and pixel shader combination into one program; the programs of the game's precache list (`portal/glbaseshaders.cfg`, 219 entries) ship compiled as `nxcontent/nvsb/<hash>.nvsb` (`NVSB`, version 0x65, the file's length at +8, a 0x60-byte header, then one GLSLC output holding a vertex and a pixel shader). A deflated member hides the GLSLC magic from the loose-file scan. The game compiles the rest with `glslcCompile` (its movie shaders and its own blits: 10 outputs within 30 s of boot), and the capture rule looked only for `main` holding GLSL sources and importing `glslcCompile`.
- **Fix:** `shader_split.py` reads zip archives member by member in parallel ranges (`zip_jobs` / `zip_shaders`; `zipfile`), each member scanned like a loose file, whole GLSLC outputs included. `profile_guess.py` `shader_sources` asks for the run-time GLSLC capture when an NRO the game loads imports `glslcCompile` and a subsdk defines it (`glslc_capture`). Portal's map: 242 keys from the `.nvsb` files and 13 from the capture, 229 vertex and pixel pairs.
- **Found on:** [Portal](../../dm/games/01007BB017812000/README.md). No other profile's RomFS has a zip archive holding shaders.

### GLSLC outputs in Criterion's Bundle 2 resources

- **Symptom:** an NVN game's split finds a few dozen shaders (`main`'s and a loose file's) and the screen stays black but for the loading indicator, with `[RED] shader … is not in the shader map` and `program has no vertex shader`; the RomFS holds `.BNDL` files that start with `bnd2`.
- **Cause:** Criterion's Bundle 2 (Burnout Paradise's format; version 3 here) keeps each resource zlib-compressed: the header gives the entry count and table offset (+0x10, +0x14), four block data offsets (+0x18) and the flags (+0x28, bit 0: compressed); a 0x50-byte entry gives the resource ID, four sizes on disk (+0x20, alignment in the top nibble), four offsets (+0x30) and the type (+0x44). Need for Speed: Hot Pursuit Remastered's `SHADERS.BNDL` holds 2,112 whole GLSLC outputs in block 1 of its type 0x8 resources, none of which a raw scan of the file can see.
- **Fix:** `shader_split.py` reads a file that starts with `bnd2` resource by resource (`bnd2_resources`) and scans each inflated block for GLSLC outputs (`bnd2_shaders`), by the outputs' magic rather than the resource type, whose numbering is per game. Hot Pursuit Remastered: 89 → 2,154 shaders, the 10 its first frames asked for among them.
- **Resource validation:** the reader believes no size in the file. The header and the entry table must lie inside the file, each block's data after the table and inside the file, and each block's declared size (alignment bits masked) must match what is stored; a block that declares a size and stores nothing is an error too. Inflation stops one byte past the declared size and needs a complete stream, so a size that lies cannot grow without bound. A bundle that fails is named in a red line and its shaders are left out, like any file no reader can read (`scan_file`). The platform field (+0x08) is not checked: nothing here reads it, and a big-endian bundle already fails on its version. Measured on the game: [What it needs](../../dm/games/010029B0118E8000/README.md#what-it-needs).
- **Found on:** [Need for Speed: Hot Pursuit Remastered](../../dm/games/010029B0118E8000/README.md).

### G2S containers in KTGL RDB archives

- **How it is stored:** each material's shaders are a G2S container (`_S2G3000`, type KTID `0x7bcd279f`; 23,108 entries in `asset/system.rdb.bin*`, 2.3 GB inflated) of GLSLC outputs, as in 13 Sentinels' `main`, inside zlib-chunked IDRK entries ([01](../agents/01-extract.md)). One `_DOK0000` container (type `0x20a6a0bb`, the engine's system resources) holds 151 more bare GLSLC outputs, 21 of them drawn in the first battle.
- **Reader:** `shader_split.py` maps the RDB files and inflates only entries that start with those two magics (`idrk_shaders`), and the 117 outputs in `main` as in every game. Many outputs carry 0x850-byte control sections and the compute ones have no code header, so `glslc_shaders` accepts controls from 0x760 bytes and takes the stage of headerless code from the thread group size at control +0x748. Pairs come from the single-stage cover (`separable_pairs`).
- **Found on:** [Hyrule Warriors: Age of Calamity](../../dm/games/01002B00111A2000/README.md) (Koei Tecmo's KTGL): 40,726 shaders (37,043 pixel, 3,440 vertex, 241 compute, 2 geometry), Mario's texture handle layout (`0x20 + 8 × slot`). [ROMANCE OF THE THREE KINGDOMS 8 REMAKE](../../dm/games/0100C2A01A730000/README.md)'s newer KTGL keeps them (`_S2G5000` / `6000` / `7000`, and `_DOK0000`) in `.fdata` files whose IDRK entries are LZ4 or zlib in chunks of another form ([01](../agents/01-extract.md)): 59,214 outputs there and 149 in `main`, all of them built.
- **Note:** the older KTGL of Three Houses packs the same G2S containers differently ([SPKG packages in KTGL LINKDATA](#spkg-packages-in-ktgl-linkdata)).

### SPKG packages in KTGL LINKDATA

- **How it is stored:** shader packages (`SPKG`: `patch*/nx/shader/*.spkg` and 132 entries of `DATA1.bin`) of G2S containers as in the [RDB entry](#g2s-containers-in-ktgl-rdb-archives), one `_SPK4100` package without G2S (the post effects, `DATA1.bin` entry 549), and map and model archives carrying packages of their own.
- **Reader:** `kt_shaders` inflates every LINKDATA entry and loose `*.gz` / `*.spkg` file ([01](../agents/01-extract.md)), takes every GLSLC output that passes the triple check, and scans its SARCs for BNSH; `DATA1.bin` is split into 64 scan jobs by entry range (`linkdata_jobs`).
- **Symptom:** 24 shaders before the control-length fix, 1,355 after.
- **Cause:** its GLSLC (output version 0x1d.2) writes the 0x848 / 0x850-byte controls and headerless compute code described in the RDB entry, which `glslc_shaders` used to drop. A few outputs of version 0x1d.0 keep something else at control +0x20 (odd values such as 0x136b), so a control length must be a multiple of 4; the section-header reading now takes 8 of them ([GLSLC outputs read through their section headers](#glslc-outputs-read-through-their-section-headers)).
- **Found on:** [Fire Emblem: Three Houses](../../dm/games/010055D009F78000/README.md) (the older KTGL).
- **Note:** nine more shaders exist only at run time ([Built-in shaders compiled with the game's own GLSLC](#built-in-shaders-compiled-with-the-games-own-glslc)).

### GLSLC outputs in base64 in Heaps' shader cache

- **Symptom:** the NSP script stops at `the shader split found no shaders in main or the RomFS` on a game whose `main` imports `glslcCompile` and whose RomFS is one or two `.pak` files.
- **Cause:** Heaps keeps its files in a PAK archive (`PAK`, a version byte, i32 header size, i32 data size; version 1 adds a 64-character hash; then a tree of named entries, `DATA`, and the file data, which starts at the header size). Its `hxsl.CacheFile` for the Switch, `shaders.cache.nx`, holds a table of programs (`!<signature>\t<vertex hash>\t<pixel hash>`), then each stage's GLSL; the GLSL of each program's pixel stage ends with the program's GLSLC output in base64, `//BIN=<base64>#`. The output's magic is hidden by the base64, and GLSL strings are not in `main`, so neither the loose-file scan nor the GLSLC capture rule sees it.
- **Fix:** `shader_split.py` lists a file that starts with `PAK` by its own header (`heaps_pak_files`; a file that does not parse as one, ending in `DATA`, goes on to the other readers), reads only its `shaders.cache.nx`, and decodes each comment (`heaps_shaders`). Each output holds both stages, so they pair with each other ([GLSLC outputs holding several stages](#glslc-outputs-holding-several-stages)).
- **Found on:** [Dead Cells](../../dm/games/0100646009FBE000/README.md): `res1.pak`, 1,222 outputs, 1,262 unique shaders and 1,222 pairs.

### Gzip-wrapped shader archives

- **Symptom:** only movies render, while the 3D scene reports missing shaders and programs without a vertex stage; the shader split finds loose programs but misses those in `Shaders/Shaders.arc.gz`.
- **Cause:** the archive is standard gzip around a Wii U8 archive containing GLSLC outputs. Its shader signatures are hidden by compression, so the loose-file scan discards it before the existing GLSLC reader can discover its programs.
- **Fix:** `shader_split.py` recognizes gzip by its magic, regardless of the filename, and decodes it before the shader scan. `ungzip` validates the stream's CRC and end, including concatenated members, and bounds the total decoded bytes to the greater of 1 GiB and `NEST` times the compressed size. Deflate reaches about 1000x on repetitive data, and an unreadable file stops a registered game's build, so the floor lets a small, highly compressible file without shaders through and still caps a gzip bomb. A damaged or oversized source is named as unreadable and contributes no partial shader output. The existing reader retains each program's vertex/pixel pairing; Gust's differently encoded `.gz` files keep their own reader below.
- **Found on:** [The Legend of Zelda: Skyward Sword HD 1.0.1, issue #497](https://github.com/seedds/nxnative/issues/497#issuecomment-6550286909), reported with a source-only patch by zatteogit. Measured on that build's own file: 4.9 MB of gzip decode to a 22.8 MB U8 archive with 2,096 GLSLC programs in `cache/` (one a byte-identical copy of another) and 2,096 fixed 632-byte records in `desccache/` that hold no shaders. The split takes 2,412 unique shaders (384 vertex, 2,028 pixel) and 2,095 pairs from it, and 2,470 shaders for the whole game with `main`. No registered game profile or local gameplay validation is added here: the game has not been booted with these shaders. The separate movie-upload and UNORM rendering changes in that report remain outside this fix.
- **Check:** `test_split_gzip.py` uses an authored U8/GLSLC fixture: raw and gzip inputs produce byte-identical package/table files through the full CLI, with two shader stages and one pair. It also covers magic-based discovery, timestamp bytes resembling other compression headers, concatenated members, corrupt CRCs, truncated footers, trailing garbage and the decoded-byte budget.

### Gust's chunked zlib `.gz` files

- **Symptom:** every frame black, the title screen included, while the game runs at full rate and draws (about 19,000 NVN calls a frame), with `[RED] shader … is not in the shader map` (26 keys in the first minute) and `program has no vertex shader`; the dumped shaders (`DM_NVN_SHDUMP`) match no file byte for byte, and some differ from a split shader only in the control section's compile information (+0x7d0), which is a sign of another GLSLC compile of the same source, not of a shared section.
- **Cause:** Gust's `*.elixir.gz` files come in two forms: most are `CRAE` archives stored as they are, which the split reads like any loose file, and the rest are chunked zlib, each chunk a u32 stored length and one zlib stream (16 KB inflated, the last shorter), ended by a zero length. The post effects' archive `Data/NX/PostEffect/pe_pack.elixir.gz` is one of them, with 162 GLSLC outputs (128 keys found nowhere else, among them the frame composite and the copy passes), which a raw scan cannot see.
- **Fix:** `shader_split.py` inflates a `.gz` file in that form (`gust_gz`: every chunk must be one whole zlib stream and the zero length must end the file, so another format is left alone) and scans it as a loose file. Atelier Ryza DX: 921 → 1,049 shaders, every key of the opening, title, story and field runs.
- **Found on:** [Atelier Ryza: Ever Darkness & the Secret Hideout DX](../../dm/games/0100199023B2E000/README.md) (Gust's engine on KTGL): 156 of its 4,112 `.gz` files are chunked zlib.

### GLSLC outputs in FromSoftware's BND3 and dvdbnd archives

- **How it is stored:** RomFS is two archives, `dvdbnd0.bdt` and `dvdbnd1.bdt` (3.6 GB), each indexed by the `.bhd5` beside it: `BHD5`, 0xff (little-endian), u32 1, u32 file size, u32 bucket count, u32 0, u64 bucket table offset; each 16-byte bucket `{u32 count, u32, u64 offset}` points at 16-byte entries `{u32 path hash, u32 size, u64 offset}` (paths only hashed). 5,184 entries are DCX containers (`DCX\0`, big-endian header, `DFLT` = a zlib stream at +0x4c), most of them BND3 archives (28-byte entries `{u8 flags, i32 size, u64 offset, i32 ID, u32 name offset, i32 unpacked size}` in formats 0x5c / 0x7c). Six BND3 archives in `dvdbnd0.bdt` hold the shaders, one GLSLC output per entry (`.vpo` vertex, `.fpo` pixel, named `Shader/FRPG_Sfx/NVNShaderBinary/ps/…` or by base name in the two FLVER material archives, 2,602 + 486 entries).
- **Reader:** `bhd5_entries`, `undcx`, `bnd3_members`, `bdt_shaders` in `shader_split.py`, in 64 jobs by entry range, about 10 s. Only archives that contain a GLSLC magic are parsed (the others use format 0x78).
- **Note:** SDK 4's GLSLC needs its section headers read ([GLSLC outputs read through their section headers](#glslc-outputs-read-through-their-section-headers)).
- **Found on:** [DARK SOULS: REMASTERED](../../dm/games/01004AB00A260000/README.md).

### MSHD in MercurySteam's bshdat files

- **How it is stored:** `system/shd/*.bshdat`, the only place the game's shaders live (nothing in `main` or the `.pkg` packs). Each stage is a control section, whose length is at control +0x20, sometimes followed by a small compiler-option table (`fastmath true`), then the code area at the next code header. Every file holds vertex / pixel stages in alternating order (1, 2 or 4 programs) and is paired two by two.
- The control section carries NV assembly text after its first 0x880 bytes, and the SASS disassembly follows each code area, so the binding layout can be read straight from the files (`TEX` / `TEXS` handle immediates).
- **Reader:** `mshd_shaders`.
- **Found on:** [Metroid Dread](../../dm/games/010093801237C000/README.md) (MercurySteam's engine): 478 files, 1,342 stages, 1,010 unique keys (359 vertex, 651 pixel), 659 pairs.

### SDF definitions in RE Engine PAKs

- **How it is stored:** shader definitions (`SDF\0`, 2,685 files in the PAK view, [01](../agents/01-extract.md)): u16 version 1, u16 entry count, u64, then 0x108-byte entries from +0x10. Entry +0x10 holds six u64 offsets (vertex, hull, domain, geometry, pixel, compute; 0 = none), each to `{u32 control length, u32 code length}` followed by the control section (0x880 bytes) and the code area; compute code has no header. An entry's stages are one program, so its vertex / pixel stages are a pair. Empty definitions are 16 bytes of version 0.
- **Reader:** `kpka_jobs` maps the PAKs, takes the newest copy of each path, peeks at 4 decompressed bytes and inflates only SDF entries (`sdf_shaders`).
- **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md): 63,111 shaders (30,233 vertex, 21,475 pixel, 11,271 compute, 53 geometry, 33 hull, 46 domain) and 102,345 pairs, in 5 s with 12 jobs.

### Shader records in Snowdrop SDF data slices

- **Symptom:** the split finds no shaders in a RomFS made of `.sdfdata`, `.sdftoc` and `.sdfver` files.
- **Cause:** Snowdrop stores the shaders in individually compressed records inside its `BERG` data slices. A cache record can itself be wrapped in an SDF zstd chunk. These SDF archives are a different format from RE Engine's shader definitions above. A shader record has a 20-byte header `{u32 codec, flags, 0, raw length, stored length}`: `0xabcd00c5` precedes zlib data, `0xabcd00c6` a zstd frame; reflection data precedes the NVN sections inside it. Both supplied games use SDF version 22, so that version does not select the codec.
- **Fix:** `shader_split.py` recognizes `BERG` by magic and reads the length-delimited compressed records (`snowdrop_shaders`). Every control section is followed immediately by its code, with lengths from control +0x20 and +0x704; they are byte-packed, so a control need not be aligned. A single record can hold more than one shader. Each section must pass the control / code / stage checks. Programs hold one stage, and the vertex shaders take the [single-stage cover](#single-stage-programs-paired-with-a-cover). The reader also peeks at each zstd frame's decoded header and finishes only frames beginning with a cache record. This finds nested records without treating a coincidental frame magic or a bare frame containing control-like bytes as a shader asset.
- **DLC:** additional shader records ship in the DLC's data slices. The NSP script includes each extracted DLC RomFS as an extra shader root, without an ExeFS module, and fingerprints its index and extraction timestamps for split reuse. Vulkan and OpenGL builds include the DLC roots too. When installing or removing optional DLC changes an NVN shader union, the script uses an external profile with new shader tables only after every non-DLC shader row and fact matches the registered tables; package indices may change. The registered module builds, settings and build-gated cuts are preserved. A cached proof of this comparison permits repeat builds to reuse the split.
- **Found on:** [Mario + Rabbids Sparks of Hope](../../dm/games/0100317013770000/README.md) (zstd); [Mario + Rabbids Kingdom Battle](../../dm/games/010067300059A000/README.md) (zlib records, some wrapped in zstd chunks; shorter control sections and headerless compute code).

### CompiledBytecode resources in Frostbite superbundles

- **Symptom:** the survey finds no shader container in a RomFS of `.toc` and `.cas` files (`Data/NX`, `Patch/NX`, `Data/layout.toc`, `Data/initfs_NX`); the split reads every file loosely and finds nothing.
- **Cause:** Frostbite keeps its data in bundles that superbundle tables (`.toc`, magic `00 D1 CE 01`) list, the bundles' manifests and items in the install chunks' cas files, each item in blocks of at most 256 KB compressed with Oodle Leviathan (`8C 0C`). A toc entry gives each item's cas file by an 8-byte cas id (Data or Patch, the install chunk's `persistentIndex` from `layout.toc`, the cas number), its offset and stored size; items of one bundle need not lie together. The NVN shaders are the bundles' `CompiledBytecode` resources (`shaders/bytecode/<GUID>`, res type `0x0e2b02f7`, the type name's hash): a 0x20-byte header, then a `VNVN` record (`fb::NvnCustomShaderData` of the renderer NRO) with one control section and code area at offsets its trailer gives. `ShaderBlockDepot` resources hold parameter blocks only.
- **Fix:** `dm/gfx/tools/frostbite.py` reads `layout.toc`'s install chunks (a DbObject), every toc's bundle entries, the manifests and the blocks; `shader_split.py` (`frostbite_jobs`, `frostbite_shaders`) decodes each `CompiledBytecode` resource with the game's own Oodle in the Mac host's tool mode, as for Trinity packs, and leaves the `.toc` and `.cas` files out of the loose scan. `profile_guess.py` turns the Oodle decoder on for a RomFS with `Data/layout.toc`. Data and Patch are read alike: keys are content hashes, so a base copy the update replaced only adds an unused shader. The resources name no vertex / pixel pairs (Frostbite pairs them in its pipeline data), so the tables have none.
- **Found on:** [EA SPORTS FC 27](../../dm/games/0100C49025D3E000/README.md) (Frostbite 2025.1.2): 83,325 bundles, 10,401 resources, 10,342 shaders (3,738 vertex, 3,306 pixel, 3,298 compute), 35 MB stored; the split takes about 6 s.

### SPK shader package in an MT Framework ARC

- **How it is stored:** the engine's one shader package, `sc\NX_64\root` (`SPK\0`, type hash `0x02358e1a`) in `nativeNXx64/sa/NX_64/root.arc` ([01](../agents/01-extract.md)). After the package's tables every shader is a code area followed directly by its control section, back to back to the end of the file: 230 vertex shaders, then 230 pixel shaders.
- **Reader:** each control section names its code length (+0x704), so the code is the bytes right before it, and every block must start where the previous one ended (`mtspk_shaders`). The tables that say which vertex and pixel shaders the engine combines are not read: the vertex shaders are paired with the single-stage cover (one pixel shader reads every input any of them reads).
- **Note:** `CustomShaderPackage.mfx` (`MFX\0`, the effect descriptions) holds no shader code, and neither `main` nor any other ARC entry does.
- **Found on:** [The Great Ace Attorney Chronicles](../../dm/games/010036E00FB20000/README.md) (Capcom's MT Framework): all 460 build, with Mario's binding layout; every binding decodes.

### PKZ version 1 archives

- **Symptom:** `[RED] ... PKZ version 0x10000: only 0x20000 is known` stops the shader split.
- **Cause:** the PKZ reader accepted only the version first found on Bayonetta 3. Version `0x10000` uses the same 0x20-byte header, 0x20-byte entries and relative name offsets; the last header word is zero rather than a v2 alignment. The named compression method still determines how each entry is decoded.
- **Fix:** `shader_split.py` accepts both `0x10000` and `0x20000`, retaining rejection of unknown versions. The existing Zstandard decoder and DAT/shader scanners handle the contents. Synthetic directory fixtures cover both versions in `dm/gfx/tests/test_pkz.py`.
- **Found on:** [ASTRAL CHAIN](../../dm/games/01007300020FA000/README.md).

### PlatinumGames' shader objects

- PlatinumGames' shader objects: `VSO\0` / `PSO\0` / `CSO\0` / `HSO\0` / `DSO\0` (the stage), u32 compiler version, u32 control length, u32 control offset (0x18), u32 code length, u32 code offset, and nothing else. The engine makes a program of each and binds them per stage at draw time, so they are single-stage programs, paired with a cover pixel shader (`shader_split.py` `pg_shader`; the archives: [01](../agents/01-extract.md)). It finds its built-in ones by name (`hwfillcolor.vso`, …) and checks each object's GPU code version against the driver's (the NVN layer reports the translator's supported GPU-code range: [Newest GPU code version reported by the driver](nvn-objects.md#newest-gpu-code-version-reported-by-the-driver), SDK 13.3.2).
- **Found on:** [Bayonetta 3](../../dm/games/01004A4010FEA000/README.md): 8,024 objects in `root_000.pkz` (`shader.dat`, `shader2.dat`, `shadereff*.dat`, `autoshadereff.dat`); 6,424 distinct keys (4,288 pixel, 2,047 vertex, 78 compute, 6 domain, 5 hull).

### Size-delimited PlatinumGames shader objects

- **Symptom:** the split finds zero shaders and refuses the shader build, although the CPK's DAT archives contain the objects passed to NVN at boot.
- **Cause:** another PlatinumGames wrapper stores a binding table before the code and control sections, rather than the stage-tagged object above or a BNSH file: u32 prefix length (including that word), binding bytes, u32 code length, code, u32 control length, control. The observed prefix is `0x88` bytes and the control section `0x880` bytes.
- **Fix:** `shader_split.py` `pg_sized_shader` recognizes the self-described section boundaries and NVN magics (`0x12345678` for graphics code, `0x12345679` for compute, `0x98761234` for control). The prefix must be aligned and in bounds, both sections must fit, the record must end with its control section, and the control's code length at `+0x704` must agree. Graphics stages come from SPH; the compute magic identifies compute. The reader emits `#PGSO` records so the existing single-stage cover pairing applies. Recognition needs no title or filename rule, and the shader translation chain is unchanged.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): the objects in five DAT shader archives inside `data000.cpk`; the game note records the census and the runtime-key checks.

### BNSH files in PlatinumGames DAT archives inside CRI CPK archives

- PlatinumGames' engine in its Wii U form keeps its shaders as BNSH files of one stage each, the `.vso` / `.pso` members of PlatinumGames DAT archives (`shader.dat`, `shader2.dat`, `shadereff.dat`) inside a CRI CPK archive, all CRILAYLA-compressed ([01](../agents/01-extract.md) has the CPK layout). `dm/gfx/tools/cpk.py` reads the @UTF tables (and the masked ones of Unicorn Overlord's CPKs, the same tables as before once unmasked), the TOC and CRILAYLA (a decoder over a bit string, about 29 MB/s); `shader_split.py` `cpk_jobs` splits the CPKs into jobs and scans each entry (a CPK it cannot read is scanned as one file, with a `[RED]` line).
- Each file is one variation with one stage, so the split pairs nothing; as in Hades II, the vertex shaders write every output their pixel shaders read and no pipeline failed to link.
- **Found on:** [Bayonetta 2](../../dm/games/01007960049A0000/README.md): 1,856 distinct keys (1,547 pixel, 309 vertex) in `shader.dat`, `shader2.dat` and `shadereff.dat` of `data000.cpk` (the split about 2.5 minutes); [Bayonetta](../../dm/games/010076F0049A2000/README.md): 458 distinct keys (236 pixel, 222 vertex), the members of `shader.dat` in `data00.cpk` (about 1.7 minutes). All build for both.
- **Not checked:** 13 Sentinels and Catherine: Full Body's CPKs (no files on the drives).

### Monolith ARH1, XBC1 and HCPS/SLCT shader containers

- **Symptom:** archive or loose asset shaders are absent from the map; only boot shaders in `main` are found, sometimes reported as raw controls with no identified code area.
- **Cause:** Monolith uses ARH1/ARD archives and XBC1 streams. HCPS shader containers live inside models, map streams and standalone shader files; a standalone SLCT in `main` omits the HCPS header.
- **Fix:** `shader_split.py` reconstructs paths from the ARH1 XOR trie, reads bounded archive entries in parallel, and decodes stored/zlib/zstd XBC1 streams recursively. The ARH stored size excludes an XBC1 wrapper's 48 bytes. Decoded members are searched for both HCPS and complete GLSLC outputs (effect asset packs); compressed literals are visited only after decoding. HCPS's SLCT/NVSD descriptors locate each variant's control and code ranges; a vertex bundle can include geometry, so stages are split by each control's own lengths. Text-bearing controls advance at eight-byte alignment. Standalone SLCT sections use offsets relative to their own start. Source-only HCPS containers with empty code/control ranges hold no NVN binaries. Pairs stay within their program variant. Synthetic fixtures cover wrapper bounds, embedded geometry, variants, standalone offsets and compressed streams containing literal container magics.
- **Nested effect compression:** Firehawk WIPAC resources can wrap their `BlShaderArchive` GLSLC programs in `cmp\0zlib`: an eight-byte signature, little-endian stored/raw u32 lengths, then zlib. This layer lives inside the decoded XBC1 entry. The recursive reader checks both lengths and the end of the stream; it excludes these compressed ranges from raw container scans too. The four initially missing geometry/pixel keys and two additional vertex/pixel keys from Xenoblade Chronicles 3's opening match `/effect/fld/ma01a.wipac` after this decoding. Synthetic tests cover nested literal streams and incorrect sizes.
- **Loose containers:** HCPS programs, XBC1 streams, DRSM model streams and compressed effect resources also occur outside ARH/ARD archives. Both the file scanner and nested-container scanner dispatch their headers to the same recursive decoder, before the generic shader-magic filter. Otherwise even a complete loose HCPS program is discarded. Synthetic file scans cover direct programs, compressed programs and several embedded streams, independent of filename suffix.
- **Found on:** [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md#files-and-shaders), whose note records the census and runtime key checks; [Xenoblade Chronicles 2](../../dm/games/0100E95004038000/README.md#files-and-shaders), with the same archive and standalone SLCT layouts on SDK 3.5.0; [Xenoblade Chronicles: Definitive Edition](../../dm/games/0100FF500E34A000/README.md#files-and-shaders), with loose containers.
- **Format references:** [ARH1 layout](https://github.com/roccodev/ard-tools/blob/master/docs/arh1.md), [xc3_lib shader structures](https://github.com/ScanMountGoat/xc3_lib/blob/main/xc3_lib/src/spch.rs). The reader's lengths and keys were checked against the supplied game files.

### Needle shader containers in main and GFD archives

- **Symptom:** the split finds no shaders, although `main` contains NVN control sections; it reports that no data record or code-address pair identifies their code areas.
- **Cause:** Needle keeps independent shader programs in `HNSHV002` containers, embedded in `main` and in GFD's GSC files inside CRI CPK archives. These are neither GLSLC outputs nor BNSH files. A container can hold many variants separated by reflection metadata.
- **Fix:** `shader_split.py` `needle_shaders` uses the container's big-endian whole length at +8, then the little-endian NVN section descriptors: 0x90 bytes immediately before each control section, with control offset, code offset, code length and control length at +0x30. Both regions must fit within that container, and the control/code headers and lengths must agree. Graphics stages come from SPH; headerless compute stages from the control's thread-group dimensions. Each program is paired with the single-stage cover, since its draw-time partner is not recorded by this wrapper. Synthetic fixtures check embedded variants, compute, length mismatches and container boundaries.
- **Found on:** [Persona 5 Royal](../../dm/games/01005CA01580E000/README.md#shaders), with 0x850-byte control sections. The game note records counts and runtime coverage.

### Shaders in a zlib stream in main (YEBIS .tfxo)

- **Symptom:** shaders the game draws are missing from the map (`DM_NVN_SHDUMP`), in no RomFS file and no module, stored or LZ4, and not compiled at run time.
- **Cause:** Silicon Studio's YEBIS post-effect library (bloom, glare, depth of field, vignette; `C:/workspace/YEBIS_draft/src/GPU/…` in `main`) keeps its shaders in a zlib stream in `main`: YEBIS's `.tfxo file`, a text header (the shader count, `hlsl`, `TexUtil_NVN_NX`, the technique list) and the shaders, each a control section (its own length 0x850 on SDK 7) directly followed by its code area.
- **Fix:** `shader_split.py` inflates it and reads the shaders as single-stage programs (`tfxo_streams`, `tfxo_shaders`).
- **Found on:** [Pokémon Sword](../../dm/games/0100ABF008968000/README.md): after the [GFPAK reader](#bnsh-entries-of-gfpak-archives), 38 shaders the opening drew were still missing; none was a patched copy of a split shader (the nearest ones of the same length differ in hundreds of bytes). `main` imports `glslcCompile` from `subsdk1` but never calls it and holds no GLSL source. The only zlib stream in `main` that inflates (`main+0x1c98ed9`, 39.5 MB) holds 3,155 shaders (3,015 distinct, 3,112 pixel and 43 vertex); they covered all 38.
- **Check:** a temporary trace in `nvnProgramSetShaders` (the key's code and control addresses, a backtrace and the heap around them) showed control sections packed back to back with technique names such as `ps_BrightPassT1_THRESHOLD_LUMINANCE_S1` beside them; the names are in `main`, and so is the data.

### Loose control sections and code areas in main (Bink 2)

- **Symptom:** two `[RED] shader … is not in the shader map` lines and `program has no vertex shader` as the opening movie starts; every movie drew white.
- **Cause:** the movie player (Bink 2's Unity plugin) builds its programs from 32 control sections and code areas in `main`'s read-only segment that are in no GLSLC output and no BNSH file. 26 are 0x20-byte records in two tables in `.data` (`{code area pointer, control section pointer, u32 code length, u32 control length}`, both pointers relative relocations); the code takes the other 6 by address (`adrp` + `add` to the control section and, a few instructions away, to the code area it copies with `memcpy`), the code area right after the control section's own length (+0x20, 16-byte aligned) or right before it, and two control sections share one code area. The 8 compute shaders' code areas have no header (SDK 10).
- **Fix:** the reader (`raw_shaders`) takes them. Both neighbours can pass the header and length checks, so it takes only a code area the code builds the address of, the nearest one to the control section's.
- **Found on:** [Advance Wars 1+2: Re-Boot Camp](../../dm/games/0100300012F2A000/README.md) (Unity 2020.2 with Bink 2's Unity plugin). Of 25 `main`s on hand no other has such control sections.

### SNVN program records in Retro Studios' RFRM forms

- **Symptom:** the provisional split finds only `main`'s 294 GLSLC outputs; the 3D scene and the title logo are magenta, with `[RED] shader … is not in the shader map` lines.
- **How it is stored:** Retro's files are RFRM forms (`RFRM`, big-endian u64 data length, u64, then the form type and two u32 versions; chunks inside are a fourcc, u64 length, two u32 versions, then the data). `MaterialArchive.arc` is one `MTRL` form holding an `SDX ` chunk (the same materials' D3D bytecode, 3,468 DXBC programs) and an `SNVN` chunk (the NVN ones). The `.pak` files are `PACK` forms whose `ADIR` chunk lists the resources (u32 count, then 36-byte entries: fourcc type, 16-byte GUID, big-endian u64 offset and length); each `MTRL` resource is an RFRM form whose data is one zlib stream holding another `MTRL` form with an `SNVN` chunk. An SNVN program record is a big-endian u32 code length, u32 control length, a byte, u32 stream length, then a zlib stream that inflates to the program's code areas back to back followed by their control sections in the same order (0x840 bytes each on SDK 4). The chunk's tables differ between files, so the reader finds the records by their own lengths: a stream whose header ends where it says and that inflates to exactly code + control bytes.
- **Startup archive in `main`:** the first programs the game builds (the `CGraphicsUtilCache` quads that draw the title logo) come from a startup material archive compressed into `main` (`NMaterialStartup::get_startup_archive_compressed_buffer`, 5,870 bytes; `NMaterialArchive::ImportStartupArchive` inflates it at boot). Its `CBufferedCompression` header is a type byte (1: LZSS on 8-bit units, `CLZSSDecompressor::DecompressBuffer8`) and three zero bytes; the data inflates to an `MTRL` form like `MaterialArchive.arc`'s. Its programs share code areas with `MaterialArchive.arc`'s but differ in control +0x7d0 (a per-program value), so their keys are their own.
- **Little-endian forms (Metroid Prime Remastered):** the same forms with every length, count and record header little-endian (`rfrm_order` tells the order from the u64 length: the reading that is a length is the small one). A pack's first chunk is a `TOCC` form holding the `ADIR` chunk, whose entries are 52 bytes: fourcc type, GUID, two u32 versions, u64 offset, u64 length and u64 stored length (shorter for a compressed resource; none of the 778 `MTRL` resources is). An `MTRL` resource's zlib stream inflates to an `EXPT` chunk and then the `MTRL` form. Control sections are SDK 12's 0x880 bytes, and four vertex stages of `MaterialArchive.arc` carry 0xc0 bytes more than their own length (+0x20) says: GLSLC's extra data at +0x7e8 (offset 0x880) / +0x7ec (length) / +0x7f0 (16 entries of 12 bytes), which the reader steps over to reach the next stage's section.
- **Reader:** `rfrm_shaders`, `snvn_records`, `retro_main_forms` and `retro_lzss8` in `shader_split.py`; each record is one program, so its stages pair with each other.
- **Found on:** [Donkey Kong Country: Tropical Freeze](../../dm/games/0100C1F0051B6000/README.md) (SDK 4.4.0): 24,098 shaders (11,644 vertex, 11,734 pixel, 720 geometry, which run as compute kernels): 3,460 first seen in `MaterialArchive.arc` (its 3,468 stages), the rest in the 795 `MTRL` resources of the packs, 8 in the startup archive and 294 GLSLC outputs in `main`; about 8 s.
- **Also found on:** [Metroid Prime Remastered](../../dm/games/010012101468C000/README.md) (SDK 12.3.1, little-endian): 40,483 shaders (20,233 vertex, 20,233 pixel, 12 geometry, 5 compute): 1,141 stages in `MaterialArchive.arc`, the rest in the 778 `MTRL` resources of the world, preload and UI packs, and 8 in the startup archive (`main+0x1d0e7a4`, the same `CBufferedCompression` type 1); `main` holds no GLSLC output. Tropical Freeze's split gives the same 80,212 stages from its 76 RFRM files with the byte-order-aware reader.

### MGFX effects in MonoGame XNB files

- **How it is stored:** the 13 XNB files in `Content/Effects` (XNB platform `S`, uncompressed, read by `EffectReader`) each hold one MGFX effect, and `main` holds SpriteBatch's built-in one; no other file and no GLSLC output. A compressed XNB (MonoGame's LZX flag 0x80 or LZ4 flag 0x40, as in Stardew Valley) is decompressed first (`xnb_decompress`).
- **Layout:** MGFX version 8 is MonoGame's own layout: `MGFX`, u8 version, u8 profile, u32 effect key, the constant buffers, then each shader as u8 vertex (1) / pixel (0), u32 length, the bytecode (u32 control length, the control section, u32 code length, the code area: GLSLC's output without its header), the samplers, the constant buffer indices and an attribute list (as MonoGame's OpenGL backend writes it), then the parameters and the techniques with their passes (`mgfx_effect`).
- **Reader:** the reader parses every effect to its end, which checks the layout; in `main`, four bytes `MGFX` count only when version 8 and a control section follow.
- **Pairing:** a pass names its vertex and pixel shader by index, 255 for none, and 10 of the 13 effects have no vertex shader: SpriteBatch binds its own and the effect's pass only replaces the pixel shader. Vertex and pixel shaders therefore combine across effects, and they are paired as single-stage programs.
- **Found on:** [Celeste](../../dm/games/01002B30028F6000/README.md) (MonoGame for the Switch, its C# compiled to C++ by Sickhead's BRUTE): 26 shaders (20 pixel, 6 vertex), all with Mario's binding layout; every binding decodes. [Stardew Valley](../../dm/games/0100E65002BB8000/README.md) (compressed XNB).

### MGFX version 5 effects (MonoGame 3.0)

- **Symptom:** the NSP script's split stops with `[RED] 24 file(s) could not be read`, each `an XNB effect without an MGFX version 8 effect of NVN shaders`.
- **Cause:** MonoGame 3.0's EffectReader (`MonoGame.Framework, Version=3.0.0.0` in the XNB) writes MGFX version 5: `MGFX`, u8 version, u8 profile and no effect key; constant buffers without names (i16 size, u8 count, {u8 parameter, u16 offset}); each shader's bytecode is a whole GLSLC output of one stage (magic 0x19866891, its +0x44 length equal to the shader's), not a control section and a code area; and the shaders are in MonoGame's DirectX layout: samplers without names and no attribute list. The parameters and techniques are version 8's. The length in front of `MGFX` in the XNB runs about 0xff0 bytes past the effect, all zeros.
- **Fix:** `mgfx_effect` reads version 5 too, taking each shader through `glslc_shaders`; an XNB effect may end in zeros. The 7 `.dx11.mgfxo` files in `Resources` (MonoGame's stock effects) are version 5 as well: 5 hold GLSLC outputs and are read the same way, `EnvironmentMapEffect` and `SkinnedEffect` hold only DirectX bytecode and give no shader.
- **Found on:** [Bastion](../../dm/games/010038600B27E000/README.md): 72 shaders, 24 pixel shaders in `Content/Effects` (no vertex shader: SpriteBatch's own pairs with them) and 48 in `Resources`.

### Unity 2020.3 Shader objects in UnityFS bundles

- **How it is stored:** Unity 2020.3 keeps each shader's compiled programs in its `Shader` object (class 48), in UnityFS bundles (`StreamingAssets/aa`) and in the player's own serialized files (`data.unity3d`, `Resources/unity default resources`).
- **Reader:** `dm/gfx/tools/unityfs.py` reads them: UnityFS blocks (LZ4 / LZMA), the serialized file, the object through its type tree, then `compressedBlob` per platform (LZ4 segments; each segment starts with a u32 count and 12-byte `{offset, length, segment}` entries).
- **Type trees:** the player's files have no type trees; they borrow the bundles' `Shader` tree with the same type hash (same Unity build, same layout). Type-tree strings with the high bit set index Unity's common string table, which `unityfs.common_strings` reads from the game's own `main` (`AABB\0AnimationClip\0…`).
- **Program code:** on Switch every subprogram is one whole NVN program. Its code is a 0x60-byte header of four `i32 ×6` arrays indexed by NVN stage (0 vertex, 1 fragment, 2 geometry, 3 tess control, 4 tess evaluation, 5 compute): an unused / scratch size, the block offset from +0x60 (−1 = no stage), the control size and the code size. Each block is the control section followed by the code area. The vertex and fragment stages of one program form a pair.
- **Found on:** [Fire Emblem Engage](../../dm/games/0100A6301214E000/README.md): 209 `Shader` objects in 188 bundles plus the engine files, 9,269 stage shaders (4,634 vertex / pixel pairs and one geometry shader), no compute shaders.

### Legacy Unity stages lose their pass partners

- **Symptom:** `failed to create render pipeline` with a fragment `user(locn1)` input missing from the vertex output, while the split found all shaders but almost no vertex / pixel pairs.
- **Cause:** before Unity 2019.3 each blob entry contains one stage. The vertex and pixel variants of a parsed pass have different blob indices, so pairing by the source's `#var<platform>.<index>` connects none of them. Without its pixel partner, output padding does not know which unwritten vertex attributes the fragment shader reads. Return of the Obra Dinn's `OneBit/Rope` has vertex blob 0 with pixel blob 8 and vertex blob 28 with pixel blob 35; both fragment stages read generic locations 0–4.
- **Fix:** `unity_pass_pairs` in `shader_split.py` reads the stages listed by each parsed pass, separately per platform, and pairs its vertex stages with a deterministic pixel-input cover in the supplemental `pso_pairs.tsv`. This covers every input used by the pass's pixel variants without depending on equal stage keyword lists or generating every combination. Existing evidence-pair tables, combined-stage Unity programs and source labels stay unchanged; profiles without a stored supplemental table take the generated one beside the shader package.
- **Found on:** [Return of the Obra Dinn](../../dm/games/010032E00E6E2000/README.md) (Unity 2017.4.28f1, 1.0.1): the two `OneBit/Rope` pipeline failures reported in PR #427.

### Unity 2022 Shader objects with parameter blobs

- **How it is stored:** the program code keeps the 0x60-byte stage header ([Unity 2020.3 Shader objects in UnityFS bundles](#unity-20203-shader-objects-in-unityfs-bundles)), but two things changed. From Unity 2022 the blob also holds each program's parameter tables as entries of their own (the parsed form's `m_ParameterBlobIndices`; their bytes start `PGlobals`), and the player's programs are listed in `m_PlayerSubPrograms[tier][].m_BlobIndex` (`m_SubPrograms` is empty).
- **Reader:** `unity_code_blobs` keeps only those player programs. UnityPy's type trees for 2020.3 (Engage) and 2021.3 have neither field, so those read every entry as before. Subprogram version 202012090 (from Unity 2021.2) also has no local keyword list, so `unityfs.subprogram` reads no second keyword array (with one it read the code length as a keyword count).
- **Type trees:** the player's serialized files have no type trees and the Addressables bundles hold no `Shader`, except the `unitybuiltinshaders` bundle, whose tree has the same hash and is borrowed in the second pass.
- **Found on:** [Ace Attorney Investigations Collection](../../dm/games/010005501E68C000/README.md) (Unity 2022.3.17f1): 4,951 stage shaders: 2,475 vertex / pixel pairs (4,294 in the player's files, 517 built-ins, 140 in bundles) and one geometry shader, which Metal cannot express (its build fails).
- **Note:** armyja's issue #17 proposed UnityPy and a reader of its own; neither was needed.

### Unity Shader layouts written out for players without bundles

- **Symptom:** the split stopped with `no bundle has a Shader type tree of the same hash to borrow`.
- **Cause:** a Unity player with no bundle to borrow a tree from: its serialized files have no type trees, and no bundle carries a `Shader` tree of the same hash.
- **Fix:** `unityfs.py` carries the layouts written out (`KNOWN_TREES`; `tree_from_layout` builds the tree a file with type trees would have), and the split's second pass gives one to the stripped files when no bundle lent a tree of that hash. Only there: taken in the first pass, Ace Attorney Investigations Collection's player files came before its `unitybuiltinshaders` bundle and the built-ins' source column and package indices moved, while with the second pass its tables come out byte-identical. The structs both the 2020.3 and 2022.3 layouts share are written once (`_SHADER_COMMON`), and the 2022.3 layout comes out byte-identical.
- **Found on:**
  - [Brain Memory Training Academy](../../dm/games/0100F5D01DA3A000/README.md) (Unity 2022.3.5f1: only `Data/*.assets`, `level*` and the two engine files in `Data/Resources`, none with type trees). The `Shader` type hash there (`ea424a5f…`) is the one in Ace Attorney Investigations Collection's bundles (2022.3.17f1): the same layout, about 30 structs. Checked against Ace Attorney's own tree: the same 2,324 nodes with the same alignment flags, and its 40 bundle `Shader` objects decode identically with either. 496 stage shaders: 439 engine built-ins, 46 of the game's, 11 of the player's in `main`, and the built-in VR blit's geometry shader, which does not build.
  - [Patrick's Parabox](../../dm/games/0100E44016C22000/README.md) (Unity 2020.3.30f1, built-in shaders only, no bundles) stopped there with hash `a30e3dcf…`, which one 18 KB shader bundle of Fire Emblem Engage (2020.3.18f1) carries; its tree is now in the table too (`SHADER_2020_3`: per subprogram global and local keyword indices, every subprogram in `m_SubPrograms`, no parameter blobs, keyword names or stage counts), equal to Engage's node for node (2,179 nodes, the alignment flags included), and each of the game's 48 `Shader` objects reads to its exact end.
  - [Raging Loop](../../dm/games/010020400E1C2000/README.md) (Unity 2018.3.10f1, no bundles: 48 `Shader` objects in `Data/*.assets` and `Data/Resources`) stopped there with hash `8db76007…`, and no game of that Unity version with a bundle was at hand, so `SHADER_2018_3` is written from AssetRipper's TypeTreeDumps (`StructsDump/release/2018.3.10f1.dump`), checked node for node against the dump (1,325 nodes, the alignment flags included). It is a layout of its own, older than the common structs: no conservative raster state, `BufferBinding` without `m_ArraySize`, `ConstantBuffer` without `m_IsPartialCB`, the parameters in each subprogram and no `m_CommonParameters`, one segment per platform (the pre-2019.3 form Cuphead's trees gave `shader_programs`). Each of the 48 objects reads to its exact end: 450 subprograms of version 201802150 (older than the local keyword list, which starts at 201806140).
  - [Return of the Obra Dinn](../../dm/games/010032E00E6E2000/README.md) (Unity 2017.4.28f1): hash `266d5311…`. `SHADER_2017_4` derives from the 2018.3 layout by removing `m_HasProceduralInstancingVariant` and `m_NonModifiableTextures`, with the instancing flag retaining the alignment. It matches all 1,313 nodes and alignment flags in [AssetRipper's release type tree](https://github.com/AssetRipper/TypeTreeDumps/blob/main/StructsDump/release/2017.4.28f1.dump); the game note records the object census and exact-end checks.
  - [Donkey Kong Country Returns HD](../../dm/games/01009D901BC56000/README.md) (Unity 6000.0.28f1, no bundles): hash `2ec7223d…` matches BALL x PIT's Unity 6000.0.62f1 bundles. `SHADER_6000_0` derives from the 2022.3 layout by removing each pass's `m_EditorDataHash` and `m_Platforms` arrays and appending the shader's `m_AssetGUID`. The generated tree equals the bundle tree in types, field order and alignment flags; the game note records the object census and build results.
  - LEGO Builder's Journey (Unity 2021.2.13f1, provisional profile: every serialized file is stripped and no bundle carries a `Shader` tree): hash `d8236c9b…`. `SHADER_2021_2` is written from [AssetRipper's 2021.2.13f1 release tree](https://github.com/AssetRipper/TypeTreeDumps/blob/main/StructsDump/release/2021.2.13f1.dump) and matches it node for node (2,174 nodes, the alignment flags included): its subprograms drop the local keyword list (as Unity 2022 does), keep their parameters and use signed 64-bit requirements; each pass appends `m_SerializedKeywordStateMask`, and the parsed form adds keyword names and flags. All 195 `Shader` objects read to their exact ends, and the split finds 9,410 stage shaders (3,567 vertex, 3,571 pixel, 2,271 compute, one geometry) in the player's files and `main`.
  - Dodgeball Academia (Unity 2019.4.21f1, provisional profile: every serialized file is stripped and no bundle carries a `Shader` tree): hash `ed2a06ce…`. `SHADER_2019_4` derives from the 2018.3 layout, checked node for node against [AssetRipper's 2019.4.21f1 release tree](https://github.com/AssetRipper/TypeTreeDumps/blob/main/StructsDump/release/2019.4.21f1.dump) (1,481 nodes, the alignment flags included) and equal to the tree [Subnautica](../../dm/games/0100429011144000/README.md)'s Unity 2019.4.36f1 bundles carry under the same hash: global and local keyword indices in each subprogram, a ray tracing program in each pass, and 2019.3's segments (one `offsets` / `compressedLengths` / `decompressedLengths` list per platform); the parameters stay in each subprogram, without `m_CommonParameters`. All 62 `Shader` objects read to their exact ends, and the split finds 7,996 stage shaders (4,162 vertex, 3,778 pixel, 55 compute, one geometry) where 11 came from `main` alone.
- **Note:** a Unity version whose `Shader` hash is not in the table still stops with the same message; read its tree from any bundle of that Unity version, or from AssetRipper's TypeTreeDumps for that exact version (compared node for node, the alignment flags included), and add it with its hash to `KNOWN_TREES` and to `dm/gfx/tests/test_unity_known_trees.py`. To take one file out of an NSP without extracting the RomFS, filter `nx_extract`'s `Copy` by path.

### Unity 6 Shader objects

- Read with no change: Hollow Knight: Silksong's 2,068 Addressables bundles carry their type trees (they report the engine version `0.0.0`, Unity's stripped version string), so no layout had to be written out.
- BALL x PIT keeps most of its shaders in the player's loose serialized files (`resources.assets`, `sharedassets1.assets`, the built-in resources), which have no type trees; they borrow the `Shader` tree its 10 shader bundles carry (`shader_split.py` defers a file that raises `NoTree` to a second pass), so it needed no layout either.
- **Found on:** [Hollow Knight: Silksong](../../dm/games/010013C00E930000/README.md) (Unity 6000.0.50f1): 2,450 stage shaders: 1,189 vertex, 1,193 pixel, 67 compute and one geometry shader, which does not build; 11 of them are the player's own in `main`. [BALL x PIT](../../dm/games/010086A022444000/README.md) (6000.0.62f1): 38,099 shaders, most of them in the loose files. [Silver Pines](../../dm/games/0100D5B026844000/README.md) (6000.3.15f1): 56,386 stage shaders in its Addressables bundles, read with no change.

### Unity ComputeShader objects

- **Symptom:** the 3D battle scene was black behind a correct UI, with `[RED] shader 854a8afc9d07f88d is not in the shader map` and `the dispatched program has no compute shader`.
- **Cause:** `DM_NVN_SHDUMP` saved the program; its code bytes were in `Data/sharedassets2.assets`, uncompressed, inside a `ComputeShader` (class 72): the post-processing stack's colour grading LUT baker (`Lut3DBaker#KGenLut3D_AcesTonemap`), so the LUT the scene is graded through stayed black. The split read only `Shader` objects. A `ComputeShader` holds per platform variant a list of kernels, each with its unique variants; a variant's `code` is the program code of a `Shader` subprogram (the 0x60-byte stage header, only the compute block present) and its `threadGroupSize` matches control +0x748.
- **Fix:** `unity_compute` in `shader_split.py` reads them beside the `Shader` objects of every serialized file. The game has no bundle with a class 72 tree, so `unityfs.py` carries Unity 2022.3.54f1's layout (`COMPUTE_2022_3`, written from AssetRipper's TypeTreeDumps for that version; all 12 objects decode, `resourcesResolved` true at each end). 85 compute shaders came in (86 with the player's own).
- **Found on:** [Monster Train 2](../../dm/games/010051701FB46000/README.md) (Unity 2022.3.54f1).
- **Older layouts:** Unity 2019.4 stores each kernel's code directly, without `uniqueVariants` or `variantMap`. `unity_compute` accepts that form too, and `COMPUTE_2019_4` supplies the stripped-file layout. Found on [Pokémon Brilliant Diamond](../../dm/games/0100000011D90000/README.md) (2019.4.27f1); its missing runtime shader is the built-in GPU skinning kernel. The layout matches the release dump's 116 nodes and alignment flags, and both built-in compute objects decode to their exact ends.
- **Unity 2021.3:** `COMPUTE_2021_3` covers stripped class 72 hash `abd9135b8ce83d043fef4e9ec7f53366`, matching all 144 nodes and alignment flags in [AssetRipper's 2021.3.18f1 release tree](https://github.com/AssetRipper/TypeTreeDumps/blob/main/StructsDump/release/2021.3.18f1.dump). Kernel parents use `variantMap`, global and local keywords; each kernel holds byte-vector code, a `staticvector` thread group size and signed 64-bit requirements. Unlike 2022.3 there are no `uniqueVariants`, `variantIndices` or dynamic keywords; both versions store code as bytes. `resourcesResolved` aligns after every platform variant, including when another platform follows. The synthetic tests cover two platforms, multiple keyword variants, byte-code padding, exact object consumption and shader provenance. Found on Shin chan: Shiro and the Coal Town (Unity 2021.3.18f1): [PR #440](https://github.com/seedds/nxnative/pull/440) reports recovered GPU skinning kernels and correctly rendered characters on the iPhone; that gameplay result has not been independently repeated here.
- **Unity 2020.3:** `COMPUTE_2020_3` covers stripped class 72 hash `d0171b849bb064b2424893e7d4809be5`. It is the 2021.3 layout without the kernel's trailing `SInt64 requirements`: kernel parents use `variantMap` (not aligned) with global and local keywords, code is a byte vector and `threadGroupSize` a staticvector (not aligned). All 143 nodes and alignment flags match [AssetRipper's 2020.3.27f1 release tree](https://github.com/AssetRipper/TypeTreeDumps/blob/main/StructsDump/release/2020.3.27f1.dump); the synthetic tests cover the 2021.3 fixture without `requirements`. Found on MADiSON (`0100C3F0181C8000`, 1.1, Unity 2020.3.27f1, SDK 13, provisional profile): `unity default resources` and `sharedassets0.assets` skipped their ComputeShader objects, the split had 1 compute shader, and the run logged 6 `shader … is not in the shader map` and one `the dispatched program has no compute shader`. The 3D scene was black on an iPhone 16 while the crosshair, menus and sound worked. With the layout the split takes 91 compute shaders (8,795 in all), the Mac run logs no red line and the scene draws; those results are the author's, and the fixed build has not run on the iPhone. The hash is not MADiSON's alone: Unity's type hash (MD4 over each node's type, name, byte size, array flag, version and align flag) recomputed from the release trees of 2020.3.0f1, 18f1, 27f1, 30f1 and 48f1 gives it each time, and the same formula gives 2021.3.18f1's `abd9135b…`. So the layout can also reach the registered 2020.3 games. [AI: THE SOMNIUM FILES - nirvanA Initiative](../../dm/games/0100713016CCC000/README.md) already has its compute kernels, since its bundles carry type trees. The tables of [Fire Emblem Engage](../../dm/games/0100A6301214E000/README.md) and [Patrick's Parabox](../../dm/games/0100E44016C22000/README.md) have no Unity compute kernels: if their stripped files hold ComputeShader objects, their next NSP pack adds them through the shader-only derived profile (below). Neither game's files were available to check this.
- **Unity 2017.4:** [Return of the Obra Dinn](../../dm/games/010032E00E6E2000/README.md) needs hash `20291e02…` for stripped `internalskinning` objects. `COMPUTE_2017_4` restores the nested buffer-counter fields; all 120 nodes and alignment flags match the release dump. The game note records the object census and exact-end checks.
- **Unity 6000.3:** `COMPUTE_6000_3` covers stripped class 72 hash `43a0948a3f403068c0b31bff90ba0c9a`: the 2022.3 layout with an `int resType` after each resource's `texDimension` and `threadGroupSize` a `staticvector` again. All 159 nodes and alignment flags match [AssetRipper's 6000.3.11f1 release tree](https://github.com/AssetRipper/TypeTreeDumps/blob/main/StructsDump/release/6000.3.11f1.dump) and its 6000.3.0b4, 6000.3.0f1 and 6000.3.20f1 trees; a synthetic test reads a fixture to its exact end and into the split. Found on Hello Kitty Party Land (Unity 6000.3.11f1, provisional profile): its characters drew as huge stretched triangles, with `[RED] … unity default resources: ComputeShader objects that no bundle lends a type tree to and unityfs.py has no layout for` in the split and `the dispatched program has no compute shader` in the run, since the built-in `Internal-Skinning` kernels were left out. [PR #857](https://github.com/seedds/nxnative/pull/857) reports that all 11 objects read to their exact ends (`Internal-Skinning` 40 kernels, `Internal-BlendShape` 9), the split finds 92 compute shaders instead of 1, and the Mac app draws the characters right; [PR #791](https://github.com/seedds/nxnative/pull/791) reports the same hash on another game. Neither result has been repeated here, since no game with this hash is available. [Silver Pines](../../dm/games/0100D5B026844000/README.md) (6000.3.15f1) already has these kernels in its table without this layout.
- **Registered tables after a reader update:** a newly supported layout can add kernels to an existing game's shader map. For registered module builds, the NSP script permits a shader-only derived profile only when every original shader's package, stage and source, every original fact and every evidence pair remain unchanged (package indices may move), and every added shader is compute with matching key/source evidence from `unity_compute`. Each added kernel must also have its fact row. Registered settings, build rows and cuts are retained; the repository tables are not edited. A cached comparison pins the registered tables, split inputs and comparison code. Deleted or changed original rows, graphics additions and additions without reader evidence still stop the build. The incremental-build tests cover acceptance, rejection and cache invalidation.
- **Note:** a file whose `ComputeShader` has no tree still goes to the second pass; if it has none there either, the split prints a `[RED]` line and leaves its kernels out instead of stopping. The original bring-up omitted stripped objects in [Cuphead](../../dm/games/0100A5C00D162000/README.md) (2), [Fire Emblem Engage](../../dm/games/0100A6301214E000/README.md) (4) and [Raging Loop](../../dm/games/010020400E1C2000/README.md) (2 built-ins; no compute dispatched). Their type hashes were not recorded, so their compatibility with the newer layouts remains unchecked. Unity 2017.4.9f1, 2017.4.17f1 and 2017.4.28f1's published class 72 release dumps are identical, but that does not establish the hashes in Cuphead or AI: THE SOMNIUM FILES.

### GLSLC outputs in a Unity TextAsset

- **Symptom:** black, jagged holes in the ground of a map (whole tiles of the terrain mesh missing while others draw), with about 25 distinct `[RED] shader … is not in the shader map` keys and `program has no vertex shader`, on a Unity game whose `Shader` objects all split cleanly.
- **Cause:** the game draws its maps and effects with programs of its own, not Unity's: `DM_NVN_SHDUMP` saved 24 of them, and none was in any `Shader` subprogram of the 1,606 `Shader` objects in the RomFS (all of them decoded); their code areas were in `Data/resources.assets`, uncompressed, inside the `TextAsset` (class 49) `shader.switch.vfxsb` (11 MB, first word 0x20160817): 1,160 whole GLSLC outputs of a vertex and a pixel stage each. Their control sections carry GLSLC's full header (control +0x20 = 0xa10 and more), unlike the 0x850-byte sections Unity 2018 stores in its subprograms. The split read only `Shader` and `ComputeShader` objects, and the fallback for loose GLSLC outputs skips Unity files.
- **Fix:** `unity_text_glslc` in `shader_split.py` reads the `TextAsset`s of a serialized file that holds the GLSLC magic (`unityfs.text_assets`: name and bytes, the two leading fields every Unity version shares, so no type tree is needed) through `glslc_shaders`, which keeps only outputs that pass its checks. A file without the magic is not decoded, so other games' tables do not change. This one TextAsset added 1,207 distinct shaders (the game has 31,635 in all), with 1,156 vertex / pixel pairs; they covered all 57 keys the runs had reported missing.
- **Found on:** [FINAL FANTASY CRYSTAL CHRONICLES Remastered Edition](../../dm/games/0100CE4010AAC000/README.md) (Unity 2018.4.2f1): the intro's overworld shot.

### Unreal Engine 5 shader code archives (Oodle)

- **How it is stored:** `ShaderArchive-Global-SF_GLSL_SWITCH_MOBILE-SP_GLSL_SWITCH_MOBILE.ushaderbytecode` (1,219 shaders) and `ShaderArchive-Octopath_Traveler0_Chunk0-…` (9,377), both stored in the pak ([01](../agents/01-extract.md)). Version 2 archives of shader map hashes, shader hashes, shader map entries, shader entries `{u64 offset, u32 size, u32 raw size, u8 frequency}` (17 bytes packed; frequency 0 vertex, 3 pixel, 5 compute), preload entries and shader indices, then the code.
- **Compression:** every shader is Oodle-compressed (Mermaid, block header `8C 0A`). There is no Oodle decoder in this repository, so the game's own `OodleLZ_Decompress` (statically linked into `main`; the host finds it as the one function that takes the address of the library's own name string "OodleLZ_Decompress", its profiler scope: Octopath Traveler 0 `main+0x4f77e10`, Pokemon Legends: Z-A `main+0x1a9c9e0`) decodes them in the Mac host: `DM_OODLE_JOB` (`dm/host/hle_oodle.c`), the NSP script's `decode_ue_shaders` for any game with Unreal paks.
- **Record:** a decoded shader is an `LSLG` record: `LSLG`, a stage tag (`SV`, `SP`, `SC`), a header with the source name (`/Engine/Private/…usf:MainPS`), then `u32` code length, the code area, `u32` control length (0x880), the control section and a 24- or 29-byte tail, all at unaligned offsets (`uepak.lslg`).
- **Check:** the host keys came out exactly: 37 of 37 shaders a first run reported missing were in the two archives.
- **Found on:** [Octopath Traveler 0](../../dm/games/01005270232F2000/README.md): 10,324 unique shaders (1,457 vertex, 6,346 pixel, 2,521 compute); vertex and pixel shaders combine per draw, so every vertex shader is paired with the single-stage cover. Bindings: separate textures and samplers, as Engage's (`0x168` / `0x568`); 61 vertex shaders pick a texture slot at run time (profile key `tex_dyn_slots`, [Texture slots picked at run time](#texture-slots-picked-at-run-time)), 2,084 shaders do atomics on buffer images ([Image atomics](#image-atomics)).
- **Note:** 16 compute shaders do atomics on a 3D `r32ui` image, all Unreal's paged global distance field (`GlobalDistanceFieldCompositeObjects.usf:CompositeObjectsIntoPagesCS`), which the mobile renderer does not dispatch (none in any run). They build since image atomics are native ([Image atomics](#image-atomics)).

### Unreal Engine 4.27 shader code archives (LZ4)

- **How it is stored:** the desktop renderer's `ShaderArchive-Global-SF_GLSL_SWITCH.ushaderbytecode` (1,764 shaders) and `ShaderArchive-DOLL_Chunk0-SF_GLSL_SWITCH.ushaderbytecode` (63,185), stored in the pak: the same version 2 archives and the same `LSLG` records as [Unreal Engine 5's](#unreal-engine-5-shader-code-archives-oodle), but every shader is an LZ4 block (the first bytes are LZ4 tokens followed by `LSLG`; LZ4 is Unreal's default shader compression since 4.26).
- **Reader:** `uepak.py` decodes them itself (`lz4_shaders`): every compressed shader must decode to its raw size as an `LSLG` record, else the archive goes to the Oodle job as before.
- **Found on:** [DRAGON QUEST VII Reimagined](../../dm/games/0100A9D01C446000/README.md) (Unreal Engine 4.27): 62,901 unique shaders (10,114 vertex, 50,756 pixel, 1,968 compute, 63 geometry).
- **Note:** its fragment shaders carry no output types either: one global shader writes integers bit-cast into a float output to an `R8Uint` target, listed in the profile's `shader_facts_manual.tsv` ([Integer fragment outputs declared by hand](#integer-fragment-outputs-declared-by-hand)).

### Unreal Engine 5.6 IoStore shader libraries

- **How it is stored:** the desktop renderer (`SF_GLSL_SWITCH`) with the same `LSLG` records, but no archive in the pak: the two libraries are IoStore shader libraries in `FFRS-Switch.ucas`, whose shader groups are Oodle-compressed as a whole ([01](../agents/01-extract.md)). Each shader entry packs its frequency (bits 0–3), group (4–33) and offset in the group's decoded bytes (from bit 34), and a shader runs to the next one's offset.
- **Reader:** `uepak.py` writes them out as version 2 archives, so the split reads them as any other.
- **Found on:** [FINAL FANTASY RESONANCE](../../dm/games/0100606026F12000/README.md) (Unreal Engine 5.6.1): 73,496 unique shaders (9,253 vertex, 61,305 pixel, 2,934 compute, 4 geometry).

### IoStore patch containers keep their shader groups in the base container

- **Symptom:** `ValueError: …/Glow-Switch_0_P.utoc: shader group 0 has no chunk` from `uepak.py io_archive` while the NSP script decodes the shaders (`decode_ue_shaders`, after `[4/8] UE shader decode`).
- **Cause:** an update's IoStore container is a patch: its shader libraries have the base game's chunk IDs (so they replace their tables), and an unchanged shader group is listed with the base container's chunk ID. `io_archive` looked every group up in the container the library chunk came from.
- **Fix:** `uepak.py` reads every `.utoc` under the RomFS (`io_tocs`, in mount order: the base game's containers, then `_P` patches and numbered `_0_P`, `_1_P`, … patches, `mount_order`, case-insensitive `_P` marker) and resolves a group's chunk ID in any of them; a library chunk ID two containers share is taken from the later (patch) one, whose tables are the current set. TOC metadata is read without a key; `Toc.read` refuses encrypted data without a key only when the selected library or group needs it, so an encrypted group-only container is refused and a fully superseded encrypted library is never read.
- **Found on:** LEGO Horizon Adventures (update v262144, Unreal Engine 5; provisional profile): the update's `Glow/Content/Paks/Glow-Switch_0_P.utoc` over the base's `Glow-Switch.utoc`; 20,889 shaders in its two libraries, 20,880 unique after the split.

### An Unreal pak key split into two 16-byte literals

- **Symptom:** `[RED] …/Nicola-Switch_0_P.pak: the index is encrypted and no window of main decrypts it` from the NSP script before `[5/8] UE shader decode` (the message now reads `no key in main (a 32-byte window or two 16-byte literals its code stores together) decrypts it`).
- **Cause:** `uepak.py find_key` tried every 32-byte window of `main`'s read-only and data segments. The engine gets the key from a callback that copies a 32-byte array into its argument (`UE_REGISTER_ENCRYPTION_KEY`), and clang compiled the copy as two 16-byte constants loaded from separate places: `adrp`, `adrp`, `ldr q0`, `ldr q1`, `stp q0, q1, [x0]`, `ret`, so the key's halves are not side by side and no window holds it.
- **Fix:** `find_key` first tries the keys the code builds that way (`code_keys`: each `stp qA, qB, [x0]` with both registers loaded by `ldr q` from an `adrp` page within the 8 instructions before it, a register written by anything else in between dropped), then the windows as before. The right key is the one that decrypts the index start to a short mount point, as for a window. Covered by `dm/tools/tests/test_uepak_key.py` (a synthetic `main` and pak, no game files).
- **Found on:** [DRAGON QUEST III HD-2D Remake](../../dm/games/01003E601E324000/README.md) (Unreal Engine 4.27): 8 candidate pairs in `main`, the key found in about 2 s.

### Unreal Engine 4.18 shaders inline in cooked packages

- **How it is stored:** legacy version 4 paks with LZ5-compressed blocks ([01](../agents/01-extract.md)) and no shader code archive. Each material's shader map is serialized in its `.uexp` (38 to 87 shaders in the ones looked at) and the engine's global shaders in `Engine/GlobalShaderCache-GLSL_SWITCH.bin` (945: 597 pixel, 104 vertex, 241 compute, 3 geometry). Each shader's code is a byte array (`i32` length, then the bytes) holding one zlib stream that inflates to the same `LSLG` record as Unreal Engine 5's.
- **Reader:** `shader_split.py` (`ue_legacy_jobs`, `ue_inline_shaders`) decodes every file of the newest copy of each path with the game's own LZ5 and takes each zlib stream that inflates whole from exactly the length in front of it and starts with `LSLG` (every such stream in the files looked at). Records stored uncompressed are taken too, but none were seen.
- **Found on:** [DRAGON QUEST XI S](../../dm/games/01006C300E9F0000/README.md): 38,892 unique shaders (4,654 vertex, 33,981 pixel, 234 compute, 23 geometry) from 1,740 files; vertex and pixel shaders are paired as single-stage programs, as for Octopath Traveler 0.

### Unreal Engine 4.27 shaders inline in cooked packages (LZ4)

- **Symptom:** `the RomFS has Unreal paks, but none … holds a shader code archive`, then `the shader split found no shaders in main or the RomFS`.
- **How it is stored:** a version 11 pak (Zlib blocks, index not encrypted) with no `.ushaderbytecode`: the project turned `bShareMaterialShaderCode` off, so each material keeps its shaders in its own `.uexp` and the engine's global shaders are in `Engine/GlobalShaderCache-SF_GLSL_SWITCH.bin`, as in [Unreal Engine 4.18](#unreal-engine-418-shaders-inline-in-cooked-packages). Unreal 4.27 serializes each one as an entry of `FShaderMapResourceCode`: the code (`i32` length, then an LZ4 block, the engine's shader compression since 4.26), the raw size (`i32`) and the frequency (`u8`).
- **Reader:** `uepak.py` `inline_shaders` tells such a pak from one with archives by its directory index (cooked packages and no `.ushaderbytecode`; an encrypted index is taken for one with archives), and `shader_split.py` `ue_legacy_jobs` reads it file by file, Zlib and LZ4 blocks decoded in Python (`Pak.read`), so no host runs. `ue_inline_shaders` finds each LZ4 block by its first literals, the record's `LSLG`, with the token and literal length bytes in front of it, and takes it when it decodes to exactly the raw size after the block.
- **Found on:** [OCTOPATH TRAVELER II](../../dm/games/0100A3501946E000/README.md): 13,980 unique shaders (1,910 vertex, 11,346 pixel, 694 compute, 30 geometry) from the 3.3 GB pak in 13 s.

### Unreal Engine 3 shader caches (Oodle packages)

- **Symptom:** `the shader split found no shaders in main or the RomFS`; the survey lists the RomFS as cooked packages (`.xxx`) and `.usm` movies only.
- **How it is stored:** the engine's global shaders are in `GlobalShaderCache-Switch.bin` (magic `BMSG`, not compressed), the materials' in `CombinedShaderCache-Switch.xxx`, an Unreal Engine 3 package (tag `0x9E2A83C1`, a positive file version after it: 864; Unreal Engine 4's is a negative legacy version). After the package summary come compressed chunks, each {tag, block size, compressed and raw size of the chunk, then each block's two sizes} and the blocks: 0x20000 bytes of Oodle Mermaid each (`8C 0A`), a block stored whole having its two sizes equal. Each shader is a record {`u32` length, `u32` code area length, `u32` control length (0x880), in the global cache one more `u32`}, then its code area and right after it its control section.
- **Reader:** `shader_split.py` `ue3_package` / `ue3_body` decode a package's chunks (Zlib in Python, Oodle with the game's own decoder through the Mac host's tool mode, `DM_OODLE_SERVE`), and `ue3_shaders` finds each record by its control section's magic, its code area length (+0x704) back to the code area's magic, and the two lengths in front. `profile_guess.py` counts the packages whose first block is Oodle, so the NSP script gives the split the host. All of the game's 4,058 packages are decoded (the whole split took about a minute); none of 68 packages sampled besides the combined cache holds a shader.
- **Found on:** [Batman: Arkham Knight](../../dm/games/0100ACD0163D0000/README.md): 1,847 records in the global cache and 22,874 in the combined one, 21,923 shaders with the profile's [`shader_key_skip`](#control-sections-shared-among-shaders-that-differ-only-in-their-compile-information). Vertex and pixel shaders are paired as single-stage programs.

### The NDSR pack's record start

- **Symptom:** `draw skipped: no render pipeline (vertex - fragment -)` in a game's final passes; shaders "not in the shader map", all from `agl_ndsr.pck`.
- **Cause:** the pack's records start at +0x10 in one SDK's version and +0x14 in another's, by the words in its header (Donkey Kong Bananza 0x300, 0x6efeff, 0xc0000; [Splatoon Raiders](../../dm/games/0400380020056000/README.md) 0x40100, 0x80300, 0x90, 0x20c00).
- **Fix:** `ndsr_shaders` in `shader_split.py` takes the start at which more of the first 16 records point at an NVN2 control section.
- **Found on:** [Splatoon Raiders](../../dm/games/0400380020056000/README.md).

## Control sections and code

### Code area and control section layout

- **Code area:** graphics = 0x30-byte header (magic `0x12345678`) + 0x50 SPH + instructions (at +0x80); compute = 0x30-byte header (magic `0x12345679`) + **zero padding, instructions at +0x100** (all 71 of Mario's).
- **Control section:** magic `0x98761234`; +0x700 = c1 offset in the code area, +0x704 = code size, +0x748 = compute thread-group size (u32 ×3), +0x754 = compute shared memory bytes.
- **Found on:** [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md).

### Control sections shorter than 0x880 bytes

- **Cause:** the control section stores its own length at +0x20. Mario's, Engage's, Dread's and 13 Sentinels' are all at least 0x880 bytes; 3D World has 2,888 of 0x850 bytes (its ShaderData bfsha and six layouts) and 0x848 in BAHS version 9, each followed directly by other data.
- **Fix:** the key hashes the first `min(0x880, control +0x20)` bytes (`ctl_len` in `shader_split.py`, `keys_hash` in `hle_nvn.c`); the other games' tables came out byte-identical.
- **Found on:** [Super Mario 3D World + Bowser's Fury](../../dm/games/010028600EBDA000/README.md) (SDK 10).

### Program keys computed inside `nvnProgramSetShaders`

- **Symptom:** shaders reported missing at random, and their dumped control sections are garbage (Unity).
- **Cause:** NVN lets the caller free a program's control section once `nvnProgramSetShaders` returns; Unity does, and the host used to hash it at the first draw, so keys computed later hashed freed memory.
- **Fix:** the keys are now computed inside `SetShaders` (`hle_nvn.c` `prog_key`). Mario calls `SetShaders` for some programs with a code address that is not mapped yet, so a stage whose code is not in a registered memory pool at that moment is hashed at its first draw, as before.
- **Found on:** [Fire Emblem Engage](../../dm/games/0100A6301214E000/README.md) and Unity games in general (the source names no other); [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md) (the unmapped code addresses).

### Control sections shared among shaders that differ only in their compile information

- **Symptom:** `[RED] shader … is not in the shader map` for shaders whose control section and code area are both in the game files, but in different records: dumped (`DM_NVN_SHDUMP`), the control section is one record's and the code area another's.
- **Cause:** the engine shares one control section among shaders whose sections differ only in GLSLC's compile information (+0x778..+0x791 and +0x7d0..+0x7e7): `DM_NVN_PROGTRACE` showed the same control pointer set first with its own code area and later with another shader's. The GPU does not read those bytes, so on hardware it makes no difference; the host's key hashes them. Masked, Batman: Arkham Knight's 24,721 control sections are 5,615; the tools read nothing there (`dmsh_build.py` reads up to +0x75f).
- **Fix:** the profile key `shader_key_skip <offset>+<length>` (`0x770+0x78`): the key hashes that range as zeros, in `shader_split.py` (`key_ctl`, and `same_shader` ignores it) and in the host (`hle_nvn.c` `keys_hash`). Without the key, every game's keys are as before. It is a profile input of the shader chain, so the libraries follow it.
- **Found on:** [Batman: Arkham Knight](../../dm/games/0100ACD0163D0000/README.md): 40 shaders missing in the first 45 s before the key. With `0x7d0+0x18` alone the opening cutscene still met one whose section differed at +0x778, as [pull request #340](https://github.com/seedds/nxnative/pull/340) had found, so the range starts at +0x770.
- **Also found on:** [EA SPORTS FC 27](../../dm/games/0100C49025D3E000/README.md) (`0x7d0+0x18`): Frostbite's NVN renderer keeps one copy of each control section in a cache keyed by a hash that leaves out the compile information (`fb::ShaderPlatformDataNvn`, `s_controlBlockCacheMutex` in `Engine.Render.Core2.PlatformNvn.nro`); 10 programs were missing in the first menus, and the 10,342 keys are 7,873 shaders with the key. Its sections differed only at +0x7d0..+0x7e7 (+0x778 is zero in all of them).

### Compute code with and without a header

- **Cause:** compute code areas have a header in Mario (`0x12345679`), not in D2R. 3D World's particle compute shaders (SDK 10) start their instructions at code +0 (like D2R), not after a `0x12345679` header.
- **Fix:** `shader_split.py` accepts both; `dmsh_build.py` strips only a header that is there.
- **Found on:** [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md) (header); [Super Mario 3D World + Bowser's Fury](../../dm/games/010028600EBDA000/README.md) (no header, SDK 10).

### Exact program start for the decoder

- **Symptom:** 12 compute shaders (particles, post effects) translated into an endless loop that hung the GPU at the first gameplay scene; SPIRV-Cross failed on 7 shaders instead of 2 ([SPIRV-Cross parse failures](#spirv-cross-parse-failures)).
- **Cause:** Ryujinx's decoder treats every 32-byte boundary *relative to the start* as a scheduling word. Mario's compute code was first fed from +0x30: the groups were misaligned by 0x10, and instructions were dropped or misread.
- **Fix:** the program start must be exact (compute instructions at +0x100). `dmsh_build.py` now refuses compute code whose padding is not zero.
- **Check:** for a new title, the first non-zero word after the header must be a control word, and the trailing `EXIT` (`e30000000007000f`) / `BRA`-to-self (`e2400fffff87000f`) must sit in instruction slots of 32-byte groups counted from your chosen start.
- **Found on:** [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md).

### GLSLC outputs read through their section headers

- **Symptom:** GLSLC outputs fail the reader's check and are dropped (44 of DARK SOULS: REMASTERED's `main`'s 72 outputs and most of its archives' pixel shaders).
- **Cause:** SDK 4's GLSLC outputs (version 0x20.4): pixel shaders' control sections store at +0x20 a value 12 bytes short of their length, so the code area is not right after it.
- **Fix:** an output's own section headers say where everything is: the output's +0x4c is the section count, section i's header is at +0x90 + 0x90 i (`{size, data offset from the output, type}`, type 0 = GPU code), and a GPU code header adds +0x30 control offset, +0x34 code offset (both from the section's data), +0x38 code length, +0x3c control length. `glslc_section` reads them for outputs the old rule rejects.
- **Found on:** [DARK SOULS: REMASTERED](../../dm/games/01004AB00A260000/README.md) (SDK 4.5.0).
- **Check:** on every output the old rule keeps (1,073 in 13 Sentinels, Three Houses, FINAL FANTASY TACTICS, Engage and Age of Calamity) both give the same key; the new reading also takes 8 outputs of [Fire Emblem: Three Houses](../../dm/games/010055D009F78000/README.md) and 3 of [Hyrule Warriors: Age of Calamity](../../dm/games/01002B00111A2000/README.md) (GLSLC version 0x1d.0, whose control +0x20 holds odd values) that used to be dropped, so those two profiles' tables gained them.

### VertexA and VertexB parts merged into one vertex shader

- **Symptom:** `shader_split.py` finds no stage for vertex shaders and stops at the first captured output (`no GLSLC output`). Translated, every post pass and the window draw one texel's colour (a black screen with sound); every varying at (0, 0, 0, 1).
- **Cause:** shader type 0 in every vertex shader's SPH (word 0 = `0x0061`; every other game's vertex shaders have `0x0461`). Shader type 0 is Maxwell's first vertex slot (VERTEX_CULL_BEFORE_FETCH), and this GLSLC (GPU code version 1.15) writes every vertex shader as two programs in one code area: part A at code +0 (the type-0 SPH) loads the position attribute and writes only the position; part B, at control +0x700 with a 0x30-byte header of zeros and an SPH of type 1, reads the other attributes and writes the other outputs. Control +0x748 lists both as {SPH offset, length from the SPH, ?}. The hardware runs A, then B. Translated alone, A left every varying at (0, 0, 0, 1), so the full-screen passes sampled their sources at (0, 0) and the 2D draws their atlases' first texel.
- **Fix:** Ryujinx's `ShaderHeader` reads a graphics header's type 0 as a vertex shader; `sph_stage` does too. `dmsh_build.py` (`vertex_ab`) merges the two parts as Ryujinx's `TranslatorContext.Translate(VertexA)` does: A's instructions with each unconditional `EXIT` turned into a `BRA` to B's first instruction, then B's, under B's SPH with A's flags and input / output maps ORed in and the larger local memory sizes. c1 is what follows part B. A type-0 shader with no part B listed is translated as it is.
- **Found on:** [Catherine: Full Body](../../dm/games/0100BF00112C0000/README.md): in all 1,244 vertex shaders A writes no generic output, A ends in an `EXIT` predicated on CC FCSM_TR (which Ryujinx never takes; left as it is), an unconditional `EXIT` and a `BRA` to itself, and neither part has an absolute or table jump, so moving B keeps every branch; c1 (what follows part B) is zero padding in all of them. A census of the other games' packages and of every `main` on hand found no other type-0 vertex shader.

## Translation

### Tested tools and compiler-specific shader caches

- **Symptom:** a shader build accepts tools that differ from `dmsh_chain.lock`; after an override build, returning to the tested Metal compiler can still restore the override's libraries.
- **Cause:** `own_compilers.py` previously accepted any present Metal compiler or SPIRV-Cross. The shader and fetch-variant caches hashed the recorded Metal version and flags, without the actual compiler or SDK. Prefill checked only a sample's compiled bytes, so a mixed folder's other libraries could be attributed to the wrong compiler. The NSP stage key recorded Xcode's version, which can stay unchanged when its separately downloaded Metal toolchain changes, and did not follow `SPIRV_CROSS` overrides.
- **Fix:** ship a universal SPIRV-Cross built from a pinned commit and require its fingerprint for NVN builds. Require both tested Metal compilers before an NVN game's shader steps; `--allow-other-metal` (`DM_ALLOW_OTHER_METAL=1` for `own_compilers.py`) explicitly permits another present Metal compiler. Shader and fetch-variant keys include the actual Metal compiler and SDK identities. Earlier-library prefill requires that identity on every library's reflection (`chain_compiler`); matching a few sample compiles cannot establish the unsampled libraries' compiler. Unmarked legacy libraries are not imported, and the Metal cache format is bumped to discard entries prefill might have misclassified. NSP and regression keys hash the selected converter's bytes, including overrides; NSP keys also include both SDKs and Metal tools.
- **Found on:** [Breath of the Wild](../../dm/games/01007EF00011E000/README.md#found-on-this-game), [issue #242](https://github.com/seedds/nxnative/issues/242). The report used an untested compiler, but whether that caused the corrupted picture remains **open**. The cache failures were reproduced with stand-in libraries; this does not establish the rendering cause.
- **Check:** `test_own_compilers.py` checks enforcement and the explicit Metal override. `test_shader_toolchain_cache.py` first restored an untested library after switching compiler identities, for both main shaders and fetch variants, and now rejects it. It also checks stage invalidation on actual Metal, SDK and selected converter changes. `test_spirv_cross_pin.py` checks the shipped binary, pin, lock, architectures and licences.

### Metal compiler of a newer Xcode

- **Symptom:** after an Xcode update the build stops with `[RED] this Mac's Metal compiler is not the one the shaders were tested with (dm/gfx/tools/dmsh_chain.lock): metal_macosx Apple metal version 32023.922 …`; `xcodebuild -downloadComponent MetalToolchain`, as the message said, downloads the same untested compiler again.
- **Cause:** since Xcode 26 the Metal toolchain is a system asset that each Xcode downloads in its own build and `xcrun` takes for the selected Xcode (`xcodebuild -showComponent MetalToolchain` names it), whatever the macOS version. When #717 was reported, macOS 27.2 with Xcode 27.1 selected gave the then-pinned 32023.921, with Xcode 27.2 selected 32023.922. The message named the Xcode to install but not that it must be the selected one, and README linked the App Store, which offers only the newest Xcode.
- **Fix:** the stop message names the selected Xcode and its folder, and says to install the tested Xcode from Apple's download page (beside a newer one as `Xcode-<version>.app`), select it (`xcode-select -s` or `DEVELOPER_DIR`), then download its Metal toolchain; README says the same. The lock stays at the one tested compiler: another one needs a regression run before it is recorded, and the lock is part of the chain version. Whether two installed toolchain assets each follow their own Xcode is inferred from `-showComponent`, not yet tried with both installed (**open**).
- **Found on:** no game: [GitHub issue #717](https://github.com/seedds/nxnative/issues/717), on macOS 27.2 with Xcode 27.2 (Metal Toolchain 27B47).
- **Check:** `test_own_compilers.py` (the selection step before the download), `test_ci_toolchain.py` (the message's and README's Xcode are CI's).

### Idle MSBuild workers block regression measurements

- **Symptom:** `regress.py` waits for other jobs to end while the only reported jobs are sleeping `MSBuild.dll /nodemode:1 /nodeReuse:true` workers using 0% CPU.
- **Cause:** the busy-machine check treated every `dotnet` process as an active build. Reusable MSBuild workers can remain asleep after their build exits, and later builds can recreate them after a build-server shutdown.
- **Fix:** exclude only sleeping, zero-CPU MSBuild reuse workers whose parent has exited. Active workers, I/O-blocked workers, workers still attached to a parent, and build commands remain busy jobs.
- **Check:** `test_regress.py` supplies the process census that blocked the measurement, with active, I/O-blocked and parented controls.
- **Found on:** no game: Xcode 27.2 toolchain validation on an M4 Pro.

### SPIRV-Cross parse failures

- **Symptom:** SPIRV-Cross can't parse the SPIR-V (`nullptr` exception).
- **Cause:** 2 agl post-effect compute shaders still produce such SPIR-V. Before the compute start-offset fix ([Exact program start for the decoder](#exact-program-start-for-the-decoder)) it was 7: misdecoded code produced the bad SPIR-V.
- **Found on:** [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md).

### Pixel shader inputs the vertex shader does not write

- **Symptom:** `Fragment input(s) user(locnN) mismatching vertex shader output type(s) or not written by vertex shader` and the draw is skipped; the host reports `[RED] failed to create render pipeline (vertex … fragment …)`. Offline packing fails with `pso_offline: metal-tt failed` for the same reason.
- **Cause:** every pixel shader input must be a vertex shader output. Maxwell attributes are untyped slots: a pixel shader that reads a slot the vertex shader never wrote gets a default, but Metal refuses to build the pipeline. Translating one shader at a time cannot know what the partner reads. The offline archive exposed a second omission: prebuild reads both `shader_pairs.tsv` and BAHS program variations in `pso_pairs.tsv`, but padding read only the first table.
- **Fix:** `dmsh_build.py` reads both pair tables and gives each vertex shader the outputs its partners read but it doesn't write, as `float4` (0, 0, 0, 1) (`vs_pad`). A profile without `pso_pairs.tsv` takes the split's copy beside the package, as the archive builder does. The union controls pixel-before-vertex build order and shader cache dependencies ([Shader cache key covering the pairs](#shader-cache-key-covering-the-pairs)); the profile input stamp includes its `pso_pairs.tsv`, and the NSP script's shader reuse key includes the generated fallback. Padding precedes the depth, shadow-input, pulling and geometry-capture variants. Rebuild shaders and the host for the new chain. The host caches failed pipelines and reports each once with its vertex / fragment keys; before that, every draw retried the failing compile.
- **Found on:** [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md): 331 of 9,043 pairs failed; the opening course used none of them, later levels did (about 11,500 failed builds in one 11-minute session). The missing BAHS table was reproduced on [Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md#files-and-shaders) ([#283](https://github.com/seedds/nxnative/issues/283)) and [Super Mario Odyssey](../../dm/games/0100000000010000/README.md#files-and-shaders) ([#276](https://github.com/seedds/nxnative/issues/276)); their notes keep the pair counts and validation.
- **Check:** `dm/gfx/tests/test_dmsh_pairs.py` covers both table locations, union/deduplication, defaults and cache dependencies; Metal compiles the generated sources for Mac and iOS and builds an archive with regular and pulling vertex entries in both depth modes. `test_nsp_mac.py` checks fallback invalidation and progress ordering; `test_nsp_incremental.py` regenerates a missing fallback. Check all pairs offline after a game build (recipe in [05](../agents/05-debugging.md)).

### Single-stage programs paired with a cover

- **Cause:** a game that builds one NVN program per GLSLC output and combines vertex and pixel programs per draw with `BindProgram` stage bits has its real pairs only at draw time, so the offline build cannot know which pixel shaders a vertex shader must feed.
- **Fix:** the composite-program path ([Programs bound per stage](rendering.md#programs-bound-per-stage)) does all the pairing at replay. When every pixel shader input is a plain `float4` at fixed locations, `shader_split.py` pairs every vertex shader with a *cover*: the fewest pixel shaders whose inputs together are every input any pixel shader reads. `vs_pad` then gives each vertex shader all of them, and any vertex / pixel combination links.
- **Found on:** [13 Sentinels: Aegis Rim](../../dm/games/01003FC01670C000/README.md): 838 `ProgramSetShaders` at startup; every pixel shader input is a plain `float4` at locations 1–9 (the SPH input map, words 6–13, matches the translated MSL inputs for all 375 pixel shaders), so the cover is one shader.

### Unwritten components of a written output start as (0, 0, 0, 1)

- **Symptom:** outdoors dim and nothing casts a sun shadow (every shadow-map lookup fails).
- **Cause:** a vertex shader may write only part of an output (Link's Awakening writes a world position's x, y and z) while the pixel shader reads the rest (it divides by w). On the hardware the unwritten components read as (0, 0, 0, 1); Ryujinx gets the same by initializing the components the next stage reads when it links stages. The translator's MSL starts every vertex shader with `main0_out out = {};`, all zeros, so w was 0, and a fragment shader dividing by it got infinite compare values.
- **Fix:** general: `dmsh_build.py` (`vs_init`) sets every `float4` output to (0, 0, 0, 1) right after it, and the translated body overwrites what it writes.
- **Found on:** [The Legend of Zelda: Link's Awakening](../../dm/games/01006BB00C6F0000/README.md) (Mac and iPhone): 12,281 of 18,410 vertex outputs never write w. The sun's shadow map (2048×2048, depth 32-bit float) was correct; the surface shaders compute the lookup from a world position the vertex shader passes in `locn24`, divided by its w. Found with `DM_NVN_TEXRAW` on the shadow map, the light matrices from `DM_NVN_CBDUMP`, and a hand-edited copy of one surface shader that wrote the compare value and the stored depth as its colour. Compared with forcing the compare to pass (`compareFunction` always): the same brightness, plus shadows from the fences, the house and the trees.

### Early fragment tests from the control section

- **Symptom:** rows of dark dots along contour lines of alpha-tested surfaces drawn after a depth pre-pass (Tears of the Kingdom: the stone lattice in front of the Great Sky Island start, as reported by MacNativeGames ae8f4f4; not reproduced here).
- **Cause:** byte +0x718 of a pixel shader's NVN control section is its early fragment tests flag: the NVN driver's getter is `ldrb w0, [x8, #0x718]` in the `sdk` of all 21 games on hand, and Ryujinx receives it as the 3D engine's EarlyZForce. The translator always passed `earlyZForce` false. Without the attribute, Metal tests depth after a fragment shader that writes memory. In the lattice's G-buffer pass (EQUAL, no depth write, after a LESS pre-pass) that late depth differed by an ulp from the pre-pass along contour lines, the test failed there and the G-buffer kept its clear value.
- **Fix:** general: `dmsh_build.py` (`early_fragment_tests`) sets `early_fragment_tests` in the translation job of every pixel shader whose control byte +0x718 is set. `dmsh_xlate` passes it to Ryujinx as `earlyZForce`, and the MSL entry point is declared `[[ early_fragment_tests ]]`. A shader whose SPH writes depth (word 19 bit 1) is left as it is, since a test made early would not use the depth it computes.
- **Found on:** [Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md#found-on-this-game): 1,677 of its pixel shaders set the flag, and none of them writes depth. Rebuilding the chain changed exactly those 1,677 shaders' MSL, by the attribute alone. Elsewhere the attribute alone changed 699 of Wonder's 16,268 shaders and 51 of Scarlet's 30,142. Breath of the Wild and ANIMAL WELL set the flag in no shader.
- **Check:** `dm/gfx/tests/test_early_fragment_tests.py`.
- **Open:** with the attribute, a fragment shader that writes memory gives a picture that varies from run to run on Apple's GPU: [Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md#open)'s `8f470f5ce5864df7`, a per-fragment visibility flag (2026-10-30; the frame's replays: one digest without the attribute, 3 to 7 in 8 with it).

### Dead-code removal dropped the inputs of atomics

- **Cause:** Ryujinx's `Optimizer.RemoveNode` removes a dead node, then follows each source whose last use went away into the node that produced it. It did so even when that producer has side effects, such as an atomic or a call whose result is unused. Such a producer stays in the block, but its sources lost their uses, so the next dead-code pass deleted the chain that feeds it, and the atomic read undefined values.
- **Fix:** general: `RemoveNode` follows only into producers without side effects (`HasSideEffects`). This follows MacNativeGames' translator patch 0002.
- **Found on:** no game on hand so far. Rebuilding the chain changed no shader for this reason in Tears of the Kingdom, Breath of the Wild, Wonder, Scarlet or ANIMAL WELL (94,717 shaders); every MSL change was the early fragment tests attribute.

### Position invariance across shaders

- **Symptom:** the ground kept stale 32×32 blocks where the G-buffer draw failed the depth test: blocks of a stale image where a depth pre-pass and a later pass with other vertex shaders meet (Odyssey: 32×32 blocks on every fourth frame).
- **Cause:** a game may draw the same geometry with two programs and rely on identical depth, for example a depth pre-pass and then a less-or-equal G-buffer pass. Maxwell evaluates identical position math identically; Metal only does for outputs marked `[[invariant]]` (the translation marks `gl_Position`) when compiled with `-fpreserve-invariance`. Without it, Metal compiled two vertex shaders' identical position math to slightly different depths, so fragments of the second pass failed a less-or-equal depth test against the pre-pass and those pixels kept the previous content.
- **Fix:** `dmsh_build.py` (Mac) and `ios_metal.py` (iOS) pass `-fpreserve-invariance` to every compile.
- **Found on:** [Super Mario Odyssey](../../dm/games/0100000000010000/README.md) (Mac): the Cap Kingdom ground. 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. 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.
- **Note:** 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 there and draws.

### iOS math mode (relaxed)

- **Cause:** the translated shaders reproduce Maxwell's NaN behaviour with `isnan` selects (8,337 of 12,809 use them), and the default fast mode allows the GPU compiler to assume there is no NaN.
- **Fix:** the iOS libraries are compiled with `-fmetal-math-mode=relaxed` (`dm/ios/tools/ios_metal.py`; the Vulkan and GL libraries in `vk_shaders.py` and `gl_shaders.py` too). Relaxed keeps fast math's reassociation but not that assumption.
- **Found on:** listed among the mechanisms first met in [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md); the source gives the 8,337 of 12,809 count without naming the game.

### Front end at -O1: Apple's control-flow lowering

- **Symptom:** making a pipeline archive prints `Max Unstruct ID 758 in function agc.main in shader 13941 exceeds supported limit`, and one pipeline takes many minutes, in `metal-tt` offline or in the host's prebuild on the device (`Preparing pipelines` holding one count for minutes). The line is a note, not an error: `metal-tt` goes on, exits 0 and writes that pipeline into the archive.
- **Cause:** at the front end's default `-O2`, `CorrelatedValuePropagation` and the `SimplifyCFG` after it fold a translated shader's repeated tests of one value into switches whose cases jump past their own join, so nested `if`s become an unstructured graph. Apple's GPU backend lowers unstructured control flow with per-function IDs; past its budget (255 on `applegpu_g13s`, the check in `libapplegpu-nt.dylib` of AIR-TT 32023.921 picks 255 or 65,535 by a capability not yet characterized), it prints the note and lowers the function another way that is far slower and larger. At `-O1` those two passes do not run; the backend still optimizes the AIR itself, so ordinary shaders come out the same.
- **Fix:** `dmsh_build.py` (Mac) and `ios_metal.py` (iOS) compile with `-O1`.
- **Found on:** [Super Mario Odyssey](../../dm/games/0100000000010000/README.md#files-and-shaders), [issue #276](https://github.com/seedds/nxnative/issues/276): the float body of `alRenderMaterialUber` (fragment `007d7e411ef14ebd`, 16,227 `if`s in one function) took 24.6 minutes and 17.3 MB of `g13s` code at `-O2`, 31 s and 1.3 MB at `-O1`. Its `-O1` graph still goes past the budget (642), at that small cost.
- **Check:** random samples of archive bases from seven games (Odyssey, Tears of the Kingdom, Breath of the Wild, Kirby, Luigi's Mansion 3, Wonder, Pikmin 4; 22 to 400 each) compile at `-O1` to vertex and fragment code within a few percent of `-O2` (Odyssey 400 bases: −0.9% / +0.01%; TotK: +0.1% / −1.0%; Kirby: −3.1% / +0.07%), with no `Max Unstruct ID` note at either level, and the front end is faster. Odyssey's full `g13s` archive of 32,864 bases builds in 15.6 minutes (exit 0); its Cap Kingdom frames 4,200 to 6,600 match `-O2` except for scattered differences of at most 3 of 255 in three of nine frames, and their GPU time per frame (`DM_NVN_GPUPROF=every:150`, three runs each on an M-series `g18g`) has medians of 4.76 to 5.05 ms at `-O1` and 4.75 to 4.83 ms at `-O2`, within run-to-run noise.

### min, max and clamp with a NaN operand

- **Symptom:** the half of Mario facing the camera's left was white (GitHub issue 72, on an iPhone).
- **Cause:** spirv-cross writes `FMin` / `FMax` / `FClamp` as `fast::min` / `max` / `clamp`, which leave a NaN operand's result to the GPU. Maxwell's `FMNMX` returns the other operand and its saturation gives 0, as Metal's `precise::` functions (IEEE `fmin` / `fmax`) do. Plain `max` / `min` / `clamp` compile to the `fast_` intrinsics in every Metal math mode, relaxed included ([iOS math mode (relaxed)](#ios-math-mode-relaxed)). Galaxy's GX lighting normalizes light directions the game leaves at (0, 0, 0) (the key and fill lights of its characters) and takes `max(0, dot(ldir, NaN))` as the spot term: 0 on the M4 Pro with either function, so the attenuation is 1. The fill light sits at the camera's left in view space, so its terminator runs down the middle of a character facing the camera.
- **Fix:** `dmsh_build.py` (`precise`), `vk_shaders.py` and `gl_shaders.py` write the `precise::` functions instead (every game).
- **Found on:** [Super Mario Galaxy 2](../../dm/games/0100FD8022DAA000/README.md) (GitHub issue 72). The NVN chain had about 124,000 `fast::max`, 98,000 `fast::clamp` and 41,000 `fast::min` in [Paper Mario: The Thousand-Year Door](../../dm/games/0100ECD018EBE000/README.md), none guarded by its `isnan` selects.

### Integer fragment outputs declared by hand

- **Symptom:** a pipeline red line for a float output meeting an integer attachment (the red line prints the row to add); a BC encoder's pipeline fails.
- **Cause:** with no PC version the translator has no output types and declares every output float, and Metal refuses a pipeline whose float output meets an integer attachment. A scan for outputs built from integer operations was not trusted: the translated code keeps float bits in integer variables (register reuse), so such a scan cannot tell an integer output from raw float bits stored into a float target.
- **Fix:** the profile's `shader_facts_manual.tsv` declares such outputs `uint`. When another one fails, the pipeline red line prints the row to add. Shaders that write an output from integer bits now also get integer entry points ([Integer entry points for integer colour outputs](#integer-entry-points-for-integer-colour-outputs)).
- **Found on:** [The Legend of Zelda: Breath of the Wild](../../dm/games/01007EF00011E000/README.md) (agl's BC4 encoder `8e2cc2671cae6020` and its twin for a 2D source `93aa5e22a8202409`, whose failure showed up in an iPhone run); [The Legend of Zelda: Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md) (the BC4 encoder `0544282c162305db` and the Ultrahand mask `e3fb0bb58a691e10`, R8Uint, found by PR #11; the Mac route does not reach it); [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md) (two system pixel shaders write an ID into an `R8_UI` target on the character creation screen); [DRAGON QUEST VII Reimagined](../../dm/games/0100A9D01C446000/README.md) (one global shader writing an `R8Uint` target, [Unreal Engine 4.27 shader code archives](#unreal-engine-427-shader-code-archives-lz4)).
- **Note:** the compressed render target side is in [Compressed textures as render targets through an uncompressed shadow](rendering.md#compressed-textures-as-render-targets-through-an-uncompressed-shadow).

### Integer entry points for integer colour outputs

- **Symptom:** `failed to create render pipeline … output of type float4 is not compatible with a MTLPixelFormatR32Uint color attachment` (or `MTLPixelFormatRGBA16Uint`); every draw into the target skipped.
- **Cause:** with no PC version the translator types every fragment output float, and the shaders write bit-packed integers through `as_type<float>`. A hand-written row per shader (`shader_facts_manual.tsv`, Breath of the Wild's way, [Integer fragment outputs declared by hand](#integer-fragment-outputs-declared-by-hand)) does not scale.
- **Fix:** `dmsh_build.py` `fs_uint` gives a pixel shader a second entry point, `main0_u`, with every colour output the unsigned integer vector of the same bits, when it writes an output from integer bits or nothing but constants (a clear drawn into an R32Uint target); the hardware writes the output's bits to an integer target whatever the shader computed, so the copy is exact. The host takes it for a pipeline whose written colour attachments are all unsigned integer formats (`pso_all_uint`). A shader whose outputs are of both kinds also gets a third entry point, `main0_m`, that returns only the outputs written from integer bits as unsigned integers, and records their colour indices (`uint_mask_m` in its reflection); the host takes it for a pipeline whose integer attachments are exactly those (`pso_mixed_uint`). Otherwise a pipeline with a float target as well keeps `main0` and the red line.
- **Found on:** [Luigi's Mansion 3](../../dm/games/0100DCA0064A6000/README.md): `main0_u`; about 1,400 pixel shaders writing its G-buffer of two integer targets (R32Uint, RGBA8Uint), so every scene draw was skipped. [Pokémon Sword](../../dm/games/0100ABF008968000/README.md): `main0_m`; 10 pipelines as the opening's stadium loads; the scene shaders write colour 0 and 2 as floats and colour 1 from integer bits into an RGBA16Uint target (colour 1 in all 336 of Sword's). [Silver Pines](../../dm/games/0100D5B026844000/README.md): `main0_m`; URP's rendering-layers mask, an R8Uint target beside float outputs (colour 1 forward, 4 and 5 deferred), in 9,720 pixel shaders; without it only the fog drew.

### Integer textures and images through float declarations keep their bits

- Integer textures and images read and written through float declarations (an `RGBA8_UINT` texture read by `texture3d<float>` and bit-cast, `R16_UINT` / `R32_UINT` buffers through `texture_buffer<float>`, integer stores into a float-declared `texture3d` image) were checked on an Apple GPU: all keep the integer bits, so they need no facts.
- **Found on:** [DRAGON QUEST VII Reimagined](../../dm/games/0100A9D01C446000/README.md) (the volumetric lightmap's indirection texture, the light grid buffers).

### Integer-typed vertex inputs fed normalized formats

- **Symptom:** `failed to create render pipeline … Cannot convert attribute …`; draws missing (a pipeline that fails to build).
- **Cause:** the translation declares a vertex input `uint4` when the shader only moves its bits (x bit-cast to an integer, y copied as a float; [Integer vertex inputs](rendering.md#integer-vertex-inputs)), and a game that feeds it a 16-bit normalized format gets a pipeline Metal will not build (no integer input pairs with a normalized format).
- **Fix:** the profile's `shader_facts_manual.tsv` declares such inputs float (the bits Maxwell would see), which changes only that profile's chain version. The guess is not wrong in general, so overrides are added only for shaders seen failing with float formats. The general replay-side answers are the CPU conversion and the shadow vertex inputs ([CPU conversion of vertex attributes the input type cannot take](rendering.md#cpu-conversion-of-vertex-attributes-the-input-type-cannot-take), [Shadow vertex inputs](rendering.md#shadow-vertex-inputs)).
- **Found on:** [The Legend of Zelda: Link's Awakening](../../dm/games/01006BB00C6F0000/README.md): missing palm tree and flowers (the pipeline `147c829e79fb322f` + `0ebba5ff6a44aa52`, `shadowmap_pbr.bfsha` variant 641, vertex inputs 1–5): the foliage's shadow-map and depth pre-pass draws; the main pass shades with depth test EQUAL, so without the pre-pass the plants never appeared. Four outdoor shaders with the same loop (`03c43da3e3f46da8`, `079700f3577e311f`, `15c02683a10f8e9d`, `2344874045b1e758`) failed the same way on the iPhone in Mabe Village (Half2, RG8 / RG16 normalized) and are listed too. 1,437 of the game's vertex shaders have integer-typed inputs, many of them bone indices the shader multiplies as integers.

### Depth mode near = 0 needs the per-mode entry points

- **Symptom:** outdoors the whole screen blurred (depth of field out of focus everywhere).
- **Cause:** a game that sets depth mode near = 0 had every vertex shader translated for GL, so all depths sat in 0.5..1.
- **Fix:** the per-mode entry points of [04, section 5](../agents/04-graphics.md#5-depth-convention).
- **Found on:** [The Legend of Zelda: Link's Awakening](../../dm/games/01006BB00C6F0000/README.md) (Mac and iPhone): the depth-of-field composite (`52fcbb4ca2a78907`, `postprocess.bfsha`) blends the sharp frame with two blurred levels by a blur amount computed from the half-size depth copy: `pow(clamp((c5[2].y / (depth + c5[2].x) − c5[3].x) · c5[3].y), c5[3].z)` above a screen line, by screen position below it. Found with a save file from the iPhone (Mabe Village outside the house), `DM_NVN_PASSDUMP` / `DM_NVN_TEXDUMP` of that frame, and LASRecomp, which rewrites z only for the GL mode.

### Open: the window Y-direction register `SR_Y_DIRECTION`

- The register is −1 under a lower-left origin. The Ryujinx recompiler answers a constant +1 (a TODO in `InstEmitMove.cs`; LASRecomp, a static recompilation of Link's Awakening, answers −1). Answering it per profile means a translator change, i.e. every game's chain.
- **Found on:** [The Legend of Zelda: Link's Awakening](../../dm/games/01006BB00C6F0000/README.md): 2,487 of its pixel shaders read it; in the room shader the value only scales a depth slope under `abs()`; other shaders may use its sign (`dFdy`-based effects).

### Subgroup operations rewritten for a one-lane subgroup

- **Symptom:** `SPIRV-Cross threw an exception: Subgroup builtins are not available in this type of function`.
- **Cause:** Metal has no SIMD-group functions in vertex functions, and SPIRV-Cross's own subgroup emulation covers only `Elect`. RE Engine scalarizes uniform loads: ballot the active lanes, shuffle the first lane's value to all and take a fast path when all agree.
- **Fix:** with a one-lane subgroup that is exact, so `dmsh_build.py` rewrites the vertex shaders' SPIR-V that way before SPIRV-Cross (`spv_single_lane`: the ballot is the lane's own bit, shuffles and broadcasts return its own value, all / any are the predicate, all-equal and elect are true, lane ID 0, size 1). A mask component read through `OpAccessChain` / `OpInBoundsAccessChain` becomes an extraction from the constant mask, preserving the index.
- **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md) (about 4 % of its vertex shaders); [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md) (vertex `4b98320a0f512efe`, a prefix ballot that reads `SubgroupLtMask` through a component pointer).
- **Note:** the rewrite covers vertex shaders only: two of Bayonetta 3's domain shaders still fail this way ([Tessellation stages are not drawn](#tessellation-stages-are-not-drawn)).

### Geometry shaders that only pick the layer

- **Symptom:** the tone mapper read an empty LUT and the screen came out a flat blue.
- **Cause:** Unreal Engine 4.18 builds its 32×32×32 colour-grading LUT with `WriteToSliceMainVS` + `FWriteToSliceGS`, one slice per instance. Metal has no geometry stage.
- **Fix:** `dmsh_build.py` recognises such a shader in its SPIR-V (`gs_slice`: triangles in, at most 3 vertices out, positions and other inputs copied, `gl_Layer` taking the bits of one input component). It builds no library for it and records the component in its reflection. Each vertex shader that writes integer bits from the instance index into an output component gets two more entry points (`vs_layer_entry`: `main0_l` / `main0_lz` write that component as `[[render_target_array_index]]`). At bind time `prog_fold_gs` (`nvn_mtl.m`) swaps the vertex shader for its layer variant and drops the geometry stage when the components match. The pass and pipeline side of layer-writing vertex shaders: [Layered passes for vertex shaders that write the layer](rendering.md#layered-passes-for-vertex-shaders-that-write-the-layer), [Layer-writing vertex shaders get the triangle class](#layer-writing-vertex-shaders-get-the-triangle-class).
- **Note:** Unreal binds each stage from a program of its own, so a vertex shader is often paired for a moment with the previous draw's geometry shader. A stub that does not fit stays in the program, and only a draw that uses it prints a red line (`draw_gs_check`). The other 22 geometry shaders in DQ11 (one-pass cube shadows that emit transformed triangles) still do not build ([Geometry shaders with no Metal stage](#geometry-shaders-with-no-metal-stage)).
- **Found on:** [DRAGON QUEST XI S](../../dm/games/01006C300E9F0000/README.md).

### Geometry shaders with no Metal stage

- **Symptom:** a geometry shader fails to compile with `unknown main0_out`: SPIRV-Cross writes `unknown main0_out main0()` for the geometry stage, and the shader does not build.
- **Cause:** Metal has no geometry stage, so a geometry shader that does more than pick the layer has nothing to translate to. A geometry shader that only picks the layer is folded into the vertex shader instead ([Geometry shaders that only pick the layer](#geometry-shaders-that-only-pick-the-layer): `gs_slice`, `prog_fold_gs`; DRAGON QUEST XI S).
- **Found on:** [Fire Emblem Engage](../../dm/games/0100A6301214E000/README.md) (one geometry shader, in `unity_builtin_extra`; nothing in the opening scenes uses it); [Super Mario 3D World + Bowser's Fury](../../dm/games/010028600EBDA000/README.md) (fur fins, `RenderFinFur`, its one failing shader); [Captain Toad: Treasure Tracker](../../dm/games/01009BF0072D4000/README.md) (the fin fur geometry shader, as in Super Mario 3D World, is the one shader of its 7,258 that does not build); [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md) (53 geometry shaders); [FINAL FANTASY TACTICS - The Ivalice Chronicles](../../dm/games/010038B015560000/README.md) (one, `bahamuteffect.tec`); [Big Brain Academy: Brain vs. Brain](../../dm/games/0100620012D6E000/README.md) (Bezel's line renderer, below); [Luigi's Mansion 3](../../dm/games/0100DCA0064A6000/README.md) (10 geometry shaders, one expands a box from two corners; no draw used one from the title through the opening cutscene to the lobby). The Unity games' built-in geometry shaders do not build either (the container entries above).
- **Note:** a program whose geometry shader has no library used to draw silently without that stage, with the vertex shader's positions taken as clip coordinates; `draw with geometry shader … which has no library` now names it (Luigi's Mansion 3).
- **Note:** Big Brain Academy's 12 geometry shaders (`_BezelSystemResources/line/shader/line.bfsha`) do not build (`unknown main0_out`) and their draws are skipped (`[RED] primitive type 11 not implemented`). They take a line strip with adjacency (primitive type 11) and emit up to 12 vertices per segment as triangle strips: thick lines with joins and caps (28 emit sites). From the main menu on, every frame draws one such line (width 10 in its uniform block) into the top-left 512x512 of the frame and copies that region into a texture, so the menu's chart loses its outlines. Not done.
- **Note:** geometry shaders that do more than pick the layer are now built as compute kernels from the vendored translator's geometry-to-compute transforms ([geometry shaders on Metal](../geometry-metal-research.md), [Passthrough geometry shaders](#passthrough-geometry-shaders)).

### Passthrough geometry shaders

- **Symptom:** the models a passthrough geometry shader draws are missing: its own draw (a depth prepass) shows nothing, and so do the draws that test against that depth for equality. Before the shader loaded at all, the same models showed as black silhouettes.
- **Cause:** bit 24 of the shader header's first word marks a Maxwell passthrough geometry shader (NV_geometry_shader_passthrough): the input primitive goes on to the rasterizer unchanged, and the shader only writes per-primitive outputs and memory. A viewport mask (`a[0x3a0]`) of 0 sends the primitive to no viewport, which drops it. Ryujinx emits the primitive itself only when the host reports no passthrough support. The translator's accessor kept the default answer, "supported", so the compute kernel emitted no vertex and nothing was rasterized.
- **Fix:**
  - `Program.cs` answers `QueryHostSupportsGeometryShaderPassthrough` with false for the compute geometry jobs;
  - the vendored translator then emits the input primitive at each return of the shader (`EmitterContext.PrepareForReturn`) with every attribute the shader does not write: it is translated without the next stage, against the fixed 144-word input map. It skips the primitive when the shader wrote a viewport mask of 0, and the mask starts at all viewports (`Translator.EmitOutputsInitialization`);
  - `geometry_finish` in `dmsh_build.py` takes the output vertex count and primitive class of such a shader from the input topology (the header declares one vertex of a triangle strip).
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): `IndexBuildIndependent_{b,x}{cw,ccw}.gso` are the depth prepass, one draw per mesh, and each also writes the mesh's culled index list and its count for the shading pass's indexed indirect draws.

### Geometry passthrough and viewport masks in compute

- **Symptom:** geometry reader compilation fails with undeclared `gl_BuiltIn_5253` and a scalar `.elements` access. Compiling only the compute kernel also leaves native passthrough geometry without emitted vertices.
- **Cause:** the accessor advertises NV geometry passthrough, but Metal needs explicit compute emission. The reader tries to expose the unsupported NV viewport-mask array, and the raw SPH output bound can be one even though passthrough emits a complete input primitive. Geometry storage buffers also need their own driver-table base, and `BufferUsageFlags.Write` is bit 0.
- **Fix:** compute jobs report native geometry passthrough unsupported, activating the translator's explicit passthrough emission. `geometry_vertices` records the translated bound per input topology. `GeometryToCompute` applies viewport-zero mask visibility to emitted indices before strip compaction; the vertex reader skips the mask output. This covers the current replay's viewport-zero path, not replication across multiple viewports. Geometry storage-buffer slots decode from `c0 + 0x410` in reflection and replay, and writable reflection follows the translator's actual flag definition.
- **Found on:** [NieR:Automata US 1.0.3](../../dm/games/010056B015FE8000/README.md) (four index-building passthrough shaders; all 1,225 compile; visible/rejected primitive GPU fixtures pass).

### Geometry capture indices deduplicated with a bounded bitset

- **Symptom:** sorting repeated mesh indices consumes much of the render thread's CPU time.
- **Cause:** every indexed geometry draw copied and comparison-sorted its full input index list to construct sorted unique vertex capture IDs, even for static meshes rendered every frame.
- **Fix:** mark indices relative to the validated minimum in a bitset, then enumerate set bits in ascending order. The existing vertex-capture bound limits the bitset to about 57 KiB; primitive assembly retains the original index order and restart boundaries. Non-indexed captures need no unique index list. No cross-frame cache or new hazard assumptions are involved.
- **Found on:** [NieR:Automata](../../dm/games/010056B015FE8000/README.md): a Mac CPU sample attributes 316 of 687 replay samples to `qsort`. The geometry tests cover repeated indices, a sparse range spanning multiple bitset words, restart, base vertex and instancing (121 GPU checks, also passing with forced conservative waits).

### Geometry runs retain external buffer dependencies

- **Symptom:** later geometry uses stale shader-written data or appears in the earlier draw's position when draws are combined into a geometry run.
- **Cause:** a run captures all vertex inputs, dispatches all geometry shaders, then rasterizes all draws. That reorders a later capture before an earlier geometry write. Shader writes can also communicate with other geometry or fragment work in the run. Separate scratch outputs do not prove independence of the game's own buffers.
- **Fix:** record external vertex-buffer, uniform and storage-buffer access ranges for each job, including write access and read-only storage metadata. A read/write or write/write overlap stops the run before the next capture; read-only sharing and disjoint writes can still batch. Unknown sizes, range overflow, indirect storage pointers, image accesses and typed buffers conservatively stop batching. Existing encoder barriers then order the separate jobs. `t_geometry_hazards` exercises an actual geometry write followed by vertex capture through a storage binding and through a uniform alias, checking both the counter and literal covered pixels. All 14 checks pass with early submission enabled or disabled. Removing the dependency guard causes six failures: both counters are stale and both dependent triangles disappear from their expected positions.
- **Found on:** investigation of [NieR:Automata US shadow jitter](../../dm/games/010056B015FE8000/README.md#library-shadow-sequence-and-geometry-ordering). This is a reproduced general replay bug; its contribution to that game's reported shadow motion is not yet established.

### Geometry dependency checks skip read-only pairs

- **Symptom:** the CPU spends substantial replay time checking dependencies among geometry jobs that share many read-only buffers.
- **Cause:** the original dependency predicate compares every pair of access ranges, including read/read pairs that cannot communicate. Repeating that quadratic scan across a run can dominate the check even when each job has only one write. Even without read/read pairs, checking each new job against every earlier job of the run stays quadratic in the run's length: in Donkey Kong Bananza (runs of about 56 jobs, 19 ranges and 2 writes a job, hundreds of geometry draws a frame) that loop alone was about 8 % of the submitting thread's samples.
- **Fix:** maintain a 64-bit write mask as ranges are added or promoted from read to write. A run keeps two sorted sets of merged half-open intervals (`geo_run_access`, `nvn_geometry.inc`): the union of its jobs' writes and of their reads. A new job conflicts with some earlier job exactly when one of its ranges meets the write union or one of its writes meets the read union. Overlap with a union is overlap with one of its members, so merging touching and overlapping intervals loses nothing. Each check is a branch-free binary search, and a range already covered is not inserted again. Unknown access, overflow, half-open bounds and the 64-range limit retain their conservative behavior; the run's ordering policy is unchanged.
- **Check:** `t_geometry_hazards` compares 4,096 generated runs of 1 to 63 jobs with the exhaustive scalar pairwise oracle (both outcomes occur), and covers bit 63, duplicate promotion, adjacent bounds, reads across two merged writes and overflow on either side. All 26 checks, including the existing real-GPU storage/uniform dependency cases, pass with early submission on and off. On the same scripted Donkey Kong Bananza route (`sample`, 25 s), the submitting thread's `geo_draw_` self time fell from 562 to 81 samples, plus 42 in the new set code, and the game's 60/30 interval switches fell from 30 to 12. In an isolated CPU benchmark of a 64-job run with 32 shared reads and one private write each, the run check takes 13 us against 54 us for the pairwise check.
- **Found on:** [NieR:Automata iPhone CPU capture](../../dm/games/010056B015FE8000/README.md#iphone-resolution-and-frame-rate-controls) (the read/read pairs); [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md#found-on-this-game) (the quadratic run).

### An empty vertex stream kept every geometry draw out of a run

- **Symptom:** a game whose frames are mostly geometry-shader draws stays GPU-bound by attachment traffic: a `DM_NVN_GPUPROF` frame shows hundreds of two-draw render passes (each geometry draw, then the reopened pass of its raster) and gigabytes of attachment loads and stores, and no `geometry hoist` line ever appears in the log.
- **Cause:** the game leaves a vertex stream bound at address 1 with size 0. Three checks looked at every one of the 16 streams, used or not: the lookahead's alignment test in `geo_prep` (a run may not prepare an unaligned stream ahead of intervening draws) returned -1 for it; `geo_access_of` added it as a zero-length range, which `geo_access_add` marks unknown, so every pair of geometry draws counted as a dependency; and `csh_draw_reads` could not list the draw's reads, which stopped compute dispatches from moving out of the pass. `vbuf_prepare` and `bind_vbufs` also sent it to `vbuf_unaligned`, which logged a `[RED]` that it could not be copied although there were no bytes to copy.
- **Fix:** all of them skip a stream of size 0. It holds no guest memory that a draw could depend on: a draw that reads past it gets zeros, from a copy filled on the CPU with no ordering to keep (`vbuf_unaligned`), so neither the ordering nor the dependency proof changes for a stream that holds data. Geometry captures do not bound their streams (`geo_capture_one` clears `g_vbuf_need`), so a vertex shader that really reads an empty stream there still reads the memory at its address, or a stale binding if that address is not aligned, instead of zeros, with or without a run. That is open.
- **Check:** `t_geometry_hazards` forms a run across an empty stream at address 1, keeps a stream with data at an unaligned address out of one, and checks the lists of `geo_access_of` and `csh_draw_reads`; `t_vbuf` checks that an empty stream logs no red line and still reads zeros when a draw reads it. The frame rate below is the reporter's measurement, not reproduced here.
- **Found on:** FINAL FANTASY XII THE ZODIAC AGE (`0100EB100AB42000`, no profile yet): about 180 geometry draws a frame from two triangle-in, three-vertex-out shaders. On an M1 in handheld mode, gameplay went from 264 render passes and 2.9 GB load + 2.3 GB store a frame (GPU 52–68 ms, 12–17 fps) to 6,053 draws prepared in 153 runs by frame 1,500, GPU about 15 ms and a steady 30 fps, the game's own rate there.

### Parallel geometry index compaction

- **Symptom:** geometry-heavy rendering spends substantial GPU time compacting the geometry shader's emitted indices.
- **Cause:** the original `geo_pack_indices` kernel walks every invocation group in a single GPU thread.
- **Fix:** use upstream commit `16f4889`'s ordered count, offset-scan and write dispatches. The count kernel computes exclusive offsets within each 64-group block using SIMD scans and a threadgroup barrier; the middle kernel scans block totals, and the write kernel adds local and block offsets. Scratch holds one word per group plus one per block. The middle scan now uses 128 GPU lanes, each walking a contiguous chunk before scanning lane totals across SIMD groups. Draws with at most 1,024 invocation groups instead compact in one 128-thread workgroup, without global scratch or two extra dispatches. Both paths preserve exact invocation order, strip winding, restart boundaries and whole-primitive clipping at the output capacity. `DM_NVN_GEOMETRY_PACK=legacy` restores the earlier three-dispatch policy with its single-thread block scan.
- **Found on:** upstream OCTOPATH TRAVELER II; the blocked scan was measured on [NieR:Automata](../../dm/games/010056B015FE8000/README.md). On an M4 Pro, a 65,536-group compaction benchmark with alternating warmed runs measured 3,918 us median for the upstream scan and 81 us for the blocked scan, with identical output. Production geometry checks pass 157 checks, including 257-group partial blocks, empty groups, restart boundaries, output capacity and pending index writes. Both normal and forced-wait policies pass. Full-game timing remains scene-dependent.
- **Further check:** the newer fused/parallel paths pass 914 geometry checks in normal, legacy and split-compute modes, including 0–1,024-group boundaries, points/lines/triangles, partial chunks and a 16,387-block prefix stress case. An alternating warmed 256-draw query workload improves 2.998 to 2.832 ms GPU median with identical query results; the large block-scan stress benchmark improves 1,221 to 45.8 us with byte-identical prefixes and indirect counts. The latter is an isolated stress workload, not a measured game frame. Game FPS with these additional changes remains to be measured.

### Geometry reader viewport masks

- **Symptom:** the geometry kernel compiles, but its `_bgv.metal` vertex reader fails with `use of undeclared identifier 'gl_BuiltIn_5253'` and `member reference base type 'int' is not a structure or union`.
- **Cause:** `ResourceReservations` packs `ViewportMask` into one output word; `GenerateVertexPassthroughForCompute` stores that word as a scalar although SPIR-V `IoMap` declares builtin 5253 (`ViewportMaskNV`) as `int[1]`. SPIRV-Cross emits an invalid array-copy expression for it, and Metal has no viewport-mask output. The mask still matters: zero discards a primitive, while bit zero selects viewport zero ([NV_viewport_array2](https://registry.khronos.org/OpenGL/extensions/NV/NV_viewport_array2.txt)).
- **Fix:** `dmsh_geometry.reader` recognizes that generated scalar-array store, retains the mask's integer bits and moves vertices beyond the X clip plane when bit zero is clear. This runs before `vs_depth` makes both depth entries. Using the X plane preserves culling with depth clamp too; moving vertices beyond the far depth plane allowed masked points to render when depth clamp was enabled. Unknown mask expressions fail the build instead of silently losing culling. No pinned translator change is required.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): its four geometry shaders' readers.
- **Check:** `dm/gfx/tests/test_dmsh_geometry.py` uses synthetic generated-reader MSL. Both Mac and iOS compiler checks reproduce the original error before the fix. Mac GPU readback covers six masks, points and triangles, depth clip and clamp, and both production depth entries (48 cases); zero-mask points under depth clamp fail with the earlier depth-plane cull. All four original readers also pass both platform syntax checks.
- **Limits:** replay binds viewport zero only. Multi-viewport fanout is not implemented, and per-vertex-disagreeing masks are not validated; the tested primitives use one mask for every vertex. These tests establish compilation and viewport-zero culling, not the appearance of those effects in gameplay.

### Passthrough geometry output bounds follow input topology

- **Symptom:** a passthrough geometry kernel writes beyond its vertex and raw-index scratch regions, and compaction reads the wrong groups; corruption can reach the following primitive-map buffers without a red diagnostic.
- **Cause:** the builder used the SPH's raw maximum output count for every topology. Ryujinx's `ShaderDefinitions` overrides that count for `GpPassthrough` (SPH word 0 bit 24) with the input primitive's vertex count excluding adjacency. A header can declare one vertex while the triangle variant writes three vertices and six raw index words per group. Allocating two index words per group then overlaps later scratch regions.
- **Fix:** `dmsh_geometry.output_layout` mirrors the translator's override for each input variant: points, lines, lines with adjacency, triangles, triangles with adjacency use vertex bounds `1, 2, 2, 3, 3` and index strides `2, 4, 4, 6, 6`. Passthrough also inherits the input primitive: native GLSL omits the SPH's output primitive declaration, so `ShaderDefinitions` and metadata use point, line strip or triangle strip accordingly. `geometry_finish` writes those bounds and topology into each variant's metadata, which replay uses for allocation and compaction. Ordinary geometry shaders retain their declared bound and topology. Rebuild shaders and the host for the new chain.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): its passthrough index-building geometry shaders; the triangle kernel's six-word stores disagreed with its two-word metadata while rendering flight. The game note tracks the visual result after correction.
- **Check:** `test_dmsh_geometry.py` covers all five variants, ordinary bounds and malformed headers. With `DM_TEST_GEOMETRY` pointing at `t_geometry`, it passes production-generated layouts to a native emitter with independent topology bounds. GPU canaries protect the nominal vertex and raw-index allocations, and compaction must preserve the emitted count and order. The old rule produced 34 failures; the corrected metadata passes all 170 native checks with Metal validation.

### Passthrough geometry kernels lose raster vertices

- **Symptom:** a geometry kernel compiles and runs its storage-buffer side effects, but contains no `geometry_vb_output` binding and fills its index stream with restart markers; its following raster pass draws no geometry.
- **Cause:** the offline accessor inherited native geometry-passthrough support. Native NV passthrough forwards input vertices without an explicit emit, but Metal's compute stage must emit them. Disabling native support alone still loses generic attributes: independently compiled stages have no linked next-stage usage to populate `PassthroughAttributes`.
- **Fix:** fixed-input compute jobs report no native geometry-passthrough support. `EmitterContext` explicitly emits input positions and every captured generic attribute the GS did not replace, removing adjacency vertices as the input topology requires. This uses the complete Maxwell capture map; ordinary linked translation retains its existing usage mask. Copying all 128 generic components also exposes a Metal backend compiler failure at normal optimization (`Internal compiler error` when creating the pipeline, despite successful offline compilation). `dmsh_geometry.capture_input` keeps capture reads volatile, preserving their values and normal kernel optimization while preventing that failure. No game IDs or shader keys participate.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): its passthrough index-building shaders run before depth-writing raster passes during flight. The game note tracks rendering after rebuilding.
- **Check:** `PassthroughCompiler` translates a synthetic EXIT-only Maxwell passthrough shader for all five input topologies, deliberately declaring triangle output in every raw header. Both Mac and iOS compile checks run. With `DM_TEST_GEOMETRY`, production reflection and the argument encoder dispatch the actual kernels; readback checks positions, all 32 generic attributes, adjacency selection, two instances, raw indices, restart padding, primitive compaction and allocation canaries (9,066 native checks). The original pinned translator fails all five variants because raster output is absent. Removing volatile capture reads makes the 128-component synthetic pipeline fail to compile on the M4 Pro; reducing to 96 components or compiling at `-O0` independently removes that failure.

### Geometry capture waits for every indexed draw

- **Symptom:** geometry-heavy frames have a large submission wall/CPU gap despite short GPU execution; stack samples spend most of the main thread in `geo_draw` → `mdi_wait` → shared-event wait, with index sorting another substantial cost.
- **Cause:** every indexed geometry draw unconditionally committed and drained the GPU before assembling its primitives on the CPU. It also waited for vertex streams consumed only by native GPU capture. The capture list sorted every index reference to remove duplicates, even though the scratch bound already limits the vertex span.
- **Fix:** `nvn_geometry.inc` checks each actual CPU-read range against pending copies/updates/fills, current shader writes and previous live submission writes. Unknown or overflowing write sets retain the wait. Indirect records are checked before decoding; their selected index ranges are checked afterwards. CPU-converted vertex bindings retain the same hazard checks, while native capture follows GPU encoder ordering. A successful drain covers the remaining reads of that draw. Capture uniqueness uses a bounded vertex-span bitset and ascending set-bit enumeration; allocation failure or a wider span retains an in-place sort fallback. Non-indexed captures need no unique index list. The primitive map, restart handling, base vertex and instance semantics are unchanged.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): its flight geometry path. The game note records matched frame-time and visual validation.
- **Check:** `t_geometry` first reproduced five unnecessary drains for clean static and indirect inputs. Its 226 native checks cover clean versus pending/current/prior GPU-written ranges, unknown write sets, GPU-only and converted streams, indirect arguments selecting a different dirty index range, 8/16/32-bit indices, restart, base vertex and instances. Actual copy/compute writes and raster readback validate synchronization; converted data must see the new GPU-written value. Sparse duplicates, ascending output, short destinations, fallback and allocation canaries cover the capture list. These checks establish behavior; throughput is measured separately in gameplay.

### Geometry compaction serial prefix scan

- **Symptom:** geometry-heavy frames remain GPU-bound after CPU readback waits are removed; encoder timestamps attribute most compaction time to the offset scan rather than the geometry shader or index writes.
- **Cause:** count and write dispatches ran in parallel, but one GPU thread walked all group counts to compute ordered offsets. Geometry replay also created a separate compute encoder for each of the geometry, count, scan and write dispatches, with full queue barriers at every boundary.
- **Fix:** `geo_pack_scan` divides the groups into contiguous chunks across at most 256 threads. Each thread sums its chunk; a threadgroup prefix scan gives its starting offset; a second chunk pass writes the ordered offsets. The group count does not need to fit in the threadgroup. Totals use 64 bits; unrepresentable 32-bit offsets saturate, output writes use subtraction-based capacity checks, and indirect counts retain complete-primitive truncation. `geo_pack_encode` keeps the geometry and three compaction dispatches in one encoder, with explicit dispatch-to-dispatch device-memory barriers. Full barriers remain at its boundaries. The primitive order and large-draw algorithm are preserved; no shader-library chain change is required.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): its flight geometry passes. The game note records the baseline encoder timings and matched result.
- **Check:** `t_geometry` passes 1,066 checks with Metal validation. A GPU producer and the production compaction stages share one encoder; readback compares every offset and ordered index with an independent CPU result, including 129,032 groups. Cases cover points, lines and triangle strips, empty groups, zero/odd/thread-boundary group counts, 1/32/192/256-thread scans, short capacities, integer overflow and allocation canaries. The actual translated passthrough integration also passes all five input topologies through the same encoder path. Performance is checked separately with matched gameplay.

### Geometry draws above the capture scratch bound

- **Symptom:** `[RED] geometry draw needs 261 MiB of scratch, above the 256 MiB bound` (or `exceeds the capture or invocation bound`); the draw is skipped.
- **Cause:** the geometry path captures the vertex shader's outputs for every vertex between the draw's lowest and highest index, for every instance, and the geometry kernel's outputs take room per primitive (`nvn_geometry.inc` `geo_draw`). Unreal's one-pass point light shadow draws a whole level mesh through its geometry shader.
- **Fix:** such a draw is drawn in runs in the hardware's order (`geo_split`): one instance at a time, then runs of at most 4,096 primitives of a list topology whose index range fits 64 MB of capture. The index list is read once and the runs skip the GPU wait the outer draw took (a first version halved the draw recursively, each part reading the indices again and waiting: that frame took 805 ms). The index compaction after the geometry kernel was one thread walking every group; it is a count, an offset scan and a write dispatch now (`nvn_geometry.h`, `t_geometry`). Scratch arenas above 32 MB are kept while such draws go on and released after 600 idle submissions (`nvn_mtl.m` `arena_begin`), and `arena_new` takes the smallest free arena that fits.
- **Found on:** [OCTOPATH TRAVELER II](../../dm/games/0100A3501946E000/README.md): 129,032 triangles of one mesh (vertex `9214254d203a4e45`, geometry `cb17a9453d4789d0`, both from the default material) at the start of the first battle and in the scene after it; with the one-thread compaction each such frame took about 300 ms, now only the first is long (395 ms on the Mac, with two new pipelines and the scratch allocation).
- **Note:** strips, fans and primitive restart are not split. Guest primitive numbering retains the original draw's offset across supported splits: [Split geometry primitive IDs](#split-geometry-primitive-ids).

### Split geometry primitive IDs

- **Symptom:** a geometry shader that reads `PrimitiveID` repeats its visibility decisions or storage addresses after a large draw is split into compute dispatches.
- **Cause:** fixed-input `GeometryToCompute` lowered the guest ID to dispatch-local `GlobalId.x`. Each child dispatch starts at zero. That local index is correct for capture, topology remapping and scratch output addresses, but loses the guest draw's primitive offset.
- **Fix:** replay passes an explicit primitive base through recursive splits, keeping it unchanged when splitting instances and adding each primitive run's start when splitting the range. The counter resets between instances, as specified for [geometry primitive IDs](https://docs.vulkan.org/refpages/latest/refpages/source/PrimitiveId.html). `vb_info.geometry_counts.y` carries it to the shader. Only the guest `PrimitiveID` load adds that base; helper addressing remains local, and ordinary draws start at zero. The rebuilt translator and shared chain stamp prevent older libraries from using the new ABI.
- **Found on:** review of [OCTOPATH TRAVELER II](../../dm/games/0100A3501946E000/README.md)'s large-draw splitter following the NieR integration. The primitive-ID failure is reproduced with a synthetic Maxwell shader; no new Octopath visual failure is attributed to it.
- **Check:** `test_geometry_split.py` checks 160 nested range/ID cases over all five supported list topologies, all index widths, non-indexed draws, two instance counts and nonzero ancestor bases. `ComputeSplitPrimitiveId` in `test_dmsh_geometry.py` compiles a real Maxwell `ALD PrimitiveID` shader for both Metal targets. Its native `t_geometry_primitive_id` replay checks every emitted ID and captured vertex/instance attribute in 40 GPU cases, with 1/8,199 primitives, five topologies, 1/2 instances and 1/3 shader invocations. The old translator repeats ID zero at primitive 4,096; the corrected translator preserves it. A GPU-independent ABI assertion catches stale translator binaries too. Feedback counts and an output canary also check that the offset never leaks into local scratch addresses.

### Recursive geometry splits retain their original draw range

- **Symptom:** a large geometry draw crashes in `geo_split` / `memcpy` while reading indices, sometimes preceded by `geometry draw exceeds the capture or invocation bound` for an unexpectedly small child draw.
- **Cause:** `geo_prep` passes a draw descriptor stored in the reusable `g_geo_run` job to the splitter. A child draw prepares its own job in the same slot, overwriting the parent's borrowed descriptor. Later chunks add their offsets to the previous child's index address or first vertex instead of the original draw's range. The job-based geometry integration exposed this lifetime error; the earlier splitter received a local descriptor.
- **Fix:** `geo_split` snapshots the input descriptors before issuing any child draw. Every chunk and recursive instance split derives its range from that immutable copy. This repairs ownership for indexed and non-indexed draws without changing split limits or skipping geometry.
- **Found on:** [OCTOPATH TRAVELER II](../../dm/games/0100A3501946E000/README.md), the 129,032-triangle shadow draw after Hikari's first attack, during the NieR geometry integration's cross-game review.
- **Check:** `test_geometry_split.py` compiles the production splitter with a callback that reuses the caller's descriptor exactly as `geo_prep` does. The expanded 160-case test checks each chunk's range, instance and primitive base for 8/16/32-bit indices and non-indexed draws, including 129,032 primitives and nested splits. Removing the snapshot reproduces the range drift.

### Geometry jobs preserve transform-feedback replay

- **Symptom:** geometry draws with transform feedback report `geometry draw cannot capture this pipeline` and write no feedback records after the geometry-job integration.
- **Cause:** the new preparation path rejected active feedback and omitted main's gather of the geometry shader's emitted vertices. Replaying a geometry reader through the ordinary vertex-feedback branch would also capture its synthetic draw a second time.
- **Fix:** full geometry jobs gather their packed output into feedback before rasterization or raster discard. Active feedback remains in draw order, and feedback state/buffer commands end lookahead. The subsequent synthetic raster draw does not capture vertex feedback again. Lean and mesh eligibility still exclude active feedback.
- **Found on:** final review of the [NieR geometry integration](../../dm/games/010056B015FE8000/README.md#geometry-optimization-review), against main's existing geometry-feedback implementation; no game-specific condition is involved.
- **Check:** `t_geometry_feedback` replays actual capture, geometry, compaction, feedback pause/resume and rasterization with GPU readback. Its oracle checks every emitted vertex and saved count, excludes the draw made while paused, and runs with raster discard both enabled and disabled.

### Passthrough geometry shaders drawn by the vertex shader

- **Symptom:** geometry-heavy frames are GPU-bound on the emulation itself. A `DM_NVN_GPUPROF` frame shows, for every geometry draw, a capture pass (`0x0 0c`), a compute encoder and a one-draw reopening of the game's pass, with most of the frame's GPU time in the first two.
- **Cause:** a passthrough geometry shader (SPH word 0 bit 24, NV_geometry_shader_passthrough) cannot change the vertices it passes on; it only decides per primitive (the viewport mask) and may write storage. The emulation still captured all 144 words of every vertex, ran a kernel that copied 134 words per output vertex (volatile loads of the whole capture map, [Passthrough geometry kernels lose raster vertices](#passthrough-geometry-kernels-lose-raster-vertices)), and drew the copies through the reader.
- **Fix:** for point, line and triangle input (no adjacency), `dmsh_build.py` adds a lean kernel `<key>_bl<topology>` (`dmsh_geometry.lean`): every output store but the viewport mask's word goes, the mask is written one word per output vertex, what only fed the removed stores is dropped (a worklist over the generated straight-line code), and the capture reads that remain (literal words of the record) are renumbered into a record of just those words (`capture_reads`). The capture writes each word once through a slot table and a record stride the host passes (`dm_geometry_capture`; the full kernel gets every word in order). At replay (`nvn_geometry.inc`) `geo_lean_cull` turns the vertices of masked primitives into restarts before the usual compaction, `geo_lean_index` maps the packed output vertices to the draw's own vertex indices (output vertex `g × n + k` is input vertex `k` of primitive `g`, the primitive map's word) and sets the base instance, and the program's own vertex and fragment stages draw them indexed and indirect. Instanced or multi-invocation draws, vertex shaders with storage writes (they would run twice), transform feedback and adjacency input keep the full path; so does `DM_NVN_GEO_LEAN=0`. The CPU side reads indices with typed loops and copies list topologies without restart directly into the primitive map; a `memcpy` per index was most of the translation time of such frames.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): its four `IndexBuildIndependent_*` shaders, about 80 draws a frame in the opening flight, most of which keep no primitive (they build culled index lists for later passes). The lean kernels read 4 of the 144 capture words and are about 970 lines instead of 10,500. On an M1 Mac at 720p: GPU work per frame 33–43 ms → about 10 ms, translation 75 M → 42 M instructions per frame, frame time 50 → 27 ms where the Mac was otherwise idle.
- **Check:** `test_dmsh_geometry.py` covers capture slots, lean rewriting, both Metal targets and mesh composition. The native geometry suite covers capture, compaction and dependency ordering. All M1 measurements in this entry were reported for the original implementation in PR #416; they are not measurements of the reconciled host.
- **Fix:** the shared geometry run (`geo_run_walk`, `nvn_geometry.inc`) batches independent captures, geometry kernels and rasters, retaining #336's buffer dependency proof, exact-byte index cache, query positions and parallel compaction. A storage write that could affect another capture or raster ends the run. Lean triangle output uses fixed triangle slots and skips compaction; point and line output retains ordered compaction, then maps output vertices back to the draw's own indices. `DM_NVN_GEO_HOIST=0` disables batching for comparison. No game identity participates in these decisions.
- **Found on:** the same scene, M1 at 720p: render passes a frame 205 → 39, encoders 290 → 45, attachment traffic 860 → 190 MB, GPU busy 7.6 ms, translation 4.6 ms, frame time p50 / p95 16.66 / 17.26 ms: 60 fps held through four minutes of the flight, factory and ocean combat (frame time 16.6–16.8 ms in every 300-frame window), where it was 50 ms at first.
- **Fix:** the compute path still costs the CPU a capture, a kernel and a compaction per draw, and the GPU the vertex shader twice. A pair listed in the profile's `geo_mesh.tsv` gets a mesh library `<vertex key>_<geometry key>_m3` at the end of `dmsh_build.py` (`dmsh_mesh.py`): the vertex shader's pulling entry ([vertex pulling](../../dm/gfx/tools/dmsh_pull.py)) and the lean triangle kernel, each in its own namespace, under one mesh entry per depth convention (`dm_mesh`, `dm_mesh_z`). A threadgroup takes 32 triangles of a list: 96 threads run the vertex shader once per index, write the vertex, and keep the words the kernel reads in threadgroup memory; the first 32 run the kernel body (capture reads from that memory, the topology map as the identity, the mask store into a variable, index stores dropped) and cull their triangle when viewport 0's bit is clear (`[[primitive_culled]]`). The kernel's storage writes and the subgroup operations that order them run as before, one triangle per thread, 32 consecutive triangles per SIMD group. The host (`geo_mesh_draw`, `DM_NVN_GEO_MESH=0` off) takes it for a non-instanced, non-indirect triangle list without restart whose vertex shader has a pulling entry and no storage writes: the draw stays in the open pass, with the pulling layout, the vertex and fragment arguments as for any draw, the kernel's arguments and the draw's index source in two free slots of the vertex table (which the mesh stage shares), and `drawMeshThreadgroups`. The mesh pipeline (the draw's colour state, as `pso_desc` gives it) is built in the background; until it is ready, the draw takes the compute path. The vertex shader runs per index instead of per vertex.
- **Found on:** the same game: its 24 shadow-map vertex shaders (`zmap_vs_*.vso`) with the four kernels, 96 pairs. In the Resistance Camp on an M1 at 720p, every geometry draw became a mesh draw: translation 184 M → 122 M instructions (16.3 → 12.0 ms) a frame, GPU 12.3 → 11.3 ms, attachment traffic 463 → 415 MB a frame, frame time 21.2 → 16.6–17.4 ms.

### Mesh geometry primitive IDs restart per group

- **Symptom:** a passthrough geometry shader that uses `PrimitiveID` makes different visibility decisions or storage writes after the first 32 triangles when composed into a mesh shader.
- **Cause:** `GeometryToCompute` uses `GlobalId.x` for both the guest primitive ID and capture-buffer addressing. The mesh composer supplied a group-local lane as that ID to address local capture records, repeating guest IDs in every mesh group.
- **Fix:** `dmsh_mesh.compose` supplies the draw-wide primitive ID and rebases only capture loads to the mesh group's first record. Guest calculations and storage writes retain the original ID. A translated Maxwell fixture reads `PrimitiveID`, shifts it by five and combines it with a captured attribute to control viewport zero. Its GPU test covers 1–97 primitives, both depth conventions and non-indexed, 16-bit and 32-bit indexed draws with a negative base vertex. The original mapping fails at 33 triangles; the corrected mapping passes all 48 cases.
- **Found on:** review of [NieR:Automata](../../dm/games/010056B015FE8000/README.md) and PR #416. Its four actual `IndexBuildIndependent_*` shaders load generic attributes only, with no `PrimitiveID` or physical attribute loads. This is a general mesh-composition defect exposed by the synthetic shader, not an observed NieR rendering failure.

### Mesh geometry groups sized to the vertex output

- **Symptom:** after a pair is added to `geo_mesh.tsv`, `[RED] geometry mesh pipeline: ... Total mesh size (37248) exceeds the maximum mesh size allowed (32768)`, and the frames that draw the pair get much slower than on the compute path (Donkey Kong Bananza: translation from about 10 to 22–31 ms a frame).
- **Cause:** a mesh threadgroup took 32 triangles, 96 vertices, whatever the vertex shader outputs. Metal limits a mesh threadgroup's output to 32 KB, and 96 vertices of 388 bytes (24 `float4` varyings, the position and the point size) are 37,248 bytes. The pipeline failed in the background, but the pair kept its mesh function, so `geo_run_walk` kept ending its runs at each of the pair's draws, and every draw went through the compute path alone.
- **Fix:** `dmsh_mesh.prims_for` takes the most triangles of 32, 16, 8, ... whose vertices and primitives fit 32 KB, from the size of `main0_out` (`out_bytes`, Metal's packed count). Any count other than 32 is in the entry's name (`dm_mesh_p16`, `dm_mesh_z_p16`), so libraries of 32 keep their names and the host reads the count from the name (`geo_mesh_prims`). A triangle's kernel invocations then share a SIMD group with fewer active lanes, which the guest's subgroup operations already allow. A mesh pipeline that fails sends its pair back to the compute path for good (`geo_mesh_pso`). A mesh draw's storage writes are noted as its arguments are encoded, so it keeps a drain's list of later writes ([CPU reads after a drain wait only for later writes](rendering.md#cpu-reads-after-a-drain-wait-only-for-later-writes)). With every geometry draw a mesh draw, two per-draw lookups showed: the pair's mesh function (a key string and a dictionary lookup) and whether the pass's colour attachments can be left out (`geo_color_omit`, resolving each attachment's Metal texture). Both now keep the last answer, the second keyed by the targets with their serials, views, clears and colour mask: translation 7.20 → 6.64 ms a frame on the same route.
- **Check:** `test_dmsh_geometry.py` checks the sizes (388 bytes takes 16 triangles, the test shader 32, arrays and half types, an unknown member type and a vertex too large for one triangle refused) and composes and compiles a 16-triangle mesh shader. On Donkey Kong Bananza's scripted route, its one geometry pair as mesh draws against the compute path (`DM_NVN_GEO_MESH=0`), each run alone: translation 7.20 against 9.94 ms a frame, GPU p50 9.0 against 11.8 ms, no red lines. Frames through gameplay show the terrain, rocks, vines and props.
- **Found on:** [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md#found-on-this-game): its `geo_mesh.tsv` pair, about 600 geometry draws a frame.

### Tessellation stages are not drawn

- **Symptom:** `primitive type 14 not implemented`; patch draws skipped.
- **Cause:** SPIRV-Cross already writes the hull shaders as kernels and the domain shaders as post-tessellation vertex functions, but Metal 4's render encoder and pipelines have no tessellation (no `drawPatches`, no factor buffer). The hull kernels also write quad factors (`MTLQuadTessellationFactorsHalf`) for triangle patches.
- **Fix:** not done. The draws would need a tessellator in a compute kernel and domain shaders rewritten as plain vertex functions.
- **Found on:** [Luigi's Mansion 3](../../dm/games/0100DCA0064A6000/README.md): 15 hull and 18 domain shaders. Two tessellated programs (vertex `28a39709d0adfcfa` / hull `51d2d156bb12c260` / domain `0b95742e863bec50`, and `d4d89931b5a9a514` / `8aa6feeb7d33e9f6` / `f254b4dce2d6c012`) draw twice a frame for about 30 frames, in the title scene behind the file select as it slides out before a new game. [Bayonetta 3](../../dm/games/01004A4010FEA000/README.md): two domain shaders do not build (`DS_Vanish05_*.dso`: `SPIRV-Cross threw an exception: Subgroup builtins are not available in this type of function`); the one-lane rewrite (`spv_single_lane`, [Subgroup operations rewritten for a one-lane subgroup](#subgroup-operations-rewritten-for-a-one-lane-subgroup)) covers vertex shaders only, and tessellated draws are skipped anyway.
- **Note:** the vertex shaders that feed a hull or geometry shader can be found from the shader headers (SPH): the generic attribute map at +0x18 of the consumer's header (inputs) equals the one at +0x36 of the vertex shader's (outputs). In Luigi's Mansion 3, 40 vertex shaders match a hull shader and 49 a geometry shader.

### Ampere shader code (Switch 2)

- **Symptom:** a Switch 2 game's shaders do not translate (the translator reads Maxwell: 64-bit instructions behind scheduling words, a 0x50-byte SPH), or, rewritten as Maxwell code, UI text draws as solid boxes or not at all.
- **Cause:** NVN2 ships Ampere (sm_86) code: 128-bit instructions, a 0x80-byte SPH (version 4; its first 0x50 bytes keep version 3's layout), uniform registers UR0–UR62 and uniform predicates, convergence barriers (BSSY / BSYNC) instead of the SSY / SYNC stack, CALL / RET with the return offset in a register, and scoreboards: a variable-latency instruction (one that sets a write barrier, control bits 110..113) writes its registers when its result arrives, and an instruction that does not wait on that barrier (wait mask, bits 116..122) reads the old values. A rewrite into Maxwell code got several meanings wrong: IPA (0x326, 0x526) is IPA.PASS and the shader multiplies by 1 / a[0x7c] itself; IPA's indexed form takes its index from a uniform register; TEX's bit 79 (`.F16`) packs the results as half pairs; an `LDG.128 R0, [R0.64]` must read its address once.
- **Fix:** the translator decodes Ampere natively (`TranslationFlags.Ampere`, job `arch: sm86`; `Decoders/DecoderAmpere.cs`, `Instructions/InstEmitAmpere*.cs`) into the shared IR. Uniform registers and predicates are register types of their own; a CALL into code with no reachable RET is a jump; BRX targets come from the c1 jump table its LDC reads (BRXU's from the uniform LDC of its uniform register), up to the next table another BRX of the program reads: the compiler puts a shader's switch tables back to back in c1, and read on to the end, Donkey Kong Bananza's `SilverStarMaterial` vertex shaders took every later switch's targets, 400 and more nested comparisons, past the Metal compiler's 256 nested brackets (11 shaders). An instruction with a write barrier writes into shadow registers that are copied to its destinations before the first instruction that waits on the barrier, at DEPBAR and at the end of the block (`Translator.EmitOpsAmpere`). Field positions follow Mesa NAK's SM70 encoder, CuAssembler's sm_86 tables and the facts of other projects (Eglantine). Maxwell code takes the old path, unchanged (byte-identical SPIR-V on 7,168 shaders).
  - Semantics that differ from Maxwell: `MUFU.SIN` / `COS` take the angle in turns (the compiler multiplies by 1 / 2π first; Maxwell's came through RRO); bit 76 of FMUL / FFMA is .FMZ (a product with a zero factor is zero).
  - Warp-cooperative instructions (`InstEmitAmpereTensor.cs`): HMMA (m16n8k8 / m16n8k16, F16 sources, F16 or F32 accumulator; in a compute shader, m16n8k16 with an F16 accumulator goes through Apple's matrix hardware instead: [Ampere HMMA through simdgroup_matrix](#ampere-hmma-through-simdgroup_matrix)), LDSM (1, 2 or 4 8×8 matrices, plain or transposed) and MOVM run per thread: each lane gathers the fragments it needs from the lanes that hold them with subgroup shuffles (Metal's SIMD-groups are 32 wide, as the warp), in the layouts NVIDIA documents for PTX `mma.m16n8k16`, `ldmatrix` and `movmatrix`. HSETP2, HSET2 and HMNMX2 are the packed-half compares; MATCH (.ANY / .ALL, 32- or 64-bit) compares a value across the warp with a ballot of the active lanes and 32 shuffles, and REDUX reduces one into a uniform register the same way.
  - Geometry stage: OUT (0x124, emit / cut; none is OUTFINAL) and ISBERD (the vertex index passes through, as on Maxwell). A geometry shader that ends through an exit stub sets it up with LEPC and 0x357, and the stub runs OUTFINAL and 0x957 before its EXIT; per thread there is nothing to finish, so LEPC writes its own address and 0x357 / 0x957 do nothing.
  - NVN2 compute lays out `c[0]` per shader, from the control section: from 0x10 its push constants (+0xa9, in words), then its uniform buffer entries (+0xb1 of them) and storage buffer entries (+0xb2), 0x10 bytes each, then its texture handles (+0xb3) and image handles, 8 bytes each; the u16 at +0xba is the total (checked on all 205 of Donkey Kong Bananza's compute shaders). `dmsh_build.py` passes the counts (job `nvn2_cs`) and the frontend maps each range onto the tables the host already fills (`AMapC0`): push constants to `c[24]` (`DMSH_PUSH_BANK`), uniform entries to the compute uniform table, storage entries to the compute storage table, handles to `c[2]`. Graphics stages keep NVN's fixed tables. The push constants themselves come from `nvnCommandBufferPushGraphicsConstants` / `PushComputeConstants` (`OP_PUSH_CONST`, a `DMSH_PUSH_SIZE` block a stage group).
  - A storage address added from two constant buffer values (a `c[0]` entry plus an offset from `c[1]`, by UIMAD.WIDE) put the `c[1]` term second, and `GlobalToStorage` took the second operand as the base; it now prefers the driver segment's operand when only one is from it (Maxwell output unchanged, 28,669 shaders).
- **Found on:** [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md#shaders): 22,751 of 22,822 shaders translate; its menu renders, and its gameplay, including agl's upscaler network (`agl_ndsr.pck`, 24 compute shaders with HMMA).
- **Resolved:** 0x593 is SUATOM, an image atomic through a handle (`ASuatom`, `InstEmitAmpereMem.cs`): SUST's handle (40..59) and dimension (61..64), ATOMG's operation (87..91), size (73..76) and predicate result (81..83, PT here). Its two shaders, agl's `agl_technique_pfx` compute shaders, take an append counter with it (SUATOM.ADD.U32 on a 1D image, the old value shuffled to the warp, then their SUST), and now translate: all 22,852 of the package's shaders do. Neither has been met in play so far, so the picture is not checked; `test_dmsh_ampere_mem.py` translates the encoding. A handle in another bank or a predicate result is refused.
- **Open:** loads and stores through addresses read from memory, which had no storage buffer to become, go through the memory pools now ([Global memory at an address loaded from memory](#global-memory-at-an-address-loaded-from-memory)).

## Bindings and handles

- **Note:** the layout of NVN's driver segments `c0` and `c2` (the handle tables, the storage and uniform buffer tables, the words the host fills) and the binding kinds of the reflection JSON have one definition, `dm/gfx/nvn/dmsh_abi.h`. The host includes it (`nvn_mtl.m`, `nvn_refl.m`, `nvn_obj.h`); `dmsh_build.py` reads it through `dm/gfx/tools/dmsh_abi.py`, which stops on any line it cannot read. Both files are part of the chain version. `dm/tools/tests/test_dmsh_abi.py` checks that the tables fit together, that the host converts every kind the build may write, that the translator's own copy of the `c2` layout (`Program.cs`) agrees, and that the values do not come back as literals in the host or the build. The entries below cite the offsets.

### A store only a dispatch's first invocation makes

- **Symptom:** objects from a compute-built list flicker: each is in one of two states, on and off irregularly and independently of the others, while the draws are the same frame to frame and the camera is still. Shots of consecutive frames show the change per frame jumping between two levels.
- **Cause:** a compute shader's preamble `if (-local.x == workgroup.x << log2(local size)) buffer[k] = 0` resets an append counter on the dispatch's first invocation, while the others add to it hundreds of instructions later. Nothing orders them in the API; it works where the first workgroup starts ahead of the rest. Apple's GPUs start workgroups in no order, so a workgroup that appended before the reset lost its entries to the next one's.
- **Fix:** `first_store` in `dmsh_build.py` recognises exactly that condition, for a flat local size and constant values, and takes the store out of the shader; the reflection's `first_store` has the host make it in order before the dispatch (`nvn_mtl.m`).
- **Found on:** [Splatoon Raiders](../../dm/games/0400380020056000/README.md): `029d9c3ec1fcb39f` (`Hoian_Proc`), the plants' instance lists, the only one of the game's 9 compute shaders with atomics that has it. The change per frame over the plants: 1.1, 2.4 or 4.5 at random, to a flat 1.0-1.3; pixels that jumped and returned the next frame: up to 88 % of them, to at most 21 %. The buffers the shader reads alternate frame by frame, which turned out unrelated.
- **Check:** `python3 dm/gfx/tests/test_dmsh_first_store.py`.

### Goto elimination forcing loops around a whole program

- **Symptom:** `metal` runs for hours on a shader of about 100 MB of MSL, mostly indentation, stuck in LICM (`hoistRegion` → `ScalarEvolution::forgetValue`), then the GPU compiler's connection breaks; even `-O0` and finished libraries crash it. Hundreds of forced one-pass loops nest around the program.
- **Cause:** `MoveOutward` in `GotoElimination.cs` encloses what follows an if in an else when it moves a goto out. When another goto jumps into that range, it needs a loop of its own, and each one encloses the others.
- **Fix:** the else is skipped when a goto would have to jump into it (`HasGotoInto`, the goto's own label included). Only shaders that got a forced loop change.
- **Found on:** [Splatoon Raiders](../../dm/games/0400380020056000/README.md): 26 such loops around the whole program in its uber shaders: the mid-size shader went from 176 loops to 0, the largest from 364 to 5, and 8 workers stalled for 5 h 40 before that. Three test shaders then compile in 89 s.
- **Check:** `DMSH_TEST_XLATE=<translator> python3 dm/gfx/tests/test_dmsh_control_flow.py`.

### Ampere jump tables read past their end

- **Symptom:** `BRX with no target`, or code about 20 % larger than it should be.
- **Cause:** a BRX's table in c1 stores no length. The decoder read entries until one failed, so a table ran on into the next one; a target of 0 counted as a jump to the start of the program; and it looked only on one path for the `LDC` that loads the table's base, so a base loaded before an if / else join was not found.
- **Fix:** `FindBrxTargetsAmpere` reads a table one new target per pass, so the next table is known before this one overruns it, and stops at the next known base; an entry of 0 is filler for unused cases. `FindBrxTableAmpere` walks every predecessor and takes the base only when all paths reach the same `LDC`.
- **Found on:** [Splatoon Raiders](../../dm/games/0400380020056000/README.md): 18 `Hoian_UBER` vertex shaders, 5 of them used at boot.
- **Check:** `DMSH_TEST_XLATE=<translator> python3 dm/gfx/tests/test_dmsh_control_flow.py`.

### MUFU.F16

- **Symptom:** `Ampere instruction 308` refused; a neural upscaler's programs missing, so the frame never reaches the window (a black 3D scene).
- **Cause:** MUFU has an F16 form (bit 73): the source is the half that bits 60..61 select (bit 62 its abs, as in every B slot), and the result goes back as the low half of the register, the high half zero. MUFU is the multi-function unit (op 2 EXP2, 3 LOG2, …), not a multiply.
- **Fix:** `AMufu` in `InstEmitAmpereAlu.cs`.
- **Found on:** [Splatoon Raiders](../../dm/games/0400380020056000/README.md): 10 compute shaders of `agl_ndsr.pck`.
- **Check:** `DMSH_TEST_XLATE=<translator> python3 dm/gfx/tests/test_dmsh_ampere_mem.py`.

### Ampere HMMA through simdgroup_matrix

- **Symptom:** a Switch 2 game's frame rate swings between 60 and 30 fps (Donkey Kong Bananza's own controller, `main+0x204a18`, drops to 30 after two frames of 20 ms from their start to the GPU's end, and comes back after six), and the GPU profile (`DM_NVN_GPUPROF`) shows one compute dispatch of 2-6.5 ms a frame: agl's upscaler network (`agl_ndsr.pck`, shader `7629d857a0a93833`).
- **Cause:** its 144 HMMA.16816 a thread (F16 sources and accumulator) were emulated lane by lane: 32 subgroup shuffles and 64 FMAs a lane for each, 4,922 shuffles in the shader. Without its shuffles (each returning its own value; a wrong picture, measured only) the dispatch took 0.24-0.40 ms. On this GPU shuffles cost more than threadgroup memory: one HMMA a SIMD-group took 5.42 ns emulated, 1.08 ns through threadgroup memory and `simdgroup_matrix`, 0.65 ns with the tiles loaded straight from where LDSM reads them, and 2.73 ns moving the register fragments into `simdgroup_matrix` elements with shuffles (Apple's 8×8 layout gives each lane a pair of columns of one row too, but in another order). Metal's matrix types have no SPIR-V form spirv-cross converts (its cooperative matrix support is GLSL only).
- **Fix:** in a compute shader, HMMA.16816 with an F16 accumulator becomes `MatrixMultiplyAdd` (`AHmma`), then a call of the helper `dmsh_hmma16816_f16` (`TensorPass`, `HelperFunctionManager.GenerateHmma16816F16`), named in the SPIR-V. Each SIMD-group has 256 words of a shared memory of its own (`hmma_scratch`, 1 KB a SIMD-group): the lanes store their fragment words as they are, A row-major, B transposed and C row-major, and after a SIMD-group barrier (the new IR instruction `SubgroupBarrier`, Metal `simdgroup_barrier`) the helper's body computes D from them lane by lane. `dmsh_build.py` (`hmma_msl`) replaces that body in the Metal source with `simdgroup_load` (B with its transpose flag), four `simdgroup_multiply_accumulate` and a `simdgroup_store` into the same block; the name, the arguments and the block are the contract, and the translator's body stays correct where the rewrite does not apply. Other HMMA forms (m16n8k8, an F32 accumulator) keep the per-lane emulation. A chain change.
- **Check:** `dm/gfx/tests/test_dmsh_hmma.py`: the rewrite of the helper as spirv-cross writes it, and on the Mac's GPU the rewritten body against a CPU reference of PTX `mma.m16n8k16` on random fragments (4 SIMD-groups a threadgroup; largest error 0.009 against a reference that rounds once: the matrix hardware rounds the F16 accumulator as it goes, as Ampere's does). The translator's body agreed with the same reference to 0.002.
- **Found on:** [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md): its 24 NDSR compute shaders. On one save and route, the upscaler took 1.1-1.6 ms a frame instead of 0.1-5.1 ms (2.5-4.5 typically), a frame's GPU span 8.9-19.4 ms instead of 13.8-23.3, and the game's 60 ↔ 30 switches fell from 87 to 68 in the same run length; the picture is unchanged.
- **Note:** three of four of these HMMAs start an output tile, with C = 0. A variant for those calls that neither stores nor loads C (`simdgroup_multiply` first) was measured and not kept: 0.767 against 0.820 ns a warp-HMMA (6 %) in an isolated loop of the build's two bodies on an M4 Pro, about 0.04 ms of the upscaler's 1.6-1.9 ms a frame, and no measurable change in the game.
- **Open:** LDSM and MOVM (27 and 72 a shader) are still emulated with shuffles (314 left in the upscaler, whose floor without them is about 0.3 ms); loading A and B straight from LDSM's source and keeping D as matrices between HMMAs (0.65 ns an HMMA) needs the translator to follow LDSM → MOVM → HMMA → HMMA chains.

### Driver-segment reads outside fragment shaders

- **Symptom:** a vertex or compute shader reads zeros from NVN's driver segment `c0`.
- **Cause:** `dmsh_build.py`'s `drv_reads` scan matches only fragment-shader names (`fp_c0` / `fp_c2`), so vertex and compute shaders that read `c0` were never marked, and the host passed them an all-zero `c0`; these reads are not listed in the JSON.
- **Fix:** the host treats a declared driver-segment binding as read (`nvn_refl.m`). The words the host fills are in the entries below.
- **Found on:** [The Legend of Zelda: Link's Awakening](../../dm/games/01006BB00C6F0000/README.md), [Super Mario Odyssey](../../dm/games/0100000000010000/README.md), [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md).

### Compute `c0`: dispatch size and uniform buffer table

- **Symptom:** flickering hard-edged white tiles (tiled lighting with every tile row getting row 0's or other rows' lights).
- **Cause:** a compute shader that computes its tile index as workgroup y × `c0[0x200]` + workgroup x reads the dispatch's workgroup count from the driver segment, which the host left zero.
- **Fix:** for compute, `c0` carries the dispatch size at 0x200 (x, y, z) and the uniform buffer table at 0x210 + 16 r (address, size) in front of the storage buffer table at 0x310 (`encode_args_`).
- **Found on:** [The Legend of Zelda: Link's Awakening](../../dm/games/01006BB00C6F0000/README.md) (Mac and iPhone): the step of Marin and Tarin's doorway. The room shaders use tiled forward lighting: the fragment position picks a 64×64-pixel tile (`trunc(y / 64) × ceil(width / 64) + trunc(x / 64)`), whose light count and light indices come from two storage buffers that a compute shader (`light_culling_compute.bfsha`, one thread per workgroup, one workgroup per tile) fills every frame. Every row's workgroups wrote row 0's entries, the last writer won (a different one each frame), and the bottom tile row, where the step is, got the lights of other rows; a dump of the tile lists showed only the first of 17 rows filled. After the fix every tile row has its lights and the doorway shows the game's soft light from outside.
- **Check:** Breath of the Wild (101 compute and 4 vertex shaders read `c0`) and Super Mario 3D World (11 and 1) render the same as before on the Mac.

### Geometry-stage storage buffer table

- **Symptom:** `[RED] shader ... reflection conversion failed: B binding gp_s0 has no resolved slot number`; a geometry shader reading driver segment `c0` reports that its storage-buffer table base was never measured.
- **Cause:** geometry shaders access their storage buffers through `c0 + 0x410 + 0x10 r`, but the builder's `SSBO` map and replay's `c0_ssbo_base` covered only vertex, pixel and compute. Running the lowered geometry shader as a Metal kernel does not change its logical NVN stage or driver table.
- **Fix:** geometry uses base `0x410` in both reflection and driver-buffer population. `encode_args` keeps NVN stage 2 for these kernels, so storage slots and driver-table rows come from the geometry bindings. Hull and domain bases remain unmeasured and unresolved.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): the original shader's constant-buffer operands and the pinned translator's metadata agree on addresses in rows 0 and 2.
- **Check:** `dm/gfx/tests/test_dmsh_reflect.py` resolves synthetic geometry descriptors to slots 0 and 2 and rejects unknown or unaligned table addresses. `t_geometry` passes synthetic metadata through native reflection and the production argument encoder; GPU readback checks both rows' guest addresses and sizes with different compute-stage bindings, plus a write through the geometry storage pointer.

### Storage-buffer write flag in translation metadata

- **Symptom:** a storage buffer the shader stores to or uses atomically has `rw: false` in its generated reflection; replay does not record its GPU writes for dependent buffer reads.
- **Cause:** Ryujinx's `BufferUsageFlags.Write` is bit 0 (`1`), and `Program.cs` serializes that enum unchanged. The builder tested bit 1 (`2`), incorrectly marking every such descriptor read-only. Native reflection honors the explicit `rw` field, so its older fallback to the MSL pointer qualifier could not correct it.
- **Fix:** `reflect` tests bit 0 for both driver-table storage buffers and addresses carried in guest uniform buffers. Replay then uses its existing write-access and pending-write tracking paths. Rebuild shaders to replace the old metadata.
- **Found on:** [NieR:Automata The End of YoRHa Edition](../../dm/games/010056B015FE8000/README.md): its geometry counter atomic and output stream both emitted translator flag `1` but reached reflection as read-only.
- **Check:** `test_dmsh_reflect.py` exercises writable and read-only translator descriptors, including an address in a guest uniform. `t_geometry` verifies that native reflection preserves explicit true and false flags and that the writable geometry output enters the GPU-write tracker before its dispatch.

### Graphics `c0 + 0x610`: the largest point size

- **Symptom:** point sprites of size 0 (Odyssey's Cap Kingdom fog never parted around Mario).
- **Cause:** a vertex shader that writes `gl_PointSize = min(size, c0[0x610])` read 0, because the host had nothing at that offset of the driver segment.
- **Fix:** `c0 + 0x610` holds the largest point size, and the host writes 2047 there for the graphics stages. Writing the `SetPointSize` value instead (always 1 in Odyssey) left one-pixel splats whose disc weight came out 0 (the fields held `-0.0`).
- **Found on:** [Super Mario Odyssey](../../dm/games/0100000000010000/README.md) (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 simulation fields stayed exactly zero (raw dumps with `DM_NVN_TEXDUMP` + `DM_NVN_SNAPRAW`; the 8-bit dump shows velocities as black either way). After the fix the fields hold velocities up to ±44, and the fog clears around Mario.

### Graphics `c0 + 0x644`: the draw's base instance

- **Symptom:** occlusion queries that always return 0 samples (Z-A: characters' legs folded back from the hips once they walked); every instanced draw's transform taken for instance 0.
- **Cause:** a shader that picks its transform with (instance index − base instance) + `c0[0x644]`: Maxwell's instance ID leaves the base instance out, and the NVN driver keeps it in its constant segment, where GLSLC reads it for `gl_InstanceID + gl_BaseInstance`. The host left `0x644` at 0.
- **Fix:** `encode_args` writes the draw's base instance there for graphics stages and does not reuse a stage's argument block across draws with different base instances. The word before it, `0x640`, is added to the vertex ID by 58 of Z-A's shaders and stays 0 until what the driver puts there is known.
- **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md): from the first step on, the player's legs stuck straight out behind her at hip height, and the rival's forward in a later cutscene; standing still beforehand looked right. The bone palette the game computes on the CPU was already bent. The cause was six occlusion queries a frame, one 36-vertex box each (`a793f436dcabd350`, base instance 61, 62, …), drawn against mip level 2 of the depth pyramid: every box was box 0, every query came back with 0 samples, and the game took its characters for hidden and posed their legs that way. Forcing every result to visible fixed the legs, and depth testing with ALWAYS did not. 793 of the game's vertex shaders read the word.
- **Note:** how it was found: the raycasts were logged with scratch observation gates (all correct), then `DM_HLECOUNT` showed nothing walking-specific, then the counter reports.

### The base instance of indirect draws comes from Metal

- **Symptom:** flat grey or background-coloured triangles with straight edges across part of the scene, changing shape from frame to frame while the camera stands still; dropping the colour pass's fragment shaders in groups (`DM_NVN_SKIP_FS`) did not remove them, as the instances are misplaced in the depth pre-pass too, whose depth the colour pass tests for equality.
- **Cause:** for an indirect draw (`DrawElementsIndirect` / `DrawArraysIndirect`) the base instance is in the argument record, which the game's GPU culling writes earlier in the same submission; `encode_args` could only write the command's base instance into `c0 + 0x644`, 0 for an indirect draw. A vertex shader whose instance is (instance index − base instance) + `c0[0x644]` then read its transforms from the first records for every indirect draw.
- **Fix:** the translation reads the word from Metal's `[[base_instance]]`, which holds every draw's base instance, direct, indirect or masked (`dmsh_build.py` `vs_base_instance`: in vertex functions `c0[100].y` becomes the base instance's bits, and `main0` gets the parameter when it lacks it). `encode_args` still writes the word for the other stages.
- **Found on:** [Super Mario Party Jamboree](../../dm/games/0100965017338000/README.md): every board. Its culling (`2b5485a2a731a231`, 62 points with rasterization off) writes 80-byte transforms per visible instance into a 30 KB buffer and counts them, `276504dc5c987ecf` writes nine indirect records from the count, and the instanced draws (`5a805bbb08421b86` in the depth pre-pass, `4b05556737e9535f` and `eaa3a44ba28f0744` in the colour pass) draws them, both reading record = instance + `c0[0x644]`. 5,249 of its vertex shaders read the word.

### Pixel `c0 + 0x620`: the sample index and GLSLC's sample position table

- **Symptom:** `[RED] shader … reflection conversion failed: reads driver-bank locations at run time that the replay layer cannot fill: c0 dynamic index c0+0x630`, and that shader's draws are missing; the build notes `Shader instruction Pixld is not implemented` for pixel shaders; `[RED] NVN unimplemented function called: nvnProgramSetSampleShading`.
- **Cause:** GLSLC reads the sample index with `PIXLD.MYINDEX` (RZ, offset 0; nouveau reads `gl_SampleID` the same way), which Ryujinx's translator left unwritten, so the register read as 0. For `gl_SamplePosition` (or interpolation at a sample) it then indexes a byte table the NVN driver keeps at `c0 + 0x620`: index = ((((y & w) << z) + (x & z)) << log2 samples) + (sample & mask), with {mask, log2 samples, z, w} at `c0 + 0x630` and x, y the pixel (`gl_FragCoord` over the render scale); each byte holds the sample's x (low nibble) and y in 1/16 pixel. The replay filled neither row, and the build refused any computed index into `c0`.
- **Fix:** the translator writes the sample index for `PIXLD.MYINDEX` (`IoVariable.SampleId`, SPIR-V `SampleId` with `SampleRateShading`; MSL `[[sample_id]]`, so the fragment function runs per sample on a multisampled target, and SPIRV-Cross moves `gl_FragCoord` to the sample); the other modes still log. `encode_args` fills both rows for the pixel stage from the pass's sample count (`c0_sample_rows`: Metal's default positions, the same for every pixel, so the grid masks are 0; one sample is 0x88, the pixel centre), and `dmsh_build.py` counts them as filled for pixel shaders (`SAMPLE_ROWS`), with a computed `c0` index accepted when the shader reads the 0x630 row. `nvnProgramSetSampleShading` returns true: Metal has no switch to force per-sample shading, which a shader gets by reading the sample index.
- **Found on:** [Metroid Prime Remastered](../../dm/games/010012101468C000/README.md): 438 of the 562 pixel shaders of `MaterialArchive.arc` read the sample index; in 30 the translated shader keeps it, and there it reaches the table, whose byte the G-buffer shaders pack into an integer target (`04023899305e93e1`). Its targets are single-sample (1600×900), so sample 0 was right and only the table mattered: 0 there put every sample in the pixel's corner. No other game here reads either row (58 games' libraries checked). The nibble order for several samples is the hardware's, not seen in a game.

### Texture and sampler pools with overlapping IDs

- **Symptom:** a menu's labels drawn as plain bars, other menus empty except for the cursor (one pool's textures overwritten by another's); a black screen when IDs past a fixed table read as empty slots.
- **Cause:** a game that binds two texture pools and two sampler pools whose IDs overlap; the replay had one ID table.
- **Fix:** each pool 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` in `hle_nvn.c`, the base kept in `cmd_t`, [next entry](#bases-of-bound-handles-kept-in-the-command-buffer); log line `pool #2 (…): numbered from N`). Handles read from uniform buffers take the base of the pool bound at replay (DARK SOULS: REMASTERED, `ubo_tex_id` / `ubo_samp_id`, [Handles from uniform buffers take the bound pools' bases](#handles-from-uniform-buffers-take-the-bound-pools-bases)). Deferred binds are not remapped (red line; [Deferred texture binds read at execution](#deferred-texture-binds-read-at-execution)). The table's capacity is the sum of the pools.
- **Found on:** [Fire Emblem: Three Houses](../../dm/games/010055D009F78000/README.md): 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 scene's default 8×8 white texture overwrote the UI's textures (the title menu's labels drew as plain bars). 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.

### Bases of bound handles kept in the command buffer

- **Symptom:** red lines from the list of command buffers that bound a non-first pool (6,734 in one DARK SOULS: REMASTERED run).
- **Cause:** that list had 64 entries and was never cleared, and the game binds a non-first pool in more than 64.
- **Fix:** the bases of bound handles live in the command buffer (`cmd_t` `tex_base` / `samp_base`).
- **Found on:** [DARK SOULS: REMASTERED](../../dm/games/01004AB00A260000/README.md).

### Handles from uniform buffers take the bound pools' bases

- **Symptom:** distant walls and floors smeared sideways (DARK SOULS: REMASTERED): tiling textures took a clamp-to-edge sampler and repeated their edge column.
- **Cause:** the map materials read their combined texture / sampler handles from a uniform buffer (`fp_t_cb13_*`, resolved at replay by `ubo_handle`), and the game keeps several sampler pools, each with its own samplers under the same IDs, one per command buffer that records the scene. The replay resolved uniform-buffer handles in the first pool.
- **Fix:** the replay keeps the texture and sampler pools bound (`OP_SET_TEX_POOL` / `OP_SET_SAMP_POOL`, which it used to skip; a command buffer that binds none keeps the queue's, as on the GPU) and adds their bases to handles read from uniform buffers (`ubo_tex_id` / `ubo_samp_id` in `nvn_mtl.m`).
- **Found on:** [DARK SOULS: REMASTERED](../../dm/games/01004AB00A260000/README.md): the map materials `FRPG_Phn_*Lit*`; 12 sampler pools of 2,304 entries (ID 261 is clamp-to-edge with LOD 0..0 in one pool, repeat with mips in others); the walls' tiling base textures (u ≈ −3.45).
- **Check:** found by highlighting the six most common fragment shaders of the scene pass (`DM_NVN_HILITE`), then the next six, and reading the wall draw's binding lines (`DM_NVN_PASSDUMP_ALL`).

### Deferred texture binds read at execution

- **Symptom:** the 3D scene black, only particles and the HUD drawn; pass dump `B` lines all `tid=0 dummy`.
- **Cause:** `nvnCommandBufferBindTexturesDeferred` passes a pointer to the handle, and "deferred" means the GPU reads the handle from that address *when the command executes*. A game that records its command blocks once, submits them every frame and writes the handles afterwards got the handles the host copied at recording, mostly zeros, so draws read the dummy texture 0.
- **Fix:** the bind record keeps the source address (`p_bind_handles.defer / src`) and every replay-time reader refreshes the handles first (`bh_fresh` in `nvn_mtl.m`).
- **Found on:** [The Legend of Zelda: Breath of the Wild](../../dm/games/01007EF00011E000/README.md): the engine binds almost every texture this way (`main+0xec2e64`, one handle, the pointer being the handle field inside the texture-sampler object); the deferred-lighting composite (one full-screen triangle per material class, blending into the HDR target) read texture 0 for all of its inputs.
- **Note:** how it was found: a frame data-flow table built from `DM_NVN_PASSDUMP_ALL` (per pass: attachments, shader keys, sampled serials) showed the half-resolution lighting outputs never being read, and the composite pass's `B` lines were all `tid=0 dummy`. Without `_ALL` the dump lists only the first draw of each pass, which hid the composite. Deferred binds are not remapped by pool base ([Texture and sampler pools with overlapping IDs](#texture-and-sampler-pools-with-overlapping-ids)).

### Handles in uniform buffers written earlier in the same submission

- **Symptom:** bindless handles read from a uniform buffer name the wrong texture when the game writes them earlier in the same submission.
- **Cause:** the replay reads handles on the CPU, and a write queued earlier in the submission has not landed yet.
- **Fix:** inline uniform updates keep where their data sits, so the handle is read from the newest one covering it (`pend_read`). The old limit of 256 writes per submission, after which every address counted as written, is gone: pending writes are now indexed by 4 KB page (8,192 per submission).
- **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md) (it queues more than 256 writes per submission).

### Storage buffer addresses in the game's own uniform buffer

- **Symptom:** a storage buffer reported with `cb_slot` 3, which had no NVN slot.
- **Cause:** the shader reads a storage buffer's address from the game's own uniform buffer (c3 + 0x38), which Ryujinx reports with `cb_slot` 3.
- **Fix:** the build marks it `sbuf_ubo` (uniform slot `cb_slot - 3`, byte offset), and the replay reads the 64-bit address there at bind time, as it reads bindless handles.
- **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md) (about 30 compute shaders).

### Global memory at an address loaded from memory

- **Symptom:** whatever a shader writes through a pointer it read from memory never happens. The build summary counted `Failed to find storage buffer for global memory operation "Store"`, and the shader's MSL had a store function with an empty body. In Donkey Kong Bananza the tunnel's walls and floor were black and only the loose rocks drew.
- **Cause:** Ryujinx turns a global memory access into a storage buffer access by tracing its address back to a constant buffer entry (`GlobalToStorage`). An address the shader loaded from memory traces to none, and Ryujinx dropped the access: a store did nothing, a load read 0. Donkey Kong Bananza's gsys vertex shader `c9a56f796ea4a4cd` draws 96 points every frame. Each point loads one indirect draw record's 64-bit address from its storage buffer and writes {0, 1, 0, 0, 0} there (count 0, one instance) before the terrain's culling geometry shader adds its triangles to the count. Without those stores the records kept instance count 0, and their counts grew frame after frame. NVN2 draws instance count 0 as nothing. Its indirect draw (`sdk+0x6049b0`) calls macro 27 (driver constants at 0x640) and macro 40 with the record's five words. Macro 40, decoded with Mesa's MME decoder after running the driver's macro upload (`sdk+0x60fd40`) in an emulator, skips its draw loop when the instance count is 0 and never writes the record.
- **Fix:** general, in the chain. Accesses traced to one or more storage buffers translate as before; only those traced to none change:
  - **Translator** (`GlobalToStorage.TryGenerateRawStorageOp`): reserves the host's pool table as a storage buffer under a sentinel constant buffer slot (`DMSH_KIND_GMEM_MAP`, `dmsh_abi.h` `DMSH_GMEM_MAP_*`). It finds the pool holding the address by a branchless binary search over the table's starts, adds that pool's delta and accesses the result through a physical storage buffer pointer. In MSL that is `reinterpret_cast<device uint*>(ulong)`.
  - **Why a delta:** a no-copy Metal buffer's `gpuAddress` is not its CPU address, so the guest address has to move by the pool's delta (GPU address minus guest address).
  - **Host table** (`nvn_mtl.m` `gmem_map_gpu`): built from the memory pools whenever they change. Every pool is made resident while such a shader is in use.
  - **Write log:** where a store through a loaded pointer lands is known only once the GPU has run it. Counting it as a write anywhere refused every uniform shadow of the submission: Donkey Kong Bananza's tunnel ran at about 70 ms a frame, with 4,533 passes ended to land deferred updates. So each store or atomic also puts its guest page into a write log the host hands each draw (`DMSH_KIND_GMEM_LOG`, `gml_take`). The log is a hash set of pages filled with compare-and-swap, not a list: one of those draws stores 7,325 times, into a few dozen pages.
  - **Learned pools:** the host maps each logged page to the memory pool holding it. The stores count as a write anywhere until the learned pools have stayed the same over 8 logs, then as writes over those pools. A later log that shows a page in another pool adds that pool, without going back to a write anywhere (only a log that ran out of slots starts the learning again): a write anywhere makes every later geometry draw of the submission wait for the GPU (`geo_cpu_hazard`), and going back for 8 logs after each miss gave Donkey Kong Bananza streaks of 300-2,000 ms frames (38 such frames in 5,400, 6 after the change, the geometry wait in one 25 s sample 3,490 → 264 ms). The host then checks whether the miss could have mattered: every uniform shadow taken from that submission on (its own and those of the later submissions encoded while it ran, each recorded with its fence) is compared with the pages the stores wrote outside the prediction. Only an overlap, or a record that has wrapped before the check, prints a `[RED]` line; a miss no shadow saw is logged with the pools' ages.
  - **Why pools, not pages:** the game allocates new records as the level streams in, in the same pools. Learned page by page, the tunnel still met a new page 213 logs in.
  - **Result in Donkey Kong Bananza:** 4 pools learned over 9 logs at first, 60 fps (p50 16.64 ms), and no pass ended to land deferred updates. Further into the level the stores keep reaching pools the set does not hold yet: 12 in a 17,700-frame session, 10 in a scripted 13,500-frame run. Those pools are not new: they were added 3.5 to 294 s before the miss (the game streams records into pools it made earlier), so a new pool would not have announced them. None of the 7 misses checked had a page under a shadow (3 to 5 shadows each).
  - **Still dropped:** an address in no pool. Atomics on 8- and 16-bit values are not implemented.
- **Open:** the prediction is checked, not proven: a miss that does reach a shadow is reported (`[RED]`) after the frame, not prevented. The records' pointers come from memory, so the pools they will reach are not known ahead of the draw.
- **Found on:** [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md): its terrain, and agl's compute shader `49bc2268925e25b6` (one load); rebuilding the chain changed only these two of Donkey Kong Bananza's shaders. Vertex shader `588e242a52f5f7aa` (five stores) has one too, the same shader in [The Legend of Zelda: Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md) and [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md), so likely one the SDK ships; neither drew it in its regression route.

### Bindless handles through a loop phi

- **Symptom:** `no resolved slot number`, the shader does not load; the translation has a 32-entry texture-pool array (`fp_t2d_cb0_40000000`, handle type Direct).
- **Cause:** the handle is `c2[0x1A0] | 0x10500000` (separate texture slot 7, constant sampler 0x105), kept in a loop variable. Ryujinx's bindless elimination compared only the phi's direct inputs, gave up, and turned the access into an index into a 32-entry texture-pool array, which the replay cannot bind.
- **Fix:** the vendored translator now follows the phis that feed a phi, and treats two `constant buffer | constant` handles as the same value ([third_party/ryujinx/README.md](../../dm/third_party/ryujinx/README.md)). Those accesses become ordinary type-3 handles.
- **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md) (218 pixel shaders).

### Bindless reads whose handle phi also carries other values

- **Symptom:** none in the log: texture reads come back 0. Particle meshes collapse to a point (their scale reads 0), so effects are missing or show as dark patches. The translator's own message (`Failed to find handle source for bindless access of type "textureBuffer"`) used to be dropped, because `dmsh_xlate`'s accessor had no `Log`; its messages now go into the shader's notes, marked `translator: `, and the build summary counts them apart from unanswered queries.
- **Cause:** the handle reaches the read through a phi. Some inputs are c2 handle words, the others are not handles: undefined, or a value the compiler kept in the same register on a path where the predicated read does not run (float arithmetic, a texture result, the absolute value of a uniform word). Some phis merge two or more different c2 handles too (a float and a half-precision buffer). `TryConvertBindless` cannot trace one handle, and `GenerateBindlessAccess` refuses a phi with such an input, so bindless elimination deleted the read and wrote 0 to its results.
- **Fix:** only for a phi `GenerateBindlessAccess` refuses (`TryTableHandle` in the vendored `BindlessElimination.cs`, marked `nxnative:`), the read takes the c2 handles among the phi's inputs (a word, or a texture word ORed with a sampler word or a constant sampler). With one handle, it takes the phi's place. With several, one read is made per handle and the result is the read whose handle equals the phi's value at run time: the host fills c2 with the bound handles, or with slot indices for `tex_dyn_slots`. A candidate that other reads of the shader use as a texture of another shape is left out, so those reads are converted in a second sweep once every other read has registered its texture (`ResourceManager.GetTextureTypes`). An image store or atomic with several candidates is left alone, since one per candidate would write each of them. A phi made only of handle words keeps `GenerateBindlessAccess` (the dynamic arrays of [Texture slots picked at run time](#texture-slots-picked-at-run-time)), so no other shader changes: in FINAL FANTASY RESONANCE, the SPIR-V of the 69,449 other shaders is byte for byte what it was. A chain change.
- **Found on:** [FINAL FANTASY RESONANCE](../../dm/games/0100606026F12000/README.md): 58,122 reads removed (57,861 of them texture-buffer reads, mostly in Unreal's Niagara particle vertex factories), of which 58,096 in 3,197 shaders are now resolved: 21,380 with one handle, the rest choosing among 2 to 5 at run time. 26 reads in 14 shaders, whose handle comes from no c2 word, are still removed. Reported in GitHub issue #168.

### Texture slots picked at run time

- **Symptom:** the shaders' bindless arrays have no slot and used not to load; the translator falls back to a 32-entry array such as `array<texture_buffer<float>, 32> vp_tb_cb0_40000000`.
- **Cause:** the shader reads one of several c2 texture-table words, chosen by branches. With separate textures (Octopath Traveler 0) it ORs the word with a sampler-table word; with combined ones (Paper Mario) the handle is read from the combined texture table (c2 `0x20 + 8 × slot`, constant offsets only). The translator cannot trace that to one slot and falls back to Ryujinx's bindless arrays, `array<texture_buffer<float>, 32> vp_tb_cb0_40000000` and `array<sampler, 32> vp_s_cb0_40000000`, indexed by the handle's texture ID (low 20 bits) and sampler ID (high 12), capped at 31; the build leaves them without a slot.
- **Fix:** for buffer textures (and images). A sampled texture of another kind indexes the whole texture pool instead, whatever the handle's source ([Texture handles from game memory index the pools](#texture-handles-from-game-memory-index-the-pools)), which took over Paper Mario's cube and 2D arrays. With the profile key `tex_dyn_slots` the reflection reads them as `NVN_RB_TEXD` / `NVN_RB_SAMPD` arrays (`nvn_refl.m`), the replay writes slot indices into that shader's c2 tables (texture slot r → r, sampler slot r → r << 20; in the combined table `r | r << 20` in slot r) and binds slot k at element k, so the shader's own arithmetic lands on the binding: the handle names element r, which holds the texture bound in slot r. The array's other elements hold whatever the other slots do, mostly 2D textures where the array is of cubes; an array element bound to a 2D texture gets the zero cube without a type mismatch report (`tex_fit_elem`), since the shader does not pick it there. The host logs each such shader's first encode with the slots that held textures. Host only: the chain version does not change.
- **Check:** check first that the handles come only from c2's tables: in Octopath Traveler 0 every integer read of c2 in those shaders is a table slot (texture slots 0-10, samplers 0-4).
- **Found on:** [Octopath Traveler 0](../../dm/games/01005270232F2000/README.md): 61 Unreal vertex shaders, particles reading their data from buffer textures (separate textures and samplers). [FINAL FANTASY RESONANCE](../../dm/games/0100606026F12000/README.md): 2,077 particle vertex shaders. [Paper Mario: The Thousand-Year Door](../../dm/games/0100ECD018EBE000/README.md): 2,386 pixel shaders (5,476 cube samples, 487 2D), slots 0, 7, 8, 13 and 28 (combined table), until the pools took them over.

### A texture chosen at run time between two uniform-buffer handles

- **Symptom:** Ryujinx's fallback arrays (`fp_t2d_cb0_40000000`) in pixel shaders whose handle comes from the game's own uniform buffers; the shaders do not load and their draws are skipped.
- **Cause:** the handle is one of two words of the game's own uniform buffers, chosen by a flag, packed as texture ID | sampler ID << 20. `tex_dyn_slots` covers only handles from the driver's c2 tables ([Texture slots picked at run time](#texture-slots-picked-at-run-time)).
- **Fix:** the pools: whichever word the flag picks, its texture ID indexes the texture pool ([Texture handles from game memory index the pools](#texture-handles-from-game-memory-index-the-pools)). Not checked on Big Brain Academy, whose files are not on the Mac the change was made on.
- **Found on:** [Big Brain Academy: Brain vs. Brain](../../dm/games/0100620012D6E000/README.md): ten pixel shaders (six in `03GamePerception.bnbshpk`, two PBR, two in `main`); in `e10f24835bbaa5f6` the handle is c3 +0x130 or c12 +0x1d0, chosen by a flag in c12. None of the activities checked uses them; a long run of activities met one once.

### Texture handles from game memory index the pools

- **Symptom:** `[RED] shader … reflection conversion failed: a texture handle the translator could not trace to one slot (profile key tex_dyn_slots off)`, then `failed to load`, and the draws are skipped: in Mario Party Superstars every character and NPC in a lit scene is a flat grey silhouette (the depth pre-pass draws, the colour pass does not), and the opening scene's pipe is grey.
- **Cause:** a lit pixel shader loops over the lights, reading each light's shadow map, cube map and sampler handles from the game's storage buffer (`fp_s1`) at a computed offset. Ryujinx cannot trace such a handle to a slot and falls back to arrays indexed by the handle's IDs (`array<texturecube<float>, n> fp_tcube_cb0_40000000`, `array<sampler, n> fp_s_cb0_40000000`, texture ID in bits 0–19, sampler ID from bit 20), clamped to `n − 1`. The translator's accessor gave `n` = 32 and the replay had no pool to bind there: only `tex_dyn_slots` arrays, whose handles must come from c2.
- **Fix:** a chain change. The accessor sizes these arrays to the whole ID ranges (`Program.cs` `QueryTextureArrayLengthFromPool` 2²⁰, `QuerySamplerArrayLengthFromPool` 2¹²) for sampled textures other than buffer textures; buffer textures and images keep 32-entry slot arrays (vendored `BindlessElimination.cs`, `NvnSlotArrayLength`). `dmsh_build.py` `pool_arrays` turns each such array into a pointer to `struct spvPool_<name> { uint n; array<element, 1> e; }`, read as `(&e[0])[min(index, n)]`, and renumbers the argument IDs (a pointer takes one). The struct declares one element because Metal's argument encoder walks every element of a declared array (`AGX::IndirectArgumentLayoutGen2::parseArray`): with 2²⁰ textures each pool shader's encoder took 57 ms and 66 MB, and the prewarm of 733 pipelines left the Mac host at 12.5 GB resident (9.7 GB of large mallocs); the elements are 8-byte resource IDs, so the pointer reaches the others (checked on the GPU with a test kernel); the reflection names them `tex_pool` (with the element type) and `samp_pool`. The host keeps one buffer per element type and pool base (`nvn_mtl.m` `pool_gpu`): the bound n, then from `DMSH_POOL_HEAD` the resource ID of each texture (sampler) ID of the pool bound at the draw, the dummy where an ID holds nothing or a texture of another type and at n. It is rebuilt when a registration or a texture's storage changes (`g_reg_gen`, `g_ser_gen`), the old buffer released after its submissions, and an argument block holding a pool address is not reused past a rebuild. Handles read from game memory are the bound pool's IDs, so the pool takes that pool's base (as `ubo_tex_id`), and a pool shader's c2 holds the bound handles without the base `bind_handles` added, as the game wrote them. Every texture is already in the residency set. A walk over every ID also reaches IDs no draw binds, which a game leaves registered after it is done with them: the first walk once crashed in `nvn_sparse_protect` making such a texture (a garbage `pool->vp`; the ID was not caught). So a texture no draw has made is made only while its storage exists (`tex_storage_live`: the memory pool object still holds its magic, and an ordinary pool's storage is in a pool the host maps), and an ID whose memory pool the game finalized is skipped until it is registered again (`treg_t.gone`, set in `treg_drop_pool_views`). Every other route is covered too: [A texture read after its memory pool object was freed](nvn-objects.md#a-texture-read-after-its-memory-pool-object-was-freed).
- **Found on:** [Mario Party Superstars](../../dm/games/01006FE013472000/README.md): 6,370 of its 82,790 shaders, every character and NPC material (`_chara/pc/*.bnbshpk`, `_chara/npc/*.bnbshpk`) and some board variants (`hsmn_bg.bnbshpk`); a 2,304-entry texture pool (1,286–1,314 2D textures, 65 cubes) and a 512-entry sampler pool, about 0.1 ms per rebuild once the textures exist.

### An array of handles in a uniform buffer

- **Symptom:** the host stops at the shader's first load: `unexpected argument buffer layout: encodedLength=640 maxAid=48`.
- **Cause:** a loop reads one handle per light from the game's uniform buffer at a computed offset; Ryujinx's texture array from a buffer indexes it by the handle's byte offset ÷ 8 (an NVN handle is 64-bit, the shader computes `(word >> 2) << 1 | (word & 3) >> 1`). The build reads the name as one bindless texture (`tex_ubo`, offset 0), so the argument buffer was 32 entries longer than the reflection said.
- **Fix:** the host's reflection reader takes the array length from the MSL type, and the replay fills element k from the handle at byte 8k of that uniform slot (`NVN_RB_TEXU` / `SAMPU` with `n`); elements whose bytes are not a registered texture or sampler ID read the dummy, silently. No chain change: the build already wrote the array type.
- **Found on:** [Paper Mario: The Thousand-Year Door](../../dm/games/0100ECD018EBE000/README.md): the reflection shaders (`reflection.bfsha`: `array<texture2d<float>, 32> fp_t2d_cb14_0` and its samplers), c14 (uniform slot 11). The first run showed elements 0–3, 6 and 8–12 holding textures, among them the shader's fixed handles at +8, +16 and +24 and the lights' from +64.

### An array of handles from the driver table c2

- **Symptom:** `shaders with a binding whose slot was not decoded 588`; shadow cascades read through `array<texture2d_array<float>, 32> fp_t2das_cb2_0`.
- **Cause:** Ryujinx's texture array from a buffer, as Paper Mario's uniform-buffer arrays ([An array of handles in a uniform buffer](#an-array-of-handles-in-a-uniform-buffer)), but over the driver table c2, indexed by (0x80 + a cascade offset) ÷ 8.
- **Fix:** `dmsh_build.py` `c2_array` names it `tex_c2a` / `samp_c2a` with the c2 byte offset of element 0; the replay fills element k from the slot whose c2 word is at that offset + 8k (combined, else separate), and the dummy elsewhere (`NVN_RB_TEXA` / `NVN_RB_SAMPA`, `c2_handle`).
- **Found on:** [Luigi's Mansion 3](../../dm/games/0100DCA0064A6000/README.md) (586 pixel shaders, the shadow cascades).

### Arrays of c2 handles cover the whole table, with a separate sampler word

- **Symptom:** `[RED] shader … reflection conversion failed: B binding fp_t2d_cb196610_10570000 has no resolved slot number`, then, with that resolved, `… reads driver-bank locations at run time that the replay layer cannot fill: c2 dynamic index`; the shader does not load and its draws are skipped.
- **Cause:** two gaps in [An array of handles from the driver table c2](#an-array-of-handles-from-the-driver-table-c2). The translator sized every texture array from a buffer at 32 elements (`Program.cs` `QueryTextureArrayLengthFromBuffer`) and clamps the index to the last one, but over c2 the index is any handle word's byte offset ÷ 8: EA SPORTS FC 27's pixel shader `dc446d7f52c14ea6` loops over separate texture slots i = 1, 2, …, reads each handle word at c2 byte 0x168 + 8i and samples element 45 + i, so every element it picked was clamped to 31 (a combined slot). And the handle is type 1: `BindlessToArray` ORs one fixed separate sampler word (c2 byte 0x570, separate sampler slot 1) into every element's handle, which `c2_array` did not decode (it took only `_cb2_` with type 0), so the texture array had no slot and its sampler array took that one slot.
- **Fix:** a chain change. `Program.cs` gives an array over c2 one element per handle word up to the end of the separate sampler table (`C2HandleWords`, 301; other buffers keep 32), `c2_array` decodes type 1 with the sampler word in c2 too (`soff` in the reflection), and the replay ORs that word into each element's handle for the samplers (`NVN_RB_SAMPA`; `c2_handle` now also answers the separate sampler table). `nvn_rbind.n` is 16 bits and `NVN_REFL_MAXARG` 768 (such a pair takes 602 argument entries). The same loop reads each handle word itself at a computed c2 index (`fp_c2._m0[int(_237)]`, to test the light's handle), which `unfillable` refused for any bank; the replay fills c2's handle tables whole (`DRV2` in `encode_args`), so a computed index into c2 is accepted, and one into c0, filled only in its tables, still refused. Luigi's Mansion 3's cascade arrays become 301 long as well; the host fills the elements past the handles with the dummy.
- **Found on:** [EA SPORTS FC 27](../../dm/games/0100C49025D3E000/README.md) (its lights in a match).
- **Check:** `test_dmsh_reflect.py` (`DriverTableArrays`, `DriverReads`) and `test_dmsh_abi.py` (the translator's table length agrees with the header).

### A texture slot that shares its c2 word with an image slot

- **Symptom:** characters and props glowing white, about 10⁴ too bright, in DRAGON QUEST VII Reimagined's opening cutscene only, while walls and gameplay look right.
- **Cause:** a size query (`GetDimensions` on a `RWTexture3D`) is a TXQ on the image's handle, c2 `0x120` (image slot 0): the build numbers it combined texture slot 32, the slot that shares that word, and with no type in the facts table the translator read it as 2D. The host bound the empty texture slot 32, and the 1×1 dummy's size passed only thread (0, 0, 0).
- **Fix:** two general fixes: a texture binding of slot 32–39 reads the image bound in slot 0–7 when there is one, as c2 holds it (`nvn_mtl.m` `tex_slot_h`; Hades II's texture slot 34 without an image still reads the texture), and a query on an image handle the facts do not type is translated again with the type of that image from the first translation (`dmsh_xlate` `ImageTypesOf`).
- **Found on:** [DRAGON QUEST VII Reimagined](../../dm/games/0100A9D01C446000/README.md): the per-pixel volumetric lightmap read brick (0, 0, 0) of the brick atlas for almost every cell. Unreal merges each streamed level's lightmap into one atlas on the GPU: brick copies (`0a7ff4171bd1a82a`), then a patch of the shared 16³ `RGBA8_UINT` indirection texture (`7b958a7b21e72801`) whose bounds come from `GetDimensions` on the very `RWTexture3D` it writes. Only cell (0, 0, 0) was written, every other cell kept size 0, and the shader fell back to the atlas origin: the first daytime level's bricks, baked about 2¹⁵ brighter (the same bricks in gameplay hold 0.5–1.4). [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md) has four such size queries (`cp_t_tcb_48` / `_4A`).
- **Check:** how it was found: the raw scene colour (`texraw`) put the characters at 1,000–2,000 and the walls at 0.05; a debug build of one character shader that wrote its terms in alternating 8-pixel columns isolated the volumetric lightmap ambient term; dumps of the atlas (`texlist` / `texraw`), of the uploaded level bricks (the 3D upload dump) and of the indirection texture, with the cell computed from the primitive's origin and View `VolumetricLightmapWorldToUVScale / Add` (View 220 / 221), gave the rest.

### More than 64 bindings in one shader

- **Symptom:** `[RED] shader … reflection conversion failed: 66 bindings, over the limit 64`, then `failed to load`, and the draws with it are dropped (a black map).
- **Cause:** the reflection is converted into fixed arrays of `NVN_REFL_MAXB` bindings (`nvn_refl.h`), which was 64. NVN gives a stage 32 textures and 32 samplers beside its uniform and storage buffers and images, so a shader can have more.
- **Fix:** `NVN_REFL_MAXB` is 128; the argument buffer entries keep their own limit (`NVN_REFL_MAXARG`, 768, see [Arrays of c2 handles cover the whole table](#arrays-of-c2-handles-cover-the-whole-table-with-a-separate-sampler-word)). It was raised from 96 when every sampled sampler got a twin binding ([Border colours Metal has no sampler state for](rendering.md#border-colours-metal-has-no-sampler-state-for)): that shader then has 97 bindings.
- **Found on:** [OCTOPATH TRAVELER II](../../dm/games/0100A3501946E000/README.md): the 3D world map of the traveler selection (`UiMM_WorldMap3D`, pixel shader `f7c8373b7323687d`: 30 textures, 30 samplers, 3 uniform buffers and the c1 / support / driver segments).

### Argument buffer members with attributes after their id

- **Symptom:** the host stops at a shader's first load: `FATAL shader … has an unexpected argument buffer layout: encodedLength=88 maxAid=8` (crash in `c1_store` / `pull_load`), with a storage buffer or image missing from its reflection json.
- **Cause:** a pixel shader that writes a storage buffer and images inside a fragment interlock (the translator's FSI, `PixelInterlockOrderedEXT`) gets those members as `[[id(n), raster_order_group(0)]]` from SPIRV-Cross. `dmsh_build.py` `reflect` matched only `[[id(n)]]`, so it left them out, and the host's argument buffer came out longer than the reflection.
- **Fix:** `reflect` reads attributes after the id, and a member it still cannot read fails that shader's build (`reflection failed`) instead of loading with a short reflection. A chain change.
- **Found on:** [Mario Party Superstars](../../dm/games/01006FE013472000/README.md): `light_primitive_point_cluster.bnsh` (`56071a5cf62dccc3`, the engine's clustered point lights), first drawn on the way into a party.

### Image atomics

- **Symptom:** the replay put a texture resource ID where the shader wanted an address; with the old 64 KB dummy buffer, pointers in the game's objects and in the host's program objects got their upper 4 bytes replaced and the game or the replay crashed, or `kIOGPUCommandBufferCallbackErrorPageFault`, after which macOS ignores the process's command buffers (GitHub issue #101). Later, `[RED] image atomics on a block-linear texture … not followed, the atomics go to a scratch buffer`, and compute shaders with atomics on a 3D image did not build (`array subscript is not an integer`).
- **Cause:** below MSL 3.1, SPIRV-Cross emulates an image atomic with a raw device pointer beside the texture (`texture_buffer<uint> cp_i_cb2_4C_r32ui` and `device atomic_uint* cp_i_cb2_4C_r32ui_atomic`), indexed by the bound image's width (`spvImage2DAtomicCoord`), so the host had to bind memory in that packed layout. A buffer image's pool memory has it. A block-linear image's Metal texture is a deswizzled copy whose pool memory has another layout, so its atomics went to a scratch buffer that nothing read, and a pointer shorter than the image (or an invalid coordinate) wrote over other GPU memory. For a 3D image SPIRV-Cross indexes the pointer with the `int3` coordinate itself.
- **Fix:** image atomics are Metal's native texture atomics (MSL 3.1: `texture3d<uint, access::read_write>::atomic_fetch_max`). `dmsh_atomic.py` has SPIRV-Cross write shaders with image atomics (an `OpImageTexelPointer`) at MSL 3.1, and every other shader at MSL 2.3 as before, so their MSL is unchanged and their libraries come from the compile cache; a shader whose MSL still has the pointer form fails the build. The atomics act on the texture every other access of the image uses (a buffer image's view of its pool, a linear image's pool memory, a block-linear image's Metal texture), for buffer, 2D, 2D-array and 3D images. Metal runs texture atomics only on an R32Uint / R32Sint texture made with `MTLTextureUsageShaderAtomic`: NVN images of those formats take that usage when they are made (`nvn_mtl.m` `img_usage`; a view keeps working through its base), the build marks the slots a shader uses atomically (`"atomic"` in the reflection, `nvn_rbind.atom`), and the host binds an atomic-capable write sink, logged red, in place of any other texture there (`dummy_img_atomic`). The translator's coordinate guard stays ([Bounds of image atomics](#bounds-of-image-atomics)).
- **Found on:** [Octopath Traveler 0](../../dm/games/01005270232F2000/README.md): 2,084 shaders with atomics on buffer images (Unreal's GPU particles; three ran in the first 100 seconds) and 16 on a 3D image. [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md): an atomic max into a 480x270 r32ui block-linear 2D array (fragment shader `2bd5037003980716`, the scene after the story opening's first cutscene), which ran about 1.4 MB past the 64 KB dummy buffer into zero-copy pools over the game's heap and into unmapped addresses (GitHub issue #101), then went to a scratch buffer. [FINAL FANTASY RESONANCE](../../dm/games/0100606026F12000/README.md) (8 shaders) and [Ghost Trick](../../dm/games/010029B018432000/README.md) (5, one of them a compare-and-swap): atomics on 3D images.
- **Check:** `dm/tools/tests/test_image_atomic_bounds.py` with the rebuilt translator, `DMSH_TEST_METAL=1 DMSH_TEST_GPU=1`: synthetic Maxwell atomics on buffer, 2D, 2D-array and 3D images, compare-and-swaps among them, run on native atomic textures of a Mac GPU under the Metal validation layer (`t_atomic_return.m`). `t_refl.m` checks the atomic flag of a reflected image.

### Buffer image reflection keeps its texture type

- **Symptom:** `image atomics: unsupported or overflowing IMG extent` at the first dispatch, despite a valid four-byte counter buffer.
- **Cause:** `nvn_refl.m` classified `texture_buffer<uint>` as the default 2D texture. The atomic admission check of the pointer emulation then rejected the actual Metal buffer texture against that 2D descriptor.
- **Fix:** the shared reflection type decoder recognizes `texture_buffer` as `MTLTextureTypeTextureBuffer`, so a buffer image keeps its buffer-image shape. `t_refl.m` checks the reflection converter with a buffer image.
- **Found on:** [SHIN MEGAMI TENSEI V: Vengeance](../../dm/games/010069C01AB82000/README.md), compute shader `2d2093c1f1d3a7e0`, image slot 2 (one R32Uint texel).

### Bounds of image atomics

- **Symptom:** Metal shader validation reports invalid image-atomic stores even with an image-sized scratch buffer; without validation, GPU page faults can freeze rendering.
- **Cause:** the translator lost the surface instruction's `IGN` bounds mode before lowering an image atomic, so a coordinate outside the image reached memory. Neither a pointer nor a native texture atomic outside its image is defined on Metal.
- **Fix:** the translator (`TranslationFlags.MetalImageAtomics`, Metal only; other backends keep their behaviour) keeps the bounds mode and whether the result is used, queries the bound image's size and tests every original coordinate, the layer included, before the atomic. Negative coordinates fail an unsigned comparison; coordinates are not clamped or flattened. Invalid coordinates leave memory unchanged and return zero ([ARB_shader_image_load_store, issue 30](https://registry.khronos.org/OpenGL/extensions/ARB/ARB_shader_image_load_store.txt)): a returning `IGN` atomic joins the old texel from the valid branch and zero from the other with an SPIR-V phi. A compare-and-swap (`SUATOM.CAS`: compare with Rb, swap in Rb+1) is Metal's texture compare-exchange behind the same guard; SPIRV-Cross loops `atomic_compare_exchange_weak` until it succeeds or the texel differs. Predicate-returning, `TRAP` and unknown modes fail translation, as do 64-bit, floating-point and half2 sources (checked before the descriptor format can erase their size), descriptors that are not R32 integers, bindless handles that do not resolve statically, and indexed or multisample images. Source 1D and 1D-array surfaces keep the translator's one-row normalization to 2D.
- **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md#rendering) ([PR #178](https://github.com/seedds/nxnative/pull/178)); [SHIN MEGAMI TENSEI V: Vengeance](../../dm/games/010069C01AB82000/README.md), whose two buffer-image counter shaders dispatched at boot consume their returned values.
- **Check:** `dm/tools/tests/test_image_atomic_bounds.py`. Portable tests check the MSL version choice and the rejection of the pointer form. With `DMSH_TEST_XLATE` set to a rebuilt translator (and `spirv-val` on `PATH`; spirv-cross is the in-repo build) the synthetic shaders are translated, the guard's dataflow is checked in the SPIR-V (the size query and comparison dominate the atomic, the phi carries zero on the invalid edge), and the MSL is produced; `DMSH_TEST_METAL=1` compiles and links it (`DMSH_TEST_METAL_SDKS=macosx,iphoneos` for both SDKs) and `DMSH_TEST_GPU=1` runs the returning fixtures (above). Changing this path changes the chain version: rebuild the translator and update the lock before building shaders.

## Shaders compiled at run time

### A GLSL language preamble is not a complete shader

- **Symptom:** the NSP script stops with `captured only 0 run-time shaders` although the game reaches its frame loop with precompiled shaders.
- **Cause:** the detection counted every `#version` occurrence. Snowdrop carries a GLSL header containing extensions and type macros, without an entry point; importing `glslcCompile` alongside it does not establish a boot shader compilation.
- **Fix:** `profile_guess.py` counts embedded source strings containing both `#version` and `void main(...)` before their terminating NUL. The survey, build and split share that detector. Three Houses still has its nine complete boot sources; NRO imports retain their separate capture rule.
- **Found on:** [Mario + Rabbids Kingdom Battle](../../dm/games/010067300059A000/README.md), whose only `#version` string is the 646-byte language preamble also carried by Sparks of Hope.

### Built-in shaders compiled with the game's own GLSLC

- **Symptom:** shaders missing from the map that no game file holds.
- **Cause:** a game whose shaders come from its own GLSLC at boot has no file holding their output. KTGL compiles 9 built-in GLSL shaders at boot with the game's own GLSLC (`glslcCompile` from `main+0x9c89f8`; the sources are strings in `main`).
- **Fix:** they are captured once on the Mac (`DM_GLSLC_CAPTURE`): `DM_GLSLC_CAPTURE=<dir>` saves every compile's output (`hle_glslc.c`; the compile object's first word points at {status, output}), and `shader_split.py` reads `<version dir>/glslc_capture` (the NSP script makes it for any game whose `main` imports `glslcCompile` and holds GLSL sources or, with none there, whose RomFS holds `.vert` / `.frag` files). The outputs are identical on every run. `shader_split.py` asks for captured GLSLC outputs only when a module of the profile's defines `glslcCompile`, as `profile_guess.py` does.
- **Found on:**
  - [Fire Emblem: Three Houses](../../dm/games/010055D009F78000/README.md) (SDK 6.5.0): the 2D primitive, text and loading-feather shaders.
  - [Catherine: Full Body](../../dm/games/0100BF00112C0000/README.md) (SDK 8.2.1): its two built-in programs are captured on the Mac the same way; they compile only once the host reports the GLSLC's GPU code version ([next entry](#glslc-gpu-code-version-against-the-drivers-range)).
  - [The Legend of Zelda: Link's Awakening](../../dm/games/01006BB00C6F0000/README.md): `subsdk0` is NVN's runtime GLSL compiler and `main` imports `glslcCompile*`, but up to the house, the beach and the menus every shader the game loaded was in the offline tables.
  - [Captain Toad: Treasure Tracker](../../dm/games/01009BF0072D4000/README.md) (the rule that a module must define `glslcCompile`): `main` imports `glslcCompile` and holds agl's GLSL sources, but no module defines it, so nothing is compiled at run time.

### GLSLC GPU code version against the driver's range

- **Symptom:** a null read at `main+0x3316dc` in `nnMain`, right after the window setup (the program object's interface pointers at +0xd0 / +0xd8 both null).
- **Cause:** the game compiles two built-in programs at boot with its own GLSLC (`subsdk0`, the same build as Link's Awakening's), and first checks that GLSLC's GPU code version against `nvnDeviceGetInteger` items 0x2E–0x31 (the driver's supported range, `main+0x330dcc`). The host reported Super Mario Odyssey's SDK 7 driver range, 1.5–1.14; this GLSLC writes 1.15, so nothing was compiled and the first uniform lookup read the empty program.
- **Fix:** every SDK's driver runs the code of the GLSLC in the same SDK, so the host now reports at least the version the game's own `glslcGetVersion` returns (`hle_glslc.c` `dm_glslc_gpu_code`, `hle_nvn.c`). The host also advertises the translator's supported range for precompiled shaders ([Newest GPU code version reported by the driver](nvn-objects.md#newest-gpu-code-version-reported-by-the-driver)).
- **Found on:** [Catherine: Full Body](../../dm/games/0100BF00112C0000/README.md) (SDK 8.2.1, Atlus's GFD engine, CRI middleware; one boot).

### GLSL ES 1.00 sources without a `#version` line

- **Symptom:** `gl_shaders.py` failed on the captured sources: `ES shaders for SPIR-V require version 310 or higher`.
- **Cause:** GameMaker's runner writes each shader variant as **GLSL ES 1.00 with the version statement left out** (`attribute`, `varying`, `texture2D`, `gl_FragColor`, `precision mediump float;`), and uses `sample` as a variable name, a reserved word since ES 3.00.
- **Fix:** `gl_shaders.py`'s `es100()` pass, run when the source has no `#version` line or says `#version 100` (that line is dropped): drop `precision <p> <t>;` lines (the pattern has to accept `\r`, the captured files mix CRLF and LF), strip `lowp`/`mediump`/`highp` (GameMaker has `#define LOWPREC lowp`, so a macro would otherwise expand to nothing useful), rename `attribute`→`in`, `varying`→`out` (vertex) / `in` (fragment), declare `out vec4 gm_FragColor;` and rename `gl_FragColor`, rename `sample`→`gm_sample`, then prepend `#version 310 es` with `precision highp float; precision highp int;`. Texture functions stay with the existing `fixups()` so the rename happens after `glslangValidator -E` and a macro cannot hide one.
- **Found on:** [Well Dweller](../../dm/games/0100217023F6C000/README.md) (93 run-time shaders; see its [`gl_capture.tsv`](../../dm/games/0100217023F6C000/gl_capture.tsv)).

### Struct uniforms become member rows

- **Symptom:** `uniform LightEnv { … };` failed with `uniform u_LightEnv: type LightEnv (structs are not implemented)`, so every program that used one was missing from the library.
- **Cause:** GameMaker declares several uniform blocks as struct uniforms, and the layer's `glGetUniformLocation` compares the name the game asks for (minus a `[n]` suffix) against the table's rows, whose name field is 48 bytes — the game asks for `u_LightEnv.m_dirLights_dir` and similar dotted member names.
- **Fix:** `gl_shaders.py` gained `uni_members()`, which walks a struct type's members and emits one row per member with the dotted name (`u:u_LightEnv.m_dirLights_dir[0]`), adding each member's offset to its parent's (the reflection gives a nested member's offset relative to its parent); arrays of structs return an error.
- **Found on:** [Well Dweller](../../dm/games/0100217023F6C000/README.md), whose lighting programs were the missing ones.

## Pipelines and caches

### A compile that does not return stalls every flush and draw

- **Symptom:** the pipeline prebuild stops at the same line on every start (`pipeline prebuild: 7000 of 20740 pairs, 1987 built, 5013 already archived, 0 left to the draws, 0 failed`), behind *Preparing pipelines* or, with `DM_PSO_PREPARE=0`, as a black screen while the sound plays on; the NSP script's `[7/7] Pipeline prebuild` does not end (`dmhost` at 0% CPU, `MTLCompilerService` at 100%), and `MTLCompilerService` crash reports (`EXC_CRASH`, `llvm::report_bad_alloc_error`) come every 6 to 10 minutes. The bases built since the last file are never written, so the next start builds them again and stops at the same pair.
- **Cause:** one base Apple's GPU compiler cannot build, and a host lock that let its compile hold everything else. The compiler grows `MTLCompilerService` past several GB until it aborts in `llvm::report_bad_alloc_error`, and Metal hands the request to the relaunched service; only after four aborts (33.7 minutes on an M6 with 16 GB, macOS 27.0) does `newRenderPipelineStateWithDescriptor` return, with `Compilation failed due to an interrupted connection: XPC_ERROR_CONNECTION_INTERRUPTED. This error occurred after multiple retries.` Metal's serializer cannot be written while it captures: the write (`-[_MTLBinaryArchive airntSerializeToURL:options:error:]`) runs on the archive's `_serializeQueue` and then empties the archive (its entries, AIR objects, library cache and file mapping) on its `_queue`, while each function a compile captures adds itself from the compile's completion (`__122-[MTLCompiler compileFunctionRequestInternal:…]_block_invoke`: a new `MTLBinaryEntry` from the request's AIR script and the function's bitcode, mapped by `newArchiverIdWithBinaryKey` before it takes `_queue`, then added under it, before the compile reports back). A compile under way across a write came back with both nil and crashed in `dispatch_data_create_map` (seen in lldb at the crash, the write's thread waiting for `_serializeQueue`). So every compile the archive captured held a read lock (`g_arch_rw`) and a flush its write lock. The flush at the next 1,000 bases waited behind the compile, and every other compile, draw and the present path waited behind the flush. The prepare step also waited for the last compile to return.
- **Fix:** flush periods (`ArchEpoch` in `nvn_mtl.m`). Each period has its own compiler and serializer; a compile the archive captures joins the current period (`arch_capture`) and leaves it when it returns. A flush (`arch_flush`) gives the compiles after it a new period and returns at once; the old period's serializer is written (`arch_write`, on the archive queue) once every compile in it has returned. No compile waits for a flush, and a compile that does not return holds back only its own period's file. The prebuild's flushes every 1,000 bases no longer wait; its last one (`arch_flush_wait`) waits for every period before `prebuilt.tsv` is written. Specializations go through a compiler without a serializer (`g_comp4`). The prepare step waits until every pair is handed to a worker; compiles still under way then (at most one per worker) finish in the background, and the log says `pipeline prepare: every shader pair handed out after … s; N compiles still under way finish in the background`. The prebuild tool still waits for every compile, since it writes them all. The base itself was made buildable by [Front end at -O1](#front-end-at--o1-apples-control-flow-lowering) (11.4 s at run time on the same M6). A base that fails no longer keeps `prebuilt.tsv` from being written (every start compiled it again in the background, half an hour at `-O2`): the pair goes into `pso_archive/prebuild_failed.tsv` with the system build and the stamp it failed under (`pre_fail_load`), and a later start of that system build with that stamp leaves it to the draws without compiling it; another system build (a new GPU compiler) or stamp tries it again. A failure Metal reports has already been retried by Metal itself. The log says `pipeline prebuild: N pairs whose base failed are left to the draws (…prebuild_failed.tsv; M of them failed on an earlier start of this system build, not built again)`. **Open:** a compile still under way when the game exits leaves no result, so a base that takes longer than a session to fail is compiled again in the background on the next start.
- **Found on:** [Super Mario Odyssey](../../dm/games/0100000000010000/README.md#files-and-shaders), update 1.4.1, GitHub issue #294: the base of pair 5641, `3c3435a62e05975b` + `007d7e411ef14ebd` (`alRenderMaterialUber.bfsha` variation 0, 444,013 lines of MSL), found by a reporter's bisect of the pair list with `DM_PSO_PREBUILD=tool`, with that chain's `-O2` front end. On an M6 with 16 GB (macOS 27.0), the prebuild tool on the whole list with the `-O2` library and 8 workers: the old host stopped at 7,000 of 20,740 at 65 s and wrote nothing more in 5 minutes; with flush periods 20,000 pairs were done at 262 s, and 17 files held 18,739 of the other 20,739 bases, only the 2,000 in that compile's period waiting for it (three compiler aborts in 40 minutes). With the `-O1` library (all 20,740 built): 253 s cold and 227 s with the old host, 213 and 225 s with flush periods; the files written gave all 20,740 back, none compiled.
- **Check:** `t_pso_stall`: a compiler that holds one prebuild compile; while it is held, the prebuild hands out every pair and builds the others, a draw's compile returns and a later period's flush writes its file, and a base of a written period still specializes; once released, its period is written and the prebuild stamps. A pair whose compile fails: stamped all the same and listed; not compiled again on a later start of the same system build, also without the stamp; tried again with a list from another system build, and the list removed once it builds. On Odyssey's pair 5641 alone with the `-O2` library (M6, macOS 27.0): Metal's error after 36.6 minutes, the pair listed and the stamp written; the next start had nothing to do (1 s), and one without `prebuilt.tsv` left the pair to the draws without compiling it. The Metal behaviour, by a probe outside the tree on an M6 (macOS 27.0), eight threads compiling 1,500 pipelines: one serializer written every 150 ms while they compile crashed 3 runs of 3 in `newArchiverIdWithBinaryKey`; a write that first waits for the compiles under way but lets new ones start, 6 of 6; one that lets no compile start until it is done, 1 run of 111, during the write itself, with [issue #342](#a-compaction-flush-that-crashes-then-crashes-every-start)'s signature (`dispatch_retain(NULL)` in `materializeAirNTFromOnlineForSlice`); a compiler and serializer per period, none of 118, every pipeline found in the files, also with 121 periods in one run (no memory growth), and making a period took about 30 µs.

### Shader-pair tables in a Mac app

- **Symptom:** a self-contained Mac app (`--platform mac-app`) logs `pipeline prebuild: no .../Resources/dm/shader_pairs.tsv, nothing to build` although every shader library is there; it draws no hybrid stand-ins, and each pipeline seen for the first time compiles on the submit thread (frames of 100–550 ms).
- **Cause:** `dm/mac/pack_app.py` copied the shader map, the prewarm manifest, `c1_full.tsv` and `cb_extent.tsv` beside `DM_NVN_SHMAP`, but not `shader_pairs.tsv` and `pso_pairs.tsv`, which `build_ipa.sh` copies into the IPA and the host's pipeline prebuild reads there (`pso_prebuild`). Without them the app makes no pulling bases. For a profile without `pso_pairs.tsv` a Mac run (`mac_run`) passes the shader split's copy as `DM_PSO_PAIRS`; the app had neither. A full shader-library build does not build the pipeline bases by itself.
- **Fix:** the packer copies both tables when the profile has them, and the split's `pso_pairs.tsv` (`--pso-pairs`, which the NSP script passes) when the profile has none. They stay optional. An app packed before the fix needs repacking. An NVN game's app also carries its pipeline archive made offline for this Mac's GPU, so the first start builds no base: [A Mac app](../development.md#a-mac-app).
- **Found on:** [The Legend of Zelda: Echoes of Wisdom](../../dm/games/01008CF01BAAC000/README.md#rendering-and-shaders), update 1.1.0: a 308.4 ms frame, 283.5 ms of it pipeline acquisition and 5.9 ms GPU work; with the tables the offline prebuild built all 7,495 bases. [The Legend of Zelda: Breath of the Wild](../../dm/games/01007EF00011E000/README.md#pipelines-on-a-cold-start): frames of 100–546 ms before, none of its bases compiled on the first start after.
- **Check:** `test_nsp_mac_app`: `App.test_pipeline_pair_tables`, `App.test_split_pso_pairs` and `App.test_layout` fail without the tables in the packer; `Script.test_pack_command` checks the split's copy is passed only when the profile has no `pso_pairs.tsv`.

### Offline pipeline archive missing for the device GPU

- **Symptom:** `pipeline archive bundled with the app: none for this GPU (...)`: the device builds the bases itself although the app carries an offline archive.
- **Cause:** each archive targets one GPU architecture. The iOS packers defaulted to `g18p` for profiles with `pso_mtl4 on`; another GPU family needed an explicit build override.
- **Fix:** both packers also read the optional profile key `pso_gpu` behind the existing `pso_mtl4 on` gate. It names the families to build for; explicit build overrides still win. Without the new key, opted-in profiles retain `g18p`; profiles with Metal 4 off retain no archive. See [archive target selection](../development.md#the-pipeline-archive-made-offline).
- **Found on:** [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md#offline-pipeline-archive-probe). Its note records the contributor's Mac build probe and the outstanding device checks.
- **Check:** `dm/tools/tests/test_game_profile.py` validates the key; `dm/tools/tests/test_pso_offline.py` exercises both packers' selection and override precedence with synthetic profiles.

### Offline pipeline archive reuse and fallback

- **Symptom:** a truncated `.mtl4a` reaches `MTL4ArchiveImpl::loadFromURL`; the bundle reports coverage with `0 of 0 probed bases found`; rebuilt shader libraries reuse an archive of their old binaries; or packing fails with `pso_offline: no base to build` when every eligible pair has function constants, or when the profile's `shader_pairs.tsv` holds only its header (`--platform mac-app`, [GitHub issue #338](https://github.com/seedds/nxnative/issues/338)).
- **Cause:** the archive guard checked only four magic bytes; coverage accepted `found == probed` even when every sampled pair was ineligible; the offline cache key included library paths and profile stamps but omitted library bytes; and a zero-base plan aborted before writing its device remainder list, also when there was no pair at all, so `mac_pso()` never reached its no-archive branch.
- **Fix:** `arch_file_ok` checks the GPU Mach-O load commands and payload ranges, universal slice bounds and the serializer's bundled Metal library ranges before opening the archive. It does not validate the compiler's private payload. `pso_bundle_covers` requires at least one actual lookup and no misses. `pso_offline.py` hashes every consumed library's bytes into the cache key; when all eligible pairs need device specialization, it removes old archives and writes `rest_pairs.tsv` without invoking `metal-tt`. A pair list with no pair takes the same path (old archives removed, stamp written, exit 0), and the app is packed without an archive; pairs that the host would leave out entirely (wrong or missing libraries) still fail.
- **Found on:** [The Legend of Zelda: Breath of the Wild](../../dm/games/01007EF00011E000/README.md#found-on-this-game), during review of its offline archive path with synthetic libraries and pair lists; no new gameplay or iPhone run. The empty pair list on [ANIMAL WELL](../../dm/games/010020D01AD24000/README.md#found-on-this-game).
- **Check:** `dm/tools/tests/test_pso_offline.py` changes vertex and fragment library bytes without changing entries or stamps, and exercises the all-constant plan, the header-only pair list and pairs all left out. `t_pso_bundle` rejects bare headers, truncated archives, malformed load commands and out-of-range universal slices, accepts real offline and serialized archives, and rejects zero-probe coverage.

### Shader extraction reruns after translation-only updates

- **Symptom:** an unchanged NSP spends minutes in `shader split` after a shader-chain update, on both Mac and iOS builds, although the extracted package is identical.
- **Cause:** the split key included all graphics Python tools, every profile file and the whole decoder host executable. Translation tools, profile notes and the host's compiled `DMSH_CHAIN` label therefore invalidated extraction.
- **Fix:** the NSP script follows the extractor's transitive repository Python imports and fingerprints its profile settings, manual facts, checked/generated tables, modules, RomFS/DLC identities and capture directories. A managed decoder host records its build inputs with only the chain label normalized, bound to its executable's bytes; unknown or replaced binaries use their binary identity. All other host build inputs and rewritten module contents still invalidate conservatively. Package and supplemental-pair digests reject damaged cached outputs. The common split serves Mac and iOS, while translation keeps its own target-specific keys. Existing split caches require one warm-up.
- **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md#build-performance).
- **Check:** `test_nsp_incremental.py` covers translation-only hits, transitive readers, extraction settings, auxiliary modules, explicit captures, decoder changes, corrupt outputs and the existing registered/derived, DLC and damaged-input paths. `test_host_msl.py` verifies executable-bound provenance, chain-only reuse and invalidation for decoder/profile/compiler/header changes. `test_nsp_mac.py` checks separate Mac/iOS translation outputs.

### Offline pipeline archive rebuilt after a chain-only update

- **Symptom:** rebuilding an unchanged NSP after a shader-chain update runs `metal-tt` for minutes even when all referenced Metal libraries and pipeline settings remain identical.
- **Cause:** the offline archive key included the chain label, profile stamp, whole builder source and absolute library paths inside `b-<chain>`. These can change without changing the compiler's inputs.
- **Fix:** `pso_offline.py` keys each archive by the complete generated pipeline configuration with library locations normalized to their labels, every referenced library's name and bytes, GPU architecture, platform, minimum OS and Metal translator version. It still validates reflection against the current chain before planning. After a cache hit it writes the current `prebuilt.tsv`, remainder list and archive metadata, so host validation remains current. macOS and iOS use the same builder with separate target keys. Legacy cache keys rebuild once; a changed shader, entry point, pipeline descriptor or compiler target still rebuilds the archive. The cache format has an explicit version for future changes not represented by these inputs.
- **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md#build-performance), while measuring repeat Mac app builds. Timings and platform-validation limits are in the game note.
- **Check:** `dm/tools/tests/test_pso_offline.py` exercises chain relocation and metadata refresh on both platforms, changed shader bytes and pipeline settings, target and compiler changes, missing archives, stale reflection and legacy keys.

### Offline pipeline archive and profile depth disagree

- **Symptom:** the host logs `pipeline archive bundled with the app: made for ... this run needs ...`, naming different depths, and pipeline bases compile during play despite a successful archive build. There need not be a `[RED]` line.
- **Cause:** offline builds use the profile's `depth_mode` (`gl` when omitted), while runtime vertex entries follow `nvnDeviceSetDepthMode`. An archive built with `main0` cannot supply a near-zero base using `main0_z`. The runtime stamp comparison correctly rejects the mismatch; the missing input is in the profile.
- **Fix:** declare the observed mode in the profile and rebuild the archive. Both vertex entries already exist, so this does not require a shader-chain change. `scripts/regress.py` records the runtime depth and compares it with the tested revision's profile on every NVN smoke run, including quick runs without an archive. It also records bundled archive stamp rejection for any chain, depth or pair-list mismatch. New mismatches fail the check; unchanged baseline mismatches are reported as known and prevent `--record` from marking the game validated. An absent depth log is unknown, not evidence for GL. This detects mismatches on exercised routes; it does not infer unobserved games' modes or edit profiles automatically. For a new game the NSP script does: the heap and address region checks of a provisional or derived profile set the key from the game's own `nvnDeviceSetDepthMode` call in their boot, before the shader steps ([Super Mario RPG](../../dm/games/0100BC0018138000/README.md) before this: a 24 minute prebuild at the default mode, then 0 of 34 archive lookups found). A game with no check boot, or one that stops before the call, keeps the default.
- **Found on:** [The Legend of Zelda: Echoes of Wisdom](../../dm/games/01008CF01BAAC000/README.md#rendering-and-shaders), update 1.1.0, [GitHub issue #758](https://github.com/seedds/nxnative/issues/758): its note already recorded near-zero depth, but its profile omitted the key. The game note records archive validation and its limits; castle-entry performance remains open.
- **Also confirmed:** missing depth keys on [SteamWorld Heist II](../../dm/games/010089C01E9CA000/README.md#shaders), [NieR:Automata](../../dm/games/010056B015FE8000/README.md#files-and-shaders), [Portal](../../dm/games/01007BB017812000/README.md#rendering), [DRAGON QUEST III HD-2D Remake](../../dm/games/01003E601E324000/README.md#files-and-shaders), [Raging Loop](../../dm/games/010020400E1C2000/README.md#files-and-shaders), [Subnautica](../../dm/games/0100429011144000/README.md#files-and-shaders), [Subnautica: Below Zero](../../dm/games/010014C011146000/README.md#files-and-shaders), [Batman: Arkham Knight](../../dm/games/0100ACD0163D0000/README.md#what-it-needs), [Blasphemous 2](../../dm/games/010040801A4C8000/README.md#what-it-needs), [EA SPORTS FC 27](../../dm/games/0100C49025D3E000/README.md#what-it-needs) and [Paper Mario: The Origami King](../../dm/games/0100A3900C3E2000/README.md#what-it-needs). Recorded routes, short isolated boots and the NSP script's check boots select near-zero depth on their registered module builds; their notes record the evidence. The profile corrections have not been measured as rebuilt-archive performance fixes.
- **Check:** `dm/tools/tests/test_pso_offline.py` runs the archive builder with the affected games' real profiles and synthetic libraries, checking the stamp and the vertex entry. `dm/tools/tests/test_regress.py` replays runtime depth and rejection logs through a smoke run, checking failure classification and reporting.

### Offline pipeline archive made for another major version

- **Symptom:** `pipeline archive bundled with the app: this system cannot open …/applegpu_<gpu>.mtl4a (deployment target for architecture … found in archive … is not compatible with current OS …): the bases are built on the device`. On macOS 26.6 with Xcode 27.1, `t_pso_bundle` failed 10 checks, from `bundle: opened` on, and a Mac app packed there compiled every base on its first start, as without the archive.
- **Cause:** a Metal 4 pipeline archive opens only on the major system version it was made for. Measured on macOS 27.2 (M4 Pro, `newArchiveWithURL:`): archives that `metal-tt` made with `-macos-version-min` 26.0 or 26.6 were refused, from libraries built for 26.0 or 26.6 alike; those made for 27.0 and 27.2 opened. `pso_offline.py` took the SDK's major version for every platform. That suits the iPhone, whose system is the SDK's, but the Mac app is built for the Mac that builds it, and Xcode 27.1's SDK is 27.0 on macOS 26.6. Also, the Mac host had no deployment target, so clang took the SDK's 27.0, and `pack_app.py` copied it into the app's `LSMinimumSystemVersion`. `metal-tt` also refuses libraries built for a later system than the archive's (`image AIR version (2.9) is bigger than the one of the target (2.8)`, macOS 27.0 libraries for a 26.6 archive); the chain's Mac libraries are built for the building Mac's system, so that Mac's major version is never below them.
- **Fix:** for macOS, `pso_offline.py`'s default `--os-min` is this Mac's major version (`sw_vers`); iOS keeps the SDK's. The Mac host's minimum is `MACOS_MIN` in `dm/mac/min_os.mk` (26.6, as the README asks), on the host's compile lines (`-mmacosx-version-min`, libvpx's too), which the NSP script and `pack_app.py` follow. The host builds with the same warnings at 26.6 as at 27.0 (no unguarded newer API). An archive is still for one major version: an iPhone on iOS 26 cannot open the IPA's 27.0 archive and builds its bases itself (open: one archive per major version would serve both).
- **Found on:** no game: [GitHub issue #673](https://github.com/seedds/nxnative/issues/673), `t_pso_bundle` on an M4 Pro with macOS 26.6.2 and Xcode 27.1.
- **Check:** `dm/tools/tests/test_pso_offline.py` (`test_macos_takes_this_macs_major_version`). `t_pso_bundle` builds its archive with the default and opens it on macOS 27.2: 52 checks pass with its libraries built for 27.0 and for 26.6 (`MACOSX_DEPLOYMENT_TARGET=26.6`), the mirror of the reporter's Mac; the same libraries in an archive made for 26.6 failed the 10 checks. Open: no run on macOS 26.x here.

### Offline pipeline archive on a full disk

- **Symptom:** `[RED] pso_offline: metal-tt failed for applegpu_<gpu> (exit 1)` with only `applegpu-nt command failed` in `metal-tt.<arch>.log`.
- **Cause:** `metal-tt` writes one `.airp` per base and its `translate-*` folders in `TMPDIR` (the work folder), and it does not say when the volume is full. Cruis'n Blast's 33,858 bases failed with 366 MB free. A failed run also left its temporary files in `pso/mac-work`.
- **Fix:** `pso_offline.py` clears those temporary files and stops before `metal-tt` when the work folder's volume has less than about 96 KB per base free (2 GB at least), and names the free space. When `metal-tt` still fails, it clears the temporary files and says when the volume is nearly full.
- **Found on:** Cruis'n Blast (0100B41013C82000, version 1.0.7; no profile in this repository), on a Mac with 366 MB free.
- **Check:** `dm/tools/tests/test_pso_offline.py` (`test_low_disk_space_stops_before_metal_tt`).

### The device's prewarm list unread without a seed list

- **Symptom:** `[RED] pipeline prewarm: seed list does not exist: …/pso_manifest.tsv` at every start, no `pipeline prewarm list …/pso_user.tsv` line, and the same compile stalls every session in places already visited, while `pso_user.tsv` keeps growing; the pipeline archive gains files every start and is never compacted (Donkey Kong Bananza: 88 files, 7 s to open them).
- **Cause:** `warm_start` (`dm/gfx/nvn/nvn_mtl.m`) returned as soon as the profile had no seed list, before it added the device list, so neither list was warmed, and the prebuild of every shader pair's pulling base and the archive's compaction, which run after the prewarm, never ran either. Draws then waited for their pipelines on the submit thread with no hybrid stand-in.
- **Fix:** a missing seed list is logged and left out; the device list is still warmed, and the prewarm returns only when neither list exists. Host only.
- **Found on:** [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md): before, a 13,200-frame session had 154 frames over 100 ms (122 mostly compiling, the worst 7.1 s), 87 s of compiling on the submit thread and 95 draws with a hybrid pipeline; after, the start warmed the 2,484-line device list in 7.6 s (1,002 bases from the archive), compacted 98 archive files into one and ran the prebuild once (20,647 pairs, 116 s before the first frame), and a 17,700-frame session had 68 such frames (40 compiling, the worst 1.9 s), 12 s of compiling and 5,320 hybrid draws. The game has a seed list now.

### Shader libraries loaded when the program is made

- **Symptom:** the first frame that draws with new programs loads their shader libraries on the submit thread (`[longframe] ... shader libraries N, X ms`): tens of milliseconds with the files cached, seconds when the drive is slow.
- **Cause:** a library was loaded when a draw first resolved its program (`prog_get`), or, from the first pipeline miss of a submission, for the whole rest of the submission at once and in parallel (`shader_prefetch_rest`), but still on the submit thread's time. The game names a program's shaders earlier, when it makes the program (`nvnProgramSetShaders`), often while it loads a scene.
- **Fix:** `mtl_prog_prefetch` (called from `nvnProgramSetShaders`) loads each stage's library on a background queue, with a vertex shader's pulling entry as the prewarm does, and the next submission takes it into the shader cache (`warm_claim`), writing it to the prewarm lists as `shader_get` would. A set of the libraries loaded or under way (`progpf_first`, which `shader_get` and `shader_prefetch_rest` also mark) loads each once. A draw that comes before its library has loaded loads it itself, as before. The shader map is read once from any thread (`shmap_read_once`); the prewarm still starts on the submit thread, once (`shmap_load`). A first version put the prewarm inside the same `dispatch_once`, and the pipeline prebuild's own call of `shmap_load` from inside it re-entered it: libdispatch aborted Tears of the Kingdom and Pokémon Scarlet at boot (the regression found it; Donkey Kong Bananza runs no prebuild). `DM_NVN_PREFETCH=0` turns it off.
- **Check:** Donkey Kong Bananza's scripted route, both runs with the drive idle: the first frame's 4 libraries (22 ms on the submit thread) no longer load there; the level opening's 3 still do (0.5-3.9 ms), its programs made just before their draws.
- **Found on:** [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md#found-on-this-game): 3 to 7 libraries a new area, which took up to 4 s a frame in a session with the drive slowed by other work.

### The prepare step inside the shader map's once

- **Symptom:** at boot, `CRASH sig=5 … pc=_dispatch_once_wait.cold.1` right after the `pipeline bases so far` line, with `bt: pso_prebuild __shmap_load_block_invoke … mtl_prog_prefetch h_nvnProgramSetShaders`.
- **Cause:** making `shmap_load` safe from any thread for the [prefetch](#shader-libraries-loaded-when-the-program-is-made) put its whole body in a `dispatch_once`, and that body ended in `warm_start`. With a prepare step (a window, `DM_PSO_PREPARE=1`, or the regression harness's runs), `warm_start` runs the prebuild in place, and `pso_prebuild` calls `shmap_load` again on the same thread inside the once, which libdispatch traps. The prefetch also moved the start of the prepare step from the first submission to the game's first `nvnProgramSetShaders`, on the game's thread. Separately, `warm_start` created the list that finished items go on, so a library the prefetch finished before it, or in a run with no prewarm list, was dropped, and its first draw loaded it again.
- **Fix:** `shmap_read_once` is the `dispatch_once` around reading the map; `shmap_load` reads through it and then starts the prewarm once, outside it. The prefetch only reads the map, so the step runs at the first submission again ([the prepare step](../development.md#the-prepare-step)). `warm_enqueue` makes its own list. Host only.
- **Check:** `dm/gfx/tests/t_shmap_warm.m` needs no game: the map's read leaves the prewarm unstarted, an item finished earlier is kept, and the first submission gets through the prebuild's second call. Before the fix the same sequence trapped with exit 133.
- **Found on:** general to games that run a prebuild: [The Legend of Zelda: Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md) and [Pokémon Scarlet](../../dm/games/0100A3D008C5C000/README.md) at boot, [Blasphemous](../../dm/games/0100698009C6E000/README.md) under `regress.py`, [Hollow Knight](../../dm/games/0100633007D48000/README.md#found-on-this-game) with a window. [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md), where the prefetch came in, runs no prebuild and never showed it.

### Metal's build folders left by pipeline archive flushes

- **Symptom:** the user cache folder (`getconf DARWIN_USER_CACHE_DIR`, then `com.apple.gpuarchiver/`) fills with `gpuarchiver-*` folders of 0.3–5 MB each, one per pipeline archive flush (456 folders, 2.5 GB after two days). Later folders are larger, up to about 127 MB (an M4 Pro, 2026-10-29: 19 GB in 423 folders after a day of Mac runs and regression checks); why they grew is **open**. A later compile can fail with `cannot open '…/gpuarchiver-…/PersistentState': No such file or directory`.
- **Cause:** the Metal 4 serializer (`MTL4PipelineDataSetSerializer`) creates private working folders. The former cleanup inferred ownership from the sole new directory between two snapshots. The mutex protected only this process: if its serializer already had a folder and another compiler created one during the capture or flush, the foreign folder was claimed and deleted. Serialization success concerns the requested output archive, not the lifetime of Metal's private files. The public serializer/compiler API exposes neither a scratch-directory selector nor an ownership token; an empty open-file list does not prove a folder is unused.
- **Fix:** `nvn_mtl.m` leaves Metal's working folders alone during capture, flush and exit. Only caller-owned archive files are removed by compaction. The iOS cold start can remove folders from previous runs in its app-private cache, before this run captures anything. Mac cache reclamation belongs to Metal or an explicit cleanup outside the host: `scripts/clean_metal_cache.py` removes a folder only when nothing in it changed for an hour (`--idle-min`) and its last change is older than the start of every running `dmhost`, so a running host keeps its own; `scripts/regress.py` runs it at its start and when the disk falls under its limit. A folder of another program's live compiler that stayed idle for an hour would still go; none is known. Folders kept this way may still grow between cleanups; disk growth is preferable to deleting another live compiler's state.
- **Found on:** [Pokémon Scarlet](../../dm/games/0100A3D008C5C000/README.md#pipeline-cache-ownership): the ownership defect is reproduced with an isolated foreign-folder fixture. Its longer house run had 20 missing-`PersistentState` failures, but the deleting process was not observed, so that run's attribution remains open. Earlier cache-growth report: MacNativeGames (`1eff989`, `110e6c6`, Metal 3's `MTLBinaryArchive`, 25 GB overnight on Tears of the Kingdom).
- **Check:** the ownership fixture reproduces the former deletion of a foreign `PersistentState`. `t_pso_archive` exercises base reuse, archive lookup and compaction with real Metal pipelines and preserves a foreign compiler fixture created during a compiler call through the subsequent flush. `dm/tools/tests/test_clean_metal_cache.py` keeps folders changed within the idle time or since a running host started, and removes the rest.

### A compaction flush that crashes, then crashes every start

- **Symptom:** a previously working game exits with 139 on every start after pipeline preparation; `dispatch_retain(NULL)` is called by `-[_MTLBinaryArchive materializeAirNTFromOnlineForSlice:entryList:]` under `airntSerializeToURL` and `arch_flush` ([issue #342](https://github.com/seedds/nxnative/issues/342)).
- **Cause:** more than eight run archives enables compaction. The prewarm used `MTL4CompilerTaskOptions.lookupArchives` to recapture their bases into the serializer that also captures newly compiled bases. Metal can crash while materializing that dataset for serialization. It crashes before `arch_compact_finish` removes the old files, so every restart retries it. Framework analysis recorded in issue #342 identifies a nil `MTLBinaryEntry.bitcode`, read from the captured function's library by `bitcodeDataInternal` when the bytes at its bitcode offset are not an AIR wrapper. Which function lacks AIR and why remain **open**. A lead, not a cause: a probe that reused one serializer after each write, with no compile during the write, crashed once in 111 runs with this signature, and a compile under way across a write comes back with its entry's AIR script and bitcode nil ([A compile that does not return stalls every flush and draw](#a-compile-that-does-not-return-stalls-every-flush-and-draw)); the host now never uses a serializer again after its write, and a serializer per period crashed in none of 118 runs; simple synthetic recaptures and a replay with no new bases did not reproduce it. A mixed Tears of the Kingdom replay did, with the reported null call and backtrace, on an M4 Pro with macOS 27.0.1 (26A434).
- **Fix:** work around the framework defect by capturing each compacted base from its original shader functions, with no lookup archives. Metal's ordinary compiler cache still applies. The archive threshold, ordinary archive lookups, deferred flush during prewarm and removal of old files only after a successful write remain in place. A compaction start can do more compiler work; normal starts keep their archive hits. `arch_capture` passes no task options, so captured bases cannot accidentally import archive entries. The recovery guard from the earlier workaround remains: `pso_archive/compacting` records the system build and the capture method before a compaction flush, and is removed when the flush returns. A marker for this build and method suppresses another attempt after a crash. Legacy OS-only markers from archive imports are cleared so the safer capture method can run; a later OS or capture-method change also allows another attempt.
- **Found on:** [The Legend of Zelda: Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md#found-on-this-game), reported on an M5 Max and reproduced with the local shader libraries on an M4 Pro. This is shared NVN archive code, not a game profile rule.
- **Check:** the same 19-file snapshot, 1,490 archived bases and 44 new bases failed in two of three runs with lookup imports and passed all three when only that option was removed. The implemented fix passed three further compactions and their following restarts: each compacted file supplied all 1,534 bases with zero new base compilations on restart. These are library-only manifest replays using the local `e7245308050d` libraries and a matching diagnostic chain header, not a new gameplay run or an iOS test. `t_pso_archive` rejects lookup imports at the real compiler seam; `test_pso_archive.py` exercises disk-built libraries, concurrent recaptures mixed with a fresh function, function constants, compaction and restart in separate processes (`make -C dm/gfx/tests GAME=<id> check_t_pso_archive`). The marker tests retain the recovery guard and check migration from the old OS-only marker, including a fresh-process restart with that marker present.

### Pipeline keys omit unused colour state

- **Symptom:** identical render pipelines compile repeatedly under different keys, lengthening startup prewarm and adding pipeline misses during play.
- **Cause:** NVN blend and channel-mask state persists between draws, including state for colour targets absent from the current pass. `pso_desc` skips targets without a format and reads blend factors only when blending is enabled, but the old keys retained unused state. Deduplicating manifest lines alone also misses old rows whose unused state differs.
- **Fix:** `pso_key_norm` in `nvn_pso_skey.h` clears blend factors when blending is disabled and, for normal passes, the blend-enable bit and write-mask nibble of every target without a format. Merged passes retain their logical blend-enable and mask state because `lmap` maps physical attachments to logical targets; a missing physical format does not imply the same logical target is unused. `pso_key_fill` normalizes live keys, `pso_skey_parse` normalizes historical prewarm rows, and `warm_worker` deduplicates the parsed keys across all input lists before compiling. Depth and stencil formats remain in the key because the `MTLDevice` path needs them, even though Metal 4 bases omit them.
- **Found on:** [The Legend of Zelda: Breath of the Wild](../../dm/games/01007EF00011E000/README.md#files-and-shaders) (disabled blend factors), [The Legend of Zelda: Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md#files-and-shaders) (absent-target state, reported by MacNativeGames).
- **Check:** `t_pso_skey` covers normalized round trips, historical rows differing only in unused state, and preservation of merged-pass logical masks and blend bits; `t_pso_archive` checks disabled blend factors with all targets present. `make -C dm/gfx/tests GAME=<id> check_t_pso_warm` loads two synthetic manifests through `warm_worker`, checks compilation counts and draw-time cache hits, and compares GPU pixels with the original descriptors through both MTLDevice and Metal 4. It includes a merged pass whose physical attachment 0 uses logical target 7's blend and mask state.

### Shader cache key covering the pairs

- **Symptom:** a vertex shader whose pairs changed was restored unpadded from `cas/`.
- **Cause:** `dmsh_build.py`'s shader cache was keyed without the pairs.
- **Fix:** the per-shader key covers the pairs, and a paired vertex shader's key also covers its pixel shaders' keys (`shader_key`, `pair_keyed`), so a change to a profile's pairs rebuilds what it affects ([Pixel shader inputs the vertex shader does not write](#pixel-shader-inputs-the-vertex-shader-does-not-write)).
- **Found on:** [13 Sentinels: Aegis Rim](../../dm/games/01003FC01670C000/README.md).

### Partial vertex builds missing paired pixel sources

- **Symptom:** a full build reports successful shaders, but `metal-tt` rejects its pipeline archive with fragment inputs not written by the vertex shader. Fifteen cached vertices lacked default outputs in DRAGON QUEST I & II.
- **Cause:** a partial build requested only its vertex keys. `vs_pad` skipped absent paired pixel MSL, then cached the incomplete vertex under a key that already included those pixel shaders' input digests. A later full build restored it before pixel-first finishing and never padded that cache hit again.
- **Fix:** include available paired pixel shaders in partial build keys and output cleanup, then finish pixels before vertices. Do not publish a paired vertex cache entry while any paired pixel MSL is missing. The code fingerprint excludes the old incomplete cache entries; the Metal source cache still reuses unchanged complete libraries.
- **Found on:** [DRAGON QUEST I & II HD-2D Remake](../../dm/games/0100D7C01F254000/README.md#found-on-this-game), while building its Mac pipeline archive.
- **Check:** `test_shader_toolchain_cache.py` requests a vertex alone, exercises dependency cleanup and pixel-first finishing, and checks that padding survives publication and a later full-cache restore. A missing pixel source prevents publication.

### Target selection keeps the shader analysis files in the cache

- **Symptom:** a repeated iOS-only shader build crashes with `FileNotFoundError` for `<key>_b.json` in the Mac output directory; requesting only Mac libraries with `IOSOUT` set still compiles iOS side libraries.
- **Cause:** target selection excluded the Mac directory from the per-shader cache entirely, but cache-hit processing, paired vertex padding and the profile tables still read its MSL and reflection. Side-library paths tested `IOSOUT` rather than the selected targets. Prefilling the Metal cache also compiled validation samples for unrequested targets.
- **Fix:** cache and restore the Mac MSL and reflection for every target; require and copy metallibs only for selected targets. A cache entry without its analysis source is a miss. Clear `IOSOUT` when iOS is not selected and restrict prefill validation to selected targets. The NSP script writes only the requested target's reuse key.
- **Found on:** [Celeste](../../dm/games/01002B30028F6000/README.md#shaders), while reviewing PR #275 (resolved).
- **Check:** real Metal output and cache comparisons are recorded in Celeste's shader note above. `test_shader_toolchain_cache.py` covers geometry analysis files, incomplete entries and prefill selection; `test_nsp_mac.py` checks target-specific reuse keys.

### An iOS build after a Mac-only build kept no Mac libraries

- **Symptom:** after `scripts/build_ipa_from_nsp.py <NSP folder> --platform mac`, a plain (iOS) run of the script on the same folder succeeds, but the Mac run command the first build printed then logs hundreds of `[RED] failed to load shader <key>_b: cannot open …/shader/cache/b-<chain>/<key>_b.metallib` lines, until `--platform mac` is run again.
- **Cause:** both builds write the Mac sources and reflection into `<id>_build/shader/cache/b-<chain>/`, which also holds the Mac libraries. A full `dmsh_build.py` run first removes every `<key>_b*` output there, so that a key that left the map or a shader that no longer builds keeps no old library, and then restores or compiles only the targets in `DMSH_TARGETS`. Since target selection ([Target selection keeps the shader analysis files in the cache](#target-selection-keeps-the-shader-analysis-files-in-the-cache)) the NSP script asked an iOS build for `ios` alone, so the Mac `.metallib` files were removed and only their `.metal` and `.json` came back.
- **Fix:** the NSP script (`build_shaders`) asks an iOS build for `mac,ios` when the folder holds the current chain's Mac libraries, and then writes both reuse keys, so a later `--platform mac` run is reused. The Mac libraries are rebuilt, not kept as they were, so none is stale: the Metal compile cache gives back every one whose MSL was compiled before, and only a changed MSL is compiled. The first such build of a folder may translate again, because the per-shader cache entries an iOS-only build left hold no Mac library. A folder without Mac libraries (only ever packed as an IPA) still compiles only the iPhone's. `dmsh_build.py`'s rule, that a full build leaves only what it built, is unchanged; it is part of the chain version, so the fix stays in its caller.
- **Found on:** [Team Sonic Racing](../../dm/games/010092B0091D0000/README.md#found-on-this-game) (resolved).
- **Check:** `test_nsp_mac.py` (`test_ios_after_mac_keeps_mac_libraries`, `test_ios_after_mac_of_another_chain`).

### Old shader library folders grew a game's build folder with each chain change

- **Symptom:** the build folder of a game (`<id>_build/shader/`) gains about 2 GB for Breath of the Wild with every upstream update that changes the shader chain: after two updates `shader/cache` held `b-157e1b96c8da` (2.1 GB) and `b-a78f966a7b38` (2.3 GB) besides the current one, and the next build added a third.
- **Cause:** the libraries are written to `b-<chain>/` (`shader/cache` for the Mac, `shader/ios` for iOS, with the reuse keys `shader/ios/b-<chain>.key` and `shader/mac-b-<chain>.key`) and a chain change starts a new folder. Nothing removed the old ones. The host loads only the current chain's folder, `dm/mac/pack_app.py` and `build_ipa.sh` pack only it, and `regress.py` reads the folder of its own run, so an old folder is dead weight. `dmsh_build.py` took libraries of older folders into the Metal compile cache (`dmsh_mtlcache.py`) before it translated, and `shader_regressions.py` compared the new build with the newest older folder's `_b_build.json`.
- **Fix:** the NSP script (`prune_stale_chains`, called last in `build_shaders`) removes the folders named `b-<12 hex digits>` of other chain versions in `shader/cache` and `shader/ios`, and their keys, once the libraries, their tables and the key of the current chain are in place (a reused build prunes too). It logs one line, `removed old shader libraries b-… (N GB)`. A build that stops or fails earlier keeps the old libraries. The regression check has run by then, and the next chain change compares with the folder this build keeps. The compile cache already holds what the old libraries compiled, since `dmsh_build.py` took them in at the first build of the new chain. The only loss is for a chain that comes back (a checkout of an older commit): its shaders compile again where the cache does not hold their MSL. The compile cache itself (`shader/cache/cas/metallib`) is keyed by the MSL and the compiler, not by the chain, so entries are shared across chains and it is not pruned. It grows only with MSL that is new, and no code removes an entry.
- **Found on:** [Breath of the Wild](../../dm/games/01007EF00011E000/README.md#found-on-this-game) (resolved).
- **Check:** `test_nsp_mac.py` (`test_prune_stale_chains_after_success`, `test_failed_build_keeps_stale_chains`, `test_prune_ignores_other_names`).

### Concurrent warm tasks prepare each shader once

- **Symptom:** startup tasks can prepare the same shader more than once, even though only one object enters the warm dictionary. A measured Tears of the Kingdom startup avoided 320 duplicate preparations; this is not a steady-frame-rate claim.
- **Cause:** an `S` task and an `R` task, or two shader pairs, can miss the same key before either publishes. `shader_load` and the pulling-entry preparation used to run independently for each miss. Discarding the losing `NvShader` does not undo the strong function-descriptor registrations (`fn4_note`, `fn_spec`) or constant-bank reservations (`c1_store`). These ownership mechanisms are deliberate; source inspection alone does not establish a leak or its size.
- **Fix:** `warm_shader` installs a per-key `WarmLoad` under the dictionary lock. The owner prepares the library, reflection, functions, argument encoder, c1 storage and optional pulling entry outside that lock, then publishes before releasing its waiters. A cached lookup allocates nothing. Other keys can prepare concurrently. A failed warm load stays `NSNull` only for that prewarm, matching the previous warm-call behavior; `warm_finish` drops warm dictionaries after all tasks join, and demand loading retains its existing retry behavior. This does not coalesce demand/prefetch loads racing a background prewarm (`DM_PSO_PREPARE=0`), nor merge layer variants.
- **Measurement:** a Tears of the Kingdom startup probe observed 1,537 distinct shader keys. The old racing path performed 1,857 preparations and discarded 320 successful duplicates; per-key coalescing performed 1,537 preparations and discarded none. Observed startup times were 6.63 s before and 6.51 s / 7.15 s with coalescing, so these runs establish duplicate avoidance, not a wall-time improvement or a memory leak. A Breath of the Wild 1.9.0 probe on the same Mac and chain went from 844 preparations of 819 keys to 819, at 3.07 s before and 2.78 s / 2.83 s with coalescing: the same limit applies.
- **Found on:** [Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md#found-on-this-game), through source review of the shared prewarm lookup/load/publish path and a normal-start Mac probe; [Breath of the Wild](../../dm/games/01007EF00011E000/README.md#found-on-this-game), the same probe.
- **Regression seam:** `t_pso_warm` parses actual `S` and `R` manifest tasks sharing a vertex key and gates the first real `MTLDevice` library load. A test-owned dictionary proxy observes the second same-key lookup under the production dictionary lock before the owner is released; a third key must finish independently while the owner is held. Actual device library-load counts check that the shared key is prepared once and cached hits do not reload it; both callers retain one shader identity. It also checks warm negative caching and a later successful demand retry after reflection appears. The existing `check_t_pso_warm` target runs both MTLDevice and Metal 4 paths.

### Function constants specialized at load

- **Symptom:** `fragmentFunction main0 cannot be used to build a pipeline state` (a Metal validation assertion).
- **Cause:** SPIRV-Cross declares an optional function constant, `spvLinearTextureAlignmentOverride`, in shaders that read a texel buffer, and Metal refuses any function with function constants until it is specialized.
- **Fix:** `shader_load` specializes such functions with an empty set of values. On the Metal 4 path the pipeline's descriptor specializes them, not the library: [Functions with function constants specialized by the pipeline, not the library](#functions-with-function-constants-specialized-by-the-pipeline-not-the-library).
- **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md) (one imgui pixel shader).

### Functions with function constants specialized by the pipeline, not the library

- **Symptom:** with Metal 4 pipelines on (the Mac; on iOS `pso_mtl4 on` or `DM_PSO_MTL4=1`), Apple's performance HUD says `Compiled 2 shaders. Total compilation time is 102.70ms. Consider adopting MTL4Compiler for more fine grained compilations` during play, when a shader with function constants loads.
- **Cause:** a function that declares function constants must be specialized before it goes into a pipeline ([Function constants specialized at load](#function-constants-specialized-at-load)). The host did that through `MTLLibrary`'s `newFunctionWithName:constantValues:error:` when it loaded the shader (`shader_load`) and the libraries made from it (the pulling and shadow-input vertex entries, `pull_load` and `vin_fn`; the alpha test, `at_fn`). That call compiles through the MTLDevice, synchronously, on the thread that loads the shader. A Metal 4 pipeline does not need it: `fn4_note` already gives MTL4Compiler an `MTL4SpecializedFunctionDescriptor` with the same values.
- **Fix:** `fn_spec` in `nvn_mtl.m`. With the Metal 4 compiler (`pso4_compiler`) the `MTLFunction` stays unspecialized and is noted with its constant values, so MTL4Compiler specializes it with the pipeline. Its reflection and argument encoders are the same as the specialized function's. A pipeline the MTLDevice builds after all (`pso_build`'s fallback, when a function has no Metal 4 descriptor, such as the `DM_NVN_HILITE` override) takes the function `fn_dev` specializes through the library, once per function. A stable record per function serializes overlapping first uses, so they receive the same specialized function; unrelated functions can specialize concurrently, and registering the entry again preserves its cached result. Unchanged:
  - Compute shaders are still specialized at load: the MTLDevice builds their pipelines.
  - Without the Metal 4 compiler, functions are specialized as before.
  - The variants that differ only in their constant values still specialize through the library: the logic operation entries (`lo_fn`) and the readable-tile versions (`bf_fn`, `DM_NVN_TBDR`). Their pipelines would need one function object per set of values, and `newFunctionWithName:` may hand back the same object.
- **Found on:** [Animal Crossing: New Horizons](../../dm/games/01006F8002326000/README.md#found-on-this-game) on an iPhone 17 Pro Max with Metal 4 pipelines turned on: the HUD warning when a shader loaded during play. On the Mac, with a probe that logs each MTLDevice compile, the specializations came from `shader_load` during the prepare step and again when a shader loaded during play. After the fix none were left, and the picture was the same.
- **Check:** `t_pso_archive`. A function with a constant stays unspecialized, its descriptor specializes it, and a pipeline is made through the Metal 4 compiler. With a fragment function that has no descriptor, the pipeline goes through the MTLDevice with the vertex function specialized once. A gated library holds the first specialization while a second caller requests it: both receive one shared result, an unrelated function progresses, and registering the same entry again does not specialize it again.

### Layer-writing vertex shaders get the triangle class

- **Symptom:** a pipeline for a vertex shader that writes `render_target_array_index` cannot be built.
- **Cause:** Metal builds such a pipeline only with an input primitive class, and the translator writes a point size in every vertex shader, which needs the point class: no class fits both.
- **Fix:** `vs_layer` drops the point size from vertex shaders that write the layer, and the host gives their pipelines the triangle class once Metal names the reason (`layer_out` in `pso_get_` and in the background prewarm). Their passes are layered ([Layered passes for vertex shaders that write the layer](rendering.md#layered-passes-for-vertex-shaders-that-write-the-layer)).
- **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md): two full-screen vertex shaders (`57246f9a04af8529`, `b29f08619f379a83`: one triangle per instance, layer = instance + offset, 20 pixel shaders read the layer).

### The precompile look-ahead restores the scissor

- **Symptom:** the picture bright green and white from some scene on, in runs with pipelines compiled on demand (prewarm lists off or short) and in none with `DM_NVN_PREDICT=0`.
- **Cause:** after a pipeline miss the replay looks ahead through the rest of the submission, replaying state only, to compile the next pipelines in the background (`pso_predict_rest`). It restored `S`, but `state_op` also sets the scissor, which lives in `R`: the replay went on with the last scissor of the submission until its next `OP_SCISSOR`, so the draws in between were clipped to another pass's rectangle.
- **Fix:** the scissor is saved and restored with `S`; the fusion probe (`fz_probe_cov2`) already restored all of `R`. Every game.
- **Found on:** [OCTOPATH TRAVELER II](../../dm/games/0100A3501946E000/README.md): from Castti's opening flashback and the battlefield of Hikari's opening on.
- **Check:** found by bisecting with four runs started at once on copies of one save: `DM_NVN_PREDICT=0` fixed it, a look-ahead that compiled nothing did not, and of its parts only the state ops brought it back; four runs green before the fix, four right after it.

### GPU frame captures need pipelines built through the device

- **Symptom:** a GPU frame capture (`DM_GPU_CAPTURE`, or Xcode with `MTL_CAPTURE_ENABLED=1`) fails with `Capturing Render pipeline specialization is not supported`, or the capture layer crashes on the pipeline archive; a capture of the `MTLDevice` is refused.
- **Cause:** Metal cannot capture pipelines made from Metal 4 bases (`DM_PSO_MTL4`, the default on the Mac), and its Metal 4 to 3 conversion for the capture layer crashed on the pipeline archive. A device the replay drives through Metal 4 cannot be the capture object.
- **Fix:** a run that may capture (`DM_GPU_CAPTURE` or `MTL_CAPTURE_ENABLED` set) builds its render pipelines through the `MTLDevice`, without bases or archive (`pso4_on`), and `gpu_capture_frame` captures the Metal 4 queue (`g_q4`): work on the old queue (present, some uploads) is not in the trace. Such a run builds pipelines more slowly than a normal one, so its frame times are not the game's. Every game.
- **Found on:** [The Legend of Zelda: Breath of the Wild](../../dm/games/01007EF00011E000/README.md) (Mac and iPhone).

### A cold Metal cache looks like a long load

- **Symptom:** the first boot sat on the loading clock for about a minute, the CPU idle, with every file read done after 12 s.
- **Cause:** that run compiled its Metal libraries and pipelines on first use (cold cache). With a warm cache the title is up after about 15 s; what looked like a 40-second wait in the read log (`DM_FSLOG`) was the title screen itself, which reads nothing.
- **Check:** screenshots every few seconds (`ctl.sh snap` from a script) told the two apart.
- **Found on:** [Bayonetta 3](../../dm/games/01004A4010FEA000/README.md) (SDK 13.3.2).

### Geometry runs across occlusion query boundaries

- **Symptom:** hundreds of geometry raster passes in combat, despite support for batching consecutive geometry draws.
- **Cause:** the lookahead stopped at each samples-passed query report/reset. NieR wraps most geometry draws in independent occlusion queries, so each one still captured vertices, ran geometry/compaction and reopened the depth buffer.
- **Fix:** queries remain in their original raster replay positions while lookahead prepares later geometry draws. The shared raster pass uses a separate visibility-result slot after each report/reset; helper capture and compute stages do not count raster samples. Query-result addresses remain marked as report writes for hazard checks. `DM_NVN_GEOMETRY_QUERY_RUN=0` retains the earlier split path for comparison.
- **Found on:** [NieR:Automata US 1.0.3](../../dm/games/010056B015FE8000/README.md): the captured factory combat frame has 697 geometry raster draws with 670 intervening query/event/reset groups. A GPU fixture verifies that the first visible draw reports 216 samples and the second, hidden draw reports zero, with the completion event and later color/depth coverage preserved.

### Prepared geometry indices reused after byte comparison

- **Symptom:** geometry preparation dominates rendering CPU even after geometry passes have been batched; primitive assembly and capture-index deduplication repeat for unchanged meshes.
- **Cause:** each draw decodes all source indices, assembles its primitive map and rebuilds sorted unique capture IDs, even when its source bytes and topology are identical to the previous frame.
- **Fix:** a 32 MB / 4,096-entry LRU cache keeps only the CPU primitive map, unique capture IDs and original index bytes. The key includes address, count, index width, topology and restart settings. Every reuse compares all original bytes, after the existing readability and GPU-write wait checks; there is no content-hash collision assumption or dependence on write-protection. CPU edits, GPU copies and reused storage with different contents rebuild preparation. Base vertex, instances, vertex capture, transforms and GPU outputs remain per draw. `DM_NVN_GEOMETRY_INDEX_CACHE=0` disables reuse for comparison.
- **Found on:** [NieR:Automata US 1.0.3](../../dm/games/010056B015FE8000/README.md): the outdoor flight-unit encounter has approximately 2,300 draws including 855 geometry draws, 24.9 ms rendering CPU and 12.2 ms GPU time. A controlled CPU-only preparation benchmark (64 draws, 32 meshes, 6,144 indices each), alternating warmed paths, drops median CPU preparation from 2.783 to 0.149 ms; it does not execute the GPU geometry stages or establish a gameplay FPS. Production GPU tests preserve literal pixel coverage on unchanged repeats, CPU index edits and GPU index copies, including forced-wait and disabled-cache paths. Cache budget and topology/restart separation are covered too.

### Geometry capture initializes only missing words

- **Symptom:** geometry vertex-capture passes consume substantial GPU time before geometry compute or depth rasterization.
- **Cause:** the capture entry first zeroed all 144 words of every output record, then wrote the default W components and finally overwrote the actual shader outputs. Much of the initialization was redundant GPU storage traffic.
- **Fix:** `dmsh_geometry.capture` initializes only words absent from the captured outputs. The capture slot map stores each selected word once, including defaults for absent outputs. Position W, point size, generic W, clip distances, layer and viewport retain their existing bit values. Full capture selects all 144 words; lean capture selects its consumed words. Shaders and host must be rebuilt together.
- **Check:** the synthetic throughput comparison alternates control/candidate order and verifies all 9,437,184 captured words against both the original buffer and literal expected values. Mac runs vary in timing, so their percentage gain does not establish iPhone gameplay FPS. The optional real-translator capture fixture in `t_geometry` also checks all uncaptured defaults. All 1,225 NieR shaders build for Mac and iPhone; standard geometry replay and dependency checks pass.
- **Found on:** [NieR:Automata US 1.0.3](../../dm/games/010056B015FE8000/README.md#factory-combat-capture-and-9s-texture-follow-up), iPhone City Ruins: 14 geometry groups in the captured frame spend 7.485 ms in vertex capture, 6.512 ms in compute and 1.019 ms in depth raster. The accompanying pass dump affects total frame time; its large uncovered GPU interval must not be treated as a clean gameplay hitch measurement.

### The host's own shaders compiled from source at run time

- **Symptom:** Apple's performance HUD (`DM_METAL_HUD=1`) says `Compiled 2 shaders. Total compilation time is 102.70ms. Consider adopting MTL4Compiler for more fine grained compilations` during play, with a dip in the frame-rate graph at that moment, for example when a new room first issues a scissored clear.
- **Cause:** the host compiled the shaders it writes itself from MSL source at run time (`newLibraryWithSource:`): built-in kernels, clears, copies, `nvnCommandBufferDrawTexture`, the RGB565 red / blue exchange, index conversions, geometry packing, transform feedback, the present, the GPU deswizzle and the OpenGL layer's clear quad. The thread that asked waits for the compiler. The clear triangles were one library per colour slot and type (`snprintf`), up to 24, each compiled when a clear of that kind was first met: about 50 ms each on an iPhone, on the submitting thread. `MTL4Compiler`'s `newLibraryWithDescriptor:` with a source is no way around it: it calls the device's `newLibraryWithSource:` itself (a check with a swizzled device shows the call), and the HUD counts that too.
- **Fix:** the sources are `.metal` files in `dm/gfx/msl` (the clears in one library: `clr_fs_<slot>_<float|uint|int>`). `dm/gfx/tools/host_msl.py` compiles them when the host is built: `make -C dm/host` for the Mac, and `gen_build.py` with `MSL_SDK=iphoneos` for the app's deployment target on iOS. It writes `dm/host/gen/msl-<sdk>/host_msl_gen.h`, which holds each library's metallib and source. `host_msl_lib` (`dm/gfx/nvn/host_msl.m`) loads one by name with `newLibraryWithData:`, once per device, so nothing is compiled; the pipelines are made as before. A build without the Metal toolchain carries the sources only: the host then compiles each library once and logs a `[RED]` line (`compiled from its source in … ms`). The first load logs `host shader libraries: 13, compiled for <sdk> by <metal version>`. The diagnostic overrides (`DM_NVN_HILITE`, `DM_NVN_MAGSCAN`, `DM_NVN_TEXDUMP`) still compile theirs from source.
- **Found on:** [Animal Crossing: New Horizons](../../dm/games/01006F8002326000/README.md#found-on-this-game), iPhone 17 Pro Max: the HUD's compile warning in play, from clear-triangle libraries compiled when a new room first issued a scissored clear.
- **Check:** `t_host_msl` loads every library, with a swizzled `newLibraryWithSource:` that counts 0 calls, and finds all 24 clear functions. `t_builtin`, `t_geometry` and `t_deswz_gpu` use the built libraries.
- **Build inputs:** Mac libraries use `MACOS_MIN` from `dm/mac/min_os.mk`, matching the host's deployment target; otherwise a newer SDK emits libraries older supported Macs cannot load and the host falls back to runtime compilation. Every build rechecks the compiler, SDK, deployment target and source contents, reusing the header only when all match. Installing the Metal toolchain after a source-only build therefore regenerates the libraries without a clean build.
- **Build callers:** the NSP packer and regression harness generate the Mac header before checking their cached host executables, and include it in the host fingerprint. The standalone iOS compile check generates the iPhone header before its dry run, as the app build does. A dry run alone does not generate either header.