Port Mario + Rabbids Sparks of Hope

AGENTS.md

@@ -1,6 +1,6 @@
 # Notes for agents
 
-- This repository runs Nintendo Switch games natively on Apple platforms: the game's own code runs, rewritten only where Apple's kernel requires it (`dm/tools/patch_nso.py`), on a host that answers Horizon's SVCs and the SDK's services (`dm/host`) and implements NVN, Vulkan and NVIDIA's desktop OpenGL on Metal (`dm/gfx/nvn`, `dm/gfx/vk`, `dm/gfx/gl`). Seventy-four games are playable, each tested on the Mac or the iPhone (a Mac gameplay test is enough to count); the [games table](dm/games/README.md#games) lists each with its build, SDK, engine and graphics API, and links its note, `dm/games/<id>/README.md`: what its profile needs and why, what was found on it, how to drive it on the Mac, what is open. [README.md](README.md) is the player-facing page (games, packing, playing); [docs/development.md](docs/development.md) has the internals, the script's options, building by hand, prewarm lists and power.
+- This repository runs Nintendo Switch games natively on Apple platforms: the game's own code runs, rewritten only where Apple's kernel requires it (`dm/tools/patch_nso.py`), on a host that answers Horizon's SVCs and the SDK's services (`dm/host`) and implements NVN, Vulkan and NVIDIA's desktop OpenGL on Metal (`dm/gfx/nvn`, `dm/gfx/vk`, `dm/gfx/gl`). Seventy-five games are playable, each tested on the Mac or the iPhone (a Mac gameplay test is enough to count); the [games table](dm/games/README.md#games) lists each with its build, SDK, engine and graphics API, and links its note, `dm/games/<id>/README.md`: what its profile needs and why, what was found on it, how to drive it on the Mac, what is open. [README.md](README.md) is the player-facing page (games, packing, playing); [docs/development.md](docs/development.md) has the internals, the script's options, building by hand, prewarm lists and power.
 - Where to start:
 
   | Task | Read first |

README.md

@@ -152,6 +152,9 @@ Under each game: the IPA's size, and "60 fps with 1.2.0" where the game runs at
 <td align="center" valign="top" width="33%"><img src="docs/images/zelda-echoes-of-wisdom.jpg" width="256" alt="The Legend of Zelda: Echoes of Wisdom"><br>The Legend of Zelda: Echoes of Wisdom<br><sub>~7.8 GB</sub></td>
 <td align="center" valign="top" width="33%"><img src="docs/images/donkey-kong-country-returns-hd.png" width="256" alt="Donkey Kong Country Returns HD"><br>Donkey Kong Country Returns HD<br><sub>~11.6 GB</sub></td>
 </tr>
+<tr>
+<td align="center" valign="top" width="33%"><img src="docs/images/mario-rabbids-sparks-of-hope.jpg" width="256" alt="Mario + Rabbids Sparks of Hope"><br>Mario + Rabbids Sparks of Hope<br><sub>~9.8 GB</sub></td>
+</tr>
 </table>
 
 ## Pack an IPA

dm/games/0100317013770000/README.md

