Pokémon Legends: Z-A plays on the Mac: Trinity packs through the game's Oodle, transform feedback, stencil views, the base instance in the driver segment

- Profile 0100F43008C44000 (2.0.2, SDK 19.3.11): 5 GB layout with two cuts (StandardAllocator 12 → 4 GB, node area), mem_map_limit (the pool enforced as on Horizon), shader_oodle.
- Shaders: Trinity packs (arc/data.trpfs, FlatBuffers, Oodle Mermaid) decoded with the game's own OodleLZ_Decompress in the Mac host; 28,603 shaders.
- SDK 19: GetAudioOutBufferCount, crypto random bytes, ToPosixTime with a rule, NVN state getters, CommandBufferFenceSync, the 64-bit WaitForAddress type.
- Graphics: transform feedback (vertex capture entry point, rasterizer discard), layered passes (the colour grading LUT), point size dropped beside a layer output,
  function constants specialized, a stencil-only view for STENCIL-mode texture views (striped lighting), sampler min / max reduction (Metal 26),
  the draw's base instance at c0 + 0x644 (occlusion boxes; characters' legs folded back from the hips), the zero cube for a 2D texture in a cube slot.
- Shader chain changed (dmsh_build.py): every profile's dmsh_chain.h rewritten; the other profiles' c1_full lists carried over (that change touches no c1 read).

dm/games/0100000000010000/c1_full.tsv

@@ -1,3 +1,3 @@
-# chain b-f53065c05507 2026-10-20T06:37:59+08:00
+# chain b-166a14482eae 2026-10-20T06:37:59+08:00 (list carried over from b-f53065c05507: that chain change did not touch any c1 read)
 6405f3e55facc7c9
 d9097b8aba3a7080

dm/games/0100000000010000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100000000010000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "f53065c05507"
+#define DMSH_CHAIN "166a14482eae"

dm/games/010005501E68C000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-7aa26e75458e 2026-10-20T18:17:18+08:00
+# chain b-c792d8a4c610 2026-10-20T18:17:18+08:00 (list carried over from b-7aa26e75458e: that chain change did not touch any c1 read)
 0ffd2c92087e4b54
 69d3868c8d828fe6
 8a06e3dda01235c1

dm/games/010005501E68C000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010005501E68C000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "7aa26e75458e"
+#define DMSH_CHAIN "c792d8a4c610"

dm/games/010015100B514000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-48e79badc676 2026-10-20T06:18:38+08:00
+# chain b-dde5a04eafc6 2026-10-20T06:18:38+08:00 (list carried over from b-48e79badc676: that chain change did not touch any c1 read)
 35cbc686dd8461b8
 4872bec9a5567f8b
 54a618792b2d5f0f

dm/games/010015100B514000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010015100B514000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "48e79badc676"
+#define DMSH_CHAIN "dde5a04eafc6"

dm/games/010028600EBDA000/c1_full.tsv

@@ -1,3 +1,3 @@
-# chain b-aa62503b91c1 2026-10-20T06:23:06+08:00
+# chain b-821df0347986 2026-10-20T06:23:06+08:00 (list carried over from b-aa62503b91c1: that chain change did not touch any c1 read)
 25c5829f891f9b2c
 976839d3ba92c8f3

dm/games/010028600EBDA000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010028600EBDA000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "aa62503b91c1"
+#define DMSH_CHAIN "821df0347986"

dm/games/01002B00111A2000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-cc188a5e9879 2026-10-20T07:05:13+08:00
+# chain b-2ec808c15267 2026-10-20T07:05:13+08:00 (list carried over from b-cc188a5e9879: that chain change did not touch any c1 read)
 0002d193f6c52811
 0010e810186cdc91
 001299ca51e856c2

dm/games/01002B00111A2000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 01002B00111A2000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "cc188a5e9879"
+#define DMSH_CHAIN "2ec808c15267"

dm/games/01003FC01670C000/c1_full.tsv

@@ -1 +1 @@
-# chain b-2c262d3399d1 2026-10-20T06:20:26+08:00
+# chain b-b0883f45f985 2026-10-20T06:20:26+08:00 (list carried over from b-2c262d3399d1: that chain change did not touch any c1 read)

dm/games/01003FC01670C000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 01003FC01670C000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "2c262d3399d1"
+#define DMSH_CHAIN "b0883f45f985"

dm/games/01005270232F2000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-ba76ba138e78 2026-10-20T21:58:38+08:00
+# chain b-57e1140395a0 2026-10-20T21:58:38+08:00 (list carried over from b-ba76ba138e78: that chain change did not touch any c1 read)
 006a737523e70720
 00da789b4ca18c03
 00e5f2d3b5851fab

dm/games/01005270232F2000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 01005270232F2000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "ba76ba138e78"
+#define DMSH_CHAIN "57e1140395a0"

dm/games/0100535012974000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100535012974000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "675e2f28a97a"
+#define DMSH_CHAIN "101896cbf0ed"

dm/games/010055D009F78000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-b6dcbdf4b4ba 2026-10-20T07:05:43+08:00
+# chain b-6f195117b2ca 2026-10-20T07:05:43+08:00 (list carried over from b-b6dcbdf4b4ba: that chain change did not touch any c1 read)
 1a2b65cd05ff557c
 21c7ff0d9591a637
 33e4340e5bcbc5c6

dm/games/010055D009F78000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010055D009F78000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "b6dcbdf4b4ba"
+#define DMSH_CHAIN "6f195117b2ca"

dm/games/010069401ADB8000/c1_full.tsv

@@ -1 +1 @@
-# chain b-239c170b7616 2026-10-20T21:23:08+08:00
+# chain b-a74254ba6a54 2026-10-20T21:23:08+08:00 (list carried over from b-239c170b7616: that chain change did not touch any c1 read)

dm/games/010069401ADB8000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010069401ADB8000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "239c170b7616"
+#define DMSH_CHAIN "a74254ba6a54"

dm/games/01006BB00C6F0000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-75a11052925e 2026-10-20T06:31:05+08:00
+# chain b-35936311a3c9 2026-10-20T06:31:05+08:00 (list carried over from b-75a11052925e: that chain change did not touch any c1 read)
 26c5fd7adbc12cc0
 67f23f4ad5925e06
 96a3e288731b2433

dm/games/01006BB00C6F0000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 01006BB00C6F0000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "75a11052925e"
+#define DMSH_CHAIN "35936311a3c9"

dm/games/01007EF00011E000/c1_full.tsv

@@ -1,3 +1,3 @@
-# chain b-d3a07e79011d 2026-10-20T14:40:41+08:00
+# chain b-10befbce38df 2026-10-20T14:40:41+08:00 (list carried over from b-d3a07e79011d: that chain change did not touch any c1 read)
 21fe759b8114941c
 60e2041ed375cae1

dm/games/01007EF00011E000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 01007EF00011E000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "d3a07e79011d"
+#define DMSH_CHAIN "10befbce38df"

dm/games/010093801237C000/c1_full.tsv

@@ -1,3 +1,3 @@
-# chain b-eb6b1782c78e 2026-10-20T06:20:18+08:00
+# chain b-9c3f90e7f282 2026-10-20T06:20:18+08:00 (list carried over from b-eb6b1782c78e: that chain change did not touch any c1 read)
 57eb9f0bae677756
 f826f2b59cfa39a8

dm/games/010093801237C000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010093801237C000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "eb6b1782c78e"
+#define DMSH_CHAIN "9c3f90e7f282"

dm/games/010099C022B96000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010099C022B96000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "9a08b2229283"
+#define DMSH_CHAIN "78abd7b76d0d"

dm/games/0100A00019DE0000/c1_full.tsv

@@ -1 +1 @@
-# chain b-3f6065a844fa 2026-10-20T19:40:15+08:00
+# chain b-119d78df4ead 2026-10-20T19:40:15+08:00 (list carried over from b-3f6065a844fa: that chain change did not touch any c1 read)

dm/games/0100A00019DE0000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100A00019DE0000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "3f6065a844fa"
+#define DMSH_CHAIN "119d78df4ead"

dm/games/0100A6301214E000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-6a37b6ab3705 2026-10-20T06:20:07+08:00
+# chain b-d98265054359 2026-10-20T06:20:07+08:00 (list carried over from b-6a37b6ab3705: that chain change did not touch any c1 read)
 001275671d2c9985
 0021d75a567e2122
 00a98da1158c6f14

dm/games/0100A6301214E000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100A6301214E000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "6a37b6ab3705"
+#define DMSH_CHAIN "d98265054359"

dm/games/0100B04011742000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-94b2f6efaab8 2026-10-20T05:18:53+08:00
+# chain b-6631a1590f6d 2026-10-20T05:18:53+08:00 (list carried over from b-94b2f6efaab8: that chain change did not touch any c1 read)
 00096443b71e7fef
 001540067ae4cdf5
 002e5538fc05267e

dm/games/0100B04011742000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100B04011742000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "94b2f6efaab8"
+#define DMSH_CHAIN "6631a1590f6d"

dm/games/0100D9F01D474000/c1_full.tsv

@@ -1 +1 @@
-# chain b-e42dda1cbc03 2026-10-20T22:47:54+08:00
+# chain b-22c0a0e0ae72 2026-10-20T22:47:54+08:00 (list carried over from b-e42dda1cbc03: that chain change did not touch any c1 read)

dm/games/0100D9F01D474000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100D9F01D474000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "e42dda1cbc03"
+#define DMSH_CHAIN "22c0a0e0ae72"

dm/games/0100F2C0115B6000/c1_full.tsv

@@ -1,4 +1,4 @@
-# chain b-743f7a102a72 2026-10-20T13:29:14+08:00
+# chain b-330a5957853d 2026-10-20T13:29:14+08:00 (list carried over from b-743f7a102a72: that chain change did not touch any c1 read)
 35cbc686dd8461b8
 4872bec9a5567f8b
 759c14c6d029f531

dm/games/0100F2C0115B6000/dmsh_chain.h

@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100F2C0115B6000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "743f7a102a72"
+#define DMSH_CHAIN "330a5957853d"

dm/games/0100F43008C44000/bf_want.tsv

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

dm/games/0100F43008C44000/dmsh_chain.h

@@ -0,0 +1,2 @@
+// Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100F43008C44000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
+#define DMSH_CHAIN "21d0d9d3a966"

dm/games/0100F43008C44000/game.tsv

@@ -0,0 +1,55 @@
+# Game profile of Pokémon Legends: Z-A (Switch, base game v0 + update 2.0.2 = v393216). Plays on the Mac
+# Format: see dm/games/README.md
+# Game Freak's Trinity engine on SDK 19.3.11, NVN 55, Havok 2022.2. Every file lives in the Trinity packs of arc/data.trpfs (FlatBuffers, Oodle Mermaid); started from the provisional profile
+# scripts/build_ipa_from_nsp.py wrote from its survey (dm/games/profile_guess.py)
+
+# ---- Identity ----
+title_id	0100F43008C44000
+program_id	0100F43008C44000
+
+# ---- App ----
+name	Pokémon Legends: Z-A
+ipa	PokMonLegendsZA.ipa
+app_category	public.app-category.games
+pad_layout	simple
+
+# ---- Versions ----
+build	v393216	2.0.2	main=b1f12fd919eae86ab8a978317677e64b	sdk=df7a3960c74c3a4803f826c587153252
+package	v393216
+shader_versions	v393216
+
+# ---- Files ----
+# every module main imports from
+modules	main sdk
+romfs_marker	arc/data.trpfs
+zip_store	.bnk:.pck:.trpfd:.trpfs:.metallib
+
+# ---- Graphics ----
+graphics	nvn
+
+# ---- Memory ----
+# nninitStartup (main+0x1b2e650) initializes its StandardAllocator with virtual address memory over 12 GB of address space (StandardAllocator::Initialize(nullptr, 0x300000000), main+0x1b3e3ec),
+# which the Switch's 64 GB alias region holds. Two cuts for 2.0.2 (dm/tools/patch_nso.py) make it fit the 5 GB block proven on a phone: the allocator gets 4 GB and the SDK's node area 62 MB instead
+# of 1 GB, so the alias region holds 62 MB + 4 GB (SDK 19 puts the region 4 MB aligned right after the node area)
+mem_reserve	0x140000000
+mem_alias	0x110000000
+mem_stack	0x10000000
+app_pool	0xCD500000
+# nninitStartup allocates and frees all it can (pool − 64 MB) and the SDK keeps those pages mapped until a map fails: the pool is enforced as on Horizon, so it gives them back (hle_svc.c)
+mem_map_limit	on
+
+# ---- NVN driver constant segments (Mario's values, the translator's built-in layout, unless the engine says otherwise) ----
+nvn_tex_base	0x20
+nvn_tex_slots	32
+nvn_samp_base	0x568
+nvn_img_base	0x120
+nvn_img_slots	8
+nvn_ssbo_bases	vs=0x110 ps=0x510 cs=0x310
+
+# ---- Shader split ----
+# Game Freak's Trinity engine: every shader is a BNSH inside the Oodle-compressed packs of arc/data.trpfs (shader_split.py trpfs_jobs), named by the pack's first file:
+# the material shaders in system_resource/gfx2/shader, particle effects in ik_effectdata, UI layouts (__ArchiveShader.bnsh in .truiv) in uidata
+shader_groups	#arc/system_resource=system #arc/ik_effect=effect #arc/ik_pokemon=pokemon #arc/ui=layout /arc/=arc
+shader_note	Pokémon Legends: Z-A shaders
+# The packs are Oodle Mermaid streams; the shader split decodes them with the game's own OodleLZ_Decompress in the Mac host (dm/host/hle_oodle.c), which the NSP script builds for it
+shader_oodle	on

dm/games/0100F43008C44000/pso_manifest.tsv

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

dm/games/0100F43008C44000/shader_facts.tsv

@@ -0,0 +1,128 @@
+# 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
+05f3cca65c39e259	{"local":[32,1,1],"shared":0}
+09f37d0bb3c15f08	{"local":[32,1,1],"shared":0}
+0a3d2821485d7f55	{"local":[32,1,1],"shared":0}
+0c0992d2c341d097	{"local":[32,1,1],"shared":0}
+0cee38e62a1e0bf1	{"local":[8,8,1],"shared":0}
+0d0849eee6f9ffd4	{"local":[32,1,1],"shared":0}
+110bd7e69cbc27fa	{"local":[6,6,1],"shared":0}
+124236db30c92275	{"local":[32,1,1],"shared":0}
+124259be914eb00a	{"local":[32,1,1],"shared":0}
+16f13f6a7a4e7c3e	{"local":[32,1,1],"shared":0}
+1836a0fab4b80b5b	{"local":[32,1,1],"shared":0}
+1d3d5e5ceae87a47	{"local":[32,1,1],"shared":0}
+1d9b07f10373ef16	{"local":[32,1,1],"shared":0}
+210c2c516a0328c7	{"local":[32,1,1],"shared":0}
+21ad739f2a3fc373	{"local":[4,4,1],"shared":0}
+237a86b045c4c16c	{"local":[32,1,1],"shared":0}
+24fa9091aa642a86	{"local":[32,1,1],"shared":0}
+2b13baa4bcdc0fac	{"local":[32,1,1],"shared":0}
+2c33e8336f46923d	{"local":[32,1,1],"shared":0}
+2d1ae09b3af9bc24	{"local":[32,1,1],"shared":0}
+2d35718a0a439bd5	{"local":[4,4,1],"shared":0}
+2ed56106981c01dd	{"local":[32,1,1],"shared":0}
+2f227a0f100c1a33	{"local":[32,1,1],"shared":0}
+3eba81df1b7f44ce	{"local":[32,1,1],"shared":0}
+3ec0f09890eec4c3	{"local":[20,11,1],"shared":0}
+3fe0f0e077d818f2	{"local":[32,1,1],"shared":384}
+46643f602264b7c2	{"local":[32,1,1],"shared":0}
+49666cb704482a9c	{"local":[32,1,1],"shared":0}
+4a84c56d8dea5a2c	{"local":[32,1,1],"shared":0}
+4c6ce7bad5504c43	{"local":[32,1,1],"shared":0}
+4c7f3092d9316923	{"local":[8,8,1],"shared":0}
+4db045d8cf0abb49	{"local":[32,1,1],"shared":0}
+4e2a8a909c99d822	{"local":[12,12,1],"shared":0}
+517c3014d05e2271	{"local":[32,1,1],"shared":0}
+539ce06b647a1a04	{"local":[32,1,1],"shared":0}
+53edb9e5bdc1d0f7	{"local":[6,6,1],"shared":0}
+54fbdb4620ccbc34	{"local":[32,1,1],"shared":0}
+5aa2f2c5effe71e1	{"local":[32,1,1],"shared":384}
+604a1a41c64f6a62	{"local":[6,6,1],"shared":0}
+63e9c58765eb99d2	{"local":[32,1,1],"shared":0}
+66f848f2d65e7b75	{"local":[64,1,1],"shared":0}
+690fddd101121e60	{"local":[32,1,1],"shared":0}
+69c6375ee0f50300	{"local":[32,1,1],"shared":0}
+6a993c93a57be70c	{"local":[32,1,1],"shared":384}
+6b355f0751e5f5c0	{"local":[32,1,1],"shared":0}
+6c9250dba05fc17d	{"local":[32,1,1],"shared":0}
+6d0c60eb1ac58fb9	{"local":[32,1,1],"shared":0}
+6e2e368ca3055e09	{"local":[20,11,1],"shared":0}
+7001c2db8496d1d5	{"local":[32,1,1],"shared":0}
+715d98f08d407860	{"local":[32,1,1],"shared":384}
+717affd4d81b976b	{"local":[32,1,1],"shared":384}
+72e58ba8b53dedcc	{"local":[32,1,1],"shared":0}
+7bbbacf8d69d744f	{"local":[32,1,1],"shared":0}
+7e29de1ee51ed275	{"local":[32,1,1],"shared":0}
+7e476096dcb6a260	{"local":[32,1,1],"shared":0}
+7ec5c5a81d458277	{"local":[32,1,1],"shared":0}
+7f9c09aea5a37d39	{"local":[32,1,1],"shared":0}
+82657d6bc0a65fa7	{"local":[32,1,1],"shared":0}
+840869974949b776	{"local":[12,12,1],"shared":0}
+86400095f8465eca	{"local":[32,1,1],"shared":0}
+883ba381ba27d247	{"local":[32,1,1],"shared":0}
+89f054c0e666ae69	{"local":[16,16,1],"shared":0}
+8a1ddd6b584a0b1b	{"local":[32,1,1],"shared":0}
+8a29154666f4c9a4	{"local":[32,1,1],"shared":0}
+9232f37154ba9e01	{"local":[32,1,1],"shared":0}
+93a98a79aed93292	{"local":[32,1,1],"shared":0}
+98f4158d81cc4c63	{"local":[32,1,1],"shared":0}
+9960aba5e1bfd8a1	{"local":[32,1,1],"shared":0}
+9a213b7e97710e13	{"local":[32,1,1],"shared":0}
+9ecba00a6e7deec9	{"local":[32,1,1],"shared":0}
+a06736f732fc7386	{"local":[8,8,1],"shared":0}
+a152bbe4e00b96b9	{"local":[32,1,1],"shared":0}
+a80427d2270cc171	{"local":[32,1,1],"shared":0}
+a8046980651354a6	{"local":[32,1,1],"shared":0}
+a8e5dfd9e1c6ee33	{"local":[32,1,1],"shared":0}
+a9efc7174883d4b1	{"local":[8,8,1],"shared":0}
+ab6a2fe9adf97efb	{"local":[32,1,1],"shared":0}
+ac1a16e353f36e24	{"local":[16,16,1],"shared":0}
+ac41698c630fce0e	{"local":[32,1,1],"shared":0}
+af5dadc74d0425c8	{"local":[32,1,1],"shared":0}
+afe68d67c2e17e22	{"local":[32,1,1],"shared":0}
+b3c9ff8edf51b32d	{"local":[32,1,1],"shared":0}
+b4d4e76dbaf27045	{"local":[32,1,1],"shared":640}
+b5de63d9665e8723	{"local":[32,1,1],"shared":0}
+b767103838f729ca	{"fout":{"1":"uint"}}
+bd5a246169d14959	{"local":[16,16,1],"shared":0}
+bf6a6abb5ecbecd6	{"local":[4,4,1],"shared":0}
+bfce8d4bf052df42	{"local":[12,12,1],"shared":0}
+bff4efb12891c4cb	{"local":[32,1,1],"shared":0}
+c2122510ed611ddc	{"local":[32,1,1],"shared":0}
+c2264a55bdbb651c	{"local":[32,1,1],"shared":0}
+c509e77bef5b4249	{"fout":{"1":"uint"}}
+c50b95b6b919b9f5	{"local":[32,1,1],"shared":0}
+c5c47db15a1dc14d	{"local":[32,1,1],"shared":0}
+c6b958e8cd7b68cc	{"local":[32,1,1],"shared":0}
+c6f377dd7a254b11	{"local":[32,1,1],"shared":0}
+c7e61caa514480d9	{"local":[16,16,1],"shared":0}
+cbd7dc6a9b493e91	{"local":[32,1,1],"shared":0}
+cea5256400b7246a	{"local":[32,1,1],"shared":0}
+d2514cbc70a4b18e	{"local":[32,1,1],"shared":0}
+d2fcbee2f1954836	{"local":[32,1,1],"shared":0}
+d35089e9001ce898	{"local":[32,1,1],"shared":0}
+d35131b20ef086ed	{"local":[32,1,1],"shared":0}
+d910f652494e5739	{"local":[32,1,1],"shared":0}
+db1b8b7800e55a67	{"local":[32,1,1],"shared":384}
+dcb0faa6bdec563b	{"local":[32,1,1],"shared":0}
+df4d3d5093a33caf	{"local":[32,1,1],"shared":0}
+e0bd2ce131c0d2fe	{"local":[32,1,1],"shared":0}
+e0f9eaa2d641616d	{"local":[32,1,1],"shared":0}
+e3e3108f7ae2593a	{"local":[32,1,1],"shared":0}
+e5167b0786ef8222	{"local":[32,1,1],"shared":0}
+e7dd2a9e45c705b9	{"local":[64,1,1],"shared":0}
+eb87cb5b08142632	{"local":[4,4,1],"shared":0}
+ec5ebd6d9902a9a1	{"local":[64,1,1],"shared":0}
+ee881dd573510cc2	{"local":[32,1,1],"shared":0}
+eec04b88ce82300b	{"local":[32,1,1],"shared":0}
+ef13eeecadb58c03	{"local":[32,1,1],"shared":0}
+ef59c644743cf9ff	{"local":[32,1,1],"shared":384}
+eff7ab1c363df916	{"local":[32,1,1],"shared":0}
+f00b4819a4bfaf8a	{"local":[32,1,1],"shared":0}
+f17e5c59e4e53861	{"local":[1,1,1],"shared":0}
+f52da9c7b782e6a0	{"local":[32,1,1],"shared":0}
+f6d10f02cd9bf411	{"local":[32,1,1],"shared":0}
+f7606c9e438f8163	{"local":[32,1,1],"shared":0}
+facf01381f79bb7a	{"local":[32,1,1],"shared":0}
+fc1fdf2576a1c534	{"local":[6,6,1],"shared":0}
+fd3d94ca24af15d1	{"local":[12,12,1],"shared":0}

dm/games/0100F43008C44000/shader_facts_manual.tsv

