diff --git a/dm/gfx/nvn/dmsh_chain.h b/dm/gfx/nvn/dmsh_chain.h index 3338ce799d..4b6f195fd4 100644 --- a/dm/gfx/nvn/dmsh_chain.h +++ b/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 /b-/. 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" diff --git a/dm/gfx/tests/test_dmsh_uniform_tests.py b/dm/gfx/tests/test_dmsh_uniform_tests.py new file mode 100644 index 0000000000..4b315ff7de --- /dev/null +++ b/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 +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(%(w)s.w) != 100; + if (!_11) + { + _10 = gl_FragCoord.x; + } + else + { + bool _12 = as_type(%(w)s.w) != 101; + if (!_12) + { + _10 = sin(gl_FragCoord.y); + } + else + { + bool _13 = as_type(%(w)s.w) != 102; + if (!_13) + { + _10 = cos(gl_FragCoord.x); + } + else + { + bool _14 = 103 == as_type(%(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(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(%s.w) != 100;\nbool _2 = as_type(%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(%s.w) != _5;\nbool _2 = as_type(%s.w) >= 3;\nbool _3 = _6 == as_type(%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(%s.y) == 4u;\nbool _2 = as_type(%s.y) != 5u;\nbool _3 = -1 != as_type(%s.x);\nbool _4 = as_type(%s.x) == -2;\n' % ((WORD,) * 4) + got = dmsh_rewrite.uniform_tests(src) + self.assertEqual(got.count('as_type(dm_uword(%s, 1))' % WORD), 2) + self.assertEqual(got.count('as_type(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() diff --git a/dm/gfx/tools/dmsh_build.py b/dm/gfx/tools/dmsh_build.py index 1e67117889..b94b709b58 100644 --- a/dm/gfx/tools/dmsh_build.py +++ b/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) diff --git a/dm/gfx/tools/dmsh_rewrite.py b/dm/gfx/tools/dmsh_rewrite.py index 6d19028db4..be6ece4cf2 100644 --- a/dm/gfx/tools/dmsh_rewrite.py +++ b/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((*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'(?u?int)').replace(r'(\(\*', r'(?P\(\*').replace('([xyzw])', '(?P[xyzw])'))) +_UTEST_LEFT = re.compile(r'(? static inline __attribute__((always_inline)) T dm_uword(const constant vec& 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: diff --git a/docs/agents/symptoms.md b/docs/agents/symptoms.md index 20128beeed..13758a1991 100644 --- a/docs/agents/symptoms.md +++ b/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_: 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 `.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) diff --git a/docs/reference/shaders.md b/docs/reference/shaders.md index 8d86dd9253..2b3123f9d2 100644 --- a/docs/reference/shaders.md +++ b/docs/reference/shaders.md @@ -805,6 +805,14 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b - **Found on:** [Super Mario Odyssey](../../dm/games/0100000000010000/README.md#files-and-shaders), [issue #276](https://github.com/seedds/nxnative/issues/276): the float body of `alRenderMaterialUber` (fragment `007d7e411ef14ebd`, 16,227 `if`s in one function) took 24.6 minutes and 17.3 MB of `g13s` code at `-O2`, 31 s and 1.3 MB at `-O1`. Its `-O1` graph still goes past the budget (642), at that small cost. - **Check:** random samples of archive bases from seven games (Odyssey, Tears of the Kingdom, Breath of the Wild, Kirby, Luigi's Mansion 3, Wonder, Pikmin 4; 22 to 400 each) compile at `-O1` to vertex and fragment code within a few percent of `-O2` (Odyssey 400 bases: −0.9% / +0.01%; TotK: +0.1% / −1.0%; Kirby: −3.1% / +0.07%), with no `Max Unstruct ID` note at either level, and the front end is faster. Odyssey's full `g13s` archive of 32,864 bases builds in 15.6 minutes (exit 0); its Cap Kingdom frames 4,200 to 6,600 match `-O2` except for scattered differences of at most 3 of 255 in three of nine frames, and their GPU time per frame (`DM_NVN_GPUPROF=every:150`, three runs each on an M-series `g18g`) has medians of 4.76 to 5.05 ms at `-O1` and 4.75 to 4.83 ms at `-O2`, within run-to-run noise. +### Tests of one uniform word folded into switches + +- **Symptom:** 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_: 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. +- **Cause:** a game's material shader chooses a branch by a uniform word, and the translator writes the GLSL switch as Ryujinx's structurizer gives it: one test of the word for each case, with flags for the fall-throughs and each case nested in the next (`bool _6240 = as_type((*spvDescriptorSet0.fp_c5)._m0[16].w) != 102; if (!_6240) { … }`). The Metal front end's SimplifyCFG (which runs at [-O1](#front-end-at--o1-apples-control-flow-lowering) 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, and Apple's GPU compiler (`applegpu-nt`, AIR-TT 32023.922) takes that graph for hours. Turning the switches back into compares (`metal-opt -passes=lowerswitch`) does not help: the graph is the cost, not the instruction. Before the host bounded [reads at fixed offsets](#reads-past-a-uniform-bindings-size) (chain `b7249b1c2b66`), each read's bounds test sat in the test's own block, so no chain was folded; with the reads plain, Nintendo Switch Sports Resort's pixel shader `692008b305ce264f` has 42 switches in its IR (none before) and `7c2becb054a53765` 66 (6 before). +- **Fix:** every game, a chain change. `dmsh_rewrite.uniform_tests`, called by `dmsh_build.py` after the other rewrites (and on the geometry stage and its reader), reads a word that the shader tests for equality with an integer constant more than once through `dm_uword`, a volatile load of that word for each test (`as_type(dm_uword((*spvDescriptorSet0.fp_c5)._m0[16], 3)) != 102`). The front end can neither merge one test's load with the next nor move it out of its block, so it folds no chain. A uniform load per test is less than the bounds test each read had before. A word tested once, a test against a variable, and an ordering test are left as they are. +- **Found on:** Nintendo Switch Sports Resort (`04008E001F098000`, no profile yet): with chain `b7249b1c2b66` the Mac app's archive ran 15 hours in its first group (vertex `1f0649e274faa626`, `1552eae0ea641bcd` and `24ce8042ec8edfff` with pixel `7c2becb054a53765`, vertex `2265ca4db20439bf` with pixel `692008b305ce264f`, each left out after about 30 minutes alone), where chain `592225da745b` had built 1,000 bases in about 3 minutes. Measured alone with `metal-tt` 32023.922 (`applegpu_g14s`, macOS 27.2, M2 Max with 32 GB), `2265ca4db20439bf` + `692008b305ce264f`: chain `592225da745b` 7.4 s and 0.25 GB; chain `b7249b1c2b66` stopped at 10 minutes and 13.4 GB; that chain's vertex library with the old pixel library 4.2 s, the old vertex library with the new pixel library stopped at 5 minutes and 10 GB; the old MSL with the fixed-offset bounds tests taken out stopped at 5 minutes; the same with the tests kept only on the 910 reads that feed an equality test, 3.1 s; with `uniform_tests`, 6.3 s and 0.21 GB. `1f0649e274faa626` + `7c2becb054a53765` with `uniform_tests`: 11.5 s and 0.45 GB. +- **Check:** `python3 dm/gfx/tests/test_dmsh_uniform_tests.py` (the forms it rewrites and leaves; with Xcode, the plain shader's IR has a switch with `dmsh_build.py`'s and `ios_metal.py`'s commands and the rewritten one none). + ### Bracket nesting past the front end's 256 - **Symptom:** the shader build prints `fatal error: bracket nesting level exceeded maximum of 256` and `.metal: does not compile`, and the shader is left out.