Integer reads of game uniforms are volatile loads: bare, they made Metal's backend take minutes and run out of memory on Switch Sports Resort's pipelines

dm/gfx/nvn/dmsh_chain.h

@@ -1,3 +1,3 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header): the B shader library translation chain version this source tree expects (default options), one for every game;
 // the host loads only <library dir>/b-<V>/. Run dmsh_chain.py --header after changing the chain's tools and commit this file with the change
-#define DMSH_CHAIN "7c609af36b0f"
+#define DMSH_CHAIN "3695a7c258bb"

dm/gfx/tests/test_dmsh_flag_reads.py

@@ -0,0 +1,60 @@
+#!/usr/bin/env python3
+# dmsh_rewrite.flag_reads: a game uniform read as an integer at a fixed offset goes through dm_vl (a volatile load), so Metal's backend keeps each read a load of its own
+# (docs/reference/shaders.md, Integer reads of game uniforms and Metal's compile time). Float reads, the driver's banks and computed offsets are left as they are
+from pathlib import Path
+import shutil
+import subprocess
+import sys
+import tempfile
+import unittest
+
+TOOLS = Path(__file__).resolve().parent.parent / 'tools'
+sys.path.insert(0, str(TOOLS))
+from dmsh_rewrite import flag_reads   # noqa: E402
+
+HEAD = '#include <metal_stdlib>\nusing namespace metal;\n'
+SHADER = HEAD + '''struct _10 { float4 _m0[4096]; };
+struct spvDescriptorSetBuffer0 { constant _10* fp_c0 [[id(0)]]; constant _10* fp_c5 [[id(1)]]; constant _10* fp_c12 [[id(2)]]; };
+struct main0_out { float4 m_0 [[color(0)]]; };
+fragment main0_out main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
+{
+    main0_out out = {};
+    bool _1 = 0 != as_type<int>((*spvDescriptorSet0.fp_c5)._m0[40].x);
+    uint _2 = as_type<uint>((*spvDescriptorSet0.fp_c12)._m0[3].w);
+    float _3 = (*spvDescriptorSet0.fp_c5)._m0[41].y;
+    int _4 = as_type<int>((*spvDescriptorSet0.fp_c0)._m0[2].z);
+    int _5 = as_type<int>((*spvDescriptorSet0.fp_c5)._m0[int(_2)].x);
+    out.m_0 = float4(_1 ? _3 : 0.0, float(_2), float(_4), float(_5));
+    return out;
+}
+'''
+
+
+class FlagReads(unittest.TestCase):
+    def test_integer_reads_of_game_slots_are_volatile(self):
+        out = flag_reads(SHADER)
+        self.assertIn('as_type<int>(dm_vl((*spvDescriptorSet0.fp_c5)._m0[40]).x)', out)
+        self.assertIn('as_type<uint>(dm_vl((*spvDescriptorSet0.fp_c12)._m0[3]).w)', out)
+        self.assertEqual(out.count('dm_vl('), 3)   # the two reads and the helper
+        self.assertLess(out.index('using namespace metal;'), out.index('T dm_vl(constant T& r)'))
+
+    def test_other_reads_stay(self):
+        out = flag_reads(SHADER)
+        self.assertIn('float _3 = (*spvDescriptorSet0.fp_c5)._m0[41].y;', out)                 # a float read
+        self.assertIn('as_type<int>((*spvDescriptorSet0.fp_c0)._m0[2].z)', out)                 # a driver bank
+        self.assertIn('as_type<int>((*spvDescriptorSet0.fp_c5)._m0[int(_2)].x)', out)           # a computed offset
+
+    def test_a_shader_with_none_is_unchanged(self):
+        src = HEAD + 'fragment float4 main0() { return float4(0.0); }\n'
+        self.assertEqual(flag_reads(src), src)
+
+    @unittest.skipUnless(shutil.which('xcrun') and subprocess.run(['xcrun', '-sdk', 'macosx', '-f', 'metal'], capture_output=True).returncode == 0, 'no Metal compiler')
+    def test_compiles(self):
+        with tempfile.TemporaryDirectory() as d:
+            f = Path(d, 's.metal'); f.write_text(flag_reads(SHADER))
+            r = subprocess.run(['xcrun', '-sdk', 'macosx', 'metal', '-c', '-O1', '-fpreserve-invariance', '-fwrapv', str(f), '-o', str(Path(d, 's.air'))], capture_output=True, text=True)
+            self.assertEqual(r.returncode, 0, r.stderr)
+
+
+if __name__ == '__main__':
+    unittest.main()

