LDSM is fenced by subgroup barriers: Ocarina of Time's movie filter read other lanes' shared-memory stores before they landed by MetaIX · Pull Request #1376 · seedds/nxnative

MetaIX · 2026-11-05T20:32:12Z

What and why

Symptom: Ocarina of Time's title movie (0400486024DEC000 v131072, profile in #1299) shows dotted 1-pixel streaks: every other pixel of a few hundred rows per frame, mostly green and magenta, in different places every run. The movie's YUV→RGB texture (6230) is clean. The streaks are in the output of the movie library's NDVE network (10 compute shaders, texture 6231), before agl's upscaler.

Cause: these shaders store tiles to shared memory (STS) and read them back with LDSM. Five of them (ba079f68ae288dba, a8b4fb48fee0ccf8, e6be82d835fbe4d3, 7f7ff0f36b54d90f, 5a7d69b067a7f011) have no BAR between the two. That is safe on the Switch, where a warp's shared-memory accesses are in order. Metal orders another lane's threadgroup store and this lane's load only through a barrier, and its compiler moved them past each other. The wrong pixels are mostly at odd x, the high halves of packed pairs: about 2.3 times as many as at even x.

Change: InstEmitAmpereTensor.ALdsm emits a subgroup barrier (Instruction.SubgroupBarrier, simdgroup_barrier(mem_flags::mem_threadgroup)) before LDSM's loads and another after them. The second keeps any lane's next store over the rows from landing before all lanes have read them. LDSM is warp-collective (the whole warp runs it), so the whole SIMD-group reaches both barriers.

Relation to #1275: its tg_fence (a SIMD-group-scope memory fence after each threadgroup store, no execution barrier) is complementary but not enough here. With fences only, the streak count fell to about 3,800; with the barriers, to the noise floor.

Reference entry: LDSM reads rows other lanes stored.

Checked

