Port Xenoblade Chronicles 3 2.2.1

Add Monolith archive and shader extraction, a build-gated compact memory layout, vertex storage writes under rasterizer discard, primitive ordering and interleaved geometry feedback. Correct native decoder/audio wait accounting and the oe calls used after movies. Validated opening combat and movement on Mac with real audio, the full Mac/iOS shader build, GPU regressions, and the unsigned IPA archive. iOS runtime, later areas and DLC gameplay remain untested.
2026-10-27 07:58:05 +08:00
1 parent 6f58a7d4f6
commit 9bdef355ab
@@ -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-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.
+- 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-six 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 |
@@ -155,6 +155,11 @@ Under each game: the IPA's size, and "60 fps with 1.2.0" where the game runs at
 <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>
+<tr>
+<td align="center" valign="top" width="33%"><img src="docs/images/xenoblade-chronicles-3.jpg" width="256" alt="Xenoblade Chronicles 3"><br>Xenoblade Chronicles 3<br><sub>~19.2 GB</sub></td>
+<td></td>
+<td></td>
+</tr>
 </table>

 ## Pack an IPA
@@ -0,0 +1,168 @@
+# Xenoblade Chronicles 3
+
+Base title ID `010074F013262000`, update 2.2.1 (`v720896`), SDK 12.3.9,
+Monolith's engine, NVN. Modules: `main`, `subsdk0` (NintendoSDK_movie), `sdk`.
+The build IDs are in [game.tsv](game.tsv).
+
+## Current status
+
+Opening gameplay tested on an Apple-silicon Mac (2026-10-27). The compact layout
+renders the title, terrain, characters, lighting and effects. New-game setup,
+opening movies, the first enemy defeat, destination tutorials, walking and the
+engagement-area collision boundary work. All 27,055 base/update shaders compile
+for Mac and iOS. The unsigned IPA is built and its archive checks pass.
+
+iOS runtime testing, DLC gameplay and later-game allocation peaks remain open.
+The supplied four DLC packs are included; their additional shaders have not been
+integrated into the base/update shader census, so DLC rendering is unvalidated.
+
+## Profile and memory
+
+- The NPDM declares a 32 MB system resource. Memory is committed through address
+  regions, without `SetHeapSize`. The profile uses a 5 GB guest block, including
+  a 4.375 GB alias region, and the Switch application pool of `0xCD500000` bytes.
+- `mem_map_limit on`: startup attempts 1 MB physical commits until allocation
+  fails, then frees its probe. An unlimited host lets it map all 60 GB; enforcing
+  the hardware pool makes that probe stop at the expected physical limit.
+- **Address-space adaptation:** `nninitStartup` calls `main+0x128e58c`, reserving 60 GB.
+  Its first arena reserves 4 GB + 128 KB (`main+0x1241860`), so reducing only the
+  containing region cannot work. [patches.py](patches.py) changes that arena's
+  bump increment, subtraction and capacity together, to 512 MB + 128 KB, and its
+  containing region to 4 GB. All changes are gated by the 2.2.1 main build ID.
+  The SDK build takes the existing node-area and large-region alignment cuts.
+- A 40-second diagnostic run with the original arenas had 2,092.4 MB committed:
+  about 203 MB in the first arena, about 1.8 GB in subsequent windows, and a 64 MB
+  allocation near the end of the reservation. Original-layout opening battle
+  (`boot29wide`) commits 225.44 MB in the first arena. That leaves about 286 MB
+  of capacity in the compact first arena; later-game peaks remain unvalidated.
+- Compact runs stayed at the title for over ten minutes without an allocator
+  failure. This validates the title allocation pattern, not later gameplay peaks.
+- `pad_layout full` supplies both sticks and the D-pad.
+
+## Files and shaders
+
+- The merged RomFS is 14.15 GB: 238 WebM movies, nine Wwise PCKs, and `bf3.arh` /
+  `bf3.ard`. The ARH1 index has 76,445 files; the ARD is 9.05 GB. Archive entries
+  and model/map streams use XBC1 (zlib and zstd).
+- The full base/update split finds 27,055 shader keys: 8,201 vertex, 18,533 pixel,
+  308 geometry and 13 compute, with 17,648 variation pairs. `main` supplies 204;
+  archive assets supply 26,851. All compile for Mac and iOS with feedback entries
+  (chain `7fbca6dd477d`, profile inputs `2342f43192c3`, after integration with main).
+- **Resolved, extraction:** six boot shader stages live in one standalone SLCT
+  at `main+0x14e01c0`, without an HCPS header. Control and code offsets are relative
+  to the SLCT itself. Its control sections include assembly text and are aligned
+  to eight bytes. The first pair's keys, `58e2e00d24c6f3ea` and
+  `5d36d4144f49ec9a`, match runtime dumps exactly.
+- Twelve more keys captured during startup match the extracted archive programs.
+  Equal code can have different compiler identity bytes at control +0x7d0; all
+  such keys are retained, with their original control bytes.
+- **Resolved, extraction:** `/effect/cmn_assets.wipac` contains complete GLSLC
+  outputs alongside HCPS programs. Reading them adds 914 base-game keys. Captured
+  `6babc936d58fcd8f` (vertex), `7a7eb30dfb6187b5` (geometry) and
+  `b540e3f8e28bd41e` (pixel) match these outputs exactly.
+- **Resolved, translation:** vertex `4b98320a0f512efe` reads `SubgroupLtMask`
+  through a component pointer; the general one-lane lowering now covers that form.
+- **Resolved, title rendering:** the black foreground model had valid G-buffer
+  colour and normals, but the separate IMAGE stencil descriptor read empty
+  storage. The general sharing fix restores lighting (`boot17shared`, frame 1200).
+- **Resolved, opening deadlock:** the output thread held the host decoder mutex
+  while signalling the higher-priority input thread; the handoff waited for that
+  input thread, which was blocked on the same native mutex. Annotating contended
+  native waits restores playback (`boot18movie`, `mv01005100.webm`, 140 seconds).
+- **Current, native audio setup:** a later real-audio run (`boot25full`) stopped
+  presenting while `AKEventManager` joined the Apple audio-unit setup helper.
+  A native join must release the guest core, just like a contended codec lock;
+  `dm_host_thread_join` now preserves that scheduler accounting. Gated join,
+  nesting and audio-cycle regressions pass. The sample does not establish that
+  Apple audio-unit construction itself would remain blocked indefinitely.
+  The corrected compact run continues for over ten minutes with real output
+  `#0`, without the diagnostic watchdog or reported audio underruns. It does
+  not reopen the output, so the earlier construction stall is not isolated.
+- **Resolved, nested effect compression:** those four geometry/pixel keys, and
+  vertex `d021f6a3763a9327` / pixel `1c6c6a015f6f3acf`, match GLSLC programs inside
+  `/effect/fld/ma01a.wipac` after its additional `cmp\0zlib` layer is decoded.
+  The expanded base/update census is 27,055 keys (8,201 vertex, 18,533 pixel,
+  308 geometry, 13 compute), with 17,648 variation pairs.
+- **Resolved, vertex work under discard:** `cf66cef403c57845` and
+  `36710e69a07e849d` write bone palettes through storage buffers. Running their
+  vertex work restores geometry (`boot21discard`), initially with large distorted
+  surfaces. `boot22full` reproduced this with the complete shader set. A native
+  post-vertex capture of the first `4f6b050575d4de0b` draw gives 2,080 finite
+  vertices inside the clip volume. The remaining corruption was the primitive
+  ordering failure below.
+- **Resolved, feedback:** effect geometry `c01e3ade8e15a842` captures eleven
+  interleaved generic outputs (172 bytes per vertex). The general feedback path
+  now decodes these byte component counts, captures geometry output and saves
+  GPU draw counts. Native regression checks pass; `boot23feedback` runs this path
+  without missing feedback calls or capture-entry errors.
+- **Resolved, exploding meshes:** the game orders bone-palette work with
+  `ORDER_PRIMITIVES` barriers (`0x21`). The host only split render encoders for
+  fragment barriers after storage writes. Enforcing primitive ordering removes
+  the exploding model parts (`boot24barrier`).
+- **Resolved, terrain allocation:** `boot27draw` frame 4000 replays every
+  recorded program bind/draw, with no map shader or rejected draw. CPU probes
+  (`boot28cpu`) find the map owner rejected before visibility because its core
+  buffer at `+0xd18` is null. The 390 MB allocation (`boot30alloc`) needs a
+  512 MB virtual window. `boot32region` at `main+0x12429a4` records bump
+  `0xe46c0000` → `0x1046c0000`, beyond the compact 4 GB container, before any
+  physical mapping. The heap's byte limit is disabled (`heap31.bin`). Reducing
+  the first arena's reservation from 1 GB to 512 MB frees that window; this
+  retains the Switch physical-memory pool and all map work. The original-layout
+  battle's reservation sequence predicts 159.56 MB of remaining outer space;
+  corrected compact gameplay (`boot33compact512`) confirms bump `0xf6070000`,
+  a 512 MB + 128 KB first-arena capacity and about 225 MB committed there.
+
+## Mac validation and remaining work
+
+`boot33compact512` uses the complete shader set, the corrected compact modules
+and real audio for over ten minutes. It defeats the first enemy, advances through
+the destination and tips tutorials, then walks across Everblight Plain to the
+engagement boundary.
+`run/compact33_terrain.png`, `compact33_move.png` and `compact33_walk.png` retain
+these milestones. Alias mappings settle near 2,653–2,655 MB during this sequence.
+Later areas, saves/reloads and DLC have not been exercised.
+
+The unsigned iOS app compiles with signing disabled. The IPA is 20,629,976,490
+bytes (19.21 GiB), contains all 27,055 base/update shader libraries and four
+mounted DLC directories, and passes every ZIP CRC. It has no code signature or
+provisioning profile. No device installation or runtime test was performed.
+The delivered IPA is the original validated `bab889ccd129` build; integration
+with main rebuilt both platforms' libraries but did not replace that artifact.
+
+After integrating main at `6f58a7d`, `boot34main` repeats the first enemy defeat
+and destination tutorial with the combined host and shader chain. The full Metal
+suite, all 48 Python test files, all 178 iOS source compile checks and the pinned
+lint/type ratchet pass. A fresh launch after shader metadata generation accepts
+the new library/table stamps. The all-local-game regression harness was started
+but stopped during its first baseline compilation, before any comparison
+completed; no cross-game comparison coverage is claimed for that sweep.
+
+## Mac development files
+
+The supplied folder is `/Volumes/S790/nsp/Xenoblade Chronicles 3/`. Extraction,
+the shader packages and isolated hosts are in `010074F013262000_build/`; analysis,
+private keys, Python tools and isolated runs are in `010074F013262000-work/`.
+`run/boot.py` accepts a log name, duration and environment assignments. Diagnostic
+logs `boot01` / `boot02` show the original allocator failure; `boot03wide` through
+`boot07wide` retain the wider-reservation experiments. `boot07wide` ends cleanly
+after 40 seconds; it uses an empty shader map and is not a visual test.
+
+Use `patched_512/` with `tools/dmhost_512` and `shader/fx_cache/` for the validated
+Mac layout; the earlier `patched_compact/` has the insufficient 1 GB first-arena
+reservation. `DM_NVN_SHMAP` points to this profile's `shader_map.tsv`.
+On the title/new-game screens press A. The auto-attack tip advances with R and
+closes with A; stand near an enemy to auto-attack and use the left stick to walk.
+During a movie, Plus then a held X skips it. The diagnostic watchdog is off when
+`DM_WATCH` is absent; setting it to `0` still enables its default stall threshold.
+
+## Found on this game
+
+- [Monolith ARH1, XBC1 and HCPS/SLCT shader containers](../../../docs/reference/shaders.md#monolith-arh1-xbc1-and-hcpsslct-shader-containers).
+- [Subgroup operations rewritten for a one-lane subgroup](../../../docs/reference/shaders.md#subgroup-operations-rewritten-for-a-one-lane-subgroup).
+- [IMAGE usage does not split texture storage](../../../docs/reference/rendering.md#image-usage-does-not-split-texture-storage).
+- [Rasterizer discard preserves vertex storage writes](../../../docs/reference/rendering.md#rasterizer-discard-preserves-vertex-storage-writes).
+- [Primitive barriers order vertex storage writes](../../../docs/reference/rendering.md#primitive-barriers-order-vertex-storage-writes).
+- [Transform feedback captured by a vertex entry point](../../../docs/reference/rendering.md#transform-feedback-captured-by-a-vertex-entry-point): interleaved outputs and geometry capture.
+- [A host lock held across a same-core handoff](../../../docs/reference/threads.md#a-host-lock-held-across-a-same-core-handoff): contended native decoder waits must release a handoff.
+- [oe calls answered without an applet proxy](../../../docs/reference/services.md#oe-calls-answered-without-an-applet-proxy): enabling album-image notifications dereferenced the absent SelfController.
+- [The SDK's address-region node area](../../../docs/reference/memory.md#the-sdks-address-region-node-area) and [1 GB alignment of regions of 1 GB and more](../../../docs/reference/memory.md#1-gb-alignment-of-regions-of-1-gb-and-more).
@@ -0,0 +1 @@
+# generated
@@ -0,0 +1,39 @@
+# Xenoblade Chronicles 3 2.2.1 (SDK 12.3.9, Monolith engine, NVN). Opening gameplay tested on Mac.
+# See README.md for measured status and the build-gated allocator changes.
+
+# ---- Identity ----
+title_id	010074F013262000
+program_id	010074F013262000
+
+# ---- App ----
+name	Xenoblade Chronicles 3
+ipa	XenobladeChronicles3.ipa
+app_category	public.app-category.role-playing-games
+pad_layout	full
+
+# ---- Versions ----
+build	v720896	2.2.1	main=82d187fe9ef9be921b8e1b1cddf7bd28	subsdk0=be1ee79a9968a35a6b63fc8f1036a6d4	sdk=e8cdceb1e77a8465fe20f21928e8d296
+package	v720896
+shader_versions	v720896
+
+# ---- Files ----
+# every module main imports from
+modules	main subsdk0 sdk
+romfs_marker	bf3.ard
+zip_store	.ard:.arh:.pck:.webm:.metallib
+
+# ---- Memory ----
+# the NPDM declares 32 MB of system resource (virtual address memory) and main never sets a heap: a 5 GB block with a 4.375 GB alias region, where the game maps its memory itself
+mem_reserve	0x140000000
+mem_alias	0x118000000
+mem_stack	0x10000000
+app_pool	0xCD500000
+# the NPDM declares 32 MB of system resource
+sys_resource	0x2000000
+
+# ---- Shader split ----
+shader_groups	/movie/=movie /sound/=sound
+shader_note	Xenoblade Chronicles 3 shaders
+
+# The allocator probes the physical-memory budget by mapping until failure.
+mem_map_limit	on
@@ -0,0 +1,23 @@
+# Xenoblade Chronicles 3: changes to its own modules, checked by NSO build ID.
+# Nintendo SDK module changes belong to dm/tools/sdk_patches.py.
+CUTS = {
+    # 2.2.1: nninitStartup reserves 60 GB for Monolith's allocator. Its first
+    # arena reserves 4 GB + 128 KB, maps the first 128 KB for bookkeeping, and
+    # takes additional pages on demand. iOS cannot reserve that address space.
+    # The original arena commits 203 MB at the title and 225 MB in the opening
+    # battle. A 1 GB first arena inside a 4 GB container leaves too little room
+    # for the battle's 390 MB core buffer: its 512 MB virtual window ends
+    # 70.75 MB past the container (12429a4, before any physical mapping).
+    # Reserve 512 MB for the first arena and 4 GB for the container, changing
+    # the bump increment, subtraction and stored capacity together. This is
+    # an address-space adaptation; physical limits still follow the Switch
+    # pool. Later-game allocation peaks remain to test.
+    '82d187fe9ef9be921b8e1b1cddf7bd28': [
+        (0x128e58c, 'e001c0d2', '2000c0d2', 'malloc address region: 60 GB to 4 GB'),
+        (0x1241870, 'e98340b2', 'e9ff9fd2', 'first arena bump size, low part: mov x9,#0xffff'),
+        (0x1241880, '4900a0f2', '4900a4f2', 'first arena bump size: 512 MB + 128 KB + 64 KB - 1'),
+        (0x12418d8, 'e97b5fb2', 'e98361b2', 'first arena subtraction, sign extension'),
+        (0x12418dc, 'c9ffbff2', 'c9ffbbf2', 'first arena subtraction: -(512 MB + 128 KB)'),
+        (0x12418e8, '2800c0f2', '4800a4f2', 'first arena stored capacity: 512 MB + 128 KB'),
+    ],
+}
@@ -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).
@@ -0,0 +1 @@
+# kind: evidence  vertex / pixel shader pairs of agl program variations (BAHS) that shader_pairs.tsv does not hold, for the pipeline prebuild only (generated by dm/gfx/tools/shader_split.py, do not edit): vertex key <TAB> pixel key
@@ -0,0 +1,14 @@
+# 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
+0279b7ec721373c9	{"local":[16,16,1],"shared":0}
+0aa01a2d1f8423dc	{"local":[64,1,1],"shared":0}
+120831ee717bbb78	{"local":[64,1,1],"shared":0}
+121d68f829de34e5	{"local":[64,1,1],"shared":0}
+1f10fd089f87ce2b	{"local":[16,16,1],"shared":0}
+20ddcfb262852fec	{"local":[64,1,1],"shared":0}
+5abc9844d234bcae	{"local":[16,16,1],"shared":11280}
+8cabe1ed395bddbd	{"local":[16,1,1],"shared":80}
+9417c75269b574ae	{"local":[64,1,1],"shared":0}
+ac9ae141f38b78fb	{"local":[16,16,1],"shared":11280}
+cbadd1e4769f8138	{"local":[16,16,1],"shared":4112}
+d5b5be30dd68a08a	{"local":[16,16,1],"shared":13328}
+f625a4f33c78facd	{"local":[16,16,1],"shared":7184}
@@ -0,0 +1,2 @@
+# kind: evidence  control-section key conflict rulings: empty
+dm_ctl_sha16	chosen_pc	class	candidates	records
@@ -115,6 +115,7 @@ One row per profile; the notes are each profile's `README.md`.
 | [The Legend of Zelda: Link's Awakening](01006BB00C6F0000/README.md) | `01006BB00C6F0000` | 1.1.0 | 8.2.0 | Grezzo's engine on NintendoWare | NVN |
 | [The Legend of Zelda: Tears of the Kingdom](0100F2C0115B6000/README.md) | `0100F2C0115B6000` | 1.4.3 | 20.5.22 | Nintendo EPD (sead / agl / gsys) | NVN |
 | [Unicorn Overlord](010069401ADB8000/README.md) | `010069401ADB8000` | 1.05 | 16.2.4 | Vanillaware's engine | NVN |
+| [Xenoblade Chronicles 3](010074F013262000/README.md) | `010074F013262000` | 2.2.1 | 12.3.9 | Monolith's engine | NVN |

 ## Keys

@@ -1,3 +1,3 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header): the B shader library translation chain version this source tree expects (default options), one for every game;
 // the host loads only <library dir>/b-<V>/. Run dmsh_chain.py --header after changing the chain's tools and commit this file with the change
-#define DMSH_CHAIN "a4e2088c2e69"
+#define DMSH_CHAIN "7fbca6dd477d"
@@ -617,10 +617,20 @@ NVN(nvnProgramSetShaders, NVNboolean, NVNprogram *p, int n, const NVNshaderData
    mtl_prog_release(o); o->nstage = n; memcpy(o->sd, sd, n * sizeof *sd); o->serial = nvn_new_serial();
    // The keys are computed here (see prog_key in nvn_obj.h): the control section is only guaranteed valid during this call; the code area lives in a memory pool and stays valid while the program is in use
    static int dump = -1; if (dump < 0) dump = getenv("DM_NVN_SHDUMP") != NULL;
+    static const char *trace; static int trace_init;
+    if (!trace_init) { trace = getenv("DM_NVN_PROGTRACE"); trace_init = 1; }
    prog_keys_t *k = pk_get(p, 1);
    k->ok = 0;
    for (int i = 0; i < n; i++) {
        if (keys_hash(&sd[i], 1, k->key[i])) k->ok |= (uint8_t)(1u << i);
+        // Diagnostic: trace the loader of selected shader keys (comma-separated hex).
+        if (trace && (k->ok & (1u << i))) {
+            char key[17]; for (int j = 0; j < 8; j++) snprintf(key + j * 2, 3, "%02x", k->key[i][j]);
+            if (strstr(trace, key)) {
+                LOG("[nvn] ProgramSetShaders key=%s control=%p code=%#llx", key, sd[i].control, (unsigned long long)sd[i].data);
+                dm_backtrace((uint64_t)__builtin_frame_address(0));
+            }
+        }
        // DM_NVN_SHDUMP (debugging): keep a copy of the control section so the replay layer can save shaders missing from the map (never freed)
        if (dump && sd[i].control) { void *c = malloc(0x880); memcpy(c, sd[i].control, 0x880); o->sd[i].control = c; }
    }
@@ -1133,7 +1143,7 @@ static _Atomic uint64_t g_nvn_ns;   // Host time spent in submit / present this
 static pthread_mutex_t g_susp_mu = PTHREAD_MUTEX_INITIALIZER;
 static pthread_cond_t g_susp_cv = PTHREAD_COND_INITIALIZER;
 static int g_susp, g_inflight_api;
-__thread _Atomic int dm_host_work;   // see hle.h (defined here: the graphics unit tests compile this file but not hle_thread.c)
+__thread _Atomic int dm_host_work;   // see dm.h (defined here: the graphics unit tests compile this file but not hle_thread.c)
 static void gfx_enter(void) { pthread_mutex_lock(&g_susp_mu); while (g_susp) pthread_cond_wait(&g_susp_cv, &g_susp_mu); g_inflight_api++; pthread_mutex_unlock(&g_susp_mu); dm_host_work++; }
 static void gfx_leave(void) { dm_host_work--; pthread_mutex_lock(&g_susp_mu); g_inflight_api--; pthread_cond_broadcast(&g_susp_cv); pthread_mutex_unlock(&g_susp_mu); }
 void dm_gfx_enter(void) { gfx_enter(); }   // the same gate for the Vulkan layer's submits and presents (dm/gfx/vk)
@@ -2165,7 +2175,9 @@ NVN(nvnCommandBufferBeginTransformFeedback, void, NVNcommandBuffer *c, NVNbuffer
 NVN(nvnCommandBufferEndTransformFeedback, void, NVNcommandBuffer *c, NVNbufferAddress a) { *REC(c, OP_TF_STATE, p_tf_state) = (p_tf_state){1, 0, a}; }
 NVN(nvnCommandBufferPauseTransformFeedback, void, NVNcommandBuffer *c, NVNbufferAddress a) { *REC(c, OP_TF_STATE, p_tf_state) = (p_tf_state){2, 0, a}; }
 NVN(nvnCommandBufferResumeTransformFeedback, void, NVNcommandBuffer *c, NVNbufferAddress a) { *REC(c, OP_TF_STATE, p_tf_state) = (p_tf_state){3, 0, a}; }
-NVN(nvnCommandBufferDrawTransformFeedback, void, NVNcommandBuffer *c, NVNdrawPrimitive m, NVNbufferAddress a) { (void)c; (void)m; (void)a; ONCE_LOG("[RED] DrawTransformFeedback: drawing the captured vertex count is not implemented, nothing drawn"); }
+NVN(nvnCommandBufferDrawTransformFeedback, void, NVNcommandBuffer *c, NVNdrawPrimitive m, NVNbufferAddress a) {
+    *REC(c, OP_DRAW, p_draw) = (p_draw){.prim=m,.instances=1,.ind=a+16};
+}

 // Device: combined texture handle = texture id | sampler id << 20 (same encoding as the separate handles)
 NVN(nvnDeviceGetTextureHandle, NVNtextureHandle, const NVNdevice *d, int tex, int samp) { (void)d; return (NVNtextureHandle)((uint64_t)(uint32_t)tex | (uint64_t)(uint32_t)samp << 20); }
@@ -0,0 +1,21 @@
+// The first 64 bytes are shared with offline main0_tf (dmsh_build.py).
+// Control buffers are opaque to the game: four byte offsets followed by Metal
+// array-draw arguments. All updates and saved/restored counts follow GPU order.
+typedef struct {
+    uint32_t offset[4], args[4]; uint64_t order;
+    uint32_t first, count, base_instance, pad; uint64_t reserved;
+    uint64_t state, source_count; uint32_t stride[4], capacity[4];
+    uint32_t requested, primitive; uint64_t reserved2;
+} nvn_tf_job;
+_Static_assert(offsetof(nvn_tf_job, state) == 64 && sizeof(nvn_tf_job) == 128, "feedback capture ABI");
+static const char k_nvn_tf_msl[] =
+"#include <metal_stdlib>\nusing namespace metal;\n"
+"struct capture {uint4 offset,args; device const uint* order; uint first,count,base_instance,pad; ulong reserved;};\n"
+"struct job {capture c; device uint* state; device const uint* source_count; uint4 stride,capacity; uint requested,primitive; ulong reserved;};\n"
+"kernel void tf_prepare(device job& j [[buffer(0)]]) {\n"
+" uint n=j.source_count?j.source_count[0]:j.requested;\n"
+" for(uint b=0;b<4;b++){j.c.offset[b]=j.state[b];if(j.stride[b]) n=min(n,(j.capacity[b]-min(j.capacity[b],j.state[b]))/j.stride[b]);}\n"
+" n-=n%j.primitive; j.c.args=uint4(n,1,0,j.c.base_instance);\n"
+" for(uint b=0;b<4;b++)if(j.stride[b])j.state[b]+=n*j.stride[b];\n"
+" j.state[4]+=n;j.state[5]=1;j.state[6]=j.state[7]=0;\n"
+"}\n";
@@ -0,0 +1,84 @@
+// Transform feedback uses a void vertex entry. For GS output the generated
+// reader gathers packed primitive indices; its draw count stays on the GPU.
+enum { TF_SLOT = 8 };
+static id<MTLBuffer> g_tf_state;
+static id<MTLComputePipelineState> g_tf_prepare;
+static void tf_state_gpu(const p_tf_state *b) {
+    rp_end(); ubd_land();
+    if (!g_tf_state) { g_tf_state = [g_dev newBufferWithLength:32 options:MTLResourceStorageModeShared]; res_add(g_tf_state); }
+    if (b->what == 0) {
+        id<MTLBuffer> buf; size_t off = arena_alloc(32, &buf);
+        uint32_t initial[8] = {0,0,0,0,0,1,0,0}; memcpy((uint8_t *)buf.contents + off, initial, 32);
+        enc_b_copy_buffer(blit(), buf, off, g_tf_state, 0, 32);
+    } else if (b->ctl) {
+        void *buffer; uint64_t off, gpu;
+        if (!mtl_gpu_addr(b->ctl, 32, &buffer, &off, &gpu)) { RED_ONCE(103, "transform feedback control buffer is outside a memory pool"); return; }
+        if (b->what == 3) enc_b_copy_buffer(blit_buf(b->ctl,b->ctl+32,0,0), (__bridge id<MTLBuffer>)buffer, off, g_tf_state, 0, 32);
+        else { enc_b_copy_buffer(blit_buf(0,0,b->ctl,b->ctl+32), g_tf_state, 0, (__bridge id<MTLBuffer>)buffer, off, 32); pend_add(b->ctl,32,4); }
+    }
+    g_wr_epoch++; g_mdx_epoch++;
+}
+static id<MTLRenderPipelineState> tf_pso(NvShader *vs) {
+    pso_key k; static const uint32_t cf[8]; pso_key_fill(&k, vs, nil, cf, 0, 0, 1, 0); k.fsfn = TF_MARK;
+    id<MTLRenderPipelineState> cached = hm_get_id(&g_pso, &k);
+    if (cached || hm_get_id(&g_pso_fail, &k)) return cached;
+    NSError *e = nil; id<MTLRenderPipelineState> ps = pso_build(pso_desc(&k, vs, nil, nil, 0), &e);
+    if (!ps) { MLOG("[RED] transform feedback pipeline %s: %s", vs->key, e.localizedDescription.UTF8String); hm_put_id(&g_pso_fail, &k, NSNull.null); }
+    else hm_put_id(&g_pso, &k, ps);
+    return ps;
+}
+static uint64_t tf_prepare(NvShader *vs, uint64_t source_count, uint64_t order, uint32_t count, uint32_t primitive,
+                           uint32_t first, uint32_t per_instance, uint32_t base_instance, uint64_t dst[4]) {
+    if (!g_tf_state || !vs || !vs->tf_fn) { RED_ONCE(100, "transform feedback shader %s has no capture entry", vs ? vs->key : "-"); return 0; }
+    nvn_tf_job j = {.order=order,.first=first,.count=per_instance,.base_instance=base_instance,
+                    .state=g_tf_state.gpuAddress,.source_count=source_count,.requested=count,.primitive=primitive};
+    for (int b = 0; b < 4; b++) if (vs->tf_stride[b]) {
+        void *buf; uint64_t off;
+        if (!S.tf[b].addr || S.tf[b].size > UINT32_MAX || !mtl_gpu_addr(S.tf[b].addr,S.tf[b].size,&buf,&off,&dst[b])) {
+            RED_ONCE(102, "transform feedback buffer %d of %s is unbound or outside its memory pool", b, vs->key); return 0; }
+        j.stride[b]=vs->tf_stride[b];j.capacity[b]=(uint32_t)S.tf[b].size;
+    }
+    if (!g_tf_prepare) {
+        NSError *e=nil; id<MTLLibrary> lib=[g_dev newLibraryWithSource:@(k_nvn_tf_msl) options:nil error:&e];
+        if (lib) g_tf_prepare=[g_dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"tf_prepare"] error:&e];
+        if (!g_tf_prepare) { MLOG("[RED] transform feedback prepare: %s",e.localizedDescription.UTF8String); return 0; }
+    }
+    rp_end(); ubd_land();
+    if (g_blit) { enc_b_end(g_blit);g_blit=nil; }
+    id<MTLBuffer> buf;size_t off=arena_alloc(sizeof j,&buf);memcpy((uint8_t *)buf.contents+off,&j,sizeof j);
+    uint64_t job=buf.gpuAddress+off;
+    id<MTL4ComputeCommandEncoder> ce=enc_c_begin(NULL);enc_c_pso(ce,g_tf_prepare);enc_c_args(ce,job,0);
+    enc_c_dispatch(ce,MTLSizeMake(1,1,1),MTLSizeMake(1,1,1));enc_c_end(ce);
+    return job;
+}
+static void tf_note(NvShader *vs) {
+    for(int b=0;b<4;b++)if(vs->tf_stride[b])pend_add(S.tf[b].addr,S.tf[b].size,4);
+    args_dirty_all();g_enc_pso=nil;g_enc_dss=nil;g_ed.ok=0;g_wr_epoch++;g_mdx_epoch++;
+}
+static void tf_draw(NvProg *pr,const p_draw *dw,const p_draw_idx *di) {
+    NvShader *vs=pr?pr->sh[0]:nil;
+    if(!dw||dw->ind||dw->mdi_cnt||!(dw->prim==0||dw->prim==1||dw->prim==4)) {
+        RED_ONCE(101,"transform feedback array capture requires a point, line or triangle list (vertex %s)",vs?vs->key:"-");return; }
+    if(dw->count<=0||dw->instances<0)return;
+    unsigned instances=MAX(dw->instances,1);uint64_t count=(uint64_t)dw->count*instances;
+    if(count>UINT32_MAX)return;
+    id<MTLRenderPipelineState> ps=vs&&vs->tf_fn?tf_pso(vs):nil;if(!ps)return;
+    uint64_t dst[4]={0};uint64_t job=tf_prepare(vs,0,0,(uint32_t)count,dw->prim==0?1:dw->prim==1?2:3,
+                                              dw->first,dw->count,dw->base_instance,dst);if(!job)return;
+    MTL4RenderPassDescriptor *rd=[MTL4RenderPassDescriptor new];rd.renderTargetWidth=rd.renderTargetHeight=rd.defaultRasterSampleCount=1;
+    id<MTL4RenderCommandEncoder> en=enc_r_begin(rd);enc_r_pso(en,ps);bind_vbufs(en);vconv_bind(en,vs,1,dw,NULL);
+    args_dirty_all();encode_args(en,vs,0,1);enc_r_vbuf(en,14,job);
+    for(int b=0;b<4;b++)if(vs->tf_stride[b])enc_r_vbuf(en,TF_SLOT+b,dst[b]);
+    enc_r_draw_instanced(en,MTLPrimitiveTypePoint,dw->first,dw->count,instances,dw->base_instance);enc_r_end(en);
+    swrite_note(pr);tf_note(vs);
+    PD("T %llu tf capture vs=%s first=%d count=%d instances=%u",g_stat_pass,vs->key,dw->first,dw->count,instances);
+}
+static void tf_geometry(NvShader *reader,uint64_t indices,uint64_t args,uint32_t primitive,uint32_t base_instance) {
+    if(!reader->tf_fn){RED_ONCE(104,"geometry feedback reader %s has no capture entry",reader->key);return;}
+    id<MTLRenderPipelineState> ps=tf_pso(reader);if(!ps)return;
+    uint64_t dst[4]={0};uint64_t job=tf_prepare(reader,args,indices,0,primitive,0,0,base_instance,dst);if(!job)return;
+    MTL4RenderPassDescriptor *rd=[MTL4RenderPassDescriptor new];rd.renderTargetWidth=rd.renderTargetHeight=rd.defaultRasterSampleCount=1;
+    id<MTL4RenderCommandEncoder> en=enc_r_begin(rd);enc_r_pso(en,ps);args_dirty_all();encode_args(en,reader,0,1);enc_r_vbuf(en,14,job);
+    for(int b=0;b<4;b++)if(reader->tf_stride[b])enc_r_vbuf(en,TF_SLOT+b,dst[b]);
+    enc_r_draw_indirect(en,MTLPrimitiveTypePoint,job+16);enc_r_end(en);tf_note(reader);
+}
@@ -80,6 +80,52 @@ static id<MTLRenderPipelineState> geo_capture_pso(NvShader *vs) {
    else { hm_put_id(&g_pso_fail, &k, NSNull.null); MLOG("[RED] geometry vertex capture %s: %s", vs->key, e.localizedDescription.UTF8String); }
    return ps;
 }
+// One diagnostic snapshot through the same native vertex fetch as geometry
+// capture. Refuse shaders with writes: observing them must not run their side
+// effects twice. Records retain the documented 144-word geometry capture ABI.
+static void vertex_dump(NvShader *vs, const p_draw *dw, const p_draw_idx *di) {
+    static const char *key, *path; static int configured, captured, limit; static uint64_t frame;
+    if (!configured) { key = getenv("DM_NVN_VSDUMP_KEY"); path = getenv("DM_NVN_VSDUMP");
+        const char *f = getenv("DM_NVN_VSDUMP_FROM"), *n = getenv("DM_NVN_VSDUMP_MAX");
+        frame = f ? strtoull(f,NULL,0) : 0; limit = n ? atoi(n) : key && !strcmp(key,"*") ? 256 : 1; configured = 1; }
+    int all = key && !strcmp(key,"*");
+    if (captured >= limit || !key || !path || atomic_load(&g_cur_frame) < frame || !vs || (!all && strncmp(vs->key, key, 16)) || !vs->geo_fn ||
+        (dw ? dw->ind || dw->mdi_cnt : di->ind || di->mdi_cnt) || S.rast_discard) return;
+    for (int i = 0; i < vs->rf.nb; i++) if (vs->rf.b[i].rw) return;
+    int count = dw ? dw->count : di->count, instances = dw ? dw->instances : di->instances;
+    if (count <= 0 || instances > 1 || (di && (di->itype < 0 || di->itype > 2))) return;
+    unsigned number = captured++; lc_land(); rp_flush(); mdi_wait();
+    uint32_t lo = dw ? (uint32_t)dw->first : UINT32_MAX, hi = dw ? lo + count - 1 : 0;
+    if (di) {
+        int size = 1 << di->itype;
+        if (!nvn_host_readable(di->ibuf, (uint64_t)count * size)) return;
+        for (int i = 0; i < count; i++) { uint32_t n = 0; memcpy(&n, (void *)(uintptr_t)(di->ibuf + (uint64_t)i * size), size);
+            if (S.restart.on && n == (S.restart.index & (size == 1 ? 255u : size == 2 ? 65535u : UINT32_MAX))) continue;
+            lo = MIN(lo, n); hi = MAX(hi, n); }
+        if (lo == UINT32_MAX || (int64_t)lo + di->base_vertex < 0 || (int64_t)hi + di->base_vertex > UINT32_MAX) return;
+        lo += di->base_vertex; hi += di->base_vertex;
+    }
+    uint64_t span = (uint64_t)hi - lo + 1, bytes = span * 144 * 4;
+    if (!span || bytes > (64u << 20)) return;
+    id<MTLRenderPipelineState> ps = geo_capture_pso(vs); if (!ps) return;
+    id<MTLBuffer> output; size_t offset = arena_alloc(bytes + 256, &output);
+    uint64_t gpu = output.gpuAddress + offset;
+    uint64_t info[4] = {gpu + 256, 0, (uint64_t)lo | span << 32, (uint32_t)(dw ? dw->base_instance : di->base_instance) | 1ull << 32};
+    memcpy((uint8_t *)output.contents + offset, info, sizeof info);
+    if (g_blit) { enc_b_end(g_blit); g_blit = nil; }
+    MTL4RenderPassDescriptor *rd = [MTL4RenderPassDescriptor new]; rd.renderTargetWidth = rd.renderTargetHeight = rd.defaultRasterSampleCount = 1;
+    id<MTL4RenderCommandEncoder> en = enc_r_begin(rd); enc_r_pso(en, ps); bind_vbufs(en);
+    p_draw range = {.first = (int32_t)lo, .count = (int32_t)span, .instances = 1, .base_instance = (int32_t)(info[3] & UINT32_MAX)};
+    vconv_bind(en, vs, 1, &range, NULL); args_dirty_all(); encode_args(en, vs, 0, 1); enc_r_vbuf(en, 13, gpu);
+    enc_r_draw_instanced(en, MTLPrimitiveTypePoint, lo, span, 1, (uint32_t)info[3]); enc_r_end(en);
+    mdi_wait();
+    char filename[1200]; if (all) snprintf(filename,sizeof filename,"%s_%04u_%s.bin",path,number,vs->key); else strlcpy(filename,path,sizeof filename);
+    FILE *f = fopen(filename, "wb"); if (f) { uint32_t header[4] = {lo, (uint32_t)span, all ? 4u : 144u, (uint32_t)info[3]}; fwrite(header, sizeof header, 1, f);
+        if (all) for (uint64_t i=0;i<span;i++) fwrite((uint8_t *)output.contents + offset + 256 + i*144*4,1,16,f);
+        else fwrite((uint8_t *)output.contents + offset + 256, 1, bytes, f); fclose(f); }
+    MLOG("vertex dump %s: first %u, count %llu, %s", vs->key, lo, (unsigned long long)span, filename);
+    args_dirty_all(); g_enc_pso = nil; g_enc_dss = nil; g_ed.ok = 0;
+}
 static void geo_draw(NvProg *pr, const p_draw *dw0, const p_draw_idx *di0);
 // A draw whose capture would need more scratch than the bound, drawn as smaller draws in the hardware's order: one instance at a time, then runs of a list topology's primitives
 // (points, lines, triangles, lines and triangles with adjacency), each with its own index range. The capture lays the vertices out by vertex number from the lowest index to the highest for
@@ -138,7 +184,7 @@ static int geo_split(NvProg *pr, const p_draw *dw, const p_draw_idx *di, int cou
 }
 static void geo_draw(NvProg *pr, const p_draw *dw0, const p_draw_idx *di0) {
    NvShader *vs = pr->sh[0], *base = pr->sh[2];
-    if (!vs || !vs->geo_fn || !base || !base->geo_vertices || pr->sh[3] || pr->sh[4] || S.tf_on) {
+    if (!vs || !vs->geo_fn || !base || !base->geo_vertices || pr->sh[3] || pr->sh[4]) {
        MLOG("[RED] geometry draw cannot capture this pipeline (vertex %s, geometry %s, tessellation %d, feedback %d)", vs ? vs->key : "-", base ? base->key : "-", !!(pr->sh[3] || pr->sh[4]), S.tf_on); return;
    }
    p_draw dw = dw0 ? *dw0 : (p_draw){0}; p_draw_idx di = di0 ? *di0 : (p_draw_idx){0};
@@ -262,6 +308,7 @@ static void geo_draw(NvProg *pr, const p_draw *dw0, const p_draw_idx *di0) {
    geo_pack_encode(ce, pack, job, (uint32_t)groups); enc_c_end(ce);
    args_dirty_all(); g_enc_pso = nil; g_enc_dss = nil; g_ed.ok = 0;
    swrite_note(pr); udc_note_draw(pr); g_wr_epoch++; g_mdx_epoch++;
+    if (S.tf_on) tf_geometry(reader, packed, args, gs->geo_output == 1 ? 1 : gs->geo_output == 6 ? 2 : 3, first_instance);
    if (S.rast_discard) return;
    NvProg *raster = pr->geo_raster;
    if (!raster || raster->sh[0] != reader || raster->sh[1] != pr->sh[1]) {
@@ -2476,6 +2476,9 @@ static id<MTLTexture> tex_mtl_(const tex_t *ct) {
        return mt;
    }
    alias_key ak; memset(&ak, 0, sizeof ak); ak.pool = (uint64_t)(uintptr_t)t->pool; ak.off = t->off; ak.d = t->d;
+    // IMAGE permits shader writes; it does not change storage. Xenoblade Chronicles 3's depth and stencil-sampling objects differ only by this flag. Separate Metal textures
+    // made the stencil read zero and rejected every deferred-lighting pass. Actual image writes add Metal usage before opening their encoder (image_prepare).
+    ak.d.flags &= ~NVN_TEXTURE_FLAGS_IMAGE;
    pthread_mutex_lock(&g_smu); __unsafe_unretained id shared = hm_get_id(&g_alias, &ak); if (shared) nvn_hm_put(&g_alias_seen, &ak, (void *)(uintptr_t)(g_fence + 1)); pthread_mutex_unlock(&g_smu);
    if (shared) { tex_put(t->serial, shared); return shared; }
    id<MTLTexture> mt = tex_linear(t);
@@ -3392,7 +3395,7 @@ static void shmap_load(void) {
    warm_start();
 }

-@interface NvShader : NSObject { @public id<MTLFunction> fn, tf_fn, fn_u, fn_m, geo_fn; uint8_t mask_m;   // fn_u: the integer-output entry point (dmsh_build.py fs_uint), for pipelines whose colour attachments are all unsigned integer (pso_desc); fn_m: the mixed one, its outputs in mask_m unsigned integer
+@interface NvShader : NSObject { @public id<MTLFunction> fn, tf_fn, discard_fn, fn_u, fn_m, geo_fn; uint8_t mask_m;   // fn_u: the integer-output entry point (dmsh_build.py fs_uint), for pipelines whose colour attachments are all unsigned integer (pso_desc); fn_m: the mixed one, its outputs in mask_m unsigned integer
    uint16_t tf_stride[4]; uint32_t cb_ext[16], sb_ro; nvn_refl rf; id<MTLArgumentEncoder> ae; NSUInteger arg_len; int stage, range_ro, layer_out; const char *builtin; uint64_t logged; char key[24]; id<MTLBuffer> c1buf; uint64_t c1_gpu; size_t c1off;
    uint32_t geo_vertices, geo_invocations, geo_output, geo_stride, geo_index_stride;
    id<MTLComputePipelineState> geo_pso;
@@ -3408,6 +3411,7 @@ static void shmap_load(void) {
@end
 #include "nvn_pso_skey.h"
 #include "nvn_geometry.h"
+#include "nvn_feedback.h"
@interface NvProg : NSObject { @public NvShader *sh[6]; char gs_miss[17]; NvProg *geo_raster; } @end   // geo_raster: GS output reader paired with this program's fragment stage
@implementation NvProg @end
 // Warm-up entry: built in the background, waiting to be taken by the submit thread. vs / fs are strong references (in pso they are only raw pointers, and the background dictionary is released when the task ends)
@@ -3594,6 +3598,14 @@ static NvShader *shader_load(NSString *pc, int stage, int layer) {
    // A function that declares function constants cannot go into a pipeline until it is specialized, even when every constant is optional. SPIRV-Cross declares one for shaders that read a
    // texel buffer (spvLinearTextureAlignmentOverride, used only through is_function_constant_defined; one of Pokémon Legends: Z-A's imgui pixel shaders, 2bd5037003980716): an empty set of
    // values keeps its default. Without it the first pipeline with that shader failed Metal's validation (`fragmentFunction main0 cannot be used to build a pipeline state`)
+    if (j[@"entry_discard"]) {
+        o->discard_fn = fn4_note([lib newFunctionWithName:j[@"entry_discard"]], lib, j[@"entry_discard"], nil);
+        if (o->discard_fn.functionConstantsDictionary.count) {
+            MTLFunctionConstantValues *cv = [MTLFunctionConstantValues new]; NSError *error = nil;
+            o->discard_fn = fn4_note([lib newFunctionWithName:j[@"entry_discard"] constantValues:cv error:&error], lib, j[@"entry_discard"], cv);
+        }
+        if (!o->discard_fn) MLOG("[RED] shader %s: no rasterizer-discard entry point %s in its library", pc.UTF8String, [j[@"entry_discard"] UTF8String]);
+    }
    if (j[@"entry_tf"]) {   // transform feedback capture entry point (dmsh_build.py vs_tf) and the bytes per vertex of each feedback buffer
        o->tf_fn = fn4_note([lib newFunctionWithName:j[@"entry_tf"]], lib, j[@"entry_tf"], nil); NSArray *ts = j[@"tf_strides"];
        for (NSUInteger b = 0; b < 4 && b < ts.count; b++) o->tf_stride[b] = (uint16_t)[ts[b] unsignedIntValue];
@@ -4202,7 +4214,7 @@ static struct {
    uint32_t st_vmask[2], st_wmask[2], st_ref[2];
    p_poly_offset po; float vp[16][4]; float dr[2]; int has_vp; int depth_clamp; uint8_t vp_sw[16][4];   // vp_sw: viewport swizzles XOR the identity (p_vp_swizzle)
    float blend_color[4]; int64_t copy_stride, copy_img_stride;
-    p_tf_bind tf[4]; uint64_t tf_off[4]; int tf_on, rast_discard;   // transform feedback buffers, bytes written into each since Begin, capture active; rasterizer discard (tf_draw)
+    p_tf_bind tf[4]; int tf_on, rast_discard;   // capture buffers and enables; byte offsets and captured count stay on the GPU (nvn_feedback.inc)
    int tex_base, samp_base;   // host-table bases of the texture / sampler pool bound (OP_SET_TEX_POOL / OP_SET_SAMP_POOL), for handles read at replay from uniform buffers (encode_args)
 } S;
 // The handle c2 holds at byte b, as encode_args fills the table (DRV2): a combined slot (past slot 31 the image there, tex_slot_h), a separate texture, else none
@@ -4273,7 +4285,7 @@ static void bare_dims(NSUInteger *w, NSUInteger *h) {
    *w = (NSUInteger)MIN(MAX(ceil(x1), 1.0), 16384.0); *h = (NSUInteger)MIN(MAX(ceil(y1), 1.0), 16384.0);
 }
 // The open draw pass has a draw whose shaders write an image or a storage buffer. NVN orders later draws after those writes when the game asks with nvnCommandBufferBarrier
-// (ORDER_FRAGMENTS / ORDER_FRAGMENTS_TILED); Metal gives no order between draws of one render pass for shader writes, so such a barrier ends the pass and the next draw opens one
+// (ORDER_PRIMITIVES / ORDER_FRAGMENTS / ORDER_FRAGMENTS_TILED); Metal gives no order between draws of one render pass for shader writes, so such a barrier ends the pass and the next draw opens one
 // behind the full consumer barrier. FINAL FANTASY TACTICS - The Ivalice Chronicles composites its UI into the window texture in software, a read-modify-write of the image by draw after
 // draw in one attachment-less pass with barrier bits 0x12 between them: without the split a later element read the window before the earlier one had landed and wrote the old pixels
 // back (the dialog boxes had holes the size of their text)
@@ -4722,11 +4734,12 @@ static NSMutableDictionary<NSString *, id<MTLFunction>> *g_bf_fn_cache;
 static NSMutableDictionary<NSString *, id<MTLLibrary>> *g_bf_lib_cache;
 static nvn_hmap g_pso = {.klen = sizeof(pso_key)}, g_dss;
 static nvn_hmap g_pso_fail = {.klen = sizeof(pso_key)};   // pipelines that failed to build: remembered and never retried (previously every draw recompiled it, over 10,000 times in one iPhone session)
-static const char g_tf_mark, g_amp_mark;
+static const char g_tf_mark, g_amp_mark, g_discard_mark;
 static char g_geo_mark;
 #define GEO_MARK ((const void *)&g_geo_mark)
 #define AMP_MARK ((const void *)&g_amp_mark)   // pso_key.fsfn of a slice-merged pass's pipeline (vertex amplification, sm_amp_get)
 #define TF_MARK ((const void *)&g_tf_mark)   // pso_key.fsfn of a transform feedback capture pipeline (tf_pso)
+#define DISCARD_MARK ((const void *)&g_discard_mark)   // vertex work still runs, without rasterization or a fragment stage
 static int mtl_fmt_uint(MTLPixelFormat f) {
    switch (f) { case MTLPixelFormatR8Uint: case MTLPixelFormatR16Uint: case MTLPixelFormatR32Uint: case MTLPixelFormatRG8Uint: case MTLPixelFormatRG16Uint: case MTLPixelFormatRG32Uint:
        case MTLPixelFormatRGBA8Uint: case MTLPixelFormatRGB10A2Uint: case MTLPixelFormatRGBA16Uint: case MTLPixelFormatRGBA32Uint: return 1; default: return 0; }
@@ -4781,6 +4794,7 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS
    if (vs->layer_out && k->fsfn != GEO_MARK) d.inputPrimitiveTopology = vs->geo_output == 1 ? MTLPrimitiveTopologyClassPoint : vs->geo_output == 6 ? MTLPrimitiveTopologyClassLine : MTLPrimitiveTopologyClassTriangle;
    if (k->fsfn == AMP_MARK) d.maxVertexAmplificationCount = 2;   // a slice-merged pass's pipeline: the view mapping's layer offset applies (see g_sm_plan)
    if (k->fsfn == TF_MARK || k->fsfn == GEO_MARK) { d.vertexFunction = k->fsfn == GEO_MARK ? vs->geo_fn : vs->tf_fn; d.fragmentFunction = nil; d.rasterizationEnabled = NO; }   // transform feedback capture (tf_draw)
+    if (k->fsfn == DISCARD_MARK) { d.vertexFunction = vs->discard_fn; d.fragmentFunction = nil; d.rasterizationEnabled = NO; }
    d.rasterSampleCount = k->samples ? k->samples : 1;
    d.alphaToCoverageEnabled = k->a2c;
    for (int i = 0; i < 8; i++) {
@@ -4805,7 +4819,7 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS
    }
    d.depthAttachmentPixelFormat = (MTLPixelFormat)k->dfmt; d.stencilAttachmentPixelFormat = (MTLPixelFormat)k->sfmt;
    MTLVertexDescriptor *vd = [MTLVertexDescriptor vertexDescriptor];
-    const nvn_refl *rf = &vs->rf; int nconv = 0, on = vin_on(vs, k, k->fsfn == TF_MARK || k->fsfn == GEO_MARK);
+    const nvn_refl *rf = &vs->rf; int nconv = 0, on = vin_on(vs, k, k->fsfn == TF_MARK || k->fsfn == GEO_MARK || k->fsfn == DISCARD_MARK);
    if (on) { d.vertexFunction = vs->vfn; g_vin_psos++; }   // the shadow-input entry point (see vin_on)
    #define ZERO_ATTR(i, ct_) do { vd.attributes[i].format = (ct_) == NVN_CT_UINT ? MTLVertexFormatUInt4 : (ct_) == NVN_CT_INT ? MTLVertexFormatInt4 : MTLVertexFormatFloat4; vd.attributes[i].offset = 0; vd.attributes[i].bufferIndex = 30; \
                                  vd.layouts[30].stride = 16; vd.layouts[30].stepFunction = MTLVertexStepFunctionConstant; vd.layouts[30].stepRate = 0; } while (0)   // reads the zero buffer
@@ -5574,6 +5588,17 @@ static __unsafe_unretained id<MTLRenderPipelineState> pso_get_(NvProg *pr, MTLRe
    __unsafe_unretained NvShader *vs = pr->sh[0], *fs = pr->sh[1];
    if (!vs) { RED_ONCE(13, "program has no vertex shader (or it is not mapped)"); return nil; }
    pso_key k; pso_key_fill(&k, vs, fs, g_pass.cfmt, g_pass.dfmt, g_pass.sfmt, g_pass.samples, g_pso_lmap != NULL);
+    if (S.rast_discard) {
+        k.fsfn = DISCARD_MARK;
+        id<MTLRenderPipelineState> cached = hm_get_id(&g_pso, &k);
+        if (cached) return cached;
+        if (hm_get_id(&g_pso_fail, &k)) return nil;
+        NSError *error = nil;
+        id<MTLRenderPipelineState> ps = pso_build(pso_desc(&k, vs, nil, nil, 0), &error);
+        if (ps) hm_put_id(&g_pso, &k, ps);
+        else { MLOG("[RED] rasterizer-discard pipeline (vertex %s): %s", vs->key, error.localizedDescription.UTF8String); hm_put_id(&g_pso_fail, &k, NSNull.null); }
+        return ps;
+    }
    if (g_pso_lmap) { memcpy(k.lmap, g_pso_lmap, 8); k.fsfn = (__bridge void *)g_pso_fsfn; }
    id<MTLFunction> ms_fn = ms_fs_override(fs); if (ms_fn) k.fsfn = (__bridge void *)ms_fn;   // DM_NVN_MAGSCAN phases 2 / 3
    if (g_pso_amp) { if (k.fsfn || g_pso_lmap) return nil; k.fsfn = AMP_MARK; return sm_amp_get(&k, vs, fs); }   // a slice-merged pass's pipeline: built in the background, nil until then
@@ -6509,6 +6534,7 @@ static int attr_read(int a, const p_draw *dw, const p_draw_idx *di, float P[64][
    return 0;
 }
 static int first_draw_covers(int i, id<MTLTexture> t, NSUInteger level) {
+    if (S.rast_discard) return g_fd_why = 2, 0;
    const p_draw *dw = g_fd_dw; const p_draw_idx *di = g_fd_di;
    int n = dw ? dw->count : di->count, prim = dw ? dw->prim : di->prim;
    __unsafe_unretained NvProg *pr = S.prog ? prog_get(S.prog) : nil; if (!pr) return g_fd_why = 1, 0;
@@ -6587,10 +6613,9 @@ static int state_op(uint32_t op, const void *a) {
        case OP_BLEND_COLOR: memcpy(S.blend_color, ((const p_blend_color *)a)->rgba, 16); break;
        case OP_COPY_STRIDE: { const p_copy_stride *cs = a; if (cs->image) S.copy_img_stride = cs->stride; else S.copy_stride = cs->stride; break; }
        case OP_TF_BIND: { const p_tf_bind *b = a; if (b->index >= 0 && b->index < 4) S.tf[b->index] = *b; break; }
-        case OP_TF_STATE: { const p_tf_state *b = a;   // Begin writes from the start of each buffer; Pause / Resume keep the offsets (the host keeps them itself, the game's control buffer is not used)
-            if (b->what == 0) { S.tf_on = 1; memset(S.tf_off, 0, sizeof S.tf_off); } else S.tf_on = b->what == 3; break; }
+        case OP_TF_STATE: { const p_tf_state *b = a; S.tf_on = b->what == 0 || b->what == 3; break; }
        case OP_PRIM_RESTART: S.restart = *(const p_prim_restart *)a; break;
-        case OP_RAST_DISCARD: S.rast_discard = ((const p_rast_discard *)a)->on; break;
+        case OP_RAST_DISCARD: S.rast_discard = ((const p_rast_discard *)a)->on; g_pso_dirty = 1; break;
        default: return 0;
    }
    return 1;
@@ -6669,7 +6694,7 @@ static void pso_predict_rest_(void) {
            df = dfull; sf = sfull; continue;
        }
        if (hq.op == OP_DRAW || hq.op == OP_DRAW_IDX) {
-            if (S.rast_discard) continue;   // rasterizes nothing: transform feedback capture or nothing (tf_draw)
+            if (S.rast_discard) continue;   // vertex work uses its own non-rasterizing pipeline, built on demand
            draws++;
            __unsafe_unretained NvProg *pr = S.prog ? prog_get(S.prog) : nil; if (!pr || !pr->sh[0]) continue;
            __unsafe_unretained NvShader *vs = pr->sh[0], *fs = pr->sh[1];
@@ -7629,7 +7654,7 @@ static int zskip_can_draw(NvProg *pr) {
    // Skip conditions (all must hold): the fragment key is in the fact table, the view bound to every texture slot declared by its reflection covers only subresources known to be all zero, the coefficient is finite,
    // no depth write and no stencil, not counting an occlusion query, and every color attachment to be written is already 0 (a candidate cleared entirely to 0 in this render pass, or known all zero) -- writing 0 over 0.
    // Reasons for not holding (the N line of PASSDUMP): 1 depth / stencil 2 occlusion query counting with a report before the reset 3 coefficient unreadable or non-finite 4 image slot 5 sampled view not known all zero 6 no target 7 target is not 0
-    if (!zskip_on() || !pr) return 0;
+    if (S.rast_discard || !zskip_on() || !pr) return 0;
    NvShader *fs = nil; for (int st = 0; st < 5; st++) if (pr->sh[st] && pr->sh[st]->stage == 1) fs = pr->sh[st];
    zskip_fact f; float coeff; int why = 0;
    if (!fs || !zskip_fact_for(fs->key, &f)) return 0;
@@ -8258,60 +8283,7 @@ static int mdx_next(const p_draw *dw, const p_draw_idx *di) {
    g_stat_draw++; g_mdx_fast++; g_mdxc.next = g_exec_next; g_mdxc.i = i;
    return 1;
 }
-// ---------- Transform feedback ----------
-// A draw under rasterizer discard rasterizes nothing; with feedback active it captures its vertex shader's outputs into the bound feedback buffers. Pokémon Legends: Z-A skins its characters
-// so: a vertex-only skinning program writes positions (12 bytes a vertex) and packed normal / tangent bits (4 bytes each), and the character draws read those buffers as vertex streams.
-// Metal has no feedback stage: the shader's main0_tf entry point (dmsh_build.py vs_tf) runs in a pipeline with rasterization off, in an encoder of its own with no attachments (the game draws
-// these while its frame target is bound; a pass of that target would load and store it for nothing), and stores the outputs at TF_SLOT + b. Each draw writes from the buffer's current offset on.
-// Only plain array draws are captured (all this game makes); the encoder's barriers order the writes before the draws that read them
-enum { TF_SLOT = 8 };   // = dmsh_build.py TF_SLOT
-static __unsafe_unretained id<MTLRenderPipelineState> tf_pso(NvShader *vs) {
-    pso_key k; static const uint32_t cf[8]; pso_key_fill(&k, vs, nil, cf, 0, 0, 1, 0); k.fsfn = TF_MARK;
-    __unsafe_unretained id<MTLRenderPipelineState> cached = hm_get_id(&g_pso, &k);
-    if (cached) return cached;
-    if (hm_get_id(&g_pso_fail, &k)) return nil;
-    NSError *e = nil; id<MTLRenderPipelineState> ps = pso_build(pso_desc(&k, vs, nil, nil, 0), &e);
-    if (!ps) { MLOG("[RED] failed to create the transform feedback pipeline (vertex %s): %s", vs->key, e.localizedDescription.UTF8String); hm_put_id(&g_pso_fail, &k, NSNull.null); return nil; }
-    hm_put_id(&g_pso, &k, ps);
-    return ps;
-}
-static void tf_draw(NvProg *pr, const p_draw *dw, const p_draw_idx *di) {
-    __unsafe_unretained NvShader *vs = pr ? pr->sh[0] : nil;
-    if (!S.tf_on) { PD("Z %llu draw skipped: rasterizer discard", g_stat_pass + !g_enc); return; }
-    if (!vs || !vs->tf_fn) { RED_ONCE(100, "transform feedback with vertex shader %s, which has no capture entry point: nothing captured", vs ? vs->key : "-"); return; }
-    if (!dw || dw->ind || dw->mdi_cnt || dw->instances > 1 || dw->prim == 6 || dw->prim == 7) { RED_ONCE(101, "transform feedback draw that is indexed, indirect, instanced or a fan / quad list (vertex %s): not captured", vs->key); return; }
-    if (dw->count <= 0) return;
-    uint64_t gpu[4] = {0};
-    for (int b = 0; b < 4; b++) {
-        if (!vs->tf_stride[b]) continue;
-        uint64_t n = (uint64_t)dw->count * vs->tf_stride[b], at = S.tf[b].addr + S.tf_off[b]; void *mb; uint64_t mo;
-        if (!S.tf[b].addr || S.tf_off[b] + n > S.tf[b].size || !mtl_gpu_addr(at, n, &mb, &mo, &gpu[b])) {
-            RED_ONCE(102, "transform feedback buffer %d of vertex %s is unbound, too small or in no memory pool (%#llx + %llu of %llu): not captured", b, vs->key, (unsigned long long)S.tf[b].addr, (unsigned long long)(S.tf_off[b] + n), (unsigned long long)S.tf[b].size);
-            return; }
-    }
-    __unsafe_unretained id<MTLRenderPipelineState> ps = tf_pso(vs);
-    if (!ps) return;
-    rp_end(); ubd_land();   // the deferred inline updates first: this draw comes before draw's conflict check (ubd_conflict) and writes its capture buffers
-    // ubd_land copies them in the copy encoder and leaves it open, as does a copy or dispatch just before: a render encoder opened beside it is invalid (Metal validation: "Command buffer already
-    // has an open command encoder"). Pokémon Legends: Z-A crashed in the driver when that copy encoder ended, two minutes into a new game on the Mac (GitHub issue #101)
-    if (g_blit) { enc_b_end(g_blit); g_blit = nil; }
-    MTL4RenderPassDescriptor *d = [MTL4RenderPassDescriptor new]; d.renderTargetWidth = 1; d.renderTargetHeight = 1; d.defaultRasterSampleCount = 1;
-    id<MTL4RenderCommandEncoder> en = enc_r_begin(d);
-    enc_r_pso(en, ps);
-    bind_vbufs(en);
-    vconv_bind(en, vs, 1, dw, NULL);
-    args_dirty_all(); encode_args(en, vs, 0, 1);
-    for (int b = 0; b < 4; b++) if (vs->tf_stride[b]) {   // vertex v of the draw (first + i) stores at v × stride: the pointer starts first × stride before the write offset
-        enc_r_vbuf(en, TF_SLOT + b, gpu[b] - (uint64_t)dw->first * vs->tf_stride[b]);
-        uint64_t n = (uint64_t)dw->count * vs->tf_stride[b]; pend_add(S.tf[b].addr + S.tf_off[b], n, 4); S.tf_off[b] += n;
-    }
-    enc_r_draw(en, MTLPrimitiveTypePoint, (NSUInteger)dw->first, (NSUInteger)dw->count);
-    m4_end(en);
-    args_dirty_all(); g_enc_pso = nil; g_enc_dss = nil; g_ed.ok = 0;
-    { static int once; if (!once++) MLOG("transform feedback: first capture (vertex %s, %d vertices)", vs->key, dw->count); }
-    PD("T %llu tf capture vs=%s first=%d count=%d written %llu/%llu/%llu/%llu", g_stat_pass, vs->key, dw->first, dw->count,
-       (unsigned long long)S.tf_off[0], (unsigned long long)S.tf_off[1], (unsigned long long)S.tf_off[2], (unsigned long long)S.tf_off[3]);
-}
+#include "nvn_feedback.inc"

 #include "nvn_geometry.inc"

@@ -8501,6 +8473,7 @@ static void cs_hoist_try(NvProg *pr0, const p_draw *dw0, const p_draw_idx *di0)
            PD("H %llu dispatch %s moved ahead of the pass%s", g_stat_pass + 1, cs->key, depth ? " (samples the pass's depth, which no draw before it wrote)" : "");
            dispatch(dp); g_csh_op[g_csh_n++] = q; g_csh_moved++; g_csh_depth += depth; moved++;
            continue; }
+        if (hq.op == OP_TF_STATE) { g_csh_stop[5]++; break; }   // saves/restores counters in GPU memory
        if (state_op(hq.op, aq)) continue;
        g_csh_stop[5]++; break;   // anything else may end the pass, write memory or depend on the order (copies, updates, clears, targets, queries)
    }
@@ -8511,10 +8484,16 @@ static void cs_hoist_try(NvProg *pr0, const p_draw *dw0, const p_draw_idx *di0)
 void mtl_cshoist_stats(uint64_t *moved, uint64_t *depth, uint64_t *stops) { *moved = g_csh_moved; *depth = g_csh_depth; for (int i = 0; i < 8; i++) stops[i] = g_csh_stop[i]; }
 static void draw(const p_draw *dw, const p_draw_idx *di) {
    g_draw_binst = (uint32_t)(dw ? dw->base_instance : di ? di->base_instance : 0);   // c0 + 0x644 of the stages that read the driver segment (encode_args)
+    if (passdump()) { NvProg *p = g_geo_program ? g_geo_program : S.prog ? prog_get(S.prog) : nil;
+        PD("j %llu draw-entry vs=%s fs=%s gs=%s prim=%d count=%d instances=%d indexed=%d indirect=%#llx discard=%d feedback=%d", g_stat_pass + !g_enc,
+           p && p->sh[0] ? p->sh[0]->key : "-", p && p->sh[1] ? p->sh[1]->key : "-", p && p->sh[2] ? p->sh[2]->key : "-",
+           dw ? dw->prim : di->prim, dw ? dw->count : di->count, dw ? dw->instances : di->instances, di != NULL,
+           (unsigned long long)(dw ? dw->ind : di->ind), S.rast_discard, S.tf_on); }
    if (!g_geo_program && S.prog) {
        NvProg *geometry_program = prog_get(S.prog);
        if (geometry_program && geometry_program->sh[2] && geometry_program->sh[2]->geo_vertices) {
            if ((dw && dw->mdi_cnt && !mdi_live(dw->mdi_cnt, dw->mdi_i)) || (di && di->mdi_cnt && !mdi_live(di->mdi_cnt, di->mdi_i))) return;
+            image_prepare(geometry_program, (1u << 0) | (1u << 2));
            geo_draw(geometry_program, dw, di); return;
        }
    }
@@ -8526,11 +8505,14 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
    __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
-    // (a vertex-only program, colour mask 0, `7dcdaed991c74bfe`), and its fur shells read them; the capture runs first, then the draw rasterizes as on the GPU
+    if (S.rast_discard && S.tf_on) { tf_draw(pr, dw, di); return; }
+    // Feedback captures rasterized draws too (Donkey Kong Country: Tropical Freeze's fur).
    if (S.tf_on && pr && pr->sh[0] && pr->sh[0]->tf_fn) tf_draw(pr, dw, di);
+    // Rasterizer discard suppresses fragments, not vertex execution or its storage writes.
+    // Xenoblade Chronicles 3 uses vertex-only work before its skinned-model draws.
+    if (S.rast_discard && (!pr->sh[0] || !pr->sh[0]->discard_fn)) { PD("Z %llu draw skipped: rasterizer discard with no vertex storage writes vs=%s", g_stat_pass + !g_enc, pr->sh[0] ? pr->sh[0]->key : "-"); return; }
    if (g_nubd && (g_enc || ubd_lazy()) && ubd_conflict(pr, dw, di)) { g_ubd_conflicts++; rp_end(); ubd_land(); }   // land the deferred inline updates before this draw (see ubd_conflict)
+    vertex_dump(pr->sh[0], dw, di);
    if (zskip_can_draw(pr)) {   // Zero-source skip: the output is always 0, writing 0 over 0, so this draw changes no result (see zskip_can_draw)
        g_zskip_count++; { static int once; if (!once++) MLOG("zero-source skip: first case pass %llu", g_stat_pass + 1); }
        PD("N %llu zero-source draw skipped", g_stat_pass + 1); return; }
@@ -8557,7 +8539,7 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
    { static int small = -1; if (small < 0) { const char *e = getenv("DM_NVN_SKIP_SMALL"); small = e ? atoi(e) : 0; if (small) MLOG("DM_NVN_SKIP_SMALL: dropping draws whose target width is ≤ %d (diagnostic; the picture is wrong)", small); }
      if (small > 0 && R.has_rt && !g_defer_replaying) { int k = R.rt.n && R.rt.color[0] ? 0 : R.rt.depth ? 8 : -1;
        if (k >= 0) { uint64_t s_; int lv, sl; rt_sub(&R.rt, k, &s_, &lv, &sl); const tex_t *t = k == 8 ? R.rt.depth : R.rt.color[0]; if ((int)MAX(t->d.width >> lv, 1u) <= small) return; } } }
-    if (S.poly.cull == NVN_FACE_FRONT_AND_BACK) { PD("Z %llu draw skipped (before the pass begins): both faces culled", g_stat_pass + !g_enc); return; }
+    if (S.poly.cull == NVN_FACE_FRONT_AND_BACK && !S.rast_discard) { PD("Z %llu draw skipped (before the pass begins): both faces culled", g_stat_pass + !g_enc); return; }
    // The continued render pass may have been opened with fewer depth / stencil attachments (open modes 2 / 3 / 4): if the plane this draw needs is not attached, split before it (on reopen, attach per the new lookahead)
    if (g_rt_merged && g_enc_mode != 0 && (((S.ds.dt || S.ds.dw) && !g_pass.hd) || (S.ds.st && !g_pass.hs))) {
        static int once2; if (!once2++) MLOG("same-target continuation: a draw in the continued pass needs a depth / stencil plane that is not attached; split here (first case pass %llu)", g_stat_pass);
@@ -8570,7 +8552,7 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
        static int once3; if (!once3++) MLOG("readable-tile pass merge: a draw of the later pass is not provable; back to separate passes (first case pass %llu)", g_stat_pass);
        g_fz_fallback++; rp_end(); }
    fb_prepare(pr);   // same-pass feedback sampling: copy the snapshot first (see fb_prepare)
-    if (R.has_sc && !g_enc) {
+    if (R.has_sc && !g_enc && !S.rast_discard) {
        NSUInteger tw0 = 0, th0 = 0; rt_dims(&tw0, &th0);
        long x0 = R.sc.x < 0 ? 0 : R.sc.x, y0 = R.sc.y < 0 ? 0 : R.sc.y, x1 = (long)R.sc.x + R.sc.w, y1 = (long)R.sc.y + R.sc.h;
        if (x1 > (long)tw0) x1 = tw0; if (y1 > (long)th0) y1 = th0;
@@ -8634,14 +8616,14 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
        else if (sw) RED_ONCE(108, "viewport swizzle %d,%d,%d,%d not implemented (drawn unswizzled)", S.vp_sw[0][0], S.vp_sw[0][1] ^ 2, S.vp_sw[0][2] ^ 4, S.vp_sw[0][3] ^ 6); }
      if (!(g_ed.ok & ED_VP) || memcmp(&g_ed.vp, &v, sizeof v)) { enc_r_viewport(en, v); g_ed.vp = v; g_ed.ok |= ED_VP; } }
    MTLScissorRect sc = {0, 0, tw, th};
-    if (R.has_sc) {
+    if (R.has_sc && !S.rast_discard) {   // raster state cannot suppress vertex storage writes
        long x0 = R.sc.x < 0 ? 0 : R.sc.x, y0 = R.sc.y <= 0 ? 0 : R.sc.y, x1 = (long)R.sc.x + R.sc.w, y1 = (long)R.sc.y + R.sc.h;
        if (x1 > (long)tw) x1 = tw; if (y1 > (long)th) y1 = th;
        if (x1 <= x0 || y1 <= y0) { PD("Z %llu draw skipped: empty scissor (%d,%d %dx%d, target %lux%lu)", g_stat_pass, R.sc.x, R.sc.y, R.sc.w, R.sc.h, (unsigned long)tw, (unsigned long)th); return; }
        sc = (MTLScissorRect){(NSUInteger)x0, (NSUInteger)y0, (NSUInteger)(x1 - x0), (NSUInteger)(y1 - y0)};
    }
    if (!(g_ed.ok & ED_SC) || memcmp(&g_ed.sc, &sc, sizeof sc)) { enc_r_scissor(en, sc); g_ed.sc = sc; g_ed.ok |= ED_SC; }
-    if (S.poly.cull == NVN_FACE_FRONT_AND_BACK) { PD("Z %llu draw skipped: both faces culled", g_stat_pass); return; }
+    if (S.poly.cull == NVN_FACE_FRONT_AND_BACK && !S.rast_discard) { PD("Z %llu draw skipped: both faces culled", g_stat_pass); return; }
    { int cu = S.poly.cull == NVN_FACE_FRONT ? MTLCullModeFront : S.poly.cull == NVN_FACE_BACK ? MTLCullModeBack : MTLCullModeNone;
      int wi = (S.poly.front == NVN_FRONT_FACE_CCW) != atomic_load_explicit(&g_origin_ll, memory_order_relaxed) ? MTLWindingCounterClockwise : MTLWindingClockwise;   // the y flip of lower-left mode reverses the winding
      int fi = S.poly.mode == NVN_POLYGON_MODE_LINE ? MTLTriangleFillModeLines : MTLTriangleFillModeFill, cl = !S.depth_clamp;   // depth clip off = SetDepthClamp
@@ -8655,11 +8637,11 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
    if (cur_pull) { pull_bind(en); g_hy_draws++; }   // a pulling pipeline reads the streams itself (no conversion either: dm_fetch decodes every format)
    else {
        bind_vbufs(en);
-        if (pr->sh[0]) vconv_bind(en, pr->sh[0], 0, dw, di);   // integer inputs fed float-type formats or 8 / 16-bit integers of the other sign (see vconv_ncomp, vconv_int)
+        if (pr->sh[0]) vconv_bind(en, pr->sh[0], S.rast_discard, dw, di);   // integer inputs fed float-type formats or 8 / 16-bit integers of the other sign (see vconv_ncomp, vconv_int)
    }
    g_kd_force = kd_draw(pr, dw, di);
    if (pr->sh[0]) encode_args(en, pr->sh[0], 0, 1);
-    if (pr->sh[1]) encode_args(en, pr->sh[1], 1, 0);
+    if (pr->sh[1] && !S.rast_discard) encode_args(en, pr->sh[1], 1, 0);
    // Parameter binding captures the actual resources under the registry lock; if the version changed between the eligibility check and the binding, split and retry before any GPU draw.
    if (g_rt_range_active && (g_pass.reg_gen != g_reg_gen || g_pass.ser_gen != g_ser_gen || g_pass.range_dirty != nvn_dirty_now())) { rp_end(); draw(dw, di); return; }
    int prim = dw ? dw->prim : di->prim;
@@ -8726,6 +8708,9 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
               dw ? dw->count : di->count, dw ? dw->instances : di->instances, S.color.blend_en, S.chmask.mask, (int)S.ds.dt, (int)S.ds.dw, (int)S.ds.df, (int)S.ds.st,
               (unsigned long)sc.x, (unsigned long)sc.y, (unsigned long)sc.width, (unsigned long)sc.height, (unsigned long)tw, (unsigned long)th,
               S.has_vp ? S.vp[0][0] : 0, S.has_vp ? S.vp[0][1] : 0, S.has_vp ? S.vp[0][2] : (float)tw, S.has_vp ? S.vp[0][3] : (float)th, S.has_vp ? S.dr[0] : 0.f, S.has_vp ? S.dr[1] : 1.f, (unsigned)S.color.logic_op);
+            PD("s %llu front cmp=%u fail=%u zfail=%u pass=%u ref=%#x read=%#x write=%#x back cmp=%u fail=%u zfail=%u pass=%u ref=%#x read=%#x write=%#x cull=%u winding=%u", g_stat_pass,
+               S.ds.sf_f, S.ds.fail_f, S.ds.zfail_f, S.ds.zpass_f, S.st_ref[0], S.st_vmask[0], S.st_wmask[0],
+               S.ds.sf_b, S.ds.fail_b, S.ds.zfail_b, S.ds.zpass_b, S.st_ref[1], S.st_vmask[1], S.st_wmask[1], S.poly.cull, S.poly.front);
        }
    }
    g_kd_force = 0;
@@ -9559,13 +9544,18 @@ void mtl_exec(const uint8_t *p) {
            case OP_DRAW: { if (g_rc_skip) break; p_draw t; g_exec_cur = p; g_exec_next = p + h.size; draw(draw_rec(a, h.size, &t), NULL); g_exec_next = NULL; break; }   // Lookahead start point
            case OP_DRAW_IDX: { if (g_rc_skip) break; p_draw_idx t; g_exec_cur = p; g_exec_next = p + h.size; draw(NULL, draw_idx_rec(a, h.size, &t)); g_exec_next = NULL; break; }
            // Do not mark dirty if the state did not change (the game often re-sends the same program / state block / vertex format before each draw; the dirty flag only decides whether to re-fetch the pipeline / depth-stencil object, and equal values mean the same object)
-            case OP_SCISSOR: case OP_BIND_PROGRAM: case OP_BIND_STATE: case OP_BIND_VATTR: case OP_BIND_VSTREAM: case OP_BIND_VBUF: case OP_BIND_VBUF1:
+            case OP_BIND_PROGRAM: { const p_bind_program *b = a; state_op(h.op,a);
+                if (passdump()) { NvProg *pr = S.prog ? prog_get(S.prog) : nil;
+                    PD("Y bind_program cmd=%p program=%p stages=%#x vs=%s fs=%s gs=%s",p,(void *)b->prog,b->stages,
+                       pr && pr->sh[0] ? pr->sh[0]->key : "-", pr && pr->sh[1] ? pr->sh[1]->key : "-",pr && pr->sh[2] ? pr->sh[2]->key : "-"); }
+                break; }
+            case OP_SCISSOR: case OP_BIND_STATE: case OP_BIND_VATTR: case OP_BIND_VSTREAM: case OP_BIND_VBUF: case OP_BIND_VBUF1:
            case OP_BIND_UBO: case OP_BIND_UBO_SMALL: case OP_BIND_SSBO: case OP_BIND_TEX: case OP_BIND_TEX1: case OP_BIND_SAMP: case OP_BIND_IMG: case OP_STENCIL:
            case OP_POLY_OFFSET: case OP_VIEWPORTS: case OP_DEPTH_RANGE: case OP_DEPTH_CLAMP: case OP_VP_SWIZZLE: case OP_BLEND_COLOR: case OP_COPY_STRIDE:
                state_op(h.op, a); break;
            case OP_TF_BIND: { const p_tf_bind *b = a; PD("Y tf bind %d addr=%#llx size=%llu", b->index, (unsigned long long)b->addr, (unsigned long long)b->size); state_op(h.op, a); break; }
-            case OP_TF_STATE: { const p_tf_state *b = a; static const char *const w[4] = {"begin", "end", "pause", "resume"}; PD("Y tf %s ctl=%#llx written %llu/%llu/%llu/%llu", w[b->what & 3], (unsigned long long)b->ctl,
-                (unsigned long long)S.tf_off[0], (unsigned long long)S.tf_off[1], (unsigned long long)S.tf_off[2], (unsigned long long)S.tf_off[3]); state_op(h.op, a); break; }
+            case OP_TF_STATE: { const p_tf_state *b = a; static const char *const w[4] = {"begin", "end", "pause", "resume"}; PD("Y tf %s ctl=%#llx", w[b->what & 3], (unsigned long long)b->ctl);
+                state_op(h.op, a); tf_state_gpu(b); break; }
            case OP_PRIM_RESTART: state_op(h.op, a); break;
            case OP_RAST_DISCARD: PD("Y rasterizer discard %d", ((const p_rast_discard *)a)->on); state_op(h.op, a); break;
            case OP_COPY_B2T: copy_b2t(a); break;
@@ -9575,7 +9565,7 @@ void mtl_exec(const uint8_t *p) {
            case OP_DISPATCH: if (g_csh_n && csh_take(p)) { PD("H %llu dispatch already run ahead of the pass", g_stat_pass); break; }
                dispatch(a); break;
            case OP_BARRIER: { int32_t bits; memcpy(&bits, a, 4);   // ordered at the next draw pass (barrier_game); inside a pass with shader writes, the pass ends first (swrite_note)
-                int split = g_enc && g_pass_swrite && (bits & (NVN_BARRIER_ORDER_FRAGMENTS_BIT | NVN_BARRIER_ORDER_FRAGMENTS_TILED_BIT));
+                int split = g_enc && g_pass_swrite && (bits & (NVN_BARRIER_ORDER_PRIMITIVES_BIT | NVN_BARRIER_ORDER_FRAGMENTS_BIT | NVN_BARRIER_ORDER_FRAGMENTS_TILED_BIT));
                PD("I %llu barrier bits %#x%s", g_stat_pass, bits, split ? ": ends the pass (shader writes)" : g_enc ? " (render pass open)" : "");
                if (split) { rp_end(); static int once; if (!once++) MLOG("a game barrier (bits %#x) after shader writes in a draw pass: the pass ends there (first case pass %llu)", bits, (unsigned long long)g_stat_pass); }
                if (bits) g_game_barrier = 1; break; }
@@ -68,8 +68,8 @@ t_vbuf
 t_vbuf.dSYM/
 t_geometry
 t_geometry.dSYM/
-t_feedback
-t_feedback.dSYM/
+t_transform_feedback
+t_transform_feedback.dSYM/
 t_dirty_race
 t_dirty_race.dSYM/
 t_vin
@@ -130,3 +130,7 @@ t_present_snap
 t_present_snap.dSYM/
 t_link_ticks
 t_link_ticks.dSYM/
+t_rast_discard
+t_rast_discard.dSYM/
+t_feedback
+t_feedback.dSYM/
@@ -20,6 +20,12 @@ FW=-framework Metal -framework MetalFX -framework Foundation -framework QuartzCo
 # Tests that #include nvn_mtl.m link the rest of the NVN layer beside it (hle_nvn.c: memory pools, sparse pools, serials) and the host's large stack, with weak stubs of what they need of the host (t_host_stubs.c)
 MTL_T_DEP=../nvn/hle_nvn.c ../../host/crash.c ../../host/hle_bigstack.c t_host_stubs.c $(GFX) $(GFX_H)
 MTL_T_SRC=../nvn/hle_nvn.c ../../host/crash.c ../../host/hle_bigstack.c t_host_stubs.c $(filter-out ../nvn/nvn_mtl.m,$(GFX))
+t_transform_feedback: t_transform_feedback.m $(MTL_T_DEP)
+	clang $(CFLAGS) -fobjc-arc -o $@ t_transform_feedback.m $(MTL_T_SRC) $(FW)
+.PHONY: check_t_transform_feedback
+check_t_transform_feedback: t_transform_feedback
+	DM_TEST_TRANSFORM_FEEDBACK=$(CURDIR)/t_transform_feedback MTL_DEBUG_LAYER=1 python3 -m unittest test_transform_feedback.py
+run: check_t_transform_feedback
 t_nvn: t_nvn.c ../nvn/hle_nvn.c ../../host/crash.c ../../host/hle_bigstack.c $(GFX) $(GFX_H)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_nvn.c ../../host/crash.c ../../host/hle_bigstack.c $(GFX) $(FW)
 t_dirty_slots: t_dirty_slots.c ../nvn/nvn_dirty.c ../../host/crash.c ../nvn/nvn_dirty.h
@@ -141,6 +147,9 @@ t_geometry: t_geometry.m $(MTL_T_DEP)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_geometry.m $(MTL_T_SRC) $(FW)
 t_feedback: t_feedback.m $(MTL_T_DEP)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_feedback.m $(MTL_T_SRC) $(FW)
+t_rast_discard: t_rast_discard.m $(MTL_T_DEP)
+	clang $(CFLAGS) -fobjc-arc -o $@ t_rast_discard.m $(MTL_T_SRC) $(FW)
+run: t_rast_discard
 .PHONY: check_t_feedback
 check_t_feedback: t_feedback
 	MTL_DEBUG_LAYER=1 ./t_feedback
@@ -291,5 +300,6 @@ 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_upgrade_alias
 	MTL_DEBUG_LAYER=1 ./t_image_usage
 	MTL_DEBUG_LAYER=1 ./t_stencil_view
+	MTL_DEBUG_LAYER=1 ./t_rast_discard
 	DM_GAME=$(GAME) python3 test_tbdr_budget.py
 .PHONY: run
@@ -0,0 +1,80 @@
+// Discard suppresses fragments while array, indexed and instanced vertex writes still execute.
+#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, fails;
+#define CHECK(c, label) do { checks++; if (!(c)) { fails++; fprintf(stderr,"FAIL: %s\n",label); } } while (0)
+static uint8_t *cmd; static size_t used;
+static void emit(unsigned op, const void *p, size_t n) {
+    size_t z = (8 + n + 7) & ~7ull; op_hdr h = {op, (uint32_t)z};
+    memcpy(cmd + used, &h, 8); if (n) memcpy(cmd + used + 8, p, n); used += z;
+}
+static void discard(int on) { p_rast_discard d = {.on=on}; emit(OP_RAST_DISCARD,&d,sizeof d); }
+static void clear_blue(void) { p_clear_color c = {.rgba={0,0,1,1},.mask=15}; emit(OP_CLEAR_COLOR,&c,sizeof c); }
+static void execute(void) {
+    emit(OP_END,NULL,0); mtl_sub_begin(); uint64_t handle=(uint64_t)cmd; mtl_sub_index(&handle,1); mtl_exec(cmd);
+    uint64_t fence=mtl_sub_commit(); CHECK(mtl_wait(fence,5000000000ull),"GPU submission completes");
+    CHECK(!mtl_last_commit_error(),"GPU reports no error"); used=0;
+}
+static void pixels(const tex_t *t, int red) {
+    mtl_sub_begin(); id<MTLBuffer> b=[g_dev newBufferWithLength:256 options:MTLResourceStorageModeShared]; res_add(b);
+    enc_b_copy_texture_buffer(blit(),tex_mtl(t),0,0,MTLOriginMake(0,0,0),MTLSizeMake(8,8,1),b,0,32,256);
+    uint64_t fence=mtl_sub_commit(); CHECK(mtl_wait(fence,5000000000ull),"readback completes");
+    uint8_t *p=b.contents; int good=1;
+    for(int i=0;i<64;i++) good &= p[i*4]==(red?255:0) && p[i*4+1]==0 && p[i*4+2]==(red?0:255) && p[i*4+3]==255;
+    CHECK(good,red?"normal draw rasterizes every pixel red":"discard leaves every pixel at the blue clear");
+}
+int main(void) { @autoreleasepool {
+    setenv("DM_NVN_TBDR","0",1); mtl_init();
+    void *memory=NULL; if(posix_memalign(&memory,getpagesize(),65536))abort(); memset(memory,0,65536); mtl_pool_add(memory,65536);
+    cmd=memory; uint32_t *output=(uint32_t *)((uint8_t *)memory+32768);
+    uint32_t *indices=(uint32_t *)((uint8_t *)memory+36864); indices[0]=7;indices[1]=8;indices[2]=9;
+    NSError *error=nil; id<MTLLibrary> lib=[g_dev newLibraryWithSource:@"#include <metal_stdlib>\nusing namespace metal;\n"
+        "struct A{device uint* p[[id(0)]];};\n"
+        "vertex float4 v(uint n[[vertex_id]],uint i[[instance_id]],constant A&a[[buffer(0)]]){a.p[n+16*i]=0x500+n+16*i;float2 p[3]={float2(-1,-1),float2(3,-1),float2(-1,3)};return float4(p[n%3],0,1);}\n"
+        "vertex void vd(uint n[[vertex_id]],uint i[[instance_id]],constant A&a[[buffer(0)]]){a.p[n+16*i]=0x500+n+16*i;}\n"
+        "vertex float4 p(uint n[[vertex_id]]){return float4(0,0,0,1);}\n"
+        "vertex float4 c(uint n[[vertex_id]]){return float4(0,0,0,1);}\n"
+        "vertex void pd(uint n[[vertex_id]],constant A&a[[buffer(0)]]){a.p[n]=n+0x1200;}\n"
+        "vertex void cd(uint n[[vertex_id]],constant A&a[[buffer(0)]]){a.p[n+4096]=a.p[n];}\n"
+        "fragment float4 f(){return float4(1,0,0,1);}\n" options:nil error:&error];
+    if(!lib){fprintf(stderr,"%s\n",error.localizedDescription.UTF8String);return 1;}
+    static prog_t pg={.serial=910501}; NvProg *program=[NvProg new];
+    NvShader *vs=[NvShader new];vs->fn=[lib newFunctionWithName:@"v"];vs->ae=[vs->fn newArgumentEncoderWithBufferIndex:0];
+    vs->discard_fn=[lib newFunctionWithName:@"vd"];
+    vs->rf.nb=1;vs->rf.b[0]=(nvn_rbind){.kind=NVN_RB_SBUFP,.rw=1,.reg=0,.aid=0};
+    program->sh[0]=vs;program->sh[1]=[NvShader new];program->sh[1]->fn=[lib newFunctionWithName:@"f"];program->sh[1]->stage=1;serial_put(pg.serial,program);
+    tex_t t={.magic=MAGIC('T','E','X',0),.serial=910502};t.d=(nvn_texdesc){.target=NVN_TEXTURE_TARGET_2D,.format=NVN_FORMAT_RGBA8,.width=8,.height=8,.depth=1,.levels=1,.samples=1};
+    state_reset();memset(&R,0,sizeof R);
+    p_set_rt rt={.n=1};rt.color[0]=&t;emit(OP_SET_RT,&rt,sizeof rt);
+    p_bind_program bp={.prog=&pg,.stages=3};emit(OP_BIND_PROGRAM,&bp,sizeof bp);
+    p_bind_buf buf={.stage=0,.index=0,.addr=(uint64_t)output,.size=256};emit(OP_BIND_SSBO,&buf,sizeof buf);
+    ds_st ds={.df=NVN_DEPTH_FUNC_ALWAYS};p_bind_state bs={.kind=5};memcpy(bs.bytes,&ds,sizeof ds);emit(OP_BIND_STATE,&bs,sizeof bs);
+    clear_blue();execute();
+    p_scissor sc={0};emit(OP_SCISSOR,&sc,sizeof sc);
+    poly_st poly={.cull=NVN_FACE_FRONT_AND_BACK};bs.kind=4;memcpy(bs.bytes,&poly,sizeof poly);emit(OP_BIND_STATE,&bs,sizeof bs);
+    discard(1);p_draw arrays={.prim=NVN_DRAW_PRIMITIVE_TRIANGLES,.first=4,.count=3,.instances=1};emit(OP_DRAW,&arrays,sizeof arrays);
+    p_draw_idx indexed={.prim=NVN_DRAW_PRIMITIVE_TRIANGLES,.itype=2,.count=3,.ibuf=(uint64_t)indices,.instances=2};emit(OP_DRAW_IDX,&indexed,sizeof indexed);
+    execute();pixels(&t,0);
+    CHECK(output[4]==0x504 && output[5]==0x505 && output[6]==0x506,"discarded array draw writes its three vertex IDs");
+    CHECK(output[7]==0x507 && output[9]==0x509 && output[23]==0x517 && output[25]==0x519,"discarded indexed draw writes both instances");
+    sc.w=sc.h=8;emit(OP_SCISSOR,&sc,sizeof sc);poly.cull=0;memcpy(bs.bytes,&poly,sizeof poly);emit(OP_BIND_STATE,&bs,sizeof bs);
+    discard(0);arrays.first=0;emit(OP_DRAW,&arrays,sizeof arrays);execute();pixels(&t,1);
+    clear_blue();discard(1);emit(OP_DRAW,&arrays,sizeof arrays);execute();pixels(&t,0);
+    // ORDER_PRIMITIVES must make one vertex shader's storage writes visible
+    // to the next vertex shader, even with rasterization disabled throughout.
+    static prog_t pp={.serial=910503},cp={.serial=910504};NvProg *producers[2]={[NvProg new],[NvProg new]};
+    for(int i=0;i<2;i++){NvShader *s=[NvShader new];s->fn=[lib newFunctionWithName:i?@"c":@"p"];s->discard_fn=[lib newFunctionWithName:i?@"cd":@"pd"];
+        s->ae=[s->discard_fn newArgumentEncoderWithBufferIndex:0];s->rf.nb=1;s->rf.b[0]=(nvn_rbind){.kind=NVN_RB_SBUFP,.rw=1,.reg=0,.aid=0};producers[i]->sh[0]=s;}
+    serial_put(pp.serial,producers[0]);serial_put(cp.serial,producers[1]);
+    buf.size=32768;emit(OP_BIND_SSBO,&buf,sizeof buf);bp.prog=&pp;emit(OP_BIND_PROGRAM,&bp,sizeof bp);
+    arrays=(p_draw){.prim=0,.count=4096,.instances=1};emit(OP_DRAW,&arrays,sizeof arrays);
+    int32_t barrier=NVN_BARRIER_ORDER_PRIMITIVES_BIT;emit(OP_BARRIER,&barrier,sizeof barrier);
+    bp.prog=&cp;emit(OP_BIND_PROGRAM,&bp,sizeof bp);emit(OP_DRAW,&arrays,sizeof arrays);execute();
+    int ordered=1;for(int i=0;i<4096;i++)ordered&=output[i+4096]==(uint32_t)(i+0x1200);
+    CHECK(ordered,"primitive barrier orders vertex storage writes before the consuming draw");
+    printf("rasterizer discard: %d checks, %d failed\n",checks,fails);return fails!=0;
+} }
@@ -32,6 +32,8 @@ int main(void) { @autoreleasepool {
        s.d = (nvn_texdesc){.target=NVN_TEXTURE_TARGET_2D, .format=formats[i], .width=W, .height=H, .depth=1, .levels=1, .samples=1};
        nvn_tex_layout(&s.d, &s.L);
        tex_t sb = s, d = s; sb.serial++; sb.stencil = 1; d.serial += 2;   // sb: the same storage built in stencil mode; d: the copy's destination
+        sb.d.flags |= NVN_TEXTURE_FLAGS_IMAGE;   // usage-only difference, as in Xenoblade Chronicles 3
+        d.off = 0x100000;   // a separate destination, not another descriptor of the source storage
        mtl_tex_init(&s); mtl_tex_init(&sb); mtl_tex_init(&d);
        id<MTLTexture> base = tex_mtl(&s), dst = tex_mtl(&d);
        mtl_drain();   // finish production initialization before the fixture writes on its independent classic queue
@@ -42,6 +44,12 @@ int main(void) { @autoreleasepool {
        for (int y = 0; y < H; y++) for (int x = 0; x < W; x++) { dz[x + y * W] = depth_at(x, y); st[x + y * W] = stencil_at(x, y); }
        plane_io(base, 0, 1, dz); plane_io(base, 1, 1, st);

+        mtl_sub_begin();
+        id<MTLTexture> image = tex_mtl(&sb);
+        mtl_sub_commit(); mtl_drain();   // adding IMAGE to a sampling descriptor retains both populated planes
+        base = tex_get(s.serial);
+        CHECK(image == base, "IMAGE usage and render descriptors share their populated storage");
+
        // A stencil-mode view, and the texture built in stencil mode, are X32_Stencil8 views of the depth-stencil storage itself (no upgraded copy)
        view_t v = {.set_dsmode=1, .dsmode=NVN_TEXTURE_DEPTH_STENCIL_MODE_STENCIL};
        mtl_reg_tex(800 + 2 * i, &s, &v); mtl_reg_tex(801 + 2 * i, &sb, NULL);
@@ -0,0 +1,89 @@
+// Production feedback replay, append/resume, overflow and GPU-count gathering.
+#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, fails;
+#define CHECK(c,label) do { checks++; if(!(c)){fails++;fprintf(stderr,"FAIL: %s\n",label);} }while(0)
+static void finish(void) { CHECK(mtl_wait(mtl_sub_commit(),5000000000ull),"submission completes");CHECK(!mtl_last_commit_error(),"no GPU error"); }
+static NvShader *shader(NSString *path) {
+    NSError *e=nil;NSString *src=[NSString stringWithContentsOfFile:path encoding:NSUTF8StringEncoding error:&e];
+    id<MTLLibrary> lib=[g_dev newLibraryWithSource:src options:nil error:&e];
+    if(!lib){fprintf(stderr,"%s\n",e.localizedDescription.UTF8String);abort();}
+    NvShader *vs=[NvShader new];vs->fn=[lib newFunctionWithName:@"main0"];vs->tf_fn=[lib newFunctionWithName:@"main0_tf"];vs->tf_stride[0]=20;
+    return vs;
+}
+static void state(int what,void *control) {p_tf_state s={.what=what,.ctl=(uint64_t)control};state_op(OP_TF_STATE,&s);tf_state_gpu(&s);}
+static void bind(void *output,size_t size) {p_tf_bind b={.index=0,.addr=(uint64_t)output,.size=size};state_op(OP_TF_BIND,&b);}
+static void record(const float *p,unsigned n) {const float want[5]={n,n+10,n+40,n+50,n+60};CHECK(!memcmp(p,want,sizeof want),"interleaved captured values and order");}
+static uint8_t *commands;static size_t used;
+static void emit(unsigned op,const void *p,size_t n) {size_t z=(8+n+7)&~7ull;op_hdr h={op,(uint32_t)z};memcpy(commands+used,&h,8);if(n)memcpy(commands+used+8,p,n);used+=z;}
+int main(int argc,char **argv) { @autoreleasepool {
+    if(argc!=2){fprintf(stderr,"pass the fixture made by test_transform_feedback.py --fixture\n");return 2;}
+    setenv("DM_NVN_TBDR","0",1);mtl_init();state_reset();memset(&R,0,sizeof R);
+    void *memory=NULL;if(posix_memalign(&memory,getpagesize(),65536))abort();memset(memory,0,65536);mtl_pool_add(memory,65536);
+    float *output=(float *)memory;uint32_t *ctl=(uint32_t *)((uint8_t *)memory+4096),*ctl2=ctl+16,*ctl3=ctl+32;
+    NvProg *pr=[NvProg new];pr->sh[0]=shader(@(argv[1]));
+    p_draw dw={.prim=0,.first=2,.count=3,.instances=1};S.rast_discard=1;
+    mtl_sub_begin();bind(output,100);state(0,ctl);tf_draw(pr,&dw,NULL);state(2,ctl);finish();
+    CHECK(ctl[0]==60&&ctl[4]==3&&ctl[5]==1,"pause saves byte offsets and draw count");
+    for(unsigned i=0;i<3;i++)record(output+5*i,i+2);
+    mtl_sub_begin();state(3,ctl);dw.first=9;tf_draw(pr,&dw,NULL);state(1,ctl2);finish();
+    CHECK(ctl2[0]==100&&ctl2[4]==5,"resume appends only the vertices that fit");
+    record(output+15,9);record(output+20,10);CHECK(output[25]==0,"overflow leaves the next record untouched");
+    // Source count is zero on the CPU until an earlier GPU copy supplies five.
+    // Whole-triangle truncation accepts three, in packed-index order.
+    uint32_t *indices=(uint32_t *)((uint8_t *)memory+8192),*args=indices+32,*source=args+16;
+    const uint32_t order[5]={6,2,4,1,5};memcpy(indices,order,sizeof order);source[0]=5;
+    NvShader *reader=shader([@(argv[1]) stringByAppendingString:@".geo"]);
+    NSError *error=nil;id<MTLLibrary> drawlib=[g_dev newLibraryWithSource:@"#include <metal_stdlib>\nusing namespace metal;\nstruct A{device uint* p[[id(0)]];};\nvertex float4 v(uint n[[vertex_id]],constant A&a[[buffer(0)]]){return float4(0,0,0,1);}\nvertex void vd(uint n[[vertex_id]],constant A&a[[buffer(0)]]){a.p[n]=0x800+n;}\n" options:nil error:&error];
+    CHECK(drawlib!=nil,"indirect-count consumer compiles");
+    NvShader *consumer=[NvShader new];consumer->fn=[drawlib newFunctionWithName:@"v"];consumer->discard_fn=[drawlib newFunctionWithName:@"vd"];
+    consumer->ae=[consumer->discard_fn newArgumentEncoderWithBufferIndex:0];consumer->rf.nb=1;consumer->rf.b[0]=(nvn_rbind){.kind=NVN_RB_SBUFP,.rw=1,.reg=0,.aid=0};
+    NvProg *drawpr=[NvProg new];drawpr->sh[0]=consumer;static prog_t pg={.serial=910601};serial_put(pg.serial,drawpr);
+    tex_t target={.magic=MAGIC('T','E','X',0),.serial=910602};target.d=(nvn_texdesc){.target=NVN_TEXTURE_TARGET_2D,.format=NVN_FORMAT_RGBA8,.width=8,.height=8,.depth=1,.levels=1,.samples=1};
+    memset(output,0,256);mtl_sub_begin();bind(output,100);state(0,ctl3);
+    void *mb;uint64_t off,gpu;mtl_gpu_addr((uint64_t)args,4,&mb,&off,&gpu);
+    void *sb;uint64_t so,sg;mtl_gpu_addr((uint64_t)source,4,&sb,&so,&sg);
+    enc_b_copy_buffer(blit(),(__bridge id<MTLBuffer>)sb,so,(__bridge id<MTLBuffer>)mb,off,4);
+    void *ib;uint64_t io,ig;mtl_gpu_addr((uint64_t)indices,sizeof order,&ib,&io,&ig);
+    tf_geometry(reader,ig,gpu,3,0);state(1,ctl3);
+    // Consume the saved count in the same submission, without a CPU readback.
+    commands=(uint8_t *)memory+32768;uint32_t *written=(uint32_t *)((uint8_t *)memory+16384);
+    p_set_rt rt={.n=1};rt.color[0]=&target;emit(OP_SET_RT,&rt,sizeof rt);
+    p_bind_program bp={.prog=&pg,.stages=1};emit(OP_BIND_PROGRAM,&bp,sizeof bp);
+    p_bind_buf ssbo={.stage=0,.index=0,.addr=(uint64_t)written,.size=64};emit(OP_BIND_SSBO,&ssbo,sizeof ssbo);
+    p_rast_discard rd={.on=1};emit(OP_RAST_DISCARD,&rd,sizeof rd);
+    p_draw draw={.prim=0,.instances=1,.ind=(uint64_t)ctl3+16};emit(OP_DRAW,&draw,sizeof draw);emit(OP_END,NULL,0);
+    mtl_exec(commands);finish();
+    CHECK(ctl3[0]==60&&ctl3[4]==3,"geometry uses the GPU count and keeps complete primitives");
+    record(output,6);record(output+5,2);record(output+10,4);CHECK(output[15]==0,"incomplete final primitive is not captured");
+    CHECK(written[0]==0x800&&written[1]==0x801&&written[2]==0x802&&written[3]==0,"saved GPU count controls the following draw");
+    // The earlier single-output, three-stream skinning layout retains its
+    // independent byte strides and offsets with the GPU counter implementation.
+    pr->sh[0]=shader([@(argv[1]) stringByAppendingString:@".split"]);
+    pr->sh[0]->tf_stride[0]=12;pr->sh[0]->tf_stride[1]=pr->sh[0]->tf_stride[2]=4;
+    float *second=output+128,*third=output+160;memset(output,0,1024);
+    mtl_sub_begin();bind(output,36);
+    p_tf_bind b1={.index=1,.addr=(uint64_t)second,.size=12},b2={.index=2,.addr=(uint64_t)third,.size=12};state_op(OP_TF_BIND,&b1);state_op(OP_TF_BIND,&b2);
+    state(0,ctl);dw.first=4;dw.count=3;tf_draw(pr,&dw,NULL);state(1,ctl);finish();
+    CHECK(ctl[0]==36&&ctl[1]==12&&ctl[2]==12&&ctl[4]==3,"three streams advance by their own byte strides");
+    const float position[9]={4,14,24,5,15,25,6,16,26},packed[3]={44,45,46};
+    CHECK(!memcmp(output,position,sizeof position),"separate position stream");
+    CHECK(!memcmp(second,packed,sizeof packed)&&!memcmp(third,packed,sizeof packed),"separate packed streams");
+    // Tropical Freeze records feedback with rasterization enabled and no colour
+    // writes. Exercise the command replay branch as well as the capture helper.
+    state_reset();memset(&R,0,sizeof R);memset(&S.chmask,0,sizeof S.chmask);
+    pr->sh[0]=shader(@(argv[1]));static prog_t tfpg={.serial=910603};serial_put(tfpg.serial,pr);
+    memset(output,0,256);used=0;mtl_sub_begin();emit(OP_SET_RT,&rt,sizeof rt);
+    bp.prog=&tfpg;emit(OP_BIND_PROGRAM,&bp,sizeof bp);
+    p_tf_bind capture={.index=0,.addr=(uint64_t)output,.size=100};emit(OP_TF_BIND,&capture,sizeof capture);
+    p_tf_state begin={.what=0,.ctl=(uint64_t)ctl},end={.what=1,.ctl=(uint64_t)ctl};
+    emit(OP_TF_STATE,&begin,sizeof begin);dw.first=1;dw.count=3;emit(OP_DRAW,&dw,sizeof dw);
+    emit(OP_TF_STATE,&end,sizeof end);emit(OP_END,NULL,0);mtl_exec(commands);finish();
+    CHECK(ctl[0]==60&&ctl[4]==3,"rasterized draw records feedback through command replay");
+    for(unsigned i=0;i<3;i++)record(output+5*i,i+1);
+    printf("feedback: %d checks, %d failed\n",checks,fails);return fails!=0;
+} }
@@ -16,6 +16,14 @@ int main(void) { @autoreleasepool {
    mtl_tex_init(&a); mtl_tex_init(&b);
    id<MTLTexture> old = tex_mtl(&a);
    CHECK(tex_mtl(&b) == old && tex_mtl(&ghost) == old, "live and finalized serials initially share storage");
+    tex_t image = a; image.serial += 3; image.d.flags |= NVN_TEXTURE_FLAGS_IMAGE; mtl_tex_init(&image);
+    CHECK(tex_mtl(&image) == old, "an IMAGE sampling descriptor shares storage without a usage change");
+    mtl_reg_tex(950, &image, NULL); S.img[5][0] = 950;
+    NvShader *writer = [NvShader new]; writer->rf.nb = 1; writer->rf.b[0] = (nvn_rbind){.kind=NVN_RB_IMG, .rw=1};
+    NvProg *program = [NvProg new]; program->sh[5] = writer; image_prepare(program, 1u << 5);
+    old = tex_mtl(&image);
+    CHECK(old.usage & MTLTextureUsageShaderWrite, "IMAGE descriptor adds shader-write usage");
+    CHECK(tex_mtl(&a) == old && tex_mtl(&b) == old && tex_mtl(&ghost) == old, "IMAGE usage does not split storage across existing descriptors");
    // Exercise replacement at the hash table's growth threshold: reinserting while iterating would resize and reshuffle it.
    uint64_t extra = 99400;
    while ((g_by_serial.n + 1) * 4 <= g_by_serial.cap * 3) { hm_put_id(&g_by_serial, &extra, old); extra++; }
@@ -0,0 +1,43 @@
+"""Discard entry points preserve vertex side effects and compile on both platforms."""
+import ast
+from pathlib import Path
+import re
+import shutil
+import subprocess
+import tempfile
+import unittest
+
+path = Path(__file__).resolve().parents[1] / 'tools/dmsh_build.py'
+node = next(n for n in ast.parse(path.read_text()).body if isinstance(n, ast.FunctionDef) and n.name == 'vs_discard')
+scope = {'re': re}
+exec(compile(ast.Module(body=[node], type_ignores=[]), str(path), 'exec'), scope)
+
+SOURCE = '''#include <metal_stdlib>
+using namespace metal;
+struct main0_out { float4 p [[position]]; };
+struct Args { device uint* values [[id(0)]]; };
+vertex main0_out main0(uint n [[vertex_id]], constant Args& a [[buffer(0)]])
+{
+    main0_out out = {};
+    a.values[n] = n + 17;
+    if (n == 0) return out;
+    a.values[n] = n + 29;
+    return out;
+}
+'''
+
+
+class DiscardTests(unittest.TestCase):
+    @unittest.skipUnless(shutil.which('xcrun'), 'requires Metal compilers')
+    def test_side_effects_before_both_returns_compile_as_void_vertex(self):
+        src = SOURCE + scope['vs_discard'](SOURCE)
+        with tempfile.TemporaryDirectory() as directory:
+            root = Path(directory); msl = root / 'discard.metal'; msl.write_text(src)
+            for sdk in ('macosx', 'iphoneos'):
+                with self.subTest(sdk=sdk):
+                    r = subprocess.run(['xcrun', '-sdk', sdk, 'metal', '-c', '-fmodules-cache-path=' + str(root / sdk), str(msl), '-o', str(root / (sdk + '.air'))], capture_output=True, text=True)
+                    self.assertEqual(r.returncode, 0, r.stderr)
+
+
+if __name__ == '__main__':
+    unittest.main()
@@ -0,0 +1,130 @@
+"""Monolith archive boundaries and multi-stage programs, without game data."""
+from pathlib import Path
+import struct
+import sys
+import unittest
+import zlib
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / 'tools'))
+import shader_split as split
+
+
+def program(stages, extra=0):
+    controls, codes = [], []
+    for stage in stages:
+        ctl, code = bytearray(0x880 + extra), 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.GFX_MAGIC
+        struct.pack_into('<I', code, 0x30, (split.STAGES.index(stage) + 1) << 10 | 1)
+        controls.append(bytes(ctl)); codes.append(bytes(code))
+    # One SLCT, two identical variants; each vertex bundle includes geometry.
+    slct = bytearray(72 + 16 + 2 * 52)
+    slct[:4] = b'SLCT'
+    struct.pack_into('<2I', slct, 16, 2, 72)
+    stored = [(len(c) + 7) & ~7 for c in controls]
+    struct.pack_into('<4I', slct, 40, 0, sum(stored) * 2, 0, sum(map(len, codes)) * 2)
+    for i in range(2):
+        p = 88 + 52 * i
+        struct.pack_into('<2I', slct, 72 + 8 * i, p, 52)
+        struct.pack_into('<2I', slct, p + 24, 32, 1)
+        vs = sum(len(c) for s, c in zip(stages, codes) if s != 'ps')
+        vc = sum(n for s, n in zip(stages, stored) if s != 'ps')
+        struct.pack_into('<5I', slct, p + 32, 1, vs, len(codes[-1]), vc, stored[-1])
+    head = bytearray(80); head[:4] = b'HCPS'
+    struct.pack_into('<I', head, 4, 1010)
+    struct.pack_into('<2I', head, 8, 72, 1)
+    ctl = b''.join(c.ljust(n, b'\0') for c, n in zip(controls, stored)) * 2; code = b''.join(codes) * 2
+    struct.pack_into('<6I', head, 24, 80, len(slct), 80 + len(slct) + len(ctl), len(code), 80 + len(slct), len(ctl))
+    return bytes(head) + slct + ctl + code, controls, codes
+
+
+class MonolithTests(unittest.TestCase):
+    def test_text_controls_advance_at_eight_byte_alignment(self):
+        data, ctls, codes = program(['vs', 'ps'], extra=13)
+        out = split.hcps_shaders(data, 0, '/test')
+        self.assertEqual([(s, c, k) for _, s, c, k in out], list(zip(['vs', 'ps'], ctls, codes)) * 2)
+
+    def test_embedded_geometry_and_variant_pairs(self):
+        data, ctls, codes = program(['vs', 'gs', 'ps'])
+        out = split.hcps_shaders(data, 0, '/test')
+        self.assertEqual([(s, c, k) for _, s, c, k in out], list(zip(['vs', 'gs', 'ps'], ctls, codes)) * 2)
+        self.assertEqual(split.var_pairs(out), {(split.skey(ctls[0], codes[0]), split.skey(ctls[2], codes[2]))})
+
+    def test_xbc_payload_is_scanned_only_after_decoding(self):
+        data = program(['vs', 'ps'])[0]
+        # A stored deflate block leaves HCPS visible inside the compressed stream.
+        packed = zlib.compress(data, level=0)
+        wrapper = struct.pack('<4s4I28s', b'xbc1', 1, len(data), len(packed), 0, b'test') + packed
+        self.assertIn(b'HCPS', packed)
+        out = []; split.monolith_shaders(wrapper, '/test', out)
+        self.assertEqual(len(out), 4)
+
+    def test_standalone_slct_uses_its_own_relative_offsets(self):
+        data, ctls, codes = program(['vs', 'ps'])
+        body = bytearray(data[80:])
+        struct.pack_into('<I', body, 4, 0x1003)
+        c = len(data[80:80 + 192])
+        # The synthetic fixture's SLCT occupies 192 bytes; control and code
+        # sections now use offsets from that SLCT rather than separate bases.
+        struct.pack_into('<I', body, 40, c)
+        struct.pack_into('<I', body, 48, c + sum(map(len, ctls)) * 2)
+        out = []; split.scan_blob(bytes(body), 'main', out)
+        self.assertEqual([(s, ctl, code) for _, s, ctl, code in out], list(zip(['vs', 'ps'], ctls, codes)) * 2)
+
+    def test_effect_glslc_output_inside_xbc(self):
+        _, controls, codes = program(['vs', 'ps'])
+        data = bytearray(0x70) + controls[0] + codes[0]
+        data[:4] = bytes.fromhex('91688619')
+        struct.pack_into('<I', data, 0x44, len(data))
+        packed = zlib.compress(data, level=0)
+        wrapper = struct.pack('<4s4I28s', b'xbc1', 1, len(data), len(packed), 0, b'effect') + packed
+        out = []; split.monolith_shaders(wrapper, '/effect', out)
+        self.assertEqual([(s, c, k) for _, s, c, k in out], [('vs', controls[0], codes[0])])
+
+    def test_effect_cmp_zlib_nested_inside_xbc(self):
+        data, controls, codes = program(['vs', 'ps'])
+        packed = zlib.compress(data, level=0)
+        cmp = struct.pack('<8s2I', b'cmp\0zlib', len(packed), len(data)) + packed
+        packed = zlib.compress(cmp, level=0)
+        xbc = struct.pack('<4s4I28s', b'xbc1', 1, len(cmp), len(packed), 0, b'effect') + packed
+        out = []; split.monolith_shaders(xbc, '/effect', out)
+        self.assertEqual([(s, c, k) for _, s, c, k in out], list(zip(['vs', 'ps'], controls, codes)) * 2)
+
+    def test_effect_cmp_zlib_checks_both_sizes(self):
+        packed = zlib.compress(b'test')
+        with self.assertRaisesRegex(SystemExit, 'stored/raw sizes'):
+            split.monolith_cmp_data(struct.pack('<8s2I', b'cmp\0zlib', len(packed), 5) + packed, 0)
+
+    def test_source_only_hcps_has_no_nvn_binaries(self):
+        data = bytearray(program(['vs', 'ps'])[0])
+        struct.pack_into('<4I', data, 32, 0, 0, 0, 0)
+        self.assertEqual(split.hcps_shaders(data, 0, '/source'), [])
+
+    def test_truncated_glslc_magic_is_not_a_program(self):
+        out = []; split.monolith_shaders(bytes.fromhex('91688619'), '/truncated', out)
+        self.assertEqual(out, [])
+
+    def test_bad_section_bounds_are_rejected(self):
+        data = bytearray(program(['vs', 'ps'])[0])
+        struct.pack_into('<I', data, 36, len(data))
+        with self.assertRaisesRegex(SystemExit, 'out of bounds'):
+            split.hcps_shaders(data, 0, '/test')
+
+    def test_arh_wrapper_is_outside_stored_size(self):
+        strings = b'\0.bin\0' + struct.pack('<I', 0)
+        nodes = [(-1, -1)] * 98; nodes[0] = (0, -1); nodes[97] = (-1, 0)
+        trie = b''.join(struct.pack('<ii', *n) for n in nodes)
+        so, no = 48, 48 + len(strings)
+        fo = no + len(trie)
+        # A non-word-aligned suffix table leaves the last two ID bytes plain.
+        mask = bytes.fromhex('12345678')
+        enc = lambda b: bytes(x ^ mask[i & 3] for i, x in enumerate(b[:len(b) & ~3])) + b[len(b) & ~3:]
+        header = struct.pack('<4s9I8x', b'arh1', len(strings), len(nodes), so, len(strings), no, len(trie), fo, 1, int.from_bytes(mask, 'little') ^ 0xf3f35353)
+        data = header + enc(strings) + enc(trie) + struct.pack('<Q4I', 32, 100, 200, 2, 0)
+        self.assertEqual(split.arh_entries(data), [('a.bin', 32, 148)])
+
+
+if __name__ == '__main__':
+    unittest.main()
@@ -0,0 +1,76 @@
+"""Vertex subgroup masks reached through component pointers."""
+import ast
+from pathlib import Path
+import shutil
+import struct
+import subprocess
+import tempfile
+import unittest
+
+
+source = Path(__file__).resolve().parents[1] / 'tools/dmsh_build.py'
+tree = ast.parse(source.read_text())
+nodes = [n for n in tree.body if
+         isinstance(n, ast.FunctionDef) and n.name in ('spv_single_lane', '_ep_drop') or
+         isinstance(n, ast.Assign) and any(isinstance(t, ast.Name) and t.id in ('SG_OPS', 'SG_BUILTINS') for t in n.targets)]
+scope = {'struct': struct}
+exec(compile(ast.Module(body=nodes, type_ignores=[]), str(source), 'exec'), scope)
+
+ASSEMBLY = '''OpCapability Shader
+OpCapability GroupNonUniform
+OpCapability GroupNonUniformBallot
+OpMemoryModel Logical GLSL450
+OpEntryPoint Vertex %main "main" %mask %vertex %position
+OpDecorate %mask BuiltIn SubgroupLtMask
+OpDecorate %vertex BuiltIn VertexIndex
+OpDecorate %position BuiltIn Position
+%void = OpTypeVoid
+%fn = OpTypeFunction %void
+%bool = OpTypeBool
+%uint = OpTypeInt 32 0
+%int = OpTypeInt 32 1
+%float = OpTypeFloat 32
+%v4uint = OpTypeVector %uint 4
+%v4float = OpTypeVector %float 4
+%pmask = OpTypePointer Input %v4uint
+%pu = OpTypePointer Input %uint
+%pi = OpTypePointer Input %int
+%po = OpTypePointer Output %v4float
+%mask = OpVariable %pmask Input
+%vertex = OpVariable %pi Input
+%position = OpVariable %po Output
+%three = OpConstant %int 3
+%main = OpFunction %void None %fn
+%entry = OpLabel
+%vid = OpLoad %int %vertex
+%index = OpBitwiseAnd %int %vid %three
+%ptr = ACCESS %pu %mask %index
+%value = OpLoad %uint %ptr
+%f = OpConvertUToF %float %value
+%p = OpCompositeConstruct %v4float %f %f %f %f
+OpStore %position %p
+OpReturn
+OpFunctionEnd
+'''
+
+
+@unittest.skipUnless(shutil.which('spirv-as') and shutil.which('spirv-val') and shutil.which('spirv-cross'), 'requires SPIR-V tools')
+class SubgroupTests(unittest.TestCase):
+    def test_mask_component_pointer_becomes_valid_constant_extraction(self):
+        for access in ('OpAccessChain', 'OpInBoundsAccessChain'):
+            with self.subTest(access=access), tempfile.TemporaryDirectory() as directory:
+                root = Path(directory)
+                src, spv = root / 'mask.spvasm', root / 'mask.spv'
+                src.write_text(ASSEMBLY.replace('ACCESS', access))
+                subprocess.run(['spirv-as', '--target-env', 'spv1.3', str(src), '-o', str(spv)], check=True, capture_output=True)
+                self.assertTrue(scope['spv_single_lane'](spv))
+                result = subprocess.run(['spirv-val', '--target-env', 'spv1.3', str(spv)], capture_output=True, text=True)
+                self.assertEqual(result.returncode, 0, result.stderr)
+                result = subprocess.run(['spirv-cross', '--msl', str(spv)], capture_output=True, text=True)
+                self.assertEqual(result.returncode, 0, result.stderr)
+                self.assertNotIn('simd_', result.stdout)
+                self.assertFalse(scope['spv_single_lane'](spv))
+
+
+if __name__ == '__main__':
+    unittest.main()
@@ -0,0 +1,94 @@
+"""Decode interleaved feedback declarations and compile their generated entries."""
+import ast
+import os
+from pathlib import Path
+import re
+import shutil
+import struct
+import subprocess
+import sys
+import tempfile
+import unittest
+
+path = Path(__file__).resolve().parents[1] / 'tools/dmsh_build.py'
+nodes = [n for n in ast.parse(path.read_text()).body if isinstance(n, ast.FunctionDef) and n.name in ('tf_layout', 'vs_tf')]
+scope = dict(re=re, struct=struct, sys=sys, TF_SLOT=8)
+exec(compile(ast.Module(body=nodes, type_ignores=[]), str(path), 'exec'), scope)
+
+SOURCE = '''#include <metal_stdlib>
+using namespace metal;
+struct main0_out { float4 a [[user(locn0)]]; float4 b [[user(locn1)]]; float4 p [[position]]; };
+vertex main0_out main0(uint n [[vertex_id]]) {
+    main0_out out = {};
+    out.a = float4(n, n + 10, n + 20, n + 30);
+    out.b = float4(n + 40, n + 50, n + 60, n + 70);
+    out.p = float4(0, 0, 0, 1);
+    return out;
+}
+'''
+LAYOUT = [{'stride': 20, 'outputs': [[0, 2], [1, 3]]}, [], [], []]
+SPLIT_LAYOUT = [{'stride': 12, 'outputs': [[0, 3]]}, {'stride': 4, 'outputs': [[1, 1]]}, {'stride': 4, 'outputs': [[1, 1]]}, []]
+
+
+def fixture(geometry=False):
+    return SOURCE + scope['vs_tf'](SOURCE, LAYOUT, geometry)
+
+
+class FeedbackTests(unittest.TestCase):
+    @unittest.skipUnless(os.environ.get('DM_TEST_TRANSFORM_FEEDBACK'), 'set DM_TEST_TRANSFORM_FEEDBACK to the native t_transform_feedback executable')
+    def test_native_gpu_replay(self):
+        with tempfile.TemporaryDirectory(prefix='feedback_') as directory:
+            source = Path(directory) / 'feedback.metal'
+            source.write_text(fixture())
+            Path(str(source) + '.geo').write_text(fixture(True))
+            Path(str(source) + '.split').write_text(SOURCE + scope['vs_tf'](SOURCE, SPLIT_LAYOUT))
+            result = subprocess.run([os.environ['DM_TEST_TRANSFORM_FEEDBACK'], str(source)], capture_output=True, text=True)
+            self.assertEqual(result.returncode, 0, result.stdout + result.stderr)
+    def test_component_counts_are_bytes_and_stream_order_is_preserved(self):
+        control = bytearray(0x880)
+        struct.pack_into('<H', control, 0x6dc, 20)
+        control[0x6e8] = 2
+        control[0xdc:0xde] = bytes([32, 33])
+        control[0x4dc:0x4de] = bytes([2, 3])
+        self.assertEqual(scope['tf_layout']('test', control), LAYOUT)
+
+    def test_interleaved_stores_keep_bit_patterns_and_stride(self):
+        source = fixture()
+        for offset, value in enumerate(('a.x', 'a.y', 'b.x', 'b.y', 'b.z')):
+            self.assertIn(f'dm_linear * 5 + {offset}] = as_type<uint>(out.{value})', source)
+        self.assertIn('dm_tf.offset[0] / 4', source)
+        self.assertIn('dm_linear < dm_tf.args.x', source)
+
+    def test_three_separate_single_output_streams(self):
+        control = bytearray(0x880)
+        struct.pack_into('<4H', control, 0x6dc, 12, 4, 4, 0)
+        for b, (location, components) in enumerate(((32, 3), (33, 1), (33, 1))):
+            control[0x6e8 + b] = 1
+            control[0xdc + b * 128] = location
+            control[0x4dc + b * 128] = components
+        self.assertEqual(scope['tf_layout']('test', control), SPLIT_LAYOUT)
+
+    def test_geometry_gathers_after_calculating_linear_capture_index(self):
+        source = fixture(True)
+        self.assertLess(source.index('uint dm_linear'), source.index('n = dm_tf.order[dm_linear]'))
+
+    @unittest.skipUnless(shutil.which('xcrun'), 'requires Metal compilers')
+    def test_vertex_and_geometry_entries_compile_for_both_platforms(self):
+        with tempfile.TemporaryDirectory() as directory:
+            root = Path(directory)
+            for geometry in (False, True):
+                source = root / ('geometry.metal' if geometry else 'vertex.metal')
+                source.write_text(fixture(geometry))
+                for sdk in ('macosx', 'iphoneos'):
+                    with self.subTest(geometry=geometry, sdk=sdk):
+                        result = subprocess.run(['xcrun', '-sdk', sdk, 'metal', '-c', '-fmodules-cache-path=' + str(root / sdk), str(source), '-o', str(root / (sdk + '.air'))], capture_output=True, text=True)
+                        self.assertEqual(result.returncode, 0, result.stderr)
+
+
+if __name__ == '__main__':
+    if len(sys.argv) == 3 and sys.argv[1] == '--fixture':
+        Path(sys.argv[2]).write_text(fixture())
+        Path(sys.argv[2] + '.geo').write_text(fixture(True))
+        Path(sys.argv[2] + '.split').write_text(SOURCE + scope['vs_tf'](SOURCE, SPLIT_LAYOUT))
+    else:
+        unittest.main()
@@ -354,7 +354,7 @@ def job_for(key, name, rec, w):
        if any(code[0x30:0x100]): sys.exit(f'[RED] {key}: compute code area +0x30..+0x100 is not all zero, so the instructions do not start at +0x100')
        code = code[0x100:]
    cp = os.path.join(w, key + '.bin'); open(cp, 'wb').write(code)
-    tf = tf_layout(key, rec[ci:]) if stage == 'vs' else None
+    tf = tf_layout(key, rec[ci:]) if stage in ('vs', 'gs') else None
    j = {'key': key, 'stage': stage, 'code': cp, 'tex': {}, 'img': {}, 'attr': {}, 'fout': {}, 'c1': c1.hex(), **({'tf': tf} if tf else {}),
         'debug': os.environ.get('XLATE_DEBUG', '1') != '0', 'reduced': os.environ.get('XLATE_REDUCED', '1') == '1'}   # translation options (see dmsh_xlate/Program.cs); relaxed precision is on by default (GPU float last-bit error does not enter game state, whole frame −26% on the second Mac)   # c1: used by the translator to read indirect jump tables
    if key not in FACTS: j['no_facts'] = True; return j, c1   # not in the answer table: queries use defaults (listed separately in the summary)
@@ -363,45 +363,62 @@ def job_for(key, name, rec, w):

 TF_SLOT = 8   # Metal vertex buffer index of transform feedback buffer 0 (1-3 follow): slot 0 is the argument buffer, 16-29 the vertex streams, 30 the zero buffer (nvn_mtl.m tf_draw)
 def tf_layout(key, ctl):
-    """Transform feedback layout from a vertex shader's control section (what the NVN driver programs into the GPU's feedback state when the program is bound): +0x6dc u16 stride per buffer
-    (4 buffers), +0x6e8 u8 per buffer the number of captured outputs, +0x6ec u8 buffer count and +0x6ed 1; per buffer b, +0xdc + 0x80 b lists the captured outputs, one byte each as 0x20 + n
-    (generic output n), and +0x4dc + 0x80 b their component counts, one byte each, each captured from x on and packed in that order (stride = 4 x the components). Pokémon Legends: Z-A's
-    28 skinning shaders capture one output per buffer (position xyz, packed normal / tangent bits); Donkey Kong Country: Tropical Freeze's fur simulation captures two in one buffer
-    (`7dcdaed991c74bfe`: outputs 0 and 1, xyz each, stride 24; count 2, component bytes 03 03).
-    → [[[[output, components], …], stride] per buffer, [] for an unused one], None for a shader without transform feedback. Any other form gets a red line and no capture entry point (the host
-    then reports the draw, nvn_mtl.m tf_draw): it is not known how it maps, and one shader must not stop a game's build (Tears of the Kingdom and Ace Attorney have no feedback shader at all)"""
+    """Decode the four feedback streams in the program's control section.
+    Each stream has up to 128 output-location bytes and matching component-count
+    bytes, captured from x on and interleaved in that order. Pokémon Legends:
+    Z-A uses one output per stream, Tropical Freeze two in one stream, and
+    Xenoblade Chronicles 3 uses up to eleven.
+    Unknown layouts are reported and left without a capture entry."""
    strides = struct.unpack_from('<4H', ctl, 0x6dc)
    if not any(strides): return None
    out = []
    for b in range(4):
        if not strides[b]: out.append([]); continue
-        k = ctl[0x6e8 + b]; locs = list(ctl[0xdc + 0x80 * b:0xdc + 0x80 * b + k]); comps = list(ctl[0x4dc + 0x80 * b:0x4dc + 0x80 * b + k])
-        if (not 1 <= k <= 4 or any(ctl[0xdc + 0x80 * b + k:0x15c + 0x80 * b]) or any(ctl[0x4dc + 0x80 * b + k:0x4e0 + 0x80 * b])
-                or not all(0x20 <= x < 0x40 for x in locs) or not all(1 <= n <= 4 for n in comps) or strides[b] != 4 * sum(comps)):
-            print(f'[RED] {key}: transform feedback buffer {b} has a layout of unknown form (stride {strides[b]}, outputs {k}: {[hex(x) for x in locs]}, components {comps}): no capture entry point; see tf_layout', flush=True)
+        count = ctl[0x6e8 + b]
+        locs = list(ctl[0xdc + 0x80 * b:0xdc + 0x80 * b + count])
+        sizes = list(ctl[0x4dc + 0x80 * b:0x4dc + 0x80 * b + count])
+        if not 1 <= count <= 128 or any(not 0x20 <= x < 0x40 for x in locs) or any(not 1 <= n <= 4 for n in sizes) or strides[b] != 4 * sum(sizes):
+            print(f'[RED] {key}: transform feedback buffer {b} has a layout of unknown form (stride {strides[b]}, outputs {locs}, components {sizes}): no capture entry point; see tf_layout', flush=True)
            return None
-        out.append([[[x - 0x20, n] for x, n in zip(locs, comps)], strides[b]])
+        out.append({'stride': strides[b], 'outputs': [[o - 0x20, n] for o, n in zip(locs, sizes)]})
    return out

-def vs_tf(src, tf):
-    """The transform feedback entry point main0_tf: the main function's body, ending in stores of the captured outputs instead of returning them. Metal has no feedback stage; the host runs this
-    vertex function with rasterization off and binds buffer b (the game's feedback buffer at its current offset, less first x stride) at TF_SLOT + b, so vertex v writes at v x stride.
-    The values are stored as their bits (packed normals are integers in float outputs); a buffer's outputs follow one another within the vertex's stride, in tf_layout's order"""
+def vs_tf(src, tf, geometry=False):
+    """Capture interleaved output bits, using offsets prepared on the GPU.
+    Geometry readers gather the ordered, decomposed primitive indices before
+    reading their outputs. Their capture index remains linear, so duplicated
+    strip vertices occupy distinct feedback records."""
    m = re.search(r'\bvertex\s+main0_out\s+main0\s*\((.*?)\)\s*\{', src, re.S)
    if not m: sys.exit('[RED] vertex shader without a main0 main function, cannot add the transform feedback entry point')
    params, body = m.group(1), src[m.end():]   # the main function is the last function in the file
    vm = re.search(r'uint\s+(\w+)\s*\[\[vertex_id\]\]', params); vid = vm.group(1) if vm else 'dm_vid'
+    im = re.search(r'uint\s+(\w+)\s*\[\[instance_id\]\]', params); iid = im.group(1) if im else 'dm_iid'
    st = re.search(r'struct main0_out\s*\{(.*?)\};', src, re.S)
    names = {int(l): nm for nm, l in re.findall(r'float4 (\w+) \[\[user\(locn(\d+)\)\]\]', st.group(1) if st else '')}
-    stores = ''
+    stores = []
    for b, e in enumerate(tf):
        if not e: continue
-        words, at = e[1] // 4, 0
-        for o, n in e[0]:
-            stores += ''.join(f' dm_tf{b}[{vid} * {words} + {at + i}] = {f"as_type<uint>(out.{names[o]}.{c})" if o in names else "0u"};' for i, c in enumerate('xyzw'[:n]))
-            at += n
-    extra = ('' if vm else ', uint dm_vid [[vertex_id]]') + ''.join(f', device uint* dm_tf{b} [[buffer({TF_SLOT + b})]]' for b, e in enumerate(tf) if e)
-    return f'\nvertex void main0_tf({params}{extra})\n{{' + body.replace('return out;', '{' + stores + ' return; }')
+        offset = 0
+        for o, n in e['outputs']:
+            for i, c in enumerate('xyzw'[:n]):
+                value = f'as_type<uint>(out.{names[o]}.{c})' if o in names else ('0x3f800000u' if c == 'w' else '0u')
+                stores.append(f'dm_tf{b}[dm_tf.offset[{b}] / 4 + dm_linear * {e["stride"] // 4} + {offset + i}] = {value};')
+            offset += n
+    extra = ('' if vm else ', uint dm_vid [[vertex_id]]') + ('' if im else ', uint dm_iid [[instance_id]]')
+    extra += ''.join(f', device uint* dm_tf{b} [[buffer({TF_SLOT + b})]]' for b, e in enumerate(tf) if e)
+    extra += ', constant dm_tf_capture& dm_tf [[buffer(14)]]'
+    pro = f'uint dm_linear = ({iid} - dm_tf.base_instance) * dm_tf.count + {vid} - dm_tf.first;\n'
+    if geometry: pro += f'{vid} = dm_tf.order[dm_linear];\n'
+    decl = '\nstruct dm_tf_capture { uint4 offset; uint4 args; device const uint* order; uint first, count, base_instance, pad; ulong reserved; };\n'
+    return decl + f'\nvertex void main0_tf({params}{extra})\n{{\n' + pro + body.replace('return out;', '{ if (dm_linear < dm_tf.args.x) { ' + ' '.join(stores) + ' } return; }')
+
+
+def vs_discard(src):
+    """Metal requires a void vertex entry when rasterization is disabled.
+    Preserve all original side effects and early returns; omit only outputs."""
+    m = re.search(r'\bvertex\s+main0_out\s+main0\s*\((.*?)\)\s*\{', src, re.S)
+    if not m: raise ValueError('vertex shader has no main0 entry for rasterizer discard')
+    return '\nvertex void main0_discard(' + m.group(1) + ')\n{' + src[m.end():].replace('return out;', 'return;')

 def raster_facts(stage, src):
    """Static facts needed for the runtime proof that the load can be skipped (full-coverage first draw) (nvn_mtl.m first_draw_covers):
@@ -754,12 +771,21 @@ def spv_single_lane(path):
    ballot the active lanes, shuffle the first lane's value to all, and take a fast path when all agree. With a subgroup of one lane that is exact: the ballot is the lane's own bit, shuffles and
    broadcasts return the lane's own value, all / any are the predicate, all-equal and elect are true, the lane ID is 0, the size 1 and the masks follow. The SPIR-V is rewritten that way
    (the builtin variables, their decorations and the subgroup capabilities removed). → True when rewritten, False when there is nothing to do or an operation it does not cover"""
-    d = open(path, 'rb').read(); words = list(struct.unpack(f'<{len(d) // 4}I', d)); hdr, ins, i = words[:5], [], 5
+    with open(path, 'rb') as f: d = f.read()
+    words = list(struct.unpack(f'<{len(d) // 4}I', d)); hdr, ins, i = words[:5], [], 5
    while i < len(words): n = words[i] >> 16; ins.append(words[i:i + n]); i += n
    op = lambda x: x[0] & 0xffff
    if not any(op(x) in SG_OPS or (op(x) == 71 and len(x) > 3 and x[2] == 11 and x[3] in SG_BUILTINS) for x in ins): return False
    if any(340 <= op(x) <= 366 and op(x) not in SG_OPS for x in ins): return False   # other subgroup operations (bit counts, arithmetic): not covered
    bvar = {x[1]: SG_BUILTINS[x[3]] for x in ins if op(x) == 71 and len(x) > 3 and x[2] == 11 and x[3] in SG_BUILTINS}
+    # A mask component can be loaded through an input pointer (Xenoblade Chronicles 3's prefix ballot), rather than loading the whole vector first. Keep its index and extract
+    # from the constant vector at the load; removing only the builtin variable would leave a dangling pointer. Other access-chain forms remain unsupported.
+    bptr = {}
+    for x in ins:
+        if op(x) in (65, 66) and len(x) > 3 and x[3] in bvar:
+            if len(x) != 5 or not bvar[x[3]].endswith('4'): return False
+            bptr[x[2]] = (bvar[x[3]], x[4])
+    if any(op(x) in (65, 66) and len(x) > 3 and x[3] in bptr for x in ins): return False
    uint = next((x[1] for x in ins if op(x) == 21 and x[2:4] == [32, 0]), None); boolt = next((x[1] for x in ins if op(x) == 20), None)
    uvec4 = next((x[1] for x in ins if op(x) == 23 and x[2] == uint and x[3] == 4), None)
    if uint is None or boolt is None or uvec4 is None: return False
@@ -777,15 +803,19 @@ def spv_single_lane(path):
        if o == 15: x = _ep_drop(x, bvar)
        if o == 54 and not placed: out += consts; placed = True
        if o == 61 and x[3] in bvar: out.append([83 | 4 << 16, x[1], x[2], c[bvar[x[3]]]]); continue   # a load of a builtin → its constant
+        if o in (65, 66) and x[2] in bptr: continue
+        if o == 61 and x[3] in bptr:
+            nm, index = bptr[x[3]]
+            out.append([77 | 5 << 16, x[1], x[2], c[nm], index]); continue   # OpVectorExtractDynamic, including a constant index
        if o in SG_OPS:
            kind = SG_OPS[o]
            if kind == 'ballot': out.append([169 | 6 << 16, x[1], x[2], x[4], c['one4'], c['zero4']]); continue   # OpSelect(pred, own bit, 0)
            out.append([83 | 4 << 16, x[1], x[2], c['true'] if kind == 'true' else x[4]]); continue
        out.append(x)
    if not placed: return False
-    if any(op(x) == 65 and len(x) > 3 and x[3] in bvar for x in out): return False   # reached through an access chain: not covered
    hdr[3] = bound
-    flat = hdr + [v for x in out for v in x]; open(path, 'wb').write(struct.pack(f'<{len(flat)}I', *flat))
+    flat = hdr + [v for x in out for v in x]
+    with open(path, 'wb') as f: f.write(struct.pack(f'<{len(flat)}I', *flat))
    return True
 def _ep_drop(x, ids):   # OpEntryPoint without the given interface IDs: model, function, the name (a nul-terminated string in whole words), then the interface list
    k = 3
@@ -841,6 +871,9 @@ def geometry_finish(key, w, j, c1):
                    'stage': 'vs', 'entry': 'main0', 'entry_z': 'main0_z', 'msl_entry_kind': 'vertex',
                    'geometry_reader': True, 'geometry_output': layout['output'], 'c1': '', 'notes': [], 'drv_reads': {}, 'drv_unfillable': [],
                    'bindings': reflect(key, 'vs', analysis, {'sbufs': []}), 'vertex_inputs': [], 'raster': {}}
+            if j.get('tf'):
+                full += vs_tf(analysis, j['tf'], geometry=True)
+                meta.update(entry_tf='main0_tf', tf_strides=[e['stride'] if e else 0 for e in j['tf']])
            save('_bgv', full, meta)
    return 'ok', '', base

@@ -849,7 +882,7 @@ def finish(key, w, j, c1):
    name, k = loc[key]
    msl = os.path.join(w, key + '.metal'); lib = os.path.join(out, key + '_b.metallib')
    if j['stage'] == 'gs':   # keep the layer-only fold; other geometry runs as compute
-        s = gs_slice(os.path.join(w, key + '.spv'))
+        s = gs_slice(os.path.join(w, key + '.spv')) if not j.get('tf') else None
        if s:
            refl = {'key': key + '_b', 'route': 'B', 'chain': CHAIN, 'dm_ctl': key, 'source': f'{name}#{k}', 'pc_paired': pc_of[key], 'stage': 'gs', 'gs_slice': s,
                    'notes': [], 'drv_reads': '', 'drv_unfillable': [], 'bindings': []}   # no library: the host draws with the vertex shader's layer entry points
@@ -880,18 +913,20 @@ def finish(key, w, j, c1):
    if j['stage'] == 'vs' and PAIRS.get(key):   # outputs the pixel shader reads but the vertex shader does not write get (0,0,0,1) (see vs_pad)
        src, pad = vs_pad(key, src)
        if pad: open(msl, 'w').write(src)
-    lay = None
+    lay = None; discard = False
+    res = json.load(open(os.path.join(w, key + '.res.json')))
    if j['stage'] == 'vs':   # the two depth-convention entry points (see vs_depth); the analysis below uses the source with only main0
        capture = dmsh_geometry.capture(src)
        src, full = vs_depth(src)
        full += capture
+        discard = any(b.get('rw') for b in reflect(key, 'vs', src, res))
+        if discard: full += vs_discard(src)
        if j.get('tf'): full += vs_tf(src, j['tf'])   # transform feedback capture (see vs_tf)
        plain = full
        full, lay = vs_layer_entry(full)   # layer entry points for a slice-writing geometry shader (see vs_layer_entry)
        open(msl, 'w').write(full)
    fu, fmask = fs_uint(src) if j['stage'] == 'ps' else (None, 0)   # the integer-output entry points (see fs_uint); the analysis below uses the source with only main0
    if fu: plain = src; open(msl, 'w').write(fu)
-    res = json.load(open(os.path.join(w, key + '.res.json')))
    try: MTL.compile('mac', msl, lib)   # MSL identical to one compiled before takes that library (dmsh_mtlcache.py)
    except subprocess.CalledProcessError as e:
        if not lay and not fu: return 'metal failed', str(e)[:160], None
@@ -906,7 +941,8 @@ def finish(key, w, j, c1):
    dr = drv_reads(src)
    refl = {'key': key + '_b', 'route': 'B', 'chain': CHAIN, 'dm_ctl': key, 'source': f'{name}#{k}', 'pc_paired': pc_of[key], 'stage': j['stage'],
            'entry': entry.group(2) if entry else 'main0', 'msl_entry_kind': entry.group(1) if entry else None, **({'entry_z': 'main0_z', 'entry_geometry_capture': 'main0_gs'} if j['stage'] == 'vs' else {}),
-            **({'entry_tf': 'main0_tf', 'tf_strides': [e[1] if e else 0 for e in j['tf']]} if j.get('tf') else {}),
+            **({'entry_discard': 'main0_discard'} if discard else {}),
+            **({'entry_tf': 'main0_tf', 'tf_strides': [e['stride'] if e else 0 for e in j['tf']]} if j.get('tf') else {}),
            **({'entry_layer': 'main0_l', 'entry_layer_z': 'main0_lz', 'layer_from': lay} if lay else {}), **({'entry_u': 'main0_u'} if fu else {}), **({'entry_m': 'main0_m', 'uint_mask_m': fmask} if fu and fmask else {}),
            'threadgroup_size': j.get('local'), 'c1': c1.hex(), 'notes': res['notes'] + (['not in the answer table, queries use defaults'] if j.get('no_facts') else []), 'drv_reads': dr, 'drv_unfillable': unfillable(dr),
            'bindings': reflect(key, j['stage'], src, res),
@@ -49,7 +49,7 @@
 #   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)
-import collections, concurrent.futures as cf, hashlib, json, os, re, struct, subprocess, sys
+import collections, concurrent.futures as cf, hashlib, json, os, re, struct, subprocess, sys, zlib
 sys.path.insert(0, os.path.join(os.path.dirname(os.path.abspath(__file__)), '../../games')); import game as _game
 sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))); import uepak
 GAME = _game.load()   # the current game profile (DM_GAME, required): source group rules and the table header note
@@ -220,7 +220,7 @@ def glslc_program(d, h, size, label):
    if len(secs) < 2 or len(set(sts)) < len(sts) or any(not st or check(label, st, ctl, code) for st, (ctl, code) in zip(sts, secs)): return []
    return [(f'{label}#GLSLC@{h:#x}#var', st, ctl, code) for st, (ctl, code) in zip(sts, secs)]

-def glslc_shaders(d, label, need=True):
+def glslc_shaders(d, label, need=True, skip=()):
    """GLSLC compiler outputs (838 in 13 Sentinels' main read-only segment, one stage each): from the magic 0x19866891, +0x44 = the whole output's length; an output of one stage holds exactly one control section with the code area right after it (control +0x20 = control length),
    one of several stages is read by its section headers (glslc_program).
    The magic is only 4 bytes, so only outputs that pass the triple check are kept. Every game's main is searched: Unity's player has 11 of its own (Fire Emblem Engage, Ace Attorney Investigations
@@ -231,6 +231,8 @@ def glslc_shaders(d, label, need=True):
    something else at control +0x20 (odd values such as 0x136b), so a length that is not a multiple of 4 is not a control length"""
    out = []
    for m in re.finditer(b'\x91\x68\x86\x19', d):
+        if any(a <= m.start() < b for a, b in skip): continue
+        if m.start() + 0x48 > len(d): continue
        h = m.start(); size = struct.unpack_from('<I', d, h + 0x44)[0]
        c = d.find(CTL_MAGIC, h, h + size) if 0x760 < size <= len(d) - h else -1
        if c < 0: continue
@@ -1562,6 +1564,175 @@ def ue_inline_shaders(d, label, out):
            except Exception: continue
            if len(s) == raw and s[:4] == b'LSLG': take(s, j); break

+def arh_entries(d):
+    """Monolith ARH1 index → (path, ARD offset, stored size). The XOR trie
+    reconstructs full paths; the metadata table supplies the actual file bounds.
+    Format: https://github.com/roccodev/ard-tools/blob/master/docs/arh1.md
+    Xenoblade 3 2.2.1: 76,445 entries in bf3.arh, 9.05 GB in bf3.ard."""
+    if len(d) < 48 or d[:4] != b'arh1': return []
+    _, count, so, sn, no, nn, fo, fn, key = struct.unpack_from('<9I', d, 4)
+    if nn != count * 8 or any(o > len(d) or n > len(d) - o for o, n in ((so, sn), (no, nn), (fo, fn * 24))): red('ARH1: index table out of bounds')
+    key = (key ^ 0xf3f35353).to_bytes(4, 'little')
+    # The format encrypts whole u32 words; a final partial word stays plain.
+    dec = lambda o, n: bytes(v ^ key[i & 3] for i, v in enumerate(d[o:o + (n & ~3)])) + d[o + (n & ~3):o + n]
+    strings = dec(so, sn); nodes = list(struct.iter_unpack('<ii', dec(no, nn)))
+    names = {}
+    for i, (nxt, prev) in enumerate(nodes):
+        if nxt >= 0 or prev < 0: continue
+        off = -nxt
+        if off >= len(strings) or (end := strings.find(b'\0', off)) < 0 or end + 5 > len(strings): red('ARH1: invalid path suffix')
+        fid = struct.unpack_from('<I', strings, end + 1)[0]; prefix = []; cur = i; seen = set()
+        while cur:
+            if cur in seen or not 0 <= cur < len(nodes): red('ARH1: invalid path trie')
+            seen.add(cur); parent = nodes[cur][1]
+            if not 0 <= parent < len(nodes): red('ARH1: invalid path parent')
+            ch = cur ^ nodes[parent][0]
+            if not 0 <= ch < 128: red('ARH1: invalid path character')
+            prefix.append(ch); cur = parent
+        names[fid] = (bytes(reversed(prefix)) + strings[off:end]).decode('ascii')
+    out = []
+    for i in range(fn):
+        off, size, raw, flags, fid = struct.unpack_from('<Q4I', d, fo + i * 24)
+        # HasXbc1Header (bit 1): stored size excludes the 48-byte wrapper.
+        if size: out.append((names.get(fid, f'#{fid}'), off, size + (48 if flags & 2 else 0)))
+    return out
+
+
+def arh_jobs(files, parts=64):
+    jobs, used = [], set()
+    for path, rel in files:
+        if not path.endswith('.arh'): continue
+        with open(path, 'rb') as f: d = f.read()
+        if d[:4] != b'arh1': continue
+        ard = path[:-4] + '.ard'
+        if not os.path.isfile(ard): red(f'{rel}: ARH1 index has no ARD data file')
+        entries = arh_entries(d); length = os.path.getsize(ard)
+        if any(o > length or n > length - o for _, o, n in entries): red(f'{rel}: ARD entry out of bounds')
+        step = max(1, (len(entries) + parts - 1) // parts)
+        jobs += [(ard, rel[:-4] + '.ard', (b'ARH1', entries[i:i + step])) for i in range(0, len(entries), step)]
+        used.update((path, ard))
+    return jobs, used
+
+
+def xbc_data(d, off):
+    """An XBC1 stream: 48-byte header, stored / zlib / zstd payload. Return
+    (decoded bytes, end), or None for bytes that merely resemble a header."""
+    if off + 48 > len(d) or d[off:off + 4] != b'xbc1': return None
+    kind, raw, size = struct.unpack_from('<3I', d, off + 4)
+    if kind not in (0, 1, 3) or not size or off + 48 + size > len(d): return None
+    data = d[off + 48:off + 48 + size]
+    if kind == 1: data = zlib.decompress(data)
+    elif kind == 3: data = zstd_mod().decompress(data)
+    if len(data) != raw: red(f'XBC1 @{off:#x}: expected {raw} bytes, decoded {len(data)}')
+    return data, off + 48 + size
+
+
+def monolith_cmp_data(d, off):
+    """Firehawk effect archives: cmp\\0zlib, stored/raw sizes, zlib payload.
+    These independently compressed shader resources sit inside XBC1 entries."""
+    if off + 16 > len(d) or d[off:off + 8] != b'cmp\0zlib': return None
+    size, raw = struct.unpack_from('<2I', d, off + 8)
+    if not size or size > len(d) - off - 16: return None
+    z = zlib.decompressobj(); data = z.decompress(d[off + 16:off + 16 + size]) + z.flush()
+    if not z.eof or z.unused_data or len(data) != raw: red(f'cmp zlib @{off:#x}: payload does not match stored/raw sizes')
+    return data, off + 16 + size
+
+
+def hcps_shaders(d, base, label):
+    """HCPS shader programs (Monolith's wishp, also embedded in models/map
+    streams). SLCT gives the control/code ranges and variant descriptors; each
+    NVSD's final stage record gives their sizes. Pair stages within each variant.
+    Layout cross-checked against xc3_lib/src/spch.rs and the game's own bytes."""
+    def span(off, n):
+        if off < base or off > len(d) or n > len(d) - off: red(f'{label} HCPS @{base:#x}: section out of bounds')
+        return d[off:off + n]
+    span(base, 72)
+    table, count = struct.unpack_from('<2I', d, base + 8)
+    slct_off, slct_size, code_off, code_size, ctl_off, ctl_size = struct.unpack_from('<6I', d, base + 24)
+    span(base + table, count * 8)
+    slcts, codes, ctls = (span(base + o, n) for o, n in ((slct_off, slct_size), (code_off, code_size), (ctl_off, ctl_size)))
+    if not code_size and not ctl_size: return []   # source-only GLSL programs (DLC model streams), no NVN binaries to extract
+    out = []
+    for i in range(count):
+        s = struct.unpack_from('<I', d, base + table + 8 * i)[0]
+        if s + 72 > len(slcts) or slcts[s:s + 4] != b'SLCT': red(f'{label} HCPS: invalid SLCT {i}')
+        num, ptr = struct.unpack_from('<2I', slcts, s + 16)
+        c, cn, k, kn = struct.unpack_from('<4I', slcts, s + 40)
+        ce, ke = c + cn, k + kn
+        if s + ptr + num * 8 > len(slcts) or ce > len(ctls) or ke > len(codes): red(f'{label} HCPS SLCT {i}: variant range out of bounds')
+        for j in range(num):
+            po, pn = struct.unpack_from('<2I', slcts, s + ptr + j * 8); p = s + po
+            if pn < 32 or p + pn > len(slcts): red(f'{label} HCPS: invalid NVSD')
+            vo, vn = struct.unpack_from('<2I', slcts, p + 24)
+            if not vn or vo + vn * 20 > pn: red(f'{label} HCPS: invalid NVSD stage records')
+            _, vs, ps, vc, pc = struct.unpack_from('<5I', slcts, p + vo + (vn - 1) * 20)
+            pend, kend = c + vc + pc, k + vs + ps
+            if pend > ce or kend > ke: red(f'{label} HCPS: program exceeds SLCT range')
+            # A vertex bundle can include a geometry stage (shd_lgt.wishp:
+            # 768-byte vertex + 256-byte geometry, two 0x880 controls). NVSD's
+            # vertex size covers both; each NVN control gives its own code size.
+            while c < pend:
+                if c + 0x708 > pend: red(f'{label} HCPS: truncated control')
+                nc = struct.unpack_from('<I', ctls, c + 0x20)[0]
+                nk = struct.unpack_from('<I', ctls, c + 0x704)[0]
+                if nc < 0x760 or c + nc > pend or not nk or k + nk > kend: red(f'{label} HCPS: shader exceeds program range')
+                ctl, code = ctls[c:c + nc], codes[k:k + nk]; c += (nc + 7) & ~7; k += nk
+                st = glslc_stage(ctl, code)
+                why = check(label, st, ctl, code) if st else 'unknown stage'
+                if why: red(f'{label} HCPS SLCT {i} variant {j}: {why}')
+                out.append((f'{label}#HCPS@{base:#x}#var{i}.{j}', st, ctl, code))
+            if k != kend: red(f'{label} HCPS: unused program code bytes')
+        if c != ce or k != ke: red(f'{label} HCPS SLCT {i}: unused control/code bytes')
+    return out
+
+
+def slct_shaders(d, base, label):
+    """A standalone SLCT has its control/code offsets relative to itself.
+    Monolith embeds this form in main, without an HCPS header. Present those
+    same relative sections to the HCPS reader; no code or control bytes change.
+    Its text-bearing control sections are individually aligned to eight bytes."""
+    if base + 72 > len(d): return []
+    c, cn, k, kn = struct.unpack_from('<4I', d, base + 40)
+    size = max(c + cn, k + kn)
+    if min(c, k) < 72 or size > len(d) - base: return []
+    header = bytearray(80); header[:4] = b'HCPS'
+    struct.pack_into('<2I', header, 8, 72, 1)
+    struct.pack_into('<6I', header, 24, 80, size, 80, size, 80, size)
+    return hcps_shaders(bytes(header) + d[base:base + size], 0, f'{label}#SLCT@{base:#x}')
+
+
+def monolith_shaders(d, label, out):
+    """Scan decoded archive entries and their XBC1 streams without relying on
+    extensions: model streams and map data embed HCPS alongside other assets."""
+    pos = 0; compressed = []
+    # Separate literal scans use Python's fast substring search. Alternation's
+    # character-set search is costly across the archive's decoded texture data.
+    markers = sorted((m.start(), magic == b'xbc1', len(magic)) for magic in (b'xbc1', b'cmp\0zlib') for m in re.finditer(magic, d))
+    for off, xbc, size in markers:
+        if off < pos: continue
+        result = xbc_data(d, off) if xbc else monolith_cmp_data(d, off)
+        pos = off + size if result is None else result[1]
+        if result is not None:
+            compressed.append((off, pos))
+            monolith_shaders(result[0], f'{label}#{"XBC1" if xbc else "cmp-zlib"}@{off:#x}', out)
+    ci = 0
+    for m in re.finditer(b'HCPS', d):
+        base = m.start()
+        while ci < len(compressed) and compressed[ci][1] <= base: ci += 1
+        if ci < len(compressed) and compressed[ci][0] <= base: continue   # a zstd literal is not a decoded container
+        if base + 8 <= len(d) and struct.unpack_from('<I', d, base + 4)[0] in (1000, 1010): out.extend(hcps_shaders(d, base, label))
+    # Effect asset packs keep complete GLSLC outputs alongside HCPS programs.
+    # As above, compressed literals are visited only after decoding their stream.
+    out.extend(glslc_shaders(d, label, need=False, skip=compressed))
+
+
+def arh_shaders(path, rel, entries, out):
+    with open(path, 'rb') as f:
+        for name, off, size in entries:
+            f.seek(off); data = f.read(size)
+            monolith_shaders(data, f'/{rel}{name}', out)
+
+
 def yaz0(s):
    """Yaz0 decompression (Super Mario 3D World + Bowser's Fury: almost all of RomFS is .szs = SARC wrapped in Yaz0). Pure Python runs at about 20–45 MB/s per core, under a minute for the whole 4.7 GB RomFS in parallel, and needs no extra dependency.
    Group header bits: 1 = one literal byte; 0 = back-reference, 2 bytes {length high nibble, 12-bit distance} (distance + 1); length 0 reads one more byte + 0x12, else + 2. Overlapping copies are expanded by their period"""
@@ -1750,6 +1921,8 @@ def unity_text_glslc(x, base):
    return out

 def scan_blob(d, label, out):
+    if d[:4] == b'HCPS': out.extend(hcps_shaders(d, 0, label)); return
+    for m in re.finditer(b'SLCT\x03\x10\0\0', d): out.extend(slct_shaders(d, m.start(), label))
    if d[:4] == b'Yaz0': d = yaz0(d)
    if d[:4] == b'SARC':
        for nm, m in sarc_members(d): scan_blob(m, f'{label}/{nm}', out)
@@ -1841,6 +2014,9 @@ def scan_file(args):
    if ext in ('.data', '.patch', '.debug', '.nxpc') and os.path.exists(stem + '.dict'): return rel, {}, set(), {}   # read through its Next Level Games index (nlg_shaders)
    with open(path, 'rb') as f: head = f.read(16)
    found = []; trees = {}
+    if len(args) > 2 and args[2][0] == b'ARH1':
+        arh_shaders(path, rel, args[2][1], found)
+        return rel, uniq_keys(found), var_pairs(found), trees
    if head[:4] == NLG_DICT:
        nlg_shaders(path, '/' + rel, found)
        return rel, uniq_keys(found), var_pairs(found), trees
@@ -1984,8 +2160,9 @@ def main():
    pz_jobs, pkzs = pkz_jobs(files)   # Bayonetta 3's PKZ archives, Oodle Kraken decoded by the game's own decoder in the Mac host
    ck_jobs, cpks = cpk_jobs(files)   # Bayonetta 2's CRI CPK archives (CRILAYLA entries)
    zp_jobs, zips = zip_jobs(files)   # Portal's game.zip (Deflate members)
-    big += pak_jobs + tr_jobs + ue_jobs + bd_jobs + arc_jobs + pz_jobs + ck_jobs + zp_jobs
-    files = [(p, r) for p, r in files if not linkdata_jobs(p, r) and p not in paks and p not in trp and p not in uep and p not in bdts and p not in arcs and p not in pkzs and p not in cpks and p not in zips]
+    ah_jobs, arhs = arh_jobs(files)   # Monolith's ARH1/ARD archives and embedded XBC1/HCPS programs
+    big += pak_jobs + tr_jobs + ue_jobs + bd_jobs + arc_jobs + pz_jobs + ck_jobs + zp_jobs + ah_jobs
+    files = [(p, r) for p, r in files if not linkdata_jobs(p, r) and p not in paks and p not in trp and p not in uep and p not in bdts and p not in arcs and p not in pkzs and p not in cpks and p not in zips and p not in arhs]
    with cf.ProcessPoolExecutor(J, initializer=_pool_init, initargs=(common, {}, zdicts)) as ex:
        results += list(ex.map(scan_file, big)) + list(ex.map(scan_file, files, chunksize=16))
    trees, defer = {}, []
@@ -1,12 +1,32 @@
 // Common header of the native Switch host (multi-module).
 #pragma once
 #include <pthread.h>
+#include <errno.h>
 #include <stdint.h>
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
 #include "game_profile.h"   // constants of the current game profile (generated from dm/games/<id>/game.tsv, chosen by the Makefile's GAME)

+// Explicit host work does not occupy a guest core (Metal submission, or a blocked native codec lock). A contended host mutex must count as blocked: a signal's same-core handoff
+// can otherwise wait forever for a peer that needs the signaller's native lock. Uncontended locks retain guest ordering; the counter nests inside existing host work.
+extern __thread _Atomic int dm_host_work;
+static inline int dm_host_mutex_lock(pthread_mutex_t *mutex) {
+    int r = pthread_mutex_trylock(mutex);
+    if (r != EBUSY) return r;
+    dm_host_work++;
+    r = pthread_mutex_lock(mutex);
+    dm_host_work--;
+    return r;
+}
+// Joining a native helper releases the guest core while that helper finishes; preserve any enclosing host-work scope.
+static inline int dm_host_thread_join(pthread_t thread, void **result) {
+    dm_host_work++;
+    int r = pthread_join(thread, result);
+    dm_host_work--;
+    return r;
+}
+
 #define MAX_MOD 128   // the profile's modules (up to 8) and the NROs a game loads at run time (Super Smash Bros. Ultimate: 98)
 #define LOG(...) do { fprintf(stderr, "[dm] " __VA_ARGS__); fputc('\n', stderr); } while (0)
 #define die(...) do { LOG("FATAL " __VA_ARGS__); abort(); } while (0)
@@ -18,8 +18,6 @@ void *objmap_get(const void *guest);
 void objmap_put(const void *guest, void *host);
 void *objmap_del(const void *guest);
 uint64_t now_ns(void);
-// This thread is doing the host's own heavy work (NVN submit / present: encoding Metal commands; on hardware that is the GPU's business and does not occupy a guest core). The same-core preemption approximation (hle_thread.c) does not count it as "running"
-extern __thread _Atomic int dm_host_work;
 // A blocking native HLE wait releases the guest core. With the condition mutex
 // held, enter with 1 before waiting and balance with 0 after reacquiring it.
 void dm_thread_native_wait(int entering);
@@ -235,7 +235,7 @@ static OSStatus au_create(aout *a) {
 static void *au_create_th(void *a) { return (void *)(intptr_t)au_create(a); }
 static void au_open(aout *a) {
    pthread_t t; void *r = NULL;
-    if (pthread_create(&t, NULL, au_create_th, a) || pthread_join(t, &r)) die("AudioUnit setup: cannot start its thread");
+    if (pthread_create(&t, NULL, au_create_th, a) || dm_host_thread_join(t, &r)) die("AudioUnit setup: cannot start its thread");
    OSStatus st = (OSStatus)(intptr_t)r; if (st) die("AudioUnit setup failed (status %d)", (int)st);
 }
 // Silent sink: consumes in real time at 48 kHz every 5 ms. When more than 100 ms behind (iOS suspended the whole process in the background, SIGSTOP on the Mac), restart the count from now and do not make up the missed time --
@@ -41,6 +41,8 @@ HLE_FORCE("_ZN5movie9Extractor13SetDataSourceEPKc", "movie demuxer: the host doe
 // The size comes from Configure's MediaData, or else from the VP9 header of the first key frame. Unity allocates its output buffers itself as width × height × 1.5 (it does not call GetOutputBufferSize), so the stride is the width and the plane height the height, unaligned.
 // Decoding runs on each decoder's own host thread, and the guest thread that sends a packet or returns an output index waits for it: guest threads run on guest stacks given by the SDK, and libvpx's stack frames are large. Events are still signalled by guest threads (SignalEvent is guest code).
 // Buffer / MediaData / event operations call the SDK's own functions (guest_sym), so the host needs no knowledge of their layouts.
+// Native mutex waits use dm_host_mutex_lock: a callback can wake a higher-priority guest thread that then waits for the decoder lock held by its signaller. That thread is blocked,
+// not runnable guest work; counting it as runnable deadlocked Xenoblade Chronicles 3 at the first released output. Keep the locks across signalling for decoder/event lifetime safety.

 #define MD_NIN 4                 // input buffers
 #define MD_INCAP (4u << 20)      // 4 MB per input buffer (a 1080p VP9 packet is far smaller)
@@ -389,7 +391,7 @@ static void *md_worker(void *arg) {
    mdec_t *d = arg;
    pthread_setname_np("dm movie decoder");
    dm_host_helper_qos(QOS_CLASS_USER_INITIATED);   // under DM_QOS_MAP: decodes ahead of playback, not on the frame's critical path
-    pthread_mutex_lock(&d->wmu);
+    dm_host_mutex_lock(&d->wmu);
    for (;;) {
        while (d->job == MDJ_NONE) pthread_cond_wait(&d->wcv, &d->wmu);
        int job = d->job;
@@ -401,7 +403,7 @@ static void *md_worker(void *arg) {
            uint64_t dt = now_ns() - t0; d->dec_n++; d->dec_ns += dt; if (dt > d->dec_max) d->dec_max = dt;
        }
        else md_close(d);
-        pthread_mutex_lock(&d->wmu);
+        dm_host_mutex_lock(&d->wmu);
        d->job = MDJ_NONE; pthread_cond_broadcast(&d->wcv);
        if (job == MDJ_QUIT) break;
    }
@@ -410,7 +412,7 @@ static void *md_worker(void *arg) {
 }
 // hand one job to the host thread and wait for it (the caller holds dmu)
 static void md_job(mdec_t *d, int job) {
-    pthread_mutex_lock(&d->wmu);
+    dm_host_mutex_lock(&d->wmu);
    d->job = job; pthread_cond_broadcast(&d->wcv);
    while (d->job != MDJ_NONE) pthread_cond_wait(&d->wcv, &d->wmu);
    pthread_mutex_unlock(&d->wmu);
@@ -419,13 +421,13 @@ static void md_job(mdec_t *d, int job) {
 static void md_pump(mdec_t *d) {
    dm_host_lock(&d->dmu);
    for (;;) {
-        pthread_mutex_lock(&d->mu);
+        dm_host_mutex_lock(&d->mu);
        if (d->avc_drain) {   // H.264 end of stream: the pictures still queued leave one per free slot; the last carries the end-of-stream flag
            int s = -1; for (int k = 0; k < d->nout; k++) if (!d->out[k].used) { s = k; break; }
            if (s < 0) { pthread_mutex_unlock(&d->mu); break; }
            d->out[s].used = 1; pthread_mutex_unlock(&d->mu);
            d->jslot = s; d->jflush = 1; md_job(d, MDJ_TAKE); int got = d->jgot, left = d->avc ? mavc_queued(d->avc) : 0;
-            pthread_mutex_lock(&d->mu);
+            dm_host_mutex_lock(&d->mu);
            d->out[s].pic = got; d->out[s].pts = got ? d->jpts_out : 0; d->out[s].flags = left ? 0 : 4; d->outq[d->noutq++] = s;
            if (!left) d->avc_drain = 0;
            pthread_mutex_unlock(&d->mu);
@@ -447,7 +449,7 @@ static void md_pump(mdec_t *d) {
        if ((pk.size > 0 && (MD_VPX(d) || MD_AUDIO(d) || avc)) || (avc && eos)) {
            d->jp = d->in[pk.in] + pk.off; d->jn = pk.size; d->jslot = slot; d->jpts = pk.pts; d->jflush = avc && eos; md_job(d, MDJ_DECODE); got = d->jgot; }
        int more = avc && eos && d->avc && mavc_queued(d->avc) > 0;   // H.264 pictures still queued: they follow this output, the last of them carries the end-of-stream flag
-        pthread_mutex_lock(&d->mu);
+        dm_host_mutex_lock(&d->mu);
        d->infree[d->ninfree++] = pk.in;
        int ready = got || eos;   // an end-of-stream packet without a picture still yields an output to carry the flag back
        if (ready) { d->out[slot].pic = got; d->out[slot].pts = got && avc ? d->jpts_out : pk.pts; d->out[slot].flags = more ? pk.flags & ~4u : pk.flags; d->outq[d->noutq++] = slot; }
@@ -541,7 +543,7 @@ HLE_FORCE("_ZN5movie7Decoder9ConfigureEPNS_9MediaDataE", "movie decoder configur
        if (G.find_buf(md, "avcc", &b) && b && G.data(b) && G.size(b)) mavc_avcc(d->avc, G.data(b), G.size(b));
    }
    int w = 0, h = 0; if (md) { G.find_i32(md, "width", &w); G.find_i32(md, "height", &h); }
-    pthread_mutex_lock(&d->mu); if (w > 0 && h > 0 && !d->fmt_sent) { d->w = w; d->h = h; } pthread_mutex_unlock(&d->mu);
+    dm_host_mutex_lock(&d->mu); if (w > 0 && h > 0 && !d->fmt_sent) { d->w = w; d->h = h; } pthread_mutex_unlock(&d->mu);
    MD_TRACE(d, 0, "Configure: MediaData width %d height %d", w, h);
    return 0;
 }
@@ -555,7 +557,7 @@ HLE_FORCE("_ZN5movie7Decoder20RegisterClientEventsEPNS_13DecoderEventsE", "movie
 }
 HLE_FORCE("_ZN5movie7Decoder5StartEv", "movie decoder start: all inputs empty, notify the client", int, h_movie_Start, void **self) {
    mdec_t *d = md_of(self); if (!d) return 0x2002;
-    pthread_mutex_lock(&d->mu); d->started = 1; pthread_mutex_unlock(&d->mu);
+    dm_host_mutex_lock(&d->mu); d->started = 1; pthread_mutex_unlock(&d->mu);
    MD_TRACE(d, 2, "Start");
    md_defer(d, 0);
    return 0;
@@ -565,7 +567,7 @@ static void md_reset(mdec_t *d, int close) {
    dm_host_lock(&d->dmu);   // wait for the packet being decoded
    if (close) md_job(d, MDJ_CLOSE);
    if (d->avc) mavc_drop(d->avc);   // H.264 pictures waiting for their turn belong to the old position
-    pthread_mutex_lock(&d->mu);
+    dm_host_mutex_lock(&d->mu);
    // Every input goes back to the free list, as with MediaCodec, but one the client holds may still be sent: Big Brain Academy: Brain vs. Brain's event thread acquires inputs while its
    // control thread flushes to loop an activity preview, and sends them after the flush (SendInputBufferForDecode claims them back). Those are withheld for a while (h_movie_AcqIn)
    uint64_t now = now_ns();
@@ -600,7 +602,7 @@ HLE_FORCE("_ZN5movie7Decoder23AcquireInputBufferIndexEPiS1_", "movie decoder: ac
    if (!idx || !left) return 0x100b;
    // No input indices after Stop (as with MediaCodec): Metroid Dread's "input available" handler (main+0x222f14) posts the decoder wrapper object to the player manager's message queue once it gets an index;
    // after Stop returned every input index, a late input event could still get one, the wrapper was destroyed afterwards, and the worker thread (main+0x223bd4) read the freed object and crashed
-    pthread_mutex_lock(&d->mu);
+    dm_host_mutex_lock(&d->mu);
    int r = MD_NO_INPUT; *idx = -1; *left = 0;
    // An input the client held across a flush (md_reset) is not handed out again for MD_STALE_NS unless the client sends it first: handing it out while the client still meant to send
    // it gave the client one buffer twice (a packet dropped as sent twice, VideoToolbox -12909). A client that abandons it gets it back after that time instead of losing a buffer
@@ -631,7 +633,7 @@ HLE_FORCE("_ZN5movie7Decoder24SendInputBufferForDecodeEiiilj", "movie decoder: s
          void **self, int i, int off, int size, int64_t pts, uint32_t flags) {
    mdec_t *d = md_of(self); if (!d) return 0x1000;
    if (i < 0 || i >= MD_NIN || off < 0 || size < 0 || (uint32_t)off + (uint32_t)size > MD_INCAP) return 0x100b;
-    pthread_mutex_lock(&d->mu);
+    dm_host_mutex_lock(&d->mu);
    // An index the client acquired before a Flush and sends after it (Big Brain Academy: Brain vs. Brain loops its activity previews this way): the flush put it back in the free list,
    // so it is claimed from there now. Before, it could be handed out a second time, two packets shared one buffer (VideoToolbox -12909 on the overwritten one) and the fifth pending
    // packet wrote past pend[] into npend and dmu (the process died destroying dmu, 2 runs in 5)
@@ -661,7 +663,7 @@ HLE_FORCE("_ZN5movie7Decoder24AcquireOutputBufferIndexEPiPlPjS1_", "movie decode
          void **self, int *idx, int64_t *pts, uint32_t *flags, int *left) {
    mdec_t *d = md_of(self); if (!d) return 0x1000;
    if (!idx || !pts || !flags || !left) return 0x100b;
-    pthread_mutex_lock(&d->mu);
+    dm_host_mutex_lock(&d->mu);
    int r = MD_NO_OUTPUT; *idx = -1; *left = 0;
    if (d->noutq && d->started) {
        int s = d->outq[0]; memmove(d->outq, d->outq + 1, --d->noutq * sizeof *d->outq);
@@ -768,7 +770,7 @@ HLE_FORCE("_ZN5movie7Decoder15GetOutputBufferEiPNS_6BufferEPNS_9MediaDataE", "mo
 HLE_FORCE("_ZN5movie7Decoder24ReleaseOutputBufferIndexEi", "movie decoder: release an output index (decodes queued packets again)", int, h_movie_RelOut, void **self, int i) {
    mdec_t *d = md_of(self); if (!d) return 0x1000;
    if (i < 0 || i >= d->nout) return 0x100b;
-    pthread_mutex_lock(&d->mu);
+    dm_host_mutex_lock(&d->mu);
    d->out[i].used = 0; d->out[i].pic = 0;
    int queued = d->npend, drain = d->avc_drain;
    pthread_mutex_unlock(&d->mu);
@@ -76,7 +76,13 @@ static int opmode(void) { return dm_opmode(); }

 // Screenshot / video recording switches: the host has no system screenshot service, no-op
 HLE_FORCE("_ZN2nn2oe17FinishStartupLogoEv", "dismiss the system boot screen (through the am application proxy); the host has no boot screen", void, h_oe_FinishLogo, void) {}   // Metroid Dread calls it at main+0xadfed0
-HLE_FORCE("_ZN2nn2oe44SetUserInactivityDetectionTimeExtendedUnsafeEb", "auto screen-off time on inactivity (through am); the host does not turn the screen off", void, h_oe_InactivityExt, int on) { (void)on; }
+// am retains the requested inactivity extension even when auto-sleep is disabled.
+// SDK 12.3.9 maps the normal setter to mode 1, Unsafe to mode 2, and false to 0.
+// Xenoblade Chronicles 3 reads it after the opening movie (sdk+0x3391f0).
+static _Atomic int g_inactivity_mode;
+HLE_FORCE("_ZN2nn2oe44SetUserInactivityDetectionTimeExtendedUnsafeEb", "inactivity extension mode 2 or 0; the host never turns the screen off", void, h_oe_InactivityExt, unsigned on) { g_inactivity_mode = (on & 1) ? 2 : 0; }
+HLE_FORCE("_ZN2nn2oe38SetUserInactivityDetectionTimeExtendedEb", "inactivity extension mode 1 or 0; the host never turns the screen off", void, h_oe_InactivityExtSafe, unsigned on) { g_inactivity_mode = on & 1; }
+HLE_FORCE("_ZN2nn2oe43GetUserInactivityDetectionTimeExtensionModeEv", "inactivity extension mode requested by the application", int, h_oe_InactivityMode, void) { return g_inactivity_mode; }
 HLE_FORCE("_ZN2nn2oe15EnableRecordingEv", "video recording switch; the host has no system recording", void, h_oe_EnableRec, void) {}
 // Applet resource user ID (ARUID, the number am assigns to this process): the host has no am, and the SDK's own getter returns 0. SDK 20's OpenAudioRenderer (the CPU rendering + manual execution path, sdk+0x221730)
 //   aborts when it gets 0 (The Legend of Zelda: Breath of the Wild 1.9.0); SDK 16 did not check. Report a fixed non-zero ID, used only as an identifier
@@ -105,10 +111,7 @@ static _Atomic int g_nosleep;
 HLE_FORCE("_ZN2nn2oe29BeginAutoSleepDisabledSectionEv", "open an auto-sleep disabled section (through am); the host never auto-sleeps, counted for the query", void, h_oe_NoSleepBegin, void) { g_nosleep++; }
 HLE_FORCE("_ZN2nn2oe27EndAutoSleepDisabledSectionEv", "close an auto-sleep disabled section (through am)", void, h_oe_NoSleepEnd, void) { if (g_nosleep > 0) g_nosleep--; }
 HLE_FORCE("_ZN2nn2oe19IsAutoSleepDisabledEv", "whether an auto-sleep disabled section is open", unsigned, h_oe_NoSleepQ, void) { return g_nosleep > 0; }
-HLE_FORCE("_ZN2nn2oe37IsUserInactivityDetectionTimeExtendedEv", "whether auto screen-off on inactivity is extended; the host does not turn the screen off, reports no", unsigned, h_oe_InactivityExtQ, void) { return 0; }
-// The same am value as a mode (0–2; SDK 4: sdk+0x15b72c, the SDK aborts on 3 or more). Donkey Kong Country: Tropical Freeze asks for it whenever a cinematic camera starts
-// (CScreenDimmingUtil::DisableScreenDimming, main+0x647d0) and extends the time when it is 0; the SDK's call through the null applet proxy crashed the first level's load
-HLE_FORCE("_ZN2nn2oe43GetUserInactivityDetectionTimeExtensionModeEv", "auto screen-off extension mode (through am); the host does not turn the screen off, reports 0 (not extended)", int, h_oe_InactivityMode, void) { return 0; }
+HLE_FORCE("_ZN2nn2oe37IsUserInactivityDetectionTimeExtendedEv", "whether the application requested an inactivity extension", unsigned, h_oe_InactivityExtQ, void) { return g_inactivity_mode != 0; }
 // Hyrule Warriors: Age of Calamity (SDK 9.6.0, main+0x8d1238): the copyright frame buffer, an overlay the system draws only into screenshots and recordings. The SDK hands the buffer to am as transfer memory
 // (svcCreateTransferMemory, then a null applet proxy). The host has no system capture, so it keeps nothing
 HLE_FORCE("_ZN2nn2oe30InitializeCopyrightFrameBufferEPvm", "copyright overlay buffer (for screenshots / recording, through am); the host has no system screenshot", void, h_oe_CopyrightInit, void *buf, size_t n) { (void)buf; (void)n; }
@@ -140,6 +143,9 @@ HLE_FORCE("_ZN2nn2oe24SetAlbumImageOrientationENS_5album16ImageOrientationE", "s
 // 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.
+// Xenoblade Chronicles 3 (SDK 12.3.9) enables this at main+0x127fac4;
+// the SDK otherwise dereferences the absent am SelfController at sdk+0x33b0f0.
+// As with Disable, there is no producer: the host has no Horizon album captures.
 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) {
@@ -26,7 +26,6 @@ HLE_FORCE("_ZN2nn2oe24SetExpectedVolumeBalanceEff", "volume balance; the host ig
 HLE_FORCE("_ZN2nn2oe25GetProgramTotalActiveTimeEv", "total active time → time since the host started", int64_t, h_oe_ActiveTime, void) { return (int64_t)(now_ns() - g_t0_ns); }
 HLE_FORCE("_ZN2nn2oe31EnterExitRequestHandlingSectionEv", "exit request critical section; the host exits through SIGTERM / the lifecycle", void, h_oe_EnterExit, void) {}
 HLE_FORCE("_ZN2nn2oe31LeaveExitRequestHandlingSectionEv", "exit request critical section; the host ignores it", void, h_oe_LeaveExit, void) {}
-HLE_FORCE("_ZN2nn2oe38SetUserInactivityDetectionTimeExtendedEb", "extended auto-sleep delay; the host ignores it", void, h_oe_Inactivity, unsigned on) { (void)on; }

 // ---------- time ----------
 // CalendarTime = { int16 year; int8 month, day, hour, minute, second }, returned by value (within 8 bytes, in x0)
@@ -125,7 +125,7 @@ static thr_t *cur(void) {
 // svc_thread calls it after waking such a thread. Wait for guest blocking, not a host dispatch snapshot or a wall-clock deadline: neither establishes that the work has finished.
 void dm_thread_mark_host(void) { thr_t *t = pthread_getspecific(g_keythr); if (t) t->host_owned = 1; }
 // Holding G. wake() clears wkind before the pthread resumes; it must already count as runnable in that window, even if Mach still reports WAITING / STOPPED.
-// Explicit host GPU work does not occupy a guest core. Ordinary host dispatch delays and HLE calls do, until a guest wait / sleep / exit says otherwise.
+// Explicit host GPU work and contended native codec locks do not occupy a guest core. Ordinary host dispatch delays and HLE calls do, until a guest wait / sleep / exit says otherwise.
 static int thr_runnable(const thr_t *t) {
    return !t->wkind && !t->sleeping && !t->native_wait && !t->pause_req && !(t->host_work && *t->host_work);
 }
@@ -77,6 +77,9 @@ static void probe_heapwalk(uint64_t h) {
    }
 }   // dump: on hit print 16 bytes at [x0+dump] (-1 = do not print); fn: dedicated handler (replaces the default print)
 static int g_nprobe;
+static uint64_t g_probe_frame_first, g_probe_frame_last = UINT64_MAX;
+static unsigned g_probe_limit;
+uint64_t dm_guest_frame(void);
 static void probe_str(uint64_t p) {
    char buf[80]; vm_size_t got = 0;
    if (p < 0x1000 || vm_read_overwrite(mach_task_self(), p, sizeof buf - 1, (vm_address_t)buf, &got) || !got) return;
@@ -107,8 +110,11 @@ static void probe_dump_once(void) {
 static int probe_hit(ucontext_t *uc) {   // Registered by probe_install with the SIGTRAP branch of crash.c
    uint64_t pc = uc->uc_mcontext->__ss.__pc;
    for (int i = 0; i < g_nprobe; i++) if (g_probe[i].at == pc) {
-        probe_dump_once();
        if (g_probe[i].fn) { g_probe[i].hits++; g_probe[i].fn(uc); uc->uc_mcontext->__ss.__pc = (uint64_t)g_probe[i].slot; return 1; }
+        uint64_t frame = dm_guest_frame();
+        if (frame < g_probe_frame_first || frame > g_probe_frame_last || (g_probe_limit && g_probe[i].hits >= g_probe_limit)) {
+            uc->uc_mcontext->__ss.__pc = (uint64_t)g_probe[i].slot; return 1;
+        }
        if (g_probe[i].freg >= 0) {   // ?xN=value / ?[xN+offset]=value: pass through silently when it does not match
            uint64_t v = uc->uc_mcontext->__ss.__x[g_probe[i].freg]; vm_size_t g = 0;
            if (g_probe[i].foff >= 0 && vm_read_overwrite(mach_task_self(), v + g_probe[i].foff, 8, (vm_address_t)&v, &g)) v = ~g_probe[i].fval;
@@ -118,8 +124,9 @@ static int probe_hit(ucontext_t *uc) {   // Registered by probe_install with the
            char nm[128]; probe_heapname(uc->uc_mcontext->__ss.__x[g_probe[i].hreg], nm, sizeof nm);
            if (!strstr(nm, g_hw.want)) { uc->uc_mcontext->__ss.__pc = (uint64_t)g_probe[i].slot; return 1; }
        }
+        probe_dump_once();
        flockfile(stderr);   // One block per hit when threads hit at once
-        { extern uint64_t g_t0_ns; uint64_t now_ns(void); uint32_t thr_cur_handle(void); fprintf(stderr, "[dm] PROBE %s #%u [%.4fs] T=%x", g_probe[i].name, ++g_probe[i].hits, (now_ns() - g_t0_ns) / 1e9, thr_cur_handle()); }
+        { extern uint64_t g_t0_ns; uint64_t now_ns(void); uint32_t thr_cur_handle(void); fprintf(stderr, "[dm] PROBE %s #%u [%.4fs] F=%llu T=%x", g_probe[i].name, ++g_probe[i].hits, (now_ns() - g_t0_ns) / 1e9, (unsigned long long)frame, thr_cur_handle()); }
        for (int r = 0; r < 8; r++) { uint64_t v = uc->uc_mcontext->__ss.__x[r]; fprintf(stderr, " x%d=%#llx", r, v); probe_str(v); }
        for (int r = 19; r < 24; r++) fprintf(stderr, " x%d=%#llx", r, uc->uc_mcontext->__ss.__x[r]);
        if (getenv("DM_PROBE_ALL")) for (int r = 9; r < 19; r++) fprintf(stderr, " x%d=%#llx", r, uc->uc_mcontext->__ss.__x[r]);   // also print x9–x18
@@ -357,6 +364,13 @@ static void watch_install(void) {
 void probe_install(void) {
 #if !TARGET_OS_IPHONE
    crash_set_trap_hook(probe_hit);
+    const char *frames = getenv("DM_PROBE_FRAMES");
+    if (frames) {
+        char *end; g_probe_frame_first = strtoull(frames, &end, 10);
+        g_probe_frame_last = *end == '-' ? strtoull(end + 1, NULL, 10) : g_probe_frame_first;
+    }
+    const char *limit = getenv("DM_PROBE_LIMIT");
+    if (limit) g_probe_limit = (unsigned)strtoul(limit, NULL, 10);
    const char *e = getenv("DM_PROBE");
    if (e) {
        char spec[512]; strlcpy(spec, e, sizeof spec);
@@ -3,6 +3,7 @@
 #include "../hle_audio.c"
 #include <time.h>
 #include <sys/mman.h>
+__thread _Atomic int dm_host_work;   // defined by the graphics host in a full build
 uint64_t now_ns(void) { struct timespec t; clock_gettime(CLOCK_MONOTONIC, &t); return (uint64_t)t.tv_sec * 1000000000ull + t.tv_nsec; }
 static struct { const void *k; void *v; } om[8];
 void *objmap_get(const void *k) { for (int i = 0; i < 8; i++) if (om[i].k == k) return om[i].v; return NULL; }
@@ -1,5 +1,6 @@
 // hle_audio.c lock-free queue stress test: three threads producing / playing / reclaiming 100 thousand blocks concurrently, the byte stream and reclaim order checked one by one by sequence number, no event lost.
 #include "../hle_audio.c"
+__thread _Atomic int dm_host_work;   // defined by the graphics host in a full build
 #include <sched.h>
 #include <time.h>

@@ -59,6 +59,13 @@ static int child_burst(int kind) {
 }
 int main(void) {
    unsetenv("DM_OE_SCRIPT"); reset();
+    CHECK(h_oe_InactivityMode() == 0 && !h_oe_InactivityExtQ(), "inactivity extension starts disabled");
+    h_oe_InactivityExtSafe(1);
+    CHECK(h_oe_InactivityMode() == 1 && h_oe_InactivityExtQ(), "normal inactivity extension is mode 1");
+    h_oe_InactivityExt(1);
+    CHECK(h_oe_InactivityMode() == 2 && h_oe_InactivityExtQ(), "unsafe inactivity extension is mode 2");
+    h_oe_InactivityExtSafe(0x100);
+    CHECK(h_oe_InactivityMode() == 0 && !h_oe_InactivityExtQ(), "clearing the extension restores mode 0");
    CHECK(empty(), "empty queue at start: returns false and leaves the output alone");
    dm_focus_set(1); dm_set_opmode(1); CHECK(empty(), "same state sends no notification");
    dm_focus_set(3); CHECK(h_oe_GetFocus() == 3, "background state 3 can be queried before popping");
@@ -228,6 +228,79 @@ static int handoff_scope(int kind) {
    printf("%s: handoff scope, %s: signaller returned, payload %d, worker exited %d\n", ok ? "PASS" : "FAIL", what[kind], payload, done);
    return ok;
 }
+static pthread_mutex_t native_mutex = PTHREAD_MUTEX_INITIALIZER;
+static void *native_wait_worker(void *arg) {
+    thr_t *t = arg; pthread_setspecific(g_keythr, t); t->host_work = &dm_host_work;
+    scope_wait(1);
+    atomic_store(&scope_payload, 1);
+    assert(!dm_host_mutex_lock(&native_mutex));
+    assert(dm_host_work == 0);
+    atomic_store(&scope_payload, 2);
+    pthread_mutex_unlock(&native_mutex);
+    thread_exit(t); return NULL;
+}
+static int native_wait_handoff(void) {
+    uint64_t r[32] = {0}; svc_thread(0x45, r); assert(r[0] == 0);
+    scope_ev[1][0] = r[1]; scope_ev[1][1] = r[2];
+    atomic_store(&scope_payload, 0);
+    thr_t *w = scope_new(&scope_peer[1], 30);
+    pthread_mutex_lock(&native_mutex);
+    pthread_t pw; assert(!pthread_create(&pw, NULL, native_wait_worker, w)); scope_until_waiting(w);
+    scope_signal(1);   // must return when the higher-priority peer blocks on this caller's native mutex
+    int blocked = atomic_load(&scope_payload) == 1;
+    pthread_mutex_unlock(&native_mutex);
+    pthread_join(pw, NULL); int finished = atomic_load(&scope_payload) == 2;
+    scope_free(w);
+    for (int i = 0; i < 2; i++) { r[0] = scope_ev[1][i]; svc_thread(0x16, r); }
+    printf("%s: native mutex wait releases the signaller, then the worker finishes after unlock\n", blocked && finished ? "PASS" : "FAIL");
+    return blocked && finished;
+}
+static pthread_t native_join_target;
+static _Atomic int native_join_release;
+static void *native_join_helper(void *arg) {
+    (void)arg;
+    while (!atomic_load(&native_join_release)) usleep(100);
+    return &native_join_release;
+}
+static void *native_join_worker(void *arg) {
+    thr_t *t = arg; pthread_setspecific(g_keythr, t); t->host_work = &dm_host_work;
+    scope_wait(1);
+    atomic_store(&scope_payload, 1);
+    void *result = NULL;
+    assert(!dm_host_thread_join(native_join_target, &result));
+    assert(result == &native_join_release && dm_host_work == 0);
+    atomic_store(&scope_payload, 2);
+    thread_exit(t); return NULL;
+}
+static void *native_join_nested_helper(void *arg) {
+    _Atomic int *depth = arg;
+    while (atomic_load(depth) != 2) usleep(100);
+    return arg;
+}
+static int native_join_handoff(void) {
+    uint64_t r[32] = {0}; svc_thread(0x45, r); assert(r[0] == 0);
+    scope_ev[1][0] = r[1]; scope_ev[1][1] = r[2];
+    atomic_store(&scope_payload, 0); atomic_store(&native_join_release, 0);
+    assert(!pthread_create(&native_join_target, NULL, native_join_helper, NULL));
+    thr_t *w = scope_new(&scope_peer[1], 30);
+    pthread_t pw; assert(!pthread_create(&pw, NULL, native_join_worker, w)); scope_until_waiting(w);
+    scope_signal(1);   // must return while the higher-priority peer joins a native helper gated by this caller
+    int blocked = atomic_load(&scope_payload) == 1 && atomic_load(w->host_work) == 1;
+    atomic_store(&native_join_release, 1);
+    pthread_join(pw, NULL); int finished = atomic_load(&scope_payload) == 2;
+    scope_free(w);
+    for (int i = 0; i < 2; i++) { r[0] = scope_ev[1][i]; svc_thread(0x16, r); }
+    // The native helper observes the nested depth while join blocks; returning must retain the caller's outer scope.
+    pthread_t nested; void *result = NULL;
+    assert(dm_host_work == 0);
+    dm_host_work++;
+    assert(!pthread_create(&nested, NULL, native_join_nested_helper, &dm_host_work));
+    assert(!dm_host_thread_join(nested, &result));
+    int nesting = result == &dm_host_work && dm_host_work == 1;
+    dm_host_work--;
+    printf("%s: native join releases the signaller, returns the helper result and preserves nested host work\n", blocked && finished && nesting ? "PASS" : "FAIL");
+    return blocked && finished && nesting;
+}
 int main(int argc, char **argv) {
    (void)argv;
    alarm(10);
@@ -236,6 +309,8 @@ int main(int argc, char **argv) {
    main_handle = svc_thread_main_init();
    condition_order();
    int ok = 1;
+    ok &= native_wait_handoff();
+    ok &= native_join_handoff();
    if (argc == 1) for (int kind = 0; kind <= 3; kind++) ok &= handoff_scope(kind);   // with every runnable peer checked (core_preempt), the busy peer holds the signaller by design
    for (int kind = 0; kind <= 4; kind++) if (kind != 3) {
        ok &= delayed_wake(1000, kind, 0);
@@ -1,6 +1,14 @@
 # The rewriter's patches of Nintendo SDK module builds (sdk), which several games can ship: dm/tools/patch_nso.py applies those of a module's NSO build ID, whatever the game.
 # A patch of a game's own module (main) goes in its profile's patches.py (dm/games/<id>/patches.py). CUTS as described in patch_nso.py ("Registered patches")
 CUTS = {
+  # Xenoblade Chronicles 3 2.2.1, SDK 12.3.9: the same AddressRegionManager
+  # constructor and size-class alignment as Engage's SDK 12.3.7. Its cut 4 GB
+  # malloc reservation must fit the 4.375 GB alias region on iOS.
+  'e8cdceb1e77a8465fe20f21928e8d296': [
+    (0x346430, '08fca752', '087ca052', 'AddressRegionManager node area limit: 1 GB to 62 MB'),
+    (0x3464cc, '09fca752', '097ca052', 'AddressRegionManager occupied node area: 1 GB to 62 MB'),
+    (0x34662c, '0a00a852', '0a40a052', 'AllocateAddressRegion: align regions of 1 GB and more to 32 MB'),
+  ],
  # Pokemon Brilliant Diamond 1.3.0 (SDK 11.4.3): nninitStartup (main+0x2c1d5c) gives pool - used to StandardAllocator with virtual address memory.
  # Its first CentralHeap::Init aborts at sdk+0x2f98e4: the 1 GB node area leaves no room for the 1 GB-aligned 3.14 GB reservation in a 4.375 GB alias region.
  # The constructor (sdk+0x34e0b8) is the same mechanism as Engage's below: the limit and the occupied length must both shrink to 62 MB.
@@ -70,6 +70,7 @@ What earlier titles needed beyond sections 1–5 is filed by mechanism: textures
 | `DM_NVN_HILITE=fskey,…` | Replace those fragment shaders with solid colours: tells "geometry is wrong" from "shading is wrong" |
 | `DM_NVN_DEPTHOVR=fskey,…` | Those fragment shaders' draws compare depth with `ALWAYS`: with `HILITE`, tells "the depth test rejects every fragment" from "nothing rasterizes". With `PASSDUMP`, the `D` line shows the depth range (`dr=`) and logic op (`lo=`), and `c` lines each recorded clear (slot, serial, mask, scissor; depth / stencil values) |
 | `DM_NVN_SHDUMP=dir` | Control + code of programs missing from the shader map |
+| `DM_NVN_VSDUMP=<file>` / `DM_NVN_VSDUMP_KEY=<vertex key>` | One native post-vertex capture: a u32 header (first vertex, count, stride in words, base instance), then the 144-word geometry records. Refuses shaders with writes. With key `*`, saves position-only files by draw/key; `DM_NVN_VSDUMP_FROM` selects the first frame, `DM_NVN_VSDUMP_MAX` limits captures (default 256). These diagnostics drain GPU work |
 | `DM_NVN_ALLOCPROF=1` | Metal allocation growth per replayed op type and per function, command allocator sizes, new textures |
 | `DM_NVN_CBDUMP_KEY=vskey` (with `CBDUMP` and `PASSDUMP_ALL`) | Only that vertex shader's draws: uniform buffers plus vertex streams and index buffer (`F<frame>_p<pass>_vb<stream>_<addr>.bin`, `_ib<type>_<addr>.bin`) |
 | `DM_NVN_TEXINFO=WxH` | Storage and description (pool, offset, flags, packaged or not, storage size and how many of its bytes in guest memory are non-zero, i.e. written by the CPU) of every texture of that size on first use, each pool registration of one (`TEXINFO register`: builder channel mapping, view and view mapping), and each one as it is initialized (`TEXINFO init`, render targets included, the first 200: serial, pool and offset, target, format, flags, storage size) |
@@ -52,6 +52,9 @@ Per-probe suffixes, in this order:
 The `!xN` heap walk takes the layout of the game's sead version from `DM_PROBE_HEAP=<name>,<start>,<size>,<free list>,<used list>` (hex offsets in the heap object; default Mario's `0,28,30,b0,c8`, Breath of the Wild's `28,38,40,e0,f8`; find them in the heap's constructor and `tryAlloc`). `DM_PROBE_HEAPNAME=<text>` reports only hits whose heap name contains the text (resource heaps are named after their file), `DM_PROBE_HEAPWORDS=1` adds the first two words of each used block (a vtable names the object), `DM_PROBE_HEAPFRAME=<vtable offset in main>` prints heaps of that class as frame heaps (head, tail, free), and `DM_PROBE_HEAPRAW=1` dumps the heap object's first 0x130 bytes, to find another class's fields.

 `DM_PROBE_ALL=1` adds x9–x18; `DM_PROBE_X2=<bytes>` dumps the memory x2 points to.
+`DM_PROBE_FRAMES=3000-3001` restricts observations to those presented frames
+(a single frame is also accepted); `DM_PROBE_LIMIT=20` caps printed hits per
+probe. These controls bound a late-game observation without logging startup.

 **Who writes a word: `DM_PROBE_WATCH=<address>[+<length>]`** (hex guest address, 8 bytes by default; the bytes must lie in one 16 KB host page). The host write-protects that page as soon as it is writable and keeps it protected: the fault handler carries out every store that lands on the page through a second, writable mapping of it (STR / STUR / STP of general and SIMD registers in their immediate, pre- and post-index forms, register-offset STR, DC ZVA), so the stores of a busy page all pass through it. A store that reaches the watched bytes prints `WATCH … store of N bytes at …: old → new`, the thread, pc, lr and a backtrace. Other instructions (atomics, exclusives) get a `not emulated` line once per pc and run on an unprotected page, protected again a millisecond later. A guest object sits at the same address in every run (the guest heap is laid out the same), so take the address from a probe's register dump first. STEINS;GATE RE:BOOT's render size setting (engine object + 0x6000): the static search for stores with that offset, conditional probes on 25 candidate stores and a wrapper around the host's `memcpy` / `memset` (the guest's go straight to the host's libc, `hle_table.c`, where guest probes cannot see them) all found nothing; a watch that unprotected the page for every fault missed the write among the bursts of stores to the engine object's page; this one found it on the first run (`main+0x389e90`, a setter its start-up code calls with 0).

@@ -142,7 +142,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - SIGSEGV in `vi::SetLayerCrop`. → [vi layer and display calls without vi service objects](../reference/services.md#vi-layer-and-display-calls-without-vi-service-objects)
 - 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)
+- the first boot crashed in `oe::DisableAlbumImageTakenNotification`: a null read at `sdk+0x338070`, the first call in `nnMain` (`main+0x6b188c`); enabling those notifications also gives a null read at `sdk+0x33b0f0` (Xenoblade Chronicles 3); 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)
@@ -328,6 +328,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - streamed textures read 0 (magenta on Monster Hunter Rise's title screen's clothes). → [Sparse textures bound from their first mapped level](../reference/rendering.md#sparse-textures-bound-from-their-first-mapped-level)
 - the process ends at once (the window vanishes), with `-[MTLTextureDescriptorInternal validateWithDevice:]` failing an assertion such as `requests 9 mipmap levels, but the dimensions (1, 1, 1) can only support 1`, or `sampleCount must be > 1 for multisample textures`. → [Texture fields Metal refuses](../reference/rendering.md#texture-fields-metal-refuses)
 - ≈13 MB/frame leak. → [Transient render targets shared by storage](../reference/rendering.md#transient-render-targets-shared-by-storage)
+- a model has correct geometry, base colour and normals but appears as a black silhouette; deferred lighting's material-ID depth texture stays zero. → [IMAGE usage does not split texture storage](../reference/rendering.md#image-usage-does-not-split-texture-storage)
 - GPU memory grows with every image shown; the memory line's `shared textures` count keeps rising. → [Shared-storage entries taken over by a new texture](../reference/rendering.md#shared-storage-entries-taken-over-by-a-new-texture)
 - GPU memory grows as resolution changes or rooms load, although almost every retained texture still has a live NVN descriptor; reclaiming unused sharing-table entries cannot recover it. → [Live resolution variants share physical image storage](../reference/rendering.md#live-resolution-variants-share-physical-image-storage)
 - the whole picture came out darker and more saturated on the Mac and the iPhone: comparing with a Switch screenshot, every channel at every level equalled the Switch value decoded from sRGB once more (0.55 → 0.26, 0.75 → 0.52). → [Render-target views that change the format](../reference/rendering.md#render-target-views-that-change-the-format)
@@ -364,6 +365,8 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - a red line: the 8-bit index buffer had GPU-side writes in the submission (indices read stale). → [8-bit indices widened on the GPU after a write in the same submission](../reference/rendering.md#8-bit-indices-widened-on-the-gpu-after-a-write-in-the-same-submission)
 - indirect draws whose record the GPU writes earlier in the frame cannot be read on the CPU at replay; GPU-driven effects drawn as indirect quads, which Metal has no primitive for. → [Indirect draws and dispatches read by the GPU](../reference/rendering.md#indirect-draws-and-dispatches-read-by-the-gpu)
 - an empty G-buffer pass (the title screen showed only the UI); later, the submitting thread waiting for the GPU before each such draw (36–54 fps in Kamura Village on an iPhone). → [Indirect-count multi-draws](../reference/rendering.md#indirect-count-multi-draws)
+- the HUD and sky render but the battle's characters and terrain are absent; the replay logs `draw skipped: rasterizer discard` between model draws. → [Rasterizer discard preserves vertex storage writes](../reference/rendering.md#rasterizer-discard-preserves-vertex-storage-writes)
+- exploding meshes and displaced model parts after vertex-only work; native post-vertex captures can contain valid positions for some meshes while other palettes are incorrect. → [Primitive barriers order vertex storage writes](../reference/rendering.md#primitive-barriers-order-vertex-storage-writes)
 - characters collapsed into one thin spike (vertex streams read from feedback buffers that were never written); a character covered by a solid box while it moves (its fur shells pulled toward lagged positions that were never written); a pass dump shows draws with a vertex shader and no fragment shader (`fs=-`) before the draws that read them, and `Y tf begin … written 0/0/0/0` with no `T … tf capture` line. → [Transform feedback captured by a vertex entry point](../reference/rendering.md#transform-feedback-captured-by-a-vertex-entry-point)
 - `SIGSEGV` in `-[AGXG18GFamilyComputeContext_mtlnext performEndEncoding]` (from `rp_begin_ex_`, or in its `dealloc` on the encoding thread); Metal validation (`MTL_DEBUG_LAYER=1 MTL_DEBUG_LAYER_ERROR_MODE=nslog`) names it: "Command buffer already has an open command encoder". → [Helper encoders end the open copy encoder first](../reference/rendering.md#helper-encoders-end-the-open-copy-encoder-first)
 - the whole picture upside down (Absolum's title logo and "Press Any Button" mirrored top to bottom); a red-reporting stub for `nvnCommandBufferSetViewportSwizzles`. → [Viewport swizzles applied as a flipped viewport](../reference/rendering.md#viewport-swizzles-applied-as-a-flipped-viewport)
@@ -406,6 +409,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - 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)
+- archive shaders are absent from the map; the two boot shaders in `main` are reported as raw controls with no identified code area. → [Monolith ARH1, XBC1 and HCPS/SLCT shader containers](../reference/shaders.md#monolith-arh1-xbc1-and-hcpsslct-shader-containers)
 - 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)
 - shaders the game draws are missing from the map (`DM_NVN_SHDUMP`), in no RomFS file and no module, stored or LZ4, and not compiled at run time. → [Shaders in a zlib stream in main (YEBIS .tfxo)](../reference/shaders.md#shaders-in-a-zlib-stream-in-main-yebis-tfxo)
 - two `[RED] shader … is not in the shader map` lines and `program has no vertex shader` as the opening movie starts; every movie drew white. → [Loose control sections and code areas in main (Bink 2)](../reference/shaders.md#loose-control-sections-and-code-areas-in-main-bink-2)
@@ -105,6 +105,7 @@ DM_NVN_PASSGPU	diagnostic	`=<file>` writes the GPU time of each render and compu
 DM_NVN_PASSGPU_SERIAL	diagnostic	`=1` makes each encoder on the main command buffer wait until all earlier work has finished, so `DM_NVN_PASSGPU` times read each pass's exclusive time (totals read high)
 DM_NVN_PREDICT	switch	`0` turns off the parallel precompile look-ahead: a pipeline miss compiles only the pipeline it needs
 DM_NVN_PROF	diagnostic	`=1` wraps the implemented NVN functions in a timing wrapper and collects the time each takes during recording
+DM_NVN_PROGTRACE	diagnostic	`=<vertex keys,...>` logs matching ProgramSetShaders calls and their guest backtraces
 DM_NVN_REGLUM	diagnostic	`=x0,y0,x1,y1:<file>` (normalized coordinates): on every present records the present number, time, mean luminance and a grayscale thumbnail of that region, to check animation speed against frame rate
 DM_NVN_RT_RANGE	switch	`=1` (experimental, off by default; with `DM_NVN_TBDR=1`) keeps the pass open when the game sets the same colour and depth target again with a view that differs only in its level count
 DM_NVN_RT_RANGE_CHECK	diagnostic	With `DM_NVN_RT_RANGE=1` and `DM_NVN_PASSDUMP`, renders a candidate pass both ways in the same frame and logs and writes the comparison (`CHECK_RESULTS.tsv`, `[RED]` on a difference)
@@ -141,6 +142,10 @@ DM_NVN_VCONV_CHECK	diagnostic	`=1` scans the indices again at every reuse of a c
 DM_NVN_VCONV_CPU	switch	`=1` converts integer vertex inputs fed formats of another class on the CPU, as before, instead of through the shader's shadow inputs (`<key>_bv`)
 DM_NVN_VCONV_RANGE	switch	`0` / `1` overrides the profile's `vconv_range`: CPU vertex conversion only over the vertices each draw reads instead of the whole bound range
 DM_NVN_VLOG	diagnostic	Set: logs each NVN memory pool as it is created (address, memory, size, flags) with a backtrace
+DM_NVN_VSDUMP	path	Output file for post-vertex captures; used with DM_NVN_VSDUMP_KEY and the frame/count bounds
+DM_NVN_VSDUMP_FROM	diagnostic	First presented frame eligible for a post-vertex capture (default 0)
+DM_NVN_VSDUMP_KEY	diagnostic	Vertex key to capture, or `*` for all eligible vertex shaders; storage-writing shaders are excluded
+DM_NVN_VSDUMP_MAX	diagnostic	Maximum post-vertex captures (default 1, or 256 for the wildcard key)
 DM_NVN_WAITS	diagnostic	`=1` turns on the wait ledger: every place a host or game thread waits on the CPU for the GPU or the encode thread, counted per frame
 DM_NVN_XV	switch	`0` turns off the cross-submission attachment validity ledger, which skips write-backs of colour attachments cleared before their next use (on by default, only with `DM_NVN_TBDR` on)
 DM_NVN_ZBLOCK	diagnostic	`=1` logs (`[zb]`) the all-zero blocks per level of each ASTC texture uploaded on the CPU (magenta when drawn); `=<folder>` also writes the raw storage of the first one per format and size
@@ -165,11 +170,13 @@ DM_PRESENT_UPSCALE	player	`metalfx` scales a picture smaller than its rectangle
 DM_PROBE	diagnostic	`=module:offset,…` (Mac only): puts a `brk` at each guest address and on every hit prints the thread, registers, strings they point to and a backtrace, then runs the original instruction
 DM_PROBE_ALL	diagnostic	Set: probe hits (`DM_PROBE`) also print x9–x18
 DM_PROBE_DUMP	diagnostic	`=module:offset:length[:file]`: on the first hit of any probe, writes that range of guest memory to a file (default `run/<platform>/probe_dump.bin`)
+DM_PROBE_FRAMES	diagnostic	Presented frame or inclusive `start-end` range in which guest probes print observations
 DM_PROBE_HEAP	diagnostic	`=<name>,<start>,<size>,<free list>,<used list>`: hex offsets of the sead heap layout the probes' `!xN` heap walk uses (default Mario's `0,28,30,b0,c8`)
 DM_PROBE_HEAPFRAME	diagnostic	`=<vtable offset in main>` (hex): the probes' heap walk treats heaps of that class as frame heaps (head +0xe0, tail +0xe8)
 DM_PROBE_HEAPNAME	diagnostic	`=<text>`: the probes' `!xN` heap walk reports only hits on heaps whose name contains the text
 DM_PROBE_HEAPRAW	diagnostic	Set: the probes' heap walk also prints the heap object's first 0x130 bytes
 DM_PROBE_HEAPWORDS	diagnostic	Set: the probes' heap walk also prints each used block's first two words
+DM_PROBE_LIMIT	diagnostic	Maximum printed hits per guest probe; 0 or absent leaves the count unlimited
 DM_PROBE_POLL	diagnostic	`=<address>[+<length>]` (hex, 4 bytes by default, up to 8): a host thread spins on the bytes and on each change snapshots every thread's pc, lr and frames, for writes the watch misses
 DM_PROBE_WATCH	diagnostic	`=<address>[+<length>]` (hex, 8 bytes by default): write-protects that host page and reports who writes these bytes, without a debugger
 DM_PROBE_X2	diagnostic	`=<bytes>` (hex): probe hits also dump that many bytes of the memory x2 points to
@@ -237,6 +244,7 @@ DM_TEST_NO_METAL4	diagnostic	`=1` treats the GPU as one without Metal 4, as on a
 DM_TEST_PRESENT_LAG	diagnostic	`=<every n>,<ms>[,<in a row>][;…]` sleeps that long before every nth present (and the ones after it in the run): a slow frame on demand
 DM_TEST_SWKBD_HOLD_MS	diagnostic	`=<ms>` (with `DM_KBD`) keeps the system keyboard up that long before `DM_KBD`'s text is returned, as a player typing would
 DM_TEST_SWKBD_LAG_MS	diagnostic	`=<ms>` sleeps that long between the keyboard handing back its text and the game's thread getting the result, to check that auto-accept does not act early
+DM_TEST_TRANSFORM_FEEDBACK	test	Path to the native t_transform_feedback executable used by the Python feedback tests
 DM_TEST_UP_LEAK	diagnostic	`=1` keeps the old upload staging block instead of retiring it, as before the fix, to measure the leak
 DM_TEST_UP_LOST	diagnostic	`=<batch index>` treats that batch of uploads as dropped, to exercise the re-upload path
 DM_THERMAL_CAP	switch	`0` turns off the thermal cap, which holds the frame step at 30 fps while iOS reports a serious or critical thermal state (on by default)
@@ -70,6 +70,7 @@ The guest address space, heaps, pools, virtual address memory and the allocators
 - **Fix:** shrink the node area with a rewriter cut for that SDK build (`dm/tools/sdk_patches.py`): the two `mov w8/w9, #0x3fe00000` in the AddressRegionManager constructor → 62 MB.
 - **Finding the constructor on a new `sdk` build:** search its text for the two `mov w8 / w9, #0x3fe00000` words (`08fca752`, `09fca752`); `sdis.py --fn` names the function `AddressRegionManagerC2Emm` from its dynamic symbols. SDK 7 / 8 builds load the node area with `orr w8 / w9, wzr, #0x3fe00000` (`e8230b32`, `e9230b32`), so the `mov` search finds nothing there.
 - **Found on** (the builds that take the cut):
+  - [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md) (SDK 12.3.9): constructor `sdk+0x346420`, cuts `+0x346430` / `+0x3464cc`, the same offsets as Engage's SDK 12.3.7, beside the large-region alignment cut. The game's own allocator needs a separate compact layout.
  - [Pokémon Brilliant Diamond](../../dm/games/0100000011D90000/README.md) (Unity 2019.4.27f1, SDK 11.4.3): constructor `sdk+0x34e0b8`, node-area limit `+0x34e0c8` and occupied length `+0x34e164`. The original node area prevents StandardAllocator's pool-minus-used reservation from fitting at startup.
  - [Fire Emblem Engage](../../dm/games/0100A6301214E000/README.md): the two `mov w8/w9, #0x3fe00000` in the constructor → 62 MB; without it the alias region ran out of address space minutes into play ([Size-aligned regions exhaust the alias address space](#size-aligned-regions-exhaust-the-alias-address-space)).
  - [Ace Attorney Investigations Collection](../../dm/games/010005501E68C000/README.md) (Unity 2022.3, SDK 17.5.3): `StandardAllocator::Initialize` aborts in `nninitStartup` on SDK 17, then the alias region fills in minutes; the same node area as Engage's SDK; the constructor is at `sdk+0x139458` on that build and takes the same cut.
@@ -111,6 +112,7 @@ The guest address space, heaps, pools, virtual address memory and the allocators
  - [Super Smash Bros. Ultimate](../../dm/games/01006A800016E000/README.md) (SDK 8.2.1): alignment at `sdk+0x336d5c`, cut as DRAGON QUEST XI S's SDK 7.3.2 build.
  - [Paper Mario: The Thousand-Year Door](../../dm/games/0100ECD018EBE000/README.md) (SDK 15.3.2, Celeste's `sdk` build): `AllocateAddressRegion`'s alignment at `sdk+0x12e2e0`, as Absolum's did.
  - [Pikmin 4](../../dm/games/0100B7C00933A000/README.md) (SDK 10.4.8): the 10.4.1 build's offsets.
+  - [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md) (SDK 12.3.9): large-region alignment at `sdk+0x34662c`, 1 GB to 32 MB; its compact 4 GB allocator reservation then fits after the 62 MB node area.
 - **Note:** FINAL FANTASY RESONANCE's SDK 20.5.17 build aligns regions of 2 MB and more to 2 MB only (`VammManager::AllocateAddressRegion`, `sdk+0x145e00`), so it needs no alignment cut.

 ### Size-class alignment doubles Unity's regions
@@ -163,6 +163,14 @@ The replay on Metal (`dm/gfx/nvn/nvn_mtl.m`) and the Vulkan and OpenGL layers. E
 - **Fix:** host textures are shared by `(pool, offset, description)`.
 - **Found on:** [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md).

+### IMAGE usage does not split texture storage
+
+- **Symptom:** a model has correct geometry, base colour and normals but appears as a black silhouette; deferred lighting's material-ID depth texture stays zero.
+- **Cause:** render and stencil-sampling descriptors name the same pool, offset and layout, differing only by `NVN_TEXTURE_FLAGS_IMAGE`. Including that usage flag in the storage key made two independent Metal textures; sampling read the empty one's stencil.
+- **Fix:** the storage key excludes `IMAGE`. Sampling shares the original texture. A writable image binding prepares any needed `ShaderWrite` usage before opening its encoder; the upgrade copies the pixels and replaces every name of the old texture. Upgrading while encoding arguments would end the active encoder underneath its caller.
+- **Check:** `t_stencil_view` writes both planes, reads through an IMAGE stencil descriptor and copies to separate storage under Metal validation; `t_upgrade_alias` checks image-write preparation and every live/finalized alias.
+- **Found on:** [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md).
+
 ### Shared-storage entries taken over by a new texture

 - **Symptom:** GPU memory grows with every image shown; the memory line's `shared textures` count keeps rising.
@@ -500,17 +508,36 @@ The replay on Metal (`dm/gfx/nvn/nvn_mtl.m`) and the Vulkan and OpenGL layers. E
 - **Found on:** [Monster Hunter Rise](../../dm/games/0100B04011742000/README.md).
 - **Note:** the replay used to commit its work and wait for the GPU before the first such draw after each GPU write, then read the count on the CPU (`mdi_wait`, now only for quads, fans and 8-bit indices). In Kamura Village that wait held the submitting thread 14 of its 20 ms a frame on an iPhone (36–54 fps). On the Mac the same scene's submit took 20 ms of wall time for 6.1 ms of work, and now 5.1 ms, all of it work.

+### Rasterizer discard preserves vertex storage writes
+
+- **Symptom:** the HUD and sky render but the battle's characters and terrain are absent; the replay logs `draw skipped: rasterizer discard` between model draws.
+- **Cause:** the replay returned from every discarded draw without active transform feedback. Rasterizer discard only suppresses rasterization: vertex storage writes and atomics still execute. Metal additionally requires a void-returning vertex function when rasterization is disabled.
+- **Fix:** the offline builder adds `main0_discard` to vertex shaders with writable reflected bindings. It retains the original body and early returns and drops only the stage outputs. Replay uses that entry in a separate pipeline with no fragment stage and rasterization disabled, with the ordinary array/indexed/instanced draw machinery and GPU-write tracking. Switching discard marks the pipeline dirty; discarded draws cannot prove an attachment fully overwritten. Pure vertex shaders with neither storage writes nor active feedback can still be skipped. Transform feedback retains its capture entry.
+- **Check:** `t_rast_discard` passes 19 checks under Metal validation: array/indexed storage writes despite an empty scissor and both faces culled, two instances, unchanged clear colours, normal rendering between two discard phases, and an ordered storage-write consumer. `test_discard.py` compiles a generated entry with side effects before two different returns for Mac and iOS.
+- **Found on:** [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md#files-and-shaders); its note tracks the gameplay verification.
+
+### Primitive barriers order vertex storage writes
+
+- **Symptom:** exploding meshes and displaced model parts after vertex-only work; native post-vertex captures can contain valid positions for some meshes while other palettes are incorrect.
+- **Cause:** `ORDER_PRIMITIVES` was omitted from the pass-splitting condition after shader storage writes. A primitive barrier orders earlier vertex work before later vertex work, including its storage writes. Metal does not supply that ordering between draws in one render encoder. Xenoblade Chronicles 3 uses barrier `0x21` after vertex shaders update its bone matrices; fragment-only barrier handling left dependent vertex work unordered.
+- **Fix:** a primitive barrier after shader writes ends the render pass, as fragment barriers already do. The next encoder's full consumer barrier supplies the required visibility.
+- **Check:** `t_rast_discard` passes 19 GPU checks under Metal validation, including a 4,096-vertex storage producer followed by a primitive barrier and a vertex consumer in the same submission.
+- **Found on:** [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md#files-and-shaders); `boot24barrier` removes the exploding meshes.
+
 ### Transform feedback captured by a vertex entry point

 - **Symptom:** characters collapsed into one thin spike (vertex streams read from feedback buffers that were never written); a character covered by a solid box while it moves (its fur shells pulled toward lagged positions that were never written); a pass dump shows draws with a vertex shader and no fragment shader (`fs=-`) before the draws that read them, and `Y tf begin … written 0/0/0/0` with no `T … tf capture` line.
 - **Cause:** a game that skins on the GPU with transform feedback: a vertex-only program runs under `SetRasterizerDiscard` between `BeginTransformFeedback` / `EndTransformFeedback` and writes into feedback buffers that later draws read as vertex streams. The host ignored all of it, so the streams kept whatever was there. Metal has no feedback stage.
 - **Fix:**
-  - the layout comes from the program's control section, where the NVN driver would read it (`tf_layout` in `dmsh_build.py`): +0x6dc u16 strides, +0x6e8 per buffer the number of captured outputs, +0x6ec buffer count, and per buffer b the captured outputs (0x20 + n, one byte each) from +0xdc + 0x80 b and their component counts (one byte each) from +0x4dc + 0x80 b, each captured from x on and packed in that order. A form that does not add up (the components times 4 must make the stride) gets a red line and no capture entry point, so one shader cannot stop a game's build;
-  - such a vertex shader gets a third entry point, `main0_tf` (`vs_tf`): the same body, returning nothing, storing the captured components as their bits into `device uint *` buffers at Metal vertex buffer slots 8–11 (`TF_SLOT`), indexed by vertex ID, a buffer's outputs one after another within the vertex's stride; the reflection lists `entry_tf` and `tf_strides`;
-  - the binds, begin / end / pause / resume and rasterizer discard are records (`OP_TF_BIND`, `OP_TF_STATE`, `OP_RAST_DISCARD`). At replay a draw under rasterizer discard runs `tf_draw` in `nvn_mtl.m`, and so does a rasterized draw while feedback is active and its vertex shader has a capture entry point (feedback records every draw between Begin and End; the draw then rasterizes as before): with feedback active and a capture entry point, an encoder of its own with no attachments, a pipeline with rasterization off, the buffers bound at their current offset less first × stride, the offsets advanced and the written ranges queued as GPU writes (`pend_add` kind 4); without feedback the draw rasterizes nothing, as on the GPU (it used to draw, depth writes included). Only plain array draws are captured; indexed, indirect or instanced ones give a red line. `DM_NVN_PASSDUMP` prints the records as `Y tf …` and each capture as a `T … tf capture` line.
+  - the layout comes from the program's control section (`tf_layout` in `dmsh_build.py`): +0x6dc u16 strides, +0x6e8 per-buffer output counts, +0x6ec buffer count, and per buffer b the output locations (0x20 + n) at +0xdc + 0x80 b and matching **byte** component counts at +0x4dc + 0x80 b. Multiple outputs are interleaved in that order. Unknown forms get a red line and no capture entry;
+  - `main0_tf` (`vs_tf`) retains the body and stores the output bits into Metal vertex buffers 8–11, with the GPU-prepared capture record at buffer 14. The reflection lists `entry_tf` and `tf_strides`. Geometry programs give their output reader the same entry; it gathers the packed point/line/triangle indices while writing linear feedback records;
+  - `nvn_feedback.inc` keeps offsets and counts on the GPU. Begin clears them; End/Pause save them in the game's opaque 32-byte control buffer; Resume restores them. A compute helper bounds each capture by all active streams and complete primitives. `DrawTransformFeedback` uses the saved count as indirect array-draw arguments. Captures, counter copies and consuming draws follow encoder ordering, and captured ranges enter GPU-write tracking. Plain vertex capture supports array point/line/triangle lists with or without rasterizer discard, including instances; indexed/indirect vertex capture remains open. Geometry capture uses the existing general assembly and compute path. `DM_NVN_PASSDUMP` prints the records as `Y tf …` and array captures as `T … tf capture`.
+- **Check:** `test_transform_feedback.py` compiles generated vertex/geometry entries for Mac and iOS; `t_transform_feedback` checks interleaving, saved offsets, resume, bounded capture, GPU-produced geometry counts, complete primitives and a following draw that consumes the saved count in the same submission.
 - **Found on:** [Pokémon Legends: Z-A](../../dm/games/0100F43008C44000/README.md): the character preview. 28 such shaders (e.g. `681cc42e48c0bbc0` in the `special/imgui.trsha` pack; skin indices in attribute 12, weights in 13) write positions (location 0 xyz, 12 bytes a vertex) and packed normal and tangent bits (locations 1 and 2, x only, 4 bytes) into three feedback buffers; the character draws read them as vertex streams 0–2 (`RGB32F`, `RGB10A2SN`, `RGB10A2SN`). Every buffer in this game holds one output from x on (strides 12/4/4, 12/12/12, 12/4, 12). The game draws these while its frame target is bound. The pass dump showed a pass of 19 draws with `fs=-` before the character pass. Tears of the Kingdom's 14,766 and Ace Attorney's 2,475 vertex shaders have no feedback at all.
  - [Donkey Kong Country: Tropical Freeze](../../dm/games/0100C1F0051B6000/README.md): DK's fur. While DK moved, his fur shells (an instanced draw, `648304c43bbe3463` / `5b5ffba62bca9e0a`, three shells on the body's index buffer) covered him as one solid box. Each frame a vertex-only draw (`6979dd3cdbd62373` / `7dcdaed991c74bfe`, 3,400 vertices) captures the fur's lagged positions into one buffer, two outputs per vertex (xyz each, stride 24; output count 2, component bytes `03 03`), and the shells read them as vertex attribute 6. Two things kept the buffer at zero (`DM_NVN_KEYDUMP=5b5ffba62bca9e0a`: stream 6 held zeros and `2.3e24`): the old one-output rule rejected the layout (`[RED] … transform feedback buffer 0 has a layout of unknown form`), and the game draws the capture with rasterization on (no fragment shader, colour mask 0, no `SetRasterizerDiscard`), which the host did not capture. Four of its 11,644 vertex shaders have this form; over Z-A's and this game's 23,277 vertex shaders together, the old and new readings give identical layouts and capture code for all 28 shaders the old rule took.

+- **Found on:** [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md#files-and-shaders): effect vertex and geometry programs use interleaved outputs, including eleven outputs in a 172-byte stream, and draw from saved feedback counts.
+
 ### Helper encoders end the open copy encoder first

 - **Symptom:** `SIGSEGV` in `-[AGXG18GFamilyComputeContext_mtlnext performEndEncoding]` (from `rp_begin_ex_`, or in its `dealloc` on the encoding thread); Metal validation (`MTL_DEBUG_LAYER=1 MTL_DEBUG_LAYER_ERROR_MODE=nslog`) names it: "Command buffer already has an open command encoder".
@@ -355,11 +355,13 @@ 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:** the first boot crashed in `oe::DisableAlbumImageTakenNotification`: a null read at `sdk+0x338070`, the first call in `nnMain` (`main+0x6b188c`); enabling those notifications also gives a null read at `sdk+0x33b0f0` (Xenoblade Chronicles 3); 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).
+- **Cause:** the SDK asks am through the applet proxy the host never made.
+- **Album notification enable:** `oe::EnableAlbumImageTakenNotification` follows the existing disable takeover. The host produces no Horizon album captures, so enabling notifications produces no messages. Found on [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md) (SDK 12.3.9, `main+0x127fac4`).
+- **Inactivity extension:** `oe::GetUserInactivityDetectionTimeExtensionMode` also reads the absent SelfController (`sdk+0x3391f0`, after Xenoblade Chronicles 3's opening movie). The host retains the requested mode: 0 initially or when either setter receives false, 1 for `SetUserInactivityDetectionTimeExtended(true)`, 2 for its `Unsafe` form. The boolean query is mode != 0, matching SDK 12.3.9's implementation. This state does not make the host turn the screen off. `t_oe` checks both setters and both queries.
+- **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). 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)).

 ### The startup logo
@@ -157,6 +157,15 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b
 - **Found on:** [Bayonetta 2](../../dm/games/01007960049A0000/README.md): 1,856 distinct keys (1,547 pixel, 309 vertex) in `shader.dat`, `shader2.dat` and `shadereff.dat` of `data000.cpk` (the split about 2.5 minutes); [Bayonetta](../../dm/games/010076F0049A2000/README.md): 458 distinct keys (236 pixel, 222 vertex), the members of `shader.dat` in `data00.cpk` (about 1.7 minutes). All build for both.
 - **Not checked:** 13 Sentinels and Catherine: Full Body's CPKs (no files on the drives).

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

 - **Symptom:** the split finds no shaders, although `main` contains NVN control sections; it reports that no data record or code-address pair identifies their code areas.
@@ -423,8 +432,8 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b

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

 ### Geometry shaders that only pick the layer
@@ -45,8 +45,16 @@ Guest threads, their synchronization and how the host schedules them. Each entry
 - **Symptom:** crashes around movie changes in the game's own player threads, about one run in two: `SignalEvent` → `MultiWaitImpl::NotifyAndWakeupThread` on a null MultiWait at address 0x20; `WaitAny` walking freed holders; a worker reading a freed object at `main+0x225338`. The other way round, a handoff that waits too widely: a black screen at boot that never loads a texture (`packaged textures: first uploads 0`), or a load that never ends behind an animated loading screen, with the thread dump's `same-core handoffs N, total T ms` larger than the run itself.
 - **Cause:** Horizon schedules the highest-priority runnable thread on a core ([Mesosphere's scheduler](https://github.com/Atmosphere-NX/Atmosphere/blob/master/libraries/libmesosphere/source/kern_k_scheduler.cpp)). A lower-priority thread that wakes a higher-priority thread pinned to its core cannot continue until that work blocks. The host's pthreads run in parallel. The earlier approximation also confused host dispatch state with guest readiness: `wake()` cleared the guest wait before the pthread resumed, but `thread_info` could still report that pthread waiting. The caller then continued too soon. Its 5 ms timeout opened the same hole under a longer host delay; the recorded `yield_pending` flag was unused unless the profile enabled checks after every thread SVC. The first unbounded handoff then waited for every runnable higher-priority thread pinned to the caller's core, not only the ones its wake made runnable. The host runs a lower-priority thread beside such threads at all other times, and a core's busy loaders, job workers and pollers are seldom all blocked at the same moment, so a wake could hold its caller for minutes.
 - **Fix:** every game's thread SVC honours `yield_pending`. Before returning, it waits for the threads its wakes made runnable that are higher priority and pinned to the caller's core, and for those such a thread makes runnable in turn while the handoff lasts (Dread's dispatcher wakes the worker), until each blocks, sleeps, yields to any thread (`SleepThread(-2)`, which takes it off its core so that lower-priority threads run there), exits, is suspended, or enters explicit host GPU work (`svc_thread` / `dm_yield_to_core_peers` in `hle_thread.c`). A wake records the thread in the waker's current handoff and in those the waker belongs to (`note_runnable`, at most four); blocking ends that (`hand_leave`), and a handoff ends when its caller's thread SVC returns (`hand_epoch`). Guest runnable state is authoritative, including the wake-to-pthread-resume interval; a host dispatch snapshot cannot release the handoff, and there is no timeout that lets unfinished work through. A higher-priority peer the wake did not make runnable does not hold the caller, as at every other point on the host. Host-created helpers and threads that can run on another core are excluded.
- **Check:** `make -C dm/host/tests GAME=010093801237C000 preempt` delays a woken pthread's dispatch by 1 ms and 20 ms through real event, mutex and condition-variable SVCs. The worker's payload and exit must precede the signaller's return. The original implementation failed both delays; changing readiness alone fixed 1 ms but still failed 20 ms. Equal-priority, other-core, migratable, host-helper, suspended and host-GPU-work peers do not block the caller. Its handoff-scope cases add a busy higher-priority peer that nobody woke: the woken worker must still finish first and the signaller must return (the whole-core wait never returns), a worker woken by the woken dispatcher must finish before the signaller returns (fails without the recorded chain), and a woken thread that yields to any thread releases the signaller (never returns without that rule).
+- **Check:** `make -C dm/host/tests GAME=010093801237C000 preempt` delays a woken pthread's dispatch by 1 ms and 20 ms through real event, mutex and condition-variable SVCs. The worker's payload and exit must precede the signaller's return. The original implementation failed both delays; changing readiness alone fixed 1 ms but still failed 20 ms. Equal-priority, other-core, migratable, host-helper, suspended and host-GPU-work peers do not block the caller. Its handoff-scope cases add a busy higher-priority peer that nobody woke: the woken worker must still finish first and the signaller must return (the whole-core wait never returns), a worker woken by the woken dispatcher must finish before the signaller returns (fails without the recorded chain), and a woken thread that yields to any thread releases the signaller (never returns without that rule). The native-mutex case wakes a higher-priority pinned peer while holding the mutex it needs; the signal must return at that native block, and the peer finishes after unlock (hangs without the wait annotation).
 - **Found on:** [Metroid Dread](../../dm/games/010093801237C000/README.md), whose per-build movie request-ring gates and `core_preempt on` setting are now removed; the whole-core wait on [The Legend of Zelda: Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md) ([issue #247](https://github.com/seedds/nxnative/issues/247)) and Donkey Kong Country: Tropical Freeze ([issue #259](https://github.com/seedds/nxnative/issues/259), section 6; no profile here).
+- **Native helper joins:** `dm_host_thread_join` also releases the guest core
+  during a native helper's lifetime, preserving an enclosing host-work scope.
+  [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md) exposed an
+  audio-unit setup join in `AKEventManager` during a later stall. That sample
+  establishes the missing scheduler annotation, but does not establish an
+  unbounded Apple audio-unit construction deadlock. The preemption regression
+  gates a native helper, verifies that a same-core signaller returns while the
+  peer joins it, then checks completion, the join result and nested host work.
 - **Scope:** these are cooperative thread-SVC handoffs, not a complete asynchronous guest scheduler. The optional `core_preempt` profile key (`DM_PREEMPT=0` / `1` overrides it) additionally checks already-runnable peers after every thread SVC; it defaults off. Historically that broader policy moved [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md)'s first frame from 0.6 s to 2.9 s by waiting for busy decompression threads. Wake-triggered handoffs apply with either setting.

 ### A host lock held across a same-core handoff
@@ -56,6 +64,7 @@ Guest threads, their synchronization and how the host schedules them. Each entry
 - **Fix:** `dm_host_lock` (`hle_thread.c`) takes a host lock that HLE code holds while it runs guest code. A guest thread that has to wait for it is in a [native HLE wait](#native-hle-waits-release-the-guest-core) (`dm_thread_native_wait`) while it waits: blocked for `thr_runnable`, and out of the handoff that made it runnable, so the handoff does not wait for it. The movie decoder takes `dmu` and the deferred-event lock `g_dfmu` with it (`hle_movie.c`). Releasing `dmu` before signalling instead would let a signal land after `Stop` has returned, on events the game is tearing down (the reason `md_defer_body` signals under its lock).
 - **Check:** `make -C dm/host/tests GAME=010093801237C000 preempt`, handoff scope "woken thread waiting for the signaller's host lock": the signaller holds a host lock across its signal and the woken higher-priority worker waits for it with `dm_host_lock`. The signaller must return before the worker finishes; without the native wait around the lock the test never returns.
 - **Found on:** [Kirby and the Forgotten Land](../../dm/games/01004D300C5AE000/README.md): the first guide movie ("A Jump", 480×270 VP9) froze the game on the first walk from the beach, every time. The job worker (priority 46) and the movie threads (priority 44) are all pinned to core 0.
+- **Found on:** [Xenoblade Chronicles 3](../../dm/games/010074F013262000/README.md), whose 140-second opening VP9 movie stopped at its first output release. A live stack sample identified the input thread in `md_pump` waiting on `dmu`, and the output thread holding it inside `SignalEvent` → `dm_yield_to_core_peers`.

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