diff --git a/dm/gfx/nvn/dmsh_chain.h b/dm/gfx/nvn/dmsh_chain.h index 2e1844790f..acf9caa122 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 "7c609af36b0f" +#define DMSH_CHAIN "3695a7c258bb" diff --git a/dm/gfx/tests/test_dmsh_flag_reads.py b/dm/gfx/tests/test_dmsh_flag_reads.py new file mode 100644 index 0000000000..b7e9b805b6 --- /dev/null +++ b/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 \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((*spvDescriptorSet0.fp_c5)._m0[40].x); + uint _2 = as_type((*spvDescriptorSet0.fp_c12)._m0[3].w); + float _3 = (*spvDescriptorSet0.fp_c5)._m0[41].y; + int _4 = as_type((*spvDescriptorSet0.fp_c0)._m0[2].z); + int _5 = as_type((*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(dm_vl((*spvDescriptorSet0.fp_c5)._m0[40]).x)', out) + self.assertIn('as_type(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((*spvDescriptorSet0.fp_c0)._m0[2].z)', out) # a driver bank + self.assertIn('as_type((*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() diff --git a/dm/gfx/tools/dmsh_build.py b/dm/gfx/tools/dmsh_build.py index 169d7892b4..ae2bd49a84 100644 --- a/dm/gfx/tools/dmsh_build.py +++ b/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 diff --git a/dm/gfx/tools/dmsh_rewrite.py b/dm/gfx/tools/dmsh_rewrite.py index 44abecdd52..5cae2448a4 100644 --- a/dm/gfx/tools/dmsh_rewrite.py +++ b/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((*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 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: diff --git a/docs/agents/symptoms.md b/docs/agents/symptoms.md index beef286e3d..6c0892146d 100644 --- a/docs/agents/symptoms.md +++ b/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) diff --git a/docs/reference/shaders.md b/docs/reference/shaders.md index ef47841e8b..a50d9fad0e 100644 --- a/docs/reference/shaders.md +++ b/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(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.