dm/gfx/tools/dmsh_build.py

@@ -757,6 +757,8 @@ def finish(key, w, j, c1):
     src = dmsh_rewrite.precise(src)   # min / max / clamp with Maxwell's NaN results (see precise)
     try: src = dmsh_fetch.select_guards(src)   # bounded uniform reads as a load and a select, as the translator means them, not a branch for each (see dmsh_fetch.select_guards)
     except ValueError as e: return 'uniform bounds rewrite failed', str(e)[:160], None
+    try: src = dmsh_rewrite.flag_reads(src)   # integer reads of game uniforms as loads of their own, which Metal's backend compiles in seconds (see flag_reads)
+    except ValueError as e: return 'flag read rewrite failed', str(e)[:160], None
     src = dmsh_rewrite.mul_extended(src)   # scalar extended products without the optnone wrapper (see mul_extended)
     try: src = dmsh_rewrite.gather_offsets(src)   # four gathers instead of a wrapper's loop over their offsets (see gather_offsets)
     except ValueError as e: return 'gather offsets rewrite failed', str(e)[:160], None

dm/gfx/tools/dmsh_rewrite.py

@@ -24,6 +24,22 @@ def mul_extended(msl):
                      for t, w in (('uint', 'ulong'), ('int', 'long')))
     return msl[:found[0].end()] + scalar + msl[found[0].end():]
 
+# A game uniform read as an integer at a fixed offset (as_type<int>((*spvDescriptorSet0.fp_c5)._m0[40].x): a material's flags and counts, which its branches test) is read
+# through a volatile pointer. A fixed read is a plain load since the host bounds it once a binding (GitHub issue #2066), and Metal's backend then takes every read of a word as
+# one value: on a material shader that tests the same flags at each level of deeply nested branches (Nintendo Switch Sports Resort: 455 such reads, 36 levels) its compile
+# went from 0.6 s to 79 s and several GB (fbe6bb186641245c with 0eb245e18bea1bac, M4 Pro), and pipelines compiled side by side ran MTLCompilerService out of memory
+# (XPC_ERROR_CONNECTION_INTERRUPTED, the pipeline archive file under way lost with them). Read volatile, each read stays a load of its own, as the compared read was, with no
+# size table and no comparison: 1.1 s. Float reads are left alone (they did not slow it, and dmsh_fetch reads their text); c0 to c2 are the driver's banks
+_FLAG_READ = re.compile(r'as_type<(u?int)>\(\(\*spvDescriptorSet0\.([a-z]+_c(?:[3-9]|\d\d))\)\._m0\[(\d+)\]\.([xyzw])\)')
+_FLAG_HELPER = 'template<typename T> static inline __attribute__((always_inline)) T dm_vl(constant T& r) { return *(constant volatile T*)&r; }\n'
+def flag_reads(msl):
+    out, n = _FLAG_READ.subn(r'as_type<\1>(dm_vl((*spvDescriptorSet0.\2)._m0[\3]).\4)', msl)
+    if not n: return msl
+    at = out.find('using namespace metal;\n')
+    if at < 0: raise ValueError('no "using namespace metal;" to put dm_vl after')
+    at += len('using namespace metal;\n')
+    return out[:at] + _FLAG_HELPER + out[at:]
+
 # Ampere HMMA.16816 with an f16 accumulator (agl's upscaler networks, the NDVE video networks) reaches the MSL as calls of dmsh_hmma16816_f16(lane, block, A0..A3, B0, B1, C0, C1,
 # flags, A's LDSM address, A0..A3's offsets, A's row stride, out D1) → D0, the subgroup's 256-word block of the scratch shared memory as the last argument (InstEmitAmpereTensor.AHmma,
 # TensorPass, HelperFunctionManager.GenerateHmma16816F16). The translator's body computes D lane by lane from the block; here it becomes simdgroup_matrix code over the same block:

