Age of Calamity's last red lines, fixed for every game: buffer textures read by 1D integer fetches are bound as 2D rows, and shadow vertex inputs replace the CPU vertex conversion

- Row wrap: Maxwell reads a texture buffer and a 1D texture with the same instructions (TLD with integer coordinates only, TLDS Texture1DLodZero) and the bound descriptor decides. Without a type answer the translator read such a slot as 2D at (x, 0) and the host bound a dummy for a buffer texture, so it read 0. The vendored translator (InstEmitTexture.cs, marked nxnative:) now addresses those fetches as (x & 0x3FFF, x >> 14) and reports the slot (fetch1d); the host binds a buffer texture there as a linear 2D texture of 16,384-texel rows over the same memory (tbuf_as_2d, nvn_tbuf_rows), cached by texture serial. KTGL's light and decal lists: Age of Calamity 16,438 shaders, Three Houses 321 (torches now light its map, blood decals are lit); FFT 5, Monster Hunter Rise 3 - Shadow vertex inputs: for each vertex shader with integer inputs the build writes <key>_bv (dmsh_vin.py), which declares every integer input again as float4 and as the other-sign integer at attribute indices 16-30 and ORs them into the input; the host sends each attribute to the index of its format's class and the zero buffer to the others (vin_ct, vin_on). Age of Calamity's Mac battle: 0 vertices converted on the CPU instead of 545,000 per submission, submit + present 10.0 -> 5.7 ms a frame, CPU 3.97 -> 2.44 W, the same picture. DM_NVN_VCONV_CPU=1 gives the CPU conversion back - A texture type mismatch no view bridges is logged once per shader and slot with its texture ID, NVN target and formats (tex_mis_note); it showed Age of Calamity's remaining one: 3D slots 8 and 9 of three vertex shaders bound with 2D textures (bound as is, as before) - Chain procedure for every profile: the 18 with game files on this drive rebuilt for Mac and iOS (c1_full.tsv / cb_extent.tsv: header lines only; new cb_extent.tsv for Hades II, Cuphead and Monster Hunter Rise), the others header-only with their lists carried over. Mac smoke runs of 13 games: no new red lines - Tests: t_tbuf2d (row aliases against texture_buffer views), t_vin (shadow inputs against the CPU conversion, every format class), test_dmsh_vin.py, nvn_tbuf_rows in t_nvn
2026-10-22 00:28:29 +08:00
1 parent ca2c35c1ae
commit 87e647f5bb
@@ -1,3 +1,3 @@
-# chain b-16260256abdc 2026-10-21T07:21:04+08:00 (list carried over from b-166a14482eae: that chain change did not touch any c1 read)
+# chain b-61eff511fce3 2026-10-21T21:45:38+08:00
 6405f3e55facc7c9
 d9097b8aba3a7080
@@ -1,4 +1,4 @@
-# chain b-16260256abdc 2026-10-21T12:01:22+08:00
+# chain b-61eff511fce3 2026-10-21T21:45:49+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 00008cfe4fc1d31b	5:304 6:1856 7:64 12:16
 0004d9615081f2dd	5:304 6:1856 7:64
@@ -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 "16260256abdc"
+#define DMSH_CHAIN "61eff511fce3"
@@ -1,4 +1,4 @@
-# chain b-6c7fd55d331c 2026-10-21T07:21:04+08:00 (list carried over from b-c792d8a4c610: that chain change did not touch any c1 read)
+# chain b-b3fcfcb9ddcc 2026-10-21T23:23:14+08:00 (list carried over from b-6c7fd55d331c: that chain change did not touch any c1 read)
 0ffd2c92087e4b54
 69d3868c8d828fe6
 8a06e3dda01235c1
@@ -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 "6c7fd55d331c"
+#define DMSH_CHAIN "b3fcfcb9ddcc"
@@ -1,4 +1,4 @@
-# chain b-5ca6f7dcb1a9 2026-10-21T07:21:04+08:00 (list carried over from b-dde5a04eafc6: that chain change did not touch any c1 read)
+# chain b-1a20936f730c 2026-10-21T23:25:42+08:00
 35cbc686dd8461b8
 4872bec9a5567f8b
 54a618792b2d5f0f
@@ -1,4 +1,4 @@
-# chain b-5ca6f7dcb1a9 2026-10-21T12:01:10+08:00
+# chain b-1a20936f730c 2026-10-21T23:25:49+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 000429e5b56f54be	0:304 2:528 3:704
 0011e7066f8b040e	0:144 1:224 2:112
@@ -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 "5ca6f7dcb1a9"
+#define DMSH_CHAIN "1a20936f730c"
@@ -1,3 +1,3 @@
-# chain b-7b0126e8d949 2026-10-21T07:21:04+08:00 (list carried over from b-821df0347986: that chain change did not touch any c1 read)
+# chain b-7f751dfec0ce 2026-10-21T23:27:25+08:00
 25c5829f891f9b2c
 976839d3ba92c8f3
@@ -1,4 +1,4 @@
-# chain b-7b0126e8d949 2026-10-21T12:00:35+08:00
+# chain b-7f751dfec0ce 2026-10-21T23:27:30+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 000064cdb6305b93	5:304 7:64 12:320
 00021722b48824b8	5:304 6:1872 7:112
@@ -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 "7b0126e8d949"
+#define DMSH_CHAIN "7f751dfec0ce"
@@ -1,4 +1,4 @@
-# chain b-a9687d01ebfa 2026-10-21T09:39:35+08:00
+# chain b-8200d93ca492 2026-10-21T23:44:26+08:00
 0002d193f6c52811
 0010e810186cdc91
 001299ca51e856c2
@@ -1,4 +1,4 @@
-# chain b-a9687d01ebfa 2026-10-21T13:05:16+08:00
+# chain b-8200d93ca492 2026-10-21T23:45:10+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 0002d193f6c52811	0:1872
 00088e149f8757ca	0:1872
@@ -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 "a9687d01ebfa"
+#define DMSH_CHAIN "8200d93ca492"
@@ -52,7 +52,8 @@ fps_target	60
 # 2,800 compute dispatches a frame after its first cutscene, which each had an encoder of their own (nvn_mtl.m dispatch_enc)
 fps_hold	off
 # KTGL binds whole mesh pools (up to ~590,000 vertices) and its soldiers' Half2 / Short2 attributes feed integer shader inputs, which the replay converts on the CPU:
-# over the bound ranges that was millions of vertices a frame in battle (26 fps at any target on the Mac); only over the vertices each draw reads
+# over the bound ranges that was millions of vertices a frame in battle (26 fps at any target on the Mac); only over the vertices each draw reads. Shadow vertex inputs
+# (nvn_mtl.m vin_on) now take those attributes on the GPU; this still serves the draws that keep the CPU conversion (transform feedback, DM_NVN_VCONV_CPU=1)
 vconv_range	on

 # ---- Shader split ----
@@ -1 +1 @@
-# chain b-adfdb13db222 2026-10-21T20:02:25+08:00
+# chain b-85e8c98ac176 2026-10-22T00:21:36+08:00 (list carried over from b-adfdb13db222: that chain change did not touch any c1 read)
@@ -1,4 +1,4 @@
-# chain b-adfdb13db222 2026-10-21T20:02:26+08:00
+# chain b-85e8c98ac176 2026-10-22T00:21:36+08:00 (list carried over from b-adfdb13db222: that chain change did not touch any uniform read)
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 01705ebebe3e302b	0:64 1:16 2:304
 020b72102b82b310	0:64
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010036E00FB20000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "adfdb13db222"
+#define DMSH_CHAIN "85e8c98ac176"
@@ -1,4 +1,4 @@
-# chain b-c18727f7b8bf 2026-10-21T15:40:22+08:00
+# chain b-7820419c3539 2026-10-21T23:27:35+08:00
 1b5e5c3e653702d2
 1bfeb381d95d18c6
 280094f212199d77
@@ -1,4 +1,4 @@
-# chain b-c18727f7b8bf 2026-10-21T15:40:22+08:00
+# chain b-7820419c3539 2026-10-21T23:27:35+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 01058b659b699495	0:288
 0117c03be720951e	0:560
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 010038B015560000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "c18727f7b8bf"
+#define DMSH_CHAIN "7820419c3539"
@@ -1 +1 @@
-# chain b-601174f81d6a 2026-10-21T07:21:04+08:00 (list carried over from b-b0883f45f985: that chain change did not touch any c1 read)
+# chain b-d3964897130f 2026-10-21T23:27:41+08:00
@@ -1,4 +1,4 @@
-# chain b-601174f81d6a 2026-10-21T12:01:30+08:00
+# chain b-d3964897130f 2026-10-21T23:27:41+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 0040f948baa144f4	0:1904
 00870f86a49cf3b9	0:1904 12:64
@@ -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 "601174f81d6a"
+#define DMSH_CHAIN "d3964897130f"
@@ -1,4 +1,4 @@
-# chain b-2eb33bdda057 2026-10-21T07:21:04+08:00 (list carried over from b-57e1140395a0: that chain change did not touch any c1 read)
+# chain b-3824e28513d2 2026-10-21T23:23:14+08:00 (list carried over from b-2eb33bdda057: that chain change did not touch any c1 read)
 006a737523e70720
 00da789b4ca18c03
 00e5f2d3b5851fab
@@ -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 "2eb33bdda057"
+#define DMSH_CHAIN "3824e28513d2"
@@ -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 "b6dd6fe0cd3c"
+#define DMSH_CHAIN "4e619a4691df"
@@ -1,4 +1,4 @@
-# chain b-1866d9bc0803 2026-10-21T07:21:04+08:00 (list carried over from b-6f195117b2ca: that chain change did not touch any c1 read)
+# chain b-1977372f08b5 2026-10-21T23:31:19+08:00
 1a2b65cd05ff557c
 21c7ff0d9591a637
 33e4340e5bcbc5c6
@@ -1,4 +1,4 @@
-# chain b-1866d9bc0803 2026-10-21T12:01:30+08:00
+# chain b-1977372f08b5 2026-10-21T23:31:20+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 00388ae860457bd9	0:160
 00a24b2c0ac95e17	0:32
@@ -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 "1866d9bc0803"
+#define DMSH_CHAIN "1977372f08b5"
@@ -1 +1 @@
-# chain b-5f5ac45ebe51 2026-10-21T07:21:04+08:00 (list carried over from b-a74254ba6a54: that chain change did not touch any c1 read)
+# chain b-d605d84674ef 2026-10-21T23:23:14+08:00 (list carried over from b-5f5ac45ebe51: that chain change did not touch any c1 read)
@@ -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 "5f5ac45ebe51"
+#define DMSH_CHAIN "d605d84674ef"
@@ -1,4 +1,4 @@
-# chain b-c3f1a81d79ec 2026-10-21T07:21:04+08:00 (list carried over from b-35936311a3c9: that chain change did not touch any c1 read)
+# chain b-f9b2bbcfc576 2026-10-21T23:29:29+08:00
 26c5fd7adbc12cc0
 67f23f4ad5925e06
 96a3e288731b2433
@@ -1,4 +1,4 @@
-# chain b-c3f1a81d79ec 2026-10-21T12:01:29+08:00
+# chain b-f9b2bbcfc576 2026-10-21T23:29:36+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 00008ab75998b65a	0:16 2:16 3:16 5:16 6:16
 000be95364b95b90	0:416 1:880 4:816 6:16 s0 s1
@@ -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 "c3f1a81d79ec"
+#define DMSH_CHAIN "f9b2bbcfc576"
@@ -1,3 +1,3 @@
-# chain b-40143d368395 2026-10-21T07:21:04+08:00 (list carried over from b-10befbce38df: that chain change did not touch any c1 read)
+# chain b-15547f29143a 2026-10-21T23:06:35+08:00
 21fe759b8114941c
 60e2041ed375cae1
@@ -1,4 +1,4 @@
-# chain b-40143d368395 2026-10-21T12:00:30+08:00
+# chain b-15547f29143a 2026-10-21T23:06:45+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 0002f3adb60ad088	5:304 6:1792 7:112 10:64 13:32
 0006cbc2ca63e38b	6:1520 12:432
@@ -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 "40143d368395"
+#define DMSH_CHAIN "15547f29143a"
@@ -1 +1 @@
-# chain b-cf4190f31748 2026-10-21T21:12:44+08:00
+# chain b-d6930bbe1baa 2026-10-22T00:21:36+08:00 (list carried over from b-cf4190f31748: that chain change did not touch any c1 read)
@@ -1,4 +1,4 @@
-# chain b-cf4190f31748 2026-10-21T21:12:44+08:00
+# chain b-d6930bbe1baa 2026-10-22T00:21:36+08:00 (list carried over from b-cf4190f31748: that chain change did not touch any uniform read)
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 00298d8d603fb418	0:16
 00ba482babf36d17	0:16
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 01008BC024FDE000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "cf4190f31748"
+#define DMSH_CHAIN "d6930bbe1baa"
@@ -1,3 +1,3 @@
-# chain b-ef65926816d8 2026-10-21T07:21:04+08:00 (list carried over from b-9c3f90e7f282: that chain change did not touch any c1 read)
+# chain b-1a9b40cf13d7 2026-10-21T23:23:22+08:00
 57eb9f0bae677756
 f826f2b59cfa39a8
@@ -1,4 +1,4 @@
-# chain b-ef65926816d8 2026-10-21T12:01:11+08:00
+# chain b-1a9b40cf13d7 2026-10-21T23:23:23+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 003b38fe9b3627ba	0:192 2:64
 00500b94cf88443d	0:1168
@@ -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 "ef65926816d8"
+#define DMSH_CHAIN "1a9b40cf13d7"
@@ -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 "59873f9b6a5a"
+#define DMSH_CHAIN "6c7c083b2aa7"
@@ -1,3 +1,3 @@
-# chain b-3f938d83f0c7 2026-10-21T17:40:06+08:00
+# chain b-2542626d576f 2026-10-21T23:46:07+08:00 (list carried over from b-3f938d83f0c7: that chain change did not touch any c1 read)
 01a847b848df5e33
 5fb90d2733898803
@@ -1,4 +1,4 @@
-# chain b-3f938d83f0c7 2026-10-21T17:40:13+08:00
+# chain b-2542626d576f 2026-10-21T23:46:44+08:00 (list carried over from b-3f938d83f0c7: that chain change did not touch any uniform read)
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 000325b608b12e54	1:448 9:32 10:288
 000b783c5701c854	1:112
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 01009BF0072D4000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "3f938d83f0c7"
+#define DMSH_CHAIN "2542626d576f"
@@ -1 +1 @@
-# chain b-22e4d6ecbc13 2026-10-21T07:21:04+08:00 (list carried over from b-119d78df4ead: that chain change did not touch any c1 read)
+# chain b-6f84388c810c 2026-10-21T23:29:40+08:00
@@ -0,0 +1,57 @@
+# chain b-6f84388c810c 2026-10-21T23:29:40+08:00
+# shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
+00b66ee69b7fde95	1:16 2:80
+02d134235f7f498b	0:128 1:16
+05d651ee6dc20ccc	1:16
+0851573221ea6cfa	2:240
+0d60254e46a57b8c	2:16
+102a3746727e5029	1:240
+15bde041d415b8cb	1:16 2:48
+175952012bc68784	0:128 1:16
+1e59b871d956e5a3	0:80
+21425f3e93b44a5e	2:16
+2234ecf1a206c06d	1:16 2:64
+254a73f8b7214283	0:128
+2de1efb8564f02fa	1:80
+3cfbb93f88ce24fc	0:128
+3db1c491ba7b6ff2	2:48
+40d2414e7d69711b	1:16
+410736215a5ed337	1:32
+42f68f016c9381f2	1:32
+4570a40f26776c99	1:16
+4c41035860215e19	0:64
+4f441f9187d9c1b3	0:128
+55893025726667d3	0:128 1:16
+5715938e6de314fb	1:16
+5d0736e9a1558078	0:128
+5f93206484de3596	1:16
+60bbb69a4b411dcd	1:16
+67c6c57ab2a55a0a	1:16 2:48
+6ed9cff7efcde280	2:16
+7be530b1cc195b2c	2:16
+800fef7e9a51f10c	1:16
+84123d26a641bbcb	1:16
+89a736b2558bdfcb	1:16
+8bd6710d108c88b6	1:16
+8eb964efef64d58e	0:128
+9168db30caac6253	2:16
+92eb59ba5ba4c1ed	1:16
+9a912e9da383d799	2:256
+a4f27b9d4a0e89ad	1:16
+a57dea937fd309a7	0:128
+a83ce5b2fde6caf2	1:16
+ae6e4a655d11e89f	1:16
+b114e406f89e38a7	1:64
+b1665fd487e31e55	2:48
+b246044c1f067ed8	1:16
+b6ebda3fafc3a313	1:32
+b9ede1459ecf43bd	0:128
+c73a9fd48874c720	2:16
+df85791082d97fa3	2:272
+df9564a0b9ea7175	0:128
+e71decce4706f348	2:16
+f203b77d1681cc81	0:80
+f54b3ecef6f96a60	0:128
+f72073702fa4546e	3:64
+fa02a9cef59b0e6e	2:16
+febf2d05477f2d3d	0:128 1:16
@@ -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 "22e4d6ecbc13"
+#define DMSH_CHAIN "6f84388c810c"
@@ -1,4 +1,4 @@
-# chain b-f1a0d6cec37d 2026-10-21T09:17:11+08:00
+# chain b-08f1551e74a6 2026-10-21T23:29:58+08:00
 101c4d73dbb253e2
 a277d45a64c92ab6
 b14f0b0add21e60a
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100A5C00D162000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "f1a0d6cec37d"
+#define DMSH_CHAIN "08f1551e74a6"
@@ -1,4 +1,4 @@
-# chain b-eaf7da2d5f8f 2026-10-21T07:21:04+08:00 (list carried over from b-d98265054359: that chain change did not touch any c1 read)
+# chain b-86c71329294b 2026-10-21T23:31:03+08:00
 001275671d2c9985
 0021d75a567e2122
 00a98da1158c6f14
