A uniform word tested against constants is read by a volatile load per test, so Metal's front end folds no switch that Apple's GPU compiler takes hours over

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 "d24f05cf9d78"
+#define DMSH_CHAIN "a9bb9ce825a2"

dm/gfx/tests/test_dmsh_uniform_tests.py

@@ -0,0 +1,131 @@
+"""A game uniform word tested against constants (dmsh_rewrite.py uniform_tests): each test of a word tested more than once reads it with a volatile load (dm_uword), so the Metal
+front end folds no chain of tests into a switch, which Apple's GPU compiler took hours and tens of gigabytes over (Nintendo Switch Sports Resort's pixel shader 692008b305ce264f in
+the offline pipeline archive). The shader below is the translator's form of a GLSL switch on one word; with dmsh_build.py's and ios_metal.py's compile commands its plain form
+becomes a switch in the front end's IR and its rewritten form none."""
+import ast
+from pathlib import Path
+import shutil
+import subprocess
+import sys
+import tempfile
+import unittest
+
+ROOT = Path(__file__).resolve().parents[3]
+TOOLS = ROOT / 'dm/gfx/tools'
+sys.path.insert(0, str(TOOLS))
+sys.path.insert(0, str(ROOT / 'dm/ios/tools'))
+import dmsh_chain  # noqa: E402
+import dmsh_rewrite  # noqa: E402
+import ios_metal  # noqa: E402
+
+WORD = '(*spvDescriptorSet0.fp_c5)._m0[16]'
+
+
+def mac_metal():   # dmsh_build.py's MAC_METAL, read without running the script (its one name, dmsh_chain.MAC_MIN_FLAG, from dmsh_chain)
+    for n in ast.parse((TOOLS / 'dmsh_build.py').read_text()).body:
+        if isinstance(n, ast.Assign) and any(isinstance(t, ast.Name) and t.id == 'MAC_METAL' for t in n.targets):
+            return eval(compile(ast.Expression(n.value), 'dmsh_build.py', 'eval'), {'__builtins__': {}, 'dmsh_chain': dmsh_chain})
+    raise AssertionError('no MAC_METAL in dmsh_build.py')
+
+
+# As spirv-cross writes Ryujinx's structured switch on c5[16].w (the cases 100 to 103 of a material's mode), cut down; the last test with the constant first
+SHADER = '''#include <metal_stdlib>
+using namespace metal;
+
+struct _16
+{
+    float4 _m0[4096];
+};
+
+struct spvDescriptorSetBuffer0
+{
+    constant _16* fp_c5 [[id(0)]];
+};
+
+struct main0_out
+{
+    float4 out_var_SV_Target [[color(0)]];
+};
+
+fragment main0_out main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], float4 gl_FragCoord [[position]])
+{
+    main0_out out = {};
+    float _10 = 0.0;
+    bool _11 = as_type<int>(%(w)s.w) != 100;
+    if (!_11)
+    {
+        _10 = gl_FragCoord.x;
+    }
+    else
+    {
+        bool _12 = as_type<int>(%(w)s.w) != 101;
+        if (!_12)
+        {
+            _10 = sin(gl_FragCoord.y);
+        }
+        else
+        {
+            bool _13 = as_type<int>(%(w)s.w) != 102;
+            if (!_13)
+            {
+                _10 = cos(gl_FragCoord.x);
+            }
+            else
+            {
+                bool _14 = 103 == as_type<int>(%(w)s.w);
+                if (_14)
+                {
+                    _10 = exp2(gl_FragCoord.y);
+                }
+            }
+        }
+    }
+    out.out_var_SV_Target = float4(_10 * %(w)s.x);
+    return out;
+}
+''' % {'w': WORD}
+
+
+class UniformTests(unittest.TestCase):
+    def test_tested_words_are_volatile_loads(self):
+        got = dmsh_rewrite.uniform_tests(SHADER)
+        self.assertEqual(got.count('as_type<int>(dm_uword(%s, 3))' % WORD), 4, 'every test, the constant on either side')
+        self.assertIn('_10 * %s.x' % WORD, got, 'a read in arithmetic is left as it is')
+        self.assertEqual(got.count('#define DM_UWORD'), 1)
+        self.assertLess(got.index('dm_uword(const constant'), got.index('struct _16'), 'the helper before the shader')
+        self.assertEqual(dmsh_rewrite.uniform_tests(got), got, 'a second pass changes nothing')
+
+    def test_left_as_it_is(self):
+        once = 'using namespace metal;\nbool _1 = as_type<int>(%s.w) != 100;\nbool _2 = as_type<int>(%s.z) != 100;\n' % (WORD, WORD)
+        self.assertEqual(dmsh_rewrite.uniform_tests(once), once, 'each word tested once: no chain')
+        other = 'using namespace metal;\nbool _1 = as_type<int>(%s.w) != _5;\nbool _2 = as_type<int>(%s.w) >= 3;\nbool _3 = _6 == as_type<int>(%s.w);\n' % (WORD, WORD, WORD)
+        self.assertEqual(dmsh_rewrite.uniform_tests(other), other, 'tests against no constant, or not of equality')
+        self.assertEqual(dmsh_rewrite.uniform_tests('bool _1 = _2 != 3;'), 'bool _1 = _2 != 3;', 'no descriptor set')
+
+    def test_unsigned_and_negative(self):
+        src = 'using namespace metal;\nbool _1 = as_type<uint>(%s.y) == 4u;\nbool _2 = as_type<uint>(%s.y) != 5u;\nbool _3 = -1 != as_type<int>(%s.x);\nbool _4 = as_type<int>(%s.x) == -2;\n' % ((WORD,) * 4)
+        got = dmsh_rewrite.uniform_tests(src)
+        self.assertEqual(got.count('as_type<uint>(dm_uword(%s, 1))' % WORD), 2)
+        self.assertEqual(got.count('as_type<int>(dm_uword(%s, 0))' % WORD), 2)
+
+    def test_needs_the_namespace(self):
+        with self.assertRaisesRegex(ValueError, 'using namespace metal'): dmsh_rewrite.uniform_tests(SHADER.replace('using namespace metal;\n', ''))
+
+    def switches(self, command, source):   # the front end's IR of source with command (its -c compile, emitting LLVM IR instead): how many switch instructions main0 has
+        with tempfile.TemporaryDirectory() as tmp:
+            root = Path(tmp)
+            (root / 'test.metal').write_text(source)
+            r = subprocess.run(command + ['-w', '-S', '-emit-llvm', str(root / 'test.metal'), '-o', str(root / 'test.ll')], capture_output=True, text=True)
+            self.assertEqual(r.returncode, 0, r.stderr)
+            return (root / 'test.ll').read_text().count(' switch ')
+
+    @unittest.skipUnless(shutil.which('xcrun'), 'needs Xcode and Metal')
+    def test_no_switch_for_the_gpu_compiler(self):
+        for name, command in (('Mac', mac_metal()), ('iPhone', ios_metal.METAL)):
+            with self.subTest(platform=name):
+                self.assertGreater(self.switches(command, SHADER), 0, 'the plain tests fold into a switch: the test no longer shows the problem')
+                self.assertEqual(self.switches(command, dmsh_rewrite.uniform_tests(SHADER)), 0)
+
+
+if __name__ == '__main__':
+    unittest.main()

