[LangRef] Clarify interaction of noalias and synchronization

[nikic]

Noalias applies to accesses on other threads, but this was not very clear in existing wording, because "during the execution of the function" is somewhat ambiguous.

Explicitly mention that it applies to accesses from other threads. This means that conflicting accesses (i.e. not read-read) need to either happen-before function entry, or function exit needs to happen-before them, otherwise behavior is undefined. For accesses not based on the noalias pointer, this requires synchronization outside the function.

[llvmorg-github-actions]

@llvm/pr-subscribers-llvm-ir

Author: Nikita Popov (nikic)

Changes

Noalias applies to accesses on other threads, but this was not very clear in existing wording, because "during the execution of the function" is somewhat ambiguous.

Explicitly mention that it applies to accesses from other threads. This means that conflicting accesses (i.e. not read-read) need to either happen-before function entry, or function exit needs to happen-before them, otherwise behavior is undefined. For accesses not based on the noalias pointer, this requires synchronization outside the function.

---

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

1 Files Affected:

• (modified) llvm/docs/LangRef.md (+6)

diff --git a/llvm/docs/LangRef.md b/llvm/docs/LangRef.md
index 415d015d39a93..c1d03be56aff4 100644
--- a/llvm/docs/LangRef.md
+++ b/llvm/docs/LangRef.md
@@ -1522,6 +1522,12 @@ Currently, only the following parameter attributes are defined:
     met. For further details, please see the discussion of the NoAlias response
     in {ref}`alias analysis <Must, May,  or No>`.
 
+    `noalias` also applies to accesses from other threads, unless they
+    happen-before function entry, or function exit happens-before them. This
+    means that conflicting concurrent accesses from other threads either need
+    to be based on the noalias pointer, or else be appropriately synchronized
+    *outside* the function.
+
     Note that this definition of `noalias` is intentionally similar
     to the definition of `restrict` in C99 for function arguments.
 

[michaelselehov]

cc @arsenm

This clarification has direct relevance for GPU/SPMD code. HIP/OpenCL frontends lower restrict on kernel pointer parameters to noalias, including workgroup-shared buffers that are concurrently accessed by sibling threads and synchronized inside the kernel. Under the wording added here, those uses become UB, which affects existing and widespread GPU codegen patterns — so a note that this is the intended reading would be valuable.

It also raises a design question: SPMD code typically wants the intra-thread guarantee (no other pointer in this thread aliases the parameter) but not the cross-thread no-concurrent-access guarantee — i.e. synchronizing operations should still be treated as clobbering. Is there appetite for a weaker attribute expressing exactly that? It would let such kernels stay correct without dropping aliasing precision entirely. Context and a motivating miscompile are in #211486.

[gonzalobg]

HIP/OpenCL frontends lower restrict on kernel pointer parameters to noalias, including workgroup-shared buffers that are concurrently accessed by sibling threads and synchronized inside the kernel. Under the wording added here, those uses become UB, which affects existing and widespread GPU codegen patterns

CUDA C++ does the same, but in CUDA C++ the illegal uses described by the new text are already undefined behavior and we aggressively optimize under that assumption today, so this change LGTM.

It also raises a design question: SPMD code typically wants the intra-thread guarantee (no other pointer in this thread aliases the parameter) but not the cross-thread no-concurrent-access guarantee — i.e. synchronizing operations should still be treated as clobbering. Is there appetite for a weaker attribute expressing exactly that?

The current C restrict semantics which CUDA C++ inherits makes restrict essentially useless for a wide range of concurrent programs. What we’d like is something that composes with concurrency, e.g., this example:

__global__ void kernel(
  int* restrict p, 
  int* restrict q
) 
{ 
  {
    for ...
     p[i] = q[i]
  }
  __syncthreads(); // acq_rel barrier arrive+wait
  if (threadIdx.x == 0) q[i] = 42; // modify q[i] from other thread
  __syncthreads();
  {
     for ...
      p[i] = q[i]
  }
}

exhibits undefined behavior, because threadIdx.x == 0 modifies q[i] via a restrict pointer that is different from e.g. the q restrict pointer that threadIdx.x == 1 is using, while the lifetime of both the thread 0 and 1 kernel functions are live.

Instead what we'd want is to be able to assume that p and q don’t alias - i.e. are disjoint - before the first syncthreads() and after the second syncthreads(), so that the loops can be optimized accordingly, but in contrast to C's restrict, we want to allow other threads to modify the data through their own local pointers as long as there is no data-race (e.g. to allow threadIdx.x == 0 to modify q[i] above; so that caching q[i] before the syncthreads, and reusing the cached value after, would not be a sound optimization).

My preference would be to update C's restrict to just compose with concurrency and support the above; if C's restrict semantics were not updated when the C11 memory model was added, that may be an oversight (C++ does not have restrict, so the C++11 memory model change did not have to deal with this interaction). I don't know whether noalias should provide that, or whether we need a new LLVM IR attribute (Rust may want to be more aggressive in terms of optimizations than C here, see: rust-lang/unsafe-code-guidelines#572, so a new attribute may be needed, idk ).

[michaelselehov]

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:

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.

[RalfJung]

LGTM, thanks!

[gonzalobg]

@michaelselehov I love your idea of strengthening C's restrict at the parallel programming model level to support these cases and implementing that by lowering to alias scope metadata.

[RalfJung]

