[llvm] [LoopInfo] Don't recognize loop as parallel if it stores to out-of-loop alloca (PR #180551)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Feb 17 03:16:55 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-analysis
Author: Julius Ikkala (juliusikkala)
<details>
<summary>Changes</summary>
Fixes #<!-- -->179272. Related discussion on Discourse: https://discourse.llvm.org/t/semantics-of-llvm-loop-parallel-accesses-and-interaction-with-alloca/89714
## Background
The parallel loop metadata `llvm.loop.parallel_accesses` works in such a way that it is implicitly invalidated by unaware passes that make transformations that are incompatible with it. Currently, this works by annotating loads and stores with the `llvm.access.group` metadata. If the metadata is dropped for any reason, the loop is no longer considered parallel.
This invalidation is necessary to avoid miscompilations, as parallel-loop unaware passes may introduce transformations that are correct in sequential loops but incorrect in parallel loops. This is done this way because metadata is always optional and passes must not produce broken code even when they don't handle some metadata.
## Why `alloca` needs metadata too
This parallel loop metadata mechanism currently has a defect in that it does not cover `alloca` insts. Those instructions are hoisted out of parallel loops by several passes that are unaware of the parallel loop metadata. This transformation is illegal in a parallel loop, because the same memory allocation becomes shared between iterations, causing race conditions and ending up in issues like #<!-- -->179272. I found many other ways to trigger the same problem:
* loop body being an inlined function with a long enough local array
* `#pragma clang loop vectorize(assume_safety)` instead of `#pragma omp simd`
* `#pragma omp private(t)` where `t` is a long enough array that it generates an alloca
To fix this, this PR requires that the `alloca`s that are being stored to in a parallel loop must be annotated with the corresponding `llvm.access.group`, similarly to how load and store insts need `llvm.access.group` for the loop to be considered parallel.
This approach allows users and compilers that manually perform privatization for `alloca` (e.g. [pocl](https://github.com/pocl/pocl)) to continue with that kind of usage, by annotating the privatized `alloca` with the `llvm.access.group` corresponding to the loop it's used with.
---
Full diff: https://github.com/llvm/llvm-project/pull/180551.diff
4 Files Affected:
- (modified) llvm/docs/LangRef.rst (+47-27)
- (modified) llvm/docs/ReleaseNotes.md (+6)
- (modified) llvm/lib/Analysis/LoopInfo.cpp (+32-16)
- (added) llvm/test/Analysis/LoopInfo/annotated-parallel-alloca.ll (+57)
``````````diff
diff --git a/llvm/docs/LangRef.rst b/llvm/docs/LangRef.rst
index 50a2515f69189..563506c956785 100644
--- a/llvm/docs/LangRef.rst
+++ b/llvm/docs/LangRef.rst
@@ -8176,10 +8176,10 @@ as it is not affected by the ``llvm.loop.disable_nonforced`` metadata.
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
``llvm.access.group`` metadata can be attached to any instruction that
-potentially accesses memory. It can point to a single distinct metadata
-node, which we call access group. This node represents all memory access
-instructions referring to it via ``llvm.access.group``. When an
-instruction belongs to multiple access groups, it can also point to a
+potentially accesses or allocates memory. It can point to a single distinct
+metadata node, which we call access group. This node represents all memory
+access or allocation instructions referring to it via ``llvm.access.group``.
+When an instruction belongs to multiple access groups, it can also point to a
list of accesses groups, illustrated by the following example.
.. code-block:: llvm
@@ -8201,42 +8201,62 @@ situation that the content must be updated which, because metadata is
immutable by design, would required finding and updating all references
to the access group node.
-The access group can be used to refer to a memory access instruction
-without pointing to it directly (which is not possible in global
+The access group can be used to refer to a memory access or allocation
+instruction without pointing to it directly (which is not possible in global
metadata). Currently, the only metadata making use of it is
``llvm.loop.parallel_accesses``.
'``llvm.loop.parallel_accesses``' Metadata
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-The ``llvm.loop.parallel_accesses`` metadata refers to one or more
-access group metadata nodes (see ``llvm.access.group``). It denotes that
-no loop-carried memory dependence exist between it and other instructions
-in the loop with this metadata.
+The ``llvm.loop.parallel_accesses`` metadata is used to explicitly declare a
+loop as "trivially parallel", indicating that there are no memory dependencies
+between iterations. If a loop has this metadata but also has memory dependencies
+between iterations, the behavior is undefined.
+
+``llvm.loop.parallel_accesses`` refers to one or more access group metadata
+nodes (see ``llvm.access.group``). In a loop with this metadata, instructions
+may have the ``llvm.access.group`` metadata to denote that no loop-carried
+memory dependence exist between those instructions, as long as the access group
+is listed in ``llvm.loop.parallel_accesses``.
Let ``m1`` and ``m2`` be two instructions that both have the
-``llvm.access.group`` metadata to the access group ``g1``, respectively
-``g2`` (which might be identical). If a loop contains both access groups
-in its ``llvm.loop.parallel_accesses`` metadata, then the compiler can
+``llvm.access.group`` metadata to the access groups ``g1`` and ``g2``
+respectively (the groups can be identical). If a loop contains both access
+groups in its ``llvm.loop.parallel_accesses`` metadata, then the compiler can
assume that there is no dependency between ``m1`` and ``m2`` carried by
this loop. Instructions that belong to multiple access groups are
considered having this property if at least one of the access groups
matches the ``llvm.loop.parallel_accesses`` list.
-If all memory-accessing instructions in a loop have
-``llvm.access.group`` metadata that each refer to one of the access
-groups of a loop's ``llvm.loop.parallel_accesses`` metadata, then the
-loop has no loop carried memory dependencies and is considered to be a
-parallel loop. If there is a loop-carried dependency, the behavior is
-undefined.
-
-Note that if not all memory access instructions belong to an access
-group referred to by ``llvm.loop.parallel_accesses``, then the loop must
-not be considered trivially parallel. Additional
-memory dependence analysis is required to make that determination. As a
-fail-safe mechanism, this causes loops that were originally parallel to be considered
-sequential (if optimization passes that are unaware of the parallel semantics
-insert new memory instructions into the loop body).
+A loop is declared to be trivially parallel (as in, there are no memory
+dependencies between loop iterations) if it has the
+``llvm.loop.parallel_accesses`` metadata referring to a set of access groups
+``G`` and fulfills the following conditions:
+
+- All memory-accessing instructions in the loop must have ``llvm.access.group``
+ metadata referring to at least one access group in ``G``.
+- If the loop has one or more instructions that write to memory allocated via
+ ``alloca``, the corresponding ``alloca`` instruction(s) must have the
+ ``llvm.access.group`` metadata referring to at least one access group in ``G``.
+
+If the above conditions are not fulfilled, the loop must not be considered as
+trivially parallel without further memory dependence analysis.
+
+These conditions exist as a fail-safe mechanism to cause loops that were
+originally explicitly parallel to be demoted to normal sequential loops when
+optimization passes that are unaware of the parallel semantics perform
+transformations or insert new instructions into the loop body. Note that even
+if the loop is no longer considered trivially parallel, it may still be
+vectorizable. In such cases the loop must be analyzed as any sequential loop is.
+
+For example, if a pass adds a new memory-accessing instruction into the
+loop body without being aware of the parallel semantics, that instruction will
+not have the corresponding ``llvm.access.group`` metadata, thereby demoting the
+loop into a sequential one. Another example where the fail-safe triggers is
+when a pass hoists an ``alloca`` instruction outside of the loop body. Without
+the fail-safe, this would cause the same allocation to be shared across
+iterations, introducing race conditions in a parallel loop.
Example of a loop that is considered parallel due to its correct use of
both ``llvm.access.group`` and ``llvm.loop.parallel_accesses``
diff --git a/llvm/docs/ReleaseNotes.md b/llvm/docs/ReleaseNotes.md
index 128c19296e75e..38b7f0d6b15e8 100644
--- a/llvm/docs/ReleaseNotes.md
+++ b/llvm/docs/ReleaseNotes.md
@@ -69,6 +69,12 @@ Changes to the LLVM IR
* The `"nooutline"` attribute is now writen as `nooutline`. Existing IR and
bitcode will be automatically updated.
+* To be considered parallel, [loops with `llvm.loop.parallel_accesses` metadata
+ now require corresponding `llvm.access.group` metadata to be present on all
+ `alloca` instructions whose address range is being written to in the loop.](https://discourse.llvm.org/t/semantics-of-llvm-loop-parallel-accesses-and-interaction-with-alloca/89714)
+ If this metadata is not present, such loops are no longer considered parallel
+ and memory dependency checks are not skipped.
+
Changes to LLVM infrastructure
------------------------------
diff --git a/llvm/lib/Analysis/LoopInfo.cpp b/llvm/lib/Analysis/LoopInfo.cpp
index a364b21c64b01..4bad9381f4b38 100644
--- a/llvm/lib/Analysis/LoopInfo.cpp
+++ b/llvm/lib/Analysis/LoopInfo.cpp
@@ -591,23 +591,39 @@ bool Loop::isAnnotatedParallel() const {
if (!I.mayReadOrWriteMemory())
continue;
- if (MDNode *AccessGroup = I.getMetadata(LLVMContext::MD_access_group)) {
- auto ContainsAccessGroup = [&ParallelAccessGroups](MDNode *AG) -> bool {
- if (AG->getNumOperands() == 0) {
- assert(isValidAsAccessGroup(AG) && "Item must be an access group");
- return ParallelAccessGroups.count(AG);
- }
-
- for (const MDOperand &AccessListItem : AG->operands()) {
- MDNode *AccGroup = cast<MDNode>(AccessListItem.get());
- assert(isValidAsAccessGroup(AccGroup) &&
- "List item must be an access group");
- if (ParallelAccessGroups.count(AccGroup))
- return true;
- }
- return false;
- };
+ auto ContainsAccessGroup = [&ParallelAccessGroups](MDNode *AG) -> bool {
+ if (AG->getNumOperands() == 0) {
+ assert(isValidAsAccessGroup(AG) && "Item must be an access group");
+ return ParallelAccessGroups.count(AG);
+ }
+
+ for (const MDOperand &AccessListItem : AG->operands()) {
+ MDNode *AccGroup = cast<MDNode>(AccessListItem.get());
+ assert(isValidAsAccessGroup(AccGroup) &&
+ "List item must be an access group");
+ if (ParallelAccessGroups.count(AccGroup))
+ return true;
+ }
+ return false;
+ };
+
+ // If the loop contains a store instruction into an alloca that is outside
+ // of the loop, it is possible that the alloca was initially related to a
+ // loop-local variable but got hoisted outside during e.g. inlining or
+ // some other parallel-loop-unaware pass. However, if the alloca itself
+ // has been marked with the access group metadata, this usage has to be
+ // assumed to be valid.
+ if (StoreInst *SI = dyn_cast<StoreInst>(&I)) {
+ AllocaInst *AI = findAllocaForValue(SI->getPointerOperand());
+ if (AI) {
+ MDNode *AccessGroup = AI->getMetadata(LLVMContext::MD_access_group);
+ if (AI && !contains(AI) &&
+ (!AccessGroup || !ContainsAccessGroup(AccessGroup)))
+ return false;
+ }
+ }
+ if (MDNode *AccessGroup = I.getMetadata(LLVMContext::MD_access_group)) {
if (ContainsAccessGroup(AccessGroup))
continue;
}
diff --git a/llvm/test/Analysis/LoopInfo/annotated-parallel-alloca.ll b/llvm/test/Analysis/LoopInfo/annotated-parallel-alloca.ll
new file mode 100644
index 0000000000000..b4e5af07950c6
--- /dev/null
+++ b/llvm/test/Analysis/LoopInfo/annotated-parallel-alloca.ll
@@ -0,0 +1,57 @@
+; RUN: opt -passes='print<loops>' -disable-output %s 2>&1 | FileCheck %s
+;
+; void func(long n, long *A) {
+; #pragma clang loop vectorize(assume_safety)
+; for (long i = 0; i < n; i += 1) {
+; long t[32];
+; for (long j = 0; j < 32; j += 1)
+; t[j] = i;
+; A[i] = t[i];
+; }
+; }
+;
+; The alloca for `t` usually gets hoisted outside of the loop (either by Clang
+; itself, or by an inlining pass if the loop body is in a function, etc.) and
+; gets incorrectly shared between iterations. Check that isAnnotatedParallel is
+; blocking this kind of usage, as it will not get vectorized correctly unless
+; mem2reg converts the array.
+;
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+
+define void @func(i64 %n, ptr noalias nonnull %A) {
+entry:
+ %t = alloca [32 x i64], align 16
+ %cmp17 = icmp sgt i64 %n, 0
+ br i1 %cmp17, label %for.body, label %for.cond.cleanup
+
+for.body:
+ %i.018 = phi i64 [ %add8, %for.cond.cleanup3 ], [ 0, %entry ]
+ br label %for.body4
+
+for.body4:
+ %j.016 = phi i64 [ 0, %for.body ], [ %add, %for.body4 ]
+ %arrayidx = getelementptr inbounds nuw i64, ptr %t, i64 %j.016
+ store i64 %i.018, ptr %arrayidx, align 8, !llvm.access.group !9
+ %add = add nuw nsw i64 %j.016, 1
+ %exitcond.not = icmp eq i64 %add, 32
+ br i1 %exitcond.not, label %for.cond.cleanup3, label %for.body4
+
+for.cond.cleanup3:
+ %arrayidx5 = getelementptr inbounds nuw i64, ptr %t, i64 %i.018
+ %0 = load i64, ptr %arrayidx5, align 8, !llvm.access.group !9
+ %arrayidx6 = getelementptr inbounds nuw i64, ptr %A, i64 %i.018
+ store i64 %0, ptr %arrayidx6, align 8, !llvm.access.group !9
+ %add8 = add nuw nsw i64 %i.018, 1
+ %exitcond19.not = icmp eq i64 %add8, %n
+ br i1 %exitcond19.not, label %for.cond.cleanup, label %for.body, !llvm.loop !10
+
+for.cond.cleanup:
+ ret void
+}
+
+!9 = distinct !{}
+!10 = distinct !{!10, !11}
+!11 = !{!"llvm.loop.parallel_accesses", !9}
+
+; CHECK: Loop info for function 'func':
+; CHECK-NOT: Parallel Loop at depth 1 containing:
``````````
</details>
https://github.com/llvm/llvm-project/pull/180551
More information about the llvm-commits
mailing list