dm/gfx/tools/dmsh_build.py

@@ -692,7 +692,7 @@ def geometry_finish(key, w, j, c1):
         dmsh_geometry.flatten_sets(spv)
         r = subprocess.run([SPIRV_CROSS, *dmsh_atomic.cross_args(XC, spv), spv], capture_output=True, text=True)
         if r.returncode: return 'geometry MSL failed', r.stderr.strip()[:160], None
-        try: selected = dmsh_fetch.select_guards(dmsh_rewrite.precise(dmsh_atomic.check_msl(r.stdout)))   # bounded uniform reads as selects (see dmsh_fetch.select_guards)
+        try: selected = dmsh_rewrite.uniform_tests(dmsh_fetch.select_guards(dmsh_rewrite.precise(dmsh_atomic.check_msl(r.stdout))))   # bounded uniform reads as selects (see dmsh_fetch.select_guards), tested words as volatile loads (see uniform_tests)
         except ValueError as e: return 'geometry MSL failed', str(e)[:160], None
         plain = dmsh_geometry.topology_pointer(dmsh_reflect.pool_arrays(selected))
         src = dmsh_geometry.capture_input(plain)
@@ -724,7 +724,7 @@ def save(suffix, text, meta):
             spv = os.path.join(w, k + '.reader.spv'); dmsh_geometry.flatten_sets(spv)
             r = subprocess.run([SPIRV_CROSS, *dmsh_atomic.cross_args(XC, spv), spv], capture_output=True, text=True)
             if r.returncode: return 'geometry reader failed', r.stderr.strip()[:160], None