I don't think we have any proposal for a precise semantics of alias scope metadata so I would caution about premature optimism in that regard.^^

I would expect that it has similar requirements wrt. accesses from other threads having to happen-before the beginning or happen-after the end of the scope.

[michaelselehov]

@RalfJung fair point, I don't want to rely on unspecified semantics either. Let me be precise about the reading we actually need, since it's narrower than what your caution is about, and it doesn't rely on in-function synchronization at all.

All we want from lowering these restrict parameters is intra-thread pairwise disjointness: within one thread's execution, accesses through p and accesses through q don't alias each other. It says nothing about other threads; peers are free to access the same memory. Unlike the noalias attribute, this needs no cross-thread happens-before to hold, so we don't have to (and don't) appeal to the in-kernel barrier to justify it. That's the difference: the attribute's cross-thread clause is the thing you can't rescue with an in-function fence, and the pairwise claim never makes that clause in the first place.

The source guarantee is just "p and q are disjoint in this thread":

__global__ void kernel(int* restrict p, int* restrict q) {
  // p and q don't alias each other in this thread; nothing is claimed about other threads
  ...
}

Two behaviors we need from it, on current IR (opt -passes=gvn):

; (1) spatial disambiguation, thread-independent, no fence involved:
%a = load i32, ptr %p, !alias.scope !Sp, !noalias !Sq
store i32 %x, ptr %q, !alias.scope !Sq, !noalias !Sp   ; can't touch p
%b = load i32, ptr %p, !alias.scope !Sp, !noalias !Sq  ; forwarded to %a (p != q)

; (2) no cross-thread exclusivity: the metadata doesn't exempt synchronization:
%c = load i32, ptr %p, !alias.scope !Sp, !noalias !Sq
fence syncscope("workgroup") acq_rel
%d = load i32, ptr %p, !alias.scope !Sp, !noalias !Sq  ; kept (the barrier clobbers p)

(1) is the precision we want, and it holds regardless of other threads: a concurrent peer write to p with nothing synchronizing in between is a data race (UB anyway), and if there is synchronization then there's a barrier sitting there that stops the forward. (2) is where this differs from the noalias attribute, which forwards %d from %c across the barrier through the sync exemption. That's the miscompile we ran into. The scoped form doesn't claim anything that would let the barrier be dropped.

Ordering between producer and consumer is the barrier's job, a normal happens-before the program needs anyway, and it's separate from the aliasing claim. The aliasing claim never depends on it. That's why the pairwise scoped form is fine where the attribute isn't.

I agree the exact semantics of alias-scope metadata isn't written down, and it'd be good to pin it down. But even the minimal "intra-thread pairwise disjointness" reading is enough to make these kernels correct, and it carries no other-thread happens-before obligation because it makes no other-thread claim. Happy to write this up as a concrete proposal if that helps.

[RalfJung]

"p and q are disjoint in this thread":

I don't know what this means.

Let me try to make a guess:
noalias on a pointer p says "accesses through this pointer do not alias accesses done through any other pointer". This, naturally, includes accesses done via other pointers in other threads, so some interaction with happens-before is needed. Consequently, alias analysis can return "noalias" for p relative to any other pointer (that's not derived from p).

It seems like you just want to say "accesses through pointer p do not alias accesses done through pointer q", without making any statement about accesses done through pointers that are derived neither from p nor from q (such as the pointers that exist in other threads)? So, alias analysis should say "noalias" for p vs q, but should say "maybe alias" for p vs some other pointer r?

However you also keep talking about threads as if they were somehow part of the definition of which aliasing is permitted, and I just can't understand what that could possibly mean here. I would expect that it doesn't matter where that other pointer r is used, this thread or other threads. So if I do

%c = load i32, ptr %p, !alias.scope !Sp, !noalias !Sq
// insert here a call to some argumentless function we can't analyze (but no synchronization or concurrency)
%d = load i32, ptr %p, !alias.scope !Sp, !noalias !Sq

then could that other function have used an alias r to clobber p or not?

[michaelselehov]

@RalfJung Yes, you've got it exactly, and your phrasing is better than mine: we want NoAlias(p, q) and MayAlias(p, r) for any r not derived from p or q, with nothing about threads in the definition. Drop the "in this thread" wording, that was just confusing.

That narrow claim is all we need: it enables exactly the p/q disambiguation @gonzalobg and we want (reorder/forward p across a q access), and because it says nothing about any other pointer, it needs no happens-before. Ordering between producer and consumer stays the barrier's job, separate from aliasing.

then could that other function have used an alias r to clobber p or not?

Right, it could clobber p, so %d has to be reloaded. The scope metadata only says p doesn't alias q; it says nothing about that call, so the call is may-alias against p like any other unknown access. The only thing that could let %d forward is ordinary AA proving on its own that the call can't reach p (e.g. p is a non-escaping local), which has nothing to do with the p/q scopes.

[RalfJung]

That definitely sounds like a reasonable thing to express. It is not what noalias does (and I don't think we should change noalias); I have no idea if it is what !alias.scope is meant to do as I know very little about that.
Also coming up with a precise (operational) spec for that sounds like an interesting but so far unsolved problem. ;)

But anyway I think for this PR we just have to agree what noalias does, and the PR certainly agrees with my understanding of how LLVM uses noalias, it agrees with how we'd like noalias to work from a Rust perspective, and it agrees with my reading of the C standard for restrict.