[llvm] [LangRef] Clarify interaction of noalias and synchronization (PR #211507)

[via llvm-commits]

michaelselehov wrote:

Thanks @gonzalobg — that example captures it exactly, and I agree with your analysis of why it's UB today.

One framing I'd add for after this lands: once the cross-thread meaning of `noalias` is explicit, a frontend that lowers `restrict` → `noalias` for a kernel parameter that can be concurrently accessed by sibling threads with only *in-kernel* synchronization is emitting an attribute whose promise the program never makes. At that point it isn't really the *user* writing UB — it's an **incorrect lowering**: the frontend generates IR that its own LangRef says is violated. So the immediate correctness fix belongs in the frontend's lowering, not in reinterpreting the attribute or special-casing alias analysis.

The encouraging part is that the semantics you describe — `p` and `q` disjoint within a synchronization-free region so the loops optimize, but a barrier still clobbers so caching across `__syncthreads()` is unsound — are already expressible on today's IR by lowering such `restrict` parameters to scoped `!alias.scope`/`!noalias` metadata instead of the `noalias` attribute. That metadata is a purely intra-thread, per-execution disambiguation: it says nothing about other threads, so it does not grant the whole-function cross-thread exclusivity that causes the miscompile, yet it keeps the intra-thread pairwise disjointness that makes the loops fast. It reuses the noalias→scope conversion the inliner already performs; the frontend would just emit it directly for these parameters.

Concretely, running `opt -passes=gvn`, for two shared buffers `p1`/`p2`:

```llvm
define i32 @two_shared_scoped(ptr %p1, ptr %p2, i32 %x) {
  %a1 = load i32, ptr %p1, !alias.scope !0, !noalias !3
  store i32 %x, ptr %p2, !alias.scope !3, !noalias !0
  %a2 = load i32, ptr %p1, !alias.scope !0, !noalias !3   ; FORWARDED from %a1  -> p1 != p2 (intra-thread precision kept)
  fence syncscope("workgroup") acq_rel
  %a3 = load i32, ptr %p1, !alias.scope !0, !noalias !3   ; KEPT               -> barrier clobbers (cross-thread ordering honored)
  ...
}
```

The `noalias` *attribute*, by contrast, forwards the load across the fence (via the sync exemption) — which is exactly the miscompile.

So my read: scoped-metadata lowering is a pragmatic frontend fix that makes these kernels correct today, with no new attribute and no LangRef change. A dedicated weaker attribute — or, as you'd prefer, `restrict` itself updated to compose with concurrency — would be a cleaner long-term spelling and could be strictly more precise, but it isn't required to unblock correctness. This is how we resolved a real miscompile in AMD's Composable Kernel GEMM (dropping the over-strong promise); background and the motivating case are in #211486. And to your point that this isn't LDS-specific: it applies equally to global memory used for cross-wave synchronization — just less common in practice because it's slower.


https://github.com/llvm/llvm-project/pull/211507