@@ -0,0 +1,45 @@
+# Mario + Rabbids Sparks of Hope (`0100317013770000`)
+
+Update 1.6.2225577 (v393216) over the base game, SDK 13.3.5, Ubisoft's Snowdrop engine on NVN, with Havok and Wwise. The ExeFS loads `main` and `sdk`; the seven supplied DLC NSPs are indices 1–7. **Mac: playable opening exploration, including a saved campaign reload. iOS remains unvalidated.**
+
+## Memory
+
+- The NPDM declares 16 MB of system resource. `nninitStartup` (`main+0x4d1300`) sets a 10 MB heap; the engine takes its memory through address regions and commits pages on demand.
+- **Resolved: startup reservation failure.** [Separate virtual ranges in Snowdrop's page allocator](../../../docs/reference/memory.md#separate-virtual-ranges-in-snowdrops-page-allocator). The constructor (`main+0x1642b8`) originally gives each pool 4 GB, alongside a 1.1875 MB metadata region. `patches.py` reserves 3 GB with 2 GB for large allocations, 512 MB for 512 KB pages and 512 MB for 64 KB pages. Both uses of the combined size change together (`+0x1642fc`, `+0x164308`); `+0x1642c8` initializes the slab-size fields, and `+0x164300` / `+0x16430c` override the large range. SDK build `43211771ed39e7fcef03d88e27091831` also needs its node area cut to 62 MB (`+0x37c140`, `+0x37c1dc`) and large regions aligned to 32 MB (`+0x37c33c`). With the node cut alone, the 3 GB reservation still fails.
+- **Resolved: large allocations crossed the class boundary.** The large allocator's initialization (`main+0x168310`) separately takes a maximum 64 KB page count in `w5`, stored at object `+8`. `patches.py` changes it from 65536 to 32768 (`+0x164338`) to match the 2 GB range. The initial outer-boundary cut alone left this child allocator at its original 4 GB capacity: while loading the first campaign it returned a 12 MB compressed-file buffer exactly at the 512 KB pool's base. `free` routed that pointer to slab handling and interpreted compressed payload as its owner (`main+0x1668dc`). The corrected rewrite loads and reloads the opening campaign. The fixed 1 MB large-allocation metadata reserve (`+0x168308`) still holds more nodes than the reduced capacity needs.
+- **Current measurements; later capacity open.** In the reloaded opening quest, 64 KB pages held 291.9 MB active / 296.3 MB of bump space, 512 KB pages held 156 MB / 213.5 MB, and large allocations held 1007.4 MB within a 1280 MB managed span (2048 MB capacity). The shortage flag was zero. The slab allocators (`main+0x164aa8`, `+0x1649a8`) set the game's shortage flag (`+0x169f20`) at their derived ends and retain their original 4 GB reserve fallback. Later worlds and sustained battle play still need capacity testing.
+
+## Files and shaders
+
+- The merged RomFS holds 352 files under `moria/sdf/nx`: Snowdrop `.sdfdata` slices, a `.sdftoc`, a `.sdfver` and a `.mapping`. DLC can supply further slices and shaders.
+- **Resolved: no shaders found by the split.** The shaders are compressed records in `BERG` archives: [Shader records in Snowdrop SDF data slices](../../../docs/reference/shaders.md#shader-records-in-snowdrop-sdf-data-slices). The split reads every supplied RomFS, including DLC.
+- The registered shader union contains 23,280 stages: 20,400 vertex, 2,599 fragment and 281 compute. Base plus update contributes 18,459; DLC indices 5, 6 and 7 add 632, 2,193 and 1,996 respectively. Indices 1–4 add no unique shaders. Every stage compiles for the Mac, including the pulling and shadow-input variants. A different installed DLC set gets shader tables derived from that set while retaining the registered module fixes.
+
+## Rendering
+
+- **Resolved: compute image writes discarded.** [Image bindings supply shader-write intent](../../../docs/reference/nvn-objects.md#image-bindings-supply-shader-write-intent). The game registers its depth-conversion output through `RegisterTexture` with only `COMPRESSIBLE`; the writable image binding now upgrades its backing before dispatch. The final Mac run uses the real image and renders the lit castle grounds.
+- **Current: inactive menu placeholder warnings.** Fragment shaders `184b0b673101a2da`, `4143b4dfc3115f89`, `a5c80eb8e6380096`, `0ca739a46129f292` and `ad34a496c970dff1` declare a 3D texture in slot 3, but some menu draws bind the game's 2D placeholder (texture ID 262). All ten sampling sites are guarded by fragment CB0 float `+0x180` being nonzero or NaN. A capture spanning frames 878–909 found 148 placeholder draws, each with this float exactly zero, and 42 draws with a real 3D texture, each with 150.0. No overlapping deferred uniform updates occurred. The type warning describes an unused binding on this captured route; no shader or binding workaround was added.
+
+## Mac testing
+
+Tested on 2026-10-27 with the registered update and all seven supplied DLC packages. Title and campaign menus, creating File 1, the opening Peach dialogue, walking around the castle grounds, camera movement and collecting coins work. Reloading File 1 returns to the active “An \"Overall\" Nice Day” quest. The final validation uses all 23,280 stamped Mac shader libraries and their generated uniform tables; the earlier image-write failure is gone and no shader libraries are missing on this route.
+
+To drive the Mac test, set `DM_GAME=0100317013770000`, use the rewritten `patched_v393216` modules, the merged update RomFS, the shader output's Mac directory and `DM_AOC_DIR=<build>/aoc`. From a clean boot, `DM_KEYS=40:A,44:A,48:A,52:A` passes the title, campaign menu, File 1 and Continue Journey. Advance the first dialogue with A when starting a new file; use the left stick to walk and the right stick for the camera. Keep `DM_SAVE` in the NSP work folder. Give each boot a fresh `DM_CTL` file so old inputs and diagnostic settings do not replay.
+
+Game files, logs and analysis stay beside the NSPs, in `0100317013770000_build` and `0100317013770000-work`.
+
+## iOS build
+
+The NSP script produced `MarioRabbidsSparksOfHope.ipa` (9.8 GB) on 2026-10-27, containing the registered update, all seven DLC packages and 23,280 iOS shader libraries. The app is unsigned, for the user to sideload. It has not been installed or tested on an iOS device.
+
+## Found on this game
+
+- [Image bindings supply shader-write intent](../../../docs/reference/nvn-objects.md#image-bindings-supply-shader-write-intent): compute writes to render textures used as images, even when their builder omitted `IMAGE`.
+- [oe calls answered without an applet proxy](../../../docs/reference/services.md#oe-calls-answered-without-an-applet-proxy): `EnableAlbumImageTakenNotification` reaches the am proxy the host does not create. The host now handles it alongside the disable call.
+- [Shader records in Snowdrop SDF data slices](../../../docs/reference/shaders.md#shader-records-in-snowdrop-sdf-data-slices).
+- [Separate virtual ranges in Snowdrop's page allocator](../../../docs/reference/memory.md#separate-virtual-ranges-in-snowdrops-page-allocator).
+
+## Open
+
+- Play into a battle, check later rendering and measure memory over sustained play; the initial exploration test does not establish whole-game compatibility.
+- iOS testing is performed by the user.

dm/games/0100317013770000/bf_want.tsv

@@ -0,0 +1 @@
+# generated

dm/games/0100317013770000/game.tsv

@@ -0,0 +1,39 @@
+# Mario + Rabbids Sparks of Hope, update 1.6.2225577 (v393216), Snowdrop on NVN, SDK 13.3.5.
+# Started from the NSP script's provisional profile; see README.md for tested progress and the platform-limit allocator cuts.
+
+# ---- Identity ----
+title_id	0100317013770000
+program_id	0100317013770000
+
+# ---- App ----
+name	Mario + Rabbids® Sparks of Hope
+ipa	MarioRabbidsSparksOfHope.ipa
+app_category	public.app-category.games
+
+# ---- Versions ----
+build	v393216	1.6.2225577	main=5b76a43b231e0640521ac66203f7c178	sdk=43211771ed39e7fcef03d88e27091831
+package	v393216
+shader_versions	v393216
+
+# ---- Files ----
+# every module main imports from
+modules	main sdk
+romfs_marker	moria/sdf/nx/data/sdf_A_0000.sdfdata
+zip_store	.sdfdata:.sdftoc:.metallib
+
+# ---- Memory ----
+# 5 GB guest layout: Snowdrop commits pages within a 3 GB address region (patches.py), beside the SDK's 62 MB node area (sdk_patches.py) and a 10 MB C-library heap.
+mem_reserve	0x140000000
+mem_alias	0x118000000
+mem_stack	0x10000000
+app_pool	0xCD500000
+# the NPDM declares 16 MB of system resource
+sys_resource	0x1000000
+
+# ---- Shader split ----
+shader_groups	/moria/=moria
+shader_note	Mario + Rabbids® Sparks of Hope shaders (provisional profile)
+
+# ---- Input ----
+# Movement and camera use both sticks.
+pad_layout	full

dm/games/0100317013770000/patches.py

@@ -0,0 +1,27 @@
+# Mario + Rabbids Sparks of Hope 1.6.2225577: Snowdrop's page allocator reserves
+# three adjacent 4 GB address ranges (large allocations, 512 KB pages, 64 KB
+# pages). Only committed pages consume the application pool, but the 12 GB
+# reservation cannot fit an iPhone's guest block. The failed reservation is
+# unchecked and the first 64 KB page writes to NULL + 0x20000ffff.
+# These platform-limit cuts reserve 3 GB: 2 GB for large allocations, 512 MB for
+# 512 KB pages and 512 MB for 64 KB pages. The large allocator has a separate
+# maximum page count (w5 at +0x164338, stored at +8 by +0x168310), which must
+# shrink with its range: otherwise its original 65536-page capacity allocates
+# across the class boundary, and free interprets compressed file data as a
+# slab-owner pointer (+0x1668dc while loading the first campaign).
+# Initialize's metadata reserve is a fixed 1 MB (+0x168308), sufficient for the
+# original 4 GB / 64 KB * 16-byte nodes and thus for the reduced 2 GB range.
+# The two incoming arguments repurposed at +0x164300 / +0x16430c are discarded
+# by +0x122ff0, which overwrites x1 with the size in x0 and never reads x2.
+# Per-pool high-water validation remains open until gameplay passes; the slab
+# allocators keep the game's shortage reporting and original reserve fallback.
+CUTS = {
+    '5b76a43b231e0640521ac66203f7c178': [
+        (0x1642c8, '2800c0d2', '0800a4d2', 'Snowdrop allocator: initialize the three range-size fields to 512 MB (mov x8,#0x20000000); the large range is overridden to 2 GB below'),
+        (0x1642fc, '7600c0d2', '1600b8d2', 'Snowdrop allocator: combined virtual range end uses 3 GB instead of 12 GB (mov x22,#0xc0000000)'),
+        (0x164300, 'e1031f2a', '0100b0d2', 'Snowdrop allocator: prepare 2 GB for the large range (mov x1,#0x80000000); AllocateAddressRegion wrapper discards incoming x1'),
+        (0x164308, '6000c0d2', '0000b8d2', 'Snowdrop allocator: reserve 3 GB instead of 12 GB (mov x0,#0xc0000000)'),
+        (0x16430c, 'e2031faa', '615600f9', 'Snowdrop allocator: override the large range-size field to 2 GB (str x1,[x19,#0xa8]); AllocateAddressRegion wrapper never reads x2'),
+        (0x164338, '2500a052', '05009052', 'Snowdrop large allocator: maximum 64 KB page count 65536 -> 32768, matching its 2 GB range (mov w5,#0x8000)'),
+    ],
+}

dm/games/0100317013770000/pso_manifest.tsv

@@ -0,0 +1,2 @@
+# Prewarm seed list (shipped in the app, next to shader_map.tsv). Rows: S<TAB>shader key<TAB>stage / R<TAB>render pipeline key (pso_skey hex) / C<TAB>compute shader key.
+# Provisional profile: empty until recorded on the Mac (DM_PSO_MANIFEST=off DM_PSO_RECORD).

dm/games/0100317013770000/pso_pairs.tsv

@@ -0,0 +1 @@
+# kind: evidence  vertex / pixel shader pairs of agl program variations (BAHS) that shader_pairs.tsv does not hold, for vertex output padding and the pipeline prebuild (generated by dm/gfx/tools/shader_split.py, do not edit): vertex key <TAB> pixel key

dm/games/0100317013770000/shader_facts.tsv

@@ -0,0 +1,282 @@
+# Answer table for translation queries (generated by dm/gfx/tools/shader_split.py, do not edit): key <TAB> fact JSON. This game has no PC version to pair with: only the thread group / shared memory of compute shaders (taken from the control section); other queries get the translator defaults
+0028aa3eeaddb7bb	{"local":[8,8,1],"shared":0}
+0186bb5336cd30ac	{"local":[8,8,1],"shared":0}
+03b51cc71d1b7a3e	{"local":[8,8,1],"shared":0}
+0407745627fc5725	{"local":[8,8,1],"shared":0}
+0483cab9c1c16c94	{"local":[8,8,1],"shared":0}
+04b880c353236cdb	{"local":[8,8,1],"shared":0}
+04fb9b8969086e40	{"local":[8,8,1],"shared":0}
+07260e839e1d1ca8	{"local":[8,8,1],"shared":0}
+07a77e795c1eea1c	{"local":[32,1,1],"shared":0}
+08595e4fa3a63e5f	{"local":[8,8,1],"shared":0}
+0bc1b225590bf13e	{"local":[8,8,1],"shared":0}
+0d16b80ebd653489	{"local":[8,8,1],"shared":0}
+0e3d4bc176c708bb	{"local":[8,8,1],"shared":0}
+0fd2e7077a4931d6	{"local":[8,8,1],"shared":0}
+111cd8a80f2ec8ea	{"local":[8,4,8],"shared":0}
+11b966aff6ad7d0e	{"local":[8,8,1],"shared":0}
+12164b4049c4a027	{"local":[8,8,1],"shared":0}
+135a3d74c2b9a27b	{"local":[8,8,1],"shared":0}
+144540de4be9a5c0	{"local":[8,8,1],"shared":0}
+1473c895221f57e2	{"local":[8,8,1],"shared":0}
+1508f2a63516343b	{"local":[8,8,1],"shared":0}
+164ab2cf18142d39	{"local":[8,8,1],"shared":0}
+16998dcd41c8a6b7	{"local":[8,8,1],"shared":0}
+170f256810b6898c	{"local":[8,8,1],"shared":0}
+17265d6ac98ee5a8	{"local":[32,1,1],"shared":0}
+1729619d7917be8d	{"local":[8,8,1],"shared":0}
+17497c0eb55c8964	{"local":[8,8,1],"shared":0}
+175e02d6349df066	{"local":[8,8,1],"shared":0}
+17d5aa8834f0c68f	{"local":[8,8,1],"shared":0}
+195e587fd0b38c85	{"local":[8,8,1],"shared":0}
+19a8d50138283181	{"local":[32,1,1],"shared":0}
+19e132016766b99f	{"local":[8,8,1],"shared":0}
+1a2460dd6445c780	{"local":[8,8,1],"shared":0}
+1a2713be96a718e1	{"local":[8,8,1],"shared":0}
+1b804f702e759ab2	{"local":[8,8,1],"shared":0}
+1c0738d5d6b3b15b	{"local":[8,8,1],"shared":0}
+1c561189213c421f	{"local":[32,1,1],"shared":0}
+1d0fe434c038e6ff	{"local":[8,8,1],"shared":0}
+1e014d9299a4652f	{"local":[8,8,1],"shared":0}
+1e1a9f2829da8ddb	{"local":[8,8,1],"shared":0}
+1e404b1807a8f206	{"local":[8,8,1],"shared":0}
+1e6deacceef492fc	{"local":[8,8,1],"shared":0}
+1f0b4657f8822604	{"local":[8,8,1],"shared":0}
+1fc6cd30bd625e57	{"local":[8,8,1],"shared":0}
+20b31727ebef8b3b	{"local":[8,8,1],"shared":0}
+2186b0508337d277	{"local":[8,8,1],"shared":1536}
+2307f9faf18eb0b6	{"local":[8,8,1],"shared":0}
+23e73322c998c8e7	{"local":[8,8,1],"shared":0}
+24627e0e2ed16506	{"local":[8,8,1],"shared":0}
+26e6cef2acdea7be	{"local":[8,8,1],"shared":0}
+26e820dc12bd86cb	{"local":[8,8,1],"shared":0}
+273709ab68bd00dd	{"local":[64,1,1],"shared":0}
+2903ab25d2a7e076	{"local":[1,1,1],"shared":0}
+2a14541311e09278	{"local":[8,8,1],"shared":0}
+2a6de55e438a31e2	{"local":[8,8,1],"shared":0}
+2ab229bb28a97646	{"local":[8,8,1],"shared":0}
+2ac51fbfadd0f59c	{"local":[32,1,1],"shared":0}
+2ae4c4b68d4bdb1f	{"local":[8,8,1],"shared":0}
+2b00461e72d51005	{"local":[8,8,1],"shared":0}
+2b1fe75d3099c4f9	{"local":[8,8,1],"shared":0}
+2b56734793417f3f	{"local":[8,8,1],"shared":0}
+2b5c1d93c112b7fa	{"local":[8,8,1],"shared":0}
+2cf8eeac99e05694	{"local":[8,8,1],"shared":0}
+2d05c0f448774d88	{"local":[8,8,1],"shared":0}
+2dbd29c9fdb93a54	{"local":[32,32,1],"shared":4096}
+2e38097cb93ad835	{"local":[32,1,1],"shared":0}
+2eb32ed362b43c19	{"local":[8,8,1],"shared":0}
+2f5d791e4eb8eeb0	{"local":[8,8,1],"shared":0}
+326aa6263ff090c8	{"local":[8,8,1],"shared":0}
+3467b8e31b677fd3	{"local":[8,8,1],"shared":0}
+354f01e6dd882c30	{"local":[8,8,1],"shared":0}
+356a0731f5512735	{"local":[8,8,1],"shared":0}
+359712c3419f977a	{"local":[8,8,1],"shared":0}
+36215514c8edb1ea	{"local":[64,1,1],"shared":0}
+36f5ee8de43bf33d	{"local":[8,8,1],"shared":0}
+3894f0f9e9d98ec0	{"local":[8,8,1],"shared":0}
+39147e2c4bc7baa3	{"local":[8,8,1],"shared":0}
+39538437725cb5d1	{"local":[8,8,1],"shared":0}
+3b7812cdf5cf2a82	{"local":[8,8,1],"shared":0}
+3c139fa0548a51a7	{"local":[8,8,1],"shared":0}
+3c84cf7c9827fc7d	{"local":[8,8,1],"shared":0}
+3d94d437abc1d1b4	{"local":[8,8,1],"shared":0}
+3da6b6568eadaf7f	{"local":[8,8,1],"shared":0}
+3e5dd1fda9e8bf98	{"local":[32,1,1],"shared":0}
+40d5ebb1d93541b7	{"local":[8,8,1],"shared":0}
+42b06ec3823010d4	{"local":[8,8,1],"shared":0}
+440f637ca2dbb512	{"local":[8,8,1],"shared":0}
+44261b623ccf54dd	{"local":[8,8,1],"shared":0}
+45bc563f436c8959	{"local":[8,8,1],"shared":0}
+47aa8e13d609895e	{"local":[8,8,1],"shared":0}
+482a21f6ca3b0d20	{"local":[64,1,1],"shared":0}
+489efb0560f9f951	{"local":[16,16,1],"shared":0}
+4969697e95451fdd	{"local":[8,8,1],"shared":0}
+49afe5c8362a11b9	{"local":[8,8,1],"shared":0}
+4ae937be5af9e031	{"local":[8,8,1],"shared":0}
+4b586428d157d61e	{"local":[32,1,1],"shared":0}
+4b5b1d389cde43d9	{"local":[8,8,1],"shared":1536}
+4c67fddeeb3cc4be	{"local":[8,8,1],"shared":0}
+4cc4664aa53f6a19	{"local":[32,1,1],"shared":0}
+4ce09afcb3529bae	{"local":[8,8,1],"shared":0}
+4d9897afa1f171ad	{"local":[8,8,1],"shared":0}
+4edd84f6d36a8bd6	{"local":[8,8,1],"shared":0}
+4f1da3d7866b4f60	{"local":[8,8,1],"shared":0}
+4f7963701ef1b0ad	{"local":[8,8,1],"shared":0}
+50081e72c59bc7f5	{"local":[32,1,1],"shared":0}
+502e8eb113aa9f8c	{"local":[8,8,1],"shared":0}
+54cfb91f892931cf	{"local":[8,8,1],"shared":0}
+5605248456a5b7f3	{"local":[8,8,1],"shared":0}
+56de33f888e11df8	{"local":[4,8,1],"shared":1920}
+58a68a419ff1122d	{"local":[8,8,1],"shared":0}
+59014835e284980e	{"local":[8,8,1],"shared":0}
+594e403c16349073	{"local":[8,8,1],"shared":0}
+5b571bfe473b0489	{"local":[8,8,1],"shared":0}
+5c2ae517baaba042	{"local":[8,8,1],"shared":0}
+5daf04a461f59a3d	{"local":[8,8,1],"shared":0}
+618ac5d971d43161	{"local":[32,1,1],"shared":0}
+62861ceb940a083e	{"local":[64,1,1],"shared":1536}
+637f8be7df332b8f	{"local":[8,8,1],"shared":0}
+65469f765ee157ae	{"local":[8,8,1],"shared":0}
+666287b0d97c5d61	{"local":[8,8,1],"shared":0}
+683f38fb2fd84ff2	{"local":[8,8,1],"shared":0}
+6a4a8e7eb5506930	{"local":[8,8,1],"shared":0}
+6bc31993164d5368	{"local":[8,8,1],"shared":0}
+6c9752b7ea5a165c	{"local":[8,8,1],"shared":0}
+6e942a46eda8110a	{"local":[8,8,1],"shared":0}
+6f48f62296a72810	{"local":[8,8,1],"shared":0}
+70930aa7b7e386d4	{"local":[8,8,1],"shared":0}
+709b273f8f060927	{"local":[8,8,1],"shared":0}
+70f2cb76c6ca40f0	{"local":[128,1,1],"shared":0}
+714b25f06a30a782	{"local":[8,8,1],"shared":0}
+71c6c2fbbc82b91f	{"local":[8,8,1],"shared":0}
+73bca0b21d434f03	{"local":[8,8,1],"shared":0}
+765457ee148760b8	{"local":[32,1,1],"shared":0}
+76fa73948d6c5f75	{"local":[8,8,1],"shared":0}
+7711a1d0899c5226	{"local":[32,1,1],"shared":8448}
+778cb8a1adc6f737	{"local":[8,8,1],"shared":0}
+77bbe03644edaab7	{"local":[8,8,1],"shared":0}
+78d47a1385be625c	{"local":[8,8,1],"shared":0}
+796fe1c63fa9df94	{"local":[64,1,1],"shared":0}
+7bbbc24f8dd9f573	{"local":[8,8,1],"shared":0}
+7ceb22df3d773fdf	{"local":[8,8,1],"shared":0}
+7cfaae47e5e3f1f0	{"local":[4,8,1],"shared":1920}
+7d5bb2f818132d8b	{"local":[8,8,1],"shared":0}
+7dbd01ada9e9e9b5	{"local":[8,8,1],"shared":0}
+7de0d5f5078f1a3f	{"local":[64,1,1],"shared":0}
+7e465a75fcec085a	{"local":[8,8,1],"shared":0}
+7ee10c33e42525c3	{"local":[32,1,1],"shared":0}
+7ee6df3d1631faff	{"local":[2,1,1],"shared":0}
+8048bce9624046fe	{"local":[32,1,1],"shared":0}
+81154f26080803d0	{"local":[64,1,1],"shared":0}
+85ebec1c02b09b58	{"local":[8,8,1],"shared":0}
+86dd868cdda9b360	{"local":[64,1,1],"shared":0}
+88f7743875ce9e1a	{"local":[8,8,1],"shared":0}
+892023d7f0d93369	{"local":[8,8,1],"shared":0}
+894b72869453e851	{"local":[8,8,1],"shared":0}
+8a161d3e9ef3f5e1	{"local":[8,8,1],"shared":0}
+8ab2a983df9f7faa	{"local":[32,1,1],"shared":0}
+8b51be4f4ac51bac	{"local":[8,8,1],"shared":0}
+8bfcbe312be72246	{"local":[4,8,1],"shared":1920}
+8c5f6364d2c98901	{"local":[8,8,1],"shared":0}
+8f6fb416c1787ca4	{"local":[32,1,1],"shared":0}
+8fb3d128a1604f3e	{"local":[8,8,1],"shared":0}
+90c390d6409d5996	{"local":[8,8,1],"shared":0}
+91e44fca88c08b4e	{"local":[8,8,1],"shared":0}
+926a0bfcc0bb9560	{"local":[8,8,1],"shared":0}
+92d9fd146d428a70	{"local":[8,8,1],"shared":0}
+92fb8ceb351964d2	{"local":[8,8,1],"shared":0}
+930a997f7225bc89	{"local":[8,8,1],"shared":0}
+9431f0a69b50f0cd	{"local":[8,8,1],"shared":0}
+943b2afeac4469db	{"local":[8,8,1],"shared":0}
+94bad79a8b98b1b5	{"local":[64,1,1],"shared":0}
+94e843ef9450bcd5	{"local":[8,8,1],"shared":0}
+9581a28f248130e4	{"local":[8,8,1],"shared":0}
+9781cf3ee144d468	{"local":[8,8,1],"shared":0}
+98104a2fd30443e6	{"local":[64,1,1],"shared":0}
+9836fb8b96783d87	{"local":[32,1,1],"shared":0}
+98a696dc398b6594	{"local":[8,8,1],"shared":0}
+98b04cc2167477f2	{"local":[8,8,1],"shared":0}
+99670b82347f0d27	{"local":[8,8,1],"shared":0}
+99b033ebba255119	{"local":[8,8,1],"shared":0}
+99c0c70ec70f9a68	{"local":[8,8,1],"shared":0}
+9b2746109eb9a812	{"local":[8,8,1],"shared":0}
+9c13b8e94a78a9d0	{"local":[8,8,1],"shared":0}
+9d9a9acbe6a55ee2	{"local":[8,8,1],"shared":0}
+9dca0cf5c0d1f801	{"local":[8,8,1],"shared":0}
+9e5beae2b2d661e1	{"local":[8,8,1],"shared":0}
+9ea6f47fe316732b	{"local":[8,8,1],"shared":0}
+a1119d9c7dc9a984	{"local":[8,8,1],"shared":0}
+a21c098bb3c74d02	{"local":[8,8,1],"shared":0}
+a22297dafaef3cf2	{"local":[8,8,1],"shared":0}
+a2e28363d3d51c0b	{"local":[8,8,1],"shared":0}
+a3fdd6cc153dc880	{"local":[8,8,1],"shared":0}
+a68bd599ae5e28d6	{"local":[8,8,1],"shared":0}
+a7b99b22e2d8eca9	{"local":[32,1,1],"shared":0}
+a7c96fa99ffd67a6	{"local":[8,8,1],"shared":0}
+a85d46e5229ec68d	{"local":[8,8,1],"shared":0}
+acd84202f8a6cd32	{"local":[8,8,1],"shared":0}
+ad89bd8ec5c9e32d	{"local":[8,8,1],"shared":0}
+ad8acb40eec43197	{"local":[64,1,1],"shared":1536}
+af1401e2cb1cb6ea	{"local":[8,8,1],"shared":0}
+b124b8c55f5922e9	{"local":[32,1,1],"shared":0}
+b32f446f4b6a3b2b	{"local":[8,8,1],"shared":0}
+b3dd60a3b567b644	{"local":[8,8,1],"shared":0}
+b401cc01fab28fd8	{"local":[128,1,1],"shared":0}
+b4818ffdc25b8ee5	{"local":[8,8,1],"shared":0}
+b4b3c7c154ba5a76	{"local":[8,8,1],"shared":0}
+b7cf1ca06019d7bf	{"local":[8,8,1],"shared":0}
+b83486c0f30a2d3a	{"local":[8,8,1],"shared":0}
+b94b82fa423022be	{"local":[64,1,1],"shared":1792}
+b9c41b91c4f558d0	{"local":[8,8,1],"shared":0}
+b9cd92a4c8584eb0	{"local":[64,1,1],"shared":0}
+bb0dc681ca05c3d8	{"local":[64,1,1],"shared":768}
+bb47bb055a779e07	{"local":[32,1,1],"shared":0}
+bd331a7b689be4bf	{"local":[128,1,1],"shared":0}
+bd678ff4011e8fe0	{"local":[8,8,1],"shared":0}
+bd75202645652a15	{"local":[8,8,1],"shared":0}
+bdfca84add6904fb	{"local":[8,8,1],"shared":0}
+bf0ec331cf80cd1d	{"local":[8,8,1],"shared":0}
+bff1fa30a8a8b49a	{"local":[8,8,1],"shared":0}
+c167fa5d97ce741e	{"local":[8,8,1],"shared":0}
+c204d118c9e5211a	{"local":[8,8,1],"shared":0}
+c2189b995d2368aa	{"local":[8,8,1],"shared":0}
+c2df95e6c93a866c	{"local":[64,1,1],"shared":0}
+c51fa1e565ce071f	{"local":[8,8,1],"shared":0}
+c723c2f56c1c5b1f	{"local":[8,8,1],"shared":0}
+c7871975951e76ca	{"local":[8,8,1],"shared":0}
+c99268f92cd6ae5c	{"local":[32,1,1],"shared":0}
+ca6de06c889d3808	{"local":[8,8,1],"shared":0}
+cb225eac17f4bcb3	{"local":[8,8,1],"shared":0}
+ccc56c287960b8b2	{"local":[64,1,1],"shared":1536}
+cd0222a57d4da396	{"local":[8,8,1],"shared":0}
+ce2ae5b8817f0331	{"local":[32,1,1],"shared":0}
+ceaa373504e27441	{"local":[8,8,1],"shared":0}
+cf5ba107e29b51f4	{"local":[8,8,1],"shared":0}
+cff504a890b78d7c	{"local":[32,1,1],"shared":0}
+d038b8ed760fe198	{"local":[8,8,1],"shared":0}
+d22121b40e3450d6	{"local":[8,8,1],"shared":0}
+d395dd8123169ed6	{"local":[8,8,1],"shared":0}
+d40a2eda51ec9ebd	{"local":[8,8,1],"shared":0}
+d49f9fee7588ca71	{"local":[8,8,1],"shared":0}
+d65a46cd3b8fd5ab	{"local":[32,1,1],"shared":0}
+d6a1a9ddd5071cc8	{"local":[8,8,1],"shared":0}
+d76d95fe9e7b8c09	{"local":[8,8,1],"shared":0}
+d8558c80ff97e718	{"local":[8,8,1],"shared":0}
+d8adeb7e60ea9399	{"local":[8,8,1],"shared":0}
+d8e2c5dc2f705bf8	{"local":[8,8,1],"shared":0}
+d998e4ab7aa2cae6	{"local":[32,1,1],"shared":0}
+da1f85d4c0045e54	{"local":[8,8,1],"shared":0}
+da3a5add020a9521	{"local":[8,8,1],"shared":0}
+dc4e651608f5c15e	{"local":[8,8,1],"shared":0}
+dd013c67d0ef9044	{"local":[8,8,1],"shared":0}
+dd98e759fff62dbb	{"local":[32,1,1],"shared":0}
+e0ed5204752d9e3e	{"local":[8,8,1],"shared":0}
+e11b0d58ddc7a87d	{"local":[8,8,1],"shared":0}
+e1d1b7528487a1b3	{"local":[32,32,1],"shared":0}
+e25cf8098e3f328e	{"local":[8,8,1],"shared":0}
+e415949135d42f00	{"local":[8,8,1],"shared":0}
+e4a1f5c30542a895	{"local":[8,8,1],"shared":0}
+e554ca5de561f5ed	{"local":[32,1,1],"shared":0}
+e7a266db7c348380	{"local":[8,8,1],"shared":0}
+e885562af2f4cbfc	{"local":[8,8,1],"shared":0}
+e8ba97e2cc9dabfe	{"local":[8,8,1],"shared":0}
+eb1465404ff41241	{"local":[8,8,1],"shared":0}
+ec082260a4a47648	{"local":[8,8,1],"shared":0}
+ec38f1a530d1de85	{"local":[8,8,1],"shared":0}
+ec3f64537c02a96a	{"local":[8,8,1],"shared":0}
+ecf2bcf15f667769	{"local":[8,8,1],"shared":0}
+ecf3f6c920b1861f	{"local":[8,8,1],"shared":0}
+eed76993185c8073	{"local":[64,1,1],"shared":1536}
+ef5871de5e074147	{"local":[8,8,1],"shared":0}
+ef6e080fe4635b05	{"local":[8,8,1],"shared":0}
+ef93b88dc109fccf	{"local":[8,8,1],"shared":0}
+efcfc3bab2c680f7	{"local":[8,8,1],"shared":0}
+f0824d9801e33ba0	{"local":[8,8,1],"shared":0}
+f0eded2e082fcb4b	{"local":[8,8,1],"shared":0}
+f30ad12ae6a9171a	{"local":[32,1,1],"shared":0}
+f490ce0700c4b77f	{"local":[8,8,1],"shared":0}
+f4a579eb5d97ec5a	{"local":[8,8,1],"shared":0}
+f594bad597e7d7bc	{"local":[8,8,1],"shared":0}
+fac1a79858e5758f	{"local":[8,8,1],"shared":0}
+fb9ac0d43b4fc39c	{"local":[8,8,1],"shared":0}

dm/games/0100317013770000/shader_map.tsv.conflicts.tsv

@@ -0,0 +1,2 @@
+# kind: evidence  control-section key conflict rulings: empty
+dm_ctl_sha16	chosen_pc	class	candidates	records

dm/games/README.md

@@ -77,6 +77,7 @@ One row per profile; the notes are each profile's `README.md`.
 | [Kirby and the Forgotten Land](01004D300C5AE000/README.md) | `01004D300C5AE000` | 1.1.0 | 11.4.0 | HAL Laboratory's engine (NintendoWare) | NVN |
 | [Luigi's Mansion 3](0100DCA0064A6000/README.md) | `0100DCA0064A6000` | 1.4.0 | 7.3.2 | Next Level Games' engine | NVN |
 | [Mario Kart 8 Deluxe](0100152000022000/README.md) | `0100152000022000` | 4.0.0 | 4.4.0 | Nintendo (sead / agl / gsys, NintendoWare, Pia) | NVN |
+| [Mario + Rabbids Sparks of Hope](0100317013770000/README.md) | `0100317013770000` | 1.6.2225577 | 13.3.5 | Ubisoft Snowdrop | NVN |
 | [Mario vs. Donkey Kong](0100B99019412000/README.md) | `0100B99019412000` | 1.0.1 (cartridge base) | 15.3.0 | Nintendo (sead / agl / gsys) | NVN |
 | [Metroid Dread](010093801237C000/README.md) | `010093801237C000` | 2.1.0 | 11.4.3 | MercurySteam's engine | NVN |
 | [Monster Hunter Rise](0100B04011742000/README.md) | `0100B04011742000` | 16.0.2 | 11.4.3 | Capcom's RE Engine | NVN |

dm/gfx/nvn/hle_nvn.c

@@ -255,7 +255,7 @@ NVN(nvnTexturePoolRegisterTexture, void, const NVNtexturePool *tp, int id, const
 // Image descriptors live in the same pool index space as texture descriptors, and image handles are pool indices (nvnDeviceGetImageHandle; BindImages below), so an image
 // registers into the same table. Hyrule Warriors: Age of Calamity registers the images its compute shaders write in the first battle (earlier games registered only textures)
 NVN(nvnTexturePoolRegisterImage, void, const NVNtexturePool *tp, int id, const NVNtexture *t, const NVNtextureView *v) {
-    mtl_reg_tex(dpool_base(g_tpb, g_ntpb, tp) + id, (const tex_t *)t, (const view_t *)v);
+    mtl_reg_image(dpool_base(g_tpb, g_ntpb, tp) + id, (const tex_t *)t, (const view_t *)v);
 }
 NVN(nvnSamplerPoolInitialize, NVNboolean, NVNsamplerPool *sp, const NVNmemoryPool *p, ptrdiff_t off, int n) {
     *(dpool_t *)sp = (dpool_t){(const pool_t *)p, off, n}; OLOG("SamplerPool n=%d", n); dpool_add(g_spb, &g_nspb, &g_snext, sp, n, "sampler"); return 1;

dm/gfx/nvn/nvn_mtl.m

@@ -1838,7 +1838,7 @@ static id<MTLTexture> tex_upgrade(const tex_t *t, id<MTLTexture> base, MTLTextur
         if (!named) tex_put(t->serial, upgraded);
         defer_obj_release(base);
     }
-    MLOG("texture upgraded to allow %s ser=%llu base format %lu → view format %lu (upgrades=%llu)", add & MTLTextureUsageRenderTarget ? "render-target use" : "format views", t->serial, (unsigned long)base.pixelFormat, (unsigned long)view_pf, ++g_pfv_upgrades);
+    MLOG("texture upgraded to allow %s ser=%llu base format %lu → view format %lu (upgrades=%llu)", add & MTLTextureUsageShaderWrite ? "image writes" : add & MTLTextureUsageRenderTarget ? "render-target use" : "format views", t->serial, (unsigned long)base.pixelFormat, (unsigned long)view_pf, ++g_pfv_upgrades);
     return upgraded;
 }
 static id<MTLTexture> tex_upgrade_pfv(const tex_t *t, id<MTLTexture> base, MTLPixelFormat view_pf) { return tex_upgrade(t, base, MTLTextureUsagePixelFormatView, view_pf); }
@@ -3884,6 +3884,14 @@ void mtl_reg_tex(int tid, const tex_t *t, const view_t *v) {
                  v ? v->set_format : 0, v ? v->format : 0);
     }
 }
+// RegisterImage explicitly declares image use. Keep that intent in the host snapshot without changing the guest texture or forcing every render target to lose Metal's compression.
+// If it already has a backing, image_prepare
+// adds the usage before the first draw / dispatch that uses this descriptor, preserving the previous GPU content through tex_upgrade.
+void mtl_reg_image(int tid, const tex_t *t, const view_t *v) {
+    if (!t) { mtl_reg_tex(tid, NULL, v); return; }
+    tex_t image = *t; image.d.flags |= NVN_TEXTURE_FLAGS_IMAGE;
+    mtl_reg_tex(tid, &image, v);
+}
 void mtl_reg_samp(int sid, const samp_par *sp) {
     int id_ = sid; if (id_ < 0) return;
     os_unfair_lock_lock(&g_rmu);
@@ -5943,7 +5951,13 @@ static id<MTLTexture> img_binding(int tid, const nvn_rbind *b, uint64_t *rid, in
     void *tv = NULL; *rid = tid ? tex_fast_rid(tid, &tv) : 0;
     id<MTLTexture> t = (__bridge id<MTLTexture>)tv;
     int swz_only = t && !(t.usage & MTLTextureUsageShaderWrite) && (t.parentTexture.usage & MTLTextureUsageShaderWrite);
-    if (t && !(t.usage & MTLTextureUsageShaderWrite) && !swz_only) { RED_ONCE(32, "texture bound to an image slot lacks the shader-write usage (texture ID %d)", tid); t = nil; }
+    if (t && (b->rw || b->atom) && !(t.usage & MTLTextureUsageShaderWrite) && !swz_only) {
+        const tex_t *tt = tid > 0 && tid < g_ntreg && g_treg[tid].ok ? &g_treg[tid].t : NULL;
+        RED_ONCE(32, "texture bound to an image slot lacks the shader-write usage (texture ID %d, serial %llu, target %d, format %#x, flags %#x, Metal usage %#lx, parent usage %#lx, slot %u, writes %u)",
+                 tid, tt ? (unsigned long long)tt->serial : 0ull, tt ? (int)tt->d.target : -1, tt ? (unsigned)tt->d.format : 0u, tt ? (unsigned)tt->d.flags : 0u,
+                 (unsigned long)t.usage, (unsigned long)t.parentTexture.usage, (unsigned)b->reg, (unsigned)b->rw);
+        t = nil;
+    }
     if (t && (swz_only || (b->dpf != NVN_DPF_FLOAT && mtl_pf_dpf(t.pixelFormat) != b->dpf))) {
         uint64_t r2 = 0; id<MTLTexture> v = img_fit(tid, t, b->dpf, &r2);
         if (v) { t = v; *rid = r2; } else if (swz_only) t = nil;
@@ -5952,6 +5966,40 @@ static id<MTLTexture> img_binding(int tid, const nvn_rbind *b, uint64_t *rid, in
     if (!t) { t = dummy_img(b->dtt, b->dpf); *rid = g_dummy_img_rid[b->dtt & 7][b->dpf & 3]; }
     return t;
 }
+// Resolve shader image writes before an encoder is opened. Snowdrop's cf5ba107e29b51f4 writes a COMPRESSIBLE-only R16F texture registered with RegisterTexture, not RegisterImage.
+// An image bind may follow rendering into a texture created without IMAGE; upgrading it inside encode_args would
+// copy out of an attachment while its render pass was still open. Close that pass first, then replace the shared backing (all registered sampled views follow g_ser_gen).
+static void image_prepare(NvProg *pr, unsigned stages) {
+    if (!pr) return;
+    for (int st = 0; st < 6; st++) {
+        if (!(stages & (1u << st))) continue;
+        NvShader *sh = pr->sh[st]; if (!sh) continue;
+        for (int i = 0; i < sh->rf.nb; i++) {
+            const nvn_rbind *b = &sh->rf.b[i]; if (b->kind != NVN_RB_IMG || b->reg >= 32 || (!b->rw && !b->atom)) continue;
+            int tid = (int)(S.img[st][b->reg] & 0xfffff); if (tid <= 0) continue;
+            os_unfair_lock_lock(&g_rmu);
+            if (tid >= g_ntreg || !g_treg[tid].ok) { os_unfair_lock_unlock(&g_rmu); continue; }
+            // The actual image binding is image intent too: some engines write a descriptor with RegisterTexture and use its handle in an image slot.
+            tex_t image = g_treg[tid].t; image.d.flags |= NVN_TEXTURE_FLAGS_IMAGE;
+            id<MTLTexture> base = tex_mtl(&image); while (base.parentTexture) base = base.parentTexture;
+            int upgrade = base && !(base.usage & MTLTextureUsageShaderWrite);
+            os_unfair_lock_unlock(&g_rmu);
+            if (!upgrade) continue;
+            rp_end();
+            os_unfair_lock_lock(&g_rmu);
+            // A guest registration can change while the pass ends; resolve its current snapshot again under the descriptor lock.
+            if (tid < g_ntreg && g_treg[tid].ok) {
+                image = g_treg[tid].t; image.d.flags |= NVN_TEXTURE_FLAGS_IMAGE;
+                base = tex_mtl(&image); while (base.parentTexture) base = base.parentTexture;
+                if (base && !(base.usage & MTLTextureUsageShaderWrite)) {
+                    tex_upgrade(&image, base, img_usage(&image, base.pixelFormat), base.pixelFormat);
+                    args_dirty_all();
+                }
+            }
+            os_unfair_lock_unlock(&g_rmu);
+        }
+    }
+}
 // Image atomics run natively on the image's own texture (dmsh_atomic.py), so they and every other access of the image see the same texels: a block-linear image's Metal texture as
 // much as a linear one's pool memory. They used to go through a raw pointer beside the texture, which for a block-linear image (a deswizzled copy) could only be a scratch buffer
 // nothing read, and a 3D image's did not compile. Metal allows texture atomics only on an R32Uint / R32Sint texture made with the shader-atomic usage (img_usage); any other
@@ -8476,6 +8524,7 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
     g_mdxc.next = NULL;
     if (!S.prog) { RED_ONCE(18, "no program bound before the draw"); PD("Z %llu draw skipped: no program bound", g_stat_pass + !g_enc); return; }
     __unsafe_unretained NvProg *pr = g_geo_program ? g_geo_program : prog_get(S.prog);
+    image_prepare(pr, 0x1f);
     draw_gs_check(pr);   // a geometry stage left with no function (see prog_fold_gs)
     if (S.rast_discard) { tf_draw(pr, dw, di); return; }   // the vertex stage only: transform feedback capture, or nothing (tf_draw)
     // Transform feedback records every draw between Begin and End, rasterized or not: Donkey Kong Country: Tropical Freeze captures its fur's lagged positions with rasterization on
@@ -8842,6 +8891,7 @@ static void dispatch(const p_dispatch *dp) {
     if (!tg[0]) { RED_ONCE(29, "compute shader lacks threadgroup size reflection"); return; }
     if (g_nubd && ubd_conflict_cs(cs, dp->ind)) { g_ubd_conflicts++; ubd_land(); }   // see ubd_conflict_cs
     rp_end();
+    image_prepare(prog_get(S.cprog), 0x20);
     int own = dispatch_own();
     if (own && g_blit) { enc_b_end(g_blit); g_blit = nil; }
     id<MTLComputeCommandEncoder> ce = own ? pg_compute() : (id<MTLComputeCommandEncoder>)dispatch_enc();

dm/gfx/nvn/nvn_obj.h

@@ -229,6 +229,7 @@ void mtl_create_stats(uint64_t *lib, uint64_t *pso);   // Cumulative: number of
 void mtl_read_texels(const tex_t *t, unsigned level, unsigned layer0, const NVNcopyRegion *r, uint8_t *dst, size_t row, size_t img);
 // Descriptor pool registration (texture ID / sampler ID → object snapshot), program release
 void mtl_reg_tex(int id, const tex_t *t, const view_t *v);
+void mtl_reg_image(int id, const tex_t *t, const view_t *v);   // RegisterImage declares image use even when the builder omitted IMAGE
 // textures written by WriteTexels (recorded by serial): the replay layer treats them like packaged textures, uploaded from guest memory and re-uploaded on dirty pages
 void mtl_tex_cpu_written(const tex_t *t);
 int mtl_tex_gpu_owned(uint64_t serial);                                  // its host texture is a render target or shader-writable: GPU content, never re-uploaded from guest memory

dm/gfx/tests/.gitignore

@@ -84,6 +84,8 @@ t_upload_budget
 t_upload_budget.dSYM/
 t_view_swizzle
 t_view_swizzle.dSYM/
+t_image_usage
+t_image_usage.dSYM/
 t_vk_desc
 t_vk_desc.dSYM/
 t_gl_imports

dm/gfx/tests/Makefile

@@ -223,6 +223,8 @@ check_t_image_alias: t_image_alias
 	DM_NVN_ENC_THREAD=0 MTL_DEBUG_LAYER=1 ./t_image_alias
 	DM_NVN_ENC_THREAD=1 MTL_DEBUG_LAYER=1 ./t_image_alias
 run: check_t_image_alias
+t_image_usage: t_image_usage.m $(MTL_T_DEP)
+	clang $(CFLAGS) -fobjc-arc -o $@ t_image_usage.m $(MTL_T_SRC) $(FW)
 t_stencil_view: t_stencil_view.m $(MTL_T_DEP)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_stencil_view.m $(MTL_T_SRC) $(FW)
 # The Vulkan layer (#included) over the whole NVN layer, whose device and present path it uses
@@ -235,7 +237,7 @@ gl-imports: t_gl_imports
 	./t_gl_imports
 	./t_gl_imports --stubs
 run: gl-imports
-run: t_nvn t_dirty_slots t_dirty_race t_tex_layout t_builtin t_vin t_tbuf2d t_tbuf_align t_deswz_gpu t_pace t_timestamp t_fstep t_refl t_hmap t_cover t_fuse_check t_present_crop t_reorder_chain t_defer_chain t_zskip t_xv t_ub t_vbuf t_geometry t_mip_extent t_rt_range_shape t_rt_range_ro t_rt_range t_rt_depth_ro t_pso_skey t_pso_archive t_pso_bundle t_pull t_pull_inv t_pso_pending t_pfv t_view_levels t_view_swizzle t_vk_desc t_integer_blend t_primitive_restart t_clear_ds t_depth16_clear t_upgrade_alias t_stencil_view
+run: t_nvn t_dirty_slots t_dirty_race t_tex_layout t_builtin t_vin t_tbuf2d t_tbuf_align t_deswz_gpu t_pace t_timestamp t_fstep t_refl t_hmap t_cover t_fuse_check t_present_crop t_reorder_chain t_defer_chain t_zskip t_xv t_ub t_vbuf t_geometry t_mip_extent t_rt_range_shape t_rt_range_ro t_rt_range t_rt_depth_ro t_pso_skey t_pso_archive t_pso_bundle t_pull t_pull_inv t_pso_pending t_pfv t_view_levels t_view_swizzle t_vk_desc t_integer_blend t_primitive_restart t_clear_ds t_depth16_clear t_upgrade_alias t_image_usage t_stencil_view
 	./t_nvn
 	./t_dirty_slots
 	./t_dirty_race
@@ -287,6 +289,7 @@ run: t_nvn t_dirty_slots t_dirty_race t_tex_layout t_builtin t_vin t_tbuf2d t_tb
 	MTL_DEBUG_LAYER=1 ./t_clear_ds
 	MTL_DEBUG_LAYER=1 ./t_depth16_clear
 	MTL_DEBUG_LAYER=1 ./t_upgrade_alias
+	MTL_DEBUG_LAYER=1 ./t_image_usage
 	MTL_DEBUG_LAYER=1 ./t_stencil_view
 	DM_GAME=$(GAME) python3 test_tbdr_budget.py
 .PHONY: run

dm/gfx/tests/t_image_usage.m

@@ -0,0 +1,82 @@
+// RegisterImage on a COMPRESSIBLE-only render texture: first use adds image usage, keeps the open pass's clear and every mip, and updates sampled aliases.
+#include "../nvn/nvn_mtl.m"
+void *front_metal_layer(void) { return NULL; }
+module_t *module_of(const void *p) { (void)p; return NULL; }
+module_t g_mod[MAX_MOD]; int g_nmod;
+void svc_thread_dump(void) {}
+uint64_t now_ns(void) { return bp_now(); }
+static int checks, failed;
+#define CHECK(c, label) do { checks++; if (!(c)) { failed++; fprintf(stderr, "FAIL: %s\n", label); } } while (0)
+static void read_level(id<MTLTexture> texture, unsigned level, uint16_t *out) {
+    NSUInteger n = texture.width >> level;
+    id<MTLBuffer> readback = [g_dev newBufferWithLength:n * n * 2 options:MTLResourceStorageModeShared];
+    id<MTLCommandBuffer> command = q_cb(); id<MTLBlitCommandEncoder> copy = [command blitCommandEncoder];
+    [copy copyFromTexture:texture sourceSlice:0 sourceLevel:level sourceOrigin:MTLOriginMake(0,0,0) sourceSize:MTLSizeMake(n,n,1)
+                toBuffer:readback destinationOffset:0 destinationBytesPerRow:n * 2 destinationBytesPerImage:n * n * 2];
+    [copy endEncoding]; [command commit]; [command waitUntilCompleted];
+    CHECK(command.status == MTLCommandBufferStatusCompleted, "readback completed");
+    memcpy(out, readback.contents, n * n * 2);
+}
+static void image_case(int explicit_image) { @autoreleasepool {
+    tex_t texture = {.magic=MAGIC('T','E','X',0), .serial=99501 + 10 * explicit_image, .off=(uint64_t)explicit_image << 20};
+    texture.d = (nvn_texdesc){.target=NVN_TEXTURE_TARGET_2D, .format=NVN_FORMAT_R16F, .flags=NVN_TEXTURE_FLAGS_COMPRESSIBLE,
+                            .width=8, .height=8, .depth=1, .levels=2, .samples=1};
+    nvn_tex_layout(&texture.d, &texture.L); mtl_tex_init(&texture);
+    tex_t alias = texture; alias.serial++; mtl_tex_init(&alias);
+    id<MTLTexture> old = tex_mtl(&texture);
+    CHECK(old && !(old.usage & MTLTextureUsageShaderWrite), "ordinary render target keeps its original usage");
+    CHECK(tex_mtl(&alias) == old, "another serial shares the original backing");
+    // Seed the smaller mip before the production render pass clears level 0.
+    id<MTLCommandBuffer> seed = q_cb(); MTLRenderPassDescriptor *rp = [MTLRenderPassDescriptor renderPassDescriptor];
+    rp.colorAttachments[0].texture = old; rp.colorAttachments[0].level = 1; rp.colorAttachments[0].loadAction = MTLLoadActionClear;
+    rp.colorAttachments[0].storeAction = MTLStoreActionStore; rp.colorAttachments[0].clearColor = MTLClearColorMake(0.75,0,0,1);
+    [[seed renderCommandEncoderWithDescriptor:rp] endEncoding]; [seed commit]; [seed waitUntilCompleted];
+    CHECK(seed.status == MTLCommandBufferStatusCompleted, "smaller mip seeded");
+    mtl_reg_tex(601, &alias, NULL);
+    os_unfair_lock_lock(&g_rmu); CHECK(tex_by_id(601) == old, "sampled descriptor caches the old backing"); os_unfair_lock_unlock(&g_rmu);
+    mtl_sub_begin(); memset(&R, 0, sizeof R); R.has_rt = 1; R.rt.n = 1; R.rt.color[0] = &texture;
+    R.clear_c[0] = 1; R.cc[0] = MTLClearColorMake(0.5,0,0,1);
+    CHECK(rp_begin() != nil && g_enc != nil, "production render pass is open before image registration");
+    view_t mip = {.set_levels=1, .base_level=1, .nlevels=1};
+    if (explicit_image) mtl_reg_image(602, &texture, &mip); else mtl_reg_tex(602, &texture, &mip);
+    CHECK(texture.d.flags == NVN_TEXTURE_FLAGS_COMPRESSIBLE, "image registration leaves the guest texture flags alone");
+    NvProg *program = [NvProg new]; NvShader *shader = [NvShader new]; program->sh[5] = shader;
+    shader->rf.nb = 1; shader->rf.b[0] = (nvn_rbind){.kind=NVN_RB_IMG, .dtt=MTLTextureType2D, .dpf=NVN_DPF_FLOAT, .reg=0};
+    S.img[5][0] = 602; image_prepare(program, 0x20);
+    os_unfair_lock_lock(&g_rmu); uint64_t read_rid; int read_real;
+    id<MTLTexture> read_only = img_binding(602, &shader->rf.b[0], &read_rid, &read_real);
+    CHECK(read_real && read_only && !(read_only.usage & MTLTextureUsageShaderWrite), "load-only image accepts a read-only texture");
+    CHECK(tex_mtl(&texture) == old && g_enc != nil, "load-only image does not upgrade storage or end the render pass");
+    os_unfair_lock_unlock(&g_rmu);
+    shader->rf.b[0].rw = 1; image_prepare(program, 0x20);
+    CHECK(!g_enc, "image preparation ends the previous render pass before copying it");
+    id<MTLTexture> upgraded = tex_mtl(&texture);
+    CHECK(upgraded != old && (upgraded.usage & MTLTextureUsageShaderWrite), "explicit image use upgrades only its backing");
+    CHECK(tex_mtl(&alias) == upgraded, "sampled alias follows the upgraded backing");
+    os_unfair_lock_lock(&g_rmu);
+    id<MTLTexture> sampled = tex_by_id(601); uint64_t rid; int real;
+    id<MTLTexture> image = img_binding(602, &shader->rf.b[0], &rid, &real);
+    CHECK(sampled == upgraded, "cached sampled descriptor follows the replacement");
+    CHECK(real && image.width == 4 && image.parentRelativeLevel == 1 && (image.usage & MTLTextureUsageShaderWrite), "image descriptor keeps its mip range and writable backing");
+    CHECK(rid == image.gpuResourceID._impl, "image argument names the real view");
+    os_unfair_lock_unlock(&g_rmu);
+    mtl_sub_commit(); mtl_drain();
+    uint16_t pixels[64]; read_level(upgraded, 0, pixels);
+    int same = 1; for (int i = 0; i < 64; i++) same &= pixels[i] == 0x3800; CHECK(same, "upgrade preserves the clear from the pass it ended");
+    read_level(upgraded, 1, pixels); same = 1; for (int i = 0; i < 16; i++) same &= pixels[i] == 0x3a00; CHECK(same, "upgrade preserves the other mip");
+    NSError *error = nil;
+    id<MTLLibrary> lib = [g_dev newLibraryWithSource:@"#include <metal_stdlib>\nusing namespace metal;\nkernel void put(texture2d<float, access::write> t [[texture(0)]]) { t.write(float4(0.25),uint2(1,1)); }\n" options:nil error:&error];
+    id<MTLComputePipelineState> pipeline = [g_dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"put"] error:&error];
+    CHECK(pipeline != nil, "image writer compiled");
+    id<MTLCommandBuffer> write = q_cb(); id<MTLComputeCommandEncoder> encoder = [write computeCommandEncoder];
+    [encoder setComputePipelineState:pipeline]; [encoder setTexture:image atIndex:0];
+    [encoder dispatchThreads:MTLSizeMake(1,1,1) threadsPerThreadgroup:MTLSizeMake(1,1,1)]; [encoder endEncoding]; [write commit]; [write waitUntilCompleted];
+    CHECK(write.status == MTLCommandBufferStatusCompleted, "image write completed");
+    read_level(tex_mtl(&alias), 1, pixels); same = 1; for (int i = 0; i < 16; i++) same &= pixels[i] == (i == 5 ? 0x3400 : 0x3a00);
+    CHECK(same, "sampled alias sees the image write and untouched pixels");
+} }
+int main(void) { @autoreleasepool {
+    mtl_init(); image_case(0); image_case(1);
+    printf("explicit image usage: %d checks, %d failed\n", checks, failed);
+    return failed != 0;
+} }

dm/gfx/tests/test_snowdrop.py

@@ -0,0 +1,93 @@
+"""Snowdrop's compressed shader records, without game data."""
+from pathlib import Path
+import os
+import struct
+import subprocess
+import sys
+import tempfile
+import unittest
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / 'tools'))
+import shader_split as split
+
+
+def shader(stage):
+    ctl, code = bytearray(0x880), bytearray(0x100)
+    ctl[:4] = split.CTL_MAGIC
+    struct.pack_into('<I', ctl, 0x20, len(ctl))
+    struct.pack_into('<I', ctl, 0x704, len(code))
+    code[:4] = split.CS_MAGIC if stage == 'cs' else split.GFX_MAGIC
+    if stage != 'cs':
+        struct.pack_into('<I', code, 0x30, (split.STAGES.index(stage) + 1) << 10 | 1)
+    if stage == 'ps': struct.pack_into('<I', code, 0x48, 1)   # reads generic location 0
+    return bytes(ctl), bytes(code)
+
+
+def record(raw, flags=0x48):
+    packed = split.zstd_mod().compress(raw)
+    return struct.pack('<5I', 0xabcd00c6, flags, 0, len(raw), len(packed)) + packed
+
+
+def archive(*records):
+    return b'BERG' + struct.pack('<I', 22) + bytes(48) + b''.join(records)
+
+
+class SnowdropTests(unittest.TestCase):
+    def test_byte_packed_secondary_stages_and_separable_pairs(self):
+        # A control can be at an unaligned offset; a single asset may hold two stages.
+        raw = b'asset reflection\0x' + b''.join(shader('vs')) + b'variant' + b''.join(shader('ps'))
+        data = archive(record(raw), record(b'x' + b''.join(shader('cs')), flags=0x404))
+        found = split.snowdrop_shaders(data, '/arbitrary-name')
+        self.assertEqual([s for _, s, _, _ in found], ['vs', 'ps', 'cs'])
+        for (_, stage, ctl, code) in found:
+            self.assertEqual((ctl, code), shader(stage))
+        self.assertEqual(split.var_pairs(found), set())
+        self.assertEqual(len(split.separable_pairs(split.uniq_keys(found))[0]), 1)
+
+    def test_scan_selects_by_berg_magic(self):
+        data = archive(record(b''.join(shader('ps'))))
+        with tempfile.TemporaryDirectory() as tmp:
+            path = Path(tmp) / 'without-extension'
+            path.write_bytes(data)
+            rel, keys, pairs, trees = split.scan_file((str(path), 'without-extension'))
+        self.assertEqual(rel, 'without-extension')
+        self.assertEqual(len(keys), 1)
+        self.assertEqual(pairs, set())
+        self.assertEqual(trees, {})
+
+    def test_bare_frames_and_false_control_magics_are_not_shaders(self):
+        data = archive(split.zstd_mod().compress(b''.join(shader('ps'))), record(b'other asset' + split.CTL_MAGIC + bytes(8)))
+        self.assertEqual(split.snowdrop_shaders(data, '/other assets'), [])
+
+    def test_record_bounds_and_decoded_length_are_checked(self):
+        data = archive(record(b''.join(shader('vs'))))
+        with self.assertRaisesRegex(SystemExit, 'stored bytes out of bounds'):
+            split.snowdrop_shaders(data[:-1], '/truncated')
+        bad = bytearray(data)
+        struct.pack_into('<I', bad, 56 + 12, 1)
+        with self.assertRaisesRegex(SystemExit, 'expected 1 raw bytes'):
+            split.snowdrop_shaders(bytes(bad), '/bad raw size')
+
+    def test_dlc_romfs_has_no_main_and_joins_the_shader_union(self):
+        with tempfile.TemporaryDirectory() as tmp:
+            root = Path(tmp)
+            romfs, dlc = root / 'RomFS', root / 'aoc/7'
+            romfs.mkdir(); dlc.mkdir(parents=True)
+            (dlc / 'content').write_bytes(archive(record(b''.join(shader('ps')))))
+            # Minimal uncompressed NSO with an empty code image; all shaders are in the DLC.
+            nso = bytearray(0x1100)
+            nso[:4] = b'NSO0'
+            struct.pack_into('<3I', nso, 0x10, 0x100, 0, 0x1000)
+            struct.pack_into('<3I', nso, 0x30, 0x1100, 0x1000, 0)
+            struct.pack_into('<I', nso, 0x60, 0x1000)
+            main = root / 'main'; main.write_bytes(nso)
+            outputs = [root / n for n in ('pkg', 'map.tsv', 'facts.tsv', 'pairs.tsv', 'pso.tsv')]
+            result = subprocess.run([sys.executable, str(Path(split.__file__)), str(romfs), str(main), *map(str, outputs), str(dlc)],
+                                    env=dict(os.environ, DM_GAME=split.GAME.id, J='1'), capture_output=True, text=True)
+            self.assertEqual(result.returncode, 0, result.stdout + result.stderr)
+            self.assertIn('shaders: 1', result.stdout)
+            self.assertIn('/aoc/7/content#SNOWDROP@', outputs[1].read_text())
+
+
+if __name__ == '__main__':
+    unittest.main()

dm/gfx/tools/shader_split.py

@@ -37,6 +37,7 @@
 #   · control sections loose in main, in no GLSLC output or BNSH file (Advance Wars 1+2: Re-Boot Camp's Bink 2 shaders): paired with their code areas by data records or adjacency, see raw_shaders
 #   · SNVN (Retro Studios, Donkey Kong Country: Tropical Freeze): zlib program records in the SNVN chunks of MaterialArchive.arc, of the MTRL resources in its .pak archives
 #       and of the startup material archive LZSS-compressed in main, see rfrm_shaders and retro_main_forms
+#   · BERG (Snowdrop, Mario + Rabbids Sparks of Hope): shader objects in the individually zstd-compressed records of its SDF data slices, see snowdrop_shaders
 # Code area: 0x30 header + (graphics) 0x50 SPH + instructions; header magic 0x12345678 for graphics, 0x12345679 for compute. The control section starts with 0x98761234, +0x704 = code area length,
 #   +0x700 = start of the c1 constant segment within the code area; for compute +0x748 thread group xyz (u32 ×3), +0x754 shared memory bytes (checked one by one against PC reflection on D2R's 22 compute shaders).
 # Stage: BNSH uses the stage slot; BAHS uses bits 10..13 of SPH word 0 (1 vertex 2 hull 3 domain 4 geometry 5 pixel), and a compute header means compute.
@@ -44,7 +45,7 @@
 #   (SecredUber's 16600 stage shaders have only 3079 distinct control sections, with 6367 cases of same control, different code), so hashing the control alone picks the wrong library.
 # Output: <package dir>/<stage>_<source group>.bin, each record = 'BTCDBTCD' + header (+0x28 control length, +0x2c code length) padded to 0x38 + code area + control section, one copy per key;
 #   map rows have the same columns as D2R's shader_map.tsv (column 5 records where the key first appeared); the answer table has one row per compute shader {"local", "shared"} (shared memory rounded up to 16 bytes, as D2R did).
-# Usage: shader_split.py <RomFS dir> <ExeFS main> <output package dir> <output map> <output answer table> [<output pair table> [<output pipeline pair table (pso_pairs.tsv)>]]   (decompressing main needs lz4: run with dm's venv python)   Environment J=concurrency (default 8)
+# Usage: shader_split.py <RomFS dir> <ExeFS main> <output package dir> <output map> <output answer table> [<output pair table> [<output pipeline pair table (pso_pairs.tsv)> [<DLC RomFS dirs, comma-separated>]]]   Environment J=concurrency (default 8)
 #   With an update: give RomFS and main as comma-separated lists (base game first, e.g. data/v0/RomFS,data/v327680/RomFS); the output is the union of both versions (see main())
 #   Unity type trees need Unity's common string table, which is taken from the first main
 # Pair table: (vertex key, pixel key) of every BNSH variation (every Unity program); dmsh_build uses it to give vertex shaders the outputs their pixel shaders read but they do not write (see dmsh_build.py vs_pad)
@@ -250,6 +251,31 @@ def glslc_shaders(d, label, need=True):
     if not out and need: red(f'{label}: no GLSLC output')
     return out
 
+SNOWDROP_RECORD = bytes.fromhex('c600cdab')
+def snowdrop_shaders(d, label):
+    """Snowdrop's BERG data slices (SDF version 22): compressed shader records have a 20-byte header {0xabcd00c6, flags, 0, raw length, stored length}, then one zstd frame.
+    The shader object inside keeps each control section immediately before its code area; its own +0x20 and +0x704 give their lengths. Both sections are byte-packed, without alignment.
+    Sparks of Hope's base shader slice has 17,931 records with 18,016 stages (85 hold a second stage), and its update slice adds 443 records. Other assets use bare zstd frames or
+    uncompressed data without this header. A census of all 324,413 valid frames in every slice found no further shaders. Shader programs hold one stage; pairs are selected at draw time.
+    Scan by the record header rather than slice names or every possible zstd magic: compressed bytes can contain a coincidental frame header, which is not another asset."""
+    out = []
+    for m in re.finditer(re.escape(SNOWDROP_RECORD), d):
+        h = m.start()
+        if h + 24 > len(d) or d[h + 20:h + 24] != ZSTD_MAGIC: continue
+        nraw, nstored = struct.unpack_from('<II', d, h + 12)
+        if not nstored or h + 20 + nstored > len(d): red(f'{label} Snowdrop record @{h:#x}: stored bytes out of bounds')
+        raw = zstd_mod().decompress(d[h + 20:h + 20 + nstored])
+        if len(raw) != nraw: red(f'{label} Snowdrop record @{h:#x}: expected {nraw} raw bytes, got {len(raw)}')
+        for cm in re.finditer(re.escape(CTL_MAGIC), raw):
+            c = cm.start()
+            if c + 0x708 > len(raw): continue
+            nc, nk = struct.unpack_from('<I', raw, c + 0x20)[0], struct.unpack_from('<I', raw, c + 0x704)[0]
+            if nc < 0x760 or nc & 3 or nk < 0x34 or c + nc + nk > len(raw): continue
+            ctl, code = raw[c:c + nc], raw[c + nc:c + nc + nk]
+            st = glslc_stage(ctl, code)
+            if st and not check(label, st, ctl, code): out.append((f'{label}#SNOWDROP@{h:#x}:{c:#x}', st, ctl, code))
+    return out
+
 NEEDLE = b'HNSHV002'
 def needle_shaders(d, label):
     """Needle shader containers, embedded in main and in GFD's GSC files inside CRI CPKs (Persona 5 Royal). HNSHV002 +8 is the container's whole length, big-endian. Variant and reflection
@@ -1753,7 +1779,7 @@ def ps_inputs(code):
     """Generic attributes a pixel shader reads: SPH words 6–13 are the input map, 8 bits for each of the 32 attributes (Ryujinx ShaderHeader.ImapTypes), non-zero = read; the translated MSL inputs are exactly user(locn<attribute>)"""
     return {i for i in range(32) if struct.unpack_from('<I', code, 0x30 + 4 * (6 + (i >> 2)))[0] >> ((i & 3) * 8) & 0xff}
 
-SINGLE_STAGE = ('#GLSLC@', '#UE', '#MTSPK@', '#MGFX@', '#NLG', '#PGSO', '#RAW@', '#TFXO@', '#NEEDLE@')   # sources whose every shader is a program of one stage: GLSLC outputs, Unreal's shader records, MT Framework's shader package (BNSH variations, Unity programs, MSHD and SDF pair theirs);
+SINGLE_STAGE = ('#GLSLC@', '#UE', '#MTSPK@', '#MGFX@', '#NLG', '#PGSO', '#RAW@', '#TFXO@', '#NEEDLE@', '#SNOWDROP@')   # sources whose every shader is a program of one stage: GLSLC outputs, Unreal's shader records, MT Framework's shader package (BNSH variations, Unity programs, MSHD and SDF pair theirs);
 # MonoGame's effects, whose passes may name no vertex shader and keep the one bound (Celeste: 10 of its 13 effects draw through SpriteBatch's vertex shader, which is in main); Next Level Games'
 # archives (Luigi's Mansion 3 makes 5,317 programs of one stage each at boot); PlatinumGames' shader objects (Bayonetta 3: one stage each, pg_shader); control sections loose in main
 # (Advance Wars 1+2: Re-Boot Camp's Bink shaders, raw_shaders), which no file pairs; and YEBIS's effect file (tfxo_shaders)
@@ -1854,6 +1880,9 @@ def scan_file(args):
     if head[:4] == RFRM:   # a Retro Studios form or pack (rfrm_shaders)
         found += rfrm_shaders(open(path, 'rb').read(), '/' + rel)
         return rel, uniq_keys(found), var_pairs(found), trees
+    if head[:4] == b'BERG':   # Snowdrop's SDF data slices: shader objects in individual compressed records
+        with open(path, 'rb') as f: found += snowdrop_shaders(f.read(), '/' + rel)
+        return rel, uniq_keys(found), var_pairs(found), trees
     if head[:4] == MTARC:   # an MT Framework archive (mtarc_members): its shader package
         mtarc_shaders(open(path, 'rb').read(), '/' + rel, found)
         return rel, uniq_keys(found), var_pairs(found), trees
@@ -1907,6 +1936,7 @@ def main():
     romfs_l, main_l, pkg, mapout, factsout = sys.argv[1:6]
     pairsout = sys.argv[6] if len(sys.argv) > 6 else None
     psopairsout = sys.argv[7] if len(sys.argv) > 7 else None   # pso_pairs.tsv: the pairs the pipeline prebuild takes besides shader_pairs.tsv (the BAHS program variations)
+    dlc_v = sys.argv[8].split(',') if len(sys.argv) > 8 and sys.argv[8] else []   # extra content has no ExeFS; each root is the extracted DLC index
     # Several versions (base game + update): the RomFS and main arguments are comma-separated and paired one to one; the output is their union (keys are content hashes, shaders shared by versions appear once).
     # The first version's source column is unchanged; later versions are prefixed with their workspace directory name (e.g. v327680): /v327680/Shader/…, main@v327680. They sort after the first version, so existing rows keep their package indices.
     romfs_v, main_v = romfs_l.split(','), main_l.split(',')
@@ -1938,7 +1968,10 @@ def main():
         if not names and uepak.unreal_paks(romfs, inline=False): red(f'the RomFS has Unreal paks, but there are no decoded shader code archives in {ued} (the NSP script writes them)')
         for nm in names: found += ue_shaders(open(os.path.join(ued, nm), 'rb').read(), f'/ue_shaders/{nm}')
         results.append(('main', {skey(c, k): (w, s, c, k) for w, s, c, k in found}, var_pairs(found), {}))
-    zdicts = {k: v for r in romfs_v for k, v in zs_dicts(r).items()}   # Tears of the Kingdom's zstd dictionaries (zs_dicts)
+    for romfs in dlc_v:
+        tag = 'aoc/' + os.path.basename(os.path.normpath(romfs))
+        files += sorted((os.path.join(dp, f), tag + '/' + os.path.relpath(os.path.join(dp, f), romfs)) for dp, _, fs in os.walk(romfs) for f in fs)
+    zdicts = {k: v for r in romfs_v + dlc_v for k, v in zs_dicts(r).items()}   # Tears of the Kingdom's zstd dictionaries (zs_dicts)
     keys = {}
     # First pass: every file; Unity bundles carry type trees, and the Shader trees are collected on the way. Second pass: Unity serialized files without type trees are scanned again with the borrowed trees
     J = int(os.environ.get('J', '8'))
@@ -2002,6 +2035,7 @@ def main():
                                             ('MSHD', any('#MSHD#' in r[4] for r in rows)), ('GLSLC', any('#GLSLC@' in r[4] for r in rows)), ('RE Engine SDF', any('#SDF#' in r[4] for r in rows)),
                                             ('MT Framework SPK', any('#MTSPK@' in r[4] for r in rows)), ('MonoGame MGFX', any('#MGFX@' in r[4] for r in rows)),
                                             ('PlatinumGames shader objects', any('#PGSO' in r[4] for r in rows)),
+                                            ('Snowdrop shader objects', any('#SNOWDROP@' in r[4] for r in rows)),
                                             ('Unreal Engine shader code archives', any('#UE' in r[4] for r in rows))) if hit)
         f.write(f'# kind: evidence  {NOTE} ({kinds} in RomFS + main) → control section + code area keys (dm/gfx/tools/shader_split.py)\n')
         f.write('# column 5 records where each key first appears; this game has no PC version to pair with\n')

dm/host/hle_oe.c

@@ -139,6 +139,8 @@ HLE_FORCE("_ZN2nn2oe27GetPerformanceConfigurationENS0_15PerformanceModeE", "quer
 HLE_FORCE("_ZN2nn2oe24SetAlbumImageOrientationENS_5album16ImageOrientationE", "screenshot orientation (through am); the host has no system screenshot", void, h_oe_AlbumOrient, int o) { (void)o; }
 // The Great Ace Attorney Chronicles (SDK 10), first thing in nnMain (main+0x6b188c): the SDK's version asks am through the applet proxy the host never made (null read at sdk+0x338070)
 HLE_FORCE("_ZN2nn2oe34DisableAlbumImageTakenNotificationEv", "screenshot-taken notifications (through am); the host takes no system screenshots", void, h_oe_AlbumNotifyOff, void) {}
+// Enable uses the same am application proxy (SDK 13.3.5 sdk+0x36ec50); the host produces no system-album notifications.
+HLE_FORCE("_ZN2nn2oe33EnableAlbumImageTakenNotificationEv", "screenshot-taken notifications (through am); the host takes no system screenshots", void, h_oe_AlbumNotifyOn, void) {}
 HLE_FORCE("_ZN2nn2oe24IsInStoreDemoModeEnabledEv", "retail store demo mode (through am); never", unsigned, h_oe_StoreDemo, void) { return 0; }
 HLE_FORCE("_ZN2nn2oe27SetPerformanceConfigurationENS0_15PerformanceModeEi", "set the CPU/GPU clock config; the host does not scale clocks", void, h_oe_SetPerfCfg, int mode, int cfg) {
     LOG("oe::SetPerformanceConfiguration mode=%d cfg=%#x (ignored)", mode, cfg);

dm/tools/sdk_patches.py

@@ -60,6 +60,18 @@ CUTS = {
     (0x37c140, '08fca752', '087ca052', 'AddressRegionManager constructor: node area limit mov w8,#0x3fe00000 → mov w8,#0x3e00000'),
     (0x37c1dc, '09fca752', '097ca052', 'AddressRegionManager constructor: length of the node area registered as an occupied region mov w9,#0x3fe00000 → mov w9,#0x3e00000'),
   ],
+  # Mario + Rabbids Sparks of Hope 1.6.2225577 (SDK 13.3.5): the same constructor
+  # and placement code as nirvanA Initiative above. Snowdrop reserves a small
+  # metadata region before its three-pool range (reduced from 12 GB to 3 GB in
+  # the game's patches.py). Both the 62 MB node area and 32 MB alignment are
+  # needed: with the original 1 GB alignment the 3 GB request returns 0x1803,
+  # even after the node-area cut. Its own pools align to 64 KB / 512 KB and do
+  # not require 1 GB alignment of the enclosing address region.
+  '43211771ed39e7fcef03d88e27091831': [
+    (0x37c140, '08fca752', '087ca052', 'AddressRegionManager constructor: node area limit mov w8,#0x3fe00000 → mov w8,#0x3e00000'),
+    (0x37c1dc, '09fca752', '097ca052', 'AddressRegionManager constructor: length of the node area registered as an occupied region mov w9,#0x3fe00000 → mov w9,#0x3e00000'),
+    (0x37c33c, '0a00a852', '0a40a052', 'AllocateAddressRegion: alignment of regions of 1 GB and more mov w10,#0x40000000 → mov w10,#0x2000000 (32 MB, as from 32 MB up)'),
+  ],
   # Pokémon Legends: Z-A 2.0.2 (v393216) sdk (SDK 19.3.11): the same node area cut. The constructor is at sdk+0x1423ec (limit = alias start + 0x3fe00000 at +0x28, a unit heap of 0x30-byte
   # nodes, the area then registered as occupied at +0x142468). With main's allocator region below, the guest block fits the 5 GB proven on a phone
   'df7a3960c74c3a4803f826c587153252': [

dm/tools/tests/test_nsp_incremental.py

@@ -106,6 +106,21 @@ class DeriveProfile(Base):
         self.derive(); (self.pdir / 'pso_manifest.tsv').write_text('S\tk2\tpixel\n'); self.derive()
         self.assertEqual((self.pdir / 'pso_manifest.tsv').read_text(), 'S\tk2\tpixel\n')
 
+    def test_shader_only_derivation_preserves_builds_and_registered_patch_file(self):
+        source = self.src / 'game.tsv'
+        source.write_text(source.read_text() + f'build\tv65536\t1.1.0\tmain={BID2}\nshader_versions\tv0 v65536\n')
+        patch = self.src / 'patches.py'; patch.write_text(f'CUTS = {{"{BID}": []}}\n')
+        g = game.Game(self.src)
+        nsp.derive_profile(g, self.pdir, 'v65536', self.ex, ['main'], shader_only=True)
+        derived = game.Game(self.pdir)
+        self.assertEqual(derived.rows, g.rows)
+        self.assertEqual(derived.builds(), g.builds())
+        self.assertTrue((self.pdir / 'patches.py').is_symlink())
+        self.assertEqual((self.pdir / 'patches.py').resolve(), patch.resolve())
+        derived.table('shader_map.tsv').write_text('generated DLC shader rows\n')
+        nsp.derive_profile(g, self.pdir, 'v65536', self.ex, ['main'], shader_only=True)
+        self.assertEqual(derived.table('shader_map.tsv').read_text(), 'generated DLC shader rows\n')
+
     def test_missing_corrupt_and_duplicate_override_rows_derive_again(self):
         for change in ('missing', 'corrupt', 'override', 'indented override', 'truncated'):
             with self.subTest(change):
@@ -224,6 +239,32 @@ class DeriveProfile(Base):
         self.assertFalse(game.Game(self.pdir).values('mem_alias'))
 
 
+class DlcOnlyShaders(Base):
+    def test_only_dlc_rows_and_facts_can_change(self):
+        tables = self.tmp / 'new-tables'; tables.mkdir()
+        base = 'pkg.bin\t8\tcompute\tbase\t/base.shader\n'
+        dlc = 'pkg.bin\t0\tvertex\tdlc\t/aoc/5/new.shader\n'
+        old_facts = 'base\t{"local":[8,1,1]}\ndlc\t{"other":1}\n'
+        self.src.joinpath('shader_map.tsv').write_text(base + dlc)
+        self.src.joinpath('shader_facts.tsv').write_text(old_facts)
+        g = game.Game(self.src)
+        cases = [
+            ('removed DLC and reindexed base', base.replace('\t8\t', '\t0\t'), 'base\t{"local":[8,1,1]}\n', True),
+            ('added DLC', base + dlc + 'pkg.bin\t9\tpixel\tnew-dlc\t/aoc/7/new.shader\n', old_facts, True),
+            ('shaderless DLC', base + dlc, old_facts, False),
+            ('changed base key', base.replace('\tbase\t', '\tother\t'), old_facts, False),
+            ('changed base source', base.replace('/base.shader', '/different.shader'), old_facts, False),
+            ('changed base facts', base, old_facts.replace('[8,1,1]', '[4,1,1]'), False),
+            ('no DLC difference', base + dlc, old_facts.replace('[8,1,1]', '[4,1,1]'), False),
+            ('invalid map', 'not a shader map\n', old_facts, False),
+        ]
+        for name, shader_map, facts, want in cases:
+            with self.subTest(name=name):
+                tables.joinpath('shader_map.tsv').write_text(shader_map)
+                tables.joinpath('shader_facts.tsv').write_text(facts)
+                self.assertEqual(nsp.dlc_only_shader_change(g, tables), want)
+
+
 class SplitShaders(Base):
     NAMES = ('shader_map.tsv', 'shader_facts.tsv', 'shader_pairs.tsv', 'pso_pairs.tsv')
 
@@ -237,14 +278,70 @@ class SplitShaders(Base):
         cm = mock.patch.object(nsp, 'run', self.split); cm.__enter__(); self.addCleanup(cm.__exit__, None, None, None)
 
     def split(self, cmd, log, env=None, cwd=None, on_line=None):   # shader_split.py: the package and the four tables, deterministic for the same rows
-        self.splits += 1; out = Path(cmd[4]); out.mkdir(parents=True); (out / 'main.bin').write_bytes(self.rows.encode())
+        self.splits += 1; self.last_cmd = cmd; out = Path(cmd[4]); out.mkdir(parents=True); (out / 'main.bin').write_bytes(self.rows.encode())
         for n, t in zip(self.NAMES, cmd[5:]): Path(t).write_text(self.table(n))
         Path(log).write_text('shaders: 3 (3 vertex)\n')
 
-    def table(self, n): return f'# split header\n{n}\t{self.rows}\n'
+    def table(self, n): return self.tables[n] if hasattr(self, 'tables') else f'# split header\n{n}\t{self.rows}\n'
 
-    def go(self, g, own=True, oodle=None):
-        return nsp.split_shaders(sys.executable, self.repo, [self.ex / 'RomFS'], [self.ex / 'ExeFS/main'], self.pkg, {}, self.logs, g, own, oodle)
+    def go(self, g, own=True, oodle=None, dlc_roms=(), derive_dlc=None):
+        return nsp.split_shaders(sys.executable, self.repo, [self.ex / 'RomFS'], [self.ex / 'ExeFS/main'], self.pkg, {}, self.logs, g, own, oodle, dlc_roms, derive_dlc)
+
+    def test_dlc_only_derivation_is_incremental_and_keeps_registered_tables(self):
+        base = 'other.bin\t1\tvertex\tbase\t/base.shader\n'
+        old = {'shader_map.tsv': 'other.bin\t0\tvertex\tdlc\t/aoc/7/new.shader\n' + base,
+               'shader_facts.tsv': 'base\t{"fact":1}\ndlc\t{"fact":2}\n',
+               'shader_pairs.tsv': 'base\tdlc\n', 'pso_pairs.tsv': '# none\n'}
+        for n in self.NAMES: (self.src / n).write_text(old[n])
+        g = game.Game(self.src)
+        self.tables = {'shader_map.tsv': base.replace('\t1\t', '\t0\t'), 'shader_facts.tsv': 'base\t{"fact":1}\n',
+                       'shader_pairs.tsv': '# none\n', 'pso_pairs.tsv': '# none\n'}
+        def derive():
+            nsp.derive_profile(g, self.pdir, 'v65536', self.ex, ['main'], shader_only=True)
+            return game.Game(self.pdir)
+        promote = mock.Mock(side_effect=derive)
+        first = self.go(g, own=False, derive_dlc=promote)
+        self.assertEqual((self.splits, promote.call_count), (1, 1))
+        self.assertEqual(self.go(g, own=False, derive_dlc=promote), first)
+        self.assertEqual((self.splits, promote.call_count), (1, 2))   # proof of the strict non-DLC comparison is reused
+        self.assertEqual(game.Game(self.pdir).rows, g.rows)
+        for n in self.NAMES:
+            self.assertEqual(g.table(n).read_text(), old[n])
+            self.assertEqual(game.Game(self.pdir).table(n).read_text(), self.tables[n])
+        # A new DLC extraction, even of the same file sizes, must perform the strict comparison again.
+        dlc = self.tmp / 'aoc/7'; dlc.mkdir(parents=True); asset = dlc / 'shader'; asset.write_bytes(b'bytes')
+        self.go(g, own=False, dlc_roms=[dlc], derive_dlc=promote); self.assertEqual(self.splits, 2)
+        self.redump(asset)
+        self.go(g, own=False, dlc_roms=[dlc], derive_dlc=promote); self.assertEqual(self.splits, 3)
+        moved = dlc.with_name('8'); dlc.rename(moved)
+        self.go(g, own=False, dlc_roms=[moved], derive_dlc=promote); self.assertEqual(self.splits, 4)   # index changes count even if every file statistic is identical
+        verifier = self.tmp / 'comparison-v2.py'; verifier.write_text('changed comparison code')
+        with mock.patch.object(nsp, '__file__', str(verifier)):
+            self.go(g, own=False, dlc_roms=[moved], derive_dlc=promote); self.assertEqual(self.splits, 5)
+
+    def test_base_shader_mismatch_never_derives_for_dlc(self):
+        (self.src / 'shader_map.tsv').write_text('p.bin\t0\tvertex\tbase\t/base\np.bin\t1\tvertex\tdlc\t/aoc/7/shader\n')
+        (self.src / 'shader_facts.tsv').write_text('# none\n')
+        (self.src / 'shader_pairs.tsv').write_text('# none\n')
+        self.tables = {'shader_map.tsv': 'p.bin\t0\tvertex\twrong-base\t/base\n', 'shader_facts.tsv': '# none\n',
+                       'shader_pairs.tsv': '# none\n', 'pso_pairs.tsv': '# none\n'}
+        promote = mock.Mock()
+        with self.assertRaisesRegex(nsp.BuildError, 'differ from game profile'):
+            self.go(game.Game(self.src), own=False, derive_dlc=promote)
+        promote.assert_not_called()
+
+    def test_dlc_shaders_are_passed_and_invalidate_the_split(self):
+        g = self.own(); self.go(g); self.assertEqual(self.splits, 1)
+        dlc = self.tmp / 'work/aoc/7'; dlc.mkdir(parents=True)
+        asset = dlc / 'shader.archive'; asset.write_bytes(b'new shader')
+        self.go(g, dlc_roms=[dlc]); self.assertEqual(self.splits, 2)
+        self.assertEqual(self.last_cmd[-1], str(dlc))
+        # Extra content has no main: the base RomFS/main pairing remains unchanged.
+        self.assertEqual(self.last_cmd[2:4], [str(self.ex / 'RomFS'), str(self.ex / 'ExeFS/main')])
+        self.go(g, dlc_roms=[dlc]); self.assertEqual(self.splits, 2)
+        asset.write_bytes(b'updated DLC shader')
+        self.go(g, dlc_roms=[dlc]); self.assertEqual(self.splits, 3)
+        self.go(g); self.assertEqual(self.splits, 4)   # removing the DLC removes its shaders from the union
 
     def own(self):   # a derived profile with the empty tables, as main makes it
         self.derive(); g = game.Game(self.pdir); nsp.provisional_prep(sys.executable, self.repo, {}, self.logs, g); return g

docs/agents/symptoms.md

@@ -56,6 +56,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - nothing aborts at startup as with Unity 2022.3, but the provisional profile gave a 1 GB block and a 256 MB alias region while Unity maps 1.5 GB in the first room. → [Unity 6 takes its memory with no StandardAllocator and no heap](../reference/memory.md#unity-6-takes-its-memory-with-no-standardallocator-and-no-heap)
 - `nn::diag` abort from `StandardAllocator::Initialize` long after `nninitStartup`, right after the first Vulkan query pool (`sdk+0x12dad0` ← `main+0x62c52c`). → [The movie player's own StandardAllocator](../reference/memory.md#the-movie-players-own-standardallocator)
 - `StandardAllocator::Initialize` aborts in `nninitStartup` (Pokémon Legends: Z-A: `main+0x1b2e690`); `nn::diag` abort at `sdk+0x11ff34` (`StandardAllocator::Initialize`) in `nninitStartup`, SDK 15.3.2, called from `main+0x31370c` (Paper Mario: The Thousand-Year Door). → [StandardAllocators over 12 to 62 GB of address space](../reference/memory.md#standardallocators-over-12-to-62-gb-of-address-space)
+- a write to `0x20000ffff` in static initialization (`main+0x164990`), after `AllocateAddressRegion` returns `0x1803`; reducing the request to 3 GB still returns `0x1803` with the SDK's original large-region alignment. With only the outer ranges reduced, a later free crashes at `main+0x1668dc` after reading compressed file data as a slab-owner pointer. → [Separate virtual ranges in Snowdrop's page allocator](../reference/memory.md#separate-virtual-ranges-in-snowdrops-page-allocator)
 - `nn::diag` abort at `sdk+0x12dad0` in `nninitStartup`, called from `main+0x756310`. → [.NET NativeAOT malloc and the GC region range](../reference/memory.md#net-nativeaot-malloc-and-the-gc-region-range)
 - the game aborts in its static initialisers, before `nnMain`, when its own 62 GB address region fails (`main+0x37c9bd0`, `AllocateAddressRegion(&p, 0xf80000000)` at `main+0x37c9a5c`). → [A 62 GB page allocator behind `nu::VirtualAllocHook`](../reference/memory.md#a-62-gb-page-allocator-behind-nuvirtualallochook)
 - the game calls `exit(1)` in `nninitStartup` (Octopath Traveler 0: `main+0x12e74`); `nn::diag::Abort` from `main+0x352fbec` (Unreal's fatal error) after the first commit, `AllocateMemoryPages(0, 2 MB)` while the RHI starts (`main+0x353ff8c`, result 0xe03) (Pikmin 4); a crash in `NumberLineAllocator::Compaction` at boot (DRAGON QUEST XI S, region option 4224). → [Unreal's engine address region and its options](../reference/memory.md#unreals-engine-address-region-and-its-options)
@@ -142,6 +143,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - SIGSEGV in `android::Surface::getQueuedFrameNumber` under `nn::vi::GetLatestFrameNumber`. → [Layer frame numbers](../reference/services.md#layer-frame-numbers)
 - SIGSEGV in `vi::GetDefaultDisplayPanelIdentifier` (`main+0xa9a4`, before the window). → [Default display panel identifier](../reference/services.md#default-display-panel-identifier)
 - the first boot crashed in `oe::DisableAlbumImageTakenNotification`: a null read at `sdk+0x338070`, the first call in `nnMain` (`main+0x6b188c`); other `oe` calls that reach am, found in the imports before they were met. → [oe calls answered without an applet proxy](../reference/services.md#oe-calls-answered-without-an-applet-proxy)
+- a null read at `sdk+0x36ec60` in `oe::EnableAlbumImageTakenNotification`, called from `nnMain` (`main+0x286268`). → [oe calls answered without an applet proxy](../reference/services.md#oe-calls-answered-without-an-applet-proxy)
 - abort in `nn::sm::Initialize` under `nn::oe::FinishStartupLogo`, right after the first frames. → [The startup logo](../reference/services.md#the-startup-logo)
 - null read in `oe::InitializeCopyrightFrameBuffer`, right after the window setup. → [Copyright frame buffer overlay](../reference/services.md#copyright-frame-buffer-overlay)
 - SDK 20 calls that reach a service, each one boot; SDK 20's `OpenAudioRenderer` aborts on an ARUID of 0. → [Applet resource user ID, launch parameters and VR mode](../reference/services.md#applet-resource-user-id-launch-parameters-and-vr-mode)
@@ -302,6 +304,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - an abort in device setup when `nvnTexturePoolGetSize` / `nvnSamplerPoolGetSize` returned 0. → [Pool size getters, texture address and command memory size](../reference/nvn-objects.md#pool-size-getters-texture-address-and-command-memory-size)
 - red stub lines for NVN calls. → [NVN calls that were red stubs](../reference/nvn-objects.md#nvn-calls-that-were-red-stubs)
 - `nvnWindowGetPresentInterval`, `nvnCommandBufferSetCopyImageStride` missing. → [SDK 8 window and copy-stride entry points](../reference/nvn-objects.md#sdk-8-window-and-copy-stride-entry-points)
+- `[RED] texture bound to an image slot lacks the shader-write usage`, and compute image writes go to a dummy texture. → [Image bindings supply shader-write intent](../reference/nvn-objects.md#image-bindings-supply-shader-write-intent)
 
 ## Rendering
 
@@ -400,6 +403,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - the split found no shader in the game (0 shaders). → [GLSLC outputs holding several stages](../reference/shaders.md#glslc-outputs-holding-several-stages)
 - an NVN game's split finds none of its shipped shaders, and the ones it compiles at run time were not captured: the game's renderer is in NROs that call `glslcCompile`, and its whole RomFS is one zip. → [GLSLC outputs in a zip archive, and those NROs compile at run time](../reference/shaders.md#glslc-outputs-in-a-zip-archive-and-those-nros-compile-at-run-time)
 - 24 shaders before the control-length fix, 1,355 after. → [SPKG packages in KTGL LINKDATA](../reference/shaders.md#spkg-packages-in-ktgl-linkdata)
+- the split finds no shaders in a RomFS made of `.sdfdata`, `.sdftoc` and `.sdfver` files. → [Shader records in Snowdrop SDF data slices](../reference/shaders.md#shader-records-in-snowdrop-sdf-data-slices)
 - `[RED] ... PKZ version 0x10000: only 0x20000 is known` stops the shader split. → [PKZ version 1 archives](../reference/shaders.md#pkz-version-1-archives)
 - the split finds zero shaders and refuses the shader build, although the CPK's DAT archives contain the objects passed to NVN at boot. → [Size-delimited PlatinumGames shader objects](../reference/shaders.md#size-delimited-platinumgames-shader-objects)
 - the split finds no shaders, although `main` contains NVN control sections; it reports that no data record or code-address pair identifies their code areas. → [Needle shader containers in main and GFD archives](../reference/shaders.md#needle-shader-containers-in-main-and-gfd-archives)

docs/development.md

@@ -145,7 +145,7 @@ What the script does:
    A Vulkan game without SPIR-V files in its RomFS (Super Mario Galaxy) gets the same kind of run: the game boots in the Mac host with `DM_VK_CAPTURE` pointing at `<id>_build/extract/<version>/vk_capture` until every SPIR-V key in the profile's `vk_capture.tsv` is there (Galaxy: 3,089, about 7 s after launch), and gives up after 120 s.
    So does an OpenGL game with no GLSL source in its RomFS or `main` (Balatro: LÖVE puts its shaders together at run time, in Lua): `DM_GL_CAPTURE` at `<id>_build/extract/<version>/gl_capture` until every file in the profile's `gl_capture.tsv` is there (Balatro: 35, made before its first frame). Without the profile's tables (a provisional profile) a capture stops once nothing has been added for 20 s.
    The Vulkan and GL captures also run on a Mac whose GPU has no Metal 4, such as a virtual Mac (GitHub's `macos-26` runner). Those two layers draw on classic Metal queues of their own, so there the NVN layer starts only its classic part and drops every present, which a headless capture does not need (`nvn_mtl.m` `mtl_device_classic`; `DM_TEST_NO_METAL4=1` does the same on any Mac). The GLSLC capture boots an NVN game, which needs Metal 4.
-5. It splits out the shaders and checks them against the profile's tables (a Vulkan profile: it translates the RomFS's SPIR-V and the capture with `vk_shaders.py` instead; an OpenGL one: the GLSL sources and the capture with `gl_shaders.py`).
+5. It splits out the shaders from the game and every extracted DLC RomFS and checks them against the profile's tables (a Vulkan profile: it translates the RomFS's SPIR-V and the capture with `vk_shaders.py` instead; an OpenGL one: the GLSL sources and the capture with `gl_shaders.py`). Optional DLC that changes only the NVN shader union gets new tables in an external profile after the non-DLC shaders and facts have matched; registered module builds and fixes stay in use.
 6. It builds the iOS shader libraries, with the Mac's own Metal compiler and spirv-cross when they differ from those in `dmsh_chain.lock` (`dm/gfx/tools/own_compilers.py`; one line in the output names them). If a registered profile's `dmsh_chain.h` does not hold the version the source tree gives (a chain change left it behind), the script writes it again before any host is built, so it shows as changed in the checkout (`dmsh_build.py` would stop otherwise). Then the script makes the profile's `c1_full.tsv` and `cb_extent.tsv` from this build's Mac outputs when they are missing or name another chain. These two tables are not committed and git ignores them; `build_ipa.sh` stops without a current `c1_full.tsv`. Each line in the output names the file.
    It then compares the build with the game's previous one, the newest other `b-*/` in `<id>_build/shader/cache/` (`dm/gfx/tools/shader_regressions.py`, which reads the two `_b_build.json` reports). The translation chain is one for every game, so a translator change is tested only on the games its author owns; shaders that translated before and fail now are listed in `logs/shader_regressions.log` and stop the build ([GitHub issue #272](https://github.com/seedds/nxnative/issues/272)). Report them; `--allow-shader-regressions` (`ALLOW_SHADER_REGRESSIONS = True` in Spyder) packs the app without them.
 7. It calls `build_ipa.sh`. Xcode builds the app in the game's own build folder (`<id>_build/ipa/build/ios/pkg/dd`: `build_ipa.sh` sets `SYMROOT` and `OBJROOT`, which a custom build location in Xcode's settings does not override), so a repeat build compiles only what changed and two games built at once never write the same app. Then `dm/ios/tools/pack_ipa.py` packs it with the shader libraries, the add-on content and the RomFS straight from where they are. It reads and deflates files on all cores and writes at about the T7's write speed (300 MB/s): Breath of the Wild's 16 GB IPA takes 57 s, Info-ZIP's copy and zip took 2 min 50 s. A repeat pack reuses the previous IPA through `<ipa>.idx`, which records each packed file's size, times and inode: the unchanged leading entries stay in place in an APFS clone of the old IPA, and after the first changed entry, an unchanged file is copied compressed from it (both only when the source folder is on APFS; elsewhere every file is deflated again). `build_ipa.sh` passes the RomFS, the add-on content, the pipeline archive and the shader libraries before the app, so after a host fix only the app is written: Breath of the Wild repacks in 12-23 s instead of 105-163 s.

docs/reference/memory.md

@@ -156,6 +156,13 @@ The guest address space, heaps, pools, virtual address memory and the allocators
 - **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md) (2.0.2, SDK 19.3.11; one boot); [Paper Mario: The Thousand-Year Door](../../dm/games/0100ECD018EBE000/README.md): the allocator maps 1.5 GB at boot and 2.0 GB by the arrival in Rogueport. The `sdk` build is [Celeste](../../dm/games/01002B30028F6000/README.md)'s. Absolum's malloc StandardAllocator (40 GB, cut to 2.5 GB) is in [.NET NativeAOT malloc and the GC region range](#net-nativeaot-malloc-and-the-gc-region-range).
 - **Note:** on Pokémon Legends: Z-A, `QueryMemoryInfo` reporting more than 4 GB would take the game's big-memory path (heaps of allocatable − 4.27 GB instead of − 2.76 GB).
 
+### Separate virtual ranges in Snowdrop's page allocator
+
+- **Symptom:** a write to `0x20000ffff` in static initialization (`main+0x164990`), after `AllocateAddressRegion` returns `0x1803`; reducing the request to 3 GB still returns `0x1803` with the SDK's original large-region alignment. With only the outer ranges reduced, a later free crashes at `main+0x1668dc` after reading compressed file data as a slab-owner pointer.
+- **Cause:** the engine reserves three adjacent virtual ranges before committing pages in them. Its enclosing reservation exceeds the iPhone's available guest address space, and the engine does not check the failed reserve before deriving the page addresses. A smaller reservation also needs the SDK's node area and large-region alignment adapted: a separate metadata reservation comes first, and the 1 GB alignment wastes the remaining contiguous space. The large-allocation child has an independent maximum page count; reducing only the outer range lets it allocate across the boundary used by `free` to distinguish allocation classes.
+- **Fix:** the build-gated game cuts reduce the range sizes, their combined reservation and the large allocator's maximum page count together. The fixed metadata reserve remains large enough, and page allocation, commits and frees retain their behavior. The SDK build gets the existing [node-area cut](#the-sdks-address-region-node-area) and [32 MB alignment](#1-gb-alignment-of-regions-of-1-gb-and-more). Per-range occupancy and the original growth fallbacks remain **open** until sustained gameplay validates the smaller ranges; see the game note for the layout, cut sites and current validation.
+- **Found on:** [Mario + Rabbids Sparks of Hope](../../dm/games/0100317013770000/README.md) (1.6.2225577, SDK 13.3.5).
+
 ### .NET NativeAOT malloc and the GC region range
 
 - **Symptom:** `nn::diag` abort at `sdk+0x12dad0` in `nninitStartup`, called from `main+0x756310`.

docs/reference/nvn-objects.md

@@ -424,6 +424,13 @@ Games read this count as a position in their own command memory and size their h
 - `nvnTexturePoolRegisterImage` writes the same table as `RegisterTexture` (image handles are pool indices).
 - **Found on:** [Hyrule Warriors: Age of Calamity](../../dm/games/01002B00111A2000/README.md).
 
+### Image bindings supply shader-write intent
+
+- **Symptom:** `[RED] texture bound to an image slot lacks the shader-write usage`, and compute image writes go to a dummy texture.
+- **Cause:** Metal usage came only from the texture builder's `IMAGE` flag. A render texture may instead have only `COMPRESSIBLE`, then its descriptor be used in a shader image slot. Snowdrop writes these descriptors with `RegisterTexture`, not `RegisterImage`. Its depth-conversion compute shader `cf5ba107e29b51f4` gathers from a depth texture and unconditionally writes four pixels per invocation to a 256×256 R16F texture whose flags are `0x8`. Its Metal texture had read / render-target usage (`0x5`), so the image binding was rejected.
+- **Fix:** before a draw or dispatch uses a writable image, `image_prepare` resolves it with image intent and ensures its backing allows shader writes. `RegisterImage` also keeps this intent in its host snapshot; neither path changes the guest texture. A texture already rendered into is upgraded through `tex_upgrade`: end its render pass first, copy every mip / slice, replace all shared serials and invalidate cached sampled views. Load-only image slots accept read-only storage without upgrading, and ordinary render targets retain their original usage. `t_image_usage.m` exercises both registration entry points, a late registration while a render pass is open, the clear and other mip surviving the upgrade, and a GPU image write read through a sampled alias.
+- **Found on:** [Mario + Rabbids Sparks of Hope](../../dm/games/0100317013770000/README.md) (SDK 13.3.5, Snowdrop).
+
 ### Plural texture binds, builder-registered samplers, strided texel writes, device item 0x67
 
 - `nvnCommandBufferBindTextures` (the plural of `BindTexture`), `nvnSamplerPoolRegisterSamplerBuilder` (a sampler registered from its builder), `nvnTextureWriteTexelsStrided` (`WriteTexels` with row and image strides) and device item 0x67, which the game reads before every program it loads and uses as the alignment of its copies of the shader (answered 256, not measured).

docs/reference/services.md

@@ -356,7 +356,9 @@ Most service calls can be found before they crash, before the first boot or the
 ### oe calls answered without an applet proxy
 
 - **Symptom:** the first boot crashed in `oe::DisableAlbumImageTakenNotification`: a null read at `sdk+0x338070`, the first call in `nnMain` (`main+0x6b188c`); other `oe` calls that reach am, found in the imports before they were met.
+- **Symptom:** a null read at `sdk+0x36ec60` in `oe::EnableAlbumImageTakenNotification`, called from `nnMain` (`main+0x286268`).
 - **Cause:** the SDK asks am through the applet proxy the host never made.
+- **Fix:** `EnableAlbumImageTakenNotification` uses the same host implementation as the disable call: no system-album notifications are produced. Found on [Mario + Rabbids Sparks of Hope](../../dm/games/0100317013770000/README.md) (SDK 13.3.5); the next boot reaches the NVN frame loop.
 - **Fix:** hooked (`hle_oe.c`): `oe::DisableAlbumImageTakenNotification`; `oe::IsInStoreDemoModeEnabled` (no); `oe::GetDefaultDisplayResolution` and its change event; `oe::SendServerMaintenanceOverlayNotification`; `oe::SetUserInactivityDetectionTimeExtendedUnsafe`; `oe::GetPerformanceConfiguration`; `oe::SetAlbumImageOrientation`; `oe::Begin` / `EndAutoSleepDisabledSection` and `IsAutoSleepDisabled` (counted); `oe::ExecuteProgram` (logged red: the host runs one program); `oe::GetUserInactivityDetectionTimeExtensionMode` (0, as `IsUserInactivityDetectionTimeExtended` reports no). The ones with a symptom of their own follow: [The startup logo](#the-startup-logo), [Copyright frame buffer overlay](#copyright-frame-buffer-overlay), [Applet resource user ID, launch parameters and VR mode](#applet-resource-user-id-launch-parameters-and-vr-mode), [CPU overclock calls without an apm session](#cpu-overclock-calls-without-an-apm-session), [Focus handling mode decides whether focus changes are notified](#focus-handling-mode-decides-whether-focus-changes-are-notified).
 - **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md) (SDK 11: `GetDefaultDisplayResolution`, `SendServerMaintenanceOverlayNotification`); [13 Sentinels: Aegis Rim](../../dm/games/01003FC01670C000/README.md) (SDK 11: `SetUserInactivityDetectionTimeExtendedUnsafe`); [Super Mario 3D World + Bowser's Fury](../../dm/games/010028600EBDA000/README.md) (SDK 10: `GetPerformanceConfiguration`, `SetAlbumImageOrientation`); [DRAGON QUEST XI S](../../dm/games/01006C300E9F0000/README.md) (SDK 7.3.2: the auto-sleep calls); [Octopath Traveler 0](../../dm/games/01005270232F2000/README.md) (`oe::ExecuteProgram`); [The Great Ace Attorney Chronicles](../../dm/games/010036E00FB20000/README.md) (SDK 10.4.1, Capcom's MT Framework: `DisableAlbumImageTakenNotification`, then `IsInStoreDemoModeEnabled` from the import list); [Donkey Kong Country: Tropical Freeze](../../dm/games/0100C1F0051B6000/README.md) (SDK 4.4.0: `GetUserInactivityDetectionTimeExtensionMode`, a null read at `sdk+0x15b73c` when the first level's cinematic camera starts, `CScreenDimmingUtil::DisableScreenDimming`). After all of its fixes The Great Ace Attorney Chronicles runs its frame loop at 30 fps; the black screen it showed was its textures, never uploaded (CPU-written textures, the profile's `tex_flush_upload` and `tex_dirty_upload`: [Mapped-range flushes taken as uploads (`tex_flush_upload`)](nvn-objects.md#mapped-range-flushes-taken-as-uploads-tex_flush_upload), [CPU writes with no flush, caught by page protection (`tex_dirty_upload`)](nvn-objects.md#cpu-writes-with-no-flush-caught-by-page-protection-tex_dirty_upload)).
 

docs/reference/shaders.md

@@ -116,6 +116,14 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b
 - **Reader:** `kpka_jobs` maps the PAKs, takes the newest copy of each path, peeks at 4 decompressed bytes and inflates only SDF entries (`sdf_shaders`).
 - **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md): 63,111 shaders (30,233 vertex, 21,475 pixel, 11,271 compute, 53 geometry, 33 hull, 46 domain) and 102,345 pairs, in 5 s with 12 jobs.
 
+### Shader records in Snowdrop SDF data slices
+
+- **Symptom:** the split finds no shaders in a RomFS made of `.sdfdata`, `.sdftoc` and `.sdfver` files.
+- **Cause:** Snowdrop stores the shaders in individually compressed records inside its `BERG` data slices. These SDF archives are a different format from RE Engine's shader definitions above. A shader record has a 20-byte header `{u32 0xabcd00c6, flags, 0, raw length, stored length}`, then one zstd frame; reflection data precedes the NVN sections inside it.
+- **Fix:** `shader_split.py` recognizes `BERG` by magic and reads the length-delimited compressed records (`snowdrop_shaders`). Every control section is followed immediately by its code, with lengths from control +0x20 and +0x704; they are byte-packed, so a control need not be aligned. A single record can hold more than one shader. Each section must pass the control / code / stage checks. Programs hold one stage, and the vertex shaders take the [single-stage cover](#single-stage-programs-paired-with-a-cover). Scanning for the record header avoids mistaking a coincidental zstd magic inside compressed data for another asset.
+- **DLC:** additional shader records ship in the DLC's data slices. The NSP script includes each extracted DLC RomFS as an extra shader root, without an ExeFS module, and fingerprints its index and extraction timestamps for split reuse. Vulkan and OpenGL builds include the DLC roots too. When installing or removing optional DLC changes an NVN shader union, the script uses an external profile with new shader tables only after every non-DLC shader row and fact matches the registered tables; package indices may change. The registered module builds, settings and build-gated cuts are preserved. A cached proof of this comparison permits repeat builds to reuse the split.
+- **Found on:** [Mario + Rabbids Sparks of Hope](../../dm/games/0100317013770000/README.md) (Snowdrop, SDF version 22).
+
 ### SPK shader package in an MT Framework ARC
 
 - **How it is stored:** the engine's one shader package, `sc\NX_64\root` (`SPK\0`, type hash `0x02358e1a`) in `nativeNXx64/sa/NX_64/root.arc` ([01](../agents/01-extract.md)). After the package's tables every shader is a code area followed directly by its control section, back to back to the end of the file: 230 vertex shaders, then 230 pixel shaders.

scripts/build_ipa_from_nsp.py

@@ -708,25 +708,27 @@ def derived_key(profile_dir):
     return profile_dir.with_name(profile_dir.name + '.derived.key')
 
 
-def derive_profile(g, pdir, ver, d, mods, app=None):
+def derive_profile(g, pdir, ver, d, mods, app=None, shader_only=False):
     """The game's profile for a version it was not brought up on (the folder has no update, another update, files of other builds, or a registered build under another title ID: app),
     in pdir (outside the repository, where a provisional profile goes): its keys as they are, with this version's build row, package and shader versions, and app's title ID and
     main.npdm's program ID when app is not the profile's title. The shader tables, the chain header and the run-time capture list are
     left out: the build makes them again from this version's own files, as for a provisional profile. The hand-written tables (the prewarm list, shader_facts_manual.tsv) are copied:
     they name shaders by content key. A folder derived from the same inputs (the files copied and the game.tsv written, kept in <pdir>.derived.key) is kept as it is, with what later
-    stages wrote into it (mem_used, the heap check, the split's tables), so their own reuse holds. Returns the version's display version"""
-    src = sorted(p for p in g.dir.iterdir() if p.is_file() and p.name not in DERIVED_SKIP)
+    stages wrote into it (mem_used, the heap check, the split's tables), so their own reuse holds. shader_only preserves all registered game.tsv rows and links its existing cuts:
+    only the shader tables change with the installed DLC. Returns the version's display version"""
+    src = sorted(p for p in g.dir.iterdir() if p.is_file() and (p.name not in DERIVED_SKIP or shader_only and p.name == 'patches.py'))
     T, disp = '\t', nacp.version(d / 'control.nacp') or '?'
     row = T.join(['build', ver, disp] + [f'{m}={build_id(d / "ExeFS" / m)}' for m in mods if (d / 'ExeFS' / m).exists()])
     out, put = [f'# Derived from game profile {g.dir} by scripts/build_ipa_from_nsp.py for {ver} ({disp}), a version that profile was not brought up on:',
                 '# its keys as they are, this version\'s build row, and shader tables made from this version\'s own files'], False
+    if shader_only: out = [f'# Derived from game profile {g.dir} for the installed DLC shader content; registered module builds and settings are preserved.']
     ids = {}
     if app is not None and app != g.title_id:   # the host reports the program ID and answers save data queries for it (hle_svc.c, hle_fs.c)
         ids = {'title_id': f'title_id{T}{app:016X}', 'program_id': f'program_id{T}{(pg.npdm(d / "ExeFS/main.npdm")[1] or app):016X}'}
         out.insert(1, f'# under title ID {app:016X}, not the profile\'s {g.title_id:016X}')
     for line in (g.dir / 'game.tsv').read_text(encoding='utf-8').splitlines():
         if line.split(T, 1)[0] in ids: out.append(ids[line.split(T, 1)[0]]); continue
-        if line.split(T, 1)[0] in ('build', 'package', 'shader_versions'):
+        if not shader_only and line.split(T, 1)[0] in ('build', 'package', 'shader_versions'):
             if not put: out += [row, f'package{T}{ver}', f'shader_versions{T}{ver}']; put = True
             continue
         out.append(line)
@@ -743,7 +745,10 @@ def derive_profile(g, pdir, ver, d, mods, app=None):
     atomic_text(pdir / 'game.tsv', text)
     for p in pdir.iterdir():
         if p.name != 'game.tsv': rmtree(p)
-    for p in src: shutil.copyfile(p, pdir / p.name)
+    for p in src:
+        # Registered cuts stay the same file: patch_nso.py deduplicates real paths, so an external profile must not register copied claims for the same builds.
+        if shader_only and p.name == 'patches.py': (pdir / p.name).symlink_to(p.resolve())
+        else: shutil.copyfile(p, pdir / p.name)
     kf.write_text(f'{key}\n')
     return disp
 
@@ -1339,8 +1344,39 @@ def data_rows(b):   # a table's data rows (without # comment lines: the header c
     return [l for l in b.splitlines() if not l.startswith(b'#')]
 
 
-def split_shaders(py, repo, roms, mains, pkg, env, logs, g, provisional=False, oodle=None):
-    """The shader split into pkg, checked against the profile's tables. The package is reused when the tools, the profile and each version's main, RomFS (paths and sizes) and run-time capture
+def dlc_only_shader_change(g, tables):
+    """A different DLC shader union, with the registered non-DLC shaders and facts unchanged. Package indices may be re-numbered when an add-on is removed.
+    Source ownership is the split's /aoc/<index>/ prefix; a module or base/update shader difference never qualifies for this shader-only derivation."""
+    def rows(path):
+        out = {}
+        for line in path.read_text().splitlines():
+            if not line.strip() or line.startswith(('#', 'package\t')): continue
+            c = line.split('\t')
+            if len(c) != 5 or not c[0].endswith('.bin') or c[3] in out: return None
+            out[c[3]] = (c[0], c[2], c[4])
+        return out
+    try:
+        old, new = rows(g.table('shader_map.tsv')), rows(tables / 'shader_map.tsv')
+        if old is None or new is None: return False
+        old_dlc = {k: v for k, v in old.items() if v[2].startswith('/aoc/')}
+        new_dlc = {k: v for k, v in new.items() if v[2].startswith('/aoc/')}
+        if old_dlc == new_dlc or {k: v for k, v in old.items() if k not in old_dlc} != {k: v for k, v in new.items() if k not in new_dlc}: return False
+        dlc = old_dlc.keys() | new_dlc.keys()
+        def facts(path):
+            out, seen = [], set()
+            for line in path.read_text().splitlines():
+                if not line.strip() or line.startswith('#'): continue
+                c = line.split('\t', 1)
+                if len(c) != 2 or c[0] in seen: raise ValueError('invalid facts')
+                seen.add(c[0])
+                if c[0] not in dlc: out.append(line)
+            return out
+        return facts(g.table('shader_facts.tsv')) == facts(tables / 'shader_facts.tsv')
+    except (OSError, UnicodeError, ValueError): return False
+
+
+def split_shaders(py, repo, roms, mains, pkg, env, logs, g, provisional=False, oodle=None, dlc_roms=(), derive_dlc=None):
+    """The shader split into pkg, checked against the profile's tables. The package is reused when the tools, the profile, the DLC RomFS roots and each version's main, RomFS (paths and sizes) and run-time capture
     are as at the last successful split. Returns the package's content key, the shader build's input (None for damaged files, which never reuse them): a split that runs
     again because a reader of another engine changed writes the same package, and the iOS shaders stay reused.
     oodle: for decoder-backed packs, builds or reuses the Mac host and returns DM_OODLE_HOST / DM_OODLE_MODULES; their current files are split inputs too"""
@@ -1351,22 +1387,36 @@ def split_shaders(py, repo, roms, mains, pkg, env, logs, g, provisional=False, o
         (('files', [m]), ('stat' if provisional else 'tree', r),
          ('stat', r.parent / 'glslc_capture') if provisional else ('files', sorted((r.parent / 'glslc_capture').glob('*.glslc'))),
          ('stat', r.parent / 'ue_shaders') if provisional else ('files', sorted((r.parent / 'ue_shaders').glob('*.ushaderbytecode'))))]
+    inputs += [item for r in dlc_roms for item in (('dlc', Path(r).name), ('stat', r))]   # DLC has no registered module build to pin its contents: include its index and re-extraction times, even when file sizes stay the same
     if oodle:   # a host or rewriter fix can change what the game's decoder produces
         PROG.set(step='building the Mac host for the game\'s Oodle decoder'); env = dict(env, **oodle())
         inputs += [('files', [env['DM_OODLE_HOST']]), ('stat', Path(env['DM_OODLE_MODULES']))]
     key = None if ALLOW_DAMAGED else fingerprint(*inputs, ('files', profile_files(g)))
+    dlckf, dlckey = pkg.with_name(pkg.name + '.dlc.key'), None
+    def promote_dlc():
+        nonlocal g, provisional, env
+        assert derive_dlc is not None
+        g, provisional = derive_dlc(), True; inputs[1] = True
+        env = dict(env, DM_GAMES_DIR=os.pathsep.join([str(g.dir.parent), env.get('DM_GAMES_DIR', '')]), DM_GAMES_FIRST='1')
+    if not provisional and derive_dlc:
+        dlckey = fingerprint(key, ('files', [Path(__file__)])) if key else None   # proof pins the comparison code as well as the registered tables and every split input
+        if dlckey and reused(dlckf, dlckey) is not None:
+            promote_dlc()
+            key = fingerprint(*inputs, ('files', profile_files(g)))
     have_pairs = g.table('pso_pairs.tsv').is_file() or pkg.with_name('pso_pairs.tsv').is_file()
     if key and pkg.is_dir() and have_pairs and (note := reused(kf, key)) is not None:
         PROG.end(f'{note}, reused (inputs unchanged)'); return package_key(pkg)
     PROG.set(step='splitting')
     tmp = pkg.with_name(pkg.name + '.partial'); rmtree(tmp); (tmp / 'tables').mkdir(parents=True)
-    run([py, repo / 'dm/gfx/tools/shader_split.py', ','.join(map(str, roms)), ','.join(map(str, mains)), tmp / 'pkg'] + [tmp / 'tables' / n for n in names + ('pso_pairs.tsv',)],
+    run([py, repo / 'dm/gfx/tools/shader_split.py', ','.join(map(str, roms)), ','.join(map(str, mains)), tmp / 'pkg'] + [tmp / 'tables' / n for n in names + ('pso_pairs.tsv',)] + ([','.join(map(str, dlc_roms))] if dlc_roms else []),
         logs / 'shader_split.log', env=env)
     m = re.search(r'shaders: (\d+) \(', (logs / 'shader_split.log').read_text(errors='replace'))
     if m and m.group(1) == '0':   # dmsh_build.py would make no library, and the iOS shader step stop with no word of why
         raise BuildError('the shader split found no shaders in main or the RomFS: the game keeps them in a form shader_split.py does not read; locate them (docs/agents) and add a reader')
     # pso_pairs.tsv (agl BAHS program variations): checked like the others where the profile has one; else kept beside the
     # package, where dmsh_build.py reads it too, so padding and the pipeline prebuild use the same pairs
+    if not provisional and derive_dlc and dlc_only_shader_change(g, tmp / 'tables'):
+        promote_dlc()
     if provisional or g.table('pso_pairs.tsv').is_file(): names += ('pso_pairs.tsv',)
     else: shutil.copyfile(tmp / 'tables' / 'pso_pairs.tsv', pkg.with_name('pso_pairs.tsv'))
     if provisional:   # a provisional profile takes the tables as they come
@@ -1380,6 +1430,8 @@ def split_shaders(py, repo, roms, mains, pkg, env, logs, g, provisional=False, o
     if key:   # an own profile's tables are outputs too; the next split checks the profile after they were written
         key = fingerprint(*inputs, ('files', profile_files(g)))
         kf.write_text(f'{key}\n{note}\n')
+    if provisional and dlckey: dlckf.write_text(f'{dlckey}\nDLC-only shader content\n')
+    else: rmtree(dlckf)
     PROG.end(note)
     return key and package_key(pkg)
 
@@ -1776,7 +1828,9 @@ def report_profile_status(g, reg, pkg_ver, derived, provisional, artifact='IPA')
         known = {(m, b) for row in reg.builds().values() for m, b in row.items() if m != 'display'}
         other = [m for m, b in g.builds()[pkg_ver].items() if m != 'display' and (m, b) not in known]
         registered_builds = ", ".join(f'{v} {b["display"]}' for v, b in reg.builds().items())
-        if g.title_id != reg.title_id and not other and pkg_ver in reg.builds():   # a registered build under another title ID (same_build): every fix applies
+        if g.title_id == reg.title_id and g.builds() == reg.builds():
+            say(f'this {artifact} uses the registered {pkg_ver} ({disp}) builds and every fix for them, with shader tables derived from the installed DLC')
+        elif g.title_id != reg.title_id and not other and pkg_ver in reg.builds():   # a registered build under another title ID (same_build): every fix applies
             say(f'this {artifact} is game profile {reg.id}\'s {pkg_ver} ({disp}) under title ID {g.title_id:016X}: its settings and every fix for its builds, with shaders split from this title\'s own files')
         else: say(f'[RED] this {artifact} uses {pkg_ver} ({disp}), which game profile {reg.id} was not brought up on ({registered_builds}): '
             'its settings, with this version\'s own shaders' + (f'; {", ".join(other)} {"is a build" if len(other) == 1 else "are builds"} the profile does not register, so the fixes made for '
@@ -1931,6 +1985,7 @@ def main(args=None):
         if on: say(src[k][1])
     PROG.add(glslc + vkcap + glcap + uedec + (gfx == 'nvn'))
     aoc = extract_aoc(nx, keys, repo, work, app, fam, logs)
+    dlc_roms = sorted(p for p in aoc.iterdir() if p.is_dir()) if aoc else []
     base = next((dirs[v] for v in versions if picks[v][1] is None), dirs[versions[0]])   # the folder the base game was extracted to alone
     # the app name and bundle ID both come from the base NSP: the game name in its control data, and a bundle ID from its title ID; a setting / the command line overrides them
     app_name = args.app_name or APP_NAME or nacp.name(base / 'control.nacp')
@@ -1964,15 +2019,26 @@ def main(args=None):
     sh = work / 'shader'
     if gfx == 'vulkan':   # the packed version's merged RomFS holds the SPIR-V files the game loads; a capture holds those it generates
         cap = capture_shaders(repo, game, patched, work, env, logs, g, 'vk') if vkcap else None
-        ios = build_vk_shaders(py, repo, [game / 'RomFS'] + ([cap] if cap else []), sh, env, logs, g, reuse=not ALLOW_DAMAGED, mac=mac, provisional=bool(provisional or derived))
+        ios = build_vk_shaders(py, repo, [game / 'RomFS'] + dlc_roms + ([cap] if cap else []), sh, env, logs, g, reuse=not ALLOW_DAMAGED, mac=mac, provisional=bool(provisional or derived))
     elif gfx == 'gl':   # the GLSL sources in the packed version's archives and in its main (the game compiles them at run time; the host finds each by the hash of its text), and a capture of those it puts together
         cap = capture_shaders(repo, game, patched, work, env, logs, g, 'gl') if glcap else None
-        ios = build_gl_shaders(py, repo, [game / 'RomFS', game / 'ExeFS/main'] + ([cap] if cap else []), sh, env, logs, g, reuse=not ALLOW_DAMAGED, mac=mac, provisional=bool(provisional or derived))
+        ios = build_gl_shaders(py, repo, [game / 'RomFS', game / 'ExeFS/main'] + dlc_roms + ([cap] if cap else []), sh, env, logs, g, reuse=not ALLOW_DAMAGED, mac=mac, provisional=bool(provisional or derived))
     else:
         # Trinity packs and legacy Unreal paks: the split decodes them with the packed version's own decoder (Oodle, LZ5), in the Mac host built for the profile
         oodle = (lambda: {'DM_OODLE_HOST': str(build_mac_host(repo, work / 'tools/dmhost', env, logs, g)), 'DM_OODLE_MODULES': str(patched)}) if src['oodle_packs'][0] or src['ue_legacy'][0] else None
         own = bool(provisional or derived)   # the shader tables come from this build's split (the profile's are checked otherwise)
-        pkey = split_shaders(py, repo, [dirs[v] / 'RomFS' for v in sv], [dirs[v] / 'ExeFS/main' for v in sv], sh / 'pkg', env, logs, g, own, oodle)
+        def derive_installed_dlc():   # called only after all registered non-DLC shader rows and facts matched
+            nonlocal g, derived
+            assert g is not None
+            derived = 'the installed DLC shader content differs from the registered tables'
+            pdir = Path(os.path.expanduser(args.profile_out or PROFILE_OUT)).resolve() if (args.profile_out or PROFILE_OUT) else work / 'profile'
+            derive_profile(g, pdir / g.id, pkg_ver, game, g.words('modules'), app, shader_only=True)
+            os.environ['DM_GAMES_DIR'] = env['DM_GAMES_DIR'] = os.pathsep.join([str(pdir)] + [d for d in os.environ.get('DM_GAMES_DIR', '').split(os.pathsep) if d and d != str(pdir)])
+            os.environ['DM_GAMES_FIRST'] = env['DM_GAMES_FIRST'] = '1'
+            g = gamemod.load(g.id)
+            say(f'{derived}: using a shader-only derived profile at {g.dir}; registered module fixes are preserved')
+            return g
+        pkey = split_shaders(py, repo, [dirs[v] / 'RomFS' for v in sv], [dirs[v] / 'ExeFS/main' for v in sv], sh / 'pkg', env, logs, g, own, oodle, dlc_roms, derive_installed_dlc if not own else None)
         ios = build_shaders(py, repo, sh, xlate, env, logs, g, pkey, mac)
     icon = next((d / 'icon.jpg' for d in (game, base) if (d / 'icon.jpg').exists()), None)   # the game icon nx_extract saved from the control data (English, icon.jpg); the packed version's first, then the base game's
     if mac_pack:   # a Mac app: the Mac host, then the app packed from it and the build's files