@@ -1,4 +1,4 @@
-# chain b-eaf7da2d5f8f 2026-10-21T12:00:57+08:00
+# chain b-86c71329294b 2026-10-21T23:31:05+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 0005f969d0503763	2:160 3:16 4:96
 001275671d2c9985	0:160 3:128
@@ -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 "eaf7da2d5f8f"
+#define DMSH_CHAIN "86c71329294b"
@@ -1 +1 @@
-# chain b-a2bf485acea0 2026-10-21T17:13:12+08:00
+# chain b-0779e8984a06 2026-10-21T23:46:07+08:00 (list carried over from b-a2bf485acea0: that chain change did not touch any c1 read)
@@ -1,4 +1,4 @@
-# chain b-a2bf485acea0 2026-10-21T17:13:12+08:00
+# chain b-0779e8984a06 2026-10-21T23:46:44+08:00 (list carried over from b-a2bf485acea0: that chain change did not touch any uniform read)
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 004a663e75543071	0:192
 014cebfd5986868a	0:64
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100A83022A22000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "a2bf485acea0"
+#define DMSH_CHAIN "0779e8984a06"
@@ -1,4 +1,4 @@
-# chain b-0f1c080d3eca 2026-10-21T07:21:04+08:00 (list carried over from b-6631a1590f6d: that chain change did not touch any c1 read)
+# chain b-31f046bb1876 2026-10-21T22:57:30+08:00
 00096443b71e7fef
 001540067ae4cdf5
 002e5538fc05267e
@@ -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 "0f1c080d3eca"
+#define DMSH_CHAIN "31f046bb1876"
@@ -1 +1 @@
-# chain b-befa9a8c84ac 2026-10-21T07:21:04+08:00 (list carried over from b-22c0a0e0ae72: that chain change did not touch any c1 read)
+# chain b-4a9d20dadae5 2026-10-21T23:23:14+08:00 (list carried over from b-befa9a8c84ac: that chain change did not touch any c1 read)
@@ -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 "befa9a8c84ac"
+#define DMSH_CHAIN "4a9d20dadae5"
@@ -1,4 +1,4 @@
-# chain b-f5c63fb9b17f 2026-10-21T07:21:04+08:00 (list carried over from b-330a5957853d: that chain change did not touch any c1 read)
+# chain b-1cc65a8faa4d 2026-10-21T23:19:16+08:00
 35cbc686dd8461b8
 4872bec9a5567f8b
 759c14c6d029f531
@@ -1,4 +1,4 @@
-# chain b-f5c63fb9b17f 2026-10-21T12:00:55+08:00
+# chain b-1cc65a8faa4d 2026-10-21T23:19:34+08:00
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 0001a896bf9bd7cc	5:496 6:2256 12:432
 00027f7dc4646be2	5:496 6:2320 7:112 10:96 12:944 s4
@@ -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 "f5c63fb9b17f"
+#define DMSH_CHAIN "1cc65a8faa4d"
@@ -1,4 +1,4 @@
-# chain b-12de0799766c 2026-10-21T07:21:04+08:00 (list carried over from b-21d0d9d3a966: that chain change did not touch any c1 read)
+# chain b-975400df340b 2026-10-21T23:23:14+08:00 (list carried over from b-12de0799766c: that chain change did not touch any c1 read)
 000566ec19dc74f7
 000b591830b6fa73
 0027e56f73d8e10e
@@ -1,2 +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 "12de0799766c"
+#define DMSH_CHAIN "975400df340b"
@@ -1,2 +1,2 @@
-# chain b-70a5a3b5dbd2 2026-10-21T19:42:34+08:00
+# chain b-3e79ab909aaf 2026-10-22T00:21:36+08:00 (list carried over from b-70a5a3b5dbd2: that chain change did not touch any c1 read)
 e5ac49d4582a339c
