[llvm] [Attributor] Drop norecurse when specializing an indirect call closes a cycle (PR #218637)
Larry Meadows via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 12:22:11 PDT 2026
https://github.com/lfmeadow updated https://github.com/llvm/llvm-project/pull/218637
>From 9d0fa7e77706d7bb780e45706ceae554ae5a3322 Mon Sep 17 00:00:00 2001
From: Larry Meadows <Lawrence.Meadows at amd.com>
Date: Fri, 21 Aug 2026 22:21:40 -0500
Subject: [PATCH 1/2] [Attributor] Drop norecurse when specializing an indirect
call closes a cycle
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.
That is not cosmetic. TargetFrameLowering::isSafeForNoCSROpt uses `norecurse` to
decide a function may skip callee-saved register spills entirely, and
AMDGPUResourceUsageAnalysis uses it to decide a call graph has no recursion. With
IPRA enabled, which AMDGPU does unconditionally, a value live across the newly
recursive call can be assigned a callee-saved register that is then never spilled,
and the callee clobbers it.
Record the functions whose indirect calls were specialized, and during cleanup
drop `norecurse` from every function in the call graph SCC of one of them when
that SCC has a cycle. Doing this once at the end rather than per call site also
covers cycles that no single specialization closes by itself. The attribute
cannot be dropped during manifestation because the abstract attribute holding it
may be manifested afterwards and put it back.
Cycles through a function pointer are not modelled, because CallGraph routes
indirect calls to a sink node. That matches what the attribute is used for here:
isSafeForNoCSROpt already requires that the function's address is not taken.
Two Attributor tests cover the attribute itself: norecurse_indirect_call.ll,
where specialization closes a cycle and norecurse survives on the recursive
function, and norecurse_stale_after_specialization.ll, where it is stale after
specialization without a direct self-call.
nested_parallel_reduction.c is what the miscompile costs at runtime. Its nested
reduction returns 0 instead of 6300 before this change, with no fault and no
diagnostic; the launch geometry is identical either way, so the only symptom is
the wrong answer.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
llvm/include/llvm/Transforms/IPO/Attributor.h | 8 ++
llvm/lib/Transforms/IPO/Attributor.cpp | 37 ++++++++++
.../Transforms/IPO/AttributorAttributes.cpp | 2 +
.../Attributor/norecurse_indirect_call.ll | 73 +++++++++++++++++++
.../norecurse_stale_after_specialization.ll | 44 +++++++++++
.../offloading/nested_parallel_reduction.c | 47 ++++++++++++
6 files changed, 211 insertions(+)
create mode 100644 llvm/test/Transforms/Attributor/norecurse_indirect_call.ll
create mode 100644 llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll
create mode 100644 offload/test/offloading/nested_parallel_reduction.c
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
>From 6f2d0ce28608ae88ebe18588a3deeeb38a9afa03 Mon Sep 17 00:00:00 2001
From: Larry Meadows <Lawrence.Meadows at amd.com>
Date: Tue, 25 Aug 2026 03:38:00 -0500
Subject: [PATCH 2/2] Trim the comments in the two new tests
Co-authored-by: Cursor <cursoragent at cursor.com>
---
.../norecurse_stale_after_specialization.ll | 10 ++--------
offload/test/offloading/nested_parallel_reduction.c | 12 ++----------
2 files changed, 4 insertions(+), 18 deletions(-)
diff --git a/llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll b/llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll
index c09c83322d58a..a504ff565370d 100644
--- a/llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll
+++ b/llvm/test/Transforms/Attributor/norecurse_stale_after_specialization.ll
@@ -1,16 +1,10 @@
; 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:
+; rpo-function-attrs is what marks @dispatcher norecurse here:
; 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.
+; norecurse.
@g = external global i32
diff --git a/offload/test/offloading/nested_parallel_reduction.c b/offload/test/offloading/nested_parallel_reduction.c
index 267f17e4874d5..862aeba057326 100644
--- a/offload/test/offloading/nested_parallel_reduction.c
+++ b/offload/test/offloading/nested_parallel_reduction.c
@@ -1,19 +1,11 @@
-// 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.
+// Both parallel regions have to be nested for this to reproduce, and only at
+// -O2 and above; either one alone is correct at every optimization level.
// RUN: %libomptarget-compileopt-generic
// RUN: %libomptarget-run-generic | %fcheck-generic
// REQUIRES: gpu
-
#include <stdio.h>
#define N 5
More information about the llvm-commits
mailing list