Tests (on this branch, translator rebuilt from its sources: dotnet publish, the 5 dlls copied, tool_lock.py --update-lock, dmsh_chain.py --header):
- test_dmsh_*.py all pass (2 skipped), with the new Shared.test_ldsm_reads_between_subgroup_barriers;
- test_docs.py, test_nsp_chain.py and test_xlate_anycpu.py pass;
- lint_python.py not run here (no ruff on this Mac).
Games run: Ocarina of Time 0400486024DEC000 v131072, title movie and title screen, on our stack (main 95f9a6a + [WIP] The Legend of Zelda: Ocarina of Time (Switch 2, 0400486024DEC000, v0, SDK 22.2.8) runs on the Mac through the title movie into Kokiri Forest #1299 + Ampere F2IP and a task shader's own payload read: Ocarina of Time's untranslated shaders translate, all but one that is data #1300 + Zelda: Ocarina of Time — Restore missing mesh and texture effects on Mac #1312 with this change; the same diff).
- Streak count in 6231 at frame 700 (pixels that differ from both vertical neighbours while those agree): 6,700–14,000 over 7 runs before. Serialising passes (DM_NVN_PASSGPU_SERIAL), DM_NVN_DISPATCH_OWN, DM_NVN_BARRIER=game and DM_NVN_POOL_CHUNKS=0 did not remove them. After: 71–86 over 5 runs; the clean YUV source scores 40–58.
- GPU time: no change (long-frame median 35.3 → 35.1 ms).
Not run: regress.py check --games. The change reaches every Ampere shader with LDSM (on main, agl's NDSR upscalers, e.g. Donkey Kong Bananza's), and I have no such game built on this branch. Each LDSM gains two SIMD-group barriers and nothing else changes.
Devices and OS: MacBook Pro, Apple M1 Max, 64 GB, macOS 26.6.2, Xcode 27.2.

Docs

The finding is in its docs/reference/shaders.md entry, and python3 dm/tools/docs_symptoms.py was run (the symptom index)
A chain tool changed: dmsh_chain.py --header was run and dm/gfx/nvn/dmsh_chain.h is part of the change
No game files, keys or data extracted from a game in the change or the attachments

🤖 Generated with Claude Code

…ead other lanes' shared-memory stores before they landed; chain 554f4865aeef Its NDVE shaders store tiles with STS and read them back with LDSM, five of them with no BAR between. A warp's shared-memory accesses are in order on the Switch; Metal orders another lane's threadgroup store and this lane's load only through a barrier. ALdsm now puts a SIMD-group barrier before its loads and one after (LDSM is warp-collective, so the whole SIMD-group reaches both). Every other pixel of a few hundred rows of the title movie was wrong (6,700-14,000 a frame); after, 71-86, the clean source's level. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…be rebuilt)

…dds#1376 (float add atomics, pool table count, LDSM); chain lock and header regenerated (65ed9f9547f9) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_016qteK29T18U8p1c9LsVY8g

…ers: Ocarina of Time's movie filter read other lanes' shared-memory stores before they landed

…, seedds#1410, seedds#1411, seedds#1412, seedds#1413, seedds#1414, seedds#1329 on top of main's translator fixes): chain c9cf9613657e Publishes Ryujinx.Graphics.Shader, Ryujinx.Common, Spv.Generator and dmsh_xlate (osx-arm64, framework-dependent, the universal launcher kept) from the merged source; the PRs' own binaries were built without this tree's translator changes (float atomic adds, F2IP, the growing raw pool table). The lock and dmsh_chain.h are written again. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

) replaces the pull request heads Upstream landed Ocarina of Time in main (chain adec7d81d970: "holds 60 fps in Kokiri Forest, 4 GB less") and closed the pull requests it folded in, with reworded and partly different code. Mixing a pull request head with its landed version hunk by hunk silently broke things (two GlobalToStorage tables, a duplicate AF2ip, two mtk_pso definitions), so the rule here is one version per feature: - Pull requests upstream closed (seedds#1299, seedds#1300, seedds#1312, seedds#1343, seedds#1345, seedds#1353, seedds#1354, seedds#1359, seedds#1360, seedds#1373, seedds#1428, seedds#1430, seedds#1434): their hunks are taken out of our side before the three-way merge, so upstream's landed version stands alone. Gone with them: bindless pool table patching (reg_bump), the ghost table's per-submission hold, the growing raw pool table and its write-log sizing, seedds#1428's pool_buf spares (upstream's pool_table_take is the same idea), the ixc-based indirect mesh draw (upstream uses mdx), and their tests (t_pool_table, t_mesh_indirect, t_gmem_raw, finalized_bindless_pool, test_bnsh_nvn2, test_ndve). - Open pull requests that upstream's landing supersedes: seedds#1394 (upstream's mtl_pool_addressed) and seedds#1393 pool_skip_retired (upstream's budgeted pool walk; the profile does not use it) are dropped with their tests and docs. seedds#1368's mesh pipeline prewarm is dropped too: upstream kept the background build in nvn_geometry.inc, and the prewarm's mtk_build / mtk_claim / g_warm_mps went with it. - Kept, as upstream does not have them: NXForge's own features (GPU watchdog, present filters and anti-aliasing, aniso floor, quick settings, pipeline archive compaction, upload breakdown, memory report lines) and the still open pull requests seedds#1304, seedds#1305, seedds#1308, seedds#1320, seedds#1331, seedds#1332, seedds#1355, seedds#1356, seedds#1369, seedds#1376, seedds#1377, seedds#1379, seedds#1382, seedds#1390, seedds#1391, seedds#1392, seedds#1414, seedds#1420, seedds#1436, seedds#1439, seedds#1446, seedds#1453. - Combined where both sides changed the same lines: GW_N 1 << 16 with seedds#1320's wide-write intervals (GW_PAGES = SBW_WPAGES, 128 slots), the CPU upload log with upstream's linear upload route, arch_open_end before the compaction limit, cbn_prog in geo_gs_one (seedds#1382). - Translator: upstream's sources and binaries, plus seedds#1414's InstEmitAmpereMem diff and the LDSM SubgroupBarrier in InstEmitAmpereTensor.cs. The binaries and dmsh_chain.lock are upstream's until the Mac rebuilds them (dotnet publish, tool_lock --update-lock, dmsh_chain --header); dmsh_chain.h is regenerated for now. - RED_ONCE keys: upstream took 153 and 154, so the present anti-aliasing lines move to 156 and 157 and seedds#1420's packaged-entry line to 158 (array 159). - Host: hle_net.c and t_net_eventfd.c are upstream's (its eventfd poll waits out the whole timeout, which our 5d7920c fixed another way). - Kept ours: dm/tools/romfs_image.py and its test, so the 26 GB RomFS image is not packed again. dm/tools tests on Linux: same failures as on 875449e (Linux-only quirks), plus test_nlg4 and test_dmsh_pool_arrays, which fail on upstream main too. Not yet built on the Mac. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_016qteK29T18U8p1c9LsVY8g

…er and shader chain, the branch's host services, movie clock, packing and Ocarina of Time additions Upstream reached about 60 fps in Kokiri Forest with its own versions of most of the branch's merged PRs, so where both changed the same code upstream's is kept: - The shader chain is upstream's (translator, lock, dmsh tools, Ryujinx sources), plus the branch's bracket-depth flattening (dmsh_depth.py, test_dmsh_depth.py); chain c732699d4fdb. Not carried: the LDSM subgroup fence (seedds#1376, needs a translator rebuild; upstream's HMMA/LDSM rework reports clean movies) and the shared raw-memory lookups (seedds#1405, superseded by upstream's last-entry lookups). - The renderer (dm/gfx/nvn, dm/gfx/msl) and its tests are upstream's. Not carried: seedds#1390-seedds#1394, seedds#1404, seedds#1408, seedds#1409, seedds#1417, seedds#1418, seedds#1420, seedds#1427, seedds#1320, seedds#1331, seedds#1332, seedds#1368, seedds#1285 renderer parts, pool_skip_retired, and the branch's warning fixes and prewarm progress there. They stay on oot-integration. - Ocarina of Time's profile is upstream's, with the branch's 12 movie compute rows in pso_manifest.tsv, PLAN.md and PR-AUDIT.md. - shader_split.py is upstream's plus the BAHS record check (test_split_bahs_records.py). - Host: upstream's EventFd, PlayReport send, audio lock and vibration answers; the branch's service sessions (hle_sf, hle_prepo), sync waits, AV1 movie clock and parallel decode, diag_abort frame fix and focus control are kept. - Scripts: shader.noindex, tmp.noindex, notes outside the shader key, bench and regress changes. Checked: the Mac host builds (GAME=0400486024DEC000); host unit tests pass; the Python tool tests pass on Python 3.13 except test_dmsh_pool_arrays, test_nro_pack, test_nsp_dump and test_nsp_incremental, which fail on upstream main as well. Not run: the full shader build on this chain, gameplay.

…no longer whites out through the game's own upscaler, whose LDSMs read rows other lanes had not finished storing (chain fbd67412d49d, PR #1376, issue #1537) A warp stores tiles to shared memory and reads them back with LDSM, often with no BAR between; the Switch runs a warp's shared memory accesses in order, Metal orders another lane's threadgroup store and this lane's load only through a barrier, and its compiler moves them past each other. PR #1376 found it in the movie networks (dotted streaks, every other pixel) and put the barriers in ALdsm; since bd9f422 a compute shader's LDSM is a MatrixLoad lowered late by TensorPass, so the barriers now bracket the MatrixLoads, and so the per-lane code that replaces them. agl's NDSR upscaler takes all its LDSMs that way: on chain 844bf1a5ea8f the forest was covered in white blotches on an M4 Pro (with DM_UPSCALER=metalfx in its place, none). TensorPass does not count a subgroup barrier as a shared memory write, so HMMAs still load A where their LDSM read it. Checked: test_dmsh_ampere_mem.py Shared.test_ldsm_reads_between_subgroup_barriers (fails on the translator before), the 21 translator test files with DMSH_TEST_XLATE (test_dmsh_pairs' two NameErrors are the same with the old translator), test_nsp_chain, test_xlate_anycpu, test_docs, lint_python 0 new. Ocarina of Time on an M4 Pro (seen-shader library, Kokiri Forest autosave, scripted load and walk, headless): white blotches on 844bf1a5ea8f, none on fbd67412d49d at the bridge or on the village floor; GPU time through the title movie unchanged (12.7 / 13.6 ms medians a 300-frame window on both). Regression: Ocarina only (the game this reaches most; Donkey Kong Bananza, the other NDSR game here, has no library to rebuild from and no seen list). Not fixed: the dotted horizontal rows issue #1537 reports remain with either upscaler; chain adec7d81d970 draws the same view without them, so their cause lies between 6121fd7 and 04df08d (the game note's open line). Co-authored-by: MetaIX <125941078+MetaIX@users.noreply.github.com>

seedds · 2026-11-07T00:57:45Z

Landed reworked as 1677691, with you as co-author. Since bd9f422 a compute shader's LDSM is a MatrixLoad that TensorPass lowers late, so the barriers now bracket the MatrixLoads in ALdsm (both paths). On chain 844bf1a5 the same race made agl's NDSR cover Ocarina of Time's Kokiri Forest in white blotches on an M4 Pro; the barriers remove them, with no change in GPU time. Thank you for the diagnosis.

…no longer whites out through the game's own upscaler, whose LDSMs read rows other lanes had not finished storing (chain fbd67412d49d, PR seedds#1376, issue seedds#1537) A warp stores tiles to shared memory and reads them back with LDSM, often with no BAR between; the Switch runs a warp's shared memory accesses in order, Metal orders another lane's threadgroup store and this lane's load only through a barrier, and its compiler moves them past each other. PR seedds#1376 found it in the movie networks (dotted streaks, every other pixel) and put the barriers in ALdsm; since bd9f422 a compute shader's LDSM is a MatrixLoad lowered late by TensorPass, so the barriers now bracket the MatrixLoads, and so the per-lane code that replaces them. agl's NDSR upscaler takes all its LDSMs that way: on chain 844bf1a5ea8f the forest was covered in white blotches on an M4 Pro (with DM_UPSCALER=metalfx in its place, none). TensorPass does not count a subgroup barrier as a shared memory write, so HMMAs still load A where their LDSM read it. Checked: test_dmsh_ampere_mem.py Shared.test_ldsm_reads_between_subgroup_barriers (fails on the translator before), the 21 translator test files with DMSH_TEST_XLATE (test_dmsh_pairs' two NameErrors are the same with the old translator), test_nsp_chain, test_xlate_anycpu, test_docs, lint_python 0 new. Ocarina of Time on an M4 Pro (seen-shader library, Kokiri Forest autosave, scripted load and walk, headless): white blotches on 844bf1a5ea8f, none on fbd67412d49d at the bridge or on the village floor; GPU time through the title movie unchanged (12.7 / 13.6 ms medians a 300-frame window on both). Regression: Ocarina only (the game this reaches most; Donkey Kong Bananza, the other NDSR game here, has no library to rebuild from and no seen list). Not fixed: the dotted horizontal rows issue seedds#1537 reports remain with either upscaler; chain adec7d81d970 draws the same view without them, so their cause lies between 6121fd7 and 04df08d (the game note's open line). Co-authored-by: MetaIX <125941078+MetaIX@users.noreply.github.com> (cherry picked from commit 1677691)

…no longer whites out through the game's own upscaler, whose LDSMs read rows other lanes had not finished storing (chain fbd67412d49d, PR seedds#1376, issue seedds#1537)

…no longer whites out through the game's own upscaler, whose LDSMs read rows other lanes had not finished storing (chain fbd67412d49d, PR seedds#1376, issue seedds#1537) A warp stores tiles to shared memory and reads them back with LDSM, often with no BAR between; the Switch runs a warp's shared memory accesses in order, Metal orders another lane's threadgroup store and this lane's load only through a barrier, and its compiler moves them past each other. PR seedds#1376 found it in the movie networks (dotted streaks, every other pixel) and put the barriers in ALdsm; since 95a08cc a compute shader's LDSM is a MatrixLoad lowered late by TensorPass, so the barriers now bracket the MatrixLoads, and so the per-lane code that replaces them. agl's NDSR upscaler takes all its LDSMs that way: on chain 844bf1a5ea8f the forest was covered in white blotches on an M4 Pro (with DM_UPSCALER=metalfx in its place, none). TensorPass does not count a subgroup barrier as a shared memory write, so HMMAs still load A where their LDSM read it. Checked: test_dmsh_ampere_mem.py Shared.test_ldsm_reads_between_subgroup_barriers (fails on the translator before), the 21 translator test files with DMSH_TEST_XLATE (test_dmsh_pairs' two NameErrors are the same with the old translator), test_nsp_chain, test_xlate_anycpu, test_docs, lint_python 0 new. Ocarina of Time on an M4 Pro (seen-shader library, Kokiri Forest autosave, scripted load and walk, headless): white blotches on 844bf1a5ea8f, none on fbd67412d49d at the bridge or on the village floor; GPU time through the title movie unchanged (12.7 / 13.6 ms medians a 300-frame window on both). Regression: Ocarina only (the game this reaches most; Donkey Kong Bananza, the other NDSR game here, has no library to rebuild from and no seen list). Not fixed: the dotted horizontal rows issue seedds#1537 reports remain with either upscaler; chain adec7d81d970 draws the same view without them, so their cause lies between 29df871 and 5b9a707 (the game note's open line). Co-authored-by: MetaIX <125941078+MetaIX@users.noreply.github.com>

github-actions Bot added nvn NVN on Metal and its tests (dm/gfx/nvn, dm/gfx/tests) docs Documentation shader chain The shader translation chain (DMSH_CHAIN): a change rebuilds every game's shaders labels Nov 5, 2026

gnosticgeek pushed a commit to gnosticgeek/nxforge that referenced this pull request Nov 5, 2026

Merge upstream PR seedds#1376 (translator binaries and chain lock to …

9547176

…be rebuilt)