@@ -0,0 +1,7 @@
+# Hand-written translation facts, merged over the generated rows of shader_facts.tsv by dm/gfx/tools/shader_split.py: key <TAB> facts JSON (same format)
+# This game has no PC version, so the translator has no fragment output types and declares every output float; Metal refuses a pipeline whose float output meets an integer attachment
+# (the host's red line then prints the row to add here). Rows are added for pipelines Metal refused in a run, not by guessing from the code: a value built with integer operations can as well be
+# raw bits stored into a float target (c85de3713170e519 ORs bit fields into output 0 and has not been seen with an integer attachment).
+# c509e77bef5b4249 and b767103838f729ca (env_ikkaku system shaders, variations 0 and 3 of one BNSH in the imgui.trsha pack, pixel): output 1 is an integer ID (255, or one taken from the game's uniforms) into an R8_UI target; refused on the character creation screen
+b767103838f729ca	{"fout":{"1":"uint"}}
+c509e77bef5b4249	{"fout":{"1":"uint"}}

dm/games/0100F43008C44000/shader_map.tsv.conflicts.tsv

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

dm/games/README.md

@@ -41,6 +41,7 @@ None of them has a default game: without one they stop and list the profiles. To
 | `mem_reserve`, `mem_alias`, `mem_stack`, `app_pool` | Guest address-space partition and application pool (`hle_svc.c`) |
 | `mem_used` | Memory `GetInfo(UsedMemorySize)` reports beside the heap (default `0x4000000`, 64 MB). `SetHeapSize` checks a request against the real size of the loaded modules, so a game that sizes its heap as total − used and has more than 64 MB of modules needs at least their size here. Hyrule Warriors: Age of Calamity: 66.6 MB of modules, so the default made `SetHeapSize(pool − 64 MB)` fail and `nnMain` return at once; it sets 72 MB |
 | `game_args` | Arguments the game gets when the command line gives none (the iOS app, the NSP script's Mac runs), split at spaces; `dmhost <dir> -- …` overrides them. The SDK hands them to the game before `nninitStartup`, as on hardware (`main.c`). Octopath Traveler 0 caps the address region Unreal's `nninitStartup` reserves (62 GB by default) with `-SwitchVirtualAddressSpaceSize=2048` (MB) |
+| `mem_map_limit` | `on` / `off` (the default): `MapPhysicalMemory` fails with ResultLimitReached when it would take the process past the application pool (the pool less the used memory `GetInfo` reports), as on Horizon (`hle_svc.c`). Pokémon Legends: Z-A allocates and frees everything it can at startup and the SDK keeps the freed pages mapped until a map fails, then unmaps them and retries; without the limit they stay mapped and the game maps past the pool. Off by default: Ace Attorney Investigations Collection maps a little more than the 1 GB pool its profile reports |
 | `present_interval` | `game` (the default): frames are paced by the present interval the game sets, as on the Switch (a frame presented with interval 2 stays on screen for two display ticks; a game that keeps asking for 2 for a second gets the 30 fps target, and on iOS the display link then runs at 30 Hz; Mario 1.2.1 switches to 2 for single frames after a slow one; Engage asks for 2 on the Switch, and a rewriter cut on its 2.0.0 `main` makes it ask for 1, see `0100A6301214E000/game.tsv`). `display`: one frame per display tick whatever the game asks, which is only right for games whose logic does not count frames (D2R's runs on its own 25 Hz clock). `DM_PRESENT_INTERVAL=game` / `display` overrides it, also from `Documents/env.txt` |
 | `fps_target` | Frame rate cap of the rewriter's frame step, for builds that have one (`patch_nso.py` `FRAME_STEP`: Tears of the Kingdom 1.4.3, Breath of the Wild 1.9.0), or the engine's own frame limit the host writes into game memory (`hle_fpslock.c`: Monster Hunter Rise 16.0.2, which also needs `present_interval display`; Octopath Traveler 0 1.0.5, Unreal's `t.MaxFPS` and `rhi.SyncInterval`, so the game asks for present interval 1 itself). `0` (the default): the game's own step. `30` / `60`: the host decides each frame's step in whole vsyncs and keeps the frame on screen that long (`hle_nvn.c` `fstep_tick`, `nvn_fstep.h`): 1 at a 60 cap, 2 when 60 does not hold or at a 30 cap (and while a prerendered movie plays), plus one vsync while the game is behind real time after a late frame. Presents follow that hold instead of the interval the game sets. `DM_FPS_TARGET` overrides it, also from `Documents/env.txt`; the log line `frame step: game time … of real time` checks the game's speed |
 | `core_preempt` | `on` / `off` (the default): approximate Horizon's per-core preemption. When an SVC wakes a higher-priority thread pinned to the caller's core, the caller waits (at most 5 ms) until such threads block again (`hle_thread.c`). Metroid Dread's movie player relies on it; Mario's loading gets slower with it (first frame 0.6 → 2.9 s), so it is off unless a game needs it. `DM_PREEMPT=0` / `1` overrides |
@@ -64,6 +65,7 @@ None of them has a default game: without one they stop and list the profiles. To
 | `shader_main_glslc` | `on` / `off` (the default): `shader_split.py` also reads GLSLC compiler outputs (magic `0x19866891`) in `main`. 13 Sentinels keeps all its shaders there; Engage's `main` has 11 that it does not use, so the key stays off for other games |
 | `shader_separable` | `on` / `off` (the default): the game builds single-stage programs and combines them per draw, so the real vertex / pixel pairs exist only at run time. `shader_split.py` then pairs every vertex shader with a cover of pixel shaders that together read every input any pixel shader reads, and `dmsh_build.py`'s `vs_pad` makes every vertex shader link with any pixel shader (13 Sentinels) |
 | `shader_vk_capture` | `on` / `off` (the default), with `graphics vulkan`: the game generates its SPIR-V shaders at run time (Super Mario Galaxy: Hagi's GX shaders, 3,089 modules created from its pipeline database while it shows "Initializing system..."), so no game file holds them. They are captured on the Mac with `DM_VK_CAPTURE=<dir>` (`dm/gfx/vk/hle_vk.m`; the NSP script does it itself, booting the game in the Mac host until every key the profile's `vk_capture.tsv` lists is there), and `vk_shaders.py` translates the capture folder with the RomFS |
+| `shader_oodle` | `on` / `off` (the default): the game's archives are Oodle-compressed (Pokémon Legends: Z-A's Trinity packs, `arc/data.trpfs`), so `shader_split.py` decodes them with the game's own `OodleLZ_Decompress`, run by the Mac host in its tool mode (`dm/host/hle_oodle.c`, `DM_OODLE_SERVE`, the decoder's address per `main` build). The split needs `DM_OODLE_HOST` (the Mac host) and `DM_OODLE_MODULES` (the rewritten modules); the NSP script builds the host and sets both when the split runs. `profile_guess.py` sets the key when the RomFS has `ONEPACK` files |
 | `shader_glslc_capture` | `on` / `off` (the default): the game compiles some shaders at run time with its own GLSLC (Fire Emblem: Three Houses: 9 built-in KTGL shaders at boot), so no game file holds them. They are captured on the Mac with `DM_GLSLC_CAPTURE=<dir>` (`dm/host/hle_glslc.c`; the NSP script does this itself, booting the game in the Mac host until every capture file the profile's `shader_map.tsv` names is there), and `shader_split.py` reads `<version dir>/glslc_capture/*.glslc` (beside that version's `RomFS`) or `$DM_GLSLC_CAPTURE`; without the folder those shaders get no translation and the tables differ from the profile's |
 | `shader_ue_decode` | `on` / `off` (the default): an Unreal Engine game keeps its shaders in the shader code archives of its paks (`*.ushaderbytecode`), each shader Oodle-compressed, which only the game's own Oodle decodes. The NSP script (`decode_ue_shaders`) takes the archives out of the paks with `dm/gfx/tools/uepak.py` (the pak index is AES-encrypted with a key it finds in `main`), boots the game in the Mac host with `DM_OODLE_JOB` (`dm/host/hle_oodle.c`: the game's `OodleLZ_Decompress`, found as the one function that references the library's own name string, decodes every shader right before `nnMain`, then the host exits) and writes the decoded archives to `<version dir>/ue_shaders`, where `shader_split.py` reads them (`DM_UE_SHADERS` overrides). Octopath Traveler 0: 10,596 shaders in 2 s. Its vertex and pixel shaders combine per draw, so the profile also sets `shader_separable` |
 

dm/games/game.py

@@ -118,6 +118,8 @@ def header(g):
     if tds not in ('on', 'off'): sys.exit(f'{g.id}: tex_dyn_slots must be on or off: {tds}')
     vcr = g.get('vconv_range', 'off')   # vertex attributes converted on the CPU only over the vertices each draw reads (nvn_mtl.m vc_draw_range): on / off
     if vcr not in ('on', 'off'): sys.exit(f'{g.id}: vconv_range must be on or off: {vcr}')
+    mml = g.get('mem_map_limit', 'off')   # MapPhysicalMemory refused beyond the application pool, as on Horizon (hle_svc.c): on / off
+    if mml not in ('on', 'off'): sys.exit(f'{g.id}: mem_map_limit must be on or off: {mml}')
     gts = g.get('gpu_timestamps', 'submission')   # counter timestamps (nvn_mtl.m rep_ticks): submission = one time per submission, segment = the GPU start / end of the segment holding the report
     if gts not in ('submission', 'segment'): sys.exit(f'{g.id}: gpu_timestamps must be submission or segment: {gts}')
     om = g.get('op_mode', 'console')   # the operation mode reported at start (hle_oe.c dm_opmode; DM_OPMODE overrides): console = docked, handheld = touch screen in use (Super Mario Galaxy's pointer)
@@ -141,6 +143,7 @@ def header(g):
          f'#define DM_GAME_APP_POOL {g.int("app_pool"):#x}ull',
          f'#define DM_GAME_MEM_USED {int(g.get("mem_used", "0x4000000"), 0):#x}ull',   # used memory reported beside the heap (hle_svc.c GetInfo UsedMemorySize)
          f'#define DM_GAME_ARGS {q(g.get("game_args", ""))}',   # game arguments when the command line gives none (main.c; Octopath Traveler 0: Unreal's -SwitchVirtualAddressSpaceSize=)
+         f'#define DM_GAME_MEM_MAP_LIMIT {int(mml == "on")}',
          f'#define DM_GAME_NVN_SSBO_BASE(stage) ({sb})',
          f'#define DM_GAME_NVN_TEX_BASE {g.int("nvn_tex_base"):#x}',
          f'#define DM_GAME_NVN_TEX_SLOTS {g.int("nvn_tex_slots")}',

dm/games/profile_guess.py

@@ -17,12 +17,12 @@ sys.path.insert(0, str(HERE.parent / 'tools')); import nso_syms   # dynamic symb
 MAGICS = [(b'UnityFS\0', 'UnityFS'), (b'UnityWeb', 'UnityWeb'), (b'UnityRaw', 'UnityRaw'), (b'BNSH', 'BNSH'), (b'BAHS', 'BAHS'), (b'SARC', 'SARC'), (b'Yaz0', 'Yaz0'),
           (bytes.fromhex('28b52ffd'), 'zstd'), (b'FRES', 'bfres'), (b'BNTX', 'BNTX'), (b'BWAV', 'BWAV'), (b'FSAR', 'BFSAR'), (b'CPK ', 'CRI CPK'), (b'CRID', 'CRI USM'),
           (b'@UTF', 'CRI UTF'), (b'AKPK', 'Wwise pck'), (b'BKHD', 'Wwise bnk'), (b'FSB5', 'FMOD'), (b'\x89PNG', 'PNG'), (b'RIFF', 'RIFF'), (b'OggS', 'Ogg'), (b'PK\x03\x04', 'zip'), (b'NSO0', 'NSO'),
-          (bytes.fromhex('03022307'), 'SPIR-V')]
+          (bytes.fromhex('03022307'), 'SPIR-V'), (b'ONEPACK\0', 'Trinity pack')]
-PACKED = {'UnityFS', 'UnityWeb', 'Yaz0', 'zstd', 'BWAV', 'CRI CPK', 'CRI USM', 'Wwise pck', 'FMOD', 'PNG', 'Ogg', 'zip'}   # file headers of compressed / media formats (zip_store)
+PACKED = {'UnityFS', 'UnityWeb', 'Yaz0', 'zstd', 'BWAV', 'CRI CPK', 'CRI USM', 'Wwise pck', 'FMOD', 'PNG', 'Ogg', 'zip', 'Trinity pack'}   # file headers of compressed / media formats (zip_store)
 # IL2CPP by the metadata file every IL2CPP player opens: Unity 2022's main (Ace Attorney Investigations Collection) has no il2cpp_ API names, Engage's 2020.3 has both
 MARKERS = [(b'nvnBootstrapLoader', 'NVN'), (b'vkCreateInstance', 'Vulkan'), (b'eglGetProcAddress', 'EGL / OpenGL'), (b'global-metadata.dat', 'Unity IL2CPP'), (b'UnityEngine', 'Unity'),
            (b'sead::', 'Nintendo sead'), (b'agl::', 'Nintendo agl'), (b'gsys::', 'Nintendo gsys'), (b'Unreal', 'Unreal'), (b'UE4', 'Unreal 4'), (b'Havok', 'Havok'), (b'CRIWARE', 'CRIWARE'), (b'Wwise', 'Wwise')]
-SHADER_CONTAINERS = ('BNSH', 'BAHS', 'UnityFS')   # RomFS files that hold a game's shaders: with them, GLSLC outputs in main are not its shaders (Engage's main has 11 it does not use)
+SHADER_CONTAINERS = ('BNSH', 'BAHS', 'UnityFS', 'Trinity pack')   # RomFS files that hold a game's shaders: with them, GLSLC outputs in main are not its shaders (Engage's main has 11 it does not use)
 GLSLC_MAGIC = b'\x91\x68\x86\x19'   # GLSLC compiler outputs (shader_split.py glslc_shaders)
 APP_POOL = 0xCD500000               # the application memory pool of a 4 GB Switch (3285 MiB), what every profile but Engage's reports
 MiB = 1 << 20
@@ -144,6 +144,8 @@ def guess(F, hooks=None):
         why = f'main holds {F["glslc"]} GLSLC outputs and the RomFS has no {" / ".join(SHADER_CONTAINERS)}: the shaders are compiled into main (13 Sentinels, the KTGL games)'
         G['shader_main_glslc'] = ('on', why)
         G['shader_separable'] = ('on', 'with shader_main_glslc: single-stage programs combined per draw (every GLSLC output in main seen so far is one stage)')
+    if R['mags']['Trinity pack']:
+        G['shader_oodle'] = ('on', 'the RomFS has Trinity packs (ONEPACK, arc/data.trpfs: Pokémon Legends: Z-A), Oodle-compressed: the shader split decodes them with the game\'s own decoder in the Mac host (dm/host/hle_oodle.c, its address per build)')
     glslc_mod = [m for m in keep if 'glslcCompile' in mods[m]['exports']]
     if F['glsl_src'] and 'glslcCompile' in imp and glslc_mod:
         G['shader_glslc_capture'] = ('on', f'main holds {F["glsl_src"]} GLSL sources and calls glslcCompile in {glslc_mod[0]}: shaders compiled at run time (Fire Emblem: Three Houses), captured by booting it on the Mac')
@@ -226,7 +228,7 @@ def heap_signals(log, g):
 
 
 # ---- Checking the rules against the hand-made profiles ----
