nxnative-easy/docs/agents/04-graphics.md

04 — Graphics: NVN on Metal for a new title

The graphics layer records NVN commands into its own command format (dm/gfx/nvn/hle_nvn.c) and replays them as Metal 4 at submit time (nvn_mtl.m). Shaders are never translated at run time: the original Maxwell shader binaries are translated offline (Ryujinx's translator → SPIR-V → spirv-cross → metallib) and looked up by key. See dm/gfx/README.md for the full pipeline; this page lists what changes per title.

1. Find every shader container

Do not stop at the obvious archive. In Mario the main bfsha held 4,888 of the 12,811 shaders the game actually uses.

dm/gfx/tools/shader_split.py <RomFS> <ExeFS main> <pkg dir> <map out> <facts out> [<pairs out>] scans every RomFS file (zstd-decompressed, SARC expanded recursively) and main, and recognizes the containers below. The pairs table lists the (vertex key, pixel key) of every BNSH variation or Unity program (section 6). Unity shaders need a second pass: files whose type trees are missing are scanned again once the bundles have supplied the Shader tree. For a base game plus an update, give RomFS and main as comma-separated lists, base first: keys are content hashes, so the output is the union of both versions and one shader build serves both (Mario 1.2.1 recompiled 3,451 of its pixel shaders; v0 rows keep their package indices).

• BNSH (BNSH\0\0\0\0 + grsc block at +0x60): 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 relative to the BNSH start.
• Unity Shader objects (Engage): 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). 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). 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…). 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. Engage: 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.
• Unity 2022 Shader objects (Ace Attorney Investigations Collection, Unity 2022.3.17f1): the program code keeps the 0x60-byte stage header, 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); unity_code_blobs keeps only those. 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). 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. 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). armyja's issue #17 proposed UnityPy and a reader of its own; neither was needed.
• MSHD (Metroid Dread, MercurySteam's engine: system/shd/*.bshdat, the only place its 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). 478 files, 1,342 stages, 1,010 unique keys (359 vertex, 651 pixel), 659 pairs.
• GLSLC outputs in main (13 Sentinels, Vanillaware's engine: 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. 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. Engage's main also has 11 of them, so shader_split.py reads them only when the profile sets shader_main_glslc on (and accepts only outputs that pass the control / code / length check). The stage comes from SPH word 0, as for BAHS.
• G2S in RDB archives (Hyrule Warriors: Age of Calamity, Koei Tecmo's KTGL): 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); one _DOK0000 container (type 0x20a6a0bb, the engine's system resources) holds 151 more bare GLSLC outputs, 21 of them drawn in the first battle. shader_split.py maps the RDB files and inflates only entries that start with those two magics (idrk_shaders), with shader_main_glslc on for the 117 outputs in main. 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. 40,726 shaders (37,043 pixel, 3,440 vertex, 241 compute, 2 geometry), Mario's texture handle layout (0x20 + 8 × slot); pairs come from shader_separable on. The older KTGL of Three Houses packs the same G2S containers differently (next item).
• SPKG in LINKDATA (Fire Emblem: Three Houses, the older KTGL): shader packages (SPKG: patch*/nx/shader/*.spkg and 132 entries of DATA1.bin) of G2S containers as above, one _SPK4100 package without G2S (the post effects, DATA1.bin entry 549), and map and model archives carrying packages of their own. kt_shaders inflates every LINKDATA entry and loose *.gz / *.spkg file (01), 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). Its GLSLC (output version 0x1d.2) writes the 0x848 / 0x850-byte controls and headerless compute code described above, which glslc_shaders used to drop: 24 shaders before, 1,355 after. 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. Nine more shaders exist only at run time (Seen in Fire Emblem: Three Houses, below).
• SDF in RE Engine PAKs (Monster Hunter Rise): shader definitions (SDF\0, 2,685 files in the PAK view, 01): 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. kpka_jobs maps the PAKs, takes the newest copy of each path, peeks at 4 decompressed bytes and inflates only SDF entries (sdf_shaders): 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.
• Unreal Engine shader code archives (Octopath Traveler 0: 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): 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. 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; Octopath Traveler 0: main+0x4f77e10, found from the function name its profiler hook passes) decodes them in the Mac host: profile key shader_ue_decode, DM_OODLE_JOB (dm/host/hle_oodle.c), the NSP script's decode_ue_shaders. 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). The host keys came out exactly: 37 of 37 shaders a first run reported missing were in the two archives. 10,324 unique shaders (1,457 vertex, 6,346 pixel, 2,521 compute); vertex and pixel shaders combine per draw, so shader_separable on. Bindings: separate textures and samplers, Engage's layout (0x168 / 0x568); 61 vertex shaders pick a texture slot at run time (profile key tex_dyn_slots, below), 2,084 shaders do atomics on buffer images (below). 16 compute shaders do atomics on a 3D r32ui image, which SPIRV-Cross's MSL 2.3 image-atomic emulation indexes with an int3 (array subscript is not an integer): all are Unreal's paged global distance field (GlobalDistanceFieldCompositeObjects.usf:CompositeObjectsIntoPagesCS), which the mobile renderer does not dispatch (none in any run), so they stay unbuilt.
• BAHS (agl sharcb, 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}. Control sections are shared between programs, so pair by entry, never by searching for code near a control.
• Shorter control sections (SDK 10, Super Mario 3D World + Bowser's Fury): 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. The key therefore 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.
• Compute code without a header (SDK 10): 3D World's particle compute shaders start their instructions at code +0 (like D2R), not after a 0x12345679 header. shader_split.py accepts both; dmsh_build.py strips only a header that is there.
• 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.
• The program start must be exact. 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, instructions were dropped or misread, and 12 compute shaders (particles, post effects) translated into an endless loop that hung the GPU at the first gameplay scene. Check 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. dmsh_build.py now refuses compute code whose padding is not zero.

Validate a new title's coverage at run time: DM_NVN_SHDUMP=<dir> makes the host dump the control section and code of every program it can't find in the map; search those bytes in the game files to find the container that was missed.

2. Key shaders by control section + code

The host key is sha256(control[0:0x880] + code)[:16] (prog_get in nvn_mtl.m, shader_split.py). D2R keyed by the control section alone; in Mario, 6,367 shader pairs share a control section but differ in code. A different code length or layout in a new title must be reflected in both places.

3. Find the binding layout by census

The NVN driver puts binding tables in driver constant banks, and the layout changed between NVN versions. Don't guess: translate everything once (the translator alone takes seconds) and tabulate what the shaders reference.

Binding	D2R (NVN 53)	Mario (NVN 55)	Engage (Unity, SDK 12)	How it was measured
Bound texture handles (texture buffer)	0x168 + 8 × slot, 128 slots	0x20 + 8 × slot, 32 slots	0x168 + 8 × slot, 128 slots	Mario: immediate handles of all TEX instructions, 0x20…0xF0 in steps of 8. Engage: the translator's texture names (fp_t_cb<slots>_<handle>, Ryujinx TextureHandle packing) give word 0x5A = 0x168 for the texture and 0x15A = 0x568 for the sampler
Image handles	0x120 + 8 × slot	0x120 + 8 × slot	not measured (no compute shaders)	agl compute shaders: 0x120 / 0x128 / 0x130
Separate samplers	0x568 + 8 × slot	not used	0x568 + 8 × slot (Unity binds only separate textures and samplers)	—
Uniform buffers	c[3 + slot]	same	same
Storage buffers	c0[base + 16 × slot], base vs 0x110 / cs 0x310 / ps 0x510	same	not measured
Largest point size	not measured	not measured	not measured	Super Mario Odyssey (SDK 7): GLSLC clamps a shader-written gl_PointSize to c0 + 0x610 (alPointSpriteShader: min(size, c0[0x610])). It is not the SetPointSize value, which the game leaves at 1. The host writes 2047 there for the graphics stages

Ace Attorney Investigations Collection (Unity 2022.3, SDK 17.5.3) uses Engage's layout: built with nvn_tex_base 0x168 and nvn_samp_base 0x568, all 4,950 translatable shaders decode every binding (shaders with a binding whose slot was not decoded 0), and 190 read c2 rows as data.

Metroid Dread (SDK 11.4.3) uses Mario's texture handle layout, 0x20 + 8 × slot: the TEX / TEXS handle immediates in its shader files' own SASS text are words 0x8–0x22 in steps of 2. Its full build leaves no unresolved bindings and no unfillable driver-bank reads; the other rows are not measured separately.

Hades II (SDK 19.3.9) has Mario's layout too, but its 88 shaders bind textures at fixed, sparse slots up to 34 (fp_t_tcb_4C: c2 word 0x4C = 0x20 + 8 × 34, in SceneFade.frag). With the profile's 32 slots, dmsh_build.py reported that binding as not decoded ("shaders with a binding whose slot was not decoded 1", which prints a red line when the shader loads). Its nvn_tex_slots is 35. The census of fp_t_* names in the translated libraries gives the slots in use; the host's tables hold up to 128.

Monster Hunter Rise (RE Engine, SDK 11.4.3) uses both tables at once: combined handles from nvnCommandBufferBindTexture at 0x20 + 8 × slot and separate textures from BindSeparateTexture(s) at 0x168 + 8 × slot, with separate samplers at 0x568, and binds slot 0 of both in the same frames. A census of 23,478 translated shaders' texture names: 33,813 combined (fp_t_tcb_1E), 29,342 separate textures with a constant sampler ID (handle type 3, fp_t_cb65538_3041005C: texture at c2 word 0x5C, sampler 0x104 ORed in), 8,265 separate texture + separate sampler (type 1, fp_t_cb196610_1017815A), 5,227 separate textures whose own handle carries the sampler (fp_t_tcb_5A), about 1,500 uniform-buffer handles with a constant sampler (fp_t_cb65541_30410028), and 121 fully dynamic handles (type 4, fp_t2d_cb0_40000000, not supported). The profile keys nvn_septex_base / nvn_septex_slots give the separate table: separate binds go to slots 128 + n of the replay's table and their handles decode to 128 + n; a constant sampler is the binding kind samp_id (NVN_RB_SAMPI, reg = the sampler ID).

The profile's nvn_* keys hold these constants; dmsh_build.py decodes the bindings from them (decode_tex). Program.cs still has Mario's TexBase built in, but uses it only to look up per-slot texture / image types in the facts table, so a profile with other bases builds as long as its facts table has no such entries (Engage's has none; dmsh_build.py stops otherwise); the replay layer fills driver banks with the same layout: the texture, image and separate-sampler tables of c2 (NVN_RB_DRV2 in encode_args). Some shaders read c2 rows as data at run time (Engage: 243, the texture and sampler handles of slots 0–11); the build lists reads outside those tables as cannot fill. After a full build, the summary must say shaders with a binding whose slot was not decoded 0.

4. Bindless textures

Mario's shaders read 38 % of their textures through handles stored in the game's own uniform buffers (fp_t_cb7_2A = handle at word 0x2A of c[7] = uniform slot 4). These became binding kinds tex_ubo / samp_ubo (NVN_RB_TEXU / NVN_RB_SAMPU): at replay the host reads the 32-bit word from the bound uniform buffer (low 20 bits texture ID, high 12 bits sampler ID).

Consequence: any optimization that infers "which textures a draw reads" from bound slots is wrong for such shaders. The TBDR optimizations (store / load elision, pass merging) are therefore off by default in this repository (tbdr_on(), DM_NVN_TBDR=1 to enable).

5. Depth convention

nvnDeviceSetDepthMode near = −w (GL, z in [−w, w], NVN's default) versus near = 0 (D3D / Metal, z in [0, w]). Metal clips anything with z < 0, so a GL-mode game's vertex shaders must end with z' = z/2 + w/2. The mode is runtime state, and games differ: Mario, Metroid Dread, Super Mario 3D World and Breath of the Wild keep GL; Link's Awakening, 13 Sentinels and Engage set near = 0.

So the translator leaves z alone (QueryGraphicsState depth mode false in Program.cs), and dmsh_build.py (vs_depth) writes two entry points into every vertex library: main0 with the GL rewrite (the same arithmetic Ryujinx emitted) and main0_z without it (reflection entry_z). The host picks one when it loads a vertex shader, from the mode the game set (mtl_set_depth_zero, shader_load in nvn_mtl.m). A new game needs no setting; if a game changed its mode after vertex shaders were loaded, a red line says so.

• Rewriting a near = 0 game for GL still draws correctly (depth tests keep their order), but every depth lands in 0.5..1, so effects that turn depth back into distance are wrong: Link's Awakening's outdoor depth of field blurred the whole screen.
• Not rewriting a GL game: the full-screen copy into the window is clipped away, so the window stays black even with a solid-colour fragment shader.
• Metal function constants could pick the variant too, but any function with a function constant has to be specialized before a pipeline can use it (measured: about 2.4 ms per vertex shader on an M4 Pro); the second entry point only costs build time and about 30% more library size.

6. Other per-title surprises seen in Mario

• 8-bit index buffers: Metal has only 16/32-bit indices. Widened to 16-bit at replay into the frame arena. Mario's full-screen passes and its final copy use them.

• Texture swizzle: TextureBuilderSetSwizzle / TextureViewSetSwizzle are applied with Metal swizzled texture views (NVN and Metal channel enums have the same values). The font atlas is (1, 1, 1, R); without it glyphs render as white boxes.

• Render-target views that change the format: Mario's last pass writes the RGBA8 swapchain texture through an RGBA8_SRGB view (TextureViewSetFormat 0x38 on a 0x25 texture), so the GPU encodes the linear shader output to sRGB on write. The host used to attach the base texture, which stored the linear values unencoded. The whole picture came out darker and more saturated on the Mac and the iPhone: comparing with a Switch screenshot, every channel at every level equalled the Switch value decoded from sRGB once more (0.55 → 0.26, 0.75 → 0.52). attach now attaches a Metal view of the view's format, so the pipelines get that format too; the view is kept until the submission completes. A format change that is not an sRGB / linear pair needs MTLTextureUsagePixelFormatView on the base texture; without it the host draws in the base format and reports render target view format … cannot be attached. DM_NVN_PASSDUMP prints these views as V <pass> c<i> view format …; Mario uses only this one (1 of 80 passes in a course frame). The present step has the matching rule: the Switch scans out the stored bytes, so an sRGB-typed swapchain texture (RGBA8_SRGB / BGRA8_SRGB) is sampled through a non-sRGB view (present_raw_src); sampling it directly would decode it once more on screen only, which a DM_NVN_SHOT capture (raw bytes) does not show. Mario's swapchain is RGBA8, so this path is for other games (t_present_crop covers it).

• Transient render targets: the game builds new texture objects every frame over the same storage and never finalizes them. Host textures are shared by (pool, offset, description); otherwise each new object allocates a new private Metal texture (≈13 MB/frame leak).

• Inline uniform updates (CommandBufferUpdateUniformBuffer) are replayed as in-order GPU copies; reading the uniform buffer from the CPU at encode time gives the frame's last value, not the draw's.

• Compute code areas have a header in Mario (0x12345679), not in D2R; dmsh_build.py strips it.

• Virtual memory pools (flag 0x200): aliased onto the physical pool named by MapVirtual when all requests share one base (Mario's case); otherwise unimplemented.

• SPIRV-Cross failures: 2 agl post-effect compute shaders still produce SPIR-V that SPIRV-Cross can't parse (nullptr exception). Before the compute start-offset fix it was 7: misdecoded code produced the bad SPIR-V.

• CPU-written linear textures: a texture with NVN_TEXTURE_FLAGS_LINEAR or LINEAR_RENDER_TARGET in a CPU-visible pool (Mario's glyph cache; its four 64×1 RGBA16F sky colour ramps, without which the sky is black) is created on the pool's MTLBuffer (tex_linear), so CPU writes (Mario's glyph cache) are visible to the GPU. Block-linear CPU-written textures still need the packaged-texture re-upload path.

• Integer vertex inputs: without PC reflection the translator declares every vertex input float4. Skinned models read their bone index (R8UI / RG8UI / RGBA8UI) as raw integer bits (as_type<int>(in.m_N.x) * 48); Metal converted the integers to floats, every bone index was garbage, and all characters were stretched into streaks. dmsh_build.py (int_inputs) now re-translates a vertex shader with uint inputs when every use of an input is an integer bit cast (529 of 4,854 in Mario), and the host makes the Metal vertex format agree with the declared type (vfmt_for_ct; signedness swaps of 32-bit components and 32-bit float→int reinterpretation keep the bits, while 8- and 16-bit integers of the other sign are widened on the CPU, see Three Houses below). A remaining mismatch is reported as vertex attribute format … does not match shader input type ….

• Vertex formats: Mario uses RGB10A2SN (0x69) → MTLVertexFormatInt1010102Normalized. Unmapped formats fall back to a 4-component read and print vertex attribute format … not mapped once. The integer-read-as-float formats (*_UI2F / *_I2F, 0x51–0x68), which Metal has no vertex format for, are converted on the CPU at replay like the other conversions (vconv_i2f); Breath of the Wild has an RG8_UI2F attribute.

• Instanced array draws (DrawArraysInstanced) carry instance count and base instance in the OP_DRAW record.

• Metal 4 does not order operations inside one encoder. The host batches consecutive copies, fills and inline uniform updates (UpdateUniformBuffer) into one compute encoder. Mario writes the same uniform slot twice in a row each frame (4 bytes of −1, then 16 bytes of wave coefficients); without an intra-encoder barrier the writes occasionally landed in reverse order, the fountain's GPU wave simulation (a 512×1 RGBA16F texture updated by render → copy back → splash pass every frame) ran with −1 as its self coefficient, diverged to ±65504, and its vertex shader pushed the water surface to the top of the screen (tall blue / textured panels that flickered frame to frame). blit_buf now inserts barrierAfterEncoderStages when an operation reads or writes a guest range an earlier operation in the same encoder wrote; texture copies always get one. DM_NVN_BLIT_ORDER=0 turns the barriers off to reproduce. Render passes and compute dispatches each have their own encoder with full barriers, so they were not affected.

• A texture slot picked at run time (Octopath Traveler 0: 61 Unreal vertex shaders, particles reading their data from buffer textures) → the shader reads one of several c2 texture-table words, chosen by branches, and ORs it with a sampler-table word. 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 and they used not to load. 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) and binds slot k at element k, so the shader's own arithmetic lands on the binding. 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). The host logs each such shader's first encode with the slots that held textures. Host only: the chain version does not change.

• Image atomics → SPIRV-Cross's MSL 2.3 emulation declares the image twice: the texture (texture_buffer<uint> cp_i_cb2_4C_r32ui) and a device atomic_uint* cp_i_cb2_4C_r32ui_atomic pointer to the same memory, which the build read as a second image of the same slot, so the replay put a texture resource ID where the shader wanted an address. Octopath Traveler 0 has 2,084 such shaders (Unreal's GPU particles); three ran in the first 100 seconds. The pointer is now its own kind (NVN_RB_IMGA), bound to the storage of the image in that slot: a buffer texture is a view of its memory pool, so the atomics and the texture reads see the same words. A block-linear image is a deswizzled Metal copy whose pool memory has another layout, so its atomics go to a dummy buffer (logged red) until a game needs them followed; a 3D image does not compile at all (above).

• Buffer textures at offsets below 16 bytes → Metal reports a 16-byte alignment for texture buffer views, and the host used to leave a view at a smaller offset as a dummy (Octopath Traveler 0: one at pool offset 0x168, read by the particle shaders above). On an Apple GPU a view at any multiple of the element size reads the right elements (measured on an M6 with R32Float views at offsets 4, 8 and 12), so the host now creates it there; below the element size it stays a dummy.

• New textures start as zero. Metal leaves a new texture's content undefined; a freshly mapped Switch memory pool reads as 0, and feedback effects (such as the wave simulation above) never clear their state texture. Non-packaged colour textures are zero-filled from a shared zero buffer when created (tex_zero_fill, batched with the uploads).

• Scissored colour clears: NVN ClearColor honours the scissor. Clears are normally folded into the next pass's load action, which only works for a full-target clear; Mario clears sub-rectangles of small targets to zero (pyramid levels packed into 480×272, 256×256 and 64×36 targets). A zero clear outside a pass with a partial scissor is now a copy from the shared zero buffer into that rectangle of the attachment's level / slice (clear_rect_zero). DM_NVN_CLEARLOG=1 logs clears that are still skipped (mid-pass, partial channel mask or non-zero colour with a partial scissor).

• Inline uniform updates are snapshotted at encode time (mtl_sub_update): the data is copied into a host-owned ring instead of being read from command memory when the GPU runs.

• 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. Metal refuses to build the pipeline (Fragment input(s) user(locnN) mismatching vertex shader output type(s) or not written by vertex shader), and the draw is skipped. Translating one shader at a time cannot know what the partner reads, so shader_split.py writes the pairs of each BNSH variation (shader_pairs.tsv) and dmsh_build.py gives each vertex shader the outputs its partners read but it doesn't write, as float4 (0, 0, 0, 1) (vs_pad; vertex shaders with partners are finished after the pixel shaders, and the partner list is part of the shader's cache rule). In Mario 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). Check all pairs offline after a build (recipe in 05). The host caches failed pipelines and reports each once with its vertex / fragment keys ([RED] failed to create render pipeline (vertex … fragment …)); before that, every draw retried the failing compile.

• Unwritten components of a written output are (0, 0, 0, 1) too. 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 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. In Link's Awakening 12,281 of 18,410 vertex outputs never write w.

• Positions must be invariant across shaders. 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 (Super Mario Odyssey). Maxwell evaluates identical position math identically; Metal only does for outputs marked [[invariant]] (the translation marks gl_Position) when compiled with -fpreserve-invariance, which dmsh_build.py (Mac) and ios_metal.py (iOS) pass to every compile. Without it Odyssey's ground kept stale 32×32 blocks where the G-buffer draw failed the depth test.

• iOS math mode: the iOS libraries are compiled with -fmetal-math-mode=relaxed (dm/ios/tools/ios_metal.py). 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. Relaxed keeps fast math's reassociation but not that assumption.

• Dynamic present interval (Mario 1.2.1 on the iPhone): Mario times every frame (main+0xcfd24); after a frame that ran past one vertical sync it sets present interval 2 for the next frame and sets its own time step to the interval (a float at main+0x36ceff0), so that frame catches up 2/60 s. A user's 2-minute iPhone session switched 123 times, mostly for one frame. In display mode those frames were shown for 16.7 ms and the game visibly jumped ahead ("sometimes the speed increases when jumping"). In game mode (the default since 2026-10-15) the display-link callback keeps a frame whose period is over 25 ms on screen for its whole period before taking the next one (g_link_hold_until in nvn_mtl.m; the display link up to tick log line counts these ticks as held for the present interval), and the 30 fps target that drops the display link to 30 Hz (nvn_game_fps in hle_nvn.c) waits until interval 2 has held for a full second, so one-frame switches no longer flip the display-link mode. Fixed: the game side of that hold (BotW on the iPhone, "when pressing B, Link's move pace is speeding up"). The acquire waited for one display-link tick per frame whatever the interval, so the game kept producing frames every 16.7 ms while the screen held each interval-2 frame for 33 ms and dropped the one in between (old frames dropped by merging). BotW sprinting flaps between interval 1 and 2 for runs of 6 or more frames, and its VFR time step at interval 2 is a whole 30 fps tick, so those runs went at double speed. mtl_link_wait_tick now waits for as many ticks as the game's present interval (times 2 on the 120 Hz link, half rounded up on the 30 Hz link), as the Switch blocks the next acquire for N vertical syncs. The windowless Mac host was never affected: there the acquire sleeps until the interval's slot. DM_NVN_TSV writes each frame's interval as @itv. DM_TEST_PRESENT_LAG=90,30 (sleep 30 ms before every 90th present) makes Mario do this on the Mac; run it with DM_WINDOW=1 to get the display link.

Seen in Engage (Unity)

• Program control sections are freed after nvnProgramSetShaders: shader keys are computed inside that call (03, graphics entries).
• CPU-written block-linear textures (nvnTextureWriteTexels: video frames, engine default textures) are uploaded and re-uploaded from guest memory like packaged textures (03).
• Video: Unity's Hidden/VideoDecode samples luma as .w of an R8 texture with the swizzle (0, 0, 0, R) (Unity's Alpha8) and chroma from an RG8 texture at half size. To check the movie path apart from the decoder, fill the frames with a gradient in hle_movie.c; it should appear on screen.
• One geometry shader (in unity_builtin_extra): Metal has no geometry stage, so it fails to compile (unknown main0_out); nothing in the opening scenes uses it.
• Fixed: a crash a few minutes into the prologue battle (scripted run, about 357 s). It wrote to NULL + 0x1fffe0 at main+0x9a4664: Unity's TLSF heap allocator (main+0x9a45d0) got no address region for a new 2 MB pool. The alias region was out of address space, not memory; see 03, "A null-based write minutes into play". A 12-minute scripted run now plays through the battle.
• Fixed: GPU memory grew with every scene change (the idle title demo reached 5.4 GB of GPU memory and 3.6 GB resident in 850 s; it now stays at 1.2 to 1.4 GB). Unity frees memory pools on scene changes and builds new ones, which Mario never did in its first course. Two host structures kept freed pools alive: the residency set still held each pool's no-copy MTLBuffer after nvnMemoryPoolFinalize (now retired through deferred release), and the storage-sharing texture table (g_alias, keyed by pool object, offset and description) never evicted, so the old pools' render targets stayed, and linear textures built on a pool buffer kept the whole buffer alive. Finalizing a pool now evicts its entries (mtl_alias_drop_pool). The memory up to frame line now also reports serial objects (host objects by serial), freed but alive (finalized pool buffers still alive) and shared textures (sharing table size and evictions); growth in GPU memory with those counts flat points to an allocation outside them.
• Probably fixed: about one run in four stalled early (around the opening movie or the title; no log output, not even the time-limit report). The stalls happened before the real thread pause for IL2CPP's GC (03); since then 12 runs (eight of them 200 s runs through the movie and title, four of 330 to 850 s) have not stalled once, where one in four used to. Without the pause the collector ran while other threads kept changing the heap. A second candidate, fixed later: a race in the dirty-page tracker (nvn_dirty.c, ported from d2r-mobile) could leave a page read-only while marked writable, so a game thread writing it spun forever in the fault handler; Engage's movie decoder writes frames into texture storage from a host thread while the replay thread protects pages (tests/t_dirty_race.c reproduces it). If a stall comes back, sample the process when its log stops growing (sample <pid> 3) to see where every thread waits; the crash handler now writes one line with write() and exits if the stderr lock stays taken for 1 s, so a stall that holds that lock shows up as such.
• Fixed: the idle attract demo crashed after about 220 s. Unity's particle geometry jobs wrote through a NULL index pointer on several job threads at once (main+0xdb0ad4, 0xdaea58, 0xdadb5c). The scheduler (main+0xd91e0c) takes its vertex and index chunks from a fallback that creates and maps a fresh buffer each time (main+0x7b9c64; Unity's Switch renderer always takes it, as on hardware), and from about 180 s 32-bit mesh indices also turned up inside Unity's texture objects. The cause was the alias region running out of address space, as in the battle crash (03, "A null-based write minutes into play"): past 2 MB the index buffers got no mapping. With the node-area cut an 850 s idle run is clean, with no corrupted texture objects; with the old SDK the same run dies at 159 s in Unity's heap allocator. Reproduce the old state: boot.sh <log> 245 with no input on an SDK rewritten without the cut.

Seen in Metroid Dread

• Movies in native-texture mode: the decoder writes each frame into an RGBX8_SRGB NVNtexture that the game hands over in a movie::Buffer (03, movie entry). The host writes the texture's block-linear storage on the CPU and the replay layer re-uploads the dirtied pages, as for WriteTexels.
• WriteTexels one texel at a time: the game's set-pixel helper (main+0x160e80) writes a 1×1 region per call (probably the map / minimap mask). The host used to decode the whole level, patch it and encode it again for every call, and after a few minutes of play the main thread spent whole frames there (frames stopped at about 175 s). Non-3D regions are now written in place by block-linear address (write_region_bl in hle_nvn.c). t_nvn checks 15 regions in RGBA8, BC1 and a two-layer RGBA16F array, from 1×1 to the whole level, byte for byte against the old whole-level path.
• Scenes and characters black, only foliage, light shafts and the HUD drawn → NVN's BindProgram(program, stageBits) binds per stage, and a null program unbinds just those stages. Dread's depth prepass calls BindProgram(program, vertex | fragment), then BindProgram(null, fragment), and draws depth with the vertex shader alone. The host replaced its single bound program on every bind, so a null bind meant "no program bound", and all of the prepass's ~140 draws per frame were dropped. The depth buffer stayed empty, and the G-buffer passes, which test depth EQUAL without writing, passed no pixel. The replay now tracks the program per stage (bind_stages in nvn_mtl.m). When the stages come from different programs, or the fragment stage is unbound, it builds a composite program, which gives a vertex-only pipeline for the prepass. Mario's and Engage's binds always bind every stage from one program, so they are unchanged. The combinations Dread uses: (program, 0x3), (null, 0x2), (null, 0x7). Pass dumps now print a Z line wherever a draw is dropped (no program, no pipeline, empty scissor, both faces culled, no render pass). That is how this one was found: 144 Z … draw skipped: no program bound lines in a frame whose G-buffer pass had draws but stayed black.
• Cutscenes at 30 fps, gameplay at 60, as on the Switch: the game keeps present interval 1 throughout. In a cutscene its main loop spins in 1 ms steps until the cutscene's audio clock advances, then renders a frame; that clock moves with a 30 Hz sound thread (03, timer slack entry).
• One draw per render pass: the cutscene frames have about 180 passes with one draw each. The game updates uniforms inline (CommandBufferUpdateUniformBuffer) between draws, and every copy or update ends the host's render pass. This is correct but costs submit time (about 5 ms per frame on the Mac).

Seen in 13 Sentinels

• Single-stage programs: the game builds one NVN program per GLSLC output (838 ProgramSetShaders at startup) and combines vertex and pixel programs per draw with BindProgram stage bits, so the composite-program path from Metroid Dread does all the pairing. The real pairs exist only at draw time. 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 with shader_separable on shader_split.py pairs every vertex shader with a cover: the fewest pixel shaders whose inputs together are every input any pixel shader reads (here one shader). vs_pad then gives each vertex shader all of them, and any vertex / pixel combination links.
• dmsh_build.py's shader cache is keyed without the pairs: a vertex shader whose pairs changed is restored unpadded from cas/. After changing a profile's pairs, delete that profile's cache (build/cache/metal_shader_cache/cas and b-*) and rebuild. The NSP script always starts from its own cache, which has the same problem only when a profile's pairs change between two of its runs.
• Triangle fans and indexed quads: Metal has neither. Non-indexed fans draw through a shared (0, i, i+1) index buffer with first as the base vertex; indexed fans and quads are expanded on the CPU into a 32-bit triangle list in the frame arena at replay (nvn_mtl.m, as for 8-bit indices).
• Vertex streams 14 and 15: the host bound NVN stream n to Metal vertex buffer 16 + n, but Metal has slots 0–30 only (30 is the host's constant fallback), and stream 15 hit the assertion buffer index (31) must be < 31 a minute and a half into the prologue. Translated shaders use only buffer(0), so streams 14 and 15 now go to slots 14 and 15 (VB_SLOT).
• Packaged texture data from CPU memory: every texture is created with TextureBuilderSetPackagedTextureData; NVN copies that data into the texture storage at TextureInitialize (a compressible pool cannot be written by the CPU). The host copies it and marks the texture CPU-written, so it is de-laid-out and uploaded like a packaged texture. Linear textures (the save thumbnails) are excluded: in a CPU-visible pool they are read straight from the pool buffer.

Seen in Super Mario 3D World + Bowser's Fury

• Shader containers: bfsha (BNSH) and agl sharcb (BAHS version 9) 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.
• 1D textures: the material shaders sample 64×1 R32F lookup tables created as NVN 1D textures (and filled with WriteTexels). Translated shaders never declare texture1d (none of the five games' libraries does), so a Metal 1D texture bound to their texture2d argument is a type mismatch. NVN 1D and 1D-array textures are now created as one-row 2D / 2D-array Metal textures.
• Fixed: daylight scenes washed out to white. The game creates its deferred pre-lighting buffer (the G-buffer's fourth target, R11G11B10F) as a one-layer 2D array texture, and the lighting shader samples it as 2D. On Maxwell the sample instruction carries its own dimension and a non-array sample of an array texture reads layer 0; on Metal a 2D array bound to a texture2d argument is a type mismatch and every read returned one fixed value, so the whole picture got one pale ambient colour (bloom then blew it out further). The host now compares each texture binding's declared type (the reflection's dtt) with the texture it resolves to and binds a compatible view: 2D array ↔ 2D, 2D array → cube / cube array (tex_fit in nvn_mtl.m, cached per texture-pool entry); any other mismatch is reported once. How it was found: a blit-based pass snapshot of that target showed a correct, colourful image while a lighting shader edited to output only that input showed a uniform colour, and both resolved to the same Metal texture. Earlier suspects that turned out correct: fog, bindings, the 1D lookup tables, clear colours, the depth convention, the final sRGB view, blend states, GPU luminance readback. There is no tone-mapping pass in this game: the final composite writes scene + bloom into an RGB10A2 target.
• Targets the fragment shader does not write: volume draws after the G-buffer (50d813761ca3f1f9 and others) write only outputs 0 and 2, with blending and all channels enabled on targets 1 and 3. Maxwell leaves targets outside the shader's output map untouched; on Metal such an attachment gets undefined values blended in. Pipelines now get a zero write mask and no blending for colour outputs the shader does not write (reflection colors, color_known in nvn_refl.m). All five games' regressions were unchanged by it.
• Open: one scaled texture copy (1920×1080 → 1280×720, RGB10A2, flags ENGINE_2D) when leaving the title: Metal blits cannot scale, so it is skipped with a red line; it is probably the captured frame for the transition.
• iPhone only: solid magenta blocks in World 1-1 (fix in place; waiting for a device run). The blocks are exactly (255, 0, 255), hard-edged and aligned to the 16-pixel grid (16×16, sometimes two stacked). DM_NVN_MAGSCAN on the A20 Pro found them first in the 1920×1080 RGB10A2 composite (ser=30), in the pass whose shader 66a71a3081b91b9a renders into that texture while sampling it. They kept appearing with the sparkle particle 5e4de7561c670fee dropped, so that particle, the earlier suspect (picked by position), is not the cause. Writing and sampling one attachment within a render pass is undefined in Metal: the M4 Pro happens to return the old content, while the A20 Pro now and then returns whole 16×16 blocks as (1, 0, 1), probably blocks of the compressed render target caught mid-write. On the Switch the read returns the content from before the draw. The host now ends the pass before such a draw, copies the attachment's level and layer into a snapshot texture in the replay stream, and binds the snapshot for that draw (fb_prepare in nvn_mtl.m). It does so only for the fragment shaders the profile's feedback_snap key lists (66a71a3081b91b9a, plus 30e3ccb0bd304ead, which samples a G-buffer target of its own pass on the title screen). In 1-1 that is one 8 MB copy per frame. Doing it for every shader is too expensive: Breath of the Wild has over a thousand draws a frame that sample the G-buffer of their own pass (Mac frame time 16.6 → 36.6 ms), and Mario Wonder's submit time went from 1.4 to 4.0 ms. Those games' feedback draws (Mario Wonder a7e5c6fd501cf21f and 22a61cb1ec732aef on its RG11B10F scene; BotW 4d3d5a4408c2e8d9, 753d0169ae3ce4f0 and others on its 1600×900 G-buffer) have shown no blocks on the iPhone so far; if they do, DM_NVN_MAGSCAN names the pass and the shader goes into that profile's feedback_snap. A sampled view that spans more levels than the attachment's one is left as before with a red line: 3e8567da37eb28a1 renders level 1 of a texture while its view covers levels 0–4 (mip-chain passes).

Seen in The Legend of Zelda: Breath of the Wild

• Shader containers: 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.
• Fixed: the 3D scene black, only particles and the HUD drawn. The engine binds almost every texture with nvnCommandBufferBindTexturesDeferred (main+0xec2e64, one handle, the pointer being the handle field inside the texture-sampler object). "Deferred" means the GPU reads the handle from that address when the command executes; the game records its command blocks once, submits them every frame and writes the handles afterwards. The host copied the handles at recording, mostly zeros, so the deferred-lighting composite (one full-screen triangle per material class, blending into the HDR target) read the dummy texture 0 for all of its inputs. The bind record now keeps the source address (p_bind_handles.defer / src) and every replay-time reader refreshes the handles first (bh_fresh in nvn_mtl.m). 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.
• Fixed: the same texture attached at two colour slots. The G-buffer passes attach the albedo and another buffer twice each (c1 = c2, c3 = c4); the duplicate slots have write mask 0. On Maxwell the ROP writes only masked channels and the slots do not interfere. On Apple's TBDR each attachment has its own tile memory and the whole tile is stored at the end of the pass, so the copy stored later (loaded, never written) overwrote the other one's writes, in a different order for each 32×32 tile: the albedo held only scattered squares. A slot whose texture, level and layer repeat an earlier slot is now left unattached (rt_dup_slot), also in pipeline prediction; a draw that really writes such a slot is reported once (color slot … holds the same texture as an earlier slot). Records hold copies of texture objects (tex_snap), so slots are compared by serial, not by pointer.
• Cube textures sampled as cube arrays. 1,512 shaders sample the environment probes as texturecube_array while a single cube texture is bound. Metal can make a cube-array view only of a 2D-array base texture, so NVN cube textures are now created as 6-layer 2D arrays and bound through cube / cube-array views by tex_fit (which also maps a cube array read as a cube to its first cube).
• Fixed: a compressed texture as a render target. agl's BC4 block encoder (8e2cc2671cae6020) writes a 32×32 RGTC1 (BC4) texture through an RGBA16UI view (one texel per 8-byte block: BC4 encoded on the GPU). Metal can neither attach a compressed format nor make an uncompressed view of one, and the pipeline-descriptor validation aborts the process, so the attachment used to be left out and its draws did nothing (render target is a compressed-format texture …). Such a pass now draws into an uncompressed shadow of blocks × blocks in the view's format: the level's (or layer's, or depth plane's) blocks are copied into it through a buffer before the pass and back after it (bcrt_attach / bcrt_flush in nvn_mtl.m; the method comes from PR #11). The shadows stay resident and are reused, and they are kept out of the load-skip and store-elision bookkeeping. The encoder's output must be declared uint: the game has no PC version, so the profile's shader_facts_manual.tsv lists it (without it the pipeline fails and the red line prints the row to add), and so does its twin for a 2D source (93aa5e22a8202409), whose failure showed up in an iPhone run.
• Big-stack NVN entry points: the render thread's stack is small, and Metal's driver compiles internal programs (a fast-clear blit) on the calling thread while the host ends a render pass inside nvnQueueSubmitCommands. That overflowed the guest stack (SIGBUS in AGCDeserializedReply, address in the guest stack region). All implemented NVN functions except the nvnCommandBuffer* recording calls now run on the per-thread 512 KB host stack (03, big-stack entry; nvn_big_entry in nvn_entry.S).
• 60 fps is now the host's frame step in whole vsyncs (development.md, Breath of the Wild): a stub at the framework's setWaitVsyncInterval hands the game the host's step, and frames stay on screen exactly that long. The notes below are about the method before it, which DM_FRAME_STEP=legacy still runs: a rewriter cut on the game's frame-interval logic (03); the engine's variable frame rate (ksys::VFR) scales its time step by the present interval, so the speed is right as long as the interval matches the frames. Fixed: slow motion when the GPU is the bottleneck (GitHub issue #9). The interval controller (main+0x112660c) measured a frame as framework +0x250, the main thread's CPU span from the start of procFrame to before its draw step (main+0x1126cec → main+0x1126d20): about 2.7 ms in the shrine (probe at main+0x11266a4, @x19+0x250). Only while the object at [main+0x1e3a638] + 0x1d2b8 exists did it take max(CPU span, GPU time), in 64 of 6,046 shrine frames. So the interval stepped up only on CPU spikes: an iPhone 16 Pro Max at a 25 ms frame p50 never left interval 1 in 4,800 frames, and the game ran at about two thirds of its speed. The GPU time (framework +0x220) is the difference between the frame's two timestamp reports: TIMESTAMP_TOP at the start of the frame's command buffer (main+0xe94584) and TIMESTAMP at its end (main+0xe94664), read after the draw step has waited for the frame's fence (main+0xe949bc). It read 0–6 ticks, because the host wrote all 5 reports of the frame's one submission with one time in its completion handler. Two more cuts (main+0x112691c, main+0x1126974: cbz → b) make the controller always take the GPU path. The profile key gpu_timestamps segment makes each report take the GPU start (TIMESTAMP_TOP) or end (other types) of the sub_cut segment it was encoded in (rep_ticks in nvn_mtl.m: MTL4CommitFeedback times converted to the guest clock, written in the last segment's handler as before). The GPU time now reads the frame's GPU span, gaps where the GPU waits for the host's translation included. Measured on the Mac with two BotW runs sharing the GPU under a heavy CPU load (load average about 75, a shader rebuild running): GPU time p50 16–23 ms; the fixed run stayed at interval 2 (3 while the GPU time passed 1.9 vsyncs, 1 for under 2% of the frames) at 0.93–0.98 of real time (Σ VFR ratio / 30 per wall second, probe main+0x151d3e0 @x0+0x30: +0x30 ratio, +0x34 interval). At the same time a run with the old timestamps (DM_NVN_GPU_TS=submission) spent 10–70% of its frames at interval 1 and ran at 0.78–0.95. The same GPU time drives the game's dynamic resolution (main+0x1267070): a running average as a percentage of a 60 Hz frame scales the resolution down above 175% (29 ms) and back up below 150% (25 ms). With 0 it always sat at its maximum, and it still does whenever the game holds 30.
• Render scale (DM_RENDER_SCALE): the dynamic-resolution method (main+0x12666e4) keeps its scale in s8 (0.8–1.0 in 0.05 steps, or the forced 0.75 / 0.5 of +0x284 & 0x1800) and applies it from main+0x1266880: the views' scale (+0x28a0 / +0x28a4), from which the size function (main+0x106340c) sizes the scene buffers each frame, the shader ratio uking_dynamic_resolution_ratio and the other readers. That block runs only when the dynamic path has just recomputed the scale (+0x280 bit 4, set at main+0x1266830) or in a forced mode, else main+0x126686c skips it. The stub sets s8 at the block and, for a factor other than 1.0, takes the block every frame, so the factor holds wherever the dynamic path does not run. A multiply in the size function (Tears of the Kingdom's way) would leave the ratio and those readers at the old scale.
• Fixed: the iPhone app was killed for memory in the first shrine (GitHub issue #9, iPhone 16 Pro Max, handheld). An 8 GB iPhone gives this app about 3.4 GB (resident plus headroom to the system memory limit stayed at 3,375 MB). Resident memory grew about 0.2 MB a frame from 2.6 GB at the shrine load until the kill, while the residency set and the texture counts stayed flat. The cause was the deferred-release drain running on the guest submit thread with no autorelease pool, the Metroid Dread leak (defer_drain now has its own @autoreleasepool): on the Mac a host without that pool grew 1 MB a frame in the shrine (GPU 3.0 → 4.4 GB in 2,400 frames); with it the shrine stays at 2.7–2.9 GB resident and 1.5–1.7 GB GPU, idle or walking (vmmap at 2.8 GB: about 1.6 GB guest memory, 0.75 GB private textures, 0.17 GB IOAccelerator). Found on the way: (1) about 10 serial-table entries a frame were never deleted: the game records its transient render targets (the HUD captures, the 640×360 buffers) and finalizes them before it submits, and replay put the dead serials back into g_by_serial (276,000 after 25,000 frames). Dead serials now go to a 512-entry table (tex_put, g_ghost; the memory line's finalized textures still referenced). (2) The memory-pool tracking table had a fixed 4,096 slots and the phone reached 4,006 pools; a pool past the end would have its textures re-uploaded at every use. It now grows in blocks of 1,024 (nvn_dirty.c, up to 262,144).
• Address → memory pool lookups: BotW keeps about 1,760 memory pools alive, and mtl_addr / mtl_gpu_addr scanned them all for every buffer binding of every draw: a CPU sample of the submit thread had about 1,000 of its samples under those scans (vertex buffers in draw, uniform buffers in encode_args). They now try the last hit, then a binary search over pools sorted by start address (pool_find, rebuilt when pools are added or removed). Host submit time in the shrine went from 21 ms to 7 ms per frame (p50), the median frame from 38 ms to 21 ms.
• Submissions handed to the GPU in segments. The game does not submit frame N+1 until frame N's GPU work is done: a stack sample in the shrine had the main thread in nvnSyncWait (the fence of the frame just submitted) for 59% of the time and the render thread in nvnQueueSubmitCommands for 40%. On the Switch that submit costs nothing; here it is the host's translation (about 3,600 draws and 347 passes, 7 ms), so per frame the GPU ran 13–14 ms and then sat idle 7.7 ms (DM_NVN_TSV rows @gpu_beg_ms / @gpu_end_ms against @call_ms). The replay now commits the command buffer encoded so far every 64 draws, between two passes (sub_cut in nvn_mtl.m, called after rp_flush in OP_SET_RT; DM_NVN_CUT=<draws>, 0 = off). It cuts only with no open encoder and nothing deferred, merged or hoisted, commits the pending upload batch first, and continues on the same command allocator. The fence signal, the completion work (post ops, g_done) and the allocator's fence stay on the last segment: the queue runs command buffers in commit order, and every encoder starts with a queue barrier. GPU idle per frame went from 7.7 to 3.8 ms; most of the rest is the game's own work between its fence wait and the next submit (about 2.5 ms). The other five games' frame times were unchanged (Dread's p95 dropped from 19.6 to 16.9 ms).
• Frame rate in the Shrine of Resurrection on the Mac (M4 Pro), after the above: docked (1600×900) GPU about 13 ms, frame p50 16.7 ms (60 fps; p90 17.6 ms), was 20.8 ms before segmenting; handheld (DM_OPMODE=0, 1280×720) GPU p50 9.3 ms and a steady 60. Host translation is 8 ms per frame (22,000 instructions per draw), spread over Metal's encoder creation and draw calls, argument encoding, Objective-C retain / release and pool lookups. The per-pass GPU timer (DM_NVN_PASSGPU) is switched off under Metal 4 (pg_render returns at once), so GPU time cannot yet be attributed to passes. The frame loads 376 MB and stores 432 MB of render-target data (@rt_load_mb / @rt_store_mb; 325 of 347 passes load and store everything), which is the next thing to look at.
• Fixed: dark boxes and cream cloud rectangles behind HUD prompts, a black stamina wheel and a black System menu background (layout captures; GitHub PR #6). Panes tagged CaptureOn (about 140 layouts, e.g. PaGuideKeySmoke_00; user data A8, CaptureScale 0.5, MultiFilter) are drawn into a target, blurred at half size and copied into a capture texture that the layout's materials sample in place of the placeholder FBLayout_00. The game creates that texture with its per-format default mapping (one call site, main+0xecb0bc: R8 → (R, 0, 0, 1), RG8 → (R, G, 0, 1), and so on), then rebuilds it on the same storage with the capture's own mapping: main+0xf597dc picks (1, 1, 1, R) for an A8 capture (the material colour kept, the mask as alpha) and alpha 1 when the pane says CaptureOutputAlpha 255 (read at main+0xf58e98; an RGB10A2 capture then reads (R, G, B, 1)); main+0xecacf0 stores the mapping and rebuilds the texture (main+0xecabe8), and main+0xec33f0 replaces the registered descriptor only if nvnTextureCompare reports the two textures different. The host compared the description and the storage but not the mapping, so every rebuild was dropped and the captures kept (R, 0, 0, 1), whose alpha is always 1: hard-edged dark boxes behind every prompt (the shadow pane e6a88e12dbae9743: black, alpha = capture alpha × vertex alpha 180/255) and rectangles filled with cream cloud (the glow aa209f57b133e862: alpha = min(cloud alpha, capture alpha)). Before the cause was known, the profile key alpha_from_red listed all 26 pixel-shader variants of the layout library (main's BNSH at 0x1d83000) and bound such render targets through a view mapped (R, 0, 0, R) (tex_alpha_red in nvn_mtl.m). That fixed the alpha but zeroed green and blue, so tinted panes such as the stamina wheel came out black, and the System menu background (a capture with alpha forced to 1) stayed black. nvnTextureCompare now compares the mapping too, an unset mapping counting as (R, G, B, A), and BotW's profile no longer lists alpha_from_red. Checked on the Mac by the PR's author (M4 Pro, unchanged shader libraries): green stamina, the purple noise meter and the game's own blurred System menu background. t_nvn covers the comparison.
• BC textures are never attached or viewed raw (GitHub issue #4, reported from an iPhone 14 Pro / A16). The game encodes one 32×32 BC4 texture on the GPU, drawing its blocks through an RGBA16UI view (one texel = one 8-byte block). Metal can neither attach a compressed texture nor make an uncompressed view of one, so a pass that draws into it uses an uncompressed shadow (above), and attach, the pipeline prediction and the sampling views go by the NVN format. The A16 has no BC formats; the host used to store BC textures decompressed on such GPUs (about 3.9 GB instead of 2.8 GB for this game, over iOS's limit on a 6 GB iPhone), and that path has since been removed: only GPUs with BC are supported.

Seen in The Legend of Zelda: Link's Awakening

• Lower-left window origin. It never calls nvnDeviceSetWindowOriginMode, so it keeps NVN's default, lower-left (OpenGL), where NDC y = -1 lands on a render target's first memory row. The six games before it all switch to upper-left. Read as upper-left, every texture the game renders and samples came out flipped and cancelled out, so the scene looked right; but text it renders into textures for the title and the menus came out upside down, and it copies the frame with MIRROR_Y. Viewport and scissor rectangles, fragment coordinates, copies and texture coordinates address the same memory rows in both conventions, so the host only flips the direction inside the viewport (negative viewport height), reverses the front-face winding, and flips the window image at present and in raw screenshots (g_origin_ll in nvn_mtl.m).
• Copies a blit cannot do (copy_t2t_draw): RG11B10F → RGBA16F every frame, half-size depth-stencil copies (depth and stencil are written from a fragment shader, stencil through [[stencil]] export from an X32_Stencil8 view of the source), and a scaled MIRROR_Y copy of the frame to 1280×720 RGBA8. Each is one render pass per layer; this also covers Super Mario 3D World's scaled copy when leaving the title.
• Fixed: missing palm tree and flowers (the pipeline 147c829e79fb322f + 0ebba5ff6a44aa52, shadowmap_pbr.bfsha variant 641, failed to build). The translation declared vertex inputs 1–5 as uint4 because the shader only moves their bits (x bit-cast to an integer, y copied as a float), and the game feeds them 16-bit normalized formats, which Metal will not pair with an integer input. These are 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. The profile's shader_facts_manual.tsv declares the five inputs float (the bits Maxwell would see), which changes only this profile's chain version. 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. The guess is not wrong in general: 1,437 of the game's vertex shaders have integer-typed inputs, many of them bone indices the shader multiplies as integers, so overrides are added only for shaders seen failing with float formats (failed to create render pipeline … Cannot convert attribute …).
• Fixed: no ambient occlusion (stencil lost in the depth downsample). The occlusion pass (aa05724ffa62cb82, a 480×270 R8 target cleared to 1.0) only writes where stencil bit 1 of the quarter-size depth-stencil copy is clear. The copy left the stencil as the new texture's undefined content (new depth-stencil textures are not zero-filled), so on the Mac the mask kept 1.0 everywhere and the scene had no occlusion; on a device the result depends on whatever that memory holds. The copy now carries the stencil.
• Fixed: flickering hard-edged white tiles on the step of Marin and Tarin's doorway (Mac and iPhone). 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. The culling shader computes its tile index as workgroup y × c0[0x200] + workgroup x, reading the dispatch's workgroup count from NVN's driver segment c0. The host passed that shader an all-zero c0: 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. 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. The fix is in the host: a declared driver-segment binding now counts as read (nvn_refl.m), and 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_). After it every tile row has its lights and the doorway shows the game's soft light from outside. 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.
• Fixed: outdoors the whole screen was blurred (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. The game sets depth mode near = 0, but every vertex shader was translated for GL, so all depths sat in 0.5..1 and the computed distance put everything out of focus. Fixed by the per-mode entry points in §5. 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.
• Fixed: outdoors was dim and nothing cast a sun shadow (Mac and iPhone). The sun's shadow map (2048×2048, depth 32-bit float) was correct, but every lookup into it failed, so the whole scene was lit as if in shadow. The surface shaders compute the lookup from a world position the vertex shader passes in locn24, divided by its w; the vertex shader writes only x, y and z of it. The translation started that output at 0 instead of the hardware's (0, 0, 0, 1), so w was 0 and every compare value was infinite (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). The fix is general: every vertex shader output now starts at (0, 0, 0, 1) (see "Unwritten components" above). Compared with forcing the compare to pass (compareFunction always): the same brightness, plus shadows from the fences, the house and the trees.
• Open: SR_Y_DIRECTION. 2,487 of the game's pixel shaders read the window's Y-direction system register, which is −1 under a lower-left origin. The Ryujinx recompiler answers a constant +1 (a TODO in InstEmitMove.cs; LASRecomp, a static recompilation of this game, answers −1). In the room shader the value only scales a depth slope under abs(); other shaders may use its sign (dFdy-based effects). Answering it per profile means a translator change, i.e. every game's chain.
• Open: two pixel shaders do not translate (9d8f6b8543355374, 9338e3f1b597fe7c: variants 0 and 1 of one program in postprocess.bfsha @0x252000; SPIRV-Cross throws on a null pointer). No run so far has loaded either.
• GLSLC. 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.

Seen in Fire Emblem: Three Houses

• Two texture pools and two sampler pools with overlapping IDs: the scene binds a 12,544-entry texture pool and a 3,856-entry sampler pool, the 2D UI's command buffer binds 4,096-entry ones, and both fill IDs from 256. The replay had one ID table, so the scene's default 8×8 white texture overwrote the UI's textures: the title menu's labels drew as plain bars, other menus were empty except for the cursor. Each pool now has a base in the table, in creation order; RegisterTexture / RegisterSampler add it, and binds recorded on a command buffer add the base of the pool that command buffer last set (bind_handles, cb_pool_set in hle_nvn.c; log line pool #2 (…): numbered from N). Deferred binds and handles read from uniform buffers are not remapped (red line). The table's capacity is the sum of the pools: the scene registers each draw's textures into fresh IDs of a ring (8,729 and up a few seconds in), which used to read as empty slots and left the screen black.
• An integer vertex input fed a float-type format: int_inputs declares an input uint when its first translation only bit-casts it, and KTGL's character and soldier shaders then take it back with as_type<float>; the game binds a Half2 there (UVs). Maxwell's fetch converts by the format and the register holds the float32 bits. Metal refuses such a pipeline (Vertex attribute m_36(5) of type uint4 cannot be read using MTLAttributeFormatHalf2: the characters' shading draws were dropped, black silhouettes), and reading the raw bytes as integers gives denormals (UVs collapse onto one texel: flat tan characters, also Age of Calamity's soldiers). The replay now converts such attributes (half, unorm / snorm 8 and 16, 10:10:10:2, 11:11:10, BGRA8 into an int / uint input) on the CPU into float32 bits, four components with the fetch's (0, 0, 0, 1.0) defaults, in the frame arena, and binds them as UInt4 / Int4 on buffer slots 1–12 (vconv_bind, once per range and layout per submission; red line 89).
• Depth clamp: the game turns SetDepthClamp on; it is now recorded (OP_DEPTH_CLAMP) and replayed as Metal's depth clip mode (it used to be ignored).
• Fixed: the ground under close-up battle cameras was flat steel blue. The terrain's vertex shader (56778feb1cfdee7e) reads its grid position from an RGBA16I attribute: x and z the grid cell, y the height, which spans the whole int16 range. So does the side pass that draws the river surface into targets of its own (db725c85ca060403). int_inputs declared that input uint4, and Metal refuses a signed format for an unsigned input, so vfmt_for_ct swapped Short4 for UShort4. That keeps the 16 bits but widens them with zeros: every negative height read as 32,768–65,535, and those vertices went far up out of view. Maxwell widens by the format's signedness, whatever type the shader gives the register. Where the ground was missing nothing wrote the depth (reversed-Z, cleared to 0), and the lighting takes depth 0 for sky and paints the fog colour. The G-buffer colour targets are never cleared, so they kept an old image there: cobblestones of a mesh under the terrain, which from the map camera showed through the gaps like a road. 8- and 16-bit integer formats whose signedness differs from the input's now go through the CPU conversion like the float-type formats (vconv_int in nvn_mtl.m, red line 96 once): 32-bit values extended by the format's own sign, missing components (0, 0, 0, 1) as integers. How it was traced, since two steps misled: a HILITE colour over the gap was stale too (the never-cleared albedo), and the end-of-frame depth showed 0 there. Recomputing every draw's clip positions from CBDUMP data showed that no draw of the shader reached the gap. Then the side pass's layer, which showed only scattered tiles of the river, led to the vertex format. Ruled out on the way: depth clamp (Metal's clip and clamp modes gave the same depth), the pass without a depth attachment (142595 is a screen-space decal pass: it samples the depth and writes albedo and normals), polygon offset (none), dropped draws. Still open: 64-pixel rectangles in the albedo where decals land. The decal pass (417a523053457352) and the lighting (094a2d11253f126c, f54830a75e07e14a) read their decal and light lists from texture buffers that the translation declares texture2d, because those handles have no PC type (texture type mismatch: shader declares 2, texture is 9); skipping the decal pass removes the rectangles.
• Fixed: the first loading screen took 78 s on the iPhone (5 s on the Mac). KTGL's loader uploads each resource with a copy on the queue followed by nvnQueueFinish, hundreds of times per loading screen (effect banks, face models). mtl_finish used to call mtl_present_flush first, and with a window that waits until the display link has taken every queued present: one to two display frames per call. The windowless Mac runs present at once, so they never showed it; DM_WINDOW=1 does (one step took 20 s instead of 1.7 s). nvnQueueFinish now waits only for the GPU work submitted so far, as on the Switch, where it does not wait for the display. A swapchain texture is still not handed back before its previous present is processed (AcquireTexture waits for that ticket), and mtl_drain keeps its flush. Windowed Mac run after the fix: first load 6.1 s, frame pacing unchanged. How it was found on the device: Documents/env.txt with DM_WATCH=50 (thread table at 50 s: only the yield-polling main thread used CPU) and DM_FSLOG=<absolute path> (the reads came in bursts, with stalls of 4 / 24 / 13 s right after the effect banks).
• Run-time GLSLC: the 2D primitive, text and loading-feather shaders come from the game's own compiler at boot and are captured once on the Mac (03, SDK 6).

Seen in Super Mario Odyssey

• Layout captures as in BotW. Odyssey's layouts rebuild their captures through the same agl code (main+0x982320 replaces the registered descriptor only if nvnTextureCompare reports a difference), so the comparison fix above covers them too. The profile's alpha_from_red (all 26 layout pixel shaders, main's BNSH at 0x132c000) had turned the coin counters' solid white cards into the right shapes. With the fix it no longer fires: on the Mac smoke route to frame 20,100 it converted no texture (main's host converted its first, a 640×80 R8 target, at frame 4129), and the HUD at frames 17,300 and 17,600 is pixel-identical to main's. The list stays until the change has been seen on an iPhone.
• Fixed: the Cap Kingdom fog never parted around Mario (reported on an iPhone; fix checked on the Mac). On the Switch the ground fog opens around Mario as he walks through it. The fog's shading reads velocity fields that a small fluid simulation updates every frame: two 256×256 and two 128×128 RG16F targets, with a two-pass blur and a blended composite. Its only input is a point-sprite splat per mover, drawn with alPointSpriteShader: one vertex per mover, its size in the vertex data (Mario's main splat asks for about 18 px, others 3–5 px) and a soft disc in the fragment shader (clamp(1 - 4r², 0, 1) × velocity). The vertex shader writes gl_PointSize = min(size, c0[0x610]), and the host had nothing at that offset of the driver segment, so every splat had size 0. The simulation fields stayed exactly zero (raw dumps with DM_NVN_TEXDUMP + DM_NVN_SNAPRAW; the 8-bit dump shows velocities as black either way) and the fog stayed whole. c0 + 0x610 holds the largest point size, and the host now writes 2047 there for the graphics stages. Writing the SetPointSize value instead (always 1 in this game) left one-pixel splats whose disc weight came out 0 (the fields held -0.0). After the fix the fields hold velocities up to ±44, and the fog clears around Mario. The build's drv_reads scan in dmsh_build.py only matches fragment-shader names (fp_c0), so vertex and compute reads of the driver segment are not listed in the JSON; the host treats a declared driver-segment binding as read (nvn_refl.m).
• Fixed: the Cap Kingdom ground showed 32×32 blocks of a stale image on every fourth frame (Mac). Presents 4k were brighter and blocky, 4k+1 and 4k+2 clean, 4k+3 partly blocked, repeating to the decimal across runs. The blocks were already in the HDR target when the G-buffer pass ended, and only with the ground's material (b6384a85afd3ef94) drawn. The cause: in some frames the game first draws a depth pre-pass (alRenderDepthOnly) for objects its occlusion queries report visible (the queries rotate through three slots, hence the cycle), then draws the G-buffer with other vertex shaders and depth test less-or-equal, without clearing the colour targets. Metal compiled the two vertex shaders' identical position math to slightly different depths, the G-buffer fragments failed the test, and those pixels kept the previous content. Every Metal compile now uses -fpreserve-invariance (see "Positions must be invariant" above). Ruled out on the way: the TBDR optimizations, same-pass feedback, the present path; evaluating the game's conditional rendering on the CPU from earlier frames' results made it worse (the condition belongs to the same frame on the GPU), so the host still ignores the condition and draws.

Seen in The Legend of Zelda: Tears of the Kingdom

• Fixed: the 3D scene was black (conditional rendering). TOTK wraps nearly every ClearColor / ClearDepthStencil in SetRenderEnableConditional(RENDER_IF_NOT_EQUAL, rec) … SetRenderEnable(true). Each clear is followed by ReportCounter(TIMESTAMP_TOP, rec), the CPU sets rec to 1 when a target needs a clear, and rec + 16 is never written. The condition compares the 64-bit counters at rec and rec + 16. A timestamp report puts 0 in the counter field (the time goes in the second field), so after one clear the two are equal and later conditional clears of that target are skipped. The host ignored the condition and wrote the time into both fields, so every clear ran. The first tile-lighting round, the half-resolution scene depth the light volumes test against (back faces GREATER against a buffer just cleared to 1.0 can never pass) and the light accumulation were all cleared before the composite read them. Profile key render_cond: OP_RENDER_COND is evaluated at replay from the records in memory. A timestamp report replayed earlier in the same submission counts as 0, and a condition on a samples-passed result (an occlusion query, known only when the GPU completes) keeps drawing. A pass dump shows the switches as E lines (conditional render … → not drawn) and a skipped clear as c … skipped by conditional rendering.
• Fixed: the cave was magenta (textures filled with no NVN flush). The terrain draws sample two 121-layer ASTC 8×8 arrays (42 MB each) that TOTK creates on zeroed memory and then fills from the CPU, with no nvnMemoryPoolFlushMappedRange or FlushTexels over them (its cache maintenance runs natively). tex_flush_upload never marked them, so they stayed GPU-only with content 0, and Metal decodes an all-zero ASTC block as the error colour (1, 0, 1). A zero BC block decodes black, which looked like missing light. Profile key tex_dirty_upload: ft_add write-protects a registered texture's pages, and on first use (submit thread) a write since then makes it CPU-written. DM_NVN_TEXINFO=WxH tells the cases apart: it prints how many bytes of the texture's guest memory are non-zero at first use.
• How the tile lighting was traced: the G-buffer class byte selects per-class half-resolution passes through a depth buffer of class values (gl_FragDepth = mix(0.5, 1, nibble-swapped class / 255), then full-screen or per-tile quads with depth EQUAL), so DM_NVN_DEPTHOVR (depth test forced to ALWAYS) with HILITE tells "rejected by depth" from "never rasterized". TEXSNAP copies used to run before the pass they were meant to show (classic-queue command buffer committed ahead of the Metal 4 one); they now go through the replay's copy encoder.
• Clears inside a pass, with a partial channel mask or with a partial scissor were dropped (red lines for unimplemented mid-pass, partial-mask and partial-scissor clears). NVN clears obey the scissor and the clear's own channel or stencil write mask, and can come between draws, while a Metal load action clears a whole attachment at the start of a pass. Such clears are now drawn as one triangle over the pass in the open render encoder (clear_color_tri / clear_ds_tri in nvn_mtl.m, adapted from MacNativeGames, MIT). The pipeline has the pass's formats, the clear's mask as the write mask of its slot and no writes to the other slots. Depth is ALWAYS (and written for a depth clear), stencil REPLACE through the clear's write mask, the scissor is clamped as for draws, and visibility counting is off for the triangle. The value comes from the frame arena through the triangle's own argument tables, and the draws' tables are put back after it. A clear outside a pass opens the pass as a draw would. Clears a load action can express keep it, and two cases are still cleared whole, logged once as depth-stencil clear with a scissor or partial stencil mask, cleared whole: a depth / stencil clear with a scissor outside a pass, and a partial stencil write mask. TOTK issues the first (a depth clear scissored to 8×8 of a 64×64 target), and so does Wonder. With render_cond TOTK's scissored colour clears are all skipped; without it they are 50×50 tiles of a 500×400 target cleared to opaque black, about five a frame. Metroid Dread logged the red line once while loading on its smoke route.
• 8-bit indices written earlier in the same submission were read stale (a red line: the 8-bit index buffer had GPU-side writes in the submission). Metal has no 8-bit index type, so replay widens 8-bit indices to 16-bit on the CPU. When a copy, update or fill encoded earlier in the submission covers them (pend_hits), the CPU reads the bytes from before that write. Those draws are now widened on the GPU (wide_run): a compute pass reads the indices from a 16-byte-aligned address, with the byte shift as a constant, and writes 16-bit ones into the frame arena. As in ixc_run, the first such draw widens every such draw after it up to the next buffer write outside a render pass, so only a draw met inside an open render pass with no widening yet would split it. In TOTK about one draw per widening, up to a dozen widenings a frame on the title and load screens, no pass split (log line 8-bit index GPU expansion). Draws whose indices no write covers keep the CPU widening.
• Fixed: the stone blocks and steps of the opening cave got no torch light (Mac and iPhone). On the Switch they are grey stone lit warm on the side facing the torch; here they kept their texture and the ambient light but stayed black. Their fragment shader 54e8c08998050e79 takes its roughness from a 5120×5120 BC4 atlas (roughness = the square root of its red, packed into the upper six bits of blue in colour target 3, which the tile-light shaders read); the atlas was all zero, roughness 0 gives no point light. TOTK fills the atlas on the CPU: it fills 1024×1024 BC4 material textures from the CPU with no NVN flush, reads them back with nvnTextureReadTexels on its own threads before any draw uses them, and copies the blocks into the atlas. tex_dirty_upload only marked such a texture CPU-written when the submit thread first resolved it, so ReadTexels returned the never-uploaded host texture: zeros, and the atlas got zeros. Now mtl_read_texels reads any texture whose content is in guest memory there (a packaged layout, a CPU-written texture, or a tex_dirty_upload texture written since registration), on any thread, without an upload; and nvn_ft_poll (hle_nvn.c, at each submission after guest memory changed, at most 128 MB of textures per look, going on where it stopped) marks registered textures written after their first use, because the atlas itself is filled after it was first bound. Textures the GPU renders into or writes (mtl_tex_gpu_owned) leave that list without an upload. Found by running MacNativeGames' build of the same game on the same save: their atlas held data and their blocks were lit (their ReadTexels reads guest memory). How it was traced: DM_NVN_TEXDUMP (the blocks' roughness bits in colour target 3), the shader's MSL, DM_NVN_TEXRAW on the atlas and DM_NVN_TEXINFO (guest bytes), then lldb reading the atlas memory inside their process. Fast loop: the in-cave autosave (run/save, keys f2500:A,f2800:A,f3100:A,f3600:A, shot at frame 4500); DM_DETCLOCK cannot be used (TOTK stops at frame 1 waiting for guest time).
• Fixed: draws into a 32×32×32 BC4 3D texture were lost (render target is a compressed-format texture …). agl's BC4 block encoder (0544282c162305db, the same shader Mario Wonder ships) encodes the texture one depth plane per pass through an RGBA16UI view. It now draws through the shadow described under Breath of the Wild, with its output declared uint in shader_facts_manual.tsv. A raw dump at frame 1500 of the new-game route: the texture was all zeros before; now all 32 planes hold encoded blocks.
• Fixed: a draw with no render target. An agl fragment shader that only stores to a storage buffer draws one triangle into a 1×1 viewport with no colour or depth target set (it samples a 64×64 cube map, and the game reads the result back from the buffer). The host found nothing to attach and skipped the pass (no render target could get an MTLTexture); it now opens an attachment-less pass at the viewport's size (bare_dims in nvn_mtl.m, from PR #11). Compared at the same frames of the new-game route, the cave's rock and stone blocks now get ambient light they lacked, and Zelda's bloom halo in the descent cutscene is gone.
• Integer fragment outputs. 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. The profile's shader_facts_manual.tsv declares them: the BC4 encoder above and the Ultrahand mask (e3fb0bb58a691e10, R8Uint, found by PR #11; the Mac route does not reach it). When another one fails, the pipeline red line prints the row to add. nvnDeviceRegisterFastClearDepth, a hardware hint, now returns success instead of reaching the red stub.
• Texture layout (open): MacNativeGames found that NVN chooses the automatic block height by truncating height_in_blocks * 1.5 / 8, while nvn_tex.c rounds the division up, so some sizes (a 160×160 BC3 texture: H3 instead of H2) are read with the wrong layout. Rhythm Heaven Groove confirmed it on its two 40-block-row textures (below), and Touchable (a d2r-mobile port of the same game) on 16 of its texture files; nvn_tex.c now truncates for every game (pick_tile(ah * 3 / 16)). It changes layouts at 6–10, 17–21, 38–42 and 81–85 block rows, which matters only for data the game lays out itself (no packaged layout: Tears of the Kingdom's flushed textures, Monster Hunter Rise's sparse ones); those were not re-checked.

7. Diagnostic switches (host)

Switch	Output
DM_NVN_SHOT=f1,f2:path_%d.png	Window screenshots at those presented frames
DM_NVN_PASSDUMP=f or f1-f2 (+ DM_NVN_PASSDUMP_ALL=1)	Every render pass of those frames: P attachments (with mtl= the Metal texture behind them), C clear colours, K shader keys, B each binding's actual source (also with mtl=), d/D draw state (b0–b3 the blend states of targets 0–3; on d, po= the polygon offset enables / factor, units, clamp and dc= depth clamp), Y inline updates, X texture clears, T texture copies, N passes skipped because no attachment could be attached. Without _ALL the K / B / R lines cover only the first draw of each pass
DM_NVN_TEXDUMP=f:path_%d.png (+ DM_NVN_TEXSNAP=i,…)	Every colour render target of frame f, as PNG (%d = serial × 100 + layer × 10 + level; 2D arrays by the attached layer). TEXSNAP copies the attachments of in-frame pass i (2D, or slice 0 of a 2D array) in a separate command buffer, which does not run at the pass's position in the GPU timeline: a snapshot taken right after a clearing pass showed a complete image, so read it as "what this target holds around the frame", not "right after pass i"
DM_NVN_TEXDUMP=size=WxH:path_%d.png	The same dump for the first frame that opens a render pass whose first colour attachment level is W×H: for one-off work (a mask built when a prompt appears) whose frame differs between runs. With DM_NVN_SNAPRAW=1 each image also gets a .raw file with the alpha channel
DM_NVN_CBDUMP=dir (with PASSDUMP)	Raw uniform buffer contents per binding
DM_NVN_HILITE=fskey,…	Replace those fragment shaders with solid colours: tells "geometry is wrong" from "shading is wrong"
DM_NVN_DEPTHOVR=fskey,…	Those fragment shaders' draws compare depth with ALWAYS: with HILITE, tells "the depth test rejects every fragment" from "nothing rasterizes". With PASSDUMP, the D line shows the depth range (dr=) and logic op (lo=), and c lines each recorded clear (slot, serial, mask, scissor; depth / stencil values)
DM_NVN_SHDUMP=dir	Control + code of programs missing from the shader map
DM_NVN_ALLOCPROF=1	Metal allocation growth per replayed op type and per function, command allocator sizes, new textures
DM_NVN_CBDUMP_KEY=vskey (with CBDUMP and PASSDUMP_ALL)	Only that vertex shader's draws: uniform buffers plus vertex streams and index buffer (F<frame>_p<pass>_vb<stream>_<addr>.bin, _ib<type>_<addr>.bin)
DM_NVN_TEXINFO=WxH	Storage and description (pool, offset, flags, packaged or not, storage size and how many of its bytes in guest memory are non-zero, i.e. written by the CPU) of every texture of that size on first use, and each pool registration of one (TEXINFO register: builder channel mapping, view and view mapping)
DM_NVN_TEXRAW=info (with TEXDUMP)	Raw bytes of the non-render-target textures TEXINFO picked out, at the TEXDUMP frame
DM_NVN_TEXRAW=ser,… (with TEXDUMP)	Raw bytes of those textures (by serial, as PASSDUMP prints them) at the TEXDUMP frame; a Depth32Float_Stencil8 texture gives its 32-bit float depth plane (a shadow map)
DM_NVN_T2T_TRACE=WxH:file (+ DM_NVN_T2T_FROM=f)	After each texture copy into an RGBA16F texture of that size, a GPU-timeline copy of the result into a trace buffer; 240 entries written to file (header line, 240 frame numbers, then the texels)
DM_TEST_PRESENT_LAG=n,ms	Sleep ms before every nth present: a slow frame on demand (Mario then switches to present interval 2)
DM_NVN_BLIT_ORDER=0	No barriers between dependent operations in the copy encoder (reproduces the fountain bug)
DM_NVN_KEYDUMP=<fs key>[:n[:every]]	For up to n draws of that fragment shader (default 6; at most one every every frames), the PASSDUMP binding lines plus [kd] lines: vertex attribute formats, each vertex stream's data (all instances of per-instance streams) and every value of both stages' uniform buffers. Meant for devices: no frame number to guess; compare with the same switch on the Mac
DM_NVN_MAGSCAN=<fs key>[:every[:phase]]	Device self-check for solid colour blocks. Every every frames (default 30) a compute kernel reads each render pass's colour attachments as the pass ends and logs [mag] lines for 16×16 blocks that hold one magenta value throughout, or non-finite pixels, naming the pass relative to the one holding that shader. Every phase frames (default 600) it rotates: 0 draws as usual, 1 drops that shader's draws, 2 / 3 keep its geometry with a constant grey / bright magenta fragment output. Each phase ends with a summary line, including the magenta pixel count in that shader's own pass
DM_NVN_FB_SNAP=0 / all / <keys>	Which fragment shaders get the snapshot for same-pass feedback sampling (a draw sampling a colour attachment of its own render pass reads a copy taken just before it): none, every shader (slow, see BotW above), or a comma-separated key list instead of the profile's feedback_snap
DM_NVN_ALPHA_RED=0 / <keys>	Which fragment shaders read single-channel render targets mapped (R, 0, 0, 1) with alpha = red: none, or a comma-separated key list instead of the profile's alpha_from_red (Odyssey's layout shaders; BotW no longer needs it since the nvnTextureCompare fix, see above)
DM_NVN_CUT=<draws>	Commit the replay's command buffer every that many draws, between passes (default 64; 0 = one command buffer per submission). Growing segments (64, 128, 256 … draws) were tried for fewer commits: Breath of the Wild fell from 60 to 30 fps in the same scene (GPU idle while the long last segments were encoded), and the commits cost little once the residency set stopped churning
DM_NVN_UB_DEFER=0	End the render pass at every inline uniform update (nvnCommandBufferUpdateUniformBuffer), as before. By default an update inside an open pass is deferred to the pass end; the draws after it read a shadow made of the deferred updates when they cover the whole bound range, and the pass still ends first for a draw that needs more (partial cover, the data read as vertices, indices, indirect arguments or a storage buffer). The residency set line counts deferred updates and such early ends
DM_NVN_BARRIER=game / full	How draw passes are ordered. full (default): every pass waits for all earlier GPU work. game: a pass waits only for the previous pass's fragment stage, so its vertex work overlaps the previous pass's fragments, unless the game called nvnCommandBufferBarrier in between or a compute / copy / upload ran (those keep full barriers). Only vertex-stage reads of the previous pass's output without a game barrier change order
DM_NVN_GPU_TS=segment / submission	What timestamp reports record, instead of the profile's gpu_timestamps: the GPU start / end of the segment holding each report, or one time for the whole submission
MTL_DEBUG_LAYER=1 MTL_DEBUG_LAYER_ERROR_MODE=nslog	Metal validation without aborting

A first-frame bring-up that worked for Mario: screenshot is black → TEXDUMP shows which targets have content → PASSDUMP the pass that writes the window → HILITE its fragment shader → CBDUMP its uniforms → read the translated .metal of its vertex shader.

Seen in Hyrule Warriors: Age of Calamity

• Compute has a binding of its own. BindProgram(program, stageBits) with the compute bit used to replace the one bound program, so a draw after a dispatch found the compute program (or none: no program bound before the draw, program has no vertex shader) and was dropped, and a dispatch after a graphics bind found no compute shader (the dispatched program has no compute shader). The replay now keeps the compute program apart (S.cprog in nvn_mtl.m, and g_pp_cprog in the pass analysis). The game interleaves dispatches and draws; its intro movie's NV12 → RGB draw was dropped and the movie played black under its "+ Skip" prompt.
• Images registered on their own: nvnTexturePoolRegisterImage writes the same table as RegisterTexture (image handles are pool indices).
• Indirect draws and dispatches: nvnCommandBufferDrawArraysIndirect's record {count, instances, first, base instance} has the layout of MTLDrawPrimitivesIndirectArguments (and DrawElementsIndirect's {count, instances, first index, base vertex, base instance} that of MTLDrawIndexedPrimitivesIndirectArguments), so the replay issues Metal indirect draws from the pool buffer: the GPU writes the record earlier in the frame, so the CPU cannot read it at replay. The game's GPU-driven effects are indirect quads, which Metal has no primitive for. ixc_run in nvn_mtl.m converts such records (quads, fans, and indexed ones with 8-bit indices) with a compute pass into indexed indirect arguments in the frame arena, and the draw indexes the shared quad / fan index buffers (indexed ones get a 32-bit list expanded on the GPU). A compute pass cannot run inside a render pass. So when the first such draw is met, it converts every such draw up to the next buffer write outside a render pass (copy, update, fill, texture → buffer, dispatch; they bump g_wr_epoch). Each record is then read after its last writer, and only a draw met inside an already open render pass splits it. DispatchComputeIndirect → dispatchThreadgroupsWithIndirectBuffer; when the shader reads the driver segment, the GPU copies the counts into c0 + 0x200 first. DM_NVN_IXCLOG=<n> logs the first n conversions (record → arguments) once the GPU has run them.
• Movie frames through flushed CPU writes: the decoder's NV12 output is copied by the game into a 1280×720 R8 and a 640×360 RG8 texture in a CPU-visible pool and published with nvnMemoryPoolFlushMappedRange; the profile turns on tex_flush_upload. Buffer textures (a 1,048,576-texel R32F one) stay out of that path: they are views onto the pool buffer, and an upload would have sent them past Metal's 2D size limit.

Seen in Monster Hunter Rise

• Subgroup operations in vertex shaders (about 4 % of them): SPIRV-Cross threw an exception: Subgroup builtins are not available in this type of function. Metal has no SIMD-group functions in vertex functions, and SPIRV-Cross's own subgroup emulation covers only Elect. RE Engine scalarizes uniform loads: ballot the active lanes, shuffle the first lane's value to all and take a fast path when all agree. With a one-lane subgroup that is exact, so dmsh_build.py rewrites the vertex shaders' SPIR-V that way before SPIRV-Cross (spv_single_lane: the ballot is the lane's own bit, shuffles and broadcasts return its own value, all / any are the predicate, all-equal and elect are true, lane ID 0, size 1).
• Geometry shaders (53) do not translate: SPIRV-Cross writes unknown main0_out main0() for the geometry stage, and Metal has none.
• Indirect-count multi-draws (nvnCommandBufferMultiDrawElementsIndirectCount / MultiDrawArraysIndirectCount): the scene is GPU-driven. Compute culling writes the draw records (stride 20) and a u32 count, then up to max draws (1–49 per call) are made from them. Both were empty, and the G-buffer pass had no draws: the title screen showed only the UI. Metal has no draw count read by the GPU, so each call is recorded as max ordinary indirect draws that carry the count's address and their index (capped at 1,024). The replay reads the count once the GPU has written it: at the first such draw after a GPU write it commits the work so far and waits for it (mdi_wait in nvn_mtl.m), then skips draws past the count. That costs one wait per pass that uses it; hoisting the count into a GPU mask (as ixc_run does for index ranges) would remove it.
• Textures in the sparse virtual pool hold what the CPU wrote into physical pages before mapping them: the streamed textures (flag 0x20, packaged layout) and plain sampled ones with no packaged layout. The replay took the plain ones for GPU textures, so their Metal copies stayed 0: the hunter's innerwear (ASTC 6x6) came out magenta, the ASTC error colour, and the face black. A sampled texture in a sparse pool (not a render target, image, linear or buffer texture) is now uploaded like a packaged one, through the pool's page table (03, sparse pools).
• Sparse textures are sampled only where the game has mapped them: the translator drops TEX's LOD clamp and its residency predicate, so streamed textures were sampled at levels the game had not mapped yet, which read 0 (magenta on the title screen's clothes). A sparse texture is bound as a view starting at its first level whose pages, and those of every smaller level, are all mapped (sparse_view in nvn_mtl.m). That samples exactly as a minimum LOD would; the view follows each upload.
• Bindless handles in uniform buffers written earlier in the same submission: the replay reads handles on the CPU, and a write queued earlier in the submission has not landed yet. Inline uniform updates keep where their data sits, so the handle is read from the newest one covering it (pend_read). The game queues more than the old limit of 256 writes per submission, after which every address counted as written; pending writes are now indexed by 4 KB page (8,192 per submission).
• Clear of a texture made without render-target use (a 32x80 packaged one): a whole-level clear upgrades the texture to render-target use first, as format views are upgraded (tex_upgrade).
• Storage buffer addresses in the game's own uniform buffer (about 30 compute shaders: c3 + 0x38): Ryujinx reports such a buffer with cb_slot 3, which had no NVN slot. 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.
• Bindless handles through a loop phi (218 pixel shaders): 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 (fp_t2d_cb0_40000000, handle type Direct), which the replay cannot bind (no resolved slot number, the shader does not load). 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). Those accesses become ordinary type-3 handles.

Seen in Super Mario Galaxy (Vulkan)

The game renders through Vulkan (dm/gfx/vk/hle_vk.m, written for Hades) and needed these additions; each was found from a red line or a screenshot on the route to the Star Festival:

• Shaders generated at run time: Hagi emits SPIR-V for each GX configuration and creates every pipeline of its pipeline database at boot (3,089 modules, 4,535 pipelines). DM_VK_CAPTURE=<dir> writes the modules; vk_shaders.py reads capture folders like a RomFS. Its SPIR-V puts ArrayStride on the struct member (OpMemberDecorate) instead of the array type, which NVIDIA's driver accepts and SPIRV-Cross refuses ("Struct member does not have ArrayStride set"): fix_strides copies it onto the type before translating.
• Push constants: SPIRV-Cross binds the block at buffer(0), where the first renumbered uniform buffer goes; the translator moves it to the next free buffer slot (p:0.0=<slot> in the table) and the replay sets the recorded bytes there (setVertexBytes / setFragmentBytes, the pipeline layout's range).
• Descriptor arrays: the vertex shaders read the Wii's transform memory as arrays of uniform buffers (4 position-matrix blocks, 2 normal, 2 dual-texture) and the fragment shaders 8 combined image samplers; SPIRV-Cross gives an array n consecutive slots, so the table says slot*n and the replay binds element k at slot + k.
• Dual-source blending from location 1: the GX shaders write the second blend source as an ordinary output at location 1 (no Index decoration), and pipelines blend with SRC1 factors; NVIDIA's hardware feeds MRT 1 to source 1. Metal refuses the pipeline ("Fragment shader does not write to render target color(0), index(1)"). vk_shaders.py builds a variant vk_<key>_d with that output at [[color(0), index(1)]], which the host takes for pipelines whose blend reads source 1.
• GX vertex formats: positions and colours arrive as integers read as unnormalized floats (R16G16B16_SSCALED, R8G8B8A8_USCALED, …) and in three-component formats; Metal's integer vertex formats convert to the shader's float type. A buffer layout no mapped attribute uses is left out: Metal's validation aborts on it.
• Viewport and scissor persist across render passes within a command buffer, as Vulkan says; the layer used to reset them to the render target at every pass. In handheld mode the game sets a 1280x720 viewport once and draws through several passes into its 1920x1080 swapchain, so its menus came out 1.5x too big (the file-creation dialog cut off at the right edge).
• Copies inside a render pass: the game copies from its EFB with vkCmdCopyImage between vkCmdBeginRenderPass and vkCmdEndRenderPass, which NVIDIA's driver runs. Render passes now start with deferred store actions; a copy stores every attachment and ends the encoder, the next draw resumes the pass with load actions, and the pass's own store actions are set when it really ends. Before, the rest of the pass was dropped (draw outside a render pass).
• Swapchain: the game asks for preTransform INHERIT and asserts (a stripped trap, main+0x4d26e0) when the surface does not list it. In handheld mode it keeps a 1920x1080 swapchain, draws into the top-left 1280x720 and calls vi::SetLayerCrop every frame; the host keeps that crop (hle_vi.c) and the Vulkan present shows only that rectangle.
• Linear images: staging images (VK_IMAGE_TILING_LINEAR, filled by the CPU and copied with vkCmdCopyImage) and one sampled 4x4 image alias their memory as on the Switch: the texture is made over the bound MTLBuffer with Metal's linear row alignment, which vkGetImageSubresourceLayout reports.
• Queries: two timestamp pools and a 3,072-entry occlusion pool. Timestamps are their submit's completion time; occlusion queries count into a per-submit visibility buffer, copied to the pool when the submit completes.
• Not translated: two geometry shaders (GX points and lines expanded to quads); those draws come out as one-pixel points and lines. GX's sampler LOD bias has no Metal sampler equivalent (red line).

Seen in Rhythm Heaven Groove

Nintendo's own 2D engine, SDK 21.4.0: one BNSH (shader/shader_NX.bnsh, 2 vertex and 33 pixel shaders), Mario's binding layout (every binding decodes), no packaged textures.

• Fixed: black screen, every sprite and glyph in place under HILITE. The sprite shader (40adb3021c79b155) computes texture × colour multiply + colour add and discards a fragment whose alpha is 0. The uniform buffer was right (a CBDUMP: multiply 1, add 0), so every texture read 0. The game reads each texture file into memory and makes a CPU_NO_ACCESS | GPU_CACHED pool (flags 0x21) on it: 245 pools by the first screens, each holding one texture at offset 0, 17 block-linear BC7 (flag IMAGE, no packaged layout) and 228 linear R8 (IMAGE | LINEAR). NVN keeps what a pool's storage held at nvnMemoryPoolInitialize as its content, but the host uploaded from guest memory only packaged, WriteTexels or flushed textures, and made a linear texture on its pool's buffer only in a pool the CPU can see. Now a new texture in a CPU_NO_ACCESS pool whose storage holds data starts with that content, uploaded once (tex_init_fill in nvn_mtl.m); it stays a GPU texture that can be rendered into, with no re-upload on page writes. A linear one is made on the pool's buffer in any pool (tex_linear). Left out: storage of zeros, COMPRESSIBLE textures (render targets here, in their own 0xa1 pools; their memory means nothing without the compression tags), multisample, depth / stencil and sparse ones. Found by a census of every texture's pool flags, flags, format and whether its storage held data at creation (a temporary line under DM_NVN_LOG).
• Fixed: the latency calibration's beat numbers torn and its background in blue steps. Two BC7 textures without a packaged layout were read with the wrong block height. A dump of the guest bytes deswizzled in Python with 1 to 16 GOBs and decoded with texture2ddecoder shows the right one by eye; a smoothness score does not, because a block height too tall reads padding and looks smooth. Truncating height in blocks × 1.5 / 8, as MacNativeGames found (above), matches every size checked by eye: 672×160 and 1232×160 (40 block rows: 4 GOBs, where rounding up gives 8), and 112×112 (4), 288×208 (8) and 880×432 (16), where both roundings agree and floor(rows / 8), the other rule that fits 40 rows, does not. nvn_tex.c now uses it for every game (see the Tears of the Kingdom entry).