This was referenced Nov 5, 2026

Test support for The Legend of Zelda: Ocarina of Time #1265

Open

Fix Ocarina of Time foliage jitter by preserving live shader inputs #1343

Closed

LeonardoBai12 mentioned this pull request Nov 6, 2026

Ocarina of Time on an M1 Max: the movie library's NDVE network takes 35 ms of GPU a frame, so movies and the title screen run at ~28 fps #1375

Open

fpedrolucas95 pushed a commit to fpedrolucas95/nxnative that referenced this pull request Nov 6, 2026

Merge seedds#1376 into upstream-prs: LDSM is fenced by subgroup barri…

65ba803

…ers: Ocarina of Time's movie filter read other lanes' shared-memory stores before they landed

ksolenk38ze mentioned this pull request Nov 6, 2026

The Legend of Zelda: Ocarina of Time (Switch 2, 0400486024DEC000, v0 & v131072): playable at 60 fps through Dodongo's Cavern #1444

Closed

3 tasks done

NicoPV94 mentioned this pull request Nov 6, 2026

Revert the HMMA chain change (bd9f422d): it puts dotted bands in Ocarina of Time's upscaled picture #1576

Closed

seedds closed this Nov 7, 2026

Navigation Menu

Repository navigation

LDSM is fenced by subgroup barriers: Ocarina of Time's movie filter read other lanes' shared-memory stores before they landed - #1376

LDSM is fenced by subgroup barriers: Ocarina of Time's movie filter read other lanes' shared-memory stores before they landed#1376
MetaIX wants to merge 1 commit into
seedds:mainfrom
MetaIX:pr/ldsm-subgroup-barriers

MetaIX commented Nov 5, 2026

Uh oh!

seedds commented Nov 7, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

MetaIX commented Nov 5, 2026

What and why

Checked

Docs

Uh oh!

seedds commented Nov 7, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants