[maximilianfv]

What

Under MoltenVK, agc_driver_lds_tree_reduction and agc_driver_lds_subword_store_branch read LDS before another wave had written it. This makes SubgroupControlBarrierUsable (SubgroupBarrier.hpp) return false for VK_DRIVER_ID_MOLTENVK as well, so MoltenVK takes the subgroup-scope OpMemoryBarrier path that AMD's proprietary driver already uses (#8134). WorkgroupScopeFenceUsable does not change: MoltenVK keeps workgroup-scope fences, and the subword store test passes with them.

Cause, from a MoltenVK shader dump: the recompiler orders one guest wave's LDS accesses with OpControlBarrier (Subgroup execution scope, Workgroup memory scope) inside the branch after s_cbranch_execz. SPIRV-Cross's CompilerMSL::emit_barrier takes the wider of the execution and memory scopes and emits threadgroup_barrier(mem_flags::mem_threadgroup). The tree-reduction dump has 22 threadgroup_barriers for the guest's 5 s_barriers. Waves that skip the branch pair their next barrier with another wave's extra one and read LDS before it is written: the tree reduction read 904, and the subword store read changing garbage. With this change the subgroup-scope OpMemoryBarrier becomes a simdgroup-scope atomic_thread_fence with no execution barrier. That relies on a 32-lane simdgroup running in lockstep, which Metal does not guarantee; the new debt entry "Wave LDS ordering on MoltenVK" in docs/dev/debt/shader.md records this.

The agc_subgroup_barrier driver table (tests/SubgroupBarrier.cpp) used one usable flag for both functions. MoltenVK now differs between the two (false, true), so the flag is split into controlBarrier and workgroupFence. Every other row keeps the old value in both fields.

On Linux and Windows the generated code does not change: the driver ID is never MoltenVK there.

Prior art and credit:

• Ported from Himansh Raj's (iamthehimansh) commit 50f545c "fix(agc): order wave LDS accesses with memory barriers under MoltenVK" (2026-10-21) in the closed feat: macOS host support (Apple Silicon via Rosetta 2, Vulkan on MoltenVK) #8658. The change is the same, with three differences. That version was macOS-only (#ifdef __APPLE__) in both files; this one is unconditional, because the driver ID alone already selects MoltenVK. It had code comments, which are dropped under the no-comments rule. It flipped the single usable flag, but main later added WorkgroupScopeFenceUsable, which shares the table (fix(shader): widen LDS atomic fences to device scope on AMD's proprietary driver #8489), so this version splits the flag instead. Credited with a Co-authored-by: trailer.
• None of the PRs feat: macOS host support (Apple Silicon via Rosetta 2, Vulkan on MoltenVK) #8658 was split into (build: compile the shared libc headers, thread priority and three tests on macOS #10362, build: compile FileStream, error categories, thread priority and two tests on macOS #10419) carries this commit. Macos port #7201 (marion909) and the closed feat: macOS on Apple Silicon host support (Rosetta 2 and MoltenVK) #7332 (ycarogabriel-lgtm) do not change SubgroupBarrier.hpp or the subgroup barrier test. Neither does mugurc's macOS series.
• Open PRs touching SubgroupBarrier.hpp: fix(agc): keep a DCC target's refresh proof while its keys are pending driver stores #10159 carries only the WorkgroupScopeFenceUsable addition already on main from fix(shader): widen LDS atomic fences to device scope on AMD's proprietary driver #8489, on an older base, and does not change MoltenVK's path. fix(agc): create tessellation pipelines without the pipeline cache on MoltenVK #10046 does not touch the header; its VulkanDevice.cpp diff only shows the SubgroupControlBarrierUsable call as context.

Tested

Apple M1 Pro, macOS 27.0.1, x86_64 build under Rosetta 2, Vulkan through MoltenVK 1.4.2 from the LunarG Vulkan SDK 1.4.363. main does not build on macOS by itself, so I tested on mugurc's macOS series (ci/macos-series-2, 27e9689) merged with main at 522e32e. The tree also carried other open macOS fixes (#10419 build parts, #10540, #10562, #10641, #10642, #10650, #10651, #10652) and one of my unopened test fixes (agc_driver_graphics). Only this change was toggled between the runs. I tested before rebasing onto current main; the rebased diff is the same.

