[llvm] [Attributor] Drop norecurse when specializing an indirect call closes a cycle (PR #218637)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 01:34:11 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-transforms
Author: Larry Meadows (lfmeadow)
<details>
<summary>Changes</summary>
Specializing an indirect call turns an edge that only existed through a function
pointer into a direct one. When that closes a call graph cycle, `norecurse` on
the functions in it is no longer true, and nothing removes it: attribute
inference only ever adds `norecurse`, so the Attributor could leave behind a
module in which a `norecurse` function directly recurses.
## Testing
Tested on `main` at e625da37ad1f. `llvm/test/Transforms/Attributor` and
`llvm/test/Transforms/OpenMP` pass (257 passed, 1 unsupported). On gfx90a,
`nested_parallel_reduction.c` prints `aa = 6300, expected 6300` with this
change and `aa = 0` without it.
---
Full diff: https://github.com/llvm/llvm-project/pull/218637.diff
6 Files Affected:
- (modified) llvm/include/llvm/Transforms/IPO/Attributor.h (+8)
- (modified) llvm/lib/Transforms/IPO/Attributor.cpp (+37)
- (modified) llvm/lib/Transforms/IPO/AttributorAttributes.cpp (+2)
- (added) llvm/test/Transforms/Attributor/norecurse_indirect_call.ll (+73)
- (added) llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll (+44)
- (added) offload/test/offloading/nested_parallel_reduction.c (+47)
``````````diff
diff --git a/llvm/include/llvm/Transforms/IPO/Attributor.h b/llvm/include/llvm/Transforms/IPO/Attributor.h
index 6e4d3d92d8106..ef582f15b1a7c 100644
--- a/llvm/include/llvm/Transforms/IPO/Attributor.h
+++ b/llvm/include/llvm/Transforms/IPO/Attributor.h
@@ -1883,6 +1883,11 @@ struct Attributor {
InvokeWithDeadSuccessor.insert(&II);
}
+ /// Record that an indirect call in \p Caller was replaced by direct calls.
+ void recordIndirectCallSpecialization(Function &Caller) {
+ SpecializedIndirectCallers.insert(&Caller);
+ }
+
/// Record that \p I is deleted after information was manifested. This also
/// triggers deletion of trivially dead istructions.
void deleteAfterManifest(Instruction &I) { ToBeDeletedInsts.insert(&I); }
@@ -2590,6 +2595,9 @@ struct Attributor {
SmallSetVector<WeakVH, 8> ToBeDeletedInsts;
///}
+ /// Functions in which an indirect call was replaced by direct calls.
+ SmallSetVector<Function *, 8> SpecializedIndirectCallers;
+
/// Container with all the query AAs that requested an update via
/// registerForUpdate.
SmallSetVector<AbstractAttribute *, 16> QueryAAsAwaitingUpdate;
diff --git a/llvm/lib/Transforms/IPO/Attributor.cpp b/llvm/lib/Transforms/IPO/Attributor.cpp
index 338240b2f8087..2514e138417a3 100644
--- a/llvm/lib/Transforms/IPO/Attributor.cpp
+++ b/llvm/lib/Transforms/IPO/Attributor.cpp
@@ -17,6 +17,7 @@
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/PointerIntPair.h"
+#include "llvm/ADT/SCCIterator.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/Statistic.h"
@@ -2670,6 +2671,42 @@ ChangeStatus Attributor::cleanupIR() {
Configuration.CGUpdater.removeFunction(*Fn);
}
+ // Specializing an indirect call turns an edge that only existed through a
+ // function pointer into a direct one, which can close a call graph cycle and
+ // make `norecurse` false for every function in it.
+ if (!SpecializedIndirectCallers.empty()) {
+ CallGraph CG(*SpecializedIndirectCallers.front()->getParent());
+ SmallPtrSet<Function *, 8> Handled;
+ for (Function *Caller : SpecializedIndirectCallers) {
+ if (!Handled.insert(Caller).second)
+ continue;
+ CallGraphNode *Node = CG[Caller];
+ for (scc_iterator<CallGraphNode *> It = scc_begin(Node); !It.isAtEnd();
+ ++It) {
+ if (!is_contained(*It, Node))
+ continue;
+ if (!It.hasCycle())
+ break;
+ for (CallGraphNode *N : *It) {
+ Function *Fn = N->getFunction();
+ if (!Fn || ToBeDeletedFunctions.count(Fn) || !Functions.count(Fn))
+ continue;
+ Handled.insert(Fn);
+ if (!Fn->hasFnAttribute(Attribute::NoRecurse))
+ continue;
+ Fn->removeFnAttr(Attribute::NoRecurse);
+ for (User *U : Fn->users())
+ if (auto *CB = dyn_cast<CallBase>(U))
+ if (CB->getCalledFunction() == Fn &&
+ Functions.count(CB->getFunction()))
+ CB->removeFnAttr(Attribute::NoRecurse);
+ ManifestChange = ChangeStatus::CHANGED;
+ }
+ break;
+ }
+ }
+ }
+
if (!ToBeChangedUses.empty())
ManifestChange = ChangeStatus::CHANGED;
diff --git a/llvm/lib/Transforms/IPO/AttributorAttributes.cpp b/llvm/lib/Transforms/IPO/AttributorAttributes.cpp
index 8d16f49525ca8..ce15da3074466 100644
--- a/llvm/lib/Transforms/IPO/AttributorAttributes.cpp
+++ b/llvm/lib/Transforms/IPO/AttributorAttributes.cpp
@@ -12530,6 +12530,7 @@ struct AAIndirectCallInfoCallSite : public AAIndirectCallInfo {
// Special handling for the single callee case.
if (AllCalleesKnown && AssumedCallees.size() == 1) {
auto *NewCallee = AssumedCallees.front();
+ A.recordIndirectCallSpecialization(*CB->getFunction());
if (isLegalToPromote(*CB, NewCallee)) {
promoteCall(*CB, NewCallee, nullptr);
NumIndirectCallsPromoted++;
@@ -12564,6 +12565,7 @@ struct AAIndirectCallInfoCallSite : public AAIndirectCallInfo {
continue;
}
SpecializedForAnyCallees = true;
+ A.recordIndirectCallSpecialization(*CB->getFunction());
LastCmp = new ICmpInst(IP, llvm::CmpInst::ICMP_EQ, FP, NewCallee);
Instruction *ThenTI =
diff --git a/llvm/test/Transforms/Attributor/norecurse_indirect_call.ll b/llvm/test/Transforms/Attributor/norecurse_indirect_call.ll
new file mode 100644
index 0000000000000..af8a37fc54ba6
--- /dev/null
+++ b/llvm/test/Transforms/Attributor/norecurse_indirect_call.ll
@@ -0,0 +1,73 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --check-attributes --check-globals all --version 6
+; RUN: opt -passes=attributor -attributor-assume-closed-world -S < %s | FileCheck %s
+
+; @dispatcher is norecurse because its only edge back to itself runs through the
+; function pointer. Specializing the indirect call makes the edge to @outer a
+; direct one and closes the cycle @dispatcher -> @outer -> @dispatcher, so both
+; have to lose the attribute. @inner is not on the cycle and keeps it.
+
+ at g = external global i32
+
+;.
+; CHECK: @g = external global i32
+;.
+define internal void @inner() norecurse {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(write)
+; CHECK-LABEL: define internal void @inner(
+; CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
+; CHECK-NEXT: store i32 1, ptr @g, align 4
+; CHECK-NEXT: ret void
+;
+ store i32 1, ptr @g
+ ret void
+}
+
+define internal void @outer() norecurse {
+; CHECK: Function Attrs: mustprogress nofree nosync nounwind willreturn memory(write)
+; CHECK-LABEL: define internal void @outer(
+; CHECK-SAME: ) #[[ATTR1:[0-9]+]] {
+; CHECK-NEXT: call void @dispatcher(ptr nofree noundef nonnull captures(none) @inner) #[[ATTR2:[0-9]+]]
+; CHECK-NEXT: ret void
+;
+ call void @dispatcher(ptr @inner)
+ ret void
+}
+
+define internal void @dispatcher(ptr %fn) norecurse {
+; CHECK: Function Attrs: mustprogress nofree nosync nounwind willreturn memory(write)
+; CHECK-LABEL: define internal void @dispatcher(
+; CHECK-SAME: ptr nofree noundef nonnull captures(none) [[FN:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[TMP1:%.*]] = icmp eq ptr [[FN]], @outer
+; CHECK-NEXT: br i1 [[TMP1]], label %[[BB2:.*]], label %[[BB3:.*]]
+; CHECK: [[BB2]]:
+; CHECK-NEXT: call void @outer()
+; CHECK-NEXT: br label %[[BB6:.*]]
+; CHECK: [[BB3]]:
+; CHECK-NEXT: br i1 true, label %[[BB4:.*]], label %[[BB5:.*]]
+; CHECK: [[BB4]]:
+; CHECK-NEXT: call void @inner()
+; CHECK-NEXT: br label %[[BB6]]
+; CHECK: [[BB5]]:
+; CHECK-NEXT: unreachable
+; CHECK: [[BB6]]:
+; CHECK-NEXT: ret void
+;
+ call void %fn()
+ ret void
+}
+
+define void @entry() norecurse {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(write)
+; CHECK-LABEL: define void @entry(
+; CHECK-SAME: ) #[[ATTR0]] {
+; CHECK-NEXT: call void @dispatcher(ptr nofree noundef nonnull captures(none) @outer) #[[ATTR2]]
+; CHECK-NEXT: ret void
+;
+ call void @dispatcher(ptr @outer)
+ ret void
+}
+;.
+; CHECK: attributes #[[ATTR0]] = { mustprogress nofree norecurse nosync nounwind willreturn memory(write) }
+; CHECK: attributes #[[ATTR1]] = { mustprogress nofree nosync nounwind willreturn memory(write) }
+; CHECK: attributes #[[ATTR2]] = { nofree nosync nounwind willreturn memory(write) }
+;.
diff --git a/llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll b/llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll
new file mode 100644
index 0000000000000..c09c83322d58a
--- /dev/null
+++ b/llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll
@@ -0,0 +1,44 @@
+; RUN: opt -passes='cgscc(function-attrs),rpo-function-attrs,attributor' \
+; RUN: -attributor-assume-closed-world -S < %s | FileCheck %s
+
+; Nothing in this module marks @dispatcher norecurse. rpo-function-attrs adds it:
+; addNoRecurseAttrsTopDown only walks a function's uses, never its body, so a
+; function whose body is an indirect call qualifies as long as every caller is
+; norecurse. The Attributor then replaces that indirect call with a direct call
+; to @outer, and @outer already calls @dispatcher, so @dispatcher is recursive
+; and the attribute is false.
+;
+; TargetFrameLowering::isSafeForNoCSROpt requires norecurse before a function may
+; skip saving callee-saved registers, and AMDGPUResourceUsageAnalysis uses it to
+; decide a call graph has no recursion.
+
+ at g = external global i32
+
+define internal void @inner() norecurse {
+ store i32 1, ptr @g
+ ret void
+}
+
+define internal void @outer() norecurse {
+ call void @dispatcher(ptr @inner)
+ ret void
+}
+
+define internal void @dispatcher(ptr %fn) {
+ call void %fn()
+ ret void
+}
+
+define void @entry() norecurse {
+ call void @dispatcher(ptr @outer)
+ ret void
+}
+
+; The cycle @dispatcher -> @outer -> @dispatcher exists after specialization.
+; CHECK: define internal void @outer() #[[ATTR:[0-9]+]] {
+; CHECK: call void @dispatcher(
+; CHECK: define internal void @dispatcher(ptr {{.*}}) #[[ATTR]] {
+; CHECK: call void @outer()
+
+; So neither of them may keep norecurse.
+; CHECK: attributes #[[ATTR]] = { mustprogress nofree nosync nounwind willreturn memory(write) }
diff --git a/offload/test/offloading/nested_parallel_reduction.c b/offload/test/offloading/nested_parallel_reduction.c
new file mode 100644
index 0000000000000..267f17e4874d5
--- /dev/null
+++ b/offload/test/offloading/nested_parallel_reduction.c
@@ -0,0 +1,47 @@
+// A parallel region nested inside another one gives the wrong answer once the
+// optimizer is turned up. The same source is correct at -O1 and wrong at -O2
+// and -O3, with the same launch geometry and the same Generic-SPMD execution
+// mode in both cases, so the difference is in what the optimizer does to the
+// kernel rather than in how it is launched.
+//
+// Either parallel region on its own is correct at every optimization level; it
+// takes the two of them nested to reproduce. The result does not depend on
+// thread_limit.
+
+// RUN: %libomptarget-compileopt-generic
+// RUN: %libomptarget-run-generic | %fcheck-generic
+
+// REQUIRES: gpu
+
+
+#include <stdio.h>
+
+#define N 5
+
+int main(void) {
+ long aa = 0;
+ int ng = 6, cmom = 4, nxyz = 5;
+
+#pragma omp target teams distribute num_teams(nxyz) thread_limit(4) \
+ map(tofrom : aa)
+ for (int gid = 0; gid < nxyz; gid++) {
+#pragma omp parallel for collapse(2)
+ for (unsigned g = 0; g < ng; g++)
+ for (unsigned l = 0; l < cmom - 1; l++) {
+ int a = 0;
+ for (int ii = 0; ii < N + 2; ii++) {
+#pragma omp parallel for reduction(+ : a)
+ for (int i = 0; i < N; i++)
+ a += i;
+ }
+#pragma omp atomic
+ aa += a;
+ }
+ }
+
+ long expected = (long)ng * (cmom - 1) * nxyz * (N * (N - 1) / 2) * (N + 2);
+ printf("aa = %ld, expected %ld\n", aa, expected);
+ return aa != expected;
+}
+
+// CHECK: aa = 6300, expected 6300
``````````
</details>
https://github.com/llvm/llvm-project/pull/218637
More information about the llvm-commits
mailing list