@@ -1,4 +1,4 @@
-# chain b-70a5a3b5dbd2 2026-10-21T19:42:35+08:00
+# chain b-3e79ab909aaf 2026-10-22T00:21:36+08:00 (list carried over from b-70a5a3b5dbd2: that chain change did not touch any uniform read)
 # shader key, then reg:bytes for each uniform slot whose reads stay inside its first bytes, and s<reg> for each storage slot it never writes (cb_extent.py)
 008f89e94d37f17f	1:144
 00a30a0f46c385b7	1:512
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100F5D01DA3A000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "70a5a3b5dbd2"
+#define DMSH_CHAIN "3e79ab909aaf"
@@ -1,2 +1,2 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header, game 0100FD8022DAA000): the B shader library translation chain version this source tree expects (default options); the host loads only <library dir>/b-<V>/
-#define DMSH_CHAIN "349cb1de4c49"
+#define DMSH_CHAIN "a450cde97d1b"
@@ -1929,8 +1929,11 @@ void mtl_mem_stat(unsigned long long frame, unsigned long long rss_mb) {
    snprintf(avail, sizeof avail, "%llu MB", (unsigned long long)(os_proc_available_memory() >> 20));
 #endif
    MLOG("memory up to frame %llu: resident %llu MB, GPU allocated %llu MB, deferred-release queue %lu, residency set %lu, serial objects %u, memory pools %d (resident %d; peak %d / tracking table capacity %d; freed but alive %lu, %llu MB), shared textures %u (%u of them unreferenced; evicted with pools %llu, on storage takeover %llu in total), texture re-uploads %llu in total, per-frame arena blocks created %llu in total, finalized textures still referenced %u (small table, %llu put in total); headroom to the system memory limit %s", frame, rss_mb, gpu, nd, nr, ns, live, pres, peak, cap, dead_n, dead_mb, ashared, aorph, (unsigned long long)g_alias_evicted, (unsigned long long)g_alias_taken, (unsigned long long)g_pkg_reuploads, (unsigned long long)g_arena_made, nghost, (unsigned long long)ghost_puts, avail);
-    { void mtl_vc_stats(unsigned *, uint64_t *, uint64_t *); unsigned vn; uint64_t vh, vm; mtl_vc_stats(&vn, &vh, &vm);
-      MLOG("texture views by texture (vc_get): %u textures hold cached views; %llu views reused, %llu made in total", vn, (unsigned long long)vh, (unsigned long long)vm); }
+    { void mtl_vc_stats(unsigned *, uint64_t *, uint64_t *), mtl_tb2_stats(uint64_t *, uint64_t *, uint64_t *); uint64_t mtl_tex_mis_count(void); unsigned vn; uint64_t vh, vm, tb, tm, tf;
+      mtl_vc_stats(&vn, &vh, &vm); mtl_tb2_stats(&tb, &tm, &tf);
+      MLOG("texture views by texture (vc_get): %u textures hold cached views; %llu views reused, %llu made in total; texture buffers bound as 2D rows for wrapped 1D fetches %llu times, %llu made, %llu failed "
+           "(tbuf_as_2d); texture type mismatches no view bridges %llu in total (tex_mis_note)",
+           vn, (unsigned long long)vh, (unsigned long long)vm, (unsigned long long)tb, (unsigned long long)tm, (unsigned long long)tf, (unsigned long long)mtl_tex_mis_count()); }
    if (g_allocprof > 0) { char ab[256]; int o = 0; for (int i = 0; i < g_alloc_n && o < 200; i++) o += snprintf(ab + o, sizeof ab - o, " #%d %llu MB", i, (unsigned long long)(g_alloc4[i].allocatedSize >> 20));
        for (int i = 0; i < 3 && o < 240; i++) o += snprintf(ab + o, sizeof ab - o, " up%d %llu pres%d %llu", i, (unsigned long long)(g_up_alloc4[i].allocatedSize >> 20), i, (unsigned long long)(g_present_alloc4[i].allocatedSize >> 20));
        MLOG("command allocator usage:%s", ab);
@@ -2660,7 +2663,9 @@ static void shmap_load(void) {
@interface NvShader : NSObject { @public id<MTLFunction> fn, tf_fn; uint16_t tf_stride[4]; uint32_t cb_ext[16], sb_ro; nvn_refl rf; id<MTLArgumentEncoder> ae; NSUInteger arg_len; int stage, range_ro, 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; 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)
+    id<MTLLibrary> bflib; int bflib_tried; NSMutableDictionary<NSString *, id<MTLFunction>> *bffn;   // bflib: readable-tile version <key>_bf (merge B step 2)
+    uint64_t mis[4];   // mis: texture slots 0-255 already reported for a type mismatch (tex_mis_note)
+    NSString *ventry; id<MTLFunction> vfn; int vfn_state; } @end   // vertex: the entry point loaded (by the depth mode); vfn: the same entry of the shadow-input library <key>_bv (vin_fn; state 0 untried, 1 loaded, 2 failed)
@implementation NvShader
 - (void)dealloc { free(rf.c1); }
@end
@@ -2803,7 +2808,7 @@ static NvShader *shader_load(NSString *pc, int stage) {
        atomic_fetch_add(&g_vs_loaded, 1);
        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);
+    o->fn = [lib newFunctionWithName:entry]; o->stage = stage; strlcpy(o->key, pc.UTF8String, sizeof o->key); o->ventry = entry;
    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
@@ -2948,12 +2953,15 @@ static void treg_drop_ar(treg_t *r) { if (r->ar) { defer_obj_unref((__bridge id)
 // Only live textures are cached; entries go when the texture is finalized (mtl_tex_release) or its pool is (treg_drop_pool_views), and all of them past VC_MAX textures. Under g_rmu
 typedef struct { uint32_t pf, l0, nl, s0, ns; uint8_t tt, sw[4], pad[3]; } vc_key;
 enum { VC_WAYS = 4, VC_MAX = 16384 };
-typedef struct { const void *pool; unsigned next; vc_key k[VC_WAYS]; void *view[VC_WAYS]; } vc_set;   // view: retained
+typedef struct { const void *pool; unsigned next; vc_key k[VC_WAYS]; void *view[VC_WAYS];   // view: retained
+                 void *rows, *rows_of; uint64_t rows_rid, rows_n; } vc_set;   // rows: a texture buffer as 2D rows (tbuf_as_2d; retained, resident), rows_of / rows_n: the buffer texture it was made from and its element count
 static nvn_hmap g_vc = {.klen = 8};   // texture serial → vc_set
 static uint64_t g_vc_hit, g_vc_new;
+static void vc_rows_drop(vc_set *s) { if (s->rows) { defer_obj_release((__bridge id)s->rows); CFRelease(s->rows); s->rows = s->rows_of = NULL; s->rows_rid = s->rows_n = 0; } }
 static void vc_set_free(void *p) {
    vc_set *s = p;
    for (int i = 0; i < VC_WAYS; i++) if (s->view[i]) { defer_obj_unref((__bridge id)s->view[i]); CFRelease(s->view[i]); }
+    vc_rows_drop(s);
    free(s);
 }
 static id<MTLTexture> vc_get(uint64_t ser, const vc_key *k, id<MTLTexture> base) {
@@ -2965,15 +2973,19 @@ static id<MTLTexture> vc_get(uint64_t ser, const vc_key *k, id<MTLTexture> base)
    }
    return nil;
 }
-static void vc_put(uint64_t ser, const void *pool, const vc_key *k, id<MTLTexture> v) {
+static vc_set *vc_set_for(uint64_t ser, const void *pool) {   // the texture's entry, made if needed; NULL for a texture already finalized (resolved through g_ghost): nothing would drop the entry
    pthread_mutex_lock(&g_smu); int live = nvn_hm_get(&g_tex_live, &ser) != NULL; pthread_mutex_unlock(&g_smu);
-    if (!live) return;   // finalized already (resolved through g_ghost): nothing would drop the entry
+    if (!live) return NULL;
    vc_set *s = nvn_hm_get(&g_vc, &ser);
    if (!s) {
        if (g_vc.n >= VC_MAX) nvn_hm_clear(&g_vc, vc_set_free);
-        if (!(s = calloc(1, sizeof *s))) return;
+        if (!(s = calloc(1, sizeof *s))) return NULL;
        s->pool = pool; nvn_hm_put(&g_vc, &ser, s);
    }
+    return s;
+}
+static void vc_put(uint64_t ser, const void *pool, const vc_key *k, id<MTLTexture> v) {
+    vc_set *s = vc_set_for(ser, pool); if (!s) return;
    unsigned i = s->next++ % VC_WAYS;
    if (s->view[i]) { defer_obj_unref((__bridge id)s->view[i]); CFRelease(s->view[i]); }
    s->k[i] = *k; s->view[i] = (void *)CFBridgingRetain(v); g_vc_new++;
@@ -2987,6 +2999,38 @@ static void vc_drop_pool(const void *pool) {   // caller holds g_rmu
        else i++;
    }
 }
+// A texture buffer bound to a slot whose 1D integer fetches the translator made 2D reads of rows of NVN_TBUF_ROW texels (nvn_refl f1d; Program.cs RecordTexelFetch1D): a linear 2D texture of
+// such rows over the same pool memory, so the shader reads the buffer. Maxwell reads 1D textures and texture buffers with the same instruction and the descriptor decides; without a type
+// answer the translator read them as 2D, and Age of Calamity's light and decal lists (and Three Houses' lighting) read the dummy texture, zeros. Kept with the texture's views (vc_set),
+// checked against the buffer texture it was made from. Caller holds g_rmu
+static uint64_t g_tb2_binds, g_tb2_made, g_tb2_fail;
+static id<MTLTexture> tbuf_as_2d(int tid, id<MTLTexture> t, uint64_t *rid) {
+    if (tid <= 0 || tid >= g_ntreg || !g_treg[tid].ok) return nil;
+    const tex_t *tt = &g_treg[tid].t; uint64_t ser = tt->serial;
+    vc_set *s = nvn_hm_get(&g_vc, &ser);
+    if (s && s->rows && s->rows_of == (__bridge void *)t) {   // rows_of is not retained: a texture made since at the same address must also match the buffer the rows hold, their offset, format and element count
+        id<MTLTexture> rv = (__bridge id<MTLTexture>)s->rows;
+        if (rv.buffer == t.buffer && rv.bufferOffset == t.bufferOffset && rv.pixelFormat == t.pixelFormat && s->rows_n == t.width) { g_tb2_binds++; *rid = s->rows_rid; return rv; }
+    }
+    id<MTLBuffer> buf = t.buffer; int bw, bh; uint32_t bpb = (uint32_t)nvn_fmt_bpb(tt->d.format, &bw, &bh), w, h, row;
+    NSUInteger al = [g_dev minimumLinearTextureAlignmentForPixelFormat:t.pixelFormat];
+    if (!buf || bw != 1 || bh != 1 || t.bufferOffset % al || !nvn_tbuf_rows(t.width, bpb, (uint32_t)al, &w, &h, &row)) {
+        if (!g_tb2_fail++) MLOG("[RED] texture buffer (texture ID %d, %lu elements of format %#x at pool offset %lu) has no 2D-row form for a wrapped 1D fetch: the slot reads 0", tid, (unsigned long)t.width, tt->d.format, (unsigned long)t.bufferOffset);
+        return nil;
+    }
+    while (h > 1 && t.bufferOffset + (uint64_t)row * h > buf.length) h--;   // a last row running past the pool buffer: reads beyond it are undefined anyway
+    MTLTextureDescriptor *d = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:t.pixelFormat width:w height:h mipmapped:NO];
+    d.resourceOptions = buf.resourceOptions; d.usage = MTLTextureUsageShaderRead;
+    id<MTLTexture> v = [buf newTextureWithDescriptor:d offset:t.bufferOffset bytesPerRow:row];
+    if (!v) { if (!g_tb2_fail++) MLOG("[RED] cannot create the 2D-row form of a texture buffer: %ux%u format %lu row %u bytes at pool offset %lu", w, h, (unsigned long)t.pixelFormat, row, (unsigned long)t.bufferOffset); return nil; }
+    if (!g_tb2_made++) MLOG("texture buffer read by wrapped 1D fetches: bound as %ux%u rows of %u bytes (texture ID %d, %lu elements; first case)", w, h, row, tid, (unsigned long)t.width);
+    res_add(v);
+    *rid = v.gpuResourceID._impl; g_tb2_binds++;
+    if ((s = vc_set_for(ser, tt->pool))) { vc_rows_drop(s); s->rows = (void *)CFBridgingRetain(v); s->rows_of = (__bridge void *)t; s->rows_rid = *rid; s->rows_n = t.width; }
+    else defer_obj_release(v);   // a finalized texture's last reads: released once this submission is done
+    return v;
+}
+void mtl_tb2_stats(uint64_t *binds, uint64_t *made, uint64_t *fail) { *binds = g_tb2_binds; *made = g_tb2_made; *fail = g_tb2_fail; }
 static void treg_drop_pool_views(const void *pool) {
    pthread_mutex_lock(&g_rmu);
    vc_drop_pool(pool);
@@ -3132,8 +3176,7 @@ static id<MTLTexture> tex_fit(int tid, id<MTLTexture> t, MTLTextureType want, ui
    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; }
+    if (!ns || tid <= 0 || tid >= g_ntreg) return nil;   // no view bridges it: the caller reports it (tex_mis_note)
    treg_t *r = &g_treg[tid];
    if (r->alt && r->alt_of == (__bridge void *)t && r->alt_type == (uint8_t)want) { *rid = r->alt_rid; return (__bridge id<MTLTexture>)r->alt; }
    if (r->alt) { defer_obj_unref((__bridge id)r->alt); CFRelease(r->alt); r->alt = r->alt_of = NULL; r->alt_rid = 0; }
@@ -3142,6 +3185,23 @@ static id<MTLTexture> tex_fit(int tid, id<MTLTexture> t, MTLTextureType want, ui
    r->alt = (void *)CFBridgingRetain(v); r->alt_of = (__bridge void *)t; r->alt_type = (uint8_t)want; r->alt_rid = v.gpuResourceID._impl;
    *rid = r->alt_rid; return v;
 }
+// A texture type mismatch no view bridges, reported once per shader and texture slot with what is bound: one report per run used to hide that Hyrule Warriors: Age of Calamity has
+// thousands of shaders reading its light and decal lists, buffer textures, through slots translated as 2D. For a buffer texture: its element count and how far its offset in the
+// pool buffer is from the alignment a linear texture there needs. The total goes on the memory stats line. Caller holds g_rmu
+static uint64_t g_tex_mis_n, g_tex_mis_shown;
+uint64_t mtl_tex_mis_count(void) { return g_tex_mis_n; }
+static void tex_mis_note(NvShader *sh, int stage, int slot, int tid, id<MTLTexture> t, MTLTextureType want, const char *outcome) {
+    g_tex_mis_n++;
+    if (slot < 0 || slot >= 256 || (sh->mis[slot >> 6] >> (slot & 63) & 1)) return;
+    sh->mis[slot >> 6] |= 1ull << (slot & 63);
+    if (g_tex_mis_shown++ >= 64) return;
+    const tex_t *tt = tid > 0 && tid < g_ntreg && g_treg[tid].ok ? &g_treg[tid].t : NULL;
+    char bt[96] = "";
+    if (t.textureType == MTLTextureTypeTextureBuffer) snprintf(bt, sizeof bt, ", %lu elements at pool offset %lu, %lu past the linear-texture alignment", (unsigned long)t.width, (unsigned long)t.bufferOffset,
+                                                                (unsigned long)(t.bufferOffset % [g_dev minimumLinearTextureAlignmentForPixelFormat:t.pixelFormat]));
+    MLOG("[RED] texture type mismatch: shader %s (stage %d) declares slot %d as Metal type %lu, texture ID %d is type %lu (NVN target %d, format %#x, Metal format %lu%s); %s", sh->key, stage, slot,
+         (unsigned long)want, tid, (unsigned long)t.textureType, tt ? (int)tt->d.target : -1, tt ? (unsigned)tt->d.format : 0u, (unsigned long)t.pixelFormat, bt, outcome);
+}
 // Alpha from red: the fragment shaders the game profile's alpha_from_red lists read a single-channel render target whose channel mapping is exactly (R, 0, 0, 1) as (R, 0, 0, R).
 // A workaround for NintendoWare layout captures (BotW, Odyssey): the game rebuilds each capture on the same storage with its own mapping ((1, 1, 1, R) for an A8 capture), which was lost while
 // nvnTextureCompare ignored the mapping and left the creation-time (R, 0, 0, 1), alpha always 1 (hard-edged prompt boxes, white coin cards). Since that fix it no longer fires (docs/agents/04-graphics.md).
@@ -3505,15 +3565,57 @@ static void vconv_elem(int f, const uint8_t *p, float o[4]) {
            break; }
    }
 }
+// Shadow inputs instead of the conversion: a vertex shader with integer inputs also has a library <key>_bv (dmsh_vin.py) whose entry points declare each integer input again as a float4
+// and as the integer of the other sign, at Metal attribute indices 16-30, and OR the three into the input the body reads. A pipeline sends each attribute to the declaration of its
+// format's class (half / normalized formats to the float one, 8 / 16-bit integers to the one of their own sign) and the zero buffer to the others, and Metal's fetch then gives the
+// shader Maxwell's register bits: no CPU conversion, whatever mix of formats the pipeline has (Age of Calamity's soldiers feed Half2 beside RGBA8UI and Short2 to one shader; the
+// conversion was about a quarter of its replay thread in battle). Transform feedback captures (main0_tf) keep the conversion. DM_NVN_VCONV_CPU=1 converts on the CPU as before
+static int vin_cpu(void) { static int v = -1; if (v < 0) { const char *e = getenv("DM_NVN_VCONV_CPU"); v = e && !strcmp(e, "1"); if (v) MLOG("DM_NVN_VCONV_CPU=1: integer vertex inputs fed formats of another class are converted on the CPU (no shadow inputs)"); } return v; }
+static os_unfair_lock g_vin_lock = OS_UNFAIR_LOCK_INIT;
+static uint64_t g_vin_loaded, g_vin_failed, g_vin_psos;   // shadow-input libraries loaded / not usable, pipelines built with them
+static id<MTLFunction> vin_fn(NvShader *vs) {   // the shadow-input entry point, loaded on first need; a failure is final for the shader (it keeps the conversion)
+    if (!vs->rf.has_vin || vin_cpu()) return nil;
+    os_unfair_lock_lock(&g_vin_lock);
+    if (!vs->vfn_state) {
+        vs->vfn_state = 2;
+        size_t kl = strlen(vs->key); NSString *path = [NSString stringWithFormat:@"%@/%.*s_bv.metallib", g_msl_dir, (int)(kl > 2 ? kl - 2 : kl), vs->key];
+        NSError *e = nil; id<MTLLibrary> lib = lib_from_file(path, &e); id<MTLFunction> f = lib && vs->ventry ? [lib newFunctionWithName:vs->ventry] : nil;
+        if (f && f.functionConstantsDictionary.count) f = [lib newFunctionWithName:vs->ventry constantValues:[MTLFunctionConstantValues new] error:&e];   // as shader_load
+        const char *why = !lib ? "no library" : !f ? "no entry point" : NULL;
+        if (f && vs->rf.nb && [f newArgumentEncoderWithBufferIndex:vs->rf.abi].encodedLength != vs->arg_len) { why = "its argument buffer differs from the base library's"; f = nil; }
+        if (f) { vs->vfn = f; vs->vfn_state = 1; g_vin_loaded++; }
+        else if (g_vin_failed++ < 8) MLOG("[RED] shadow-input library of %s not usable (%s%s%s): its integer inputs keep the CPU conversion", vs->key, why, e ? ": " : "", e ? e.localizedDescription.UTF8String : "");
+    }
+    id<MTLFunction> f = vs->vfn_state == 1 ? vs->vfn : nil;
+    os_unfair_lock_unlock(&g_vin_lock);
+    return f;
+}
+static inline int vin_fmt(const pso_key *k, int a) { return a >= 16 ? 0 : k ? k->attr[a].fmt : a < S.nattr ? S.attr[a].fmt : 0; }   // as pso_key_fill fills the key
+// The component type input v takes its attribute as: the class of the attribute's format when a shadow of that class exists (on), else the declared type
+static int vin_ct(const nvn_refl *rf, int v, int f, int on) {
+    int ct = rf->v[v].ct, ib, is;
+    if (!on || ct == NVN_CT_FLOAT) return ct;
+    int want = vconv_ncomp(f) ? NVN_CT_FLOAT : vconv_int(f, &ib, &is) ? (is ? NVN_CT_INT : NVN_CT_UINT) : ct;
+    return want != ct && rf->v[v].sh[want] != 0xff ? want : ct;
+}
+// Whether this pipeline (k, or the current state at bind time: both give the same answer) uses the shadow inputs: some input takes a shadow and the library loads. Asked by pso_desc
+// and vconv_bind alike, so the converted attributes and the pipeline's buffer slots agree
+static int vin_on(NvShader *vs, const pso_key *k, int tf) {
+    const nvn_refl *rf = &vs->rf;
+    if (tf || !rf->has_vin || vin_cpu()) return 0;
+    int need = 0; for (int v = 0; v < rf->nv && !need; v++) need = vin_ct(rf, v, vin_fmt(k, rf->v[v].nvn), 1) != rf->v[v].ct;
+    return need && vin_fn(vs) != nil;
+}
 // A float input fed an 8 / 16-bit integer format takes the same 32-bit register values through a Float4 attribute, which Metal passes through bit for bit: Monster Hunter Rise's
 // skinned vertex shaders (40609cd5c1443df4, 1369a37113da0012) read RGBA8UI attribute 5 as flags (as_type<int>(x) & 2 ...) and also as a float, so the build leaves the input float
-static int vconv_wanted(const nvn_refl *rf, int v, const pso_key *k) {   // does Metal input v of this pipeline take converted data
+static int vconv_wanted(const nvn_refl *rf, int v, const pso_key *k, int on) {   // does Metal input v of this pipeline take converted data (on: it uses the shadow inputs, vin_on)
    int a = rf->v[v].nvn, ct = rf->v[v].ct, ib, is;
    if (a >= 16) return 0;
    int f = k ? k->attr[a].fmt : S.attr[a].fmt;
    if (vconv_i2f(f, &ib, &is)) return 1;   // whatever the input type
    if (ct == NVN_CT_FLOAT) return vconv_int(f, &ib, &is) > 0;
    if (ct != NVN_CT_UINT && ct != NVN_CT_INT) return 0;
+    if (vin_ct(rf, v, f, on) != ct) return 0;   // a shadow input takes it
    return vconv_ncomp(f) || (vconv_int(f, &ib, &is) && is != (ct == NVN_CT_INT));
 }
 // nv elements from src (stride bytes apart) into d, four 32-bit values each. The 8 / 16-bit integer and the half formats (Age of Calamity's soldiers: Short2, Half2) have loops
@@ -3590,17 +3692,18 @@ static int vc_draw_range(const p_draw *dw, const p_draw_idx *di, int64_t *lo, in
 }
 static void vc_stat(void) {   // every 3600 submissions
    if (g_sub_serial - g_vc_stat < 3600) return;
-    if (g_vc_stat) MLOG("vertex conversion: %.0f vertices per submission (the bound ranges: %.0f), %u index ranges kept%s", g_vc_conv / 3600.0, g_vc_whole / 3600.0, g_ir.n,
-        vc_check() ? [NSString stringWithFormat:@"; checked %llu, %llu differed", (unsigned long long)g_vc_checks, (unsigned long long)g_vc_bad].UTF8String : "");
+    if (g_vc_stat) MLOG("vertex conversion: %.0f vertices per submission (the bound ranges: %.0f), %u index ranges kept%s; shadow-input libraries loaded %llu (not usable %llu), pipelines built with them %llu", g_vc_conv / 3600.0, g_vc_whole / 3600.0, g_ir.n,
+        vc_check() ? [NSString stringWithFormat:@"; checked %llu, %llu differed", (unsigned long long)g_vc_checks, (unsigned long long)g_vc_bad].UTF8String : "",
+        (unsigned long long)g_vin_loaded, (unsigned long long)g_vin_failed, (unsigned long long)g_vin_psos);
    g_vc_stat = g_sub_serial; g_vc_conv = g_vc_whole = 0;
 }
 // Per submission, one conversion per (range, format, layout): KTGL draws the same mesh several times a frame (depth, shadow, G-buffer)
 static struct { uint64_t addr, size, serial, v0, nv; uint32_t fmt, off, stride; uint64_t gpu; } g_vconv[64]; static int g_vconv_next;
-static void vconv_bind(id<MTL4RenderCommandEncoder> en, const nvn_refl *rf, const p_draw *dw, const p_draw_idx *di) {
-    int slot = VCONV_SLOT0, vc = vc_on(), rk = -1; int64_t rlo = 0, rhi = 0;   // rk: the draw's vertex range is known (-1 = not looked up yet)
-    if (vc) vc_stat();
+static void vconv_bind(id<MTL4RenderCommandEncoder> en, NvShader *vs, int tf, const p_draw *dw, const p_draw_idx *di) {   // tf: a transform feedback capture (no shadow inputs)
+    const nvn_refl *rf = &vs->rf; int slot = VCONV_SLOT0, vc = vc_on(), rk = -1, on = vin_on(vs, NULL, tf); int64_t rlo = 0, rhi = 0;   // rk: the draw's vertex range is known (-1 = not looked up yet)
+    vc_stat();
    for (int v = 0; v < rf->nv && slot < VCONV_SLOT0 + VCONV_MAX; v++) {
-        if (!vconv_wanted(rf, v, NULL) || rf->v[v].nvn >= S.nattr) continue;
+        if (!vconv_wanted(rf, v, NULL, on) || rf->v[v].nvn >= S.nattr) continue;
        int a = rf->v[v].nvn, st = S.attr[a].stream, f = S.attr[a].fmt; uint32_t off = S.attr[a].offset, stride = st < 16 ? S.stream[st].stride : 0;
        uint64_t addr = st < 16 ? S.vbuf[st].address : 0, size = st < 16 ? S.vbuf[st].size : 0, gpu = 0;
        uint64_t es = (uint64_t)vconv_size(f), v0 = 0, nv = stride == 0 ? 1 : size >= off + es ? (size - off - es) / stride + 1 : 0, whole = nv;
@@ -3711,21 +3814,21 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS
    }
    d.depthAttachmentPixelFormat = (MTLPixelFormat)k->dfmt; d.stencilAttachmentPixelFormat = (MTLPixelFormat)k->sfmt;
    MTLVertexDescriptor *vd = [MTLVertexDescriptor vertexDescriptor];
-    const nvn_refl *rf = &vs->rf; int nconv = 0;
+    const nvn_refl *rf = &vs->rf; int nconv = 0, on = vin_on(vs, k, k->fsfn == TF_MARK);
+    if (on) { d.vertexFunction = vs->vfn; g_vin_psos++; }   // the shadow-input entry point (see vin_on)
+    #define ZERO_ATTR(i, ct_) do { vd.attributes[i].format = (ct_) == NVN_CT_UINT ? MTLVertexFormatUInt4 : (ct_) == NVN_CT_INT ? MTLVertexFormatInt4 : MTLVertexFormatFloat4; vd.attributes[i].offset = 0; vd.attributes[i].bufferIndex = 30; \
+                                  vd.layouts[30].stride = 16; vd.layouts[30].stepFunction = MTLVertexStepFunctionConstant; vd.layouts[30].stepRate = 0; } while (0)   // reads the zero buffer
    for (int v = 0; v < rf->nv; v++) {
        int loc = rf->v[v].loc, a = rf->v[v].nvn;   // Metal slot loc ← NVN attribute a
        MTLVertexFormat vf = a < 16 ? vfmt(k->attr[a].fmt) : MTLVertexFormatInvalid;
+        // With the shadow inputs the attribute goes to the declaration of its format's class (at) and the input's other declarations read the zero buffer
+        int ct = on ? vin_ct(rf, v, vin_fmt(k, a), 1) : rf->v[v].ct, at = ct == rf->v[v].ct ? loc : rf->v[v].sh[ct];
+        if (on) { if (at != loc) ZERO_ATTR(loc, rf->v[v].ct); for (int q = 0; q < 3; q++) if (rf->v[v].sh[q] != 0xff && rf->v[v].sh[q] != at) ZERO_ATTR(rf->v[v].sh[q], q); }
        // When a >= S.nattr, k.attr[a] keeps the 0 from memset, vfmt(0) is Invalid, the same fallback branch is taken and the RED_ONCE condition of fmt being non-0 does not hold.
        if (a < 16 && k->attr[a].fmt && vf == MTLVertexFormatInvalid) RED_ONCE(15, "vertex attribute format %#x not mapped", k->attr[a].fmt);
-        if (vf == MTLVertexFormatInvalid) {
-            int ct = rf->v[v].ct;
-            vd.attributes[loc].format = ct == NVN_CT_UINT ? MTLVertexFormatUInt4 : ct == NVN_CT_INT ? MTLVertexFormatInt4 : MTLVertexFormatFloat4;
-            vd.attributes[loc].offset = 0; vd.attributes[loc].bufferIndex = 30;
-            vd.layouts[30].stride = 16; vd.layouts[30].stepFunction = MTLVertexStepFunctionConstant; vd.layouts[30].stepRate = 0;
-            continue;
-        }
+        if (vf == MTLVertexFormatInvalid) { ZERO_ATTR(at, ct); continue; }
        int st = k->attr[a].stream;
-        if (vconv_wanted(rf, v, k) && nconv < VCONV_MAX) {   // converted on the CPU at draw time (vconv_bind): the register's 32-bit values, 4 components, its own slot
+        if (vconv_wanted(rf, v, k, on) && nconv < VCONV_MAX) {   // converted on the CPU at draw time (vconv_bind): the register's 32-bit values, 4 components, its own slot
            int slot = VCONV_SLOT0 + nconv++;
            vd.attributes[loc].format = rf->v[v].ct == NVN_CT_UINT ? MTLVertexFormatUInt4 : rf->v[v].ct == NVN_CT_INT ? MTLVertexFormatInt4 : MTLVertexFormatFloat4; vd.attributes[loc].offset = 0; vd.attributes[loc].bufferIndex = slot;
            vd.layouts[slot].stride = 16;
@@ -3733,12 +3836,13 @@ static MTLRenderPipelineDescriptor *pso_desc(const pso_key *k, NvShader *vs, NvS
            vd.layouts[slot].stepRate = k->stream[st].stride == 0 ? 0 : k->stream[st].divisor ? k->stream[st].divisor : 1;
            continue;
        }
-        vf = vfmt_for_ct(vf, rf->v[v].ct, k->attr[a].fmt);
-        vd.attributes[loc].format = vf; vd.attributes[loc].offset = k->attr[a].offset; vd.attributes[loc].bufferIndex = VB_SLOT(st);
+        vf = vfmt_for_ct(vf, ct, k->attr[a].fmt);
+        vd.attributes[at].format = vf; vd.attributes[at].offset = k->attr[a].offset; vd.attributes[at].bufferIndex = VB_SLOT(st);
        vd.layouts[VB_SLOT(st)].stride = k->stream[st].stride ? k->stream[st].stride : 4;
        vd.layouts[VB_SLOT(st)].stepFunction = k->stream[st].stride == 0 ? MTLVertexStepFunctionConstant : k->stream[st].divisor ? MTLVertexStepFunctionPerInstance : MTLVertexStepFunctionPerVertex;
        vd.layouts[VB_SLOT(st)].stepRate = k->stream[st].stride == 0 ? 0 : k->stream[st].divisor ? k->stream[st].divisor : 1;
    }
+    #undef ZERO_ATTR
    d.vertexDescriptor = vd;
    return d;
 }
@@ -4107,9 +4211,11 @@ static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, i
                       (unsigned long)t.pixelFormat, (unsigned long)t.width, (unsigned long)t.height, (unsigned long)t.depth,
                       (unsigned long)(t.parentTexture ? t.parentRelativeLevel : 0), (unsigned long)t.mipmapLevelCount,
                       (unsigned long)(t.parentTexture ? t.parentRelativeSlice : 0), (unsigned long)t.arrayLength, (unsigned long)t.textureType, t ? (int)t.swizzle.red : 0, t ? (int)t.swizzle.green : 0, t ? (int)t.swizzle.blue : 0, t ? (int)t.swizzle.alpha : 0, (__bridge void *)(t.parentTexture ? t.parentTexture : t));
+            if (t && b->f1d && t.textureType == MTLTextureTypeTextureBuffer && (MTLTextureType)b->dtt == MTLTextureType2D) { uint64_t r2 = 0; id<MTLTexture> v = tbuf_as_2d(tid, t, &r2); if (v) { t = v; trid = r2; } }   // see tbuf_as_2d
            if (t && t.textureType != (MTLTextureType)b->dtt) { uint64_t r2 = 0; id<MTLTexture> v = tex_fit(tid, t, (MTLTextureType)b->dtt, &r2); if (v) { t = v; trid = r2; }
                // no view fits, and one side is a buffer texture (Octopath Traveler 0: a buffer texture left in a slot a shader reads as a cube): reads of it are undefined, so the slot reads 0
-                else if (t.textureType == MTLTextureTypeTextureBuffer || (MTLTextureType)b->dtt == MTLTextureTypeTextureBuffer) t = nil; }
+                else if (t.textureType == MTLTextureTypeTextureBuffer || (MTLTextureType)b->dtt == MTLTextureTypeTextureBuffer) { tex_mis_note(sh, nvn_stage, b->kind == NVN_RB_TEX ? reg : -1, tid, t, (MTLTextureType)b->dtt, "a buffer texture: the slot reads 0 (dummy)"); t = nil; }
+                else tex_mis_note(sh, nvn_stage, b->kind == NVN_RB_TEX ? reg : -1, tid, t, (MTLTextureType)b->dtt, "bound as is"); }
            if (t && nvn_stage == 1 && afr_want(sh)) { uint64_t r2 = 0; id<MTLTexture> v = tex_alpha_red(tid, t, &r2); if (v) { t = v; trid = r2; } }   // see tex_alpha_red
            if (!t) { t = dummy_tex(b->dtt, b->dpf); trid = g_dummy_tex_rid[b->dtt & 7][b->dpf & 3]; }
            slots[aid] = trid;
@@ -4138,7 +4244,7 @@ static int encode_args_(id<MTLCommandEncoder> en, NvShader *sh, int nvn_stage, i
        } else if (b->kind == NVN_RB_TBUF) {   // Typed buffer: the buffer texture (target BUFFER) in texture slot r
            int tid = reg < NTEX ? (int)(S.tex[nvn_stage][reg] & 0xfffff) : 0;
            void *tv = NULL; uint64_t trid = tid ? tex_fast_rid(tid, &tv) : 0; __unsafe_unretained id<MTLTexture> t = (__bridge id<MTLTexture>)tv;
-            if (t && t.textureType != MTLTextureTypeTextureBuffer) { RED_ONCE(40, "typed buffer slot is bound to something other than a buffer texture (texture ID %d type %lu)", tid, (unsigned long)t.textureType); t = nil; }
+            if (t && t.textureType != MTLTextureTypeTextureBuffer) { tex_mis_note(sh, nvn_stage, reg, tid, t, MTLTextureTypeTextureBuffer, "a buffer slot: it reads 0 (dummy)"); t = nil; }
            if (bl) PD("B %llu st%d tbuf reg%d tid=%d %s ser=%llu", g_stat_pass, nvn_stage, reg, tid, t ? "real" : "dummy", t && tid < g_ntreg ? g_treg[tid].t.serial : 0ull);
            if (!t) { t = dummy_tbuf(b->dpf); trid = g_dummy_tbuf_rid[b->dpf & 3]; }
            slots[aid] = trid;
@@ -6080,7 +6186,7 @@ static void tf_draw(NvProg *pr, const p_draw *dw, const p_draw_idx *di) {
    id<MTL4RenderCommandEncoder> en = enc_r_begin(d);
    enc_r_pso(en, ps);
    bind_vbufs(en);
-    vconv_bind(en, &vs->rf, dw, NULL);
+    vconv_bind(en, vs, 1, dw, NULL);
    args_dirty_all(); encode_args(en, vs, 0, 1);
    for (int b = 0; b < 4; b++) if (vs->tf_stride[b]) {   // vertex v of the draw (first + i) stores at v × stride: the pointer starts first × stride before the write offset
        enc_r_vbuf(en, TF_SLOT + b, gpu[b] - (uint64_t)dw->first * vs->tf_stride[b]);
@@ -6302,7 +6408,7 @@ static void draw(const p_draw *dw, const p_draw_idx *di) {
      if (!(g_ed.ok & ED_CLIP) || g_ed.clip != cl) { enc_r_depth_clip(en, cl); g_ed.clip = cl; g_ed.ok |= ED_CLIP; }
      if (!(g_ed.ok & ED_BLEND) || memcmp(g_ed.blend, S.blend_color, sizeof g_ed.blend)) { enc_r_blend_color(en, S.blend_color[0], S.blend_color[1], S.blend_color[2], S.blend_color[3]); memcpy(g_ed.blend, S.blend_color, sizeof g_ed.blend); g_ed.ok |= ED_BLEND; } }
    bind_vbufs(en);
-    if (pr->sh[0]) vconv_bind(en, &pr->sh[0]->rf, dw, di);   // integer inputs fed float-type formats or 8 / 16-bit integers of the other sign (see vconv_ncomp, vconv_int)
+    if (pr->sh[0]) vconv_bind(en, pr->sh[0], 0, dw, di);   // integer inputs fed float-type formats or 8 / 16-bit integers of the other sign (see vconv_ncomp, vconv_int)
    g_kd_force = kd_draw(pr, dw, di);
    if (pr->sh[0]) encode_args(en, pr->sh[0], 0, 1);
    if (pr->sh[1]) encode_args(en, pr->sh[1], 1, 0);
@@ -25,11 +25,14 @@ typedef struct {
    uint16_t reg, aid;       // D3D register number, index within the argument buffer
    uint16_t off;            // TEXU / SAMPU: byte offset of the handle within the uniform buffer; SBUFU: of the 64-bit storage buffer address
    uint8_t n;               // TEXD / SAMPD: array length (argument buffer entries aid .. aid + n - 1)
+    uint8_t f1d;             // TEX: read by 1D integer fetches as 2D rows of 16384 texels (dmsh_build.py "fetch1d"): a buffer texture here is bound as such rows (nvn_mtl.m tbuf_as_2d)
 } nvn_rbind;
 typedef struct {
    int nb, abi, nv;         // Binding count, argument buffer slot, vertex input count
    nvn_rbind b[NVN_REFL_MAXB];
-    struct { uint8_t loc, ct, nvn; } v[NVN_REFL_MAXV];   // loc: Metal vertex attribute slot (DXBC input register); nvn: NVN vertex attribute number (the n of HLSL semantic ATTRn)
+    struct { uint8_t loc, ct, nvn, sh[3]; } v[NVN_REFL_MAXV];   // loc: Metal vertex attribute slot (DXBC input register); nvn: NVN vertex attribute number (the n of HLSL semantic ATTRn);
+                                                                // sh[ct]: Metal attribute index of the input's shadow of that component type in the <key>_bv library (0xff = none; dmsh_vin.py)
+    uint8_t has_vin;         // has a shadow-input library <key>_bv (nvn_mtl.m vin_fn)
    uint16_t tg[3];          // Compute thread group size (none = 0)
    uint8_t route_b, has_c1; // Route B; has a c1 constant segment
    uint8_t bindless;        // Has TEXU / SAMPU bindings (which texture is read is known only at replay time)
@@ -54,7 +54,7 @@ int nvn_refl_from_json(NSDictionary *j, nvn_refl *o, char *err, int errlen) {
            if ([k isEqualToString:@"cbuf"]) r->kind = NVN_RB_CBUF;
            else if ([k isEqualToString:@"samp"]) r->kind = NVN_RB_SAMP;
            else if ([k isEqualToString:@"samp_tex"]) r->kind = NVN_RB_SAMPT;
-            else if ([k isEqualToString:@"tex"]) { r->kind = NVN_RB_TEX; dummy_kind(b[@"msl_type"], &r->dtt, &r->dpf); }
+            else if ([k isEqualToString:@"tex"]) { r->kind = NVN_RB_TEX; dummy_kind(b[@"msl_type"], &r->dtt, &r->dpf); r->f1d = [b[@"fetch1d"] boolValue]; }
            else if ([k isEqualToString:@"tex_ubo"]) { r->kind = NVN_RB_TEXU; dummy_kind(b[@"msl_type"], &r->dtt, &r->dpf); r->off = (uint16_t)[b[@"off"] intValue]; o->bindless = 1; }
            else if ([k isEqualToString:@"samp_ubo"]) { r->kind = NVN_RB_SAMPU; r->off = (uint16_t)[b[@"off"] intValue]; o->bindless = 1; }
            else if ([k isEqualToString:@"samp_id"]) r->kind = NVN_RB_SAMPI;
@@ -88,8 +88,23 @@ int nvn_refl_from_json(NSDictionary *j, nvn_refl *o, char *err, int errlen) {
            if (loc < 0 || loc >= NVN_REFL_MAXV || o->nv >= NVN_REFL_MAXV) FAIL("B vertex input %d out of range", loc);
            o->v[o->nv].loc = o->v[o->nv].nvn = (uint8_t)loc;
            o->v[o->nv].ct = [t hasPrefix:@"uint"] ? NVN_CT_UINT : [t hasPrefix:@"int"] ? NVN_CT_INT : NVN_CT_FLOAT;
+            memset(o->v[o->nv].sh, 0xff, sizeof o->v[o->nv].sh);
            o->nv++;
        }
+        NSDictionary *vin = j[@"vin"];   // shadow inputs of the <key>_bv library (dmsh_vin.py): per input, the Metal attribute index of its float / other-sign declaration
+        if ([vin isKindOfClass:[NSDictionary class]]) {
+            uint32_t used = 0;
+            for (NSDictionary *e in vin[@"shadow"]) {
+                int loc = [e[@"location"] intValue], v = 0;
+                while (v < o->nv && o->v[v].loc != loc) v++;
+                if (v == o->nv || o->v[v].ct == NVN_CT_FLOAT) FAIL("shadow input for location %d, which is not an integer input", loc);
+                const struct { NSString *k; int ct; } T[3] = { { @"float", NVN_CT_FLOAT }, { @"sint", NVN_CT_INT }, { @"uint", NVN_CT_UINT } };
+                for (int q = 0; q < 3; q++) { id x = e[T[q].k]; if (!x) continue; int idx = [x intValue];
+                    if (idx < 16 || idx > 30 || used >> idx & 1 || T[q].ct == o->v[v].ct) FAIL("shadow input %s of location %d at attribute %d: out of 16-30, used twice or of the input's own type", T[q].k.UTF8String, loc, idx);
+                    used |= 1u << idx; o->v[v].sh[T[q].ct] = (uint8_t)idx; }
+            }
+            o->has_vin = used != 0;
+        }
        NSDictionary *ra = j[@"raster"];
        o->pos_attr = -1; o->may_discard = 1; o->writes_smask = 1;
        if ([ra isKindOfClass:[NSDictionary class]]) {
@@ -160,3 +160,13 @@ void nvn_tex_level_tiles(const nvn_texdesc *d, unsigned level, unsigned *tileH,
    *tileH = adj_tile(L.tileH, 8, lh);
    *tileD = adj_tile(L.tileD, 1, ld);
 }
+
+int nvn_tbuf_rows(uint64_t n, uint32_t bpb, uint32_t align, uint32_t *w, uint32_t *h, uint32_t *row) {
+    if (!n || !bpb) return 0;
+    if (!align) align = 1;
+    *w = n < NVN_TBUF_ROW ? (uint32_t)n : NVN_TBUF_ROW;
+    *h = (uint32_t)((n + NVN_TBUF_ROW - 1) / NVN_TBUF_ROW);
+    uint64_t r = (uint64_t)*w * bpb; r = (r + align - 1) / align * align;
+    *row = (uint32_t)r;
+    return *h == 1 || r == (uint64_t)NVN_TBUF_ROW * bpb;   // several rows must lie back to back, as the wrapped address assumes
+}
@@ -20,3 +20,7 @@ uint32_t nvn_tex_level_gobs(const nvn_texdesc *d, unsigned level);   // GOB colu
 int nvn_fmt_bpb(int format, int *bw, int *bh);   // Bytes per block (0 = unknown format), block width and height
 // Block height / block depth actually used by level `level` (log2 of the GOB count): the per-level block height shrinks with that level's row count (block rows), and deswizzling uses it level by level
 void nvn_tex_level_tiles(const nvn_texdesc *d, unsigned level, unsigned *tileH, unsigned *tileD);
+// A texture buffer of n elements of bpb bytes read as 2D rows of NVN_TBUF_ROW texels (the translator's wrapped 1D fetch, x & 0x3FFF / x >> 14): width, height and bytes per row of the
+// linear 2D texture over the same memory. One row is n texels, its byte length rounded up to align (the linear-texture row alignment); more rows are full 16384-texel rows. 0 when n or bpb is 0
+#define NVN_TBUF_ROW 16384u
+int nvn_tbuf_rows(uint64_t n, uint32_t bpb, uint32_t align, uint32_t *w, uint32_t *h, uint32_t *row);
@@ -46,3 +46,7 @@ t_vbuf
 t_vbuf.dSYM/
 t_dirty_race
 t_dirty_race.dSYM/
+t_vin
+t_vin.dSYM/
+t_tbuf2d
+t_tbuf2d.dSYM/
@@ -24,6 +24,10 @@ t_dirty_race: t_dirty_race.c ../nvn/nvn_dirty.c ../../host/crash.c ../nvn/nvn_di
 	clang $(CFLAGS) -o $@ t_dirty_race.c ../../host/crash.c
 t_builtin: t_builtin.m ../nvn/nvn_builtin.h
 	clang -O1 -g -Wall -fobjc-arc -o $@ t_builtin.m -framework Metal -framework Foundation
+t_vin: t_vin.m
+	clang -O1 -g -Wall -fobjc-arc -o $@ t_vin.m -framework Metal -framework Foundation
+t_tbuf2d: t_tbuf2d.m ../nvn/nvn_tex.c ../nvn/nvn_tex.h
+	clang -O1 -g -Wall -fobjc-arc -o $@ t_tbuf2d.m ../nvn/nvn_tex.c -framework Metal -framework Foundation
 t_deswz_gpu: t_deswz_gpu.m ../tex/nvn_deswizzle.c ../tex/nvn_deswizzle_gpu.m ../tex/nvn_deswizzle.h ../tex/nvn_deswizzle_gpu.h ../nvn/nvn_tex.c ../nvn/nvn_tex.h
 	clang $(CFLAGS) -fobjc-arc -o $@ t_deswz_gpu.m ../tex/nvn_deswizzle.c ../tex/nvn_deswizzle_gpu.m ../nvn/nvn_tex.c -framework Metal -framework Foundation
 t_refl: t_refl.m ../nvn/nvn_refl.m ../nvn/nvn_refl.h $(GAME_DIR)/dmsh_chain.h
@@ -66,11 +70,13 @@ t_pso_skey: t_pso_skey.m ../nvn/nvn_pso_skey.h ../nvn/nvn_obj.h
 	clang $(CFLAGS) -fobjc-arc -o $@ t_pso_skey.m -framework Foundation
 t_pfv: t_pfv.m ../../host/crash.c $(GFX) $(GFX_H)
 	clang $(CFLAGS) -fobjc-arc -o $@ t_pfv.m ../../host/crash.c $(filter-out ../nvn/nvn_mtl.m ../nvn/nvn_entry.S ../nvn/nvn_thunks.S,$(GFX)) $(FW)
-run: t_nvn t_dirty_slots t_dirty_race t_builtin t_deswz_gpu t_pace t_fstep t_refl t_hmap t_cover t_fuse_check t_present_crop t_reorder_chain t_defer_chain t_zskip t_xv t_ub t_vbuf t_mip_extent t_rt_range_shape t_rt_range_ro t_rt_range t_rt_depth_ro t_pso_skey t_pfv
+run: t_nvn t_dirty_slots t_dirty_race t_builtin t_vin t_tbuf2d t_deswz_gpu t_pace t_fstep t_refl t_hmap t_cover t_fuse_check t_present_crop t_reorder_chain t_defer_chain t_zskip t_xv t_ub t_vbuf t_mip_extent t_rt_range_shape t_rt_range_ro t_rt_range t_rt_depth_ro t_pso_skey t_pfv
 	./t_nvn
 	./t_dirty_slots
 	./t_dirty_race
 	./t_builtin
+	./t_vin
+	./t_tbuf2d
 	./t_deswz_gpu
 	./t_pace
 	./t_fstep
@@ -436,6 +436,18 @@ int main(void) {
    EXPECT("unimplemented function marked as reported ([RED]) after a call", 1, g_unimpl_seen[ui] > 0);
    EXPECT("unimplemented function count 2", 2, nvn_slot[ui].cnt);

+    // 8. A texture buffer as 2D rows for wrapped 1D fetches (x & 0x3FFF, x >> 14; nvn_mtl.m tbuf_as_2d): full rows back to back, one short row rounded to the alignment
+    {
+        uint32_t w, h, row;
+        EXPECT("256 R32 elements: one row", 1, nvn_tbuf_rows(256, 4, 64, &w, &h, &row) && w == 256 && h == 1 && row == 1024);
+        EXPECT("10 R32 elements: one row rounded up to the 64-byte alignment", 1, nvn_tbuf_rows(10, 4, 64, &w, &h, &row) && w == 10 && h == 1 && row == 64);
+        EXPECT("16384 RGBA32F elements: exactly one full row", 1, nvn_tbuf_rows(16384, 16, 256, &w, &h, &row) && w == 16384 && h == 1 && row == 262144);
+        EXPECT("16385 R32 elements: two rows of 16384", 1, nvn_tbuf_rows(16385, 4, 64, &w, &h, &row) && w == 16384 && h == 2 && row == 65536);
+        EXPECT("1,048,576 R32 elements: 64 rows", 1, nvn_tbuf_rows(1048576, 4, 64, &w, &h, &row) && w == 16384 && h == 64 && row == 65536);
+        EXPECT("no elements: no rows", 0, nvn_tbuf_rows(0, 4, 64, &w, &h, &row));
+        EXPECT("several rows whose alignment would break the wrapped address: refused", 0, nvn_tbuf_rows(20000, 4, 3, &w, &h, &row));
+    }
+
    printf("== %d checks, %d failed ==\n", g_n, g_fail);
    return g_fail != 0;
 }
@@ -0,0 +1,68 @@
+// A texture buffer read through a slot the translator made a 2D read of wrapped 1D fetches (InstEmitTexture.cs nxnative: x & 0x3FFF, x >> 14): nvn_mtl.m tbuf_as_2d binds a linear
+// 2D texture of 16384-texel rows (nvn_tbuf_rows) over the same buffer bytes. Checks, for element counts from 1 to 1,048,576, pool offsets at the linear-texture alignment and two
+// formats, that every element read through the alias at the wrapped address equals the texture_buffer view's read and the memory itself; and that a one-row 2D texture (how the
+// host creates a 1D texture) reads the same at x & 0x3FFF, x >> 14 as at (x, 0)
+#import <Metal/Metal.h>
+#include <stdio.h>
+#include <string.h>
+#include "../nvn/nvn_tex.h"
+
+static const char *k_msl =
+    "#include <metal_stdlib>\nusing namespace metal;\n"
+    "kernel void ku(texture_buffer<uint> tb [[texture(0)]], texture2d<uint> t2 [[texture(1)]], device uint4 *o [[buffer(0)]], constant uint &n [[buffer(1)]], uint i [[thread_position_in_grid]])\n"
+    "{ if (i >= n) return; int x = int(i); o[2 * i] = tb.read(i); o[2 * i + 1] = t2.read(uint2(int2(x & 16383, int(uint(x) >> 14u))), 0); }\n"
+    "kernel void kf(texture_buffer<float> tb [[texture(0)]], texture2d<float> t2 [[texture(1)]], device uint4 *o [[buffer(0)]], constant uint &n [[buffer(1)]], uint i [[thread_position_in_grid]])\n"
+    "{ if (i >= n) return; int x = int(i); o[2 * i] = as_type<uint4>(tb.read(i)); o[2 * i + 1] = as_type<uint4>(t2.read(uint2(int2(x & 16383, int(uint(x) >> 14u))), 0)); }\n"
+    "kernel void k1(texture2d<uint> t2 [[texture(1)]], device uint4 *o [[buffer(0)]], constant uint &n [[buffer(1)]], uint i [[thread_position_in_grid]])\n"
+    "{ if (i >= n) return; int x = int(i); o[2 * i] = t2.read(uint2(i, 0), 0); o[2 * i + 1] = t2.read(uint2(int2(x & 16383, int(uint(x) >> 14u))), 0); }\n";
+
+static int g_n, g_fail;
+static void check(const char *what, int ok, const char *detail) { g_n++; if (!ok) g_fail++; printf("%-70s %s %s\n", what, ok ? "✓" : "✗", ok ? "" : detail); }
+
+int main(void) {
+    id<MTLDevice> dev = MTLCreateSystemDefaultDevice(); NSError *err = nil;
+    id<MTLLibrary> lib = [dev newLibraryWithSource:@(k_msl) options:nil error:&err];
+    if (!lib) { printf("MSL compile failed: %s\n", err.localizedDescription.UTF8String); return 1; }
+    id<MTLComputePipelineState> pu = [dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"ku"] error:&err], pf = [dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"kf"] error:&err],
+                                p1 = [dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"k1"] error:&err];
+    id<MTLCommandQueue> q = [dev newCommandQueue];
+    const uint32_t counts[] = { 1, 1000, 16384, 16385, 40000, 1048576 };
+    const struct { MTLPixelFormat pf; uint32_t bpb; const char *name; } fm[] = { { MTLPixelFormatR32Uint, 4, "R32Uint" }, { MTLPixelFormatRGBA32Float, 16, "RGBA32Float" } };
+    for (unsigned fi = 0; fi < 2; fi++) for (unsigned ci = 0; ci < sizeof counts / sizeof *counts; ci++) for (int oi = 0; oi < 3; oi++) {
+        uint32_t n = counts[ci], bpb = fm[fi].bpb, w, h, row; NSUInteger al = [dev minimumLinearTextureAlignmentForPixelFormat:fm[fi].pf], off = (NSUInteger)oi * al;
+        if (!nvn_tbuf_rows(n, bpb, (uint32_t)al, &w, &h, &row)) { check("rows", 0, "nvn_tbuf_rows refused"); continue; }
+        size_t len = off + (size_t)row * h + 4096;   // as a memory pool: more bytes follow the buffer texture
+        id<MTLBuffer> b = [dev newBufferWithLength:len options:MTLResourceStorageModeShared];
+        uint32_t *words = (uint32_t *)((uint8_t *)b.contents + off); for (size_t i = 0; i < (size_t)n * bpb / 4; i++) words[i] = (uint32_t)(i * 2654435761u ^ (i >> 7));
+        MTLTextureDescriptor *td = [MTLTextureDescriptor textureBufferDescriptorWithPixelFormat:fm[fi].pf width:n resourceOptions:b.resourceOptions usage:MTLTextureUsageShaderRead];
+        id<MTLTexture> tb = [b newTextureWithDescriptor:td offset:off bytesPerRow:((size_t)n * bpb + al - 1) / al * al];
+        MTLTextureDescriptor *d2 = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:fm[fi].pf width:w height:h mipmapped:NO]; d2.resourceOptions = b.resourceOptions; d2.usage = MTLTextureUsageShaderRead;
+        id<MTLTexture> t2 = [b newTextureWithDescriptor:d2 offset:off bytesPerRow:row];
+        char what[128]; snprintf(what, sizeof what, "%s %u elements at offset %lu: %ux%u rows of %u bytes", fm[fi].name, n, (unsigned long)off, w, h, row);
+        if (!tb || !t2) { check(what, 0, "texture not created"); continue; }
+        id<MTLBuffer> o = [dev newBufferWithLength:(size_t)n * 32 options:MTLResourceStorageModeShared];
+        id<MTLCommandBuffer> cb = [q commandBuffer]; id<MTLComputeCommandEncoder> ce = [cb computeCommandEncoder];
+        [ce setComputePipelineState:fi ? pf : pu]; [ce setTexture:tb atIndex:0]; [ce setTexture:t2 atIndex:1]; [ce setBuffer:o offset:0 atIndex:0]; [ce setBytes:&n length:4 atIndex:1];
+        [ce dispatchThreads:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; [ce endEncoding]; [cb commit]; [cb waitUntilCompleted];
+        const uint32_t *r = o.contents; int bad = -1;
+        for (uint32_t i = 0; i < n && bad < 0; i++) for (int c = 0; c < 4; c++) {
+            uint32_t mem = c < (int)(bpb / 4) ? words[i * (bpb / 4) + c] : (c == 3 ? (fi ? 0x3f800000u : 1u) : 0u);   // missing components of a one-channel read: 0, 0, 1
+            if (r[8 * i + c] != r[8 * i + 4 + c] || r[8 * i + c] != mem) { bad = (int)i; break; } }
+        char det[128] = ""; if (bad >= 0) snprintf(det, sizeof det, "element %d: buffer view %#x, rows %#x, memory %#x", bad, r[8 * bad], r[8 * bad + 4], words[bad * (bpb / 4)]);
+        check(what, bad < 0, det);
+    }
+    {   // a 1D texture as the host makes it (one row at most 16384 wide) reads the same at the wrapped address
+        uint32_t n = 16384; MTLTextureDescriptor *d = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatR32Uint width:n height:1 mipmapped:NO]; d.usage = MTLTextureUsageShaderRead;
+        id<MTLTexture> t = [dev newTextureWithDescriptor:d]; uint32_t *v = malloc(n * 4); for (uint32_t i = 0; i < n; i++) v[i] = i * 7919u + 3;
+        [t replaceRegion:MTLRegionMake2D(0, 0, n, 1) mipmapLevel:0 withBytes:v bytesPerRow:n * 4];
+        id<MTLBuffer> o = [dev newBufferWithLength:(size_t)n * 32 options:MTLResourceStorageModeShared];
+        id<MTLCommandBuffer> cb = [q commandBuffer]; id<MTLComputeCommandEncoder> ce = [cb computeCommandEncoder];
+        [ce setComputePipelineState:p1]; [ce setTexture:t atIndex:1]; [ce setBuffer:o offset:0 atIndex:0]; [ce setBytes:&n length:4 atIndex:1];
+        [ce dispatchThreads:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; [ce endEncoding]; [cb commit]; [cb waitUntilCompleted];
+        const uint32_t *r = o.contents; int bad = -1; for (uint32_t i = 0; i < n && bad < 0; i++) if (r[8 * i] != r[8 * i + 4] || r[8 * i] != v[i]) bad = (int)i;
+        check("a 16384-wide one-row 2D texture (a 1D texture) reads the same at the wrapped address", bad < 0, "differs");
+        free(v);
+    }
+    printf("== %d checks, %d failed ==\n", g_n, g_fail);
+    return g_fail != 0;
+}
@@ -0,0 +1,115 @@
+// Shadow vertex inputs (dm/gfx/tools/dmsh_vin.py, nvn_mtl.m vin_ct): an integer shader input also declared as a float4 and as the integer of the other sign, the three ORed, with the
+// attribute routed to the declaration of its format's class and the zero buffer to the others. Checks that Metal's fetch then gives the 32-bit register values the CPU conversion
+// (nvn_mtl.m vconv_elem) made: the converted float32 bits for half / normalized / packed formats (missing components 0, 0, 0, 1.0), the value widened by the format's own sign for
+// 8 / 16-bit integers (missing 0, 0, 0, 1), into a uint input and into an int input. The reference below follows vconv_elem's formulas; half and integer formats must match exactly,
+// normalized formats within one ULP. Negative control: the same data routed to the input's own declaration with the other class's format is refused by Metal or differs.
+#import <Metal/Metal.h>
+#include <math.h>
+#include <stdio.h>
+#include <string.h>
+
+static const char *k_msl =
+    "#include <metal_stdlib>\nusing namespace metal;\n"
+    "struct main0_in { uint4 m_a [[attribute(0)]]; int4 m_b [[attribute(1)]]; };\n"
+    "struct main0_in_v { uint4 m_a [[attribute(0)]]; int4 m_b [[attribute(1)]];\n"
+    "    float4 m_a_vf [[attribute(16)]]; float4 m_b_vf [[attribute(17)]]; int4 m_a_vo [[attribute(18)]]; uint4 m_b_vo [[attribute(19)]]; };\n"
+    "vertex void vs(main0_in_v dm_in [[stage_in]], uint vid [[vertex_id]], device uint4 *o [[buffer(1)]])\n{\n"
+    "    main0_in in = {};\n"
+    "    in.m_a = dm_in.m_a | as_type<uint4>(dm_in.m_a_vf) | as_type<uint4>(dm_in.m_a_vo);\n"
+    "    in.m_b = dm_in.m_b | as_type<int4>(dm_in.m_b_vf) | as_type<int4>(dm_in.m_b_vo);\n"
+    "    o[2 * vid] = in.m_a; o[2 * vid + 1] = as_type<uint4>(in.m_b);\n}\n";
+
+enum { NV = 64, CF = 0, CU = 1, CI = 2 };   // vertices; format classes: float-type, unsigned, signed
+typedef struct { const char *name; MTLVertexFormat mf; int cls, n, size, kind; } fmt_t;   // kind: how the reference converts (below)
+enum { K_HALF, K_UN8, K_SN8, K_UN16, K_SN16, K_U8, K_I8, K_U16, K_I16, K_1010102, K_11_11_10 };
+static const fmt_t F[] = {
+    { "Half2", MTLVertexFormatHalf2, CF, 2, 4, K_HALF }, { "Half4", MTLVertexFormatHalf4, CF, 4, 8, K_HALF }, { "UChar4Normalized", MTLVertexFormatUChar4Normalized, CF, 4, 4, K_UN8 },
+    { "Char2Normalized", MTLVertexFormatChar2Normalized, CF, 2, 2, K_SN8 }, { "UShort2Normalized", MTLVertexFormatUShort2Normalized, CF, 2, 4, K_UN16 },
+    { "Short4Normalized", MTLVertexFormatShort4Normalized, CF, 4, 8, K_SN16 }, { "UInt1010102Normalized", MTLVertexFormatUInt1010102Normalized, CF, 4, 4, K_1010102 },
+    { "FloatRG11B10", MTLVertexFormatFloatRG11B10, CF, 3, 4, K_11_11_10 },
+    { "UChar4", MTLVertexFormatUChar4, CU, 4, 4, K_U8 }, { "UShort2", MTLVertexFormatUShort2, CU, 2, 4, K_U16 },
+    { "Char4", MTLVertexFormatChar4, CI, 4, 4, K_I8 }, { "Short2", MTLVertexFormatShort2, CI, 2, 4, K_I16 }, { "Short4", MTLVertexFormatShort4, CI, 4, 8, K_I16 },
+};
+
+static float uf(uint32_t v, int mbits) { uint32_t e = v >> mbits & 31, m = v & ((1u << mbits) - 1); if (!e) return ldexpf((float)m, -14 - mbits); if (e == 31) return m ? NAN : INFINITY; return ldexpf(1.0f + (float)m / (float)(1u << mbits), (int)e - 15); }
+static void ref(const fmt_t *f, const uint8_t *p, uint32_t o[4]) {   // as vconv_elem: the register's four 32-bit values
+    if (f->cls != CF) { o[0] = o[1] = o[2] = 0; o[3] = 1;
+        for (int i = 0; i < f->n; i++) switch (f->kind) {
+            case K_U8: o[i] = p[i]; break; case K_I8: o[i] = (uint32_t)(int32_t)(int8_t)p[i]; break;
+            case K_U16: { uint16_t h; memcpy(&h, p + 2 * i, 2); o[i] = h; break; } case K_I16: { uint16_t h; memcpy(&h, p + 2 * i, 2); o[i] = (uint32_t)(int32_t)(int16_t)h; break; } }
+        return; }
+    float v[4] = { 0, 0, 0, 1 };
+    for (int i = 0; i < f->n; i++) switch (f->kind) {
+        case K_HALF: { _Float16 h; memcpy(&h, p + 2 * i, 2); v[i] = (float)h; break; }
+        case K_UN8: v[i] = p[i] / 255.0f; break; case K_SN8: v[i] = fmaxf((int8_t)p[i] / 127.0f, -1.0f); break;
+        case K_UN16: { uint16_t u; memcpy(&u, p + 2 * i, 2); v[i] = u / 65535.0f; break; } case K_SN16: { int16_t s; memcpy(&s, p + 2 * i, 2); v[i] = fmaxf(s / 32767.0f, -1.0f); break; }
+        case K_1010102: { uint32_t w; memcpy(&w, p, 4); uint32_t bits = i == 3 ? 2 : 10; v[i] = (float)(w >> (10 * i) & ((1u << bits) - 1)) / (float)((1u << bits) - 1); break; }
+        case K_11_11_10: { uint32_t w; memcpy(&w, p, 4); v[i] = uf(i == 0 ? w & 0x7ff : i == 1 ? w >> 11 & 0x7ff : w >> 22, i == 2 ? 5 : 6); break; }
+    }
+    memcpy(o, v, 16);
+}
+static uint32_t ulps(uint32_t a, uint32_t b) { int32_t x = (int32_t)a, y = (int32_t)b; if ((x < 0) != (y < 0)) return a == b ? 0 : UINT32_MAX; return (uint32_t)(x > y ? x - y : y - x); }
+
+static int g_n, g_fail;
+static void check(const char *what, int ok, const char *detail) { g_n++; if (!ok) g_fail++; printf("%-62s %s %s\n", what, ok ? "✓" : "✗", ok ? "" : detail); }
+
+int main(void) {
+    id<MTLDevice> dev = MTLCreateSystemDefaultDevice(); NSError *err = nil;
+    id<MTLLibrary> lib = [dev newLibraryWithSource:@(k_msl) options:nil error:&err];
+    if (!lib) { printf("MSL compile failed: %s\n", err.localizedDescription.UTF8String); return 1; }
+    id<MTLFunction> fn = [lib newFunctionWithName:@"vs"]; id<MTLCommandQueue> q = [dev newCommandQueue];
+    id<MTLBuffer> zero = [dev newBufferWithLength:64 options:MTLResourceStorageModeShared], out = [dev newBufferWithLength:NV * 32 options:MTLResourceStorageModeShared];
+    memset(zero.contents, 0, 64);
+    const int nf = (int)(sizeof F / sizeof *F);
+    // input a (uint) gets every format, input b (int) the next one in the list: each pipeline mixes two classes, as Age of Calamity's soldiers do
+    for (int ia = 0; ia < nf; ia++) {
+        const fmt_t *fa = &F[ia], *fb = &F[(ia + 1) % nf];
+        uint8_t da[NV * 8], db[NV * 8];
+        for (int i = 0; i < NV * 8; i++) { da[i] = (uint8_t)(i * 37 + 11 ^ i >> 3); db[i] = (uint8_t)(i * 91 + 5 ^ 0xA5); }
+        for (int i = 0; i < NV * 8; i += 2) { if (fa->kind == K_HALF && (da[i + 1] & 0x7c) == 0x7c) da[i + 1] &= 0xbb; if (fb->kind == K_HALF && (db[i + 1] & 0x7c) == 0x7c) db[i + 1] &= 0xbb; }   // no half NaN / inf
+        id<MTLBuffer> ba = [dev newBufferWithBytes:da length:sizeof da options:MTLResourceStorageModeShared], bb = [dev newBufferWithBytes:db length:sizeof db options:MTLResourceStorageModeShared];
+        MTLVertexDescriptor *vd = [MTLVertexDescriptor vertexDescriptor];
+        // the routing vin_ct makes: own class → the input's declaration, else the shadow of the format's class; every other declaration reads the zero buffer
+        const int idx_a[3] = { 16, 0, 18 }, idx_b[3] = { 17, 19, 1 };   // by class CF / CU / CI
+        const MTLVertexFormat zf[3] = { MTLVertexFormatFloat4, MTLVertexFormatUInt4, MTLVertexFormatInt4 };
+        for (int c = 0; c < 3; c++) {
+            int ia_ = idx_a[c], ib_ = idx_b[c];
+            if (c == fa->cls) { vd.attributes[ia_].format = fa->mf; vd.attributes[ia_].bufferIndex = 2; } else { vd.attributes[ia_].format = zf[c]; vd.attributes[ia_].bufferIndex = 30; }
+            if (c == fb->cls) { vd.attributes[ib_].format = fb->mf; vd.attributes[ib_].bufferIndex = 3; } else { vd.attributes[ib_].format = zf[c]; vd.attributes[ib_].bufferIndex = 30; }
+            vd.attributes[ia_].offset = vd.attributes[ib_].offset = 0;
+        }
+        vd.layouts[2].stride = 8; vd.layouts[3].stride = 8;
+        vd.layouts[30].stride = 16; vd.layouts[30].stepFunction = MTLVertexStepFunctionConstant; vd.layouts[30].stepRate = 0;
+        MTLRenderPipelineDescriptor *pd = [MTLRenderPipelineDescriptor new]; pd.vertexFunction = fn; pd.rasterizationEnabled = NO; pd.vertexDescriptor = vd;
+        id<MTLRenderPipelineState> ps = [dev newRenderPipelineStateWithDescriptor:pd error:&err];
+        char name[96]; snprintf(name, sizeof name, "%s into uint, %s into int", fa->name, fb->name);
+        if (!ps) { check(name, 0, err.localizedDescription.UTF8String); continue; }
+        memset(out.contents, 0xEE, NV * 32);
+        MTLRenderPassDescriptor *rp = [MTLRenderPassDescriptor renderPassDescriptor]; rp.renderTargetWidth = 1; rp.renderTargetHeight = 1; rp.defaultRasterSampleCount = 1;
+        id<MTLCommandBuffer> cb = [q commandBuffer]; id<MTLRenderCommandEncoder> en = [cb renderCommandEncoderWithDescriptor:rp];
+        [en setRenderPipelineState:ps]; [en setVertexBuffer:out offset:0 atIndex:1]; [en setVertexBuffer:ba offset:0 atIndex:2]; [en setVertexBuffer:bb offset:0 atIndex:3]; [en setVertexBuffer:zero offset:0 atIndex:30];
+        [en drawPrimitives:MTLPrimitiveTypePoint vertexStart:0 vertexCount:NV]; [en endEncoding]; [cb commit]; [cb waitUntilCompleted];
+        const uint32_t *o = out.contents; uint32_t worst = 0; char detail[160] = "";
+        for (int v = 0; v < NV; v++) for (int s = 0; s < 2; s++) {
+            const fmt_t *f = s ? fb : fa; uint32_t e[4]; ref(f, (s ? db : da) + 8 * v, e);
+            for (int c = 0; c < 4; c++) { uint32_t d = ulps(o[8 * v + 4 * s + c], e[c]);
+                if (d > worst) { worst = d; snprintf(detail, sizeof detail, "vertex %d %s component %d: got %#x, expected %#x (%u ULP)", v, f->name, c, o[8 * v + 4 * s + c], e[c], d); } }
+        }
+        int exact = (fa->kind == K_HALF || fa->cls != CF) && (fb->kind == K_HALF || fb->cls != CF);
+        char what[128]; snprintf(what, sizeof what, "%s: %s", name, exact ? "bit-exact" : "within 1 ULP");
+        check(what, exact ? worst == 0 : worst <= 1, detail);
+    }
+    // negative control: Half2 straight into the uint input (no shadow) is refused by Metal, which is why the shadows exist
+    {
+        MTLVertexDescriptor *vd = [MTLVertexDescriptor vertexDescriptor];
+        vd.attributes[0].format = MTLVertexFormatHalf2; vd.attributes[0].bufferIndex = 2;
+        vd.attributes[1].format = MTLVertexFormatShort2; vd.attributes[1].bufferIndex = 3;
+        for (int i = 16; i < 20; i++) { vd.attributes[i].format = i == 18 ? MTLVertexFormatInt4 : i == 19 ? MTLVertexFormatUInt4 : MTLVertexFormatFloat4; vd.attributes[i].bufferIndex = 30; }
+        vd.layouts[2].stride = 8; vd.layouts[3].stride = 8; vd.layouts[30].stride = 16; vd.layouts[30].stepFunction = MTLVertexStepFunctionConstant; vd.layouts[30].stepRate = 0;
+        MTLRenderPipelineDescriptor *pd = [MTLRenderPipelineDescriptor new]; pd.vertexFunction = fn; pd.rasterizationEnabled = NO; pd.vertexDescriptor = vd;
+        id<MTLRenderPipelineState> ps = [dev newRenderPipelineStateWithDescriptor:pd error:&err];
+        check("negative control: Half2 into a uint input is refused by Metal", ps == nil, "Metal built the pipeline");
+    }
+    printf("== %d checks, %d failed ==\n", g_n, g_fail);
+    return g_fail != 0;
+}
@@ -0,0 +1,81 @@
+"""Shadow vertex inputs (dm/gfx/tools/dmsh_vin.py): small vertex shaders in the translator's MSL shape, no game needed. The variant must give each integer input a float4 shadow and
+then other-sign shadows while attribute indices 16-30 last, OR them into the input in both entry points, leave the transform feedback entry out, and compile with the Metal compiler
+for macOS and iOS (temporary files only in a temporary folder)."""
+import os
+from pathlib import Path
+import re
+import subprocess
+import sys
+import tempfile
+import unittest
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / 'tools'))
+import dmsh_vin
+
+def shader(types, tf=False):
+    """types: the MSL type of each input in location order ('float4' / 'uint4' / 'int4')"""
+    ins = ''.join(f'    {t} m_{l} [[attribute({l})]];\n' for l, t in enumerate(types))
+    use = ''.join(f'    acc += float4(as_type<float4>(in.m_{l}));\n' if t != 'float4' else f'    acc += in.m_{l};\n' for l, t in enumerate(types))
+    main = lambda name: f'\nvertex main0_out {name}(main0_in in [[stage_in]], constant float4& u [[buffer(0)]])\n{{\n    main0_out out = {{}};\n    float4 acc = u;\n{use}    out.gl_Position = acc;\n    return out;\n}}\n'
+    src = '#include <metal_stdlib>\n#include <simd/simd.h>\n\nusing namespace metal;\n\nstruct main0_out\n{\n    float4 gl_Position [[position]];\n};\n\nstruct main0_in\n{\n' + ins.rstrip('\n') + '\n};\n'
+    src += main('main0') + main('main0_z')
+    if tf: src += '\nvertex void main0_tf(main0_in in [[stage_in]], constant float4& u [[buffer(0)]], uint dm_vid [[vertex_id]], device uint* dm_tf0 [[buffer(8)]])\n{ dm_tf0[dm_vid] = 0u; }\n'
+    return src
+
+class ShadowInputs(unittest.TestCase):
+    def compiles(self, msl):
+        with tempfile.TemporaryDirectory(prefix='dmsh_vin_test_') as d:
+            p = Path(d) / 'v.metal'; p.write_text(msl)
+            for sdk in ('macosx', 'iphoneos'):
+                r = subprocess.run(['xcrun', '-sdk', sdk, 'metal', '-c', str(p), '-o', str(Path(d) / f'v.{sdk}.air')], capture_output=True, text=True)
+                self.assertEqual(r.returncode, 0, f'{sdk}: {r.stderr[:400]}')
+
+    def test_one_uint_input(self):
+        v, sh = dmsh_vin.variant(shader(['float4', 'uint4']))
+        self.assertEqual(sh, [{'location': 1, 'float': 16, 'sint': 17}])
+        self.assertIn('    int4 m_1_vo [[attribute(17)]];', v)
+        self.assertEqual(v.count('in.m_1 = dm_in.m_1 | as_type<uint4>(dm_in.m_1_vf) | as_type<uint4>(dm_in.m_1_vo);'), 2)   # main0 and main0_z
+        self.assertEqual(v.count('    in.m_0 = dm_in.m_0;'), 2)
+        self.compiles(v)
+
+    def test_int_input_gets_a_uint_shadow(self):
+        v, sh = dmsh_vin.variant(shader(['int4']))
+        self.assertEqual(sh, [{'location': 0, 'float': 16, 'uint': 17}])
+        self.assertIn('in.m_0 = dm_in.m_0 | as_type<int4>(dm_in.m_0_vf) | as_type<int4>(dm_in.m_0_vo);', v)
+        self.compiles(v)
+
+    def test_eight_inputs_fill_every_index(self):
+        v, sh = dmsh_vin.variant(shader(['uint4'] * 8))
+        self.assertEqual([s['float'] for s in sh], list(range(16, 24)))
+        self.assertEqual([s.get('sint') for s in sh], list(range(24, 31)) + [None])   # 15 indices: the last input has only its float shadow
+        self.assertIn('in.m_7 = dm_in.m_7 | as_type<uint4>(dm_in.m_7_vf);', v)
+        self.compiles(v)
+
+    def test_more_inputs_than_indices(self):
+        v, sh = dmsh_vin.variant(shader(['uint4'] * 16))
+        self.assertEqual(len(sh), 15)   # the 16th integer input gets no shadow: its formats keep the CPU conversion
+        self.assertIn('    in.m_15 = dm_in.m_15;', v)
+        self.compiles(v)
+
+    def test_transform_feedback_entry_left_out(self):
+        v, sh = dmsh_vin.variant(shader(['uint4'], tf=True))
+        self.assertNotIn('main0_tf', v)
+        self.compiles(v)
+
+    def test_split_input_array(self):
+        src = shader(['float4', 'float4', 'uint4']).replace('m_0', 'm_7_0').replace('m_1', 'm_7_1').replace('m_2', 'm_7_2')   # SPIRV-Cross splits an input array per element
+        self.assertEqual(dmsh_vin.variant(shader(['float4', 'float4']).replace('m_0', 'm_7_0').replace('m_1', 'm_7_1'))[1], 'no integer input')
+        v, sh = dmsh_vin.variant(src)
+        self.assertEqual(sh, [{'location': 2, 'float': 16, 'sint': 17}])
+        self.assertIn('in.m_7_2 = dm_in.m_7_2 | as_type<uint4>(dm_in.m_7_2_vf) | as_type<uint4>(dm_in.m_7_2_vo);', v)
+        self.compiles(v)
+
+    def test_refusals(self):
+        self.assertEqual(dmsh_vin.variant(shader(['float4']))[1], 'no integer input')
+        no_input = re.sub(r'\nstruct main0_in\n\{\n.*?\n\};\n', '\n', shader(['uint4']), flags=re.S).replace('main0_in in [[stage_in]], ', '')
+        self.assertEqual(dmsh_vin.variant(no_input), (None, 'no integer input'))   # no stage input at all: nothing to shadow, not a refusal
+        self.assertIsNone(dmsh_vin.variant(shader(['uint4']).replace('[[attribute(0)]]', '[[attribute(16)]]'))[0])
+        self.assertIsNone(dmsh_vin.variant(shader(['uint4']).replace('uint4 m_0', 'uint2 m_0'))[0])
+
+if __name__ == '__main__':
+    unittest.main()
@@ -10,7 +10,7 @@ import ast, inspect, os, re, sys, json, struct, hashlib, shutil, subprocess, tem
 DM_ROOT = os.environ.get('DM_ROOT') or os.path.expanduser('~/dm-work')   # Workspace root (game files, build outputs and run artifacts are all under it)
 HOME = os.path.expanduser('~')
 dm_dir, out = sys.argv[1:3]; want = set(sys.argv[3:]); cache_root = out
-sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))); import dmsh_chain, dmsh_fetch, dmsh_facts, dmsh_mtlcache   # translation chain version: verify the tools against the lock file first; build output goes into b-<V>/
+sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))); import dmsh_chain, dmsh_fetch, dmsh_facts, dmsh_mtlcache, dmsh_vin   # translation chain version: verify the tools against the lock file first; build output goes into b-<V>/
 _bad = dmsh_chain.check_tools()
 if _bad: sys.exit('\n'.join('[RED] translation chain tools do not match the lock file: ' + b for b in _bad) + '\n(if the tools changed, run dmsh_chain.py --update-lock and commit the lock file; if the translator just was not rebuilt, rebuild it)')
 CHAIN = dmsh_chain.version(); out = os.path.join(out, 'b-' + CHAIN)