docs/agents/symptoms.md

@@ -1049,6 +1049,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - `[RED] image atomics on a texture Metal cannot run them on`, and whatever the shader computes through them is lost: its writes go to a sink. → [Image atomics on a 32-bit colour texture](../reference/shaders.md#image-atomics-on-a-32-bit-colour-texture)
 - Metal shader validation reports invalid image-atomic stores even with an image-sized scratch buffer; without validation, GPU page faults can freeze rendering. → [Bounds of image atomics](../reference/shaders.md#bounds-of-image-atomics)
 - the GPU hangs (Metal command-buffer timeouts, `GPU access revoked`, exit 70; on the Mac WindowServer can lose its watchdog check-in and the session restarts) or a draw reads garbage, where a shader reads a uniform slot further than the game bound it. No `[RED]` line. → [Reads past a uniform binding's size](../reference/shaders.md#reads-past-a-uniform-bindings-size)
+- a game compiles pipelines at every launch, for minutes: the prebuild's `[RED] pipeline prebuild: the base of … failed: Compilation failed due to an interrupted connection: XPC_ERROR_CONNECTION_INTERRUPTED`, `MTLCompilerService` crash reports (`operator new` under `llvm::CloneBasicBlock`), and `[RED] pipeline archive: writing … failed: Failed to generate machO:warning: units in '…/com.apple.gpuarchiver/…/air64_v29' are already packed`, so the archive file under way is lost and its pipelines are built again at the next start. → [Integer reads of game uniforms and Metal's compile time](../reference/shaders.md#integer-reads-of-game-uniforms-and-metals-compile-time)
 - on an M1 Pro (macOS 26), blue, green and white blocks for one or two frames on trees that dither-fade (OCTOPATH TRAVELER's S'warkii, walking behind a tree); a texture dump shows NaN pixels in the scene colour and holes in the normal G-buffer target, which depth of field spreads into blocks. Not seen on an iPhone or on newer Macs. On every GPU, a branch for each uniform word a shader reads. → [Bounded uniform reads written as select](../reference/shaders.md#bounded-uniform-reads-written-as-select)
 - the NSP script stops with `captured only 0 run-time shaders` although the game reaches its frame loop with precompiled shaders. → [A GLSL language preamble is not a complete shader](../reference/shaders.md#a-glsl-language-preamble-is-not-a-complete-shader)
 - shaders missing from the map that no game file holds. → [Built-in shaders compiled with the game's own GLSLC](../reference/shaders.md#built-in-shaders-compiled-with-the-games-own-glslc)

docs/reference/shaders.md

@@ -1838,6 +1838,15 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b
 - **Found on:** [Paper Mario: The Thousand-Year Door](../../dm/games/0100ECD018EBE000/README.md#rendering) ([GitHub issue #1860](https://github.com/seedds/nxnative/issues/1860)): the reporter found the variant and the short bindings, and padding those bindings with zeros let Peach's scene through on an M3 Ultra Mac and an M4 iPad Pro. [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md#found-on-this-game) ([GitHub issue #2066](https://github.com/seedds/nxnative/issues/2066)): the cost of comparing every read.
 - **Check:** `dm/tools/tests/test_dmsh_abi.py` `test_uniform_bounds_agree` (the translator's bank and slots are the header's); `dm/gfx/tests/test_dmsh_ampere_mem.py` `test_fixed_uniform_reads_are_bare_and_counted` (with `DMSH_TEST_XLATE`: two fixed reads of c[3] give no size table and `fixed` 28). On NieR:Automata's 1,225 shaders built with this change: all 148,022 reads at fixed offsets are bare, all 28,604 at computed offsets are compared (and `select_guards` writes each of them as a select, none refused), 455 shaders declare the size table, and each slot's `fixed` ends in the highest row its source reads at a constant index. Not run in a game yet (open): the host's padding has met no short binding, the potion scene's retest by the reporter of #1860 is open, and so is Bananza's frame time on the Freezer Layer. The first form's numbers, `bench.py` on NieR:Automata's opening flight (M6, two runs each, 60 s window): GPU time p50 6.43 → 6.53 ms (+1.6%), p95 9.12 → 9.70 ms (+6.4%), frame time unchanged at 60 fps. Changing this path changes the chain version: rebuild the translator, update the lock and rebuild the shaders.
 
+### Integer reads of game uniforms and Metal's compile time
+
+- **Symptom:** a game compiles pipelines at every launch, for minutes: the prebuild's `[RED] pipeline prebuild: the base of … failed: Compilation failed due to an interrupted connection: XPC_ERROR_CONNECTION_INTERRUPTED`, `MTLCompilerService` crash reports (`operator new` under `llvm::CloneBasicBlock`), and `[RED] pipeline archive: writing … failed: Failed to generate machO:warning: units in '…/com.apple.gpuarchiver/…/air64_v29' are already packed`, so the archive file under way is lost and its pipelines are built again at the next start.
+- **Cause:** a read of a game uniform at a fixed offset is a plain load since the host bounds it once a binding ([Reads past a uniform binding's size](#reads-past-a-uniform-bindings-size)). Metal's backend then takes every read of one word as one value. A material shader that tests the same integer flags at each level of deeply nested branches (Nintendo Switch Sports Resort's: 455 such reads of `c5`, branches 36 deep) sends its instruction selection into minutes and gigabytes: pixel shader `fbe6bb186641245c` with vertex shader `0eb245e18bea1bac` compiled in 0.6 s with the comparisons and in 79 s without (M4 Pro, macOS 27.2, the compiler's cache emptied before each). The same source with only its `c5` comparisons taken out is as slow; taking out those of `c6`, which it reads as floats, changes nothing. Several such compiles side by side ran the compiler service out of memory.
+- **Fix:** `dmsh_rewrite.flag_reads` reads a game uniform (`c3` and up) that the shader takes as an integer at a fixed offset through a volatile pointer (`as_type<int>(dm_vl((*spvDescriptorSet0.fp_c5)._m0[40]).x)`), so each read stays a load of its own, with no size table and no comparison. Float reads and the driver's banks are left as they are (chain change).
+- **Found on:** Nintendo Switch Sports Resort (`04008E001F098000`, v65536; no profile on main): the six pairs its prebuild failed on compile in 0.9 to 3.3 s each (1.1 s for the pair above; 3.3 s for `03e62a4a5e7ece89` with the 135,000-line `0d57ffc025e0df8f`, which had not finished after 150 s, the compiler at 3.9 GB).
+- **Not measured:** what the volatile loads cost on the GPU.
+- **Check:** `python3 dm/gfx/tests/test_dmsh_flag_reads.py`.
+
 ### Bounded uniform reads written as select
 
 - **Symptom:** on an M1 Pro (macOS 26), blue, green and white blocks for one or two frames on trees that dither-fade (OCTOPATH TRAVELER's S'warkii, walking behind a tree); a texture dump shows NaN pixels in the scene colour and holes in the normal G-buffer target, which depth of field spreads into blocks. Not seen on an iPhone or on newer Macs. On every GPU, a branch for each uniform word a shader reads.