Both sides: ANYPS5_REQUIRE_VULKAN=1 ANYPS5_NO_SHADER_CACHE=1 VK_INSTANCE_LAYERS=VK_LAYER_KHRONOS_validation ctest --test-dir build -R '^(agc_subgroup_barrier|agc_driver_lds_.*)$' -E 'audio_in|voice' -V (16 tests), three runs each.

• Without: 14/16 passed in each run; agc_driver_lds_subword_store_branch and agc_driver_lds_tree_reduction failed.
• With: 16/16 passed in each run.
• Validation: neither log has the was not found but was requested by env var VK_INSTANCE_LAYERS line, and both have 0 Validation Error lines.

Not run on Linux or Windows: the generated code does not change there, and agc_subgroup_barrier is a table check that does not depend on the host. Not tested with any title.

Checklist

• Based on current main; no other open PR implements the same functions
• One topic per PR; follow-ups go in a new PR
• No comments in code except technical debt
• Unimplemented paths throw (NotImplemented_nid_no_patch); silent stubs are listed under Silent stubs in their library's technical debt file
• No notes, investigation .md files or images added to the repository (attach them to this PR)
• New third-party code is a submodule built from source
• Generic behaviour, not specific to one title
• Depends on: none
• Fixes: none
• Unblocks: none
• GPUs tested: Apple M1 Pro, MoltenVK 1.4.2 (Vulkan SDK 1.4.363), macOS 27.0.1
• AI-assisted: yes

🤖 Generated with Claude Code

[MarcZX10]

Comment-only review of exact head 5138b1b1d25664561055b3b5b5810d9780dd93cb.

The failure mechanism is plausible and the change makes the two MoltenVK regression tests pass in the reported M1 Pro setup. I am not approving yet because the new VK_DRIVER_ID_MOLTENVK policy applies the fence-only path to every MoltenVK device, while the evidence is one M1 Pro and the debt entry itself says the path relies on 32-lane SIMD-group lockstep that Metal does not guarantee. Apple's guidance says SIMD-group size varies by GPU and advises querying it at runtime; it also distinguishes simdgroup_barrier (which waits for all threads in the SIMD group) from a memory fence. See Apple's SIMD-group and synchronization guidance.

Could you clarify whether SubgroupControlBarrierUsable is reached only when MoltenVK's reported Vulkan subgroup size matches the assumed guest wave size, and what evidence supports the fence-only lowering on other MoltenVK GPUs? If the current evidence only establishes M1 Pro, please consider limiting the policy to a supported subgroup/device condition or documenting the exact tested support boundary. I have not independently run the Metal/Vulkan tests; the three-run results and generated shader observations are author-reported.

AI-assisted review: yes.

[github-actions]

Some comparisons could not be completed because a pull request conflicts with the base branch. Review these shared files manually before concluding that there is no overlap.

Pull request	Shared files not compared
PR 5021	docs/dev/debt/shader.md
PR 5371	docs/dev/debt/shader.md
PR 6235	docs/dev/debt/shader.md
PR 6280	docs/dev/debt/shader.md
PR 6412	docs/dev/debt/shader.md
PR 6488	docs/dev/debt/shader.md
PR 6924	docs/dev/debt/shader.md
PR 7395	docs/dev/debt/shader.md
PR 7410	docs/dev/debt/shader.md
PR 7684	docs/dev/debt/shader.md
PR 7922	docs/dev/debt/shader.md
PR 8004	docs/dev/debt/shader.md
PR 8939	docs/dev/debt/shader.md
PR 9020	docs/dev/debt/shader.md
PR 9335	docs/dev/debt/shader.md
PR 9387	docs/dev/debt/shader.md
PR 9457	docs/dev/debt/shader.md
PR 9633	docs/dev/debt/shader.md
PR 9748	docs/dev/debt/shader.md
PR 9749	docs/dev/debt/shader.md
PR 9790	docs/dev/debt/shader.md
PR 9889	docs/dev/debt/shader.md
PR 9969	docs/dev/debt/shader.md
PR 10000	docs/dev/debt/shader.md
PR 10233	docs/dev/debt/shader.md
PR 10489	docs/dev/debt/shader.md
PR 10846	docs/dev/debt/shader.md
PR 10867	docs/dev/debt/shader.md
PR 10941	docs/dev/debt/shader.md
PR 11557	docs/dev/debt/shader.md

[maximilianfv]

Thanks. Answers from the code at this head (5138b1b1d2):