@@ -144,7 +144,7 @@ def publish_cache(key, digest, rule_fp, code_fp):
        for name in names:
            if name.endswith('.json'):
                src = os.path.join(out if label == 'mac' else IOSOUT, name)
-                value = json.load(open(src)); value.pop('fetch', None)
+                value = json.load(open(src)); value.pop('fetch', None); value.pop('vin', None)   # both side libraries are made again from the restored MSL (fetch_finish, vin_finish)
                atomic_json(os.path.join(d, name), value)
    atomic_json(cache_meta_path(key), {'digest': digest, 'rule_fp': rule_fp, 'code_fp': code_fp})

@@ -223,6 +223,37 @@ def fetch_finish(key, src, refl):
    write_refl(key, refl)
    return refl

+VIN_STATS = collections.Counter()
+
+def vin_finish(key, full, refl):
+    """The shadow-input library <key>_bv of a vertex shader with integer inputs (dmsh_vin.py), compiled through the MSL-keyed Metal cache; reflection 'vin' names its entry points and
+    the shadow attribute indices. Other shaders get none (stale libraries removed). Changes refl; the caller writes it"""
+    if refl['stage'] != 'vs': return refl
+    libs = [os.path.join(out, key + '_bv.metallib')] + ([os.path.join(IOSOUT, key + '_bv.metallib')] if IOSOUT else [])
+    refl.pop('vin', None)
+    def drop():
+        for p in libs:
+            try: os.unlink(p)
+            except FileNotFoundError: pass
+    v, shadows = dmsh_vin.variant(full)
+    if not v:
+        if shadows != 'no integer input': VIN_STATS['not generated'] += 1; print(f'[RED] {key}: no shadow-input library ({shadows}): its integer inputs keep the CPU conversion', flush=True)
+        drop(); return refl
+    VIN_STATS['generated'] += 1
+    msl = os.path.join(w, key + '_bv.metal'); open(msl, 'w').write(v)
+    try:
+        MTL.compile('mac', msl, libs[0])
+        if IOSOUT: MTL.compile('ios', msl, libs[1])
+    except subprocess.CalledProcessError as e:
+        VIN_STATS['failed'] += 1; print(f'[RED] {key}: the shadow-input library does not compile ({str(e)[:120]}): its integer inputs keep the CPU conversion', flush=True)
+        drop(); return refl
+    finally:
+        try: os.unlink(msl)
+        except FileNotFoundError: pass
+    refl['vin'] = {'entry': 'main0', 'entry_z': 'main0_z', 'shadow': shadows}
+    VIN_STATS['inputs'] += len(shadows); VIN_STATS['float only'] += sum(1 for s in shadows if 'sint' not in s and 'uint' not in s)
+    return refl
+
 pc_of, loc = {}, {}
 MAP_ROWS, MAP_CONFLICT_ROWS = {}, {}
 for raw in open(MAP, 'rb'):
@@ -393,6 +424,10 @@ def reflect(key, stage, msl, res):
            d = decode_tex(nm)
            if not d: e.update(kind='?')
            else: e.update(kind='img' if re.match(r'[a-z]+_i', nm) else ('tbuf' if 'texture_buffer' in ty else 'tex'), reg=d[0])
+            # a 1D integer fetch the translator read as 2D rows of 16384 texels (Program.cs RecordTexelFetch1D): the host binds a texture buffer in this slot as such rows (nvn_mtl.m tbuf_as_2d)
+            if e.get('kind') == 'tex' and e.get('reg') in res.get('fetch1d', ()):
+                e['fetch1d'] = True
+                if re.search(r'\b%s\.get_width\(' % re.escape(nm), msl): e['fetch1d_size'] = True   # also size-queried: a texture buffer of more than 16384 elements reports 16384 there
        binds.append(e)
    # Storage buffers: SPIR-V binding number → c0 offset → NVN storage buffer slot (c0 base: vertex 0x110, compute 0x310, pixel 0x510, 0x10 per slot)
    base = SSBO.get(stage)   # storage buffer base for the hull / domain stages has not been measured: a storage buffer there is reported with no slot
@@ -613,6 +648,7 @@ def finish(key, w, j, c1):
            '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)],
            'raster': raster_facts(j['stage'], src), 'vs_pad': pad}