-DEFAULTS = {'graphics': 'nvn', 'shader_vk_capture': 'off', 'shader_main_glslc': 'off', 'shader_separable': 'off', 'shader_glslc_capture': 'off', 'shader_ue_decode': 'off',
+DEFAULTS = {'graphics': 'nvn', 'shader_vk_capture': 'off', 'shader_main_glslc': 'off', 'shader_separable': 'off', 'shader_glslc_capture': 'off', 'shader_ue_decode': 'off', 'shader_oodle': 'off',
             'op_mode': 'console', 'pad_layout': 'simple', 'nvn_tex_base': '0x20', 'nvn_tex_slots': '32'}
 NO_SIGNAL = {'pad_layout', 'op_mode'}   # no static signal outside the engine presets; the player switches the pad layout in game, DM_OPMODE in env.txt sets the mode
 

dm/gfx/nvn/hle_nvn.c

@@ -717,7 +717,7 @@ NVN(nvnCommandBufferSignalEvent, void, NVNcommandBuffer *c, const NVNevent *e, N
     const ev_t *ev = (const ev_t *)e; volatile uint32_t *val = ev->val; uint32_t mm = m;
     memcpy(p + 8, &val, 8); memcpy(p + 16, &mm, 4); memcpy(p + 20, &v, 4);
 }
-// A fence recorded in a command buffer: the replay fences the sync with the submission that reaches this point (OP_FENCE_SYNC). Unreal Engine's NVN RHI fences its command buffers this way (Octopath Traveler 0)
+// A fence recorded in a command buffer: the replay fences the sync with the submission that reaches this point (OP_FENCE_SYNC). Unreal Engine's NVN RHI fences its command buffers this way (Octopath Traveler 0); Pokémon Legends: Z-A fences once a frame this way and waits on it
 NVN(nvnCommandBufferFenceSync, void, NVNcommandBuffer *c, NVNsync *s, NVNsyncCondition cond, int flags) {
     (void)cond; (void)flags;
     uint8_t *p = cmd_alloc(CB(c), 16); put_hdr(p, OP_FENCE_SYNC, 16); memcpy(p + 8, &s, 8);
@@ -1731,8 +1731,8 @@ NVN(nvnMemoryPoolMapVirtual, NVNboolean, NVNmemoryPool *p, int n, const NVNmappi
     return 1;
 }
 
-// Samplers: reduction (min / max) does not exist on Metal samplers, logged once; the getters read the parameters in the object
+// Samplers: the getters read the parameters in the object. Reduction (min / max of the texels a filtered sample covers) maps to the Metal 26 sampler reduction mode (nvn_mtl.m samp_by_id)
-NVN(nvnSamplerBuilderSetReductionFilter, void, NVNsamplerBuilder *b, NVNsamplerReduction r) { (void)b; if (r) ONCE_LOG("SamplerBuilderSetReductionFilter %d: the host samples with averaging (min/max reduction unsupported)", (int)r); }
+NVN(nvnSamplerBuilderSetReductionFilter, void, NVNsamplerBuilder *b, NVNsamplerReduction r) { SB(b)->reduction = (uint8_t)r; }
 NVN(nvnSamplerSetDebugLabel, void, NVNsampler *s, const char *l) { (void)s; (void)l; }
 #define SP(s) (&((const samp_t *)(s))->p)
 NVN(nvnSamplerGetMinMagFilter, void, const NVNsampler *s, NVNminFilter *mi, NVNmagFilter *ma) { *mi = (NVNminFilter)SP(s)->minf; *ma = (NVNmagFilter)SP(s)->magf; }
@@ -1742,11 +1742,11 @@ NVN(nvnSamplerGetLodBias, float, const NVNsampler *s) { return SP(s)->lod_bias;
 NVN(nvnSamplerGetCompare, void, const NVNsampler *s, NVNcompareMode *m, NVNcompareFunc *f) { *m = (NVNcompareMode)SP(s)->cmp_mode; *f = (NVNcompareFunc)SP(s)->cmp_func; }
 NVN(nvnSamplerGetBorderColor, void, const NVNsampler *s, float *c) { memcpy(c, SP(s)->border, 16); }
 NVN(nvnSamplerGetMaxAnisotropy, float, const NVNsampler *s) { return SP(s)->aniso; }
-NVN(nvnSamplerGetReductionFilter, NVNsamplerReduction, const NVNsampler *s) { (void)s; return (NVNsamplerReduction)0; }
+NVN(nvnSamplerGetReductionFilter, NVNsamplerReduction, const NVNsampler *s) { return (NVNsamplerReduction)SP(s)->reduction; }
 // whether two samplers have the same state (games use it to deduplicate sampler pool entries)
 NVN(nvnSamplerCompare, NVNboolean, const NVNsampler *a, const NVNsampler *b) {
     const samp_par *x = SP(a), *y = SP(b);
-    return x->minf == y->minf && x->magf == y->magf && x->ws == y->ws && x->wt == y->wt && x->wr == y->wr && x->cmp_mode == y->cmp_mode && x->cmp_func == y->cmp_func &&
+    return x->minf == y->minf && x->magf == y->magf && x->ws == y->ws && x->wt == y->wt && x->wr == y->wr && x->cmp_mode == y->cmp_mode && x->cmp_func == y->cmp_func && x->reduction == y->reduction &&
            x->lod_min == y->lod_min && x->lod_max == y->lod_max && x->lod_bias == y->lod_bias && x->aniso == y->aniso && !memcmp(x->border, y->border, sizeof x->border);
 }
 
@@ -1773,7 +1773,7 @@ NVN(nvnTextureGetTarget, NVNtextureTarget, const NVNtexture *t) { return (NVNtex
 NVN(nvnTextureSetDebugLabel, void, NVNtexture *t, const char *s) { (void)t; (void)s; }
 NVN(nvnTextureViewGetLevels, NVNboolean, const NVNtextureView *v, int *base, int *n) { const view_t *w = (const view_t *)v; *base = w->base_level; *n = w->nlevels; return w->set_levels; }
 NVN(nvnTextureViewGetTarget, NVNboolean, const NVNtextureView *v, NVNtextureTarget *t) { const view_t *w = (const view_t *)v; *t = (NVNtextureTarget)w->target; return w->set_target; }
-NVN(nvnTextureViewSetDepthStencilMode, void, NVNtextureView *v, NVNtextureDepthStencilMode m) { VW(v)->set_dsmode = 1; VW(v)->dsmode = (uint8_t)m; if (m) ONCE_LOG("TextureViewSetDepthStencilMode %d: the host reads depth", (int)m); }
+NVN(nvnTextureViewSetDepthStencilMode, void, NVNtextureView *v, NVNtextureDepthStencilMode m) { VW(v)->set_dsmode = 1; VW(v)->dsmode = (uint8_t)m; }   // STENCIL: the replay binds a stencil-only view (tex_by_id)
 NVN(nvnTextureViewSetSwizzle, void, NVNtextureView *v, NVNtextureSwizzle r, NVNtextureSwizzle g, NVNtextureSwizzle bl, NVNtextureSwizzle a) {
     view_t *w = VW(v); w->set_swizzle = 1; w->swz[0] = (uint8_t)r; w->swz[1] = (uint8_t)g; w->swz[2] = (uint8_t)bl; w->swz[3] = (uint8_t)a;
     if (r != NVN_TEXTURE_SWIZZLE_R || g != NVN_TEXTURE_SWIZZLE_G || bl != NVN_TEXTURE_SWIZZLE_B || a != NVN_TEXTURE_SWIZZLE_A) ONCE_LOG("TextureViewSetSwizzle not identity (%d %d %d %d): remapped with a Metal view when sampling", (int)r, (int)g, (int)bl, (int)a);
@@ -1782,6 +1782,85 @@ NVN(nvnTextureViewSetSwizzle, void, NVNtextureView *v, NVNtextureSwizzle r, NVNt
 // Color state: alpha test function, logic op (the fields already exist in color_st)
 NVN(nvnColorStateSetAlphaTest, void, NVNcolorState *s, NVNalphaFunc f) { ST(color_st, s)->alpha_func = (uint32_t)f; }
 NVN(nvnColorStateSetLogicOp, void, NVNcolorState *s, NVNlogicOp op) { ST(color_st, s)->logic_op = (uint32_t)op; }
+// Getters and the rest of the multisample state: Pokémon Legends: Z-A (Trinity, SDK 19) reads its state objects and builders back when it builds its pipeline states; the stub returned
+// whatever was left in x0 (the object's address), so a blend enable or depth test read back as true, a texture builder's width as an address
+NVN(nvnColorStateGetBlendEnable, NVNboolean, const NVNcolorState *s, int i) { return i >= 0 && i < 8 && (((const color_st *)s)->blend_en >> i & 1); }
+NVN(nvnColorStateGetAlphaTest, NVNalphaFunc, const NVNcolorState *s) { return (NVNalphaFunc)((const color_st *)s)->alpha_func; }
+NVN(nvnColorStateGetLogicOp, NVNlogicOp, const NVNcolorState *s) { return (NVNlogicOp)((const color_st *)s)->logic_op; }
+NVN(nvnChannelMaskStateGetChannelMask, void, const NVNchannelMaskState *s, int i, NVNboolean *r, NVNboolean *g, NVNboolean *b, NVNboolean *a) {
+    uint32_t m = i >= 0 && i < 8 ? ((const chmask_st *)s)->mask >> (4 * i) & 0xf : 0;
+    if (r) *r = (m & 1) != 0; if (g) *g = (m & 2) != 0; if (b) *b = (m & 4) != 0; if (a) *a = (m & 8) != 0;
+}
+NVN(nvnBlendStateGetBlendTarget, int, const NVNblendState *s) { return (int)((const blend_st *)s)->target; }
+NVN(nvnBlendStateGetBlendFunc, void, const NVNblendState *s, NVNblendFunc *sr, NVNblendFunc *dr, NVNblendFunc *sa, NVNblendFunc *da) {
+    const blend_st *o = (const blend_st *)s;
+    if (sr) *sr = (NVNblendFunc)o->src_rgb; if (dr) *dr = (NVNblendFunc)o->dst_rgb; if (sa) *sa = (NVNblendFunc)o->src_a; if (da) *da = (NVNblendFunc)o->dst_a;
+}
+NVN(nvnBlendStateGetBlendEquation, void, const NVNblendState *s, NVNblendEquation *er, NVNblendEquation *ea) {
+    const blend_st *o = (const blend_st *)s; if (er) *er = (NVNblendEquation)o->eq_rgb; if (ea) *ea = (NVNblendEquation)o->eq_a;
+}
+NVN(nvnPolygonStateGetCullFace, NVNface, const NVNpolygonState *s) { return (NVNface)((const poly_st *)s)->cull; }
+NVN(nvnPolygonStateGetFrontFace, NVNfrontFace, const NVNpolygonState *s) { return (NVNfrontFace)((const poly_st *)s)->front; }
+NVN(nvnPolygonStateGetPolygonMode, NVNpolygonMode, const NVNpolygonState *s) { return (NVNpolygonMode)((const poly_st *)s)->mode; }
+NVN(nvnPolygonStateGetPolygonOffsetEnables, NVNpolygonOffsetEnable, const NVNpolygonState *s) { return (NVNpolygonOffsetEnable)((const poly_st *)s)->offset_en; }
+NVN(nvnDepthStencilStateGetDepthTestEnable, NVNboolean, const NVNdepthStencilState *s) { return ((const ds_st *)s)->dt; }
+NVN(nvnDepthStencilStateGetDepthWriteEnable, NVNboolean, const NVNdepthStencilState *s) { return ((const ds_st *)s)->dw; }
+NVN(nvnDepthStencilStateGetDepthFunc, NVNdepthFunc, const NVNdepthStencilState *s) { return (NVNdepthFunc)((const ds_st *)s)->df; }
+NVN(nvnDepthStencilStateGetStencilTestEnable, NVNboolean, const NVNdepthStencilState *s) { return ((const ds_st *)s)->st; }
+NVN(nvnDepthStencilStateGetStencilFunc, NVNstencilFunc, const NVNdepthStencilState *s, NVNface f) { const ds_st *o = (const ds_st *)s; return (NVNstencilFunc)(f == NVN_FACE_BACK ? o->sf_b : o->sf_f); }
+NVN(nvnDepthStencilStateGetStencilOp, void, const NVNdepthStencilState *s, NVNface f, NVNstencilOp *fail, NVNstencilOp *zfail, NVNstencilOp *zpass) {
+    const ds_st *o = (const ds_st *)s; int b = f == NVN_FACE_BACK;
+    if (fail) *fail = (NVNstencilOp)(b ? o->fail_b : o->fail_f); if (zfail) *zfail = (NVNstencilOp)(b ? o->zfail_b : o->zfail_f); if (zpass) *zpass = (NVNstencilOp)(b ? o->zpass_b : o->zpass_f);
+}
+// Multisample state: recorded only (the replay uses the sample count, multisample enable and alpha to coverage)
+NVN(nvnMultisampleStateSetAlphaToCoverageDither, void, NVNmultisampleState *s, uint32_t e) { ST(ms_st, s)->a2c_dither = B(e) != 0; }
+NVN(nvnMultisampleStateSetRasterSamples, void, NVNmultisampleState *s, int n) { ST(ms_st, s)->raster_samples = (uint8_t)n; }
+NVN(nvnMultisampleStateSetCoverageModulationMode, void, NVNmultisampleState *s, NVNcoverageModulationMode m) { ST(ms_st, s)->cov_mod = (uint8_t)m; }
+NVN(nvnMultisampleStateSetCoverageToColorEnable, void, NVNmultisampleState *s, uint32_t e) { ST(ms_st, s)->cov2color_en = B(e) != 0; }
+NVN(nvnMultisampleStateSetCoverageToColorOutput, void, NVNmultisampleState *s, int i) { ST(ms_st, s)->cov2color_out = (uint8_t)i; }
+NVN(nvnMultisampleStateSetSampleLocationsGridEnable, void, NVNmultisampleState *s, uint32_t e) { ST(ms_st, s)->sl_grid = B(e) != 0; }
+NVN(nvnMultisampleStateGetAlphaToCoverageEnable, NVNboolean, const NVNmultisampleState *s) { return ((const ms_st *)s)->a2c; }
+NVN(nvnMultisampleStateGetAlphaToCoverageDither, NVNboolean, const NVNmultisampleState *s) { return ((const ms_st *)s)->a2c_dither; }
+NVN(nvnMultisampleStateGetRasterSamples, int, NVNmultisampleState *s) { return ((const ms_st *)s)->raster_samples; }
+NVN(nvnMultisampleStateGetCoverageModulationMode, NVNcoverageModulationMode, const NVNmultisampleState *s) { return (NVNcoverageModulationMode)((const ms_st *)s)->cov_mod; }
+NVN(nvnMultisampleStateGetCoverageToColorEnable, NVNboolean, const NVNmultisampleState *s) { return ((const ms_st *)s)->cov2color_en; }
+NVN(nvnMultisampleStateGetCoverageToColorOutput, int, const NVNmultisampleState *s) { return ((const ms_st *)s)->cov2color_out; }
+NVN(nvnMultisampleStateGetSampleLocationsEnable, NVNboolean, const NVNmultisampleState *s) { return ((const ms_st *)s)->sl_en; }
+NVN(nvnMultisampleStateGetSampleLocationsGridEnable, NVNboolean, const NVNmultisampleState *s) { return ((const ms_st *)s)->sl_grid; }
+// Builder and pool queries
+NVN(nvnBufferBuilderGetSize, size_t, const NVNbufferBuilder *b) { return (size_t)((const buf_builder_t *)b)->size; }
+NVN(nvnBufferBuilderGetMemoryOffset, ptrdiff_t, const NVNbufferBuilder *b) { return (ptrdiff_t)((const buf_builder_t *)b)->off; }
+NVN(nvnTextureBuilderGetTarget, NVNtextureTarget, const NVNtextureBuilder *b) { return (NVNtextureTarget)((const tex_builder_t *)b)->d.target; }
+NVN(nvnTextureBuilderGetWidth, int, const NVNtextureBuilder *b) { return (int)((const tex_builder_t *)b)->d.width; }
+NVN(nvnTextureBuilderGetHeight, int, const NVNtextureBuilder *b) { return (int)((const tex_builder_t *)b)->d.height; }
+NVN(nvnTextureBuilderGetDepth, int, const NVNtextureBuilder *b) { return (int)((const tex_builder_t *)b)->d.depth; }
+NVN(nvnTextureBuilderGetLevels, int, const NVNtextureBuilder *b) { return (int)((const tex_builder_t *)b)->d.levels; }
+NVN(nvnTextureBuilderGetFormat, NVNformat, const NVNtextureBuilder *b) { return (NVNformat)((const tex_builder_t *)b)->d.format; }
+NVN(nvnTextureBuilderGetSamples, int, const NVNtextureBuilder *b) { return ((const tex_builder_t *)b)->d.samples; }
+NVN(nvnTextureBuilderGetStride, ptrdiff_t, const NVNtextureBuilder *b) { return (ptrdiff_t)((const tex_builder_t *)b)->d.stride; }
+NVN(nvnTextureViewGetFormat, NVNboolean, const NVNtextureView *v, NVNformat *f) { const view_t *w = (const view_t *)v; if (f) *f = (NVNformat)w->format; return w->set_format; }
+NVN(nvnTextureViewGetDepthStencilMode, NVNboolean, const NVNtextureView *v, NVNtextureDepthStencilMode *m) { const view_t *w = (const view_t *)v; if (m) *m = (NVNtextureDepthStencilMode)w->dsmode; return w->set_dsmode; }
+NVN(nvnTextureViewGetSwizzle, NVNboolean, const NVNtextureView *v, NVNtextureSwizzle *r, NVNtextureSwizzle *g, NVNtextureSwizzle *b, NVNtextureSwizzle *a) {
+    const view_t *w = (const view_t *)v; NVNtextureSwizzle *o[4] = {r, g, b, a};
+    for (int c = 0; c < 4; c++) if (o[c]) *o[c] = (NVNtextureSwizzle)(w->set_swizzle ? w->swz[c] : NVN_TEXTURE_SWIZZLE_R + c);
+    return w->set_swizzle;
+}
+// The host keeps no depth / stencil sampling mode on the texture (TextureBuilderSetDepthStencilMode is only logged): depth, the default
+NVN(nvnTextureGetDepthStencilMode, NVNtextureDepthStencilMode, const NVNtexture *t) { (void)t; return (NVNtextureDepthStencilMode)0; }
+// Queue priority: the host has one Metal queue
+NVN_NOPROTO(nvnQueueBuilderSetQueuePriority, void, NVNqueueBuilder *b, int prio) { (void)b; (void)prio; }
+// Transform feedback (vertex outputs captured into buffers): Pokémon Legends: Z-A skins its characters this way. A vertex-only program writes skinned positions and packed normals / tangents
+// under rasterizer discard, and the character draws read those buffers as vertex streams. Recorded here and run at replay by the shader's main0_tf entry point, a vertex function that stores the
+// captured outputs into the bound buffers (dmsh_build.py vs_tf, the layout from the program's control section; nvn_mtl.m tf_* for the replay)
+NVN(nvnCommandBufferBindTransformFeedbackBuffer, void, NVNcommandBuffer *c, int i, NVNbufferAddress a, size_t n) { *REC(c, OP_TF_BIND, p_tf_bind) = (p_tf_bind){i, 0, a, n}; }
+NVN(nvnCommandBufferBindTransformFeedbackBuffers, void, NVNcommandBuffer *c, int first, int n, const NVNbufferRange *r) {
+    for (int i = 0; i < n; i++) *REC(c, OP_TF_BIND, p_tf_bind) = (p_tf_bind){first + i, 0, r[i].address, r[i].size};
+}
+NVN(nvnCommandBufferBeginTransformFeedback, void, NVNcommandBuffer *c, NVNbufferAddress a) { *REC(c, OP_TF_STATE, p_tf_state) = (p_tf_state){0, 0, a}; }
+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"); }
 
 // 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); }
@@ -1953,7 +2032,8 @@ NVN(nvnCommandBufferSetAlphaRef, void, NVNcommandBuffer *c, float r) { (void)c;
 NVN(nvnCommandBufferSetLineWidth, void, NVNcommandBuffer *c, float w) { (void)c; if (w != 1.0f) ONCE_LOG("SetLineWidth %.2f: Metal has only 1-pixel lines", w); }
 NVN(nvnCommandBufferSetPointSize, void, NVNcommandBuffer *c, float s) { (void)c; (void)s; ONCE_LOG("SetPointSize: the point size comes from the shader"); }
 NVN(nvnCommandBufferSetPrimitiveRestart, void, NVNcommandBuffer *c, uint32_t en, int idx) { (void)c; if ((en & 0xff) && idx != -1 && idx != 0xffff) ONCE_LOG("[RED] SetPrimitiveRestart index %#x: Metal only accepts all ones", idx); }
-NVN(nvnCommandBufferSetRasterizerDiscard, void, NVNcommandBuffer *c, uint32_t en) { (void)c; if (en & 0xff) ONCE_LOG("[RED] SetRasterizerDiscard on: ignored by the host for now"); }
+// Rasterizer discard: draws run their vertex stage only (with transform feedback, the capture); recorded, and at replay such a draw rasterizes nothing (nvn_mtl.m)
+NVN(nvnCommandBufferSetRasterizerDiscard, void, NVNcommandBuffer *c, uint32_t en) { REC(c, OP_RAST_DISCARD, p_rast_discard)->on = (en & 0xff) != 0; }
 NVN(nvnCommandBufferSetSampleMask, void, NVNcommandBuffer *c, int m) { (void)c; if ((m & 0xffff) != 0xffff) ONCE_LOG("SetSampleMask %#x: ignored by the host for now", m); }
 NVN(nvnCommandBufferSetConservativeRasterEnable, void, NVNcommandBuffer *c, uint32_t en) { (void)c; if (en & 0xff) ONCE_LOG("SetConservativeRasterEnable on: ignored by the host for now"); }
 NVN(nvnCommandBufferSetTiledCacheAction, void, NVNcommandBuffer *c, NVNtiledCacheAction a) { (void)c; (void)a; }        // Tegra tiled cache hint, no Metal equivalent

dm/gfx/nvn/nvn_mtl.m

@@ -5,6 +5,7 @@
 #import <ImageIO/ImageIO.h>
 #import <UniformTypeIdentifiers/UniformTypeIdentifiers.h>
 #include <mach/mach_time.h>
+#include <mach/mach.h>
 #include <dlfcn.h>
 #include <sys/stat.h>
 #include <sys/mman.h>
@@ -374,7 +375,22 @@ void mtl_pool_add(const void *mem, uint64_t size) {
     uintptr_t pg = (uintptr_t)getpagesize(), lo = (uintptr_t)mem, hi = lo + size;
     uintptr_t a = lo & ~(pg - 1), b = (hi + pg - 1) & ~(pg - 1);
     id<MTLBuffer> buf = [g_dev newBufferWithBytesNoCopy:(void *)a length:b - a options:MTLResourceStorageModeShared deallocator:nil];
-    if (!buf) MDIE("memory pool %p+%#llx: failed to create the zero-copy buffer", mem, (unsigned long long)size);
+    // Pokémon Legends: Z-A: twice on the Mac a small pool's buffer was refused although its pages were one read-write VM region, both times while another heavy job (a shader compile,
+    // a second game instance) took the machine's memory; every run on an idle machine created all of them. Wiring the pages may fail under memory pressure: retry for up to about 2 s
+    for (int t = 0; !buf && t < 12; t++) {
+        usleep(10000u << (t < 4 ? t : 4));
+        buf = [g_dev newBufferWithBytesNoCopy:(void *)a length:b - a options:MTLResourceStorageModeShared deallocator:nil];
+        if (buf) MLOG("[RED] memory pool %p+%#llx: the zero-copy buffer was refused, created at retry %d", mem, (unsigned long long)size, t + 1);
+    }
+    if (!buf) {   // say what the pages looked like: every VM region over the range, with its protection (an unmapped page in the range is the usual cause)
+        for (vm_address_t r = a; r < b; ) {
+            vm_size_t rs = 0; vm_region_basic_info_data_64_t vi; mach_msg_type_number_t cnt = VM_REGION_BASIC_INFO_COUNT_64; mach_port_t obj = MACH_PORT_NULL;
+            if (vm_region_64(mach_task_self(), &r, &rs, VM_REGION_BASIC_INFO_64, (vm_region_info_t)&vi, &cnt, &obj) != KERN_SUCCESS) { MLOG("[RED]   no VM region at or after %#llx", (unsigned long long)r); break; }
+            MLOG("[RED]   VM region %#llx+%#llx protection %d (max %d)", (unsigned long long)r, (unsigned long long)rs, vi.protection, vi.max_protection);
+            r += rs;
+        }
+        MDIE("memory pool %p+%#llx: failed to create the zero-copy buffer over %#lx..%#lx", mem, (unsigned long long)size, (unsigned long)a, (unsigned long)b);
+    }
     pthread_mutex_lock(&g_mu);
     if (g_npool == g_cap) { g_cap = g_cap ? g_cap * 2 : 64; g_pools = realloc(g_pools, g_cap * sizeof *g_pools); }
     g_pidx_ok = 0; g_plast = -1; atomic_fetch_add(&g_pool_gen, 1);   // the address index and the per-thread caches are invalid (pool_find, pc_get)
@@ -1000,7 +1016,7 @@ void mtl_drain(void) { present_drop_for_suspend(); mtl_present_flush(); mtl_init
 
 // ================= Render replay =================
 #include "../../host/front.h"
-static int g_red_once[100];
+static int g_red_once[105];
 #define RED_ONCE(k, ...) do { if (!g_red_once[k]++) MLOG("[RED] " __VA_ARGS__); } while (0)
 // NVN vertex streams 0–15 → Metal vertex buffer slots: 0–13 go to 16–29 (30 is the constant fallback buffer), 14 and 15 to 14 and 15. Metal has only the 31 slots 0–30; translated shaders use only buffer(0) (the argument buffer),
 // and the host's own vertex functions read no vertex buffers. 13 Sentinels uses stream 15, and the old 16 + stream went out of range (Metal assertion: buffer index (31) must be < 31)
@@ -1020,6 +1036,7 @@ static const pfv_pair g_pfv_srgb[] = {
 };
 static int tex_need_pfv(MTLPixelFormat base, MTLPixelFormat view) {
     if (base == view) return 0;
+    if (base == MTLPixelFormatDepth32Float_Stencil8 && view == MTLPixelFormatX32_Stencil8) return 0;   // the stencil-only view needs no format-view usage (upgrading recreated the depth buffer in use)
     for (size_t i = 0; i < sizeof g_pfv_srgb / sizeof *g_pfv_srgb; i++)
         if ((base == g_pfv_srgb[i].a && view == g_pfv_srgb[i].b) || (base == g_pfv_srgb[i].b && view == g_pfv_srgb[i].a)) return 0;
     return 1;
@@ -1746,7 +1763,8 @@ static void rr_reset(void);
 static id<MTL4RenderCommandEncoder> g_enc;
 static int enc_open(void) { return g_enc != nil; }
 // Target info of the current render pass: determined once from the attachments at rp_begin and shared by all draws in the pass (pipeline key, depth-stencil, viewport / scissor all depend only on it)
-static struct { uint32_t cfmt[8], dfmt, sfmt, samples; int hd, hs; NSUInteger tw, th; uint64_t serial, reg_gen, ser_gen, range_dirty; int range_external, range_depth_nonro; __unsafe_unretained id<MTLTexture> range_color, range_depth; } g_pass;
+static struct { uint32_t cfmt[8], dfmt, sfmt, samples; int hd, hs; NSUInteger tw, th; uint64_t serial, reg_gen, ser_gen, range_dirty; int range_external, range_depth_nonro, layered; __unsafe_unretained id<MTLTexture> range_color, range_depth; } g_pass;
+static int g_rp_layered;   // the pass about to open is for a vertex shader that writes the layer (draw → rp_begin_ex)
 static id<MTLRenderPipelineState> g_enc_pso; static id<MTLDepthStencilState> g_enc_dss;   // Already set on this encoder (not set again if identical)
 // Dynamic state already sent on the open render encoder: each draw used to send all nine (stencil reference, viewport, scissor, cull, winding, fill, depth bias, depth clip, blend colour);
 // now a value is sent only when it differs. g_ed.ok holds one bit per item and is cleared for every new encoder and after a clear triangle (which sets its own)
@@ -2296,6 +2314,18 @@ static id<MTL4RenderCommandEncoder> rp_begin_ex_(int clear_only) {
         pd.visibilityResultBuffer = g_vis;
     }
     g_enc_vis = 0; g_enc_has_vis = g_vis_on;
+    // Layered rendering: a vertex shader that writes the layer selects the attachment layer per primitive only when the pass has renderTargetArrayLength; without it every layer went to the
+    // attached one. Pokémon Legends: Z-A builds its 32x32x32 colour grading LUT this way (b29f08619f379a83, one instance per depth plane) and tone-maps through it: only plane 0 was written and the
+    // whole picture came out black. The pass gets every layer the attachments have from their attached one on (3D: depth planes); the pass is split when the next draw's shader differs (draw)
+    g_pass.layered = g_rp_layered;   // what the pass was opened for, even when the attachments have one layer (then the draws need no split)
+    if (g_rp_layered) {
+        NSUInteger nl = NSUIntegerMax;
+        for (int i = 0; i < 8; i++) { MTLRenderPassColorAttachmentDescriptor *c = pd.colorAttachments[i]; id<MTLTexture> t = c.texture; if (!t) continue;
+            NSUInteger have = t.textureType == MTLTextureType3D ? MAX(t.depth >> c.level, 1u) - c.depthPlane : t.arrayLength * (t.textureType == MTLTextureTypeCube || t.textureType == MTLTextureTypeCubeArray ? 6 : 1) - c.slice;
+            if (have < nl) nl = have; }
+        if (pd.depthAttachment.texture) { id<MTLTexture> t = pd.depthAttachment.texture; NSUInteger have = t.arrayLength * (t.textureType == MTLTextureTypeCube || t.textureType == MTLTextureTypeCubeArray ? 6 : 1) - pd.depthAttachment.slice; if (have < nl) nl = have; }
+        if (nl != NSUIntegerMax && nl > 1) { pd.renderTargetArrayLength = nl; PD("L %llu layered pass: %lu layers", g_stat_pass + 1, (unsigned long)nl); }
+    }
     pg_render(pd, (NSUInteger)R.rt.n);
     { void td_note(id pd, NSUInteger nc); td_note(pd, (NSUInteger)R.rt.n);
       void ts_note(id pd, NSUInteger nc); ts_note(pd, (NSUInteger)R.rt.n);
@@ -2393,9 +2423,9 @@ static void shmap_load(void) {
     warm_start();
 }
 
-@interface NvShader : NSObject { @public id<MTLFunction> fn; nvn_refl rf; id<MTLArgumentEncoder> ae; NSUInteger arg_len; int stage, range_ro, afr; const char *builtin; uint64_t logged; char key[24]; id<MTLBuffer> c1buf; uint64_t c1_gpu; size_t c1off;
+@interface NvShader : NSObject { @public id<MTLFunction> fn, tf_fn; uint16_t tf_stride[4]; nvn_refl rf; id<MTLArgumentEncoder> ae; NSUInteger arg_len; int stage, range_ro, afr, layer_out; const char *builtin; uint64_t logged; char key[24]; id<MTLBuffer> c1buf; uint64_t c1_gpu; size_t c1off;
     // The argument buffer encoded last time
-    id<MTLBuffer> ab_buf; uint64_t ab_gpu; size_t ab_off; uint64_t ab_pass;   // ab_pass: serial of the render pass that encoded this argument block - reused only within the same pass (the arena holding a block from another pass / another submission may already have been reclaimed and rewritten)
+    id<MTLBuffer> ab_buf; uint64_t ab_gpu; size_t ab_off; uint64_t ab_pass; uint32_t ab_binst;   // ab_pass: serial of the render pass that encoded this argument block - reused only within the same pass (the arena holding a block from another pass / another submission may already have been reclaimed and rewritten)
     id<MTLLibrary> bflib; int bflib_tried; NSMutableDictionary<NSString *, id<MTLFunction>> *bffn; } @end   // bflib: readable-tile version <key>_bf (merge B step 2)
 @implementation NvShader
 - (void)dealloc { free(rf.c1); }
@@ -2501,10 +2531,12 @@ static void pso_user_record(const char *fmt, ...) {
 }
 // Build the library from a .metallib file without permanently occupying a file descriptor: newLibraryWithURL keeps that file open for as long as the library object lives (one per game shader - of the 258 descriptors the iOS app permanently holds in a run, 213 are these,
 // while the default soft limit of an iOS process is 256; new shaders created in other acts raise it further). Here the file is mapped into memory and closed immediately, and the library is built from the in-memory data; the mapping is released when the library's data object is freed
+static __thread int t_lib_layer;
 static id<MTLLibrary> lib_from_file(NSString *path, NSError **err) {
     int fd = open(path.fileSystemRepresentation, O_RDONLY); struct stat st;
     if (fd < 0 || fstat(fd, &st) || st.st_size <= 0) { if (fd >= 0) close(fd); if (err) *err = [NSError errorWithDomain:NSPOSIXErrorDomain code:errno userInfo:@{NSLocalizedDescriptionKey: [NSString stringWithFormat:@"cannot open %@", path]}]; return nil; }
     size_t len = (size_t)st.st_size; void *m = mmap(NULL, len, PROT_READ, MAP_PRIVATE, fd, 0); close(fd);
+    t_lib_layer = m != MAP_FAILED && memmem(m, len, "render_target_array_index", 25) != NULL;   // the library's functions use the layer: for a vertex library, its output (shader_load)
     if (m == MAP_FAILED) { if (err) *err = [NSError errorWithDomain:NSPOSIXErrorDomain code:errno userInfo:@{NSLocalizedDescriptionKey: [NSString stringWithFormat:@"mmap failed %@", path]}]; return nil; }
     dispatch_data_t d = dispatch_data_create(m, len, dispatch_get_global_queue(QOS_CLASS_UTILITY, 0), ^{ munmap(m, len); });
     return [g_dev newLibraryWithData:d error:err];
@@ -2538,6 +2570,19 @@ static NvShader *shader_load(NSString *pc, int stage) {
         if (atomic_load(&g_depth_zero)) { if (j[@"entry_z"]) entry = j[@"entry_z"]; else RED_ONCE(85, "vertex shader %s has no near=0 entry (entry_z): depth stays in the GL convention", pc.UTF8String); }
     }
     o->fn = [lib newFunctionWithName:entry]; o->stage = stage; strlcpy(o->key, pc.UTF8String, sizeof o->key);
+    if (stage == 0) o->layer_out = t_lib_layer;   // writes the layer (render_target_array_index, named in the library's metadata): its passes are layered and its pipelines get an input class (pso_desc)
+    // 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_tf"]) {   // transform feedback capture entry point (dmsh_build.py vs_tf) and the bytes per vertex of each feedback buffer
+        o->tf_fn = [lib newFunctionWithName:j[@"entry_tf"]]; NSArray *ts = j[@"tf_strides"];
+        for (NSUInteger b = 0; b < 4 && b < ts.count; b++) o->tf_stride[b] = (uint16_t)[ts[b] unsignedIntValue];
+        if (!o->tf_fn) MLOG("[RED] shader %s: no transform feedback entry point %s in its library", pc.UTF8String, [j[@"entry_tf"] UTF8String]);
+    }
+    if (o->fn.functionConstantsDictionary.count) {
+        NSError *fe = nil; id<MTLFunction> f = [lib newFunctionWithName:entry constantValues:[MTLFunctionConstantValues new] error:&fe];
+        if (f) o->fn = f; else MLOG("[RED] shader %s: specializing its function constants failed: %s", pc.UTF8String, fe.localizedDescription.UTF8String);
+    }
     if (o->rf.nb) {
         o->ae = [o->fn newArgumentEncoderWithBufferIndex:o->rf.abi];
         int max_aid = -1;
@@ -2737,10 +2782,17 @@ static id<MTLTexture> tex_by_id(int tid) {   // Caller holds g_rmu
     if (r.has_v && r.v.set_swizzle) memcpy(sw, r.v.swz, 4);
     else if (r.t.set_swz) for (int c = 0; c < 4; c++) sw[c] = (uint8_t)(r.t.swz >> (3 * c) & 7);
     for (int c = 0; c < 4; c++) { if (sw[c] > MTLTextureSwizzleAlpha) sw[c] = (uint8_t)(MTLTextureSwizzleRed + c); if (sw[c] != MTLTextureSwizzleRed + c) need_sw = 1; }
-    if ((r.has_v && (r.v.set_format || r.v.set_levels || r.v.set_layers || r.v.set_target)) || need_sw) {
+    // Stencil mode (nvnTextureViewSetDepthStencilMode STENCIL): the view reads the stencil of a depth-stencil texture, through Metal's stencil-only format. The shader declares a float texture and
+    // takes the raw bits (the translator has no integer texture type): a read of the integer view gives the stencil value in those bits. Pokémon Legends: Z-A's deferred lighting tests stencil bit 6
+    // this way and discards; reading depth instead, the bit followed the depth value and the whole lit picture was striped
+    int stv = r.has_v && r.v.set_dsmode && r.v.dsmode == NVN_TEXTURE_DEPTH_STENCIL_MODE_STENCIL && base.pixelFormat == MTLPixelFormatDepth32Float_Stencil8;
+    if (r.has_v && r.v.set_dsmode && r.v.dsmode == NVN_TEXTURE_DEPTH_STENCIL_MODE_STENCIL && !stv)
+        RED_ONCE(104, "stencil-mode texture view of Metal format %lu (texture ID %d): no stencil-only view, depth is read", (unsigned long)base.pixelFormat, tid);
+    if ((r.has_v && (r.v.set_format || r.v.set_levels || r.v.set_layers || r.v.set_target)) || need_sw || stv) {
         if (!r.has_v) memset(&r.v, 0, sizeof r.v);
         uint32_t pf = (uint32_t)base.pixelFormat; uint8_t swz[4];
         if (r.v.set_format && !(nvn_fmt_to_mtl(r.v.format, &pf, swz) && pf)) pf = (uint32_t)base.pixelFormat;
+        if (stv) pf = MTLPixelFormatX32_Stencil8;
         if (r.v.set_format && nvn_fmt_bc(rp->t.d.format) && !nvn_fmt_bc(r.v.format)) {   // the raw blocks of a BC texture through an uncompressed view (see rt_view_fmt): Metal has no such view
             RED_ONCE(86, "Metal cannot create the raw-block view of a compressed texture (NVN format %#x → view %#x): treated as an empty slot", rp->t.d.format, r.v.format);
             return nil;
@@ -2769,6 +2821,7 @@ static inline __attribute__((always_inline)) void *tex_fast(int tid) {
 // The texture type the shader declares differs from the one the descriptor gives: use a compatible view. Maxwell sample instructions carry their own dimension, and a non-array instruction reading a 2D array texture reads layer 0;
 // on Metal a 2D array bound to a texture2d argument is a type mismatch with undefined reads (measured: one fixed value). Super Mario 3D World + Bowser's Fury creates its pre-lighting buffer as a one-layer 2D array
 // and the deferred lighting shader samples it as 2D: the ambient light of the whole picture became one colour, washed out white. Supports 2D array <-> 2D and 2D array -> cube / cube array; any other mismatch is reported once. Caller holds g_rmu
+static id<MTLTexture> dummy_tex(int dtt, int dpf); static uint64_t g_dummy_tex_rid[8][4];
 static id<MTLTexture> tex_fit(int tid, id<MTLTexture> t, MTLTextureType want, uint64_t *rid) {
     MTLTextureType have = t.textureType; NSUInteger ns = 0;
     if (want == MTLTextureType2D && have == MTLTextureType2DArray) ns = 1;
@@ -2776,6 +2829,11 @@ static id<MTLTexture> tex_fit(int tid, id<MTLTexture> t, MTLTextureType want, ui
     else if (want == MTLTextureTypeCube && have == MTLTextureType2DArray && t.arrayLength >= 6) ns = 6;
     else if (want == MTLTextureTypeCubeArray && have == MTLTextureType2DArray && t.arrayLength >= 6) ns = t.arrayLength / 6 * 6;
     else if (want == MTLTextureTypeCube && have == MTLTextureTypeCubeArray) ns = 6;   // a cube array read as a cube: the first cube
+    // A one-layer 2D texture in a cube slot has no cube view (Pokémon Legends: Z-A binds one to a cube sampler on its character creation screen, a placeholder for an environment map): the zero cube,
+    // a defined read, instead of a type mismatch whose reads are undefined
+    if (!ns && (want == MTLTextureTypeCube || want == MTLTextureTypeCubeArray) && have == MTLTextureType2D) {
+        RED_ONCE(103, "texture type mismatch: shader declares a cube, texture is 2D (texture ID %d); the zero cube is bound", tid);
+        id<MTLTexture> z = dummy_tex((int)want, NVN_DPF_FLOAT); *rid = g_dummy_tex_rid[want & 7][NVN_DPF_FLOAT]; return z; }
     if (!ns || tid <= 0 || tid >= g_ntreg) { RED_ONCE(76, "texture type mismatch: shader declares %lu, texture is %lu (texture ID %d); %s", (unsigned long)want, (unsigned long)have, tid,
                                                       want == MTLTextureTypeTextureBuffer || have == MTLTextureTypeTextureBuffer ? "a buffer texture: the slot reads 0 (dummy)" : "bound as is"); return nil; }
     treg_t *r = &g_treg[tid];
@@ -2840,6 +2898,15 @@ static id<MTLSamplerState> samp_by_id(int sid) {   // Caller holds g_rmu
     d.lodMinClamp = p.lod_min < 0 ? 0 : p.lod_min; d.lodMaxClamp = p.lod_max;
     d.maxAnisotropy = p.aniso < 1 ? 1 : p.aniso > 16 ? 16 : (NSUInteger)p.aniso;
     if (p.cmp_mode) d.compareFunction = (MTLCompareFunction)(p.cmp_func - 1);
+    // Reduction (NVN 1 min / 2 max): the minimum or maximum of the texels a filtered sample covers, per channel (Metal 26). Pokémon Legends: Z-A builds its depth pyramid with a min sampler
+    // (151607c7a92386cc reads the level above at an explicit LOD) and tests occlusion boxes against it; averaged, the pyramid is nearer than the scene at depth edges. Metal ignores the reduction
+    // with a nearest or no mip filter, so such samplers filter mips linearly: at the integer LODs reductions are used with that gives the same level, and a sampler without mips keeps to the base
+    // level through the LOD clamp
+    if (p.reduction == 1 || p.reduction == 2) {
+        d.reductionMode = p.reduction == 1 ? MTLSamplerReductionModeMinimum : MTLSamplerReductionModeMaximum;
+        if (d.mipFilter == MTLSamplerMipFilterNotMipmapped) { d.lodMinClamp = 0; d.lodMaxClamp = 0; }
+        d.mipFilter = MTLSamplerMipFilterLinear;
+    }
     d.borderColor = p.border[3] > 0.5f ? (p.border[0] > 0.5f ? MTLSamplerBorderColorOpaqueWhite : MTLSamplerBorderColorOpaqueBlack) : MTLSamplerBorderColorTransparentBlack;
     d.supportArgumentBuffers = YES;
     id<MTLSamplerState> s = [g_dev newSamplerStateWithDescriptor:d];
@@ -2884,6 +2951,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;
     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)
 } S;
 // Size of an attachment-less pass (rp_begin_ex_): the far corner of viewport 0, else of the scissor, clamped to Metal's 1..16384
 static void bare_dims(NSUInteger *w, NSUInteger *h) {
@@ -3227,6 +3295,8 @@ 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;
+#define TF_MARK ((const void *)&g_tf_mark)   // pso_key.fsfn of a transform feedback capture pipeline (tf_pso)
 static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvShader *fs, id<MTLFunction> fsfn, int merged) {
     MTLRenderPipelineDescriptor *d = [MTLRenderPipelineDescriptor new];
     d.vertexFunction = vs->fn; d.fragmentFunction = fsfn ? fsfn : fs ? fs->fn : nil;
@@ -3244,6 +3314,8 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS
         for (int i = 0; i < nk; i++) if (!strncmp(fs->key, keys[i], strlen(keys[i]))) {
             if (hl) d.fragmentFunction = [hl newFunctionWithName:[NSString stringWithFormat:@"hl%d", i]];
             break; } }
+    if (vs->layer_out) d.inputPrimitiveTopology = MTLPrimitiveTopologyClassTriangle;   // a vertex shader that writes the layer (pso_get_)
+    if (k->fsfn == TF_MARK) { d.vertexFunction = vs->tf_fn; d.fragmentFunction = nil; d.rasterizationEnabled = NO; }   // transform feedback capture (tf_draw)
     d.rasterSampleCount = k->samples ? k->samples : 1;
     d.alphaToCoverageEnabled = k->a2c;
     for (int i = 0; i < 8; i++) {
@@ -3371,6 +3443,13 @@ static __unsafe_unretained id<MTLRenderPipelineState> pso_get_(NvProg *pr, MTLRe
         if (g_pp_submit_mark != g_sub_serial) { g_pp_submit_mark = g_sub_serial; pso_predict_rest(); }   // First miss in this submission: first kick the later ones out to compile in parallel, then synchronously compile the one at hand
         MTLRenderPipelineDescriptor *d = pso_desc(&k, vs, fs, ms_fn ? ms_fn : g_pso_fsfn, g_pso_lmap != NULL);
         ps = [g_dev newRenderPipelineStateWithDescriptor:d error:&e];
+        // A vertex shader that writes the layer (render_target_array_index) needs the pipeline's input primitive class, which the reflection does not record and the pipeline key does not hold.
+        // Pokémon Legends: Z-A has two (b29f08619f379a83, 57246f9a04af8529, layered full-screen passes); Metal names the reason, the shader is marked and every later pipeline with it gets the
+        // triangle class (layered draws are triangle lists; a line or point draw with such a shader would fail validation)
+        if (!ps && !vs->layer_out && [e.localizedDescription containsString:@"render_target_array_index"]) {
+            vs->layer_out = 1; MLOG("vertex shader %s writes the layer: its pipelines get the triangle input primitive class", vs->key);
+            e = nil; d = pso_desc(&k, vs, fs, ms_fn ? ms_fn : g_pso_fsfn, g_pso_lmap != NULL); ps = [g_dev newRenderPipelineStateWithDescriptor:d error:&e];
+        }
     }
     g_bp_pso_n++; g_bp_pso_ns += bp_now() - bt;
     if (!ps) { char hint[200]; MLOG("[RED] failed to create render pipeline (vertex %s fragment %s): %s%s", vs->key, fs ? fs->key : "-", e.localizedDescription.UTF8String, pso_fout_hint(fs, g_pass.cfmt, hint, sizeof hint));
@@ -3428,6 +3507,8 @@ static void warm_worker(NSArray<NSString *> *paths) {
                     }
                     pso_key k = sk.k; k.vs = (__bridge void *)vs; k.fs = (__bridge void *)fs; k.fsfn = (__bridge void *)fsfn;
                     NSError *e = nil; id ps = [g_dev newRenderPipelineStateWithDescriptor:pso_desc(&k, vs, fs, fsfn, sk.merged) error:&e];
+                    if (!ps && !vs->layer_out && [e.localizedDescription containsString:@"render_target_array_index"]) {   // a vertex shader that writes the layer (see pso_get_)
+                        vs->layer_out = 1; e = nil; ps = [g_dev newRenderPipelineStateWithDescriptor:pso_desc(&k, vs, fs, fsfn, sk.merged) error:&e]; }
                     if (!ps) { MLOG("[RED] background render pipeline creation failed: %s", e.localizedDescription.UTF8String); continue; }
                     WarmItem *item = [WarmItem new]; item->kind = 1; item->vs_key = vkey; item->fs_key = fkey; item->vs = vs; item->fs = fs; item->fsfn = fsfn; item->pso = k; item->object = ps;
                     warm_enqueue(item); g_warm_rps++;
@@ -3502,7 +3583,7 @@ static id<MTLDepthStencilState> dss_get(int has_depth, int has_stencil) {
 }
 // Argument buffer: filled item by item per the B reflection (uniform slot r, separate texture / sampler slot r, storage buffer slot r, driver segments c0 / c2, c1, support_buffer)
 // Unbound resources read 0: give all-zero dummy resources (NVN descriptors 0..255 are reserved for the driver; the game binding handle 0 means 'this slot is empty')
-static uint64_t g_dummy_tex_rid[8][4], g_dummy_img_rid[8][4], g_dummy_tbuf_rid[4];
+static uint64_t g_dummy_img_rid[8][4], g_dummy_tbuf_rid[4];
 static id<MTLTexture> dummy_tex(int dtt, int dpf) {
     static __strong id<MTLTexture> cache[8][4];
     MTLTextureType tt = (MTLTextureType)dtt;
@@ -3591,13 +3672,14 @@ static void ubo_read(int stage, int reg, int off, size_t n, void *out) {
 }
 static uint32_t ubo_handle(int stage, int reg, int off) { uint32_t h; ubo_read(stage, reg, off, 4, &h); return h; }
 static uint32_t g_dispatch_dim[3];   // the current dispatch's workgroup counts (driver segment c0 + 0x200 for compute)
+static uint32_t g_draw_binst;   // the current draw's base instance (draw), for c0 + 0x644
 static __unsafe_unretained id<MTLBuffer> g_c0_dim_buf; static size_t g_c0_dim_off;   // where encode_args put them for this dispatch (nil: the shader reads no c0); an indirect dispatch copies its GPU-written counts there
 static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, int metal_vertex);
 static int encode_args(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, int metal_vertex) { AP_B; int r = encode_args_(en, sh, nvn_stage, metal_vertex); AP_E(3); return r; }
 static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, int metal_vertex) {
     if (!sh->ae) return 1;
     args_sync_dirty();
-    if (!passdump() && !g_args_dirty[nvn_stage] && sh->ab_buf && sh->ab_pass == g_pass.serial && nvn_stage != 5) {
+    if (!passdump() && !g_args_dirty[nvn_stage] && sh->ab_buf && sh->ab_pass == g_pass.serial && nvn_stage != 5 && !(sh->rf.drv_read[0] && sh->ab_binst != g_draw_binst)) {
         bind_args(en, nvn_stage, metal_vertex, sh->ab_buf, sh->ab_gpu, sh->ab_off, sh->rf.abi);
         g_stat_args_reuse++;
         return 1;
@@ -3751,6 +3833,11 @@ static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, i
                     // sets that to 1 and sizes its sprites itself (the fog splats ask for 3-18 px). Left at 0, every sprite had size 0, so the splats that push the Cap Kingdom
                     // fog aside never reached its velocity fields and the fog did not part around Mario. 2047 is Maxwell's point size limit (Metal clamps to its own)
                     if (nvn_stage != 5) { const float max_point = 2047.0f; memcpy(d + 0x610, &max_point, 4); }
+                    // Graphics stages: 0x644 holds the draw's base instance. Maxwell's instance ID leaves it out (GL's gl_InstanceID), and GLSLC adds this word where a shader needs the
+                    // absolute instance (gl_InstanceID + gl_BaseInstance): Pokémon Legends: Z-A's occlusion boxes (a793f436dcabd350, one box per draw at base instance 61, 62, …) pick their
+                    // transform this way, and 1,550 reads in 793 of its vertex shaders. Left at 0, every box was box 0: all six queries came back empty, and the game, taking its characters for
+                    // hidden, posed their legs straight back from the hips. 0x640 (read by 58 shaders, added to the vertex ID) stays 0: what the driver puts there is not established
+                    if (nvn_stage != 5) memcpy(d + 0x644, &g_draw_binst, 4);
                     if (nvn_stage == 5) {
                         memcpy(d + 0x200, g_dispatch_dim, 12); g_c0_dim_buf = db; g_c0_dim_off = doff + 0x200;
                         for (int r = 0; r < 16; r++) { uint64_t a = S.ubo[5][r].addr; uint32_t n = (uint32_t)S.ubo[5][r].size; memcpy(d + 0x210 + 16 * r, &a, 8); memcpy(d + 0x210 + 16 * r + 8, &n, 4); }
@@ -3781,7 +3868,7 @@ static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, i
     id<MTLBuffer> ab; size_t off = arena_alloc(arg_len, &ab);
     memcpy((uint8_t *)ab.contents + off, slots, arg_len);
     bind_args(en, nvn_stage, metal_vertex, ab, g_cur_arena->gpu + off, off, sh->rf.abi);
-    sh->ab_buf = ab; sh->ab_gpu = g_cur_arena->gpu + off; sh->ab_off = off; sh->ab_pass = g_pass.serial; g_args_dirty[nvn_stage] = 0;
+    sh->ab_buf = ab; sh->ab_gpu = g_cur_arena->gpu + off; sh->ab_off = off; sh->ab_pass = g_pass.serial; sh->ab_binst = g_draw_binst; g_args_dirty[nvn_stage] = 0;
     g_stat_args_enc++;
     return 1;
 }
@@ -4038,6 +4125,10 @@ static int state_op(uint32_t op, const void *a) {
         case OP_DEPTH_CLAMP: S.depth_clamp = ((const p_depth_clamp *)a)->on; break;
         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_RAST_DISCARD: S.rast_discard = ((const p_rast_discard *)a)->on; break;
         default: return 0;
     }
     return 1;
@@ -4067,6 +4158,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)
             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];
@@ -5468,10 +5560,69 @@ static int mdi_live(uint64_t cnt, int i) {
     }
     return (uint32_t)i < g_mdi_count;
 }
+// ---------- 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 = [g_dev newRenderPipelineStateWithDescriptor:pso_desc(&k, vs, nil, nil, 0) error:&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();
+    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);
+    for (int s_ = 0; s_ < 16; s_++) {
+        void *mb; uint64_t mo, g;
+        if (S.vbuf[s_].address && mtl_gpu_addr(S.vbuf[s_].address, 1, &mb, &mo, &g)) enc_r_vbuf(en, VB_SLOT(s_), g);
+    }
+    static id<MTLBuffer> zero; if (!zero) { zero = [g_dev newBufferWithLength:64 options:MTLResourceStorageModeShared]; res_add(zero); }
+    enc_r_vbuf(en, 30, zero.gpuAddress);
+    vconv_bind(en, &vs->rf, 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]);
+}
+
 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 ((di && di->mdi_cnt && !mdi_live(di->mdi_cnt, di->mdi_i)) || (dw && dw->mdi_cnt && !mdi_live(dw->mdi_cnt, dw->mdi_i))) return;   // past the multi-draw's count
     if (!S.prog) { RED_ONCE(18, "no program bound before the draw"); PD("Z %llu draw skipped: no program bound", g_stat_pass + !g_enc); return; }
     __unsafe_unretained NvProg *pr = prog_get(S.prog);
+    if (S.rast_discard) { tf_draw(pr, dw, di); return; }   // the vertex stage only: transform feedback capture, or nothing (tf_draw)
     if (g_nubd && g_enc && ubd_conflict(pr, dw, di)) { g_ubd_conflicts++; rp_end(); }   // land the deferred inline updates before this draw (see ubd_conflict)
     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); }
@@ -5524,7 +5675,10 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
     int need = g_enc || !tbdr_on() ? 3 : pass_ds_need(), mode = defer_q_draw ? 0 : need == 0 ? 2 : need == 2 ? 3 : need == 1 ? 4 : 0;   // 2 attach neither, 3 attach stencil only, 4 attach depth only
     int skip_d = mode != 0;
     if (!g_enc) { g_fd_dw = dw; g_fd_di = di; }
+    g_rp_layered = pr->sh[0] && pr->sh[0]->layer_out;
+    if (g_enc && g_pass.layered != g_rp_layered) rp_end();   // a layered pass only for layer-writing vertex shaders (see rp_begin_ex)
     id<MTLRenderCommandEncoder> en = rp_begin_ex(mode);
+    g_rp_layered = 0;
     g_fd_dw = NULL; g_fd_di = NULL;
     if (!en) { PD("Z %llu draw skipped: the render pass did not begin vs=%.16s fs=%.16s", g_stat_pass, pr && pr->sh[0] ? pr->sh[0]->key : "-", pr && pr->sh[1] ? pr->sh[1]->key : "-"); return; }   // Z lines: why a draw was dropped at these early returns
     if (defer_q_draw) defer_replay();
@@ -6376,6 +6530,10 @@ void mtl_exec(const uint8_t *p) {
             case OP_BIND_UBO: case OP_BIND_SSBO: case OP_BIND_TEX: 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_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_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;
             case OP_COPY_T2B: copy_t2b(a); break;
             case OP_COPY_T2T: copy_t2t(a); break;

dm/gfx/nvn/nvn_obj.h

@@ -45,7 +45,7 @@ FITS(view_t, NVNtextureView);
 _Static_assert(sizeof(view_t) <= 16, "view_t exceeds the newer SDKs' NVNtextureView (16 bytes)");
 
 // Sampler parameters (one copy shared by builder and object)
-typedef struct { uint8_t minf, magf, ws, wt, wr, cmp_mode, cmp_func; float lod_min, lod_max, lod_bias, aniso, border[4]; } samp_par;   // Enum values are all < 256
+typedef struct { uint8_t minf, magf, ws, wt, wr, cmp_mode, cmp_func, reduction; float lod_min, lod_max, lod_bias, aniso, border[4]; } samp_par;   // Enum values are all < 256; reduction takes the padding byte
 typedef struct { NVNdevice *dev; samp_par p; } samp_builder_t;                                FITS(samp_builder_t, NVNsamplerBuilder);
 typedef struct { uint32_t magic; samp_par p; } samp_t;                                        FITS(samp_t, NVNsampler);
 
@@ -100,7 +100,7 @@ enum {
     OP_BIND_PROGRAM, OP_BIND_STATE, OP_BIND_VATTR, OP_BIND_VSTREAM, OP_BIND_VBUF,
     OP_BIND_UBO, OP_BIND_SSBO, OP_BIND_TEX, OP_BIND_SAMP,
     OP_STENCIL, OP_POLY_OFFSET, OP_DRAW, OP_DRAW_IDX, OP_DISPATCH, OP_BARRIER,
-    OP_SET_TEX_POOL, OP_SET_SAMP_POOL, OP_COPY_B2T, OP_COPY_STRIDE, OP_BLEND_COLOR, OP_COPY_T2B, OP_BIND_IMG, OP_COPY_T2T, OP_RESET_COUNTER, OP_CLEAR_TEX, OP_CLEAR_BUF, OP_DEPTH_CLAMP, OP_RENDER_COND, OP_FENCE_SYNC, OP_COUNT_
+    OP_SET_TEX_POOL, OP_SET_SAMP_POOL, OP_COPY_B2T, OP_COPY_STRIDE, OP_BLEND_COLOR, OP_COPY_T2B, OP_BIND_IMG, OP_COPY_T2T, OP_RESET_COUNTER, OP_CLEAR_TEX, OP_CLEAR_BUF, OP_DEPTH_CLAMP, OP_RENDER_COND, OP_FENCE_SYNC, OP_TF_BIND, OP_TF_STATE, OP_RAST_DISCARD, OP_COUNT_
 };
 // Payload (immediately follows op_hdr)
 typedef struct { int32_t n; const tex_t *color[8]; view_t cview[8]; uint8_t has_cview[8]; const tex_t *depth; view_t dview; uint8_t has_dview; } p_set_rt;
@@ -133,6 +133,9 @@ typedef struct { const tex_t *tex; view_t view; uint8_t has_view, is_uint; NVNco
 typedef struct { int64_t stride; int image; } p_copy_stride;   // image: SetCopyImageStride (bytes per layer / depth slice) instead of SetCopyRowStride
 typedef struct { uint64_t addr, size; uint32_t value, _pad; } p_clear_buf;   // buffer filled with a 32-bit value (ClearBuffer)
 typedef struct { float rgba[4]; } p_blend_color;
+typedef struct { int32_t index, pad; uint64_t addr, size; } p_tf_bind;   // BindTransformFeedbackBuffer(s): one record per buffer (index 0-3)
+typedef struct { int32_t what, pad; uint64_t ctl; } p_tf_state;          // what: 0 begin, 1 end, 2 pause, 3 resume; ctl = the game's 32-byte control buffer
+typedef struct { int32_t on; } p_rast_discard;
 
 // Entry point of the Metal replay layer (nvn_mtl.m): replays a set of command handles into the current submission
 void mtl_exec(const uint8_t *handle);

dm/gfx/tools/dmsh_build.py

@@ -264,12 +264,47 @@ 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)
-    j = {'key': key, 'stage': stage, 'code': cp, 'tex': {}, 'img': {}, 'attr': {}, 'fout': {}, 'c1': c1.hex(),
+    tf = tf_layout(key, rec[ci:]) if stage == 'vs' 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['pc_missing'] = True; return j, c1   # not in the answer table: queries use defaults (listed separately in the summary)
     j.update(FACTS[key])
     return j, c1
 
+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), measured on Pokémon Legends: Z-A's
+    28 skinning shaders: +0x6dc u16 stride per buffer (4 buffers), +0x6e8 u8 per-buffer enable, +0x6ec u8 buffer count and +0x6ed 1; per buffer b, +0xdc + 0x80 b holds the captured output as
+    0x20 + n (generic output n) and +0x4dc + 0x80 b a u32 component count, captured from x on. Every buffer seen holds one output (stride = 4 x count): position xyz, packed normal / tangent bits.
+    → [[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)"""
+    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
+        locs = [x for x in ctl[0xdc + 0x80 * b:0x15c + 0x80 * b] if x]; n = struct.unpack_from('<I', ctl, 0x4dc + 0x80 * b)[0]
+        if ctl[0x6e8 + b] != 1 or len(locs) != 1 or not 0x20 <= locs[0] < 0x40 or not 1 <= n <= 4 or strides[b] != 4 * n:
+            print(f'[RED] {key}: transform feedback buffer {b} has a layout of unknown form (stride {strides[b]}, enable {ctl[0x6e8 + b]}, outputs {locs}, components {n}): no capture entry point; see tf_layout', flush=True)
+            return None
+        out.append([locs[0] - 0x20, n, strides[b]])
+    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)"""
+    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'
+    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 = ''.join(f' dm_tf{b}[{vid} * {n} + {i}] = {f"as_type<uint>(out.{names[o]}.{c})" if o in names else "0u"};'
+                     for b, e in enumerate(tf) if e for o, n, _ in [e] for i, c in enumerate('xyzw'[: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; }')
+
 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):
         Vertex: whether position x / y / z are taken unchanged from the .x / .y / .z of one and the same vertex attribute (w is the constant 1 or the same attribute's .w), with no branches and no clip distances in the main function - if so, pos_attr = attribute slot;
@@ -480,6 +515,16 @@ def vs_init(src):
     k += len('main0_out out = {};')
     return src[:k] + ''.join(f'\n    out.{n} = float4(0.0, 0.0, 0.0, 1.0);' for n in outs) + src[k:]
 
+def vs_layer(src):
+    """A vertex shader that writes the layer (render_target_array_index) loses the point size output. The translator writes gl_PointSize = 1.0 in every vertex shader; Metal builds a pipeline
+    whose vertex function writes the layer only with an input primitive class, and one that writes a point size only with the point class (or none), so with both no pipeline can be built.
+    Pokémon Legends: Z-A's layered full-screen passes (57246f9a04af8529, b29f08619f379a83: a triangle per instance, layer = instance + offset) draw triangles, where the point size is unused;
+    the host gives such pipelines the triangle class (nvn_mtl.m pso_get_). Shaders that do not write the layer are unchanged"""
+    i = src.find('struct main0_out\n{\n'); j = src.find('\n};', i)
+    if i < 0 or j < 0 or '[[render_target_array_index]]' not in src[i:j] or '[[point_size]]' not in src[i:j]: return src
+    src = src[:i] + re.sub(r'\n    float gl_PointSize \[\[point_size\]\];', '', src[i:j], count=1) + src[j:]
+    return re.sub(r'\n    out\.gl_PointSize = [^;\n]*;', '', src)
+
 def vs_depth(src):
     """Two entry points for the depth convention: the translator's vertex shader leaves z unchanged (NVN's near=0 convention); the main function is copied here - main0 gets the GL convention's (near=-w) z' = z/2 + w/2 before every return out
     (bit-identical to the formula Ryujinx used to emit in the translator), main0_z stays as is. The convention a game sets is runtime state (nvnDeviceSetDepthMode); the host picks the entry point by it, with nothing to configure per game.
@@ -559,13 +604,14 @@ def finish(key, w, j, c1):
             if r.returncode: return 'MSL failed', r.stderr.strip()[:160], None
             src = open(msl).read()
     pad = []
-    if j['stage'] == 'vs':   # outputs start as (0,0,0,1) (see vs_init)
+    if j['stage'] == 'vs':   # outputs start as (0,0,0,1) (see vs_init); no point size beside a layer output (see vs_layer)
-        src = vs_init(src); open(msl, 'w').write(src)
+        src = vs_layer(vs_init(src)); open(msl, 'w').write(src)
     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)
     if j['stage'] == 'vs':   # the two depth-convention entry points (see vs_depth); the analysis below uses the source with only main0
         src, full = vs_depth(src)
+        if j.get('tf'): full += vs_tf(src, j['tf'])   # transform feedback capture (see vs_tf)
         open(msl, 'w').write(full)
     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)
@@ -579,6 +625,7 @@ 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'} if j['stage'] == 'vs' else {}),
+            **({'entry_tf': 'main0_tf', 'tf_strides': [e[2] if e else 0 for e in j['tf']]} if j.get('tf') else {}),
             'threadgroup_size': j.get('local'), 'c1': c1.hex(), 'notes': res['notes'] + (['no paired PC reflection, queries use defaults'] if j.get('pc_missing') else []), 'drv_reads': dr, 'drv_unfillable': unfillable(dr),
             'bindings': reflect(key, j['stage'], src, res),
             'vertex_inputs': [{'location': int(m.group(2)), 'msl_type': m.group(1)} for m in re.finditer(r'(\w+)\s+\w+\s*\[\[attribute\((\d+)\)\]\]', src)],

dm/gfx/tools/shader_split.py

@@ -18,6 +18,7 @@
 #   · G2S (Koei Tecmo's KTGL, Hyrule Warriors: Age of Calamity): GLSLC compiler outputs inside the IDRK entries of the RDB data files (zlib in chunks), see idrk_shaders
 #   · SPKG (the older KTGL of Fire Emblem: Three Houses): shader packages of G2S containers, loose and inside the LINKDATA archive DATA0.bin / DATA1.bin, see kt_shaders and linkdata_jobs
 #   · SDF (Capcom's RE Engine, Monster Hunter Rise): shader definitions inside the PAK archives (KPKA, entries zstd / deflate compressed), only the newest copy of each path, see kpka_jobs and sdf_shaders
+#   · Trinity packs (Game Freak, Pokémon Legends: Z-A): BNSH inside the files of arc/data.trpfs (Oodle Mermaid, decoded by the game's own decoder in the Mac host), see trpfs_jobs and trinity_shaders
 # Code area: 0x30 header + (graphics) 0x50 SPH + instructions; header magic 0x12345678 for graphics, 0x12345679 for compute. The control section starts with 0x98761234, +0x704 = code area length,
 #   +0x700 = start of the c1 constant segment within the code area; for compute +0x748 thread group xyz (u32 ×3), +0x754 shared memory bytes (checked one by one against PC reflection on D2R's 22 compute shaders).
 # Stage: BNSH uses the stage slot; BAHS uses bits 10..13 of SPH word 0 (1 vertex 2 hull 3 domain 4 geometry 5 pixel), and a compute header means compute.
@@ -54,6 +55,14 @@ def ctl_len(ctl):
 def skey(ctl, code):
     return hashlib.sha256(ctl[:ctl_len(ctl)] + code).hexdigest()[:16]
 
+def same_shader(a, b):
+    """Two copies of one key hold the same shader: the code areas are equal, and so are the control sections up to their own length (control +0x20). Pokémon Legends: Z-A's UI layout archives
+    (__ArchiveShader.bnsh inside the .truiv packs) give 104 KB control sections whose stage descriptor length runs 456 bytes past their own length, and those trailing bytes differ between copies"""
+    (ca, ka), (cb, kb) = a, b
+    n = struct.unpack_from('<I', ca, 0x20)[0] if len(ca) >= 0x24 else len(ca)
+    n = n if 0x760 <= n <= min(len(ca), len(cb)) else max(len(ca), len(cb))
+    return ka == kb and ca[:n] == cb[:n]
+
 def check(where, st, ctl, code):
     """control section and code area consistent; returns the reason if not"""
     if ctl[:4] != CTL_MAGIC or len(ctl) < ctl_len(ctl): return 'bad control section'
@@ -284,6 +293,93 @@ def kpka_shaders(path, rel, ents, out):
         for e in ents:
             if e[3] >= 0x10 and kpka_data(d, e, 4) == b'SDF\0': out.extend(sdf_shaders(kpka_data(d, e), f'/{rel}#{e[0]:016x}'))
 
+ONEPACK = b'ONEPACK\0'
+def fb_field(d, t, i):
+    """FlatBuffers: absolute offset of field i of the table at t, None when the field is absent"""
+    vt = t - struct.unpack_from('<i', d, t)[0]
+    if 4 + 2 * i >= struct.unpack_from('<H', d, vt)[0]: return None
+    off = struct.unpack_from('<H', d, vt + 4 + 2 * i)[0]
+    return t + off if off else None
+
+def fb_vec(d, t, i):
+    """FlatBuffers: (start, count) of the vector in field i"""
+    p = fb_field(d, t, i)
+    if p is None: return 0, 0
+    p += struct.unpack_from('<I', d, p)[0]
+    return p + 4, struct.unpack_from('<I', d, p)[0]
+
+def fb_refs(d, t, i):   # the tables / strings a vector of offsets points at
+    st, n = fb_vec(d, t, i)
+    return [st + 4 * k + struct.unpack_from('<I', d, st + 4 * k)[0] for k in range(n)]
+
+def fb_root(d, at=0): return at + struct.unpack_from('<I', d, at)[0]
+
+def trpfs_jobs(files, parts=64):
+    """Game Freak's Trinity engine (Pokémon Legends: Z-A): almost everything is in arc/data.trpfs, packs of files, with arc/data.trpfd as their index, both FlatBuffers. The descriptor's root table:
+    0 file hashes [u64], 1 pack names [string], 2 file infos [{0 pack index u64}] (in file hash order), 3 pack infos [{0 pack size u64, 1 file count u64}]. The file system: 'ONEPACK\\0', u64 offset of its
+    root table {0 pack name hashes [u64], 1 pack offsets [u64]}, each pack a table of its own at its offset. A pack is named through its first file's hash (its own name hash is not a plain FNV-1a of the
+    descriptor's name). → ([(path, rel, ((pack offset, pack name), …)), …] in about equal sizes, the paths taken)"""
+    import mmap
+    jobs, taken = [], set()
+    for p, r in files:
+        if not p.endswith('.trpfs') or not os.path.exists(p[:-1] + 'd'): continue
+        with open(p, 'rb') as f:
+            if f.read(8) != ONEPACK: continue
+        fd = open(p[:-1] + 'd', 'rb').read(); fr = fb_root(fd)
+        hashes = struct.unpack_from(f'<{fb_vec(fd, fr, 0)[1]}Q', fd, fb_vec(fd, fr, 0)[0])
+        names = [fd[s + 4:s + 4 + struct.unpack_from('<I', fd, s)[0]].decode() for s in fb_refs(fd, fr, 1)]
+        pack_of = {h: struct.unpack_from('<Q', fd, fb_field(fd, t, 0))[0] if fb_field(fd, t, 0) is not None else 0 for h, t in zip(hashes, fb_refs(fd, fr, 2))}
+        sizes = [struct.unpack_from('<Q', fd, fb_field(fd, t, 0))[0] for t in fb_refs(fd, fr, 3)]
+        with open(p, 'rb') as f, mmap.mmap(f.fileno(), 0, access=mmap.ACCESS_READ) as d:
+            sr = fb_root(d, struct.unpack_from('<Q', d, 8)[0]); st, n = fb_vec(d, sr, 1)
+            packs = []
+            for off in struct.unpack_from(f'<{n}Q', d, st):
+                pr = fb_root(d, off); h0 = fb_vec(d, pr, 0)
+                if not h0[1] or struct.unpack_from('<Q', d, h0[0])[0] not in pack_of: red(f'{r}: the pack at {off:#x} has no file the descriptor lists')
+                i = pack_of[struct.unpack_from('<Q', d, h0[0])[0]]; packs.append((off, names[i], sizes[i]))
+        cur, size, step = [], 0, sum(s for _, _, s in packs) / parts
+        for off, name, s in sorted(packs):
+            cur.append((off, name)); size += s
+            if size >= step: jobs.append((p, r, tuple(cur))); cur, size = [], 0
+        if cur: jobs.append((p, r, tuple(cur)))
+        taken |= {p, p[:-1] + 'd'}
+    return jobs, taken
+
+_OODLE = None
+def oodle(comp, n):
+    """The game's own OodleLZ_Decompress (Trinity's packs are Oodle Mermaid, and Oodle is linked into main): one Mac host per worker process in its tool mode (dm/host/hle_oodle.c, DM_OODLE_SERVE),
+    fed over a pipe. DM_OODLE_HOST = the host, DM_OODLE_MODULES = the rewritten modules it loads (the NSP script sets both for profiles with shader_oodle on)"""
+    global _OODLE
+    if _OODLE is None:
+        import tempfile
+        host, mods = os.environ.get('DM_OODLE_HOST'), os.environ.get('DM_OODLE_MODULES')
+        if not host or not mods: red('Oodle-compressed packs: set DM_OODLE_HOST (the Mac host, make -C dm/host GAME=<id> dmhost) and DM_OODLE_MODULES (the rewritten modules) to decode them with the game\'s own decoder')
+        err = tempfile.TemporaryFile()
+        _OODLE = (subprocess.Popen([host, mods], stdin=subprocess.PIPE, stdout=subprocess.PIPE, stderr=err, env=dict(os.environ, DM_OODLE_SERVE='1')), err)
+    p, err = _OODLE
+    try:
+        p.stdin.write(struct.pack('<QQ', len(comp), n)); p.stdin.write(comp); p.stdin.flush()
+        got = struct.unpack('<Q', p.stdout.read(8))[0]
+    except (BrokenPipeError, struct.error): got = None
+    if got != n:
+        err.seek(0); red(f'the Oodle decoder in the Mac host failed on a {len(comp)} → {n} byte stream; ---- host output ----\n' + err.read().decode(errors='replace')[-3000:])
+    return p.stdout.read(n)
+
+def trinity_shaders(path, rel, packs, out):
+    """A range of Trinity packs (trpfs_jobs). Pack table: 0 file hashes [u64], 1 files [{1 compression i8 (3 = Oodle; stored otherwise), 2 level u8, 3 length u64, 4 data [u8]}]; each file is scanned like a loose one"""
+    import mmap
+    with open(path, 'rb') as f, mmap.mmap(f.fileno(), 0, access=mmap.ACCESS_READ) as d:
+        for off, name in packs:
+            pr = fb_root(d, off); hs = fb_vec(d, pr, 0)
+            for k, t in enumerate(fb_refs(d, pr, 1)):
+                c, ln = fb_field(d, t, 1), fb_field(d, t, 3)
+                comp, n = (struct.unpack_from('<b', d, c)[0] if c is not None else 0), (struct.unpack_from('<Q', d, ln)[0] if ln is not None else 0)
+                st, sz = fb_vec(d, t, 4)
+                if comp == 3: data = oodle(d[st:st + sz], n)
+                elif sz == n: data = d[st:st + sz]
+                else: red(f'/{rel}#{name}: file {k} has compression {comp} ({sz} → {n} bytes), which is not known')
+                scan_blob(data, f'/{rel}#{name}#{struct.unpack_from("<Q", d, hs[0] + 8 * k)[0]:016x}', 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"""
@@ -467,6 +563,9 @@ def scan_file(args):
     if head[:8] == IDRK:
         idrk_shaders(path, '/' + rel, found)
         return rel, uniq_keys(found), var_pairs(found), trees
+    if len(args) > 2 and head[:8] == ONEPACK:   # a range of Trinity packs (trpfs_jobs)
+        trinity_shaders(path, rel, args[2], found)
+        return rel, uniq_keys(found), var_pairs(found), trees
     if len(args) > 2 and head[:4] == KPKA:   # entries of an RE Engine PAK (kpka_jobs)
         kpka_shaders(path, rel, args[2], found)
         return rel, uniq_keys(found), var_pairs(found), trees
@@ -490,7 +589,7 @@ def uniq_keys(found):
     uniq = {}
     for where, st, ctl, code in found:
         key = skey(ctl, code)
-        if key in uniq and (uniq[key][2] != ctl or uniq[key][3] != code): red(f'key {key} is the same but the control section or code area differs ({where})')
+        if key in uniq and not same_shader(uniq[key][2:4], (ctl, code)): red(f'key {key} is the same but the control section or code area differs ({where})')
         uniq.setdefault(key, (where, st, ctl, code))
     return uniq
 
@@ -539,8 +638,9 @@ def main():
     J = int(os.environ.get('J', '8'))
     big = [j for p, r in files for j in (linkdata_jobs(p, r) or [])]   # Three Houses' 6.9 GB DATA1.bin: one job per range of entries, not one read of the whole file
     pak_jobs, paks = kpka_jobs(files)   # Monster Hunter Rise's PAKs: the newest copy of each path, in ranges of entries
-    big += pak_jobs
+    tr_jobs, trp = trpfs_jobs(files)   # Pokémon Legends: Z-A's Trinity packs (data.trpfs and its index data.trpfd), decoded by the game's own Oodle in the Mac host
-    files = [(p, r) for p, r in files if not linkdata_jobs(p, r) and p not in paks]
+    big += pak_jobs + tr_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]
     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 = {}, []
@@ -558,7 +658,7 @@ def main():
         pairs |= pr
         for key, v in uniq.items():
             if key in keys:
-                if keys[key][2] != v[2] or keys[key][3] != v[3]: red(f'key {key} is the same but the control section or code area differs ({v[0]} / {keys[key][0]})')
+                if not same_shader(keys[key][2:4], v[2:4]): red(f'key {key} is the same but the control section or code area differs ({v[0]} / {keys[key][0]})')
                 continue
             keys[key] = v; per_src[group(v[0])] += 1
     if SEPARABLE:

dm/host/hle_audio.c

@@ -296,6 +296,11 @@ static void au_close(void *g, void *lr) {
 }
 HLE_FORCE("_ZN2nn5audio13CloseAudioOutEPNS0_8AudioOutE", "close the audio output", void, h_au_Close, void *g) { au_close(g, __builtin_return_address(0)); }
 HLE_FORCE("_ZN2nn5audio16GetAudioOutStateEPKNS0_8AudioOutE", "playback state", int, h_au_State, const void *g) { return atomic_load_explicit(&ao(g)->state, memory_order_acquire); }
+// Pokémon Legends: Z-A (SDK 19): Wwise's sink wakeup thread (main+0x30faae8) asks after every release event and marks the sink starved at 0; the SDK's own version read the audout session the host never made
+HLE_FORCE("_ZN2nn5audio22GetAudioOutBufferCountEPKNS0_8AudioOutE", "buffers appended and not yet played to the end (the one playing included)", unsigned, h_au_BufCount, const void *g) {
+    aout *a = ao(g);
+    return atomic_load_explicit(&a->q_tail, memory_order_acquire) - atomic_load_explicit(&a->q_head, memory_order_acquire);
+}
 HLE_FORCE("_ZN2nn5audio21GetAudioOutSampleRateEPKNS0_8AudioOutE", "sample rate", int, h_au_Rate, const void *g) { (void)ao(g); return RATE; }
 HLE_FORCE("_ZN2nn5audio23GetAudioOutChannelCountEPKNS0_8AudioOutE", "channel count", int, h_au_Ch, const void *g) { return ao(g)->ch; }
 HLE_FORCE("_ZN2nn5audio15GetAudioOutNameEPKNS0_8AudioOutE", "device name; the host has only the default device", const char *, h_au_Name, const void *g) { (void)g; return "DeviceOut"; }

dm/host/hle_fs.c

@@ -5,6 +5,7 @@
 // Paths look like "rom:/data/x". FileHandle / DirectoryHandle are 8 bytes passed by value, here they store the host fd / DIR*.
 // Struct layout follows the public Horizon ABI: DirectoryEntry = {char name[0x301], pad[3], s8 type, pad[3], s64 size} (0x310 bytes).
 #include "hle.h"
+#include <mach/mach.h>
 #include "front.h"
 #include <dirent.h>
 #include <errno.h>
@@ -484,7 +485,13 @@ HLE_FORCE("_ZN2nn2fs8ReadFileEPmNS0_10FileHandleElPvm", "fsp-srv → pread(2)",
     dm_guest_write_prepare(buf, sz);   // The destination may be a texture pool page write-protected by the graphics layer: a kernel write to a read-only page returns EFAULT instead of page-faulting, so unprotect and mark dirty first
     while (got < sz) {
         ssize_t r = pread((int)h, (uint8_t *)buf + got, sz - got, off + got);
-        if (r < 0) { if (errno == EINTR) continue; return err_res("pread", "fd"); }
+        if (r < 0) { if (errno == EINTR) continue;
+            if (errno == EFAULT) {   // the destination is not writable host memory: say where, with the protection of its VM region
+                int e = errno; vm_address_t ra = (vm_address_t)((uint8_t *)buf + got); vm_size_t rs = 0; vm_region_basic_info_data_64_t vi; mach_msg_type_number_t cnt = VM_REGION_BASIC_INFO_COUNT_64; mach_port_t obj = MACH_PORT_NULL;
+                kern_return_t kr = vm_region_64(mach_task_self(), &ra, &rs, VM_REGION_BASIC_INFO_64, (vm_region_info_t)&vi, &cnt, &obj);
+                LOG("[RED] fs ReadFile into %p+%#zx failed at +%#zx (EFAULT): VM region %#lx+%#lx protection %d%s", buf, sz, got, (unsigned long)ra, (unsigned long)rs, kr == KERN_SUCCESS ? vi.protection : -1, kr == KERN_SUCCESS ? "" : " (none)");
+                errno = e; }
+            return err_res("pread", "fd"); }
         if (r == 0) break;
         got += r;
     }

dm/host/hle_misc.c

@@ -42,6 +42,7 @@ HLE_FORCE("_ZN2nn2ec24ShowShopAddOnContentListEv", "shop applet, out of scope; c
 // the host has no applet proxy (oe initialization is a no-op), and the SDK reads through the null proxy and crashes (sdk+0x38660). All return void: as if the player closed the shop at once, the game goes back to its menu.
 HLE_FORCE("_ZN2nn2ec24ShowShopAddOnContentListENS_13ApplicationIdE", "shop applet, out of scope; cut: no-op", void, h_ec_AocList1, void) { LOG("ec::ShowShopAddOnContentList (cut): the host has no shop"); }
 HLE_FORCE("_ZN2nn2ec24ShowShopAddOnContentListENS_13ApplicationIdERKNS_7account10UserHandleE", "shop applet, out of scope; cut: no-op", void, h_ec_AocList2, void) { LOG("ec::ShowShopAddOnContentList (cut): the host has no shop"); }
+HLE_FORCE("_ZN2nn2ec24ShowShopAddOnContentListERKNS_7account10UserHandleE", "shop applet, out of scope; cut: no-op", void, h_ec_AocList3, void) { LOG("ec::ShowShopAddOnContentList (cut): the host has no shop"); }
 HLE_FORCE("_ZN2nn2ec30ShowShopApplicationInformationENS_13ApplicationIdE", "shop applet, out of scope; cut: no-op", void, h_ec_AppInfo1, void) { LOG("ec::ShowShopApplicationInformation (cut): the host has no shop"); }
 HLE_FORCE("_ZN2nn2ec30ShowShopApplicationInformationENS_13ApplicationIdERKNS_7account10UserHandleE", "shop applet, out of scope; cut: no-op", void, h_ec_AppInfo2, void) { LOG("ec::ShowShopApplicationInformation (cut): the host has no shop"); }
 HLE_FORCE("_ZN2nn2ec24ShowShopSubscriptionListENS_13ApplicationIdE", "shop applet, out of scope; cut: no-op", void, h_ec_Sub1, void) { LOG("ec::ShowShopSubscriptionList (cut): the host has no shop"); }

dm/host/hle_offline.c

@@ -102,6 +102,8 @@ HLE_FORCE("_ZN2nn3err9ShowErrorENS0_9ErrorCodeE", "error display applet; the hos
 HLE_FORCE("_ZN2nn3err9ShowErrorERKNS0_18ErrorResultVariantE", "error display applet; the host only logs it", void, h_err_ShowVar, const void *v) { (void)v; LOG("game asks to show an error (ErrorResultVariant)"); }
 HLE_FORCE("_ZN2nn3err40ShowUnacceptableAddOnContentVersionErrorEv", "add-on content version mismatch error applet; the host only logs it", void, h_err_AocVer, void) { LOG("game asks to show an error: add-on content version mismatch"); }
 HLE_FORCE("_ZN2nn3err39ShowUnacceptableApplicationVersionErrorEv", "game version mismatch error applet; the host only logs it", void, h_err_AppVer, void) { LOG("game asks to show an error: game version mismatch"); }
+// Pokémon Legends: Z-A (SDK 19): the application's own error screen (an error code number and two messages in the argument) goes through the error applet as well
+HLE_FORCE("_ZN2nn3err20ShowApplicationErrorERKNS0_19ApplicationErrorArgE", "application error applet; the host only logs it", void, h_err_ShowApp, const void *arg) { (void)arg; LOG("game asks to show an application error (ApplicationErrorArg)"); }
 
 // ---------- web: web page applet (online lobby). Unavailable ----------
 HLE_FORCE("_ZN2nn3web13ShowLobbyPageEPNS0_20LobbyPageReturnValueERKNS0_16ShowLobbyPageArgE", "online lobby web applet (cut)", Result, h_web_Lobby, void *rv, const void *arg) { (void)rv; (void)arg; LOG("web::ShowLobbyPage (cut)"); return RES_FAIL; }
@@ -137,3 +139,5 @@ HLE_FORCE("_ZN2nn5album14SaveScreenshotEPKvmNS0_9ImageSizeENS0_16ImageOrientatio
 HLE_FORCE("_ZN2nn5album22SaveAndShareScreenshotEPKvmNS0_9ImageSizeENS0_16ImageOrientationENS0_17AlbumReportOptionERKNS_7account3UidEPKc", "save and share a picture: no album, fails", Result, h_album_Share,
           const void *p, size_t n, int size, int orient, int report, const void *uid, const char *text) { (void)p; (void)n; (void)size; (void)orient; (void)report; (void)uid; (void)text; return RES_FAIL; }
 HLE_FORCE("_ZN2nn5album14ShowAlbumFilesEv", "album applet: not shown", Result, h_album_Show, void) { return RES_FAIL; }
+// Pokémon Legends: Z-A (SDK 19): a screenshot of the current frame, taken by the capture service
+HLE_FORCE("_ZN2nn5album21SaveCurrentScreenshotENS0_17AlbumReportOptionE", "screenshot of the current frame (capsrv): no album, fails", Result, h_album_SaveCur, int report) { (void)report; LOG("album::SaveCurrentScreenshot (no album)"); return RES_FAIL; }

dm/host/hle_oodle.c

@@ -1,6 +1,10 @@
 // The game's own Oodle decoder, for the offline shader tools. Unreal Engine games compress each shader in their shader code archives with Oodle (Octopath Traveler 0: every entry of the
-// two .ushaderbytecode files in pakchunk0-Switch.pak, block header 8C 0A = Mermaid). Shaders are translated offline only and Oodle's decoder is not open source, so the NSP script decodes
+// two .ushaderbytecode files in pakchunk0-Switch.pak, block header 8C 0A = Mermaid), and Game Freak's Trinity engine every file of its packs (Pokémon Legends: Z-A: arc/data.trpfs, compression
-// them once on the Mac with the copy statically linked into main: DM_OODLE_JOB=<job file> runs OodleLZ_Decompress on every record right before nnMain, writes <job file>.out and exits.
+// type 3). Shaders are translated offline only and Oodle's decoder is not open source, so they are decoded on the Mac with the copy statically linked into main, in one of two modes:
+// DM_OODLE_JOB=<job file> runs OodleLZ_Decompress on every record right before nnMain, writes <job file>.out and exits (the NSP script, Unreal);
+// DM_OODLE_SERVE=1 answers requests on stdin / stdout from the guest main thread before the SDK starts, until stdin ends (shader_split.py's workers, Trinity: the packs are decoded file by file
+// while they are scanned, about 4.7 GB, so they are never written out). Request: u64 compressed length, u64 decompressed length, the compressed bytes; answer: u64 bytes decoded (0 on failure),
+// then those bytes.
 // Job file: "DMOJ", u32 count, then per record u32 compressed length, u32 raw length and the compressed bytes. Output: "DMOR", u32 count, then per record u32 decoded length
 // (0 = failed) and the decoded bytes.
 #include "hle.h"
@@ -84,3 +88,37 @@ void dm_oodle_job_maybe(void) {
     LOG("DM_OODLE_JOB: %u records decoded with the game's OodleLZ_Decompress (%u failed, %.1f MB out) → %s", n - bad, bad, raw_total / 1048576.0, outp);
     fflush(stderr); _exit(bad ? 1 : 0);
 }
+
+static int read_all(int fd, void *p, size_t n) {
+    for (uint8_t *q = p; n; ) { ssize_t k = read(fd, q, n); if (k <= 0) return 0; q += k; n -= (size_t)k; }
+    return 1;
+}
+static void write_all(int fd, const void *p, size_t n) {
+    for (const uint8_t *q = p; n; ) { ssize_t k = write(fd, q, n); if (k <= 0) die("DM_OODLE_SERVE: answer write failed"); q += k; n -= (size_t)k; }
+}
+void dm_oodle_serve_maybe(void) {
+    if (!getenv("DM_OODLE_SERVE")) return;
+    oodle_fn fn = NULL; uint64_t off = 0;
+    for (int i = 0; i < g_nmod; i++) if (!strcmp(g_mod[i].name, "main")) { off = oodle_find(&g_mod[i]); fn = (oodle_fn)(void *)(g_mod[i].base + off); }
+    if (!fn) die("DM_OODLE_SERVE: no main module");
+    enum { SLACK = 0x10000 };
+    size_t scratch_n = 16u << 20; uint8_t *scratch = malloc(scratch_n), *comp = NULL, *raw = NULL; size_t comp_cap = 0, raw_cap = 0; uint64_t n = 0, bad = 0;   // the decoder's work memory, as for a job
+    if (!scratch) die("DM_OODLE_SERVE: decoder memory");
+    LOG("DM_OODLE_SERVE: OodleLZ_Decompress at main+%#llx, requests on stdin", (unsigned long long)off);
+    for (;;) {
+        uint64_t hdr[2];
+        if (!read_all(0, hdr, sizeof hdr)) break;
+        if (hdr[0] > (1ull << 32) || hdr[1] > (1ull << 32)) die("DM_OODLE_SERVE: request too large (%#llx → %#llx)", (unsigned long long)hdr[0], (unsigned long long)hdr[1]);
+        if (hdr[0] > comp_cap) { free(comp); comp_cap = hdr[0]; if (!(comp = malloc(comp_cap))) die("DM_OODLE_SERVE: input buffer"); }
+        if (hdr[1] + SLACK > raw_cap) { free(raw); raw_cap = hdr[1] + SLACK; if (!(raw = malloc(raw_cap))) die("DM_OODLE_SERVE: output buffer"); }
+        if (!read_all(0, comp, hdr[0])) die("DM_OODLE_SERVE: request cut short");
+        int64_t got = fn(comp, (int64_t)hdr[0], raw, (int64_t)hdr[1], 1, 0, 0, NULL, 0, NULL, NULL, scratch, (int64_t)scratch_n, 3);   // fuzz safe, no CRC check, silent, unthreaded
+        uint64_t out = got == (int64_t)hdr[1] ? (uint64_t)got : 0;
+        if (!out && hdr[1] && bad++ < 8) LOG("[RED] DM_OODLE_SERVE: request %llu (%llu → %llu bytes) failed (decoder returned %lld)", (unsigned long long)n, (unsigned long long)hdr[0], (unsigned long long)hdr[1], (long long)got);
+        write_all(1, &out, sizeof out);
+        if (out) write_all(1, raw, out);
+        n++;
+    }
+    LOG("DM_OODLE_SERVE: %llu requests, %llu failed", (unsigned long long)n, (unsigned long long)bad);
+    fflush(stderr); _exit(bad ? 1 : 0);
+}

dm/host/hle_os.c

@@ -94,7 +94,8 @@ HLE("getentropy", int, g_getentropy, void *p, size_t n) {
     for (size_t i = 0; i < n; i += 8) { uint64_t v = dm_entropy64(); memcpy((uint8_t *)p + i, &v, n - i < 8 ? n - i : 8); }
     return 0;
 }
-// The SDK's own getentropy (Unicorn Overlord, SDK 16.2.4: std::random_device at nnMain) fills through nn::crypto::GenerateCryptographicallyRandomBytes, which opens the csrng service
+// The SDK's own getentropy (Unicorn Overlord, SDK 16.2.4: std::random_device at nnMain) fills through nn::crypto::GenerateCryptographicallyRandomBytes, which opens the csrng service;
+// Pokémon Legends: Z-A (SDK 19) calls it first thing in nnMain (main+0x1734b4)
 // (InitializeCsrng → sm::Initialize, the game aborted there). The call goes through the sdk's PLT, so the takeover gives it the same stream as getentropy above
 HLE_FORCE("_ZN2nn6crypto36GenerateCryptographicallyRandomBytesEPvm", "csrng service; host entropy (the getentropy stream)", void, g_csrng_bytes, void *p, size_t n) {
     g_getentropy(p, n);

dm/host/hle_svc.c

@@ -95,6 +95,18 @@ static int64_t alias_pages(uint64_t off, uint64_t n, int map) {
     }
     return map ? ch : -ch;
 }
+// Pages of the range not mapped yet (the bitmap is left alone)
+static int64_t alias_unmapped(uint64_t off, uint64_t n) {
+    int64_t c = 0;
+    for (uint64_t p = off >> 12, e = (off + n + 0xFFF) >> 12; p < e; ) {
+        uint64_t sh = p & 63, k = 64 - sh < e - p ? 64 - sh : e - p, m = (k == 64 ? ~0ull : (1ull << k) - 1) << sh;
+        c += __builtin_popcountll(~atomic_load(&g_alias_bits[p >> 6]) & m);
+        p += k;
+    }
+    return c;
+}
+static _Atomic int64_t g_alias_mapped;   // bytes mapped in the alias region (page bitmap)
+static pthread_mutex_t g_phys_mu = PTHREAD_MUTEX_INITIALIZER;   // mem_map_limit: the pool check and the map it allows are one step
 
 // Diagnostic DM_ALIAS_WINDOWS=<window MB>: every 30 s, the mapped memory of each window of that size in the alias region (windows aligned to the size in absolute addresses, counted from the page
 // bitmap; windows with nothing mapped are left out). For sizing an allocator that reserves fixed, aligned windows and commits pages from each window's start: Unreal's FMallocBinned3 takes
@@ -214,6 +226,21 @@ uint64_t svc_dispatch(int num, uint64_t *regs, uint64_t orig_addr) {
     case 0x2C: case 0x2D: {  // MapPhysicalMemory / UnmapPhysicalMemory(addr, size): only allowed inside the alias region (used by the SDK's VammManager)
         uint8_t *a = (uint8_t *)regs[0]; size_t size = regs[1];
         if (a < g_alias || a + size > g_alias + g_alias_cap) { LOG("SVC %#x outside the alias region %p+%#zx", num, a, size); regs[0] = 0xD401; return 0; }
+        // Profile mem_map_limit: a map that would take the process past the application pool fails with ResultLimitReached, as on Horizon (the pool less the used memory GetInfo reports).
+        // Pokémon Legends: Z-A's nninitStartup allocates and frees everything allocatable (pool − 64 MB, mapped at the top of its StandardAllocator region: 12 GB, 4 GB after the cut in dm/tools/patch_nso.py); the SDK keeps the freed pages mapped
+        // and maps new ones for later allocations, until a map fails: then it unmaps the cached block and retries (measured on the Mac: the refused 4 MB map at +0x43e2c000 was followed by
+        // UnmapPhysicalMemory of the 3.15 GB block and the same map again). Without the limit the mapped total went past the pool in the first minute and the freed block stayed mapped.
+        // Off by default: Ace Attorney Investigations Collection maps a little more than the 1 GB pool its profile reports
+        if (DM_GAME_MEM_MAP_LIMIT && num == 0x2C) {
+            pthread_mutex_lock(&g_phys_mu);
+            int64_t need = alias_unmapped((uint64_t)(a - g_alias), size) << 12, room = (int64_t)APP_POOL - (int64_t)(g_heap_cur + DM_GAME_MEM_USED + big_bss_bytes()) - atomic_load(&g_alias_mapped);   // the pool less the used memory GetInfo reports
+            if (need > room) {
+                pthread_mutex_unlock(&g_phys_mu);
+                static _Atomic int nref; if (atomic_fetch_add(&nref, 1) < 8) LOG("[mem] MapPhysicalMemory +%#zx refused: %.2f GB mapped, the application pool is full (ResultLimitReached, the SDK frees cached pages and retries)",
+                                                                       size, atomic_load(&g_alias_mapped) / 1073741824.0);
+                regs[0] = 0x10801; return 0;
+            }
+        }
         // Host pages are 16 KB, Horizon pages 4 KB: a map is rounded out to host pages (the extra neighboring 4K pages are inside the alias region and merely become writable early, harmless);
         // an unmap only reclaims host pages lying entirely inside the range (a partly covered page may still be used by an adjacent mapping, so it stays writable)
         uintptr_t pg = (uintptr_t)getpagesize(), lo = (uintptr_t)a, hi = (uintptr_t)a + size;
@@ -221,7 +248,7 @@ uint64_t svc_dispatch(int num, uint64_t *regs, uint64_t orig_addr) {
         int rc = 0;
         if (mhi > mlo) rc = num == 0x2C ? mprotect((void *)mlo, mhi - mlo, PROT_READ | PROT_WRITE)
                                         : (madvise((void *)mlo, mhi - mlo, MADV_FREE), mprotect((void *)mlo, mhi - mlo, PROT_NONE));
-        if (rc) die("SVC %#x mprotect %p+%#zx errno=%d", num, a, size, errno);
+        if (rc) { if (DM_GAME_MEM_MAP_LIMIT && num == 0x2C) pthread_mutex_unlock(&g_phys_mu); die("SVC %#x mprotect %p+%#zx errno=%d", num, a, size, errno); }
         if (g_flag.svc) LOG("%sPhysicalMemory %p+%#zx", num == 0x2C ? "Map" : "Unmap", a, size);
         // how deep the alias region is used (one line per 256 MB crossed): for sizing the reservation and alias region (iOS caps a single reservation at 6 GB)
         if (num == 0x2C) { static _Atomic uint64_t hw; uint64_t end = (uint64_t)(a + size - g_alias), old = atomic_load(&hw);
@@ -229,7 +256,8 @@ uint64_t svc_dispatch(int num, uint64_t *regs, uint64_t orig_addr) {
             if (end > old && end >> 28 != old >> 28) LOG("[mem] alias region mapped up to +%.2f GB (alias region %.2f GB)", end / 1073741824.0, g_alias_cap / 1073741824.0); }
         // Total actually mapped in the alias region (one line per 256 MB crossed, up or down): how much physical memory the game uses and whether it keeps growing. Counted with a page bitmap, only pages whose state really changed --
         // the SDK maps pages that are already mapped and unmaps ranges larger than what it mapped (Engage unmaps 0x1000000 and then 0x103f000 at the same address), so counting requested sizes went negative
-        { static _Atomic int64_t tot; int64_t d = alias_pages((uint64_t)(a - g_alias), size, num == 0x2C) << 12, before = atomic_fetch_add(&tot, d), after = before + d;
+        { int64_t d = alias_pages((uint64_t)(a - g_alias), size, num == 0x2C) << 12, before = atomic_fetch_add(&g_alias_mapped, d), after = before + d;
+          if (DM_GAME_MEM_MAP_LIMIT && num == 0x2C) pthread_mutex_unlock(&g_phys_mu);
           if (before >> 28 != after >> 28) LOG("[mem] alias region mapped %.2f GB (application pool %.2f GB)", after / 1073741824.0, APP_POOL / 1073741824.0); }
         regs[0] = 0; return 0;
     }
@@ -257,7 +285,9 @@ uint64_t svc_dispatch(int num, uint64_t *regs, uint64_t orig_addr) {
         case I_AslrAddr:  v = (uint64_t)g_as;    break;   case I_AslrSize:  v = RESERVE;     break;
         case I_StackAddr: v = (uint64_t)g_stack; break;   case I_StackSize: v = g_stack_cap; break;
         case I_TotalMem: case I_TotalNonSys: v = APP_POOL; break;
-        case I_UsedMem:  case I_UsedNonSys:  v = g_heap_cur + DM_GAME_MEM_USED + big_bss_bytes(); break;   // profile mem_used, 64 MB unless the game needs more (see mem_used in dm/games/README.md)
+        case I_UsedMem:  case I_UsedNonSys:  v = g_heap_cur + DM_GAME_MEM_USED + big_bss_bytes();   // profile mem_used, 64 MB unless the game needs more (see mem_used in dm/games/README.md)
+                                             // with mem_map_limit, the pages mapped in the alias region count too, as on Horizon (the limit in MapPhysicalMemory is the pool less this)
+                                             if (DM_GAME_MEM_MAP_LIMIT) v += (uint64_t)atomic_load(&g_alias_mapped); break;
         case I_SysResTotal: v = 0x2000000; break; case I_SysResUsed: v = 0x400000; break;
         case I_DebuggerAttached: v = 0; break;
         case I_RandomEntropy: { uint64_t dm_entropy64(void); v = dm_entropy64(); break; }   // a deterministic stream when DM_DETCLOCK is on (hle_os.c)

dm/host/hle_thread.c

@@ -188,6 +188,7 @@ static int64_t block(thr_t *t, int kind, int64_t ns) {
 
 // ---------- Locks ----------
 static inline uint32_t ld32(uint64_t a) { return __atomic_load_n((uint32_t *)a, __ATOMIC_SEQ_CST); }
+static inline uint64_t ld64(uint64_t a) { return __atomic_load_n((uint64_t *)a, __ATOMIC_SEQ_CST); }
 static inline void st32(uint64_t a, uint32_t v) { __atomic_store_n((uint32_t *)a, v, __ATOMIC_SEQ_CST); }
 static inline int cas32(uint64_t a, uint32_t *exp, uint32_t v) { return __atomic_compare_exchange_n((uint32_t *)a, exp, v, 0, __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST); }
 
@@ -749,11 +750,12 @@ static int svc_thread_impl(int num, uint64_t *regs) {
         regs[0] = 0; return 1;
     }
 
-    case 0x34: {  // WaitForAddress(x0 address, w1 type, w2 value, x3 timeout): 0 wait if less than, 1 decrement if less than then wait, 2 wait if equal
+    case 0x34: {  // WaitForAddress(x0 address, w1 type, x2 value, x3 timeout): 0 wait if less than, 1 decrement if less than then wait, 2 wait if equal, 3 wait if the 64-bit value is equal
+        // (Horizon 18; SDK 19's nn::os::WaitForAddress(const long*, long) uses it, which before fell into the 32-bit less-than test)
         thr_t *t = cur(); uint64_t addr = regs[0]; int type = (int)regs[1]; int32_t val = (int32_t)regs[2];
         pthread_mutex_lock(&G);
         int64_t r; uint32_t v = ld32(addr);
-        int ok = type == 2 ? (int32_t)v == val : (int32_t)v < val;
+        int ok = type == 3 ? ld64(addr) == regs[2] : type == 2 ? (int32_t)v == val : (int32_t)v < val;
         if (ok && type == 1) { while (!cas32(addr, &v, v - 1)) if (!((int32_t)v < val)) { ok = 0; break; } }
         if (!ok) r = R_BAD_STATE;
         else { t->waddr = addr; r = block(t, 3, (int64_t)regs[3]); }

dm/host/hle_time.c

@@ -97,6 +97,9 @@ HLE_FORCE("_ZN2nn4time11ToPosixTimeEPiPNS0_9PosixTimeEiRKNS0_12CalendarTimeE",
     if (n > 0 && out) out[0] = (int64_t)t;
     return RES_OK;
 }
+// With an explicit time zone rule (Pokémon Legends: Z-A, SDK 19; the rule came from the host's LoadTimeZoneRule above and stays unused): the device zone, as for ToCalendarTime
+HLE_FORCE("_ZN2nn4time11ToPosixTimeEPiPNS0_9PosixTimeEiRKNS0_12CalendarTimeERKNS0_12TimeZoneRuleE", "conversion with an explicit time zone rule; same as the device zone (host mktime)",
+          Result, h_time_ToPosixTimeRule, int *count, int64_t *out, int n, const CalendarTime *ct, const void *rule) { (void)rule; return h_time_ToPosixTime(count, out, n, ct); }
 // UTC calendar → absolute seconds (Super Mario Odyssey); no time zone involved, the host's timegm
 HLE_FORCE("_ZN2nn4time18ToPosixTimeFromUtcERKNS0_12CalendarTimeE", "UTC calendar → absolute seconds; host timegm", int64_t, h_time_ToPosixUtc, const CalendarTime *ct) {
     struct tm tm = {.tm_year = ct->year - 1900, .tm_mon = ct->month - 1, .tm_mday = ct->day, .tm_hour = ct->hour, .tm_min = ct->minute, .tm_sec = ct->second};

dm/host/main.c

@@ -15,6 +15,7 @@ static char *g_argv[64] = {"game"}; static int g_argc = 1;
 
 static void guest_main(void) {
     guest_thread_init();
+    { void dm_oodle_serve_maybe(void); dm_oodle_serve_maybe(); }   // DM_OODLE_SERVE tool mode: the game's own Oodle decoder over stdin / stdout (hle_oodle.c, exits when the input ends)
     uint32_t hmain = svc_thread_main_init();   // Main thread handle: x0 of nnosInitialize (stored in an sdk global, used by ThreadManagerImplByHorizon); x1 = pointer to the argument area
     // Copy Start step by step: _init_libc1 = __nnmusl_init_locks (initializes the stdio etc. locks as recursive locks; skipping it self-deadlocks on nested locking in putc) + __init_ssp.
     // The three callbacks passed in by rtld: A only sets an rtld flag (skipped), B = each module's DT_INIT (before nnMain), C = destructors (not done)

dm/tools/patch_nso.py

@@ -107,6 +107,18 @@ CUTS = {
     (0x139468, '08fca752', '087ca052', 'AddressRegionManager constructor: node area limit mov w8,#0x3fe00000 → mov w8,#0x3e00000'),
     (0x1394fc, '09fca752', '097ca052', 'AddressRegionManager constructor: length of the node area registered as an occupied region mov w9,#0x3fe00000 → mov w9,#0x3e00000'),
   ],
+  # Pokémon Legends: Z-A 2.0.2 (v393216) sdk (SDK 19.3.11): the same node area cut. The constructor is at sdk+0x1423ec (limit = alias start + 0x3fe00000 at +0x28, a unit heap of 0x30-byte
+  # nodes, the area then registered as occupied at +0x142468). With main's allocator region below, the guest block fits the 5 GB proven on a phone
+  'df7a3960c74c3a4803f826c587153252': [
+    (0x1423fc, '08fca752', '087ca052', 'AddressRegionManager constructor: node area limit mov w8,#0x3fe00000 → mov w8,#0x3e00000'),
+    (0x142468, '09fca752', '097ca052', 'AddressRegionManager constructor: length of the node area registered as an occupied region mov w9,#0x3fe00000 → mov w9,#0x3e00000'),
+  ],
+  # Pokémon Legends: Z-A 2.0.2 (v393216) main: nninitStartup (main+0x1b2e650) initializes its StandardAllocator with virtual address memory over 12 GB of address space (main+0x1b3e3e4), which
+  # the Switch's 64 GB alias region holds and an iPhone's single 6 GB reservation cannot. The allocator commits pages as it goes and the pool (3.2 GB) bounds what it maps (profile mem_map_limit);
+  # its largest block is nninitStartup's own probe of pool − used (3.15 GB, allocated and freed at once), so 4 GB of address space holds it
+  'b1f12fd919eae86ab8a978317677e64b': [
+    (0x1b3e3e4, '6200c0d2', '2200c0d2', 'StandardAllocator::Initialize(nullptr, size): address space mov x2,#0x300000000 → mov x2,#0x100000000 (12 GB → 4 GB)'),
+  ],
   # The Legend of Zelda: Link's Awakening 1.1.0 (v131072) main: the physics step (main+0x15730, main thread) rebuilds each world's PhysX completion task in place
   # (a PxLightCpuTask at world+0xc8) and passes it to simulate. A worker runs it (main+0x10ccf80: fetchResults, world+0x1d8 done = 1, notify, unlock) and then calls its
   # release (main+0x10cd070), which calls the task's own destructor: vtable back to the PxLightCpuTask base, mTm = 0 (main+0x10cced0). At 30 fps the game steps twice

docs/agents/01-extract.md

@@ -58,4 +58,5 @@ The work folder `$W` sits in the game's NSP folder, beside the NSP script's `<ti
 - Capcom's RE Engine (Monster Hunter Rise) keeps everything in PAK archives: `re_chunk_000.pak` plus `re_chunk_000.pak.patch_001.pak` … `patch_016` (19.6 GB in all). A PAK is `KPKA`, u8 version 4, u8 0, u16 features (0: a plain entry table), u32 count, u32 fingerprint, then one 0x30-byte entry per file: u32 murmur hashes of the lower- and upper-case UTF-16 path, u64 offset, u64 stored length, u64 length, u64 attributes (bits 0–3: 0 stored, 1 raw deflate, 2 zstd), u64 checksum. File names are only hashes. The game reads all of them as one view: a path in a later patch replaces the same hash in earlier PAKs (`kpka_jobs` in `shader_split.py` takes only the newest copy). The first bytes of each file name its format (`MESH`, `TEX\0`, `SDF\0` shader definitions, `BKHD` / `AKPK` Wwise, `USR\0`, `PFB\0`, …).
 - Koei Tecmo's KTGL (Hyrule Warriors: Age of Calamity) keeps everything in RDB archives: `asset/<name>.rdb` is the index (`_DRK0000` header, then one `IDRK0000` record per file), and `asset/<name>.rdb.bin*` plus `asset/data/0x<KTID>.file` hold the data as a run of `IDRK0000` entries: u64 entry length, u64 stored length, u64 length, u32 0, u32 file KTID, u32 type KTID, u32 flags, a few parameter words, then the stored data at the end of the entry. Entries follow each other directly, some after padding to 64 KB. Flag 0x100000 = zlib in chunks: u32 stored length + a zlib stream each, up to 16 KB inflated per chunk, a zero length ends. The first bytes of the inflated data name the format (`GT1G` textures, `_M1G` models, `_S2G` shaders, `KTSR` / `TSRS` / `ASRS` sound). The DLC folders add `.rdb.bin3_0` files of the same layout.
 - Unreal Engine (Octopath Traveler 0, UE5: the paged global distance field shaders and the LowLevelTasks task system) keeps everything in pak files: `Octopath_Traveler0/Content/Paks/pakchunk0-Switch.pak` (1.5 GB, 159,757 files) next to loose CRIWARE `.awb` sound banks and one `.usm` movie. A version 11 pak has its footer in the last 0xdd bytes (encryption key GUID, encrypted-index flag, magic `0x5A6F12E1`, version, index offset and size, SHA-1, five compression method names) and an AES-256 (ECB) encrypted index: mount point, entry count, path hash seed, path hash index, full directory index (a second encrypted block: directories, then file names with offsets into the encoded entries) and the encoded entries (one bit-packed u32 per file: block size, block count, encrypted, compression method, then offset and sizes). The key is in `main`: the one 32-byte window of its read-only or data segment that decrypts the index start to the mount point `../../../` (Octopath Traveler 0: `main+0x57c7f70`, found among 116,287 candidates in under a second). Never print it. Most files are Oodle-compressed in 64 KB blocks; the two shader code archives are stored. `dm/gfx/tools/uepak.py` reads all of this (AES through macOS CommonCrypto, no Python package) and lists a pak (`uepak.py list <pak> <main NSO> <out tsv>`).
+- Game Freak's Trinity engine (Pokémon Legends: Z-A) keeps almost everything in one archive, `arc/data.trpfs` (4.4 GB, 13,181 packs, 238,546 files, 17.7 GB unpacked), indexed by `arc/data.trpfd`; both are FlatBuffers. The descriptor's root table holds the file hashes (u64, sorted), the pack names (`arc/<the first file's path without its slashes>.trpak`), one info table per file (its pack index) and one per pack (size, file count). The file system starts with `ONEPACK\0` and the u64 offset of its own root table: the pack name hashes and the pack offsets. The name hash is not a plain FNV-1a 64 of the descriptor's name, so a pack is named through its first file's hash. A pack is a table of file hashes and file tables `{1 compression i8, 2 level u8, 3 length u64, 4 data}`; compression 3 is Oodle (every file but 15,528 small stored ones), and every Oodle stream is Mermaid (header `8C 0A`, `CC 0A` for an uncompressed chunk). Oodle is linked into `main` and there is no Mermaid decoder with a licence fit for this repository, so the shader split runs the game's own `OodleLZ_Decompress` in the Mac host's tool mode (`dm/host/hle_oodle.c`, `DM_OODLE_SERVE=1`: requests on stdin, answers on stdout). The host finds it in `main` by the one function that references its name string (as for Unreal's archives), and calls it as the game does (`main+0x1a9b7bc` on 2.0.2: fuzz safe, no CRC check, thread phase 3). The shader split's full scan (every file decoded, 17.7 GB) takes 5 minutes with 12 workers on a 12-core Apple M6.
 - The older KTGL of Fire Emblem: Three Houses uses a LINKDATA archive instead: `DATA0.bin` is an index of 32-byte entries `{u64 offset, u64 length, u64 stored length, u64 compressed}` into `DATA1.bin` (6.9 GB), and `patch1`–`patch4` hold loose overrides listed by `patchN/INFO0.bin` (0x120-byte entries ending in a `rom:/patchN/…` path). Its compressed entries and `*.gz` files are another chunked zlib: u32 chunk size (0x10000), u32 chunk count, u32 length, the stored length of each chunk, then the chunks at 0x80 alignment, each a u32 length + one zlib stream, or the plain bytes when the stored length equals the chunk's length (`kt_gz` in `shader_split.py`).

docs/agents/03-bringup.md

@@ -134,6 +134,14 @@ Each entry: symptom → root cause → fix (where the fix lives). Check these fi
   - **Addressables never start (no `StreamingAssets/aa/settings.json` in the `DM_FS=1` output at all), then a null asset crash (`main+0x31100ec`)** → Unity 2022's C# file API reaches the SDK's C library, not `nn::fs`: `File.Exists` → `SystemNative_Stat` → `stat` = `open` + `fstat` + `close` (`sdk+0x948b0`). nnmusl's path check (`sdk+0x904e0`, then the converter `sdk+0x8ef50`) asks `nn::fs::IsPathMounted`, `GetPathLengthMax`, `GetUtf16PathLengthMax`, `GetEntryNameLengthMax` and `GetUtf16EntryNameLengthMax` before it calls `nn::fs::GetEntryType` / `OpenFile`, and the SDK's own versions read the empty mount table (`IsPathMounted` false; the next one aborted with Result 0x35f202). The host answers them from its mount table, with Horizon's 0x300-byte limits and no UTF-16 limit (`hle_fs.c`). The path .NET builds is `/rom:/…`; nnmusl's converter turns it back into `rom:/…`. armyja's fork (GitHub issue #18) patched the managed stat wrappers by address instead; these hooks need no per-build table.
   - **`std::terminate` (`Break` from `sdk+0x5bfa48`) at the first real-time read** (right after the first-boot dialog, `main+0x3f3f68` → `system_clock::now()`) → `clock_gettime(CLOCK_REALTIME)` calls through a source pointer (`sdk+0xb05010`) that `nntimeStartup` sets once `time::Initialize` succeeds: 1 = `nn::time::GetStandardUserSystemClockCurrentTimespec`, 2 = the adjustable clock's (`sdk+0x8e414`). Both asked the time service; they are host take-overs now (`hle_time.c`), and `system_clock::now()` works on any build.
   - Not needed in the end: the process heap of pull request #16 (`nninitStartup` calls `SetMemoryHeapSize` / `AllocateMemoryBlock` only without virtual address memory, and nothing else in `main` calls them) and the tolerant common string table of pull request #15 (1.0.0.2's `main` has the table intact at `main+0x46169ac`, 109 entries ending at `Hash128`).
+- **SDK 19 (Pokémon Legends: Z-A, 19.3.11, Game Freak's Trinity engine)**. The import list against the hooks (the survey's `nso_syms.py` table, then `main`'s undefined `nn::` symbols that no `HLE_FORCE` names) left 86 candidates; most only fill the game's own structures (error and web argument builders, `prepo` report buffers, `swkbd` configuration, the audio resampler, `nn::async`). The ones that reach a service were hooked before they were met where the answer is clear: `time::ToPosixTime` with an explicit rule (the host's zone, as for `ToCalendarTime`), `album::SaveCurrentScreenshot` (no album: fails), `err::ShowApplicationError` (logged), and the `ShowShopAddOnContentList(UserHandle)` overload. One boot each:
+  - **`StandardAllocator::Initialize` aborts in `nninitStartup`** (`main+0x1b2e690`) → with virtual address memory (a 16 MB system resource) the game initializes its allocator over 12 GB of address space, `StandardAllocator::Initialize(nullptr, 0x300000000)` (`main+0x1b3e3ec`), as the Switch's 64 GB alias region allows. The alias region must hold the SDK's 1 GB node area and that region, 13 GB, more than a phone gives a process. Two cuts for 2.0.2 (`patch_nso.py`) make it fit the 5 GB block Tears of the Kingdom runs in: the allocator's address space 12 GB → 4 GB (`main+0x1b3e3e4`) and the node area 1 GB → 62 MB (the AddressRegionManager constructor, `sdk+0x1423fc` and `+0x142468`, Engage's cut). SDK 19's AddressRegionManager puts the region right after the node area (+0x3fe2c000 before the cut), not 1 GB aligned as Engage's SDK did, so a 4.25 GB alias region holds both (`mem_alias 0x110000000`). `QueryMemoryInfo` reporting more than 4 GB would take the game's big-memory path (heaps of allocatable − 4.27 GB instead of − 2.76 GB).
+  - **The alias region's mapped total runs past the application pool in the first minute** → `nninitStartup` allocates and frees everything allocatable (pool − used, 3.15 GB, mapped at the top of the region) and the SDK keeps the freed pages mapped; later allocations map new pages at the bottom until a map fails, and only then does it unmap the cached block and retry. The host never refused a map, so the block stayed mapped beside the 0.4 GB the game used. Profile key `mem_map_limit` (`hle_svc.c`): a map past the pool fails with ResultLimitReached, and the used memory `GetInfo` reports counts the mapped pages, as on Horizon (without that, the startup probe itself, sized as total − used, was refused). Measured: one refused 4 MB map, the 3.15 GB unmap, the same map again.
+  - **Abort in `sm::Initialize` under `crypto::GenerateCryptographicallyRandomBytes`** (the first call in `nnMain`, `main+0x1734b4`) → the bytes come from the csrng service. The host fills them from its entropy source, the same stream as `getentropy` (`hle_os.c`).
+  - **A null read in `audio::GetAudioOutBufferCount` on Wwise's sink wakeup thread** (`sdk+0x1e7ed0`, from `main+0x30faae8`) → after every release event Wwise asks how many buffers are still queued and marks the sink starved at 0; the SDK's version read the audout session the host never made. It answers from the host queue: appended and not yet played to the end (`hle_audio.c`).
+  - **NVN state read back through getters** (the red `NVN unimplemented function called` lines for `nvnColorStateGetBlendEnable`, `nvnBlendStateGetBlendFunc`, `nvnPolygonStateGetCullFace`, `nvnDepthStencilStateGetDepthTestEnable`, `nvnBufferBuilderGetSize`, `nvnTextureBuilderGetWidth`, …) → Trinity builds its pipeline states from state objects it reads back; the red-reporting stub returned whatever was in x0, the object's address, so a blend enable or depth test read as on and a texture builder's width as an address. The getters read the packed state objects (`hle_nvn.c`); the multisample setters the host lacked (alpha-to-coverage dither, raster samples, coverage modulation and coverage-to-colour) are recorded in fields `ms_st` already had.
+  - **`WaitForAddress` type 3** → Horizon 18 added a 64-bit "wait if equal" arbitration type, which SDK 19 exports as `nn::os::WaitForAddress(const long*, long)`; the host used to fall into its 32-bit "wait if less than" test, so such a wait could return at once or block when it should not. Z-A does not call it, but the SDK has it (`hle_thread.c`).
+  - **`nvnCommandBufferFenceSync` unimplemented, every `SyncWait` timed out at once** → the game fences once per frame from a command buffer and waits on it. A record (`OP_FENCE_SYNC`) fences the sync at replay with the fence value of the submission segment being encoded, so it signals once everything before it has completed.
 - **SDK 6 (Fire Emblem: Three Houses, 6.5.0)**:
   - **The first loading screen never ends; `DM_FS=1` shows `save:/system` "not found" and no `fs mount save`** → in SDK 6 `account::OpenPreselectedUser` returns a Result, and the game skips its save mount when it is not 0 (`main+0x53b7e0`). The host's hook returned void and left the handle pointer in w0. It returns 0 now; SDK 7 and later callers ignore it.
   - **No sound; red `audren: command 1 in the command list is malformed (magic 0 size 0)`** → SDK 6's DSP command header is +0x04 size (s16), +0x06 command ID (s16), +0x10 enabled, where SDK 8–20 have +0x04 enabled, +0x05 ID, +0x06 size (`CommandListProcessor::Process`, `sdk+0xaaa00`); its DeviceSink is ID 0x10. `capture` in `hle_audren.c` takes the layout whose walk finds the 0xcafebabe magic on every command of the first full list.

docs/agents/04-graphics.md

@@ -17,6 +17,7 @@ Do not stop at the obvious archive. In Mario the main `bfsha` held 4,888 of the
 - **SPKG in LINKDATA** (Fire Emblem: Three Houses, the older KTGL): shader packages (`SPKG`: `patch*/nx/shader/*.spkg` and 132 entries of `DATA1.bin`) of G2S containers as above, one `_SPK4100` package without G2S (the post effects, `DATA1.bin` entry 549), and map and model archives carrying packages of their own. `kt_shaders` inflates every LINKDATA entry and loose `*.gz` / `*.spkg` file ([01](01-extract.md)), takes every GLSLC output that passes the triple check, and scans its SARCs for BNSH; `DATA1.bin` is split into 64 scan jobs by entry range (`linkdata_jobs`). Its GLSLC (output version 0x1d.2) writes the 0x848 / 0x850-byte controls and headerless compute code described above, which `glslc_shaders` used to drop: 24 shaders before, 1,355 after. A few outputs of version 0x1d.0 keep something else at control +0x20 (odd values such as 0x136b), so a control length must be a multiple of 4. Nine more shaders exist only at run time (Seen in Fire Emblem: Three Houses, below).
 - **SDF in RE Engine PAKs** (Monster Hunter Rise): shader definitions (`SDF\0`, 2,685 files in the PAK view, [01](01-extract.md)): u16 version 1, u16 entry count, u64, then 0x108-byte entries from +0x10. Entry +0x10 holds six u64 offsets (vertex, hull, domain, geometry, pixel, compute; 0 = none), each to `{u32 control length, u32 code length}` followed by the control section (0x880 bytes) and the code area; compute code has no header. An entry's stages are one program, so its vertex / pixel stages are a pair. Empty definitions are 16 bytes of version 0. `kpka_jobs` maps the PAKs, takes the newest copy of each path, peeks at 4 decompressed bytes and inflates only SDF entries (`sdf_shaders`): 63,111 shaders (30,233 vertex, 21,475 pixel, 11,271 compute, 53 geometry, 33 hull, 46 domain) and 102,345 pairs, in 5 s with 12 jobs.
 - **Unreal Engine shader code archives** (Octopath Traveler 0: `ShaderArchive-Global-SF_GLSL_SWITCH_MOBILE-SP_GLSL_SWITCH_MOBILE.ushaderbytecode`, 1,219 shaders, and `ShaderArchive-Octopath_Traveler0_Chunk0-…`, 9,377, both stored in the pak, [01](01-extract.md)): version 2 archives of shader map hashes, shader hashes, shader map entries, shader entries `{u64 offset, u32 size, u32 raw size, u8 frequency}` (17 bytes packed; frequency 0 vertex, 3 pixel, 5 compute), preload entries and shader indices, then the code. Every shader is Oodle-compressed (Mermaid, block header `8C 0A`); there is no Oodle decoder in this repository, so the game's own `OodleLZ_Decompress` (statically linked into `main`; the host finds it as the one function that takes the address of the library's own name string "OodleLZ_Decompress", its profiler scope: Octopath Traveler 0 `main+0x4f77e10`, Pokemon Legends: Z-A `main+0x1a9c9e0`) decodes them in the Mac host: profile key `shader_ue_decode`, `DM_OODLE_JOB` (`dm/host/hle_oodle.c`), the NSP script's `decode_ue_shaders`. A decoded shader is an `LSLG` record: `LSLG`, a stage tag (`SV`, `SP`, `SC`), a header with the source name (`/Engine/Private/…usf:MainPS`), then `u32` code length, the code area, `u32` control length (0x880), the control section and a 24- or 29-byte tail, all at unaligned offsets (`uepak.lslg`). The host keys came out exactly: 37 of 37 shaders a first run reported missing were in the two archives. 10,324 unique shaders (1,457 vertex, 6,346 pixel, 2,521 compute); vertex and pixel shaders combine per draw, so `shader_separable on`. Bindings: separate textures and samplers, Engage's layout (`0x168` / `0x568`); 61 vertex shaders pick a texture slot at run time (profile key `tex_dyn_slots`, below), 2,084 shaders do atomics on buffer images (below). 16 compute shaders do atomics on a 3D `r32ui` image, which SPIRV-Cross's MSL 2.3 image-atomic emulation indexes with an `int3` (`array subscript is not an integer`): all are Unreal's paged global distance field (`GlobalDistanceFieldCompositeObjects.usf:CompositeObjectsIntoPagesCS`), which the mobile renderer does not dispatch (none in any run), so they stay unbuilt.
+- **BNSH in Trinity packs** (Pokémon Legends: Z-A, Game Freak's engine): every shader is a BNSH inside the Oodle-compressed files of `arc/data.trpfs` ([01](01-extract.md)). With the profile key `shader_oodle on`, `trpfs_jobs` / `trinity_shaders` in `shader_split.py` name the packs from `data.trpfd`, decode every file with the game's own Oodle in the Mac host (`DM_OODLE_HOST` / `DM_OODLE_MODULES`, set by the NSP script) and scan it like a loose file. 28,603 shaders (11,633 vertex, 16,845 pixel, 125 compute) and 26,051 variation pairs: the material shaders in `system_resource/gfx2/shader/NX64/env_ikkaku` (`material/standard.bnsh`, `material/transparent.bnsh`, `system/light.bnsh`, and the `special/imgui.trsha` pack with the engine's system shaders), particle effects (`.ptcl` / `.trtml` in `ik_effectdata`, 19,842 of the shaders), UI layouts (`__ArchiveShader.bnsh` in the `.truiv` packs of `uidata`) and the 40 BNSH in `main`. The layout archives' control sections are 104 KB long, and their stage descriptors run 456 bytes past the control's own length (control +0x20), with bytes that differ between copies of one shader; `same_shader` compares copies only up to the control's own length.
 - **BAHS** (agl `sharcb`, engine utility shaders): code region at header +0x1c (size +0x18). Binary entries in the table area are 16 bytes `{code offset, 0, control size (0x880/0x88c; 0x848 in BAHS version 9, Super Mario 3D World), signed offset from this entry to the control section}`. Control sections are shared between programs, so pair by entry, never by searching for code near a control.
 - **Shorter control sections (SDK 10, Super Mario 3D World + Bowser's Fury)**: the control section stores its own length at +0x20. Mario's, Engage's, Dread's and 13 Sentinels' are all at least 0x880 bytes; 3D World has 2,888 of 0x850 bytes (its ShaderData bfsha and six layouts) and 0x848 in BAHS version 9, each followed directly by other data. The key therefore hashes the first `min(0x880, control +0x20)` bytes (`ctl_len` in `shader_split.py`, `keys_hash` in `hle_nvn.c`); the other games' tables came out byte-identical.
 - **Compute code without a header (SDK 10)**: 3D World's particle compute shaders start their instructions at code +0 (like D2R), not after a `0x12345679` header. `shader_split.py` accepts both; `dmsh_build.py` strips only a header that is there.
@@ -190,6 +191,24 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog
 - **Integer fragment outputs.** With no PC version the translator has no output types and declares every output float, and Metal refuses a pipeline whose float output meets an integer attachment. The profile's `shader_facts_manual.tsv` declares them: the BC4 encoder above and the Ultrahand mask (`e3fb0bb58a691e10`, R8Uint, found by PR #11; the Mac route does not reach it). When another one fails, the pipeline red line prints the row to add. `nvnDeviceRegisterFastClearDepth`, a hardware hint, now returns success instead of reaching the red stub.
 - **Texture layout (open):** MacNativeGames found that NVN chooses the automatic block height by truncating `height_in_blocks * 1.5 / 8`, while `nvn_tex.c` rounds the division up, so some sizes (a 160×160 BC3 texture: H3 instead of H2) are read with the wrong layout. Rhythm Heaven Groove confirmed it on its two 40-block-row textures (below), and Touchable (a d2r-mobile port of the same game) on 16 of its texture files; `nvn_tex.c` now truncates for every game (`pick_tile(ah * 3 / 16)`). It changes layouts at 6–10, 17–21, 38–42 and 81–85 block rows, which matters only for data the game lays out itself (no packaged layout: Tears of the Kingdom's flushed textures, Monster Hunter Rise's sparse ones); those were not re-checked.
 
+### Seen in Pokémon Legends: Z-A
+
+Game Freak's Trinity engine on SDK 19 (NVN 55); every shader comes from the Oodle-compressed packs (section 1). Each item was found from a red line or a pass dump on the way through the language screen, character creation and the first minutes in Lumiose City:
+
+- **Characters collapsed into a spike (transform feedback skinning).** The character preview drew one thin spike. Trinity skins on the GPU with transform feedback: a vertex-only program (28 such shaders, e.g. `681cc42e48c0bbc0` in the `special/imgui.trsha` pack; skin indices in attribute 12, weights in 13) runs under `SetRasterizerDiscard` between `BeginTransformFeedback` / `EndTransformFeedback` and writes 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 then read those buffers as vertex streams 0–2 (`RGB32F`, `RGB10A2SN`, `RGB10A2SN`). The host ignored all of it, so the streams kept whatever was there. The pass dump gave it away: a pass of 19 draws with a vertex shader and no fragment shader (`fs=-`) before the character pass. Metal has no feedback stage, so:
+  - 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 enables, +0x6ec buffer count, and per buffer b the captured output (0x20 + n) at +0xdc + 0x80 b and its component count at +0x4dc + 0x80 b. Every buffer in this game holds one output from x on (strides 12/4/4, 12/12/12, 12/4, 12). Any other form gets a red line and no capture entry point, so one shader cannot stop a game's build (Tears of the Kingdom's 14,766 and Ace Attorney's 2,475 vertex shaders have no feedback at all);
+  - 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; the reflection lists `entry_tf` and `tf_strides`;
+  - the binds, begin / end / pause / resume and rasterizer discard are records now (`OP_TF_BIND`, `OP_TF_STATE`, `OP_RAST_DISCARD`). At replay a draw under rasterizer discard runs `tf_draw` in `nvn_mtl.m`: with feedback active and a capture entry point, an encoder of its own with no attachments (the game draws these while its frame target is bound), 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.
+- **A layered pass whose pipeline could not be built.** Two full-screen vertex shaders (`57246f9a04af8529`, `b29f08619f379a83`: one triangle per instance, layer = instance + offset, 20 pixel shaders read the layer) write `render_target_array_index`. Metal builds such a pipeline only with an input primitive class, and the translator writes a point size in every vertex shader, which needs the point class: no class fits both. `vs_layer` drops the point size from vertex shaders that write the layer, and the host gives their pipelines the triangle class once Metal names the reason (`layer_out` in `pso_get_` and in the background prewarm).
+- **`fragmentFunction main0 cannot be used to build a pipeline state` (a Metal validation assertion).** SPIRV-Cross declares an optional function constant, `spvLinearTextureAlignmentOverride`, in shaders that read a texel buffer (one imgui pixel shader here), and Metal refuses any function with function constants until it is specialized. `shader_load` specializes such functions with an empty set of values.
+- **Integer fragment outputs.** Two system pixel shaders write an ID into an `R8_UI` target on the character creation screen; their rows are in the profile's `shader_facts_manual.tsv`. A scan for outputs built from integer operations was not trusted: the translated code keeps float bits in integer variables (register reuse), so such a scan cannot tell an integer output from raw float bits stored into a float target.
+- **A 2D texture in a cube slot** (a placeholder for an environment map) has no cube view: the zero cube is bound instead of a type mismatch with undefined reads (`tex_fit`).
+- **Black world (layered rendering without a layered pass).** The 3D scene came out black once the opening reached Lumiose. The game builds its 32×32×32 colour grading LUT with `b29f08619f379a83`, one instance per depth plane, each writing its plane through `render_target_array_index`, and tone-maps the frame through that LUT. A Metal vertex shader selects the attachment layer per primitive only when the render pass has `renderTargetArrayLength`, so every plane went to plane 0 and the rest stayed black. A draw whose vertex shader writes the layer now opens a layered pass (`g_rp_layered` in `draw`, `g_pass.layered` in `rp_begin_ex`): the pass gets every layer the attachments have from their attached one on (3D textures: depth planes), and it is split when the next draw's shader differs. `DM_NVN_PASSDUMP` prints `L … layered pass: n layers`.
+- **Striped lighting (a stencil-mode texture view).** Every lit surface had black and white bands that followed depth, as if every pixel were in or out of shadow at random. The deferred lighting pixel shader (`ead6e2a32540b19a`, 31 bindings) first reads the depth-stencil buffer through a view whose `nvnTextureViewSetDepthStencilMode` is `STENCIL`, tests stencil bit 6 and discards. The host ignored the mode and bound depth, so the test read bit 6 of the depth value's float bits. `tex_by_id` binds an `X32_Stencil8` view for such a view. The shader declares a float texture (the translator has no integer texture type) and takes the raw bits with `as_type<uint>`, and a read of the integer view gives the stencil value in those bits. The stencil view needs no pixel-format-view usage (`tex_need_pfv`). Treating it as a format change made the host recreate the depth buffer while it was attached, and the next occlusion-counted draw crashed inside `setVisibilityResultMode`. Inputs ruled out on the way: the G-buffer albedo, the normals (12-bit pairs packed into 3 bytes, noisy by design), the screen-space shadow mask, the scaled depth copies and conditional rendering (`render_cond` changes nothing in this game and stays off).
+- **Characters' legs folded back from the hips once they walked (the base instance in the driver segment).** From the first step on, the player's legs stuck straight out behind her at hip height, and the rival's forward in a later cutscene; standing still beforehand looked right. The bone palette the game computes on the CPU was already bent, so the GPU side only drew it. The cause was six occlusion queries a frame, one 36-vertex box each (`a793f436dcabd350`, base instance 61, 62, …), drawn against mip level 2 of the depth pyramid: every query came back with 0 samples, and the game took its characters for hidden and posed their legs that way. Forcing every result to visible fixed the legs, and depth testing with ALWAYS did not, so the boxes never rasterized where they should. The box shader picks its transform with (instance index − base instance) + `c0[0x644]`: Maxwell's instance ID leaves the base instance out, and the NVN driver keeps it in its constant segment, where GLSLC reads it for `gl_InstanceID + gl_BaseInstance`. The host left `0x644` at 0, so every box was box 0. `encode_args` now writes the draw's base instance there for graphics stages and does not reuse a stage's argument block across draws with different base instances; 793 of the game's vertex shaders read it. The word before it, `0x640`, is added to the vertex ID by 58 shaders and stays 0 until what the driver puts there is known. `drv_reads` in `dmsh_build.py` only scans fragment-shader names (`fp_c0`), so these reads never showed up in the reflection; `nvn_refl.m` treats any declared driver-segment binding as read. How it was found: the raycasts were logged with scratch observation gates (all correct), then `DM_HLECOUNT` showed nothing walking-specific, then the counter reports.
+- **Sampler reduction (min / max).** The game builds that depth pyramid with a min-reduction sampler (`SamplerBuilderSetReductionFilter 1`; `151607c7a92386cc` samples the level above at an explicit LOD), which the host used to average. Metal 26 has the same thing (`MTLSamplerDescriptor.reductionMode`), but ignores it with a nearest or no mip filter, so reduction samplers filter mips linearly (`samp_by_id`): at the integer LODs reductions are used with, that is the same level.
+- **`newBufferWithBytesNoCopy` refused under memory pressure.** A memory pool created while the game maps and unmaps its 4 GB region was sometimes refused once, then accepted. `mtl_pool_add` retries up to 12 times with a short back-off and prints `[RED] memory pool … created at retry N`; on a final failure it logs the VM regions around the pool.
+
 ## 7. Diagnostic switches (host)
 
 | Switch | Output |

docs/development.md

@@ -298,6 +298,14 @@ The game's own work and its frame rate stay as they are; the host avoids work of
 - **Saves:** the first save (about 13 minutes in) needs `fs::GetSaveDataSizeMax`; four saves in a 26-minute run, and a later boot continues from them.
 - **Status:** brought up on the Mac through character creation and the story opening into free play (26 minutes from a fresh save with no red line, 2.74 GB resident), and playable on an iPhone: 30 fps through character creation, name entry and the story opening, no audio underruns, three saves, a footprint of about 2.35 GB after 3.4 minutes; it plays at 60 fps on an iPhone too (`fps_target 60`). Battles are not tested yet.
 
+**Pokémon Legends: Z-A (`0100F43008C44000`)**
+- **Status:** brought up on the Mac: language select, character creation, the opening and the first minutes of play in Lumiose City at a steady 30 fps (a 5-minute walk: GPU about 6 ms a frame, no crash; the red lines are the render condition the host does not evaluate, the placeholder cube and an occasional memory-pool retry). On an iPhone the first build (before the base-instance fix in [04](agents/04-graphics.md)) ran at 30 fps with a memory footprint of about 2 GB.
+- **Engine:** Game Freak's Trinity engine on SDK 19.3.11 (NVN 55) with Havok 2022.2, update 2.0.2 (v393216). The SDK 19 calls the host had to answer are in [03](agents/03-bringup.md).
+- **Files and shaders:** every file sits in the Trinity packs of `arc/data.trpfs` (FlatBuffers, Oodle Mermaid). The shader split decodes them with the game's own `OodleLZ_Decompress`, run in the Mac host's tool mode (`hle_oodle.c`, profile key `shader_oodle`), and finds 28,603 shaders ([01](agents/01-extract.md)).
+- **Memory:** the 5 GB guest block with two cuts for 2.0.2: the StandardAllocator's address space 12 GB → 4 GB and the SDK's node area 1 GB → 62 MB. With `mem_map_limit` the host refuses a map past the application pool, as Horizon does, so the SDK gives back the 3.15 GB it keeps mapped after `nninitStartup`.
+- **Graphics:** characters are skinned with transform feedback, the colour grading LUT is drawn as a layered pass, the deferred lighting reads stencil through a stencil-mode texture view, and the occlusion queries that decide how characters are animated need the draw's base instance in the driver segment ([04](agents/04-graphics.md)).
+- **Mac input:** a fresh save reaches play in Lumiose City with `DM_KBD=Taylor DM_KEYS="f450:A,f1200:LRight~1,f1600:A,f1900:Down,f2000:Down,f2100:A,f2400:A*80/90"`; a save made there continues with `DM_KEYS="f300:A*3/150"` (play from about frame 600), and `LUp~12` then walks along the platform.
+
 ## Tools (`dm/tools`)
 
 | Tool | What it does |

scripts/build_ipa_from_nsp.py

@@ -2,6 +2,7 @@
 # Unsigned IPA straight from the NSPs: identify the game (title ID → game profile dm/games/<id>/) → extract (nx_extract) → rewrite the code modules (patch_nso.py)
 #   → (profiles with shader_glslc_capture: capture the shaders the game compiles at run time, by booting it in the Mac host for a few seconds) → split out the shaders and check them against the profile's tables (shader_split.py)
 #   (a Vulkan game with shader_vk_capture, Super Mario Galaxy, boots the same way to capture the SPIR-V shaders it generates)
+#   (a game with shader_oodle, Pokémon Legends: Z-A, has its Oodle-compressed archives decoded for the split by its own decoder, run in the Mac host's tool mode)
 #   → iOS shader libraries (dmsh_build.py) → build_ipa.sh. A Vulkan game (profile key graphics vulkan: Hades) skips the NVN shader steps: vk_shaders.py translates its SPIR-V shaders instead. A game with no matching profile is surveyed (<NSP folder>/<title ID>_survey/survey.md), then gets a provisional profile guessed from
 #   the survey (<NSP folder>/<title ID>_build/profile/<title ID>/, outside the repository, found through DM_GAMES_DIR) and a best-effort IPA that may not boot; SURVEY_ONLY / --survey-only stops after the survey.
 # Two ways to run it: open this file in Spyder and press F5 (it uses the values under "Settings" below; an empty NSP folder or prod.keys is asked for at run time);
@@ -17,8 +18,8 @@ import collections, ctypes, getpass, hashlib, os, re, shlex, shutil, struct, sub
 from pathlib import Path
 
 # ---- Settings ----
-NSP_DIR = '/Volumes/T7/nsp/The Legend of Zelda Tears of the Kingdom/'                            # folder holding the NSPs (empty = ask at run time)
+NSP_DIR = '/Volumes/S790/nsp/The Legend of Zelda Tears of the Kingdom/'                            # folder holding the NSPs (empty = ask at run time)
-PROD_KEYS = '/Users/f2pgod/My Drive/Python_Scripts/nsp_check/prod.keys'                          # path of prod.keys (empty = ask at run time)
+PROD_KEYS = '/Users/seedds/My Drive/Python_Scripts/nsp_check/prod.keys'                          # path of prod.keys (empty = ask at run time)
 PACKAGE = ''                            # version packed into the IPA (empty = the profile's package; Mario 'v327680' i.e. 1.2.1, or 'v0')
 OUT_IPA = ''                            # output IPA (empty = <the profile's ipa> in the NSP folder)
 APP_NAME = ''                           # app name on the home screen (empty = the game name in the base NSP)
@@ -421,7 +422,8 @@ def survey(nsp_dir, app, fam, keys, repo, py, env):
     if extra: notes.append(f'modules besides main / sdk: {", ".join(f"{n} ({mods[n]["kind"]})" for n in extra)}; the provisional profile loads {G["modules"][0]} (load order main, subsdk*, sdk)')
     if rom['mags']['UnityFS'] or 'Unity IL2CPP' in F['markers']: notes.append('Unity: the shaders are Shader objects in UnityFS bundles and data.unity3d, which shader_split.py reads through unityfs.py (written for Unity 2020.3\'s Switch format, also read for 2017.4 and 2022.3; check other Unity versions first)')
     if F['glslc']: notes.append(f'main has {F["glslc"]} GLSLC compiler output magics (0x19866891): ' + ('the shaders are compiled into main (13 Sentinels, the KTGL games): shader_main_glslc on' if 'shader_main_glslc' in G else 'not its shaders, the RomFS has shader containers (Engage)'))
-    if not (rom['mags']['BNSH'] or rom['mags']['BAHS']) and not rom['mags']['UnityFS'] and not vulkan and 'shader_main_glslc' not in G:
+    if rom['mags']['Trinity pack']: notes.append('Trinity packs (Game Freak, arc/data.trpfs with its index data.trpfd): the shaders are BNSH inside Oodle Mermaid streams; shader_split.py decodes them with the game\'s own OodleLZ_Decompress in the Mac host, whose address dm/host/hle_oodle.c lists per main build (find it by its error strings, as for Pokémon Legends: Z-A)')
+    elif not (rom['mags']['BNSH'] or rom['mags']['BAHS']) and not rom['mags']['UnityFS'] and not vulkan and 'shader_main_glslc' not in G:
         notes.append('no BNSH / BAHS at the top level: the shaders are inside compressed or other containers or in main, and must be located first')
     major = {v.split('_')[0] for v in F['sdkver']}
     if not major: notes.append('no SDK version string found')
@@ -492,11 +494,11 @@ def write_profile(pdir, app, f):
          '# ---- NVN driver constant segments (Mario\'s values, the translator\'s built-in layout, unless the engine says otherwise) ----',
          *key('nvn_tex_base', '0x20'), *key('nvn_tex_slots', '32'), f'nvn_samp_base{T}0x568', f'nvn_img_base{T}0x120', f'nvn_img_slots{T}8', f'nvn_ssbo_bases{T}vs=0x110 ps=0x510 cs=0x310', '',
          '# ---- Shader split ----', *key('shader_groups', " ".join(groups) or "/=romfs"), f'shader_note{T}{name} shaders (provisional profile)',
-         *key('shader_main_glslc'), *key('shader_separable'), *key('shader_glslc_capture'), *key('shader_ue_decode')]
+         *key('shader_main_glslc'), *key('shader_separable'), *key('shader_glslc_capture'), *key('shader_ue_decode'), *key('shader_oodle')]
     (pdir / 'game.tsv').write_text('\n'.join(L) + '\n', encoding='utf-8')
     for n, text in EMPTY_TABLES.items(): (pdir / n).write_text(text)
     say(f'wrote a provisional game profile: {pdir / "game.tsv"}')
-    for k in ('graphics', 'modules', 'shader_vk_capture', 'shader_main_glslc', 'shader_glslc_capture', 'shader_ue_decode', 'mem_reserve', 'op_mode'):
+    for k in ('graphics', 'modules', 'shader_vk_capture', 'shader_main_glslc', 'shader_glslc_capture', 'shader_ue_decode', 'shader_oodle', 'mem_reserve', 'op_mode'):
         if k in G and G[k][1]: say(f'  {k} {G[k][0]}: {G[k][1]}')
 
 
@@ -914,9 +916,10 @@ def data_rows(b):   # a table's data rows (without # comment lines: the header c
     return [l for l in b.splitlines() if not l.startswith(b'#')]
 
 
-def split_shaders(py, repo, roms, mains, pkg, env, logs, g, provisional=False):
+def split_shaders(py, repo, roms, mains, pkg, env, logs, g, provisional=False, oodle=None):
     """The shader split into pkg, checked against the profile's tables. Returns its key (None for a provisional profile or damaged files, which never reuse it): the package is reused
-    when the tools, the profile and each version's main, RomFS (paths and sizes) and run-time capture are as at the last successful split"""
+    when the tools, the profile and each version's main, RomFS (paths and sizes) and run-time capture are as at the last successful split.
+    oodle: for profiles with shader_oodle on (Pokémon Legends: Z-A's Trinity packs), a function that builds the Mac host and returns the split's DM_OODLE_HOST / DM_OODLE_MODULES; called only when the split runs"""
     names = ('shader_map.tsv', 'shader_facts.tsv', 'shader_pairs.tsv')
     PROG.stage('shader split', step='checking the inputs')   # then a spinner: shader_split.py writes the package and the tables only at the end
     kf = pkg.with_name(pkg.name + '.key')
@@ -925,6 +928,8 @@ def split_shaders(py, repo, roms, mains, pkg, env, logs, g, provisional=False):
                                                    ('files', sorted((r.parent / 'ue_shaders').glob('*.ushaderbytecode'))))])
     if key and pkg.is_dir() and (note := reused(kf, key)) is not None:
         PROG.end(f'{note}, reused (inputs unchanged)'); return key
+    if oodle:   # the packs are decoded by the game's own Oodle in the Mac host (dm/host/hle_oodle.c)
+        PROG.set(step='building the Mac host for the game\'s Oodle decoder'); env = dict(env, **oodle())
     PROG.set(step='splitting')
     tmp = pkg.with_name(pkg.name + '.partial'); rmtree(tmp); (tmp / 'tables').mkdir(parents=True)
     run([py, repo / 'dm/gfx/tools/shader_split.py', ','.join(map(str, roms)), ','.join(map(str, mains)), tmp / 'pkg'] + [tmp / 'tables' / n for n in names], logs / 'shader_split.log', env=env)
@@ -1231,7 +1236,9 @@ def main(args=None):
         cap = capture_vk(repo, game, patched, work, env, logs, g) if vkcap else None
         ios = build_vk_shaders(py, repo, [game / 'RomFS'] + ([cap] if cap else []), sh, env, logs, g, reuse=not (provisional or ALLOW_DAMAGED))
     else:
-        skey = split_shaders(py, repo, [dirs[v] / 'RomFS' for v in need], [dirs[v] / 'ExeFS/main' for v in need], sh / 'pkg', env, logs, g, provisional)
+        # Trinity packs (shader_oodle): the split decodes them with the packed version's own Oodle decoder, in the Mac host built for the profile
+        oodle = (lambda: {'DM_OODLE_HOST': str(build_mac_host(repo, work / 'tools/dmhost', env, logs, g)), 'DM_OODLE_MODULES': str(patched)}) if g.get('shader_oodle', 'off') == 'on' else None
+        skey = split_shaders(py, repo, [dirs[v] / 'RomFS' for v in need], [dirs[v] / 'ExeFS/main' for v in need], sh / 'pkg', env, logs, g, provisional, oodle)
         ios = build_shaders(py, repo, sh, xlate, env, logs, g, provisional, skey)
 
     if free_gb(out_ipa.parent) < 5: raise BuildError(f'less than 5 GB free on the volume of {out_ipa.parent}')