-            try: selected = dmsh_fetch.select_guards(dmsh_rewrite.precise(dmsh_atomic.check_msl(r.stdout)))
+            try: selected = dmsh_rewrite.uniform_tests(dmsh_fetch.select_guards(dmsh_rewrite.precise(dmsh_atomic.check_msl(r.stdout))))
             except ValueError as e: return 'geometry reader failed', str(e)[:160], None
             reader = dmsh_entry.vs_layer(dmsh_entry.vs_init(dmsh_geometry.reader(dmsh_reflect.pool_arrays(selected))))
             analysis, full = dmsh_entry.vs_depth(reader)
@@ -783,6 +783,8 @@ def finish(key, w, j, c1):
     src = dmsh_rewrite.tg_fence(src)   # warp-synchronous shared memory: a simdgroup-scope fence after each threadgroup store (see tg_fence)
     src = dmsh_rewrite.border_emu(src)   # border colours Metal lacks (see border_emu)
     if j['stage'] == 'ps': src = dmsh_rewrite.frag_z_snap(src)[0]   # gl_FragCoord.z against a depth of its own surface (see frag_z_snap)
+    try: src = dmsh_rewrite.uniform_tests(src)   # a uniform word tested against constants read by a volatile load per test, so the front end folds no switch (see uniform_tests)
+    except ValueError as e: return 'uniform tests rewrite failed', str(e)[:160], None
     blend0 = 0
     if j['stage'] == 'ps': src, blend0 = dmsh_rewrite.blend_zero(src)   # a zero blend weight makes 0 of a NaN colour, as NVIDIA's blender does; before the entry points that share main0's body (see blend_zero)
     open(msl, 'w').write(src)

dm/gfx/tools/dmsh_rewrite.py

@@ -1,9 +1,10 @@
 #!/usr/bin/env python3
 # The rewrites dmsh_build.py makes in a translated shader's MSL before it compiles it, where Metal or the translator falls short of what the Switch code does: precise
-# arithmetic (precise), memory barriers as fences (membar_fence, tg_fence), 64-bit products (mul_extended), Ampere's HMMA as simdgroup matrices (hmma_msl), first stores and shared-memory stores (first_store,
-# shared_store_msl), threadgroup memory past Metal's limit (tg_spill), texture gathers with offsets (gather_offsets), border colours Metal has no sampler state for
-# (border_emu) and a pixel shader's depth snapped to its position (frag_z_snap). Functions of the MSL text and the library ABI only, with no build state, so tests import
-# them. Part of the translation chain (dmsh_chain.FILES) and of every shader's cache key (dmsh_build.py CODE_FP).
+# arithmetic (precise), memory barriers as fences (membar_fence, tg_fence), 64-bit products (mul_extended), uniform words tested against constants, which the front end would
+# fold into switches (uniform_tests), Ampere's HMMA as simdgroup matrices (hmma_msl), first stores and shared-memory stores (first_store, shared_store_msl), threadgroup memory
+# past Metal's limit (tg_spill), texture gathers with offsets (gather_offsets), border colours Metal has no sampler state for (border_emu) and a pixel shader's depth snapped to
+# its position (frag_z_snap). Functions of the MSL text and the library ABI only, with no build state, so tests import them. Part of the translation chain (dmsh_chain.FILES)
+# and of every shader's cache key (dmsh_build.py CODE_FP).
 import bisect, collections, itertools, re
 import dmsh_abi as abi, dmsh_side
 
