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

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 "223f10d34f97"
+#define DMSH_CHAIN "554f4865aeef"

dm/gfx/tests/test_dmsh_ampere_mem.py

@@ -378,6 +378,16 @@ def test_ldsm_mode_is_not_a_scale(self):
         msl = self.msl(mov(3, 0x40), inst(*self.LDSM_MT88), sts(0x20, 20))
         self.assertIn('spvSubgroupShuffle(8.96831', msl)   # the address 0x40 as a float's bits; scaled, it would be 0x100 (3.5873e-43)
 
+    # LDSM reads rows other lanes stored, often with no BAR before it (Ocarina of Time's movie filter): a SIMD-group barrier before its reads, so the stores have landed, and one after,
+    # so no lane's next store reaches the rows before every lane has read them
+    def test_ldsm_reads_between_subgroup_barriers(self):
+        msl = self.msl(mov(3, 0x40), inst(*self.LDSM_MT88), sts(0x20, 20))
+        body = msl[msl.index('kernel void main0'):]
+        barriers = [m.start() for m in re.finditer(r'simdgroup_barrier\(mem_flags::mem_threadgroup\)', body)]
+        reads = [m.start() for m in re.finditer(r'= as_type<float>\(_\d+\[int\(', body)]
+        self.assertEqual(len(barriers), 2, body)
+        self.assertTrue(reads and barriers[0] < reads[0] and reads[-1] < barriers[1], body)
+
 
 class Pokopia(Translated):
     """Pokemon Pokopia's encodings, on the register the atomic reads: F2FP.MERGE_C (22,886 shaders), F2I from a half (1,067), MEMBAR with ERRBAR (144)"""

dm/gfx/tools/dmsh_chain.lock

@@ -1,7 +1,7 @@
 # Translation chain tool fingerprints (written by tool_lock.py --update-lock; regenerate after changing tools and commit it)
-dmsh_xlate.dll	5e66d6e283045e07
-Ryujinx.Graphics.Shader.dll	8289f2a6757ed3bd
-Ryujinx.Common.dll	387eb6599652dcb0
+dmsh_xlate.dll	bfbefffd2ce37ab8
+Ryujinx.Graphics.Shader.dll	370c6a213727558f
+Ryujinx.Common.dll	5a5996b79c955d67
 xlate_source	71adcdd020e03a64
 spirv-cross	e37945d8e441d34d
 metal_macosx	Apple metal version 32023.922 (metalfe-32023.922.4)

dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Instructions/InstEmitAmpereTensor.cs

@@ -261,6 +261,12 @@ public static void ALdsm(EmitterContext context)
             int id = context.ResourceManager.SharedMemoryId;
             Operand[] results = new Operand[count];
 
+            // nxnative: LDSM reads rows that other lanes of the warp stored (STS) just before, with no BAR between (the agl networks of Ocarina of Time's movie
+            // filter); the warp's shared memory accesses are in order on the Switch, while Metal orders another lane's threadgroup store and this load only through
+            // a barrier, and the compiler moved them past each other: every other pixel of a few rows wrong (the high halves). LDSM is warp-collective (the whole warp
+            // runs it), so a subgroup barrier before its loads and one after (no lane's next store over the rows before all have read them) is safe
+            context.SubgroupBarrier();
+
             for (int m = 0; m < count; m++)
             {
                 if (mode == 0)
@@ -286,6 +292,8 @@ public static void ALdsm(EmitterContext context)
                 }
             }
 