+    if j['stage'] == 'vs': vin_finish(key, open(msl).read(), refl)   # the shadow-input library (see dmsh_vin.py), from the source with every entry point
    write_refl(key, refl)
    fetch_finish(key, src, refl)
    return 'ok', '', refl
@@ -661,7 +697,8 @@ for r in xr:
 def fin(key):
    if key in cache_hits:
        refl = json.load(open(os.path.join(out, key + '_b.json')))
-        src = open(os.path.join(out, key + '_b.metal')).read() if refl['stage'] == 'ps' else ''
+        src = open(os.path.join(out, key + '_b.metal')).read() if refl['stage'] in ('ps', 'vs') else ''
+        if refl['stage'] == 'vs': write_refl(key, vin_finish(key, src, refl))
        return key, 'ok', '', fetch_finish(key, src, refl)
    if key in xerr or not os.path.exists(os.path.join(w, key + '.spv')): return key, 'translation failed', xerr.get(key, 'no spv'), None
    result = (key,) + finish(key, w, jobs[key], c1s[key])
@@ -676,6 +713,8 @@ print(f'shader cache: hit {PRE_STATS["hit"]}, rebuilt {PRE_STATS["miss"]}')
 for reason in ('new key', 'rule changed', 'code fingerprint changed: all rebuilt'):
    if PRE_STATS[reason]: print(f'  {reason} {PRE_STATS[reason]}')
 print(f'variant cache: generated {BF_STATS["generated"]}, hit {BF_STATS["hit"]}, rebuilt {BF_STATS["miss"]}, compiled {BF_STATS["compile"]}')
