fix: compile deeply nested NS2 Metal shaders offline

THIRD_PARTY_NOTICES.md

@@ -18,6 +18,7 @@ Code in this repository that comes from other projects keeps its license; the li
 | Switch 2 (Ampere / NVN2) shaders: the decoder, emitters and control section in `games/totk/tools/totk_xlate/patches/0006-ampere-switch-2-shaders.patch`; geometry shaders as compute kernels in `patches/0007-fixed-input-geometry-as-compute.patch`, `games/totk/tools/totk_xlate/Geometry.cs` and `GeometryLean.cs` | [nxnative](https://github.com/seedds/nxnative), the commits each file names, adapted | MIT | `games/totk/third_party/licenses/nxnative.txt` |
 | geometry primitive assembly `games/totk/third_party/d2r_gfx/gfx/nvn/nvn_geometry.h` and the host kernels in `nvn_geometry_lean.inc` (from its `geo_pack.metal`) | nxnative | MIT | `games/totk/third_party/licenses/nxnative.txt` |
 | Switch 2 SDK 20 host services in `games/totk/runtime/ns2_host.c`, `ns2_audio_out.c` and `ns2_codec.c` (behaviour ported) | nxnative, the commits each file names | MIT | `games/totk/third_party/licenses/nxnative.txt` |
+| Source-specific Metal bracket depth in `build_app.py` and `games/totk/scripts/build_shader_libs.py`; synthetic nested-walk test in `tests/test_metal_bracket_depth.py` | [nxnative commit `56b73f21`](https://github.com/stan69b/nxnative/commit/56b73f21) (seedds, `oot-clean`), adapted | MIT | `games/totk/third_party/licenses/nxnative.txt` |
 | audio effects in `games/totk/runtime/audio_reverb.c`, `games/totk/runtime/audio_effects.c` (delay, reverb, compressor) and the mixer | ported from Ryujinx.Audio | MIT | `games/totk/third_party/licenses/Ryujinx-Audio.txt` |
 | `zeldaVR/streamer/basaltvr.patch` (our changes to the streamer, with lines of its source as diff context) | [BasaltVR](https://github.com/frenchbaguetteman/BasaltVR) at `4c46172` | MIT | `zeldaVR/streamer/LICENSE.BasaltVR` |
 | `windows_games/ra2/third_party/x86recomp` (x86 to arm64 recompiler and Win32 runtime) | [x86recomp](https://github.com/JTM-rootstorm/x86recomp) at `1c90fe3`, with WinNativeGames' changes (`UPSTREAM.md` lists them) | MIT | `windows_games/ra2/third_party/x86recomp/LICENSE` |

build_app.py

@@ -47,7 +47,7 @@
 Optional: Xcode's Metal toolchain (xcodebuild -downloadComponent MetalToolchain) compiles the startup shaders into the
 app, so moving or rebuilding it does not repeat that work at launch.
 """
-import argparse, concurrent.futures, glob, hashlib, json, os, platform, plistlib, shutil, struct, subprocess, sys, threading, time
+import argparse, concurrent.futures, glob, hashlib, itertools, json, os, platform, plistlib, shutil, struct, subprocess, sys, threading, time
 from pathlib import Path
 
 REPO = Path(__file__).resolve().parent
@@ -1207,6 +1207,25 @@ def metal_compiler():
 
 # The options of totk_precise_opts() in nvn_mtl.m: the translated shaders need exactly rounded float arithmetic.
 METAL_FLAGS = ['-fmetal-math-mode=safe']
+# Adapted from stan69b/nxnative oot-clean 56b73f21 (seedds, MIT); see THIRD_PARTY_NOTICES.md.
+METAL_BRACKET_LIMIT = 256
+_BRACKETS = bytes.maketrans(b'([{)]}', b'((()))')
+_NOT_BRACKETS = bytes(c for c in range(256) if c not in b'([{)]}')
+
+
+def bracket_depth(msl):
+    """Count the combined nesting of all three bracket kinds in generated MSL."""
+    # shortcut: comments/literals count too; use a lexer if this accepts arbitrary handwritten MSL.
+    marks = Path(msl).read_bytes().translate(_BRACKETS, _NOT_BRACKETS)
+    return max(itertools.accumulate(1 if c == 0x28 else -1 for c in marks), default=0)
+
+
+def metal_flags(msl):
+    """Raise the parser limit only for deeply nested generated sources (unrolled early-exit loops)."""
+    # A higher parser limit changes no code generation: keep it out of METAL_FLAGS and the shared cache stamp.
+    if bracket_depth(msl) > METAL_BRACKET_LIMIT:
+        return [*METAL_FLAGS, '-fbracket-depth=1024']
+    return METAL_FLAGS
 
 
 def variant_folder(game):
@@ -1258,7 +1277,8 @@ def step_shader_libs(tree, logs, *, version='v0'):
     def compile_one(item):
         digest, src = item
         partial = libs/f'{digest}.partial'
-        r = subprocess.run([metal, *METAL_FLAGS, '-o', partial, src], capture_output=True, text=True, env=metal_env)
+        flags = metal_flags(src) if ns2_game(version) else METAL_FLAGS
+        r = subprocess.run([metal, *flags, '-o', partial, src], capture_output=True, text=True, env=metal_env)
         if r.returncode:
             partial.unlink(missing_ok=True)
             return f'{src.name}\n{r.stdout}{r.stderr}'

games/totk/scripts/build_shader_libs.py

@@ -20,7 +20,7 @@
 
 REPO = Path(__file__).resolve().parents[3]
 sys.path.insert(0, str(REPO))
-from build_app import METAL_FLAGS, metal_compiler   # noqa: E402
+from build_app import METAL_FLAGS, metal_compiler, metal_flags   # noqa: E402
 
 
 def main():
@@ -52,7 +52,7 @@ def main():
     def compile_one(item):
         digest, src = item
         partial = a.out/f'{digest}.partial'
-        r = subprocess.run([metal, *METAL_FLAGS, '-o', partial, src], capture_output=True, text=True, env=env)
+        r = subprocess.run([metal, *metal_flags(src), '-o', partial, src], capture_output=True, text=True, env=env)
         if r.returncode:
             partial.unlink(missing_ok=True)
             return f'{src.name}\n{r.stdout}{r.stderr}'

tests/test_metal_bracket_depth.py

@@ -0,0 +1,129 @@
+"""Deep generated NS2 MSL compiles offline without invalidating ordinary shader libraries.
+
+Synthetic nested walk adapted from stan69b/nxnative, oot-clean 56b73f21 (seedds, MIT).
+"""
+import hashlib
+import importlib.util
+from pathlib import Path
+import subprocess
+import sys
+import tempfile
+import unittest
+from unittest.mock import patch
+import build_app
+
+ROOT = Path(__file__).resolve().parents[1]
+
+
+def nested_kernel(depth):
+    body = ''.join(f'uint t{i} = min(left, a[{i}]); a[{i}] -= t{i}; left -= t{i};\nif (left != 0) {{\n'
+                   for i in range(depth))
+    return ('#include <metal_stdlib>\nusing namespace metal;\n'
+            'kernel void k(device uint* a [[buffer(0)]], uint id [[thread_position_in_grid]]) {\n'
+            'uint left = id;\n' + body + '}\n' * (depth + 1))
+
+
+class MetalBracketDepth(unittest.TestCase):
+    def setUp(self):
+        tmp = tempfile.TemporaryDirectory(prefix='mng-metal-depth-')
+        self.addCleanup(tmp.cleanup)
+        self.root = Path(tmp.name)
+        self.calls = {}
+
+    def compiler(self, args, **kwargs):
+        source = Path(args[-1])
+        self.calls[source.name] = args
+        Path(args[args.index('-o') + 1]).write_bytes(b'synthetic library')
+        return subprocess.CompletedProcess(args, 0, stdout='', stderr='')
+
+    def sources(self, folder):
+        folder.mkdir(parents=True)
+        for name, depth in [('deep_b.metal', 300), ('flat_b.metal', 3)]:
+            (folder/name).write_text(nested_kernel(depth))
+
+    def check_flags(self):
+        self.assertIn('-fbracket-depth=1024', self.calls['deep_b.metal'])
+        self.assertNotIn('-fbracket-depth=1024', self.calls['flat_b.metal'])
+
+    def test_depth_boundary_and_bracket_kinds(self):
+        source = self.root/'source.metal'
+        for text, expected in [('', 0), ('f(a[b{c}]) { (x); }', 3), ('{' * 256 + '}' * 256, 256),
+                               ('{' * 257 + '}' * 257, 257)]:
+            with self.subTest(depth=expected):
+                source.write_text(text)
+                self.assertEqual(build_app.bracket_depth(source), expected)
+                self.assertEqual(build_app.metal_flags(source), ['-fmetal-math-mode=safe'] +
+                                 (['-fbracket-depth=1024'] if expected > 256 else []))
+
+    def test_app_compiler_selects_only_deep_ns2_sources_and_reuses_cache(self):
+        for version in ('pkp', 'dkb', 'v143', 'botw190'):
+            with self.subTest(version=version):
+                tree = self.root/version
+                seed = tree/f'build/shader-seed/{version}'
+                self.sources(seed)
+                libs = tree/'build/shader-libs'
+                libs.mkdir()
+                (libs/'compiler.txt').write_text('fake 1 -fmetal-math-mode=safe')
+                (libs/'existing.metallib').write_bytes(b'keep')
+                self.calls.clear()
+                with patch.object(build_app, 'metal_compiler', return_value=('metal', 'fake 1', None)), \
+                        patch.object(build_app, 'variant_sources', return_value=[]), patch.object(build_app, 'say'), \
+                        patch.object(build_app.subprocess, 'run', side_effect=self.compiler):
+                    build_app.step_shader_libs(tree, tree/'logs', version=version)
+                    if version in ('pkp', 'dkb'):
+                        self.check_flags()
+                    else:
+                        for args in self.calls.values():
+                            self.assertNotIn('-fbracket-depth=1024', args)
+                    self.assertEqual(len(self.calls), 2)
+                    self.calls.clear()
+                    build_app.step_shader_libs(tree, tree/'logs', version=version)
+                    self.assertEqual(self.calls, {})
+                self.assertEqual((libs/'compiler.txt').read_text(), 'fake 1 -fmetal-math-mode=safe')
+                self.assertEqual((libs/'existing.metallib').read_bytes(), b'keep')
+                for source in seed.glob('*.metal'):
+                    digest = hashlib.sha256(source.read_bytes()).hexdigest()
+                    self.assertTrue((libs/f'{digest}.metallib').is_file())
+                self.assertEqual((tree/'logs/shader-libs.log').read_text(), '')
+
+    def test_work_folder_compiler_selects_deep_sources_and_reuses_cache(self):
+        spec = importlib.util.spec_from_file_location('build_shader_libs', ROOT/'games/totk/scripts/build_shader_libs.py')
+        tool = importlib.util.module_from_spec(spec)
+        spec.loader.exec_module(tool)
+        cache, out = self.root/'cache', self.root/'out'
+        self.sources(cache)
+        with patch.object(tool, 'metal_compiler', return_value=('metal', 'fake 1', None)), \
+                patch.object(tool.subprocess, 'run', side_effect=self.compiler), patch('builtins.print'), \
+                patch.object(sys, 'argv', ['build_shader_libs.py', '--cache', str(cache), '--out', str(out)]):
+            for run in range(2):
+                with self.assertRaises(SystemExit) as done:
+                    tool.main()
+                self.assertEqual(done.exception.code, 0)
+                if run == 0:
+                    self.check_flags()
+                else:
+                    self.assertEqual(self.calls, {})
+                self.calls.clear()
+        self.assertEqual((out/'compiler.txt').read_text(), 'fake 1 -fmetal-math-mode=safe')
+        self.assertEqual((out/'failed.log').read_text(), '')
+
+    def test_real_metal_compiler_accepts_deep_walk_only_with_raised_limit(self):
+        found = build_app.metal_compiler()
+        if not found:
+            self.skipTest('requires the macOS Metal toolchain')
+        metal, _, env = found
+        source = self.root/'deep.metal'
+        source.write_text(nested_kernel(300))
+        output = self.root/'deep.metallib'
+        base = subprocess.run([metal, *build_app.METAL_FLAGS, '-o', output, source], env=env,
+                              capture_output=True, text=True, timeout=60)
+        self.assertNotEqual(base.returncode, 0)
+        self.assertIn('bracket nesting level exceeded', base.stderr)
+        fixed = subprocess.run([metal, *build_app.metal_flags(source), '-o', output, source], env=env,
+                               capture_output=True, text=True, timeout=60)
+        self.assertEqual(fixed.returncode, 0, fixed.stdout + fixed.stderr)
+        self.assertGreater(output.stat().st_size, 0)
+
+
+if __name__ == '__main__':
+    unittest.main()