MSL: shaders with 64-bit atomics fail Metal compilation

[lkmavi]

Summary

Two snapshot shaders — atomicCompareExchange-int64 and atomicOps-int64 — compile successfully through the Go MSL backend but are rejected by xcrun metal on macOS. Metal's support for 64-bit integer atomics is incomplete: most operations (min, max, and, or, xor, compare_exchange_weak) do not exist for long/ulong in the MSL standard library, and even add/sub require GPU Family 7+ (A15 / M2+) with metal3.0.

Discovered during M3 Max hardware verification of PR #77. The workgroup variable declarations are emitted correctly post-PR #77 — this is a separate, pre-existing limitation of Metal itself.

Affected shaders

Shader	Address spaces	Operations
atomicOps-int64	storage, workgroup, struct	store, load, add, sub, min, max, and, or, xor, exchange, compareExchangeWeak on atomic<i64> / atomic<u64>
atomicCompareExchange-int64	storage	atomicCompareExchangeWeak on atomic<i64> / atomic<u64>

Root cause

Metal supports a narrow subset of 64-bit atomic operations:

Operation	Metal support
atomic_fetch_add / atomic_fetch_sub on ulong/long	GPU Family 7+ (A15 / M2+), metal3.0 only
atomic_fetch_min / atomic_fetch_max	Not available
atomic_fetch_and / atomic_fetch_or / atomic_fetch_xor	Not available
atomic_compare_exchange_weak	Not available
64-bit atomics in threadgroup address space	Additionally restricted

The generated MSL calls metal::atomic_*_explicit intrinsics on metal::atomic_long / metal::atomic_ulong operands. xcrun metal rejects the code at overload resolution because these overloads do not exist in the active MSL standard.

Impact

• Both MSL snapshot tests are silently skipped on Darwin (t.Skipf("MSL compile failed (skipping)"), snapshot_test.go:1327).
• Any WebGPU application that uses atomic<i64> / atomic<u64> on a Metal backend will receive a shader compilation error at runtime.
• Affects MSL backend only. SPIR-V, HLSL, GLSL outputs are unaffected.

Possible approaches

1. Early backend error — return a descriptive error from msl.Compile when the IR contains 64-bit integer atomics, rather than emitting code that the Metal compiler rejects.
2. Features capability flag — add Features.Int64Atomics to msl.Options; emit an error when the flag is unset and the IR requires 64-bit atomics; lets callers gate on capability before dispatch.
3. Partial emulation — for add/sub on metal3.0+GPU-Family-7 targets, emit atomic_fetch_add_explicit; error on unsupported ops. Requires a target feature flag in options.

Option 1 is the minimal correct fix. Option 2 is the right long-term API shape.

Related

• PR fix(msl): declare workgroup vars at function-body scope in kernel entry points #77 — MSL workgroup-var function-scope fix (not the root cause here)
• snapshot/testdata/golden/msl/atomicOps-int64.msl
• snapshot/testdata/golden/msl/atomicCompareExchange-int64.msl
• snapshot/snapshot_test.go:112 — referenceAllowList for these shaders
• MSL spec §6.13 (Atomic Functions), Table 6.6