When the fence-only path is reached. Only when WaveLdsScope returns subgroup scope (SpirvEmitter.cpp#L63-L74). That needs all of these:

• a compute shader that writes LDS;
• a workgroup larger than one guest wave;
• a 32-lane host subgroup that matches a wave32 program, or a wave64 program run at two lanes per invocation on a 32-lane host.

The host subgroup size is the device's reported subgroupSize, or 32 when subgroupSizeControl lets compute pipelines require 32 (VulkanDevice.cpp#L1963-L1970, #L3792). A MoltenVK device that is not 32 lanes wide and cannot require 32 never takes the new branch, so it keeps what it had before this PR. Single-wave workgroups use workgroup scope (#L69) and are also unaffected.

Evidence. M1 Pro only (MoltenVK 1.4.2, macOS 27). I have no other Apple GPU and no Intel or AMD Mac to test on. The debt entry states this boundary and the lockstep assumption.

What Metal gets. In the shaders dumped from the M1 Pro run, every new barrier is atomic_thread_fence(mem_flags::mem_threadgroup, memory_order_seq_cst, thread_scope_simdgroup). SPIRV-Cross (CompilerMSL::emit_barrier) emits that form only for MSL 3.2 and later. For older MSL it emits the same OpMemoryBarrier(Subgroup) as simdgroup_barrier(mem_flags::mem_threadgroup). On those targets the barrier needs the whole SIMD group to reach it. That is the same subgroup-uniform control flow that the OpControlBarrier(Subgroup, …) other drivers get already requires, so this PR does not add that assumption. I have not run a pre-3.2 target.

[ycbuzi]

A Windows build and run of this head: Windows 11 Home 10.0.26200, WinLibs GCC 15.2.0 r7 (MinGW-w64 UCRT), Release. 5138b1b1d2 merged onto main 4e43321bd6 with no conflicts.

• agc_subgroup_barrier passes, 2 of 2. It's a CPU-side test (driver IDs through the recompiler, no device opened), so this checks the MinGW build and the table, not a GPU.
• Control: with main's SubgroupBarrier.hpp, it fails with driver 14: expected subgroup control barriers to be unusable, the MoltenVK case.
• The change only adds VK_DRIVER_ID_MOLTENVK to SubgroupControlBarrierUsable, so the NVIDIA and AMD Windows drivers get the same answer as before.

AI-assisted: yes (Claude Code)

[jotitadev]

A second Apple GPU (M2 Max) on a main-based tree, with both MoltenVK builds. The other runs here are Windows, where the change is a no-op.

Machine: Apple M2 Max, macOS 27.2, x86_64 Release build under Rosetta 2.
Tree: main b986c11a34 + the open macOS PRs needed to build and open a device (#10366, #7631, #10659, #10660, #11249, #11358, #10669, #10552, #7804, #10676, #11447, #11437, #11554, #11705) + this PR at 5138b1b1d2 (clean merge). Control: the same tree without this PR. ANYPS5_NO_SHADER_CACHE=1, -j1 --timeout 30.
Drivers: MoltenVK 1.4.2 (LunarG SDK 1.4.363.0) and MoltenVK main + SPIRV-Cross #2704.

	Without	With (3 runs)
agc_driver_lds_tree_reduction	fails on both drivers: "round 0: thread 32 (wave 1) read 904, expected 32896"	passes 3/3 on both drivers
agc_driver_lds_subword_store_branch	passes, 6/6 on each driver	passes 3/3 on both drivers
agc_subgroup_barrier	passes	passes

• The tree-reduction failure is the same value you saw on the M1 Pro (904), so that one reproduces on a second Apple GPU and on MoltenVK main.
• The subword store branch test did not fail here without the change (1 + 5 repeated runs per driver). On this machine it does not show the garbage reads, so it is weaker evidence for the fence-only path than the tree reduction.
• Full ^agc suite, 1037 tests, MoltenVK 1.4.2: 41 failures without, 40 with. lds_tree_reduction is the only difference, so nothing new fails.
• Validation layer (VK_INSTANCE_LAYERS=VK_LAYER_KHRONOS_validation, SDK 1.4.363.0, loader output shows it inserted) on the three tests: no VUID- on either side.

GPUs tested: Apple M2 Max (MoltenVK 1.4.2 and MoltenVK main + #2704)

Written with AI assistance (Claude Code). I ran the builds and tests above on my machine.