+            context.SubgroupBarrier();
+
             for (int m = 0; m < count && x.Rd + m < RegisterConsts.RegisterZeroIndex; m++)
             {
                 context.Copy(Register(x.Rd + m, RegisterType.Gpr), results[m]);

docs/agents/symptoms.md

@@ -827,6 +827,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - `reflection conversion failed: B binding cp_t_cb196610_10280074Smplr has no resolved slot number` at load, then `the dispatched program has no compute shader`; the shader's image has no slot either (`cp_i_tcb_AA_`). → [NVN2 compute's separate handles in c[0]](../reference/shaders.md#nvn2-computes-separate-handles-in-c0)
 - `reflection conversion failed: B binding fp_t_cb2_170 has no resolved slot number` at load (the draw is skipped), the build summary counting the binding under *slot was not decoded*; or, once the read's other inputs are known, `Failed to find handle source for bindless access of type "texture2D"` and the read answers 0. `DMSH_BINDLESS_TRACE=1` on a `dmsh_xlate` job prints what a refused handle is made of. → [A handle pair ORed sampler word first](../reference/shaders.md#a-handle-pair-ored-sampler-word-first)
 - a Switch 2 game's frame rate swings between 60 and 30 fps (Donkey Kong Bananza's own controller, `main+0x204a18`, drops to 30 after two frames of 20 ms from their start to the GPU's end, and comes back after six), and the GPU profile (`DM_NVN_GPUPROF`) shows one compute dispatch of 2-6.5 ms a frame: agl's upscaler network (`agl_ndsr.pck`, shader `7629d857a0a93833`). → [Ampere HMMA through simdgroup_matrix](../reference/shaders.md#ampere-hmma-through-simdgroup_matrix)
+- a Switch 2 movie or network output has dotted streaks: every other pixel of a few hundred rows wrong, mostly green and magenta, in other places each run, while its source is clean. Ocarina of Time's title movie after its NDVE filter (texture 6231, before agl's upscaler). → [LDSM reads rows other lanes stored](../reference/shaders.md#ldsm-reads-rows-other-lanes-stored)
 - a vertex or compute shader reads zeros from NVN's driver segment `c0`. → [Driver-segment reads outside fragment shaders](../reference/shaders.md#driver-segment-reads-outside-fragment-shaders)
 - flickering hard-edged white tiles (tiled lighting with every tile row getting row 0's or other rows' lights). → [Compute `c0`: dispatch size and uniform buffer table](../reference/shaders.md#compute-c0-dispatch-size-and-uniform-buffer-table)
 - `[RED] shader ... reflection conversion failed: B binding gp_s0 has no resolved slot number`; a geometry shader reading driver segment `c0` reports that its storage-buffer table base was never measured. → [Geometry-stage storage buffer table](../reference/shaders.md#geometry-stage-storage-buffer-table)

docs/reference/shaders.md

@@ -1192,6 +1192,14 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b
 - **Note:** three of four of these HMMAs start an output tile, with C = 0. A variant for those calls that neither stores nor loads C (`simdgroup_multiply` first) was measured and not kept: 0.767 against 0.820 ns a warp-HMMA (6 %) in an isolated loop of the build's two bodies on an M4 Pro, about 0.04 ms of the upscaler's 1.6-1.9 ms a frame, and no measurable change in the game.
 - **Open:** LDSM and MOVM (27 and 72 a shader) are still emulated with shuffles (314 left in the upscaler, whose floor without them is about 0.3 ms); loading A and B straight from LDSM's source and keeping D as matrices between HMMAs (0.65 ns an HMMA) needs the translator to follow LDSM → MOVM → HMMA → HMMA chains.
 
+### LDSM reads rows other lanes stored
+
+- **Symptom:** a Switch 2 movie or network output has dotted streaks: every other pixel of a few hundred rows wrong, mostly green and magenta, in other places each run, while its source is clean. Ocarina of Time's title movie after its NDVE filter (texture 6231, before agl's upscaler).
+- **Cause:** a warp stores tiles to shared memory (STS) and reads them back with LDSM, with no BAR between them. 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 moved them past each other. The wrong pixels are mostly at odd x, the high halves of packed pairs.
+- **Fix:** `InstEmitAmpereTensor.ALdsm` puts a subgroup barrier (`Instruction.SubgroupBarrier`, `simdgroup_barrier(mem_flags::mem_threadgroup)`) before its loads and another after them, so that no lane's next store reaches the rows before every lane has read them. LDSM is warp-collective, so the whole SIMD-group reaches both. A SIMD-group-scope fence after each store, with no execution barrier, left about 3,800 such pixels a frame.
+- **Check:** `test_dmsh_ampere_mem.py` `Shared.test_ldsm_reads_between_subgroup_barriers`. In the game, pixels of 6231 that differ from both vertical neighbours while those agree: 6,700-14,000 a frame before, 71-86 after (the YUV source: 40-58). The GPU time is unchanged.
+- **Found on:** The Legend of Zelda: Ocarina of Time (`0400486024DEC000`, profile in #1299): the movie library's NDVE shaders, five of which (`ba079f68ae288dba`, `a8b4fb48fee0ccf8`, `e6be82d835fbe4d3`, `7f7ff0f36b54d90f`, `5a7d69b067a7f011`) have no BAR between their STS and LDSM.
+
 ### Driver-segment reads outside fragment shaders
 
 - **Symptom:** a vertex or compute shader reads zeros from NVN's driver segment `c0`.