+print(f'shadow-input libraries (<key>_bv, dmsh_vin.py): {VIN_STATS["generated"]} generated, {VIN_STATS["failed"]} failed to compile, {VIN_STATS["not generated"]} not generated; '
+      f'{VIN_STATS["inputs"]} integer inputs shadowed, {VIN_STATS["float only"]} of them with no other-sign shadow (attribute indices ran out)')
 print(MTL.summary())
 for lvl in [k for k in st if k != 'ok']:
    for e, n in collections.Counter(r[2] for r in res if r[1] == lvl).most_common(10): print(f'  [{lvl}] {n:5d}  {e}')
@@ -683,6 +722,9 @@ ok = [r[3] for r in res if r[1] == 'ok']
 rd = [x for x in ok if x['drv_reads']]
 print(f'shaders reading driver segments at run time {len(rd)} (c0 {sum("0" in x["drv_reads"] for x in rd)}, c2 {sum("2" in x["drv_reads"] for x in rd)}); reading what the replay layer cannot fill {sum(bool(x["drv_unfillable"]) for x in ok)}')
 for e, n in collections.Counter(u for x in ok for u in x['drv_unfillable']).most_common(15): print(f'  cannot fill {n:5d}  {e}')
+f1 = [b for x in ok for b in x['bindings'] if b.get('fetch1d')]
+print(f'1D integer fetches read as 2D rows of 16384 texels (a texture buffer there is bound as rows, nvn_mtl.m tbuf_as_2d): {len(f1)} bindings in {sum(any(b.get("fetch1d") for b in x["bindings"]) for x in ok)} shaders, '
+      f'{sum(1 for b in f1 if b.get("fetch1d_size"))} of them also size-queried')
 SLOTK = {'cbuf', 'samp', 'samp_tex', 'tex', 'img', 'tbuf', 'sbuf', 'ubo_ptr', 'sbuf_ubo', 'tex_ubo', 'samp_ubo', 'samp_id'}
 noreg = [x for x in ok if any(b['kind'] in SLOTK and b.get('reg') is None for b in x['bindings'])]
 unk = [x for x in ok if any(b['kind'] == '?' for b in x['bindings'])]
@@ -1,8 +1,8 @@
 # Translation chain tool fingerprints (written by dmsh_chain.py --update-lock; regenerate after changing tools and commit it)
-dmsh_xlate.dll	986fad61f44dc9a4
-Ryujinx.Graphics.Shader.dll	6f995909ae6fecb5
-Ryujinx.Common.dll	917c02ac432765af
-xlate_source	3cc20c29e9528168
+dmsh_xlate.dll	b87f3970471f3b06
+Ryujinx.Graphics.Shader.dll	d0b54941cc2237d9
+Ryujinx.Common.dll	78c38eb78e9e8b4c
+xlate_source	3ac10c3660728251
 spirv-cross	d9dc197f061b42a5
 metal_macosx	Apple metal version 32023.921 (metalfe-32023.921.6)
 metal_iphoneos	Apple metal version 32023.921 (metalfe-32023.921.6)
@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 # Translation chain version of the B shader library (the build result must be determined by source alone).
-# Chain version V = first 12 digits of the sha256 of the contents of the following in-repo files + the translation options: this script, dmsh_build.py, dmsh_fetch.py (readable-tile version), the dmsh_xlate sources (Program.cs, csproj),
+# Chain version V = first 12 digits of the sha256 of the contents of the following in-repo files + the translation options: this script, dmsh_build.py, dmsh_fetch.py (readable-tile version), dmsh_vin.py (shadow vertex inputs), the dmsh_xlate sources (Program.cs, csproj),
 # from the current game profile (dm/games/<id>/, chosen by DM_GAME) the mapping table (including the conflict table), the vertex / pixel pair table (shader_pairs.tsv, used by vs_pad), the answer table for translation queries (dmsh_facts.py and shader_facts.tsv),
 # the keys of the profile's game.tsv that affect the build (nvn_*, shader_*), the iOS compile convention ios_metal.py, and the tool lock dmsh_chain.lock (measured fingerprints of the Ryujinx shader dll / spirv-cross / metal compiler). The tables differ per game, so V does too.
 # The build (dmsh_build.py) writes into <cache dir>/b-<V>/ and the reflection json records chain = V; at compile time the host takes V from the game profile's dmsh_chain.h, loads only from b-<V>/ and verifies each json.
@@ -14,7 +14,7 @@ DM_ROOT = os.environ.get('DM_ROOT') or os.path.expanduser('~/dm-work')   # Works
 HERE = os.path.dirname(os.path.abspath(__file__)); GFX = os.path.dirname(HERE); HOME = os.path.expanduser('~')
 sys.path.insert(0, os.path.join(GFX, '../games')); import game as _game   # game profiles
 GAME = _game.load()
-FILES = ['tools/dmsh_chain.py', 'tools/dmsh_build.py', 'tools/dmsh_fetch.py', 'tools/dmsh_facts.py', 'tools/dmsh_xlate/Program.cs', 'tools/dmsh_xlate/dmsh_xlate.csproj', '../ios/tools/ios_metal.py', 'tools/dmsh_chain.lock']
+FILES = ['tools/dmsh_chain.py', 'tools/dmsh_build.py', 'tools/dmsh_fetch.py', 'tools/dmsh_vin.py', 'tools/dmsh_facts.py', 'tools/dmsh_xlate/Program.cs', 'tools/dmsh_xlate/dmsh_xlate.csproj', '../ios/tools/ios_metal.py', 'tools/dmsh_chain.lock']
 TABLES = ['shader_facts.tsv', 'shader_map.tsv', 'shader_map.tsv.conflicts.tsv', 'shader_pairs.tsv']   # in the game profile folder
 LOCK = os.path.join(HERE, 'dmsh_chain.lock'); HEADER = str(GAME.table('dmsh_chain.h'))
 # the translator defaults to the in-repo build that matches the lock file (a rebuilt dll has different fingerprints)
@@ -0,0 +1,50 @@
+# Shadow vertex inputs (<key>_bv): dmsh_build.py declares a vertex shader input an integer when its first translation only bit-casts it (int_inputs), and the game may still feed it
+# a format of another class: a float-type format (half, unorm / snorm, 10:10:10:2, 11:11:10), whose converted float32 bits Maxwell's fetch puts in the register, or an 8 / 16-bit
+# integer of the other signedness, which Maxwell widens by the format's own sign. Metal's fetch reads a format only into an input of its own class, so the host converted those
+# attributes on the CPU every draw (nvn_mtl.m vconv_bind; about a quarter of Age of Calamity's replay thread in battle). This variant declares every integer input also as a float4
+# and as the integer of the other sign, at Metal attribute indices 16-30 (NVN attributes are 0-15), and ORs the three into the input the body reads before it runs. The host sends each
+# attribute to the index of its format's class and the zero buffer to the others (nvn_mtl.m vin_ct), so the shader sees Maxwell's register bits, whatever mix of formats a pipeline has.
+# Built offline into a library of its own beside <key>_b, with the same argument buffer; entry points main0 / main0_z (the transform feedback entry is left out: those draws keep
+# the CPU conversion). Usage: variant(msl) → (new MSL, [{'location', 'float', 'sint' | 'uint'}]) or (None, reason)
+import re
+
+SLOT0, SLOT_END = 16, 31   # Metal vertex attribute indices 16..30
+
+def variant(msl):
+    m = re.search(r'\nstruct main0_in\n\{\n(.*?)\n\};', msl, re.S)
+    if not m: return None, ('no main0_in struct' if '[[stage_in]]' in msl else 'no integer input')   # no stage input: the shader reads no vertex attribute
+    members = []
+    for ln in m.group(1).split('\n'):
+        mm = re.fullmatch(r'    (float4|uint4|int4) (m_\d+(?:_\d+)*) \[\[attribute\((\d+)\)\]\];', ln)   # m_28_0: an input array split per element
+        if not mm: return None, 'an input declaration not understood: ' + ln.strip()[:60]
+        members.append((mm.group(1), mm.group(2), int(mm.group(3))))
+    ints = [x for x in members if x[0] != 'float4']
+    if not ints: return None, 'no integer input'
+    if any(l >= SLOT0 for _, _, l in members): return None, 'an input at attribute 16 or above'
+    slot, shadows, extra = SLOT0, [], []
+    for t, n, l in ints:   # a float4 shadow for every integer input first (half and normalized formats are the common case) ...
+        if slot >= SLOT_END: break
+        extra.append(f'    float4 {n}_vf [[attribute({slot})]];'); shadows.append({'location': l, 'float': slot}); slot += 1
+    for e, (t, n, l) in zip(shadows, ints):   # ... then the other signedness while indices last
+        if slot >= SLOT_END: break
+        o = 'int4' if t == 'uint4' else 'uint4'
+        extra.append(f'    {o} {n}_vo [[attribute({slot})]];'); e['sint' if o == 'int4' else 'uint'] = slot; slot += 1
+    covered = {e['location']: e for e in shadows}
+    head = msl[:m.end()] + '\n\nstruct main0_in_v\n{\n' + m.group(1) + '\n' + '\n'.join(extra) + '\n};'
+    body = msl[m.end():]
+    cut = body.find('\nvertex void main0_tf(')
+    if cut >= 0: body = body[:cut] + '\n'
+    pro = '    main0_in in = {};\n'
+    for t, n, l in members:
+        e = covered.get(l)
+        ors = '' if not e else f' | as_type<{t}>(dm_in.{n}_vf)' + (f' | as_type<{t}>(dm_in.{n}_vo)' if ('sint' in e or 'uint' in e) else '')
+        pro += f'    in.{n} = dm_in.{n}{ors};\n'
+    n_entry = 0
+    for entry in ('main0', 'main0_z'):
+        hdr = re.search(r'\nvertex main0_out %s\(main0_in in \[\[stage_in\]\]([^\n]*)\)\n\{\n' % entry, body)
+        if not hdr:
+            if entry == 'main0': return None, 'main0 signature not recognized'
+            continue
+        new = f'\nvertex main0_out {entry}(main0_in_v dm_in [[stage_in]]{hdr.group(1)})\n{{\n' + pro
+        body = body[:hdr.start()] + new + body[hdr.end():]; n_entry += 1
+    return head + body, shadows
@@ -8,7 +8,8 @@
 //   Job: [{ "key", "stage": "vs|ps|cs", "code": machine code file (graphics = 0x50-byte SPH + instructions, compute = instructions),
 //           "tex": { "slot": "2d|2darray|cube|cubearray|3d|buffer|1d|1darray|2dms|2dmsarray", ... }, "img": { "slot": { "type", "fmt" } },
 //           "attr": { "location": "float|sint|uint" }, "fout": { "target": "float|sint|uint" }, "local": [x, y, z] }]
-//   Output: <key>.spv and <key>.res.json (offsets of storage buffer slots in c0, etc.; decoding texture handle → slot number is done in the script that builds the metallib)
+//   Output: <key>.spv and <key>.res.json (offsets of storage buffer slots in c0, etc.; decoding texture handle → slot number is done in the script that builds the metallib;
+//           "fetch1d": the texture slots read by 1D integer fetches translated as wrapped 2D reads, see RecordTexelFetch1D)
 using Ryujinx.Common.Memory;
 using Ryujinx.Graphics.Shader;
 using Ryujinx.Graphics.Shader.Translation;
@@ -73,6 +74,10 @@ static class SwshXlate
                _ => SamplerType.Texture2D,
            };
        }
+        // Texture slots read by a 1D integer fetch that the translator made a wrapped 2D read (x & 0x3FFF, x >> 14) for want of a type answer: they may hold a 1D texture or a texture buffer,
+        // which only the bound descriptor tells (res "fetch1d"; dmsh_build.py flags their bindings, and the host binds a texture buffer there as 16384-texel rows, nvn_mtl.m tbuf_as_2d)
+        public readonly SortedSet<int> Fetch1D = new();
+        public void RecordTexelFetch1D(int handle) { int slot = Slot(handle, -1, false); if (slot >= 0) Fetch1D.Add(slot); }
        public TextureFormat QueryTextureFormat(int handle, int cbufSlot = -1)
        {
            int slot = Slot(handle, cbufSlot, true);
@@ -155,6 +160,7 @@ static class SwshXlate
                foreach (var b in p.Info.SBuffers) res["sbufs"].AsArray().Add(new JsonObject { ["binding"] = b.Binding, ["slot"] = b.Slot, ["cb_slot"] = b.SbCbSlot, ["cb_offset"] = b.SbCbOffset, ["flags"] = (int)b.Flags });
                foreach (var (list, name) in new[] { (p.Info.Textures, "textures"), (p.Info.Images, "images") })
                    foreach (var t in list) res[name].AsArray().Add(new JsonObject { ["binding"] = t.Binding, ["type"] = (int)t.Type, ["format"] = t.Format.ToString(), ["cbuf_slot"] = t.CbufSlot, ["handle"] = t.HandleIndex, ["array"] = t.ArrayLength, ["separate"] = t.Separate, ["flags"] = (int)t.Flags });
+                if (acc.Fetch1D.Count > 0) { var f1 = new JsonArray(); foreach (int h in acc.Fetch1D) f1.Add(h); res["fetch1d"] = f1; }
                File.WriteAllText(Path.Combine(outDir, key + ".res.json"), res.ToJsonString(new JsonSerializerOptions { WriteIndented = true }));
                ok++;
            }
@@ -7,5 +7,7 @@ Three projects from Ryujinx 1.3.2: `Ryujinx.Graphics.Shader`, `Ryujinx.Common` a
  - Texture coordinates that are NaN are replaced by 0 (the `nfZero` select). This came with the first import.
 - `Ryujinx.Graphics.Shader/Translation/Optimizations/BindlessElimination.cs`, `TryGetOperation` / `IsSameOperand`: a bindless handle that reaches the texture instruction through a phi is followed through the phis that feed it (a handle kept in a loop variable comes back through the loop's own phi), and two `constant buffer | constant` handles count as the same value. Upstream compared only the phi's direct inputs, so such handles fell to the texture-pool array of 32 (`*_t2d_cb0_40000000`), which cannot be bound; 218 pixel shaders in Monster Hunter Rise (2026-10-20). Marked `nxnative:` in the code.

+- `Ryujinx.Graphics.Shader/Instructions/InstEmitTexture.cs`, `EmitTex` (TLD 1D with integer coordinates only) and `EmitTexs` (TLDS `Texture1DLodZero`), with `IGpuAccessor.RecordTexelFetch1D`: Maxwell reads a 1D texture and a texture buffer with the same instruction, and the bound descriptor decides. When the texture type query does not answer "buffer", the read is translated as 2D at `(x & 0x3FFF, x >> 14)` instead of `(x, 0)`, and the handle is reported. A 1D texture (one row at most 16384 wide on the host) reads the same; the host binds a texture buffer in such a slot as a linear 2D texture of 16384-texel rows over its memory (`nvn_mtl.m` `tbuf_as_2d`). Upstream read `(x, 0)` and the host had only a dummy texture to give: the light and decal lists of Age of Calamity (at least 15,000 shaders) and Three Houses read zeros (2026-10-21). Marked `nxnative:` in the code.
+
 The shader translation tool in `dm/gfx/tools/dmsh_xlate` references them to translate the original Switch shaders to SPIR-V.
 License: MIT, see `LICENSE.txt`. Upstream: <https://github.com/Ryubing>.
@@ -421,6 +421,15 @@ namespace Ryujinx.Graphics.Shader
            return SamplerType.Texture2D;
        }

+        /// <summary>
+        /// nxnative: a 1D integer-coordinate fetch (TLD / TLDS 1D) of the bound texture handle was translated as a 2D read at (x &amp; 0x3FFF, x &gt;&gt; 14),
+        /// because the texture type query did not say it is a texture buffer.
+        /// </summary>
+        /// <param name="handle">Texture handle</param>
+        void RecordTexelFetch1D(int handle)
+        {
+        }
+
        /// <summary>
        /// Gets the size in bytes of a bound constant buffer for the current shader stage.
        /// </summary>
@@ -210,6 +210,7 @@ namespace Ryujinx.Graphics.Shader.Instructions
            }

            bool hasLod = lodMode > Lod.Lz;
+            bool wrap1D = false;

            if (type == SamplerType.Texture1D && (flags & ~TextureFlags.Bindless) == TextureFlags.IntCoords && !(
                hasLod ||
@@ -225,6 +226,11 @@ namespace Ryujinx.Graphics.Shader.Instructions
                {
                    type = SamplerType.TextureBuffer;
                }
+                else
+                {
+                    // nxnative: the descriptor decides between a 1D texture and a texture buffer, so the 2D read below wraps x into rows of 16384 texels (see the TLDS Texture1DLodZero case)
+                    wrap1D = Sample1DAs2D;
+                }
            }

            int coordsCount = type.Dimensions;
@@ -238,7 +244,17 @@ namespace Ryujinx.Graphics.Shader.Instructions

            if (Sample1DAs2D && (type & SamplerType.Mask) == SamplerType.Texture1D)
            {
-                sourcesList.Add(ConstF(0));
+                if (wrap1D)
+                {
+                    Operand x = sourcesList[^1];
+                    sourcesList[^1] = context.BitwiseAnd(x, Const(0x3FFF));
+                    sourcesList.Add(context.ShiftRightU32(x, Const(14)));
+                    context.TranslatorContext.GpuAccessor.RecordTexelFetch1D(imm);
+                }
+                else
+                {
+                    sourcesList.Add(ConstF(0));
+                }

                type = SamplerType.Texture2D | (type & SamplerType.Array);
                is1DTo2D = true;
@@ -496,7 +512,13 @@ namespace Ryujinx.Graphics.Shader.Instructions
                        {
                            if (Sample1DAs2D)
                            {
-                                sourcesList.Add(ConstF(0));
+                                // nxnative: Maxwell fetches a 1D texture and a texture buffer with the same instruction, and the descriptor decides. With no type answer the read is
+                                // 2D at (x & 0x3FFF, x >> 14): a 1D texture (a one-row 2D texture at most 16384 wide on the host) reads (x, 0) as before, and the host binds a texture
+                                // buffer in such a slot as a linear 2D texture of 16384-texel rows over the same memory (nvn_mtl.m tbuf_as_2d)
+                                Operand x = sourcesList[^1];
+                                sourcesList[^1] = context.BitwiseAnd(x, Const(0x3FFF));
+                                sourcesList.Add(context.ShiftRightU32(x, Const(14)));
+                                context.TranslatorContext.GpuAccessor.RecordTexelFetch1D(tldsOp.TidB);

                                type &= ~SamplerType.Mask;
                                type |= SamplerType.Texture2D;
@@ -83,13 +83,15 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog
 - **Virtual memory pools** (flag 0x200): aliased onto the physical pool named by `MapVirtual` when all requests share one base (Mario's case); otherwise unimplemented.
 - **SPIRV-Cross failures**: 2 agl post-effect compute shaders still produce SPIR-V that SPIRV-Cross can't parse (`nullptr` exception). Before the compute start-offset fix it was 7: misdecoded code produced the bad SPIR-V.
 - **CPU-written linear textures**: a texture with `NVN_TEXTURE_FLAGS_LINEAR` or `LINEAR_RENDER_TARGET` in a CPU-visible pool (Mario's glyph cache; its four 64×1 RGBA16F sky colour ramps, without which the sky is black) is created on the pool's `MTLBuffer` (`tex_linear`), so CPU writes (Mario's glyph cache) are visible to the GPU. Block-linear CPU-written textures still need the packaged-texture re-upload path.
- **Integer vertex inputs**: without PC reflection the translator declares every vertex input `float4`. Skinned models read their bone index (`R8UI` / `RG8UI` / `RGBA8UI`) as raw integer bits (`as_type<int>(in.m_N.x) * 48`); Metal converted the integers to floats, every bone index was garbage, and all characters were stretched into streaks. `dmsh_build.py` (`int_inputs`) now re-translates a vertex shader with `uint` inputs when every use of an input is an integer bit cast (529 of 4,854 in Mario), and the host makes the Metal vertex format agree with the declared type (`vfmt_for_ct`; signedness swaps of 32-bit components and 32-bit float→int reinterpretation keep the bits, while 8- and 16-bit integers of the other sign are widened on the CPU, see Three Houses below). A remaining mismatch is reported as `vertex attribute format … does not match shader input type …`.
+- **Integer vertex inputs**: without PC reflection the translator declares every vertex input `float4`. Skinned models read their bone index (`R8UI` / `RG8UI` / `RGBA8UI`) as raw integer bits (`as_type<int>(in.m_N.x) * 48`); Metal converted the integers to floats, every bone index was garbage, and all characters were stretched into streaks. `dmsh_build.py` (`int_inputs`) now re-translates a vertex shader with `uint` inputs when every use of an input is an integer bit cast (529 of 4,854 in Mario), and the host makes the Metal vertex format agree with the declared type (`vfmt_for_ct`; signedness swaps of 32-bit components and 32-bit float→int reinterpretation keep the bits, while float-type formats and 8- and 16-bit integers of the other sign go through shadow inputs, below). A remaining mismatch is reported as `vertex attribute format … does not match shader input type …`.
+- **Shadow vertex inputs** (`<key>_bv`, `dm/gfx/tools/dmsh_vin.py`): an input that `int_inputs` declares an integer may still be fed a format of another class. A float-type format (half, unorm / snorm, 10:10:10:2, 11:11:10) leaves the converted float32 bits in the register; an 8- or 16-bit integer of the other signedness is widened by the format's own sign (Three Houses below). Metal reads a format only into an input of its own class, so the replay used to convert those attributes on the CPU every draw (`vconv_bind`, red lines 89 and 96): about a quarter of Age of Calamity's replay thread in battle. The build now also writes a side library for each such vertex shader. In it, every integer input is declared again as a `float4` and as the integer of the other sign, at Metal attribute indices 16–30 (NVN uses 0–15), and the body starts with `in.m_N = dm_in.m_N | as_type<T>(dm_in.m_N_vf) | as_type<T>(dm_in.m_N_vo)`. The host sends each attribute to the index of its format's class and the zero buffer to the others (`vin_ct`, `vin_on` in `nvn_mtl.m`; pipeline keys unchanged), so a pipeline with any mix of formats reads Maxwell's register bits. The float shadows are assigned first: with more than 7 integer inputs the last ones get only the float shadow (169 of 8,201 inputs in Age of Calamity). A format with no shadow keeps the CPU conversion, as do transform feedback draws (`main0_tf` has no variant) and `DM_NVN_VCONV_CPU=1`. Age of Calamity: 2,259 libraries, 77 MB each for Mac and iOS. The host loads a side library only for a pipeline that routes an attribute to a shadow, but the build cannot tell which shaders will need one (an input used only as an integer may hold bone indices or float bits), so it writes them all. Other games have 0–2,000 (at most 31 MB on iOS), except Monster Hunter Rise: 16,781 libraries, 344 MB on iOS (21% of its shader libraries, 1.5% of its IPA), and its Mac and iPhone runs so far load none. Its 8- and 16-bit integer formats feeding float inputs (red line 99, about 60,000 vertices per submission on its Mac route into the opening) still go through the CPU, since shadows cover only integer inputs. Age of Calamity's Mac battle at 60 fps against the CPU conversion: 0 vertices converted instead of 545,000 per submission, submit + present 10.0 → 5.7 ms a frame, CPU 3.97 → 2.44 W, GPU unchanged, the same picture.
 - **Vertex formats**: Mario uses `RGB10A2SN` (0x69) → `MTLVertexFormatInt1010102Normalized`. Unmapped formats fall back to a 4-component read and print `vertex attribute format … not mapped` once. The integer-read-as-float formats (`*_UI2F` / `*_I2F`, 0x51–0x68), which Metal has no vertex format for, are converted on the CPU at replay like the other conversions (`vconv_i2f`); Breath of the Wild has an `RG8_UI2F` attribute.
 - **Instanced array draws** (`DrawArraysInstanced`) carry instance count and base instance in the `OP_DRAW` record.
 - **Metal 4 does not order operations inside one encoder.** The host batches consecutive copies, fills and inline uniform updates (`UpdateUniformBuffer`) into one compute encoder. Mario writes the same uniform slot twice in a row each frame (4 bytes of −1, then 16 bytes of wave coefficients); without an intra-encoder barrier the writes occasionally landed in reverse order, the fountain's GPU wave simulation (a 512×1 RGBA16F texture updated by render → copy back → splash pass every frame) ran with −1 as its self coefficient, diverged to ±65504, and its vertex shader pushed the water surface to the top of the screen (tall blue / textured panels that flickered frame to frame). `blit_buf` now inserts `barrierAfterEncoderStages` when an operation reads or writes a guest range an earlier operation in the same encoder wrote; texture copies always get one. `DM_NVN_BLIT_ORDER=0` turns the barriers off to reproduce. Render passes have their own encoder with full barriers, so they were not affected; compute dispatches share the copy encoder and always get a barrier, as texture copies do (Age of Calamity below).
 - **A texture slot picked at run time** (Octopath Traveler 0: 61 Unreal vertex shaders, particles reading their data from buffer textures) → the shader reads one of several c2 texture-table words, chosen by branches, and ORs it with a sampler-table word. The translator cannot trace that to one slot and falls back to Ryujinx's bindless arrays, `array<texture_buffer<float>, 32> vp_tb_cb0_40000000` and `array<sampler, 32> vp_s_cb0_40000000`, indexed by the handle's texture ID (low 20 bits) and sampler ID (high 12), capped at 31; the build leaves them without a slot and they used not to load. With the profile key `tex_dyn_slots` the reflection reads them as `NVN_RB_TEXD` / `NVN_RB_SAMPD` arrays (`nvn_refl.m`), the replay writes slot indices into that shader's c2 tables (texture slot r → r, sampler slot r → r << 20) and binds slot k at element k, so the shader's own arithmetic lands on the binding. Check first that the handles come only from c2's tables: in Octopath Traveler 0 every integer read of c2 in those shaders is a table slot (texture slots 0-10, samplers 0-4). The host logs each such shader's first encode with the slots that held textures. Host only: the chain version does not change.
 - **Image atomics** → SPIRV-Cross's MSL 2.3 emulation declares the image twice: the texture (`texture_buffer<uint> cp_i_cb2_4C_r32ui`) and a `device atomic_uint* cp_i_cb2_4C_r32ui_atomic` pointer to the same memory, which the build read as a second image of the same slot, so the replay put a texture resource ID where the shader wanted an address. Octopath Traveler 0 has 2,084 such shaders (Unreal's GPU particles); three ran in the first 100 seconds. The pointer is now its own kind (`NVN_RB_IMGA`), bound to the storage of the image in that slot: a buffer texture is a view of its memory pool, so the atomics and the texture reads see the same words. A block-linear image is a deswizzled Metal copy whose pool memory has another layout, so its atomics go to a dummy buffer (logged red) until a game needs them followed; a 3D image does not compile at all (above).
 - **Buffer textures at offsets below 16 bytes** → Metal reports a 16-byte alignment for texture buffer views, and the host used to leave a view at a smaller offset as a dummy (Octopath Traveler 0: one at pool offset `0x168`, read by the particle shaders above). On an Apple GPU a view at any multiple of the element size reads the right elements (measured on an M6 with R32Float views at offsets 4, 8 and 12), so the host now creates it there; below the element size it stays a dummy.
+- **Buffer textures read by 1D integer fetches (row wrap)** → Maxwell fetches a texture buffer and a 1D texture with the same instructions (TLD with integer coordinates only, TLDS `Texture1DLodZero`), and the bound descriptor decides which it is. Without PC reflection the translator has no type answer and reads such a slot as a 2D texture at `(x, 0)`, and the host could only bind a dummy when the game put a buffer texture there: the slot read 0 (`texture type mismatch: shader declares 2, texture is 9`). Age of Calamity has 16,438 such shaders (98,970 bindings; slots 16–21 hold KTGL's light and decal lists). The vendored translator (`InstEmitTexture.cs`, marked `nxnative:`) now addresses every such fetch as `(x & 0x3FFF, x >> 14)` and reports the slot (`fetch1d` in the reflection). The host binds a buffer texture in a flagged slot as a linear 2D texture of 16,384-texel rows over the same pool memory (`tbuf_as_2d` in `nvn_mtl.m`, `nvn_tbuf_rows` in `nvn_tex.c`), cached by texture serial beside the views. A real 1D texture, which the host creates as a single row at most 16,384 wide, reads as before. Limits: a size query on a buffer of more than 16,384 elements returns 16,384 (no shader both size-queries and fetches such a slot in Age of Calamity; the build summary counts them), and the rows stop at the end of the Metal buffer. A mismatch no view can bridge is logged once per shader and slot, with its texture ID, NVN target and formats (`tex_mis_note`). Age of Calamity's 15-minute Mac battle bound 15.7 million of them through 320 row textures, none failed; the scene is darker and more saturated than with the zeros (mean luma 163 against 178 at the start of the battle), GPU time unchanged.
 - **New textures start as zero.** Metal leaves a new texture's content undefined; a freshly mapped Switch memory pool reads as 0, and feedback effects (such as the wave simulation above) never clear their state texture. Non-packaged colour textures are zero-filled from a shared zero buffer when created (`tex_zero_fill`, batched with the uploads).
 - **Scissored colour clears**: NVN `ClearColor` honours the scissor. Clears are normally folded into the next pass's load action, which only works for a full-target clear; Mario clears sub-rectangles of small targets to zero (pyramid levels packed into 480×272, 256×256 and 64×36 targets). A zero clear outside a pass with a partial scissor is now a copy from the shared zero buffer into that rectangle of the attachment's level / slice (`clear_rect_zero`). `DM_NVN_CLEARLOG=1` logs clears that are still skipped (mid-pass, partial channel mask or non-zero colour with a partial scissor).
 - **Inline uniform updates are snapshotted at encode time** (`mtl_sub_update`): the data is copied into a host-owned ring instead of being read from command memory when the GPU runs.
@@ -171,9 +173,9 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog
 ### Seen in Fire Emblem: Three Houses

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

@@ -246,6 +248,8 @@ Game Freak's Trinity engine on SDK 19 (NVN 55); every shader comes from the Oodl
 | `DM_NVN_LAZY_CLEAR=0` | Clears with no draw after them open a pass of their own when the game sets other targets, as before. By default they wait: the next pass that attaches the same texture level and layer through the same view clears it as its load action, and the rest land in passes of their own before anything else could see their textures (a draw sampling one, a pass not attaching it, any operation other than targets, clears, state, uniform updates and buffer copies or fills, the end of the command list). Off in TBDR mode. The `residency set` line counts clears taken by a later pass and those landed alone |
 | `DM_NVN_BARRIER=game` / `full` | How draw passes are ordered. `full` (default): every pass waits for all earlier GPU work. `game`: a pass waits only for the previous pass's fragment stage, so its vertex work overlaps the previous pass's fragments, unless the game called `nvnCommandBufferBarrier` in between or a compute / copy / upload ran (those keep full barriers). Only vertex-stage reads of the previous pass's output without a game barrier change order |
 | `DM_NVN_GPU_TS=segment` / `submission` | What timestamp reports record, instead of the profile's `gpu_timestamps`: the GPU start / end of the segment holding each report, or one time for the whole submission |
+| `DM_NVN_DISPATCH_OWN=1` | Every compute dispatch in a compute encoder of its own with full barriers, as before (by default dispatches share the copy encoder, `dispatch_enc`) |
+| `DM_NVN_VCONV_CPU=1` | Integer vertex inputs fed formats of another class are converted on the CPU, as before (by default the shader's `<key>_bv` library takes them through shadow inputs, `vin_on`) |
 | `MTL_DEBUG_LAYER=1 MTL_DEBUG_LAYER_ERROR_MODE=nslog` | Metal validation without aborting |

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