@@ -54,6 +55,38 @@ 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 word tested against integer constants, as the translator writes a GLSL switch (Ryujinx's structurizer: one test of the word per case, flags for the fall-throughs,
+# the cases nested in each other): bool _5 = as_type<int>((*spvDescriptorSet0.fp_c5)._m0[16].w) != 102. The Metal front end's SimplifyCFG (at -O1 too) folds a chain of blocks
+# that each hold only such a test of one value into a switch, whose cases jump into the nested blocks past their joins. Apple's GPU compiler then takes a pixel shader with dozens of
+# such switches for hours and tens of gigabytes (applegpu-nt, metal-tt 32023.922): Nintendo Switch Sports Resort's material pixel shader 692008b305ce264f (42 switches) ran past
+# 10 minutes and 13 GB alone where it had taken 7 s, and the Mac app's pipeline archive spent 15 hours on four pipelines. The reads were plain since the host pads a short binding
+# (dmsh_fetch.select_guards): before that, each read's bounds test sat in the test's block and nothing was folded. Each tested read becomes a volatile load of its word
+# (dm_uword), which the front end can neither merge with the next test's nor move out of its block, so no chain is folded; a uniform load per test is less than the bounds test
+# each read had then. Only words tested against a constant more than once: a single test is no chain. Returns the shader unchanged when there are none
+_UTEST_READ = r'as_type<(u?int)>\((\(\*spvDescriptorSet0\.\w+_c\d+\)\._m0\[\d+\])\.([xyzw])\)'
+_UTEST = re.compile(r'(?<![\w.])%s(?= [!=]= -?\d+u?(?![\w.]))|(?<=[!=]= )%s' % (_UTEST_READ, _UTEST_READ.replace('(u?int)', '(?P<t2>u?int)').replace(r'(\(\*', r'(?P<r2>\(\*').replace('([xyzw])', '(?P<c2>[xyzw])')))
+_UTEST_LEFT = re.compile(r'(?<![\w.])-?\d+u? [!=]= $')
+UWORD = ('#ifndef DM_UWORD\n#define DM_UWORD\ntemplate<typename T, int N> static inline __attribute__((always_inline)) T dm_uword(const constant vec<T, N>& v, int i) '
+         '{ return ((const volatile constant T*)&v)[i]; }\n#endif\n')
+def uniform_tests(msl):
+    if 'spvDescriptorSet0' not in msl or 'as_type<' not in msl: return msl
+    def parts(m):
+        if m.group(1): return m.group(1), m.group(2), m.group(3)
+        if not _UTEST_LEFT.search(msl, max(0, m.start() - 16), m.start()): return None   # [!=]= after something other than a constant
+        return m.group('t2'), m.group('r2'), m.group('c2')
+    found = [(m, p) for m in _UTEST.finditer(msl) for p in [parts(m)] if p]
+    tested = collections.Counter(r + c for _, (_, r, c) in found)
+    chains = {w for w, n in tested.items() if n > 1}
+    if not chains: return msl
+    out, at = [], 0
+    for m, (t, r, c) in found:
+        if r + c not in chains: continue
+        out += [msl[at:m.start()], 'as_type<%s>(dm_uword(%s, %d))' % (t, r, 'xyzw'.index(c))]; at = m.end()
+    body = ''.join(out) + msl[at:]
+    head = re.search(r'^using namespace metal;\n', body, re.M)
+    if not head: raise ValueError('uniform tests: no "using namespace metal;" to put dm_uword after')
+    return body[:head.end()] + UWORD + body[head.end():]
+
 # 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

@@ -955,6 +955,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - pixel shader compilation fails with `reference type must have explicit address space qualifier` at `dm_blend0(main0_out& dm_o, ...)`: every iOS build, and a Mac build made on macOS 26. A Mac build made on macOS 27 compiled, so the fault passed where it was written. → [The blend helper's output reference on iOS](../reference/shaders.md#the-blend-helpers-output-reference-on-ios)
 - dark specks flickering in a ground mist, a few pixels each, at different places every frame (Ocarina of Time in Kokiri Forest; the console shows none). Dropping the mist removes them, as does dropping its depth-only pass. → [gl_FragCoord.z against the depth stored for its own surface](../reference/shaders.md#gl_fragcoordz-against-the-depth-stored-for-its-own-surface)
 - making a pipeline archive prints `Max Unstruct ID 758 in function agc.main in shader 13941 exceeds supported limit`, and one pipeline takes many minutes, in `metal-tt` offline or in the host's prebuild on the device (`Preparing pipelines` holding one count for minutes). The line is a note, not an error: `metal-tt` goes on, exits 0 and writes that pipeline into the archive. → [Front end at -O1: Apple's control-flow lowering](../reference/shaders.md#front-end-at--o1-apples-control-flow-lowering)
+- the Mac app's (or the iPhone build's) pipeline archive step runs for many hours: a group of 1,000 bases runs past its time limit, and `compile_split` leaves out single pipelines that each ran past theirs (`[RED] pso_offline: applegpu_<gpu>: left out of the archive, the device builds it`), while `air-nt` grows past the Mac's memory and swap fills. Alone, such a pipeline does not finish in 10 minutes and takes 10 to 13 GB, where the build before took seconds. On the device, the same pipelines would hold `Preparing pipelines` for as long. → [Tests of one uniform word folded into switches](../reference/shaders.md#tests-of-one-uniform-word-folded-into-switches)
 - the shader build prints `fatal error: bracket nesting level exceeded maximum of 256` and `<key>.metal: does not compile`, and the shader is left out. → [Bracket nesting past the front end's 256](../reference/shaders.md#bracket-nesting-past-the-front-ends-256)
 - the half of Mario facing the camera's left was white (GitHub issue 72, on an iPhone). → [min, max and clamp with a NaN operand](../reference/shaders.md#min-max-and-clamp-with-a-nan-operand)
 - a pipeline red line for a float output meeting an integer attachment (the red line prints the row to add); a BC encoder's pipeline fails. → [Integer fragment outputs declared by hand](../reference/shaders.md#integer-fragment-outputs-declared-by-hand)