[clang] [llvm] [NVPTX] Emit bra.uni for proven uniform branches (PR #232269)
Mario Lezcano Casado via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 15 04:55:30 PDT 2026
https://github.com/lezcano updated https://github.com/llvm/llvm-project/pull/232269
>From bd9796493076ff646036cc82196a4cd7fb482291 Mon Sep 17 00:00:00 2001
From: lezcano <lezcano-93 at hotmail.com>
Date: Thu, 15 Oct 2026 11:58:29 +0200
Subject: [PATCH 1/4] [UniformityAnalysis] Propagate divergence from terminator
results
Track terminator definitions independently from control flow and
propagate divergence to their users, including definitions in divergent
cycles.
Treat callbr and invoke returns as possible divergence sources in NVPTX
and test uniform, opaque, and operand-dependent results.
---
llvm/include/llvm/ADT/GenericUniformityImpl.h | 52 ++++++-------
llvm/lib/Analysis/UniformityAnalysis.cpp | 2 -
.../lib/CodeGen/MachineUniformityAnalysis.cpp | 2 -
.../Target/NVPTX/NVPTXTargetTransformInfo.cpp | 2 +-
.../NVPTX/terminator-results.ll | 75 +++++++++++++++++++
5 files changed, 99 insertions(+), 34 deletions(-)
create mode 100644 llvm/test/Analysis/UniformityAnalysis/NVPTX/terminator-results.ll
diff --git a/llvm/include/llvm/ADT/GenericUniformityImpl.h b/llvm/include/llvm/ADT/GenericUniformityImpl.h
index c6f0bba051b609..c328734f42857a 100644
--- a/llvm/include/llvm/ADT/GenericUniformityImpl.h
+++ b/llvm/include/llvm/ADT/GenericUniformityImpl.h
@@ -473,6 +473,14 @@ template <typename ContextT> class GenericUniformityAnalysisImpl {
// Set of known-uniform values.
SmallPtrSet<const InstructionT *, 32> UniformOverrides;
+ bool isTargetUniform(const InstructionT &I) const {
+ // For custom uniformity candidates, check if the instruction can be
+ // proven uniform based on which operands are uniform/divergent.
+ // The candidate will be re-evaluated as operands become divergent.
+ return isAlwaysUniform(I) ||
+ (CustomUniformityCandidates.contains(&I) && isCustomUniform(I));
+ }
+
/// \brief Mark all nodes in \p JoinBlock as divergent and push them on
/// the worklist.
void taintAndPushAllDefs(const BlockT &JoinBlock);
@@ -821,27 +829,16 @@ auto llvm::GenericSyncDependenceAnalysis<ContextT>::getJoinBlocks(
template <typename ContextT>
void GenericUniformityAnalysisImpl<ContextT>::markDivergent(
const InstructionT &I) {
- if (isAlwaysUniform(I))
+ if (isTargetUniform(I))
return;
- // For custom uniformity candidates, check if the instruction can be
- // proven uniform based on which operands are uniform/divergent.
- // The candidate will be re-evaluated as operands become divergent.
- if (CustomUniformityCandidates.contains(&I)) {
- if (isCustomUniform(I))
- return;
- }
- bool Marked = false;
- if (I.isTerminator()) {
- Marked = DivergentTermBlocks.insert(I.getParent()).second;
- if (Marked) {
- LLVM_DEBUG(dbgs() << "marked divergent term block: "
- << Context.print(I.getParent()) << "\n");
- }
- } else {
- Marked = markDefsDivergent(I);
+ const bool MarkedTerm =
+ I.isTerminator() && DivergentTermBlocks.insert(I.getParent()).second;
+ if (MarkedTerm) {
+ LLVM_DEBUG(dbgs() << "marked divergent term block: "
+ << Context.print(I.getParent()) << "\n");
}
- if (Marked)
+ if (markDefsDivergent(I) || MarkedTerm)
Worklist.push_back(&I);
}
@@ -942,13 +939,11 @@ void GenericUniformityAnalysisImpl<ContextT>::taintAndPushAllDefs(
const BlockT &BB) {
LLVM_DEBUG(dbgs() << "taintAndPushAllDefs " << Context.print(&BB) << "\n");
for (const auto &I : instrs(BB)) {
- // Terminators do not produce values; they are divergent only if
- // the condition is divergent. That is handled when the divergent
- // condition is placed in the worklist.
- if (I.isTerminator())
- break;
-
- markDivergent(I);
+ // For terminators, taint only their definitions. Their control flow
+ // becomes divergent when a divergent condition is placed in
+ // the worklist.
+ if (!isTargetUniform(I) && markDefsDivergent(I))
+ pushUsers(I);
}
}
@@ -1165,13 +1160,12 @@ void GenericUniformityAnalysisImpl<ContextT>::compute() {
LLVM_DEBUG(dbgs() << "worklist pop: " << Context.print(I) << "\n");
- if (I->isTerminator()) {
+ if (I->isTerminator())
analyzeControlDivergence(*I);
- continue;
- }
// propagate value divergence to users
- assert(hasDivergentDefs(*I) && "Worklist invariant violated!");
+ assert((I->isTerminator() || hasDivergentDefs(*I)) &&
+ "Worklist invariant violated!");
pushUsers(*I);
}
}
diff --git a/llvm/lib/Analysis/UniformityAnalysis.cpp b/llvm/lib/Analysis/UniformityAnalysis.cpp
index c07119d7002857..af45299c21ec2a 100644
--- a/llvm/lib/Analysis/UniformityAnalysis.cpp
+++ b/llvm/lib/Analysis/UniformityAnalysis.cpp
@@ -44,8 +44,6 @@ template <>
void llvm::GenericUniformityAnalysisImpl<SSAContext>::pushUsers(
const Instruction &Instr) {
assert(!isAlwaysUniform(Instr));
- if (Instr.isTerminator())
- return;
pushUsers(cast<Value>(&Instr));
}
diff --git a/llvm/lib/CodeGen/MachineUniformityAnalysis.cpp b/llvm/lib/CodeGen/MachineUniformityAnalysis.cpp
index 21c5cd75ee799d..ecd44707aebd91 100644
--- a/llvm/lib/CodeGen/MachineUniformityAnalysis.cpp
+++ b/llvm/lib/CodeGen/MachineUniformityAnalysis.cpp
@@ -98,8 +98,6 @@ template <>
void llvm::GenericUniformityAnalysisImpl<MachineSSAContext>::pushUsers(
const MachineInstr &Instr) {
assert(!isAlwaysUniform(Instr));
- if (Instr.isTerminator())
- return;
for (const MachineOperand &Op : Instr.all_defs()) {
auto Reg = Op.getReg();
if (isDivergent(Reg))
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
index 960cd7540aa907..9071c2a3671d2b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
@@ -75,7 +75,7 @@ bool NVPTXTTIImpl::isSourceOfDivergence(const Value *V) const {
// Conservatively consider the return value of function calls as divergent.
// We could analyze callees with bodies more precisely using
// inter-procedural analysis.
- if (isa<CallInst>(I))
+ if (isa<CallBase>(I))
return true;
}
diff --git a/llvm/test/Analysis/UniformityAnalysis/NVPTX/terminator-results.ll b/llvm/test/Analysis/UniformityAnalysis/NVPTX/terminator-results.ll
new file mode 100644
index 00000000000000..e2a50917797b02
--- /dev/null
+++ b/llvm/test/Analysis/UniformityAnalysis/NVPTX/terminator-results.ll
@@ -0,0 +1,75 @@
+; RUN: opt %s -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+define ptx_kernel void @callbr_results() {
+; CHECK-LABEL: for function 'callbr_results'
+; CHECK: {{^ +}}%uniform.use =
+; CHECK: DIVERGENT: %source.use =
+; CHECK: DIVERGENT: %dependent.use =
+entry:
+ %uniform = callbr i32 asm sideeffect "", "=r,r,!i"(i32 0) #0
+ to label %source_block [label %exit]
+source_block:
+ %uniform.use = add i32 %uniform, 1
+ %source = callbr i32 asm sideeffect "", "=r,!i"()
+ to label %dependent_block [label %exit]
+dependent_block:
+ %source.use = add i32 %source, 1
+ %dependent = callbr i32 asm sideeffect "", "=r,r,!i"(i32 %source) #0
+ to label %use [label %exit]
+use:
+ %dependent.use = add i32 %dependent, 1
+ br label %exit
+exit:
+ ret void
+}
+
+define ptx_kernel void @invoke_results() personality ptr @personality {
+; CHECK-LABEL: for function 'invoke_results'
+; CHECK: {{^ +}}%uniform.use =
+; CHECK: DIVERGENT: %source.use =
+; CHECK: DIVERGENT: %dependent.use =
+entry:
+ %uniform = invoke i32 @dependent(i32 0)
+ to label %source_block unwind label %unwind
+source_block:
+ %uniform.use = add i32 %uniform, 1
+ %source = invoke i32 @source()
+ to label %dependent_block unwind label %unwind
+dependent_block:
+ %source.use = add i32 %source, 1
+ %dependent = invoke i32 @dependent(i32 %source)
+ to label %use unwind label %unwind
+use:
+ %dependent.use = add i32 %dependent, 1
+ ret void
+unwind:
+ %lp = landingpad { ptr, i32 } cleanup
+ resume { ptr, i32 } %lp
+}
+
+define ptx_kernel void @callbr_result_in_divergent_cycle(i32 %arg) {
+; CHECK-LABEL: for function 'callbr_result_in_divergent_cycle'
+; CHECK: CYCLES ASSUMED DIVERGENT:
+; CHECK: {{^ +}}%value = callbr
+; CHECK: DIVERGENT: %use =
+entry:
+ %lane = call i32 @llvm.nvvm.read.ptx.sreg.laneid()
+ %cond = icmp eq i32 %lane, 0
+ br i1 %cond, label %left, label %right
+left:
+ %value = callbr i32 asm sideeffect "", "=r,r,!i"(i32 %arg) #0
+ to label %right [label %exit]
+right:
+ br label %left
+exit:
+ %use = add i32 %value, 1
+ ret void
+}
+
+declare i32 @source()
+declare i32 @dependent(i32) #0
+declare i32 @personality(...)
+declare i32 @llvm.nvvm.read.ptx.sreg.laneid()
+attributes #0 = { nodivergencesource }
>From 7fe458d2491984bd1e84bd9f18a68a3966165117 Mon Sep 17 00:00:00 2001
From: lezcano <lezcano-93 at hotmail.com>
Date: Thu, 15 Oct 2026 11:58:30 +0200
Subject: [PATCH 2/4] [Intrinsics] Annotate intrinsics that preserve uniformity
Add IntrNoDivergenceSource and emit the existing nodivergencesource
attribute from TableGen. Annotate arithmetic and value-transforming
intrinsics through shared classes and property groups so uniformity
analysis can reuse their semantics.
---
clang/test/CodeGen/X86/math-builtins.c | 50 +++++-----
clang/test/CodeGen/builtin-sqrt.c | 2 +-
clang/test/CodeGen/libcalls.c | 2 +-
clang/test/CodeGen/math-libcalls.c | 25 ++---
.../cl20-device-side-enqueue-attributes.cl | 2 +-
llvm/include/llvm/IR/Intrinsics.td | 61 +++++++++---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 94 ++++++++-----------
.../upgrade-invariant-group-barrier.ll | 2 +-
.../AArch64/replace-with-veclib-armpl.ll | 4 +-
.../replace-with-veclib-libmvec-scalable.ll | 2 +-
.../AArch64/replace-with-veclib-libmvec.ll | 4 +-
.../replace-with-veclib-sleef-scalable.ll | 2 +-
.../AArch64/replace-with-veclib-sleef.ll | 2 +-
.../AMDGPU/amdgpu-simplify-libcall-rootn.ll | 4 +-
.../replace-with-veclib-sleef-scalable.ll | 2 +-
.../AMX/amx-low-intrinsics-no-amx-bitcast.ll | 2 +-
.../CodeGen/X86/AMX/amx-low-intrinsics.ll | 2 +-
llvm/test/Linker/drop-attribute.ll | 2 +-
llvm/test/Other/invariant.group.ll | 2 +-
.../TableGen/intrinsic-nodivergencesource.td | 25 +++++
.../AtomicExpand/AArch64/atomicrmw-fp.ll | 2 +-
.../test/Transforms/Attributor/nocapture-1.ll | 4 +-
llvm/test/Transforms/Attributor/nofree.ll | 6 +-
llvm/test/Transforms/Attributor/nosync.ll | 4 +-
llvm/test/Transforms/Attributor/willreturn.ll | 6 +-
.../EarlyCSE/replace-calls-def-attrs.ll | 2 +-
.../X86/expand-fp-convert-small.ll | 2 +-
.../X86/expand-int-convert-small.ll | 2 +-
.../X86/expand-large-fp-convert-si129tofp.ll | 2 +-
.../X86/expand-large-fp-convert-ui129tofp.ll | 2 +-
.../X86/expand-large-fp-optnone.ll | 2 +-
.../Transforms/ExpandIRInsts/X86/sdiv129.ll | 2 +-
.../Transforms/ExpandIRInsts/X86/srem129.ll | 2 +-
.../Transforms/ExpandIRInsts/X86/udiv129.ll | 2 +-
.../Transforms/ExpandIRInsts/X86/urem129.ll | 2 +-
.../Transforms/ExpandIRInsts/X86/vector.ll | 2 +-
.../Transforms/ExpandMemCmp/AArch64/memcmp.ll | 2 +-
.../Transforms/ExpandMemCmp/X86/memcmp-x32.ll | 2 +-
.../Transforms/ExpandMemCmp/X86/memcmp.ll | 4 +-
.../Transforms/InstCombine/icmp-select.ll | 2 +-
llvm/test/Transforms/InstCombine/mul.ll | 4 +-
llvm/test/Transforms/InstCombine/rem.ll | 2 +-
.../InstCombine/sdiv-canonicalize.ll | 2 +-
.../Transforms/InstCombine/select-and-or.ll | 4 +-
.../Transforms/InstCombine/select-select.ll | 2 +-
.../Transforms/InstCombine/select_meta.ll | 4 +-
.../simplify-demanded-fpclass-exp.ll | 2 +-
.../simplify-demanded-fpclass-fdiv.ll | 2 +-
.../InstCombine/sub-ashr-or-to-icmp-select.ll | 2 +-
llvm/test/Transforms/LoopIdiom/basic.ll | 2 +-
.../LoopIdiom/cyclic-redundancy-check.ll | 2 +-
.../RISCV/veclib-function-calls.ll | 2 +-
.../LowerTypeTests/aarch64-jumptable-dbg.ll | 4 +-
.../LowerTypeTests/aarch64-jumptable.ll | 2 +-
.../Transforms/LowerTypeTests/import-32bit.ll | 2 +-
llvm/test/Transforms/LowerTypeTests/import.ll | 4 +-
.../LowerTypeTests/x86-jumptable-dbg.ll | 6 +-
.../LowerTypeTests/x86-jumptable.ll | 4 +-
.../X86/switch-of-powers-of-two.ll | 4 +-
.../Transforms/SimplifyCFG/rangereduce.ll | 2 +-
.../remove-attributes-from-intrinsics.ll | 2 +-
.../TableGen/Basic/CodeGenIntrinsics.cpp | 2 +
llvm/utils/TableGen/Basic/CodeGenIntrinsics.h | 3 +
.../utils/TableGen/Basic/IntrinsicEmitter.cpp | 17 ++--
64 files changed, 238 insertions(+), 189 deletions(-)
create mode 100644 llvm/test/TableGen/intrinsic-nodivergencesource.td
diff --git a/clang/test/CodeGen/X86/math-builtins.c b/clang/test/CodeGen/X86/math-builtins.c
index c7cd9ffdd6966c..f82c28a7a032c6 100644
--- a/clang/test/CodeGen/X86/math-builtins.c
+++ b/clang/test/CodeGen/X86/math-builtins.c
@@ -140,14 +140,14 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
__builtin_frexp(f,i); __builtin_frexpf(f,i); __builtin_frexpl(f,i); __builtin_frexpf128(f,i);
-// NO__ERRNO: declare { double, i32 } @llvm.frexp.f64.i32(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare { float, i32 } @llvm.frexp.f32.i32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare { x86_fp80, i32 } @llvm.frexp.f80.i32(x86_fp80) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare { fp128, i32 } @llvm.frexp.f128.i32(fp128) [[READNONE_INTRINSIC2]]
-// HAS_ERRNO: declare { double, i32 } @llvm.frexp.f64.i32(double) [[READNONE_INTRINSIC2]]
-// HAS_ERRNO: declare { float, i32 } @llvm.frexp.f32.i32(float) [[READNONE_INTRINSIC2]]
-// HAS_ERRNO: declare { x86_fp80, i32 } @llvm.frexp.f80.i32(x86_fp80) [[READNONE_INTRINSIC2]]
-// HAS_ERRNO: declare { fp128, i32 } @llvm.frexp.f128.i32(fp128) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare { double, i32 } @llvm.frexp.f64.i32(double) [[READNONE_INTRINSIC3:#[0-9]+]]
+// NO__ERRNO: declare { float, i32 } @llvm.frexp.f32.i32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare { x86_fp80, i32 } @llvm.frexp.f80.i32(x86_fp80) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare { fp128, i32 } @llvm.frexp.f128.i32(fp128) [[READNONE_INTRINSIC3]]
+// HAS_ERRNO: declare { double, i32 } @llvm.frexp.f64.i32(double) [[READNONE_INTRINSIC3:#[0-9]+]]
+// HAS_ERRNO: declare { float, i32 } @llvm.frexp.f32.i32(float) [[READNONE_INTRINSIC3]]
+// HAS_ERRNO: declare { x86_fp80, i32 } @llvm.frexp.f80.i32(x86_fp80) [[READNONE_INTRINSIC3]]
+// HAS_ERRNO: declare { fp128, i32 } @llvm.frexp.f128.i32(fp128) [[READNONE_INTRINSIC3]]
__builtin_huge_val(); __builtin_huge_valf(); __builtin_huge_vall(); __builtin_huge_valf128();
@@ -509,10 +509,10 @@ __builtin_lgamma(f); __builtin_lgammaf(f); __builtin_lgammal(f); __builti
__builtin_llrint(f); __builtin_llrintf(f); __builtin_llrintl(f); __builtin_llrintf128(f);
-// NO__ERRNO: declare i64 @llvm.llrint.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llrint.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llrint.i64.f128(fp128) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.llrint.i64.f64(double) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llrint.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llrint.i64.f128(fp128) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @llrint(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llrintf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llrintl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -520,10 +520,10 @@ __builtin_llrint(f); __builtin_llrintf(f); __builtin_llrintl(f); __builti
__builtin_llround(f); __builtin_llroundf(f); __builtin_llroundl(f); __builtin_llroundf128(f);
-// NO__ERRNO: declare i64 @llvm.llround.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llround.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llround.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llround.i64.f128(fp128) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.llround.i64.f64(double) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llround.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llround.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llround.i64.f128(fp128) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @llround(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llroundf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llroundl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -586,10 +586,10 @@ __builtin_logb(f); __builtin_logbf(f); __builtin_logbl(f); __builtin_
__builtin_lrint(f); __builtin_lrintf(f); __builtin_lrintl(f); __builtin_lrintf128(f);
-// NO__ERRNO: declare i64 @llvm.lrint.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lrint.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lrint.i64.f128(fp128) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.lrint.i64.f64(double) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lrint.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lrint.i64.f128(fp128) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @lrint(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lrintf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lrintl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -597,10 +597,10 @@ __builtin_lrint(f); __builtin_lrintf(f); __builtin_lrintl(f); __builtin
__builtin_lround(f); __builtin_lroundf(f); __builtin_lroundl(f); __builtin_lroundf128(f);
-// NO__ERRNO: declare i64 @llvm.lround.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lround.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lround.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lround.i64.f128(fp128) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.lround.i64.f64(double) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lround.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lround.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lround.i64.f128(fp128) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @lround(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lroundf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lroundl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -803,12 +803,14 @@ __builtin_trunc(f); __builtin_truncf(f); __builtin_truncl(f); __builtin
// NO__ERRNO: attributes [[READNONE_INTRINSIC]] = { {{.*}}memory(none){{.*}} }
// NO__ERRNO: attributes [[READNONE_INTRINSIC2]] = { {{.*}}memory(none){{.*}} }
+// NO__ERRNO: attributes [[READNONE_INTRINSIC3]] = { {{.*}}memory(none){{.*}} }
// NO__ERRNO: attributes [[NOT_READNONE]] = { nounwind {{.*}} }
// NO__ERRNO: attributes [[PURE]] = { {{.*}}memory(read){{.*}} }
// NO__ERRNO: attributes [[READNONE]] = { {{.*}}memory(none){{.*}} }
// HAS_ERRNO: attributes [[NOT_READNONE]] = { nounwind {{.*}} }
// HAS_ERRNO: attributes [[READNONE_INTRINSIC2]] = { {{.*}}memory(none){{.*}} }
+// HAS_ERRNO: attributes [[READNONE_INTRINSIC3]] = { {{.*}}memory(none){{.*}} }
// HAS_ERRNO: attributes [[READNONE_INTRINSIC]] = { {{.*}}memory(none){{.*}} }
// HAS_ERRNO: attributes [[PURE]] = { {{.*}}memory(read){{.*}} }
// HAS_ERRNO: attributes [[READNONE]] = { {{.*}}memory(none){{.*}} }
diff --git a/clang/test/CodeGen/builtin-sqrt.c b/clang/test/CodeGen/builtin-sqrt.c
index 3ebf2ac91ccdfb..c0fdbc3670c8e8 100644
--- a/clang/test/CodeGen/builtin-sqrt.c
+++ b/clang/test/CodeGen/builtin-sqrt.c
@@ -11,5 +11,5 @@ float foo(float X) {
// HAS_ERRNO-NOT: attributes [[ATTR]] = {{{.*}} memory(none)
// NO_ERRNO: declare float @llvm.sqrt.f32(float) [[ATTR:#[0-9]+]]
-// NO_ERRNO: attributes [[ATTR]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+// NO_ERRNO: attributes [[ATTR]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
diff --git a/clang/test/CodeGen/libcalls.c b/clang/test/CodeGen/libcalls.c
index 923719f3ec8e43..541e67dde0e7a6 100644
--- a/clang/test/CodeGen/libcalls.c
+++ b/clang/test/CodeGen/libcalls.c
@@ -124,4 +124,4 @@ void test_builtins(double d, float f, long double ld) {
}
// CHECK-YES: attributes [[NUW]] = { nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+x87" }
-// CHECK-NO-DAG: attributes [[NUW_RNI]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+// CHECK-NO-DAG: attributes [[NUW_RNI]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
diff --git a/clang/test/CodeGen/math-libcalls.c b/clang/test/CodeGen/math-libcalls.c
index d4cd6f86b3c514..8cdeea78bba338 100644
--- a/clang/test/CodeGen/math-libcalls.c
+++ b/clang/test/CodeGen/math-libcalls.c
@@ -422,9 +422,9 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
llrint(f); llrintf(f); llrintl(f);
-// NO__ERRNO: declare i64 @llvm.llrint.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llrint.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.llrint.i64.f64(double) [[READNONE_INTRINSIC3:#[0-9]+]]
+// NO__ERRNO: declare i64 @llvm.llrint.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @llrint(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llrintf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llrintl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -434,9 +434,9 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
llround(f); llroundf(f); llroundl(f);
-// NO__ERRNO: declare i64 @llvm.llround.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llround.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.llround.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.llround.i64.f64(double) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llround.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.llround.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @llround(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llroundf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @llroundl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -506,9 +506,9 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
lrint(f); lrintf(f); lrintl(f);
-// NO__ERRNO: declare i64 @llvm.lrint.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lrint.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.lrint.i64.f64(double) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lrint.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lrint.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @lrint(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lrintf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lrintl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -518,9 +518,9 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
lround(f); lroundf(f); lroundl(f);
-// NO__ERRNO: declare i64 @llvm.lround.i64.f64(double) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lround.i64.f32(float) [[READNONE_INTRINSIC2]]
-// NO__ERRNO: declare i64 @llvm.lround.i64.f80(x86_fp80) [[READNONE_INTRINSIC2]]
+// NO__ERRNO: declare i64 @llvm.lround.i64.f64(double) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lround.i64.f32(float) [[READNONE_INTRINSIC3]]
+// NO__ERRNO: declare i64 @llvm.lround.i64.f80(x86_fp80) [[READNONE_INTRINSIC3]]
// HAS_ERRNO: declare i64 @lround(double noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lroundf(float noundef) [[NOT_READNONE]]
// HAS_ERRNO: declare i64 @lroundl(x86_fp80 noundef) [[NOT_READNONE]]
@@ -735,6 +735,7 @@ sincos(f, d, d); sincosf(f, fp, fp); sincosl(f, l, l);
// NO__ERRNO: attributes [[NOT_READNONE]] = { nounwind {{.*}} }
// NO__ERRNO: attributes [[READNONE]] = { {{.*}}memory(none){{.*}} }
// NO__ERRNO: attributes [[READONLY]] = { {{.*}}memory(read){{.*}} }
+// NO__ERRNO: attributes [[READNONE_INTRINSIC3]] = { {{.*}}memory(none){{.*}} }
// HAS_ERRNO: attributes [[NOT_READNONE]] = { nounwind {{.*}} }
// HAS_ERRNO: attributes [[READNONE_INTRINSIC2]] = { {{.*}}memory(none){{.*}} }
diff --git a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
index c93f623016660f..3555ac69950991 100644
--- a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
+++ b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
@@ -199,7 +199,7 @@ kernel void device_side_enqueue(global float *a, global float *b, int i) {
// SPIR32: attributes #[[ATTR0]] = { convergent noinline norecurse nounwind optnone denormal_fpenv(float: preservesign) "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// SPIR32: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
// SPIR32: attributes #[[ATTR2]] = { convergent noinline nounwind optnone denormal_fpenv(float: preservesign) "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
-// SPIR32: attributes #[[ATTR3:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+// SPIR32: attributes #[[ATTR3:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
// SPIR32: attributes #[[ATTR4]] = { convergent nounwind denormal_fpenv(float: preservesign) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
// SPIR32: attributes #[[ATTR5]] = { convergent nounwind "uniform-work-group-size" }
//.
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 4494bf1e2a2048..0dad5a6d8b085c 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -259,6 +259,10 @@ def IntrConvergent : IntrinsicProperty;
// This property indicates that the intrinsic is safe to speculate.
def IntrSpeculatable : IntrinsicProperty;
+// Uniform arguments produce a uniform result; the intrinsic cannot introduce
+// divergence. Parallels the nodivergencesource function attribute.
+def IntrNoDivergenceSource : IntrinsicProperty;
+
// This property can be used to override the 'has no other side effects'
// language of the IntrNoMem, IntrReadMem, IntrWriteMem, and IntrArgMemOnly
// intrinsic properties. By default, intrinsics are assumed to have side
@@ -1154,7 +1158,8 @@ def int_callbr_landingpad : Intrinsic<[llvm_any_ty], [LLVMMatchType<0>],
def int_structured_gep
: DefaultAttrsIntrinsic<[llvm_anyptr_ty],
[LLVMMatchType<0>, llvm_vararg_ty],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable,
+ IntrNoDivergenceSource]>;
def int_structured_alloca : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [],
[IntrInaccessibleMemOnly]>;
@@ -1227,6 +1232,7 @@ def int_experimental_memset_pattern
let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
in {
+ let append IntrProperties = [IntrNoDivergenceSource] in {
def int_fma : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>,
LLVMMatchType<0>]>;
@@ -1292,6 +1298,8 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
[llvm_anyfloat_ty], [LLVMMatchType<0>],
[IntrNoMem]>;
+ }
+
// If the value doesn't fit an unspecified value is returned, but this
// is not poison so we can still mark these as IntrNoCreateUndefOrPoison.
def int_lround : DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty]>;
@@ -1301,6 +1309,7 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
// TODO: int operand should be constrained to same number of elements as the
// result.
+ let append IntrProperties = [IntrNoDivergenceSource] in
def int_ldexp : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>,
llvm_anyint_ty]>;
@@ -1310,7 +1319,7 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
}
// TODO: Move all of these into the IntrNoCreateUndefOrPoison case above.
-let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoDivergenceSource] in {
// These functions do not read memory, but are sensitive to the
// rounding mode. LLVM purposely does not model changes to the FP
// environment so they can be treated as readnone.
@@ -1331,6 +1340,7 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
[llvm_anyfloat_ty]>;
}
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_minnum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
@@ -1355,6 +1365,7 @@ def int_maximumnum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
+}
// Internal interface for object size checking
def int_objectsize : DefaultAttrsIntrinsic<[llvm_any_scalar_int_ty],
@@ -1382,6 +1393,7 @@ let IntrProperties = [IntrInaccessibleMemOnly] in {
//===--------------- Floating Point Properties ----------------------------===//
//
+let append IntrProperties = [IntrNoDivergenceSource] in
def int_is_fpclass
: DefaultAttrsIntrinsic<[LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>],
[llvm_anyfloat_ty, llvm_i32_ty],
@@ -1620,17 +1632,18 @@ let IntrProperties = [IntrInaccessibleMemOnly, IntrStrictFP] in {
//===------------------------- Expect Intrinsics --------------------------===//
//
def int_expect : DefaultAttrsIntrinsic<[llvm_anyint_ty],
- [LLVMMatchType<0>, LLVMMatchType<0>], [IntrNoMem]>;
+ [LLVMMatchType<0>, LLVMMatchType<0>], [IntrNoMem, IntrNoDivergenceSource]>;
def int_expect_with_probability : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, llvm_double_ty],
- [IntrNoMem, ImmArg<ArgIndex<2>>]>;
+ [IntrNoMem, IntrNoDivergenceSource, ImmArg<ArgIndex<2>>]>;
//===-------------------- Bit Manipulation Intrinsics ---------------------===//
//
// None of these intrinsics accesses memory at all.
-let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison,
+ IntrNoDivergenceSource]
in {
def int_bswap: DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>]>;
def int_ctpop: DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>]>;
@@ -1651,7 +1664,8 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
[LLVMMatchType<0>, LLVMMatchType<0>]>;
}
-let IntrProperties = [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<1>>] in {
+let IntrProperties = [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<1>>,
+ IntrNoDivergenceSource] in {
def int_ctlz : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i1_ty]>;
def int_cttz : DefaultAttrsIntrinsic<[llvm_anyint_ty],
@@ -1732,12 +1746,12 @@ def int_ptr_annotation : DefaultAttrsIntrinsic<
[llvm_anyptr_ty],
[LLVMMatchType<0>, llvm_anyptr_ty, LLVMMatchType<1>, llvm_i32_ty,
LLVMMatchType<1>],
- [IntrInaccessibleMemOnly]>;
+ [IntrInaccessibleMemOnly, IntrNoDivergenceSource]>;
def int_annotation : DefaultAttrsIntrinsic<
[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_anyptr_ty, LLVMMatchType<1>, llvm_i32_ty],
- [IntrInaccessibleMemOnly]>;
+ [IntrInaccessibleMemOnly, IntrNoDivergenceSource]>;
// Annotates the current program point with metadata strings which are emitted
// as CodeView debug info records. This is expensive, as it disables inlining
@@ -1761,7 +1775,8 @@ def int_adjust_trampoline : DefaultAttrsIntrinsic<
//
// Expose the carry flag from add operations on two integrals.
-let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison,
+ IntrNoDivergenceSource]
in {
def int_sadd_with_overflow : DefaultAttrsIntrinsic<
[llvm_anyint_ty,
@@ -1792,6 +1807,7 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
}
//===------------------ Saturation Arithmetic Intrinsics ------------------===//
//
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_sadd_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem, IntrSpeculatable,
@@ -1880,6 +1896,7 @@ def int_ucmp : DefaultAttrsIntrinsic<
[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>],
[IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison,
Range<RetIndex, -1, 2>]>;
+}
//===------------------------- Memory Use Markers -------------------------===//
//
@@ -1916,6 +1933,7 @@ def int_invariant_end : DefaultAttrsIntrinsic<[],
def int_launder_invariant_group : DefaultAttrsIntrinsic<[llvm_anyptr_ty],
[LLVMMatchType<0>],
[IntrInaccessibleMemOnly,
+ IntrNoDivergenceSource,
IntrSpeculatable]>;
//===------------------------ Stackmap Intrinsics -------------------------===//
@@ -2136,7 +2154,8 @@ def int_pseudoprobe : DefaultAttrsIntrinsic<[],
ImmArg<ArgIndex<3>>]>;
// Saturating floating point to integer intrinsics
-let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison,
+ IntrNoDivergenceSource]
in {
def int_fptoui_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[llvm_anyfloat_ty]>;
@@ -2160,7 +2179,8 @@ def int_fake_use : DefaultAttrsIntrinsic<[], [llvm_vararg_ty],
// Intrinsic to mask out bits of a pointer.
// First argument must be pointer or vector of pointer. This is checked by the
// verifier.
-def int_ptrmask: PureIntrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_anyint_ty]>;
+def int_ptrmask: PureIntrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_anyint_ty],
+ [IntrNoDivergenceSource]>;
// Intrinsic to wrap a thread local variable.
def int_threadlocal_address : DefaultAttrsIntrinsic<[llvm_anyptr_ty],
@@ -2169,7 +2189,7 @@ def int_threadlocal_address : DefaultAttrsIntrinsic<[llvm_anyptr_ty],
IntrNoMem, IntrSpeculatable]>;
def int_stepvector : DefaultAttrsIntrinsic<[llvm_any_vector_int_ty],
- [], [IntrNoMem]>;
+ [], [IntrNoMem, IntrNoDivergenceSource]>;
def int_reloc_none : DefaultAttrsIntrinsic<[], [llvm_metadata_ty],
[IntrNoMem, IntrHasSideEffects]>;
@@ -2256,6 +2276,7 @@ def int_experimental_vector_histogram_umax : DefaultAttrsIntrinsic<[],
[ IntrArgMemOnly ]>;
// Experimental match
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_experimental_vector_match : DefaultAttrsIntrinsic<
[ LLVMScalarOrSameVectorWidth<0, llvm_i1_ty> ],
[ llvm_any_vector_int_ty,
@@ -2270,7 +2291,7 @@ def int_experimental_vector_extract_last_active:
LLVMVectorElementType<0>], [IntrNoMem, IntrSpeculatable]>;
// Operators
-let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+let append IntrProperties = [IntrNoMem, IntrSpeculatable] in {
def int_vp_merge : DefaultAttrsIntrinsic<[ llvm_anyvector_ty ],
[ LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>,
LLVMMatchType<0>,
@@ -2377,7 +2398,7 @@ def int_vp_urem : DefaultAttrsIntrinsic<[ llvm_any_vector_int_ty ],
LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>,
llvm_i32_ty], [IntrNoMem]>;
-let IntrProperties = [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<1>>] in
+let append IntrProperties = [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<1>>] in
def int_vp_cttz_elts : DefaultAttrsIntrinsic<[ llvm_any_scalar_int_ty ],
[ llvm_any_vector_int_ty,
llvm_i1_ty,
@@ -2430,6 +2451,7 @@ def int_experimental_vp_reverse:
LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>,
llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]>;
+}
//===-------------------------- Masked Intrinsics -------------------------===//
//
@@ -2470,6 +2492,7 @@ def int_masked_compressstore:
[IntrWriteMem, IntrArgMemOnly,
NoCapture<ArgIndex<1>>]>;
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_experimental_vector_compress:
DefaultAttrsIntrinsic<[llvm_anyvector_ty],
[LLVMMatchType<0>, LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>,
@@ -2503,6 +2526,7 @@ def int_masked_srem:
LLVMMatchType<0>,
LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>],
[IntrNoMem]>;
+}
// Speculatively load a value from memory.
def int_speculative_load:
@@ -2627,7 +2651,8 @@ def int_memset_element_unordered_atomic
//===------------------------ Reduction Intrinsics ------------------------===//
//
-let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison,
+ IntrNoDivergenceSource]
in {
def int_vector_reduce_fadd : DefaultAttrsIntrinsic<[LLVMVectorElementType<0>],
@@ -2674,6 +2699,7 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]
//===------ Matrix intrinsics ---------------------------------------------===//
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_matrix_transpose
: DefaultAttrsIntrinsic<[llvm_anyvector_ty],
[LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
@@ -2686,6 +2712,7 @@ def int_matrix_multiply
llvm_i32_ty],
[IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<2>>,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>]>;
+}
def int_matrix_column_major_load
: DefaultAttrsIntrinsic<[llvm_anyvector_ty],
@@ -2712,6 +2739,7 @@ def int_set_loop_iterations :
// Same as the above, but produces a value (the same as the input operand) to
// be fed into the loop.
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_start_loop_iterations : DefaultAttrsIntrinsic<
[llvm_anyint_ty],
[LLVMMatchType<0>],
@@ -2728,6 +2756,7 @@ def int_test_set_loop_iterations :
def int_test_start_loop_iterations :
DefaultAttrsIntrinsic<[llvm_anyint_ty, llvm_i1_ty], [LLVMMatchType<0>],
[IntrNoDuplicate]>;
+}
// Decrement loop counter by the given argument. Return false if the loop
// should exit.
@@ -2741,6 +2770,7 @@ def int_loop_decrement :
// executed, as usual. Any optimisations are allowed to treat it is a sub, and
// it's scevable, so it's the backends responsibility to handle cases where it
// may be optimised.
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_loop_decrement_reg :
DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>], [IntrNoDuplicate]>;
@@ -2833,6 +2863,7 @@ def int_vector_partial_reduce_add : DefaultAttrsIntrinsic<[LLVMMatchType<0>],
def int_vector_partial_reduce_fadd : DefaultAttrsIntrinsic<[LLVMMatchType<0>],
[llvm_any_vector_float_ty, llvm_any_vector_float_ty],
[IntrNoMem]>;
+}
//===----------------- Pointer Authentication Intrinsics ------------------===//
//
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 4a274fcc73b17b..1bc426ba4af783 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1173,24 +1173,6 @@ class NVVM_MMA_SP_BLOCK_SCALE_SUPPORTED<list<WMMA_REGS> frags, string kind,
);
}
-
-class SHFL_INFO<bit sync, string mode, string type, bit return_pred> {
- string Suffix = !if(sync, "sync_", "")
- # mode # "_"
- # type
- # !if(return_pred, "p", "");
-
- string Name = "int_nvvm_shfl_" # Suffix;
- bit withGccBuiltin = !not(return_pred);
- LLVMType OpType = !cond(
- !eq(type,"i32"): llvm_i32_ty,
- !eq(type,"f32"): llvm_float_ty);
- list<LLVMType> RetTy = !if(return_pred, [OpType, llvm_i1_ty], [OpType]);
- list<LLVMType> ArgsTy = !if(sync,
- [llvm_i32_ty, OpType, llvm_i32_ty, llvm_i32_ty],
- [OpType, llvm_i32_ty, llvm_i32_ty]);
-}
-
class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num,
string ElemType = "i32"> {
int shift = !cond(!eq(Shape, "16x128b"): 1,
@@ -1622,6 +1604,7 @@ let TargetPrefix = "nvvm" in {
//
// PRMT - permute
//
+ let append IntrProperties = [IntrNoDivergenceSource] in {
def int_nvvm_prmt : NVVMBuiltin,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
@@ -1635,6 +1618,7 @@ let TargetPrefix = "nvvm" in {
foreach mode = ["rc8", "ecl", "ecr", "rc16"] in
def int_nvvm_prmt_ # mode :
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
+ }
//
// Nanosleep
@@ -1654,7 +1638,7 @@ let TargetPrefix = "nvvm" in {
// Min Max
//
let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
- IntrNoCreateUndefOrPoison] in {
+ IntrNoCreateUndefOrPoison, IntrNoDivergenceSource] in {
foreach operation = ["min", "max"] in {
def int_nvvm_f # operation # _d : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
@@ -1680,19 +1664,17 @@ let TargetPrefix = "nvvm" in {
} // nan
} // variant
} // operation
- }
//
// Multiplication
//
- let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
- IntrNoCreateUndefOrPoison] in {
foreach sign = ["", "u"] in {
def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
}
}
+ let append IntrProperties = [IntrNoDivergenceSource] in {
foreach ftz = ["", "_ftz"] in
foreach sat = ["", "_sat"] in
def int_nvvm_fmul # ftz # sat :
@@ -1708,7 +1690,7 @@ let TargetPrefix = "nvvm" in {
//
// Div
//
- let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
+ let append IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
foreach ftz = ["", "_ftz"] in {
def int_nvvm_div_approx # ftz # _f : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;
@@ -1800,7 +1782,7 @@ let TargetPrefix = "nvvm" in {
//
// Exp2 Log2
//
- let IntrProperties = [IntrNoMem] in {
+ let IntrProperties = [IntrNoMem, IntrNoDivergenceSource] in {
foreach ftz = ["", "_ftz"] in
def int_nvvm_ex2_approx # ftz :
PureIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
@@ -1868,7 +1850,7 @@ let TargetPrefix = "nvvm" in {
//
// Rcp
//
- let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
+ let append IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
foreach rnd = ["rn", "rz", "rm", "rp"] in {
foreach ftz = ["", "_ftz"] in
def int_nvvm_rcp_ # rnd # ftz # _f : NVVMBuiltin,
@@ -1882,12 +1864,10 @@ let TargetPrefix = "nvvm" in {
DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
def int_nvvm_rcp_approx_ftz_d : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
- }
//
// Sqrt
//
- let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
foreach rnd = ["rn", "rz", "rm", "rp"] in {
foreach ftz = ["", "_ftz"] in
def int_nvvm_sqrt_ # rnd # ftz # _f : NVVMBuiltin,
@@ -1903,12 +1883,10 @@ let TargetPrefix = "nvvm" in {
foreach ftz = ["", "_ftz"] in
def int_nvvm_sqrt_approx # ftz # _f : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
- }
//
// Rsqrt
//
- let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
foreach ftz = ["", "_ftz"] in {
def int_nvvm_rsqrt_approx # ftz # _f : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
@@ -1921,7 +1899,7 @@ let TargetPrefix = "nvvm" in {
// FAdd
//
- let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
+ let append IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
IntrNoCreateUndefOrPoison, ImmArg<ArgIndex<2>>,
Range<ArgIndex<2>, 0, 4>,
ArgInfo<ArgIndex<2>,
@@ -2312,6 +2290,7 @@ let TargetPrefix = "nvvm" in {
def int_nvvm_bf16x2_to_ue5m3x2_rp # satfinite
: PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
}
+ }
//
// Bar.Sync
@@ -2493,7 +2472,7 @@ foreach is_shared = [true, false] in {
}
def int_nvvm_mbarrier_pending_count : NVVMBuiltin,
- NVVMPureIntrinsic<[llvm_i32_ty], [llvm_i64_ty]>;
+ NVVMPureIntrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoDivergenceSource]>;
def int_nvvm_mbarrier_check_layout :
DefaultAttrsIntrinsicFlags<[llvm_i1_ty],
@@ -2575,7 +2554,8 @@ foreach op = ["test_wait", "try_wait"] in {
// Generated within nvvm. Use for ldu on sm_20 or later. Second arg is the
// pointer's alignment.
-let IntrProperties = [IntrReadMem, IntrArgMemOnly, IntrNoCallback, IntrWillReturn, NoCapture<ArgIndex<0>>] in {
+let IntrProperties = [IntrReadMem, IntrArgMemOnly, IntrNoCallback, IntrWillReturn,
+ NoCapture<ArgIndex<0>>, IntrNoDivergenceSource] in {
def int_nvvm_ldu_global_i : Intrinsic<[llvm_anyint_ty], [llvm_anyptr_ty, llvm_i32_ty]>;
def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty], [llvm_anyptr_ty, llvm_i32_ty]>;
def int_nvvm_ldu_global_p : Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty, llvm_i32_ty]>;
@@ -2583,7 +2563,7 @@ let IntrProperties = [IntrReadMem, IntrArgMemOnly, IntrNoCallback, IntrWillRetur
// Move intrinsics, used in nvvm internally
-let IntrProperties = [IntrNoMem] in {
+let IntrProperties = [IntrNoMem, IntrNoDivergenceSource] in {
def int_nvvm_move_i16 : DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty]>;
def int_nvvm_move_i32 : DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty]>;
def int_nvvm_move_i64 : DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty]>;
@@ -2596,18 +2576,20 @@ let IntrProperties = [IntrNoMem] in {
def int_nvvm_texsurf_handle
: NVVMPureIntrinsic<[llvm_i64_ty], [llvm_metadata_ty, llvm_anyptr_ty]>;
def int_nvvm_texsurf_handle_internal
- : NVVMPureIntrinsic<[llvm_i64_ty], [llvm_anyptr_ty]>;
+ : NVVMPureIntrinsic<[llvm_i64_ty], [llvm_anyptr_ty], [IntrNoDivergenceSource]>;
/// Error / Warn
def int_nvvm_compiler_error : Intrinsic<[], [llvm_anyptr_ty]>;
def int_nvvm_compiler_warn : Intrinsic<[], [llvm_anyptr_ty]>;
+let append IntrProperties = [IntrNoDivergenceSource] in {
def int_nvvm_reflect : NVVMBuiltin, NVVMPureIntrinsic<[llvm_i32_ty], [llvm_ptr_ty]>;
// isspacep.{const, global, local, shared}
foreach space = ["const", "global", "local", "shared", "shared_cluster"] in
def int_nvvm_isspacep_ # space : NVVMBuiltin,
NVVMPureIntrinsic<[llvm_i1_ty], [llvm_ptr_ty], [NoCapture<ArgIndex<0>>]>;
+}
//
// Texture Fetch
@@ -2702,6 +2684,7 @@ foreach clamp = ["clamp", "trap", "zero"] in {
} // vec
} // clamp
+let append IntrProperties = [IntrNoDivergenceSource] in {
//===- Texture Query ------------------------------------------------------===//
foreach query = ["channel_order", "channel_data_type", "width", "height",
@@ -2721,6 +2704,7 @@ foreach query = ["channel_order", "channel_data_type", "width", "height",
foreach type = ["sampler", "surface", "texture"] in
def int_nvvm_istypep_ # type : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem]>;
+}
//===- Surface Stores -----------------------------------------------------===//
@@ -2899,32 +2883,31 @@ def int_nvvm_read_ptx_sreg_dynamic_smem_size :
//
// Sync shfl variants have synchronization side effects modeled as
// inaccessible memory accesses to prevent DCE and reordering.
-foreach is_sync = [false, true] in {
- defvar CommonIntrinsicProps = [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback];
- defvar IntrinsicProps = !if(is_sync, CommonIntrinsicProps,
- !listconcat(CommonIntrinsicProps, [IntrReadMem, IntrWillReturn]));
- let IntrProperties = IntrinsicProps in {
- foreach mode = ["up", "down", "bfly", "idx"] in {
- foreach type = ["i32", "f32"] in {
- foreach return_pred = [false, true] in {
- defvar i = SHFL_INFO<is_sync, mode, type, return_pred>;
- if i.withGccBuiltin then
- def i.Name : NVVMBuiltin, Intrinsic<i.RetTy, i.ArgsTy>;
- else
- def i.Name : Intrinsic<i.RetTy, i.ArgsTy>;
- }
- }
- }
- }
+foreach is_sync = [false, true] in
+foreach return_pred = [false, true] in
+let IntrProperties = [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback] #
+ !if(is_sync, []<IntrinsicProperty>, [IntrReadMem, IntrWillReturn]) #
+ !if(return_pred, []<IntrinsicProperty>, [IntrNoDivergenceSource]) in
+foreach mode = ["up", "down", "bfly", "idx"] in
+foreach type = ["i32", "f32"] in {
+ defvar suffix = !if(is_sync, "sync_", "") # mode # "_" # type #
+ !if(return_pred, "p", "");
+ defvar op_type = !if(!eq(type, "i32"), llvm_i32_ty, llvm_float_ty);
+ def int_nvvm_shfl_ # suffix : Intrinsic<
+ [op_type] # !if(return_pred, [llvm_i1_ty], []),
+ !if(is_sync, [llvm_i32_ty], []) # [op_type, llvm_i32_ty, llvm_i32_ty]>,
+ ClangBuiltin<!if(return_pred, "", "__nvvm_shfl_" # suffix)>;
}
//
// VOTE
//
let IntrProperties = [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback] in {
+ let append IntrProperties = [IntrNoDivergenceSource] in {
def int_nvvm_vote_all : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i1_ty]>;
def int_nvvm_vote_any : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i1_ty]>;
def int_nvvm_vote_uni : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i1_ty]>;
+ }
def int_nvvm_vote_ballot : NVVMBuiltin, Intrinsic<[llvm_i32_ty], [llvm_i1_ty]>;
}
@@ -3233,15 +3216,14 @@ def int_nvvm_movmatrix_sync_aligned_m8n8_trans_b16
IntrNoCallback, IntrNoFree]>;
// MAPA
-let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>] in {
+let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>,
+ IntrNoDivergenceSource] in {
def int_nvvm_mapa
: DefaultAttrsIntrinsic<[llvm_ptr_ty], [llvm_ptr_ty, llvm_i32_ty]>;
def int_nvvm_mapa_shared_cluster
: DefaultAttrsIntrinsic<[llvm_shared_cluster_ptr_ty], [llvm_shared_ptr_ty, llvm_i32_ty]>;
-}
// GETCTARANK
-let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>] in {
def int_nvvm_getctarank
: DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_ptr_ty]>;
def int_nvvm_getctarank_shared_cluster
@@ -3958,12 +3940,12 @@ def int_nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared
// clusterlaunchcontrol.query_cancel.is_canceled
def int_nvvm_clusterlaunchcontrol_query_cancel_is_canceled
- : NVVMPureIntrinsic<[llvm_i1_ty], [llvm_i128_ty], [],
+ : NVVMPureIntrinsic<[llvm_i1_ty], [llvm_i128_ty], [IntrNoDivergenceSource],
"llvm.nvvm.clusterlaunchcontrol.query_cancel.is_canceled">;
foreach dim = ["x", "y", "z"] in
def int_nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_ # dim
- : NVVMPureIntrinsic<[llvm_i32_ty], [llvm_i128_ty], [],
+ : NVVMPureIntrinsic<[llvm_i32_ty], [llvm_i128_ty], [IntrNoDivergenceSource],
"llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid." # dim>;
//
diff --git a/llvm/test/Bitcode/upgrade-invariant-group-barrier.ll b/llvm/test/Bitcode/upgrade-invariant-group-barrier.ll
index f03c3bd565f997..02527389e7ebdf 100644
--- a/llvm/test/Bitcode/upgrade-invariant-group-barrier.ll
+++ b/llvm/test/Bitcode/upgrade-invariant-group-barrier.ll
@@ -13,7 +13,7 @@ define void @test(ptr %p1, ptr %p16) {
ret void
}
-; CHECK: Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite)
+; CHECK: Function Attrs: nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite)
; CHECK: declare ptr @llvm.launder.invariant.group.p0(ptr)
declare ptr @llvm.invariant.group.barrier(ptr)
declare ptr @llvm.invariant.group.barrier.p0(ptr)
diff --git a/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll b/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll
index caae6ea0493d9c..55671b1c808aba 100644
--- a/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll
+++ b/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll
@@ -774,7 +774,7 @@ define <vscale x 4 x float> @llvm_tanh_vscale_f32(<vscale x 4 x float> %in) #0 {
attributes #0 = { "target-features"="+sve" }
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK: attributes #[[ATTR1]] = { "target-features"="+sve" }
-; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
diff --git a/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec-scalable.ll b/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec-scalable.ll
index 1b541d1330aaee..f42e03759fe40b 100644
--- a/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec-scalable.ll
+++ b/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec-scalable.ll
@@ -575,5 +575,5 @@ declare <vscale x 2 x double> @llvm.trunc.nxv2f64(<vscale x 2 x double>)
declare <vscale x 4 x float> @llvm.trunc.nxv4f32(<vscale x 4 x float>)
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { "target-features"="+sve" }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+sve" }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+sve" }
;.
diff --git a/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec.ll b/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec.ll
index bd2419a8a7e587..e926929376e978 100644
--- a/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec.ll
+++ b/llvm/test/CodeGen/AArch64/replace-with-veclib-libmvec.ll
@@ -573,6 +573,6 @@ declare <4 x float> @llvm.tan.v4f32(<4 x float>)
declare <2 x double> @llvm.trunc.v2f64(<2 x double>)
declare <4 x float> @llvm.trunc.v4f32(<4 x float>)
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
diff --git a/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef-scalable.ll b/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef-scalable.ll
index 1d429ece6f8101..f4113a6d2f4bcc 100644
--- a/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef-scalable.ll
+++ b/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef-scalable.ll
@@ -575,5 +575,5 @@ declare <vscale x 2 x double> @llvm.trunc.nxv2f64(<vscale x 2 x double>)
declare <vscale x 4 x float> @llvm.trunc.nxv4f32(<vscale x 4 x float>)
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { "target-features"="+sve" }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+sve" }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+sve" }
;.
diff --git a/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef.ll b/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef.ll
index ebe5e78667680d..bd3b3b73d26b21 100644
--- a/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef.ll
+++ b/llvm/test/CodeGen/AArch64/replace-with-veclib-sleef.ll
@@ -573,5 +573,5 @@ declare <4 x float> @llvm.tan.v4f32(<4 x float>)
declare <2 x double> @llvm.trunc.v2f64(<2 x double>)
declare <4 x float> @llvm.trunc.v4f32(<4 x float>)
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll
index 805e510bde37b0..a653ce34823297 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll
@@ -1979,12 +1979,12 @@ attributes #2 = { noinline }
;.
; PRELINK: attributes #[[ATTR0]] = { strictfp }
; PRELINK: attributes #[[ATTR1:[0-9]+]] = { nounwind memory(read) }
-; PRELINK: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; PRELINK: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; PRELINK: attributes #[[ATTR3]] = { noinline }
; PRELINK: attributes #[[ATTR4]] = { nobuiltin }
;.
; NOPRELINK: attributes #[[ATTR0]] = { strictfp }
-; NOPRELINK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; NOPRELINK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; NOPRELINK: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nosync nounwind strictfp willreturn memory(inaccessiblemem: readwrite) }
; NOPRELINK: attributes #[[ATTR3]] = { noinline }
; NOPRELINK: attributes #[[ATTR4]] = { nobuiltin }
diff --git a/llvm/test/CodeGen/RISCV/replace-with-veclib-sleef-scalable.ll b/llvm/test/CodeGen/RISCV/replace-with-veclib-sleef-scalable.ll
index 649d91c2d747b0..f8b5ff5c8dfcf3 100644
--- a/llvm/test/CodeGen/RISCV/replace-with-veclib-sleef-scalable.ll
+++ b/llvm/test/CodeGen/RISCV/replace-with-veclib-sleef-scalable.ll
@@ -446,5 +446,5 @@ define <vscale x 4 x float> @frem_f32(<vscale x 4 x float> %in) {
;.
; CHECK: attributes #[[ATTR0]] = { "target-features"="+v" }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+v" }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+v" }
;.
diff --git a/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics-no-amx-bitcast.ll b/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics-no-amx-bitcast.ll
index a3725ebeb93942..7fcc29bff4b247 100644
--- a/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics-no-amx-bitcast.ll
+++ b/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics-no-amx-bitcast.ll
@@ -203,7 +203,7 @@ attributes #0 = { noinline nounwind optnone }
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { noinline nounwind optnone }
; CHECK: attributes #[[ATTR1:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"branch_weights", i32 4, i32 1}
;.
diff --git a/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics.ll b/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics.ll
index 0e140f2b453de4..d3f263e2365359 100644
--- a/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics.ll
+++ b/llvm/test/CodeGen/X86/AMX/amx-low-intrinsics.ll
@@ -545,7 +545,7 @@ attributes #0 = { noinline nounwind optnone }
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { noinline nounwind optnone }
; CHECK: attributes #[[ATTR1:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"x86-lower-amx-intrinsics"}
diff --git a/llvm/test/Linker/drop-attribute.ll b/llvm/test/Linker/drop-attribute.ll
index 3d4c13c2ffc752..e3f8f07f3a6491 100644
--- a/llvm/test/Linker/drop-attribute.ll
+++ b/llvm/test/Linker/drop-attribute.ll
@@ -39,7 +39,7 @@ define void @test_nocallback_definition() nocallback {
declare void @test_nocallback_call_site()
; Test that checks that nocallback attribute on an intrinsic is NOT dropped.
-; CHECK: ; Function Attrs: nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn
+; CHECK: ; Function Attrs: nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn
; CHECK-NEXT: declare float @llvm.sqrt.f32(float) #0
declare float @llvm.sqrt.f32(float) nocallback
diff --git a/llvm/test/Other/invariant.group.ll b/llvm/test/Other/invariant.group.ll
index c8c29db4ba0748..9c3011114be98e 100644
--- a/llvm/test/Other/invariant.group.ll
+++ b/llvm/test/Other/invariant.group.ll
@@ -86,7 +86,7 @@ declare void @use(ptr readonly)
declare void @useBool(i1)
declare void @clobber(ptr)
-; CHECK: Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite){{$}}
+; CHECK: Function Attrs: nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite){{$}}
; CHECK-NEXT: declare ptr @llvm.launder.invariant.group.p0(ptr)
declare ptr @llvm.launder.invariant.group.p0(ptr)
diff --git a/llvm/test/TableGen/intrinsic-nodivergencesource.td b/llvm/test/TableGen/intrinsic-nodivergencesource.td
new file mode 100644
index 00000000000000..e48ffb1742c0cb
--- /dev/null
+++ b/llvm/test/TableGen/intrinsic-nodivergencesource.td
@@ -0,0 +1,25 @@
+// RUN: llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DTEST_INTRINSICS_SUPPRESS_DEFS %s | FileCheck %s
+
+include "llvm/IR/Intrinsics.td"
+
+def int_a_none : Intrinsic<[llvm_i32_ty], [], [Throws]>;
+def int_b_uniform : Intrinsic<[llvm_i32_ty], [], [Throws, IntrNoDivergenceSource]>;
+def int_c_uniform : Intrinsic<[llvm_i32_ty], [], [Throws, IntrNoDivergenceSource]>;
+def int_d_nounwind : Intrinsic<[llvm_i32_ty]>;
+def int_e_nounwind_uniform : Intrinsic<[llvm_i32_ty], [], [IntrNoDivergenceSource]>;
+
+// The property is sufficient on its own and participates in attribute grouping.
+// CHECK: static AttributeSet getIntrinsicFnAttributeSet(
+// CHECK: case 0: // llvm.b.uniform
+// CHECK-NEXT: return AttributeSet::get(C, {
+// CHECK-NEXT: Attribute::get(C, Attribute::NoDivergenceSource),
+// CHECK-NEXT: });
+// CHECK: case 1: // llvm.d.nounwind
+// CHECK: case 2: // llvm.e.nounwind.uniform
+// CHECK: static constexpr uint8_t IntrinsicsToAttributesMap[] = {
+// CHECK-NEXT: 3 << 0 | 0, // llvm.a.none
+// CHECK-NEXT: 0 << 0 | 0, // llvm.b.uniform
+// CHECK-NEXT: 0 << 0 | 0, // llvm.c.uniform
+// CHECK-NEXT: 1 << 0 | 0, // llvm.d.nounwind
+// CHECK-NEXT: 2 << 0 | 0, // llvm.e.nounwind.uniform
+// CHECK-NEXT: }; // IntrinsicsToAttributesMap
diff --git a/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll b/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll
index 66f1a9b5439f2b..fca9a0b90b9d17 100644
--- a/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll
+++ b/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll
@@ -549,7 +549,7 @@ define <2 x double> @atomicrmw_fminimumnum_2_x_double(ptr %ptr, <2 x double> %va
!0 = !{!"function_entry_count", i64 1000}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"atomic-expand"}
diff --git a/llvm/test/Transforms/Attributor/nocapture-1.ll b/llvm/test/Transforms/Attributor/nocapture-1.ll
index dee5cd6e85d023..3020b1fdec61a1 100644
--- a/llvm/test/Transforms/Attributor/nocapture-1.ll
+++ b/llvm/test/Transforms/Attributor/nocapture-1.ll
@@ -817,7 +817,7 @@ declare ptr @llvm.launder.invariant.group.p0(ptr)
; TUNIT: attributes #[[ATTR12]] = { nosync nounwind memory(read) }
; TUNIT: attributes #[[ATTR13:[0-9]+]] = { nounwind willreturn memory(read) }
; TUNIT: attributes #[[ATTR14]] = { mustprogress nosync nounwind willreturn }
-; TUNIT: attributes #[[ATTR15:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite) }
+; TUNIT: attributes #[[ATTR15:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite) }
; TUNIT: attributes #[[ATTR16]] = { nofree nosync nounwind willreturn memory(write) }
; TUNIT: attributes #[[ATTR17]] = { nofree nosync nounwind willreturn memory(none) }
; TUNIT: attributes #[[ATTR18]] = { nofree nosync nounwind willreturn }
@@ -842,7 +842,7 @@ declare ptr @llvm.launder.invariant.group.p0(ptr)
; CGSCC: attributes #[[ATTR15]] = { nosync nounwind memory(read) }
; CGSCC: attributes #[[ATTR16:[0-9]+]] = { nounwind willreturn memory(read) }
; CGSCC: attributes #[[ATTR17]] = { mustprogress nosync nounwind willreturn }
-; CGSCC: attributes #[[ATTR18:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite) }
+; CGSCC: attributes #[[ATTR18:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(inaccessiblemem: readwrite) }
; CGSCC: attributes #[[ATTR19]] = { nofree nounwind willreturn memory(write) }
; CGSCC: attributes #[[ATTR20]] = { nofree nosync willreturn }
; CGSCC: attributes #[[ATTR21]] = { nofree nounwind willreturn }
diff --git a/llvm/test/Transforms/Attributor/nofree.ll b/llvm/test/Transforms/Attributor/nofree.ll
index 6937238dc8c3c4..32a26076c0d762 100644
--- a/llvm/test/Transforms/Attributor/nofree.ll
+++ b/llvm/test/Transforms/Attributor/nofree.ll
@@ -238,7 +238,7 @@ define void @call_both() #0 {
; TEST 10 (positive case)
; Call intrinsic function
-; CHECK: Function Attrs: nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)
+; CHECK: Function Attrs: nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none)
declare float @llvm.floor.f32(float)
define void @call_floor(float %a) #0 {
@@ -410,7 +410,7 @@ attributes #2 = { nobuiltin nounwind }
; TUNIT: attributes #[[ATTR3]] = { mustprogress nofree noinline norecurse nosync nounwind willreturn memory(none) uwtable }
; TUNIT: attributes #[[ATTR4]] = { mustprogress nofree noinline nosync nounwind willreturn memory(none) uwtable }
; TUNIT: attributes #[[ATTR5:[0-9]+]] = { nofree noinline nounwind memory(none) uwtable }
-; TUNIT: attributes #[[ATTR6:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; TUNIT: attributes #[[ATTR6:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; TUNIT: attributes #[[ATTR7]] = { nofree nounwind }
; TUNIT: attributes #[[ATTR8]] = { nobuiltin nofree nounwind }
; TUNIT: attributes #[[ATTR9]] = { nosync memory(none) }
@@ -426,7 +426,7 @@ attributes #2 = { nobuiltin nounwind }
; CGSCC: attributes #[[ATTR3]] = { mustprogress nofree noinline norecurse nosync nounwind willreturn memory(none) uwtable }
; CGSCC: attributes #[[ATTR4:[0-9]+]] = { nofree noinline nounwind memory(none) uwtable }
; CGSCC: attributes #[[ATTR5]] = { mustprogress nofree noinline nosync nounwind willreturn memory(none) uwtable }
-; CGSCC: attributes #[[ATTR6:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CGSCC: attributes #[[ATTR6:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CGSCC: attributes #[[ATTR7]] = { nofree nounwind }
; CGSCC: attributes #[[ATTR8]] = { nobuiltin nofree nounwind }
; CGSCC: attributes #[[ATTR9]] = { nosync memory(none) }
diff --git a/llvm/test/Transforms/Attributor/nosync.ll b/llvm/test/Transforms/Attributor/nosync.ll
index bce49fcca913e6..f7af8fa329cb3f 100644
--- a/llvm/test/Transforms/Attributor/nosync.ll
+++ b/llvm/test/Transforms/Attributor/nosync.ll
@@ -487,7 +487,7 @@ define void @nosync_convergent_callee_test() {
; TUNIT: attributes #[[ATTR16:[0-9]+]] = { convergent memory(none) }
; TUNIT: attributes #[[ATTR17]] = { memory(none) }
; TUNIT: attributes #[[ATTR18]] = { nounwind }
-; TUNIT: attributes #[[ATTR19:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; TUNIT: attributes #[[ATTR19:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; TUNIT: attributes #[[ATTR20]] = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) }
; TUNIT: attributes #[[ATTR21]] = { nosync memory(none) }
; TUNIT: attributes #[[ATTR22]] = { nofree nosync nounwind }
@@ -515,7 +515,7 @@ define void @nosync_convergent_callee_test() {
; CGSCC: attributes #[[ATTR16:[0-9]+]] = { convergent memory(none) }
; CGSCC: attributes #[[ATTR17]] = { memory(none) }
; CGSCC: attributes #[[ATTR18]] = { nounwind }
-; CGSCC: attributes #[[ATTR19:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CGSCC: attributes #[[ATTR19:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CGSCC: attributes #[[ATTR20]] = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) }
; CGSCC: attributes #[[ATTR21]] = { nosync memory(none) }
; CGSCC: attributes #[[ATTR22]] = { nofree nosync nounwind }
diff --git a/llvm/test/Transforms/Attributor/willreturn.ll b/llvm/test/Transforms/Attributor/willreturn.ll
index d65480b05759a0..1dcfb504c2c381 100644
--- a/llvm/test/Transforms/Attributor/willreturn.ll
+++ b/llvm/test/Transforms/Attributor/willreturn.ll
@@ -276,7 +276,7 @@ define void @conditional_exit(i32 %0, ptr nocapture readonly %1) local_unnamed_a
; TEST 6 (positive case)
; Call intrinsic function
-; CHECK: Function Attrs: nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)
+; CHECK: Function Attrs: nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none)
declare float @llvm.floor.f32(float)
define void @call_floor(float %a) #0 {
@@ -1294,7 +1294,7 @@ attributes #1 = { uwtable noinline }
; TUNIT: attributes #[[ATTR5]] = { noreturn }
; TUNIT: attributes #[[ATTR6]] = { noinline noreturn nounwind uwtable }
; TUNIT: attributes #[[ATTR7]] = { noinline nounwind uwtable }
-; TUNIT: attributes #[[ATTR8:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; TUNIT: attributes #[[ATTR8:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; TUNIT: attributes #[[ATTR9:[0-9]+]] = { norecurse willreturn }
; TUNIT: attributes #[[ATTR10]] = { mustprogress noinline nounwind willreturn uwtable }
; TUNIT: attributes #[[ATTR11:[0-9]+]] = { noinline willreturn uwtable }
@@ -1332,7 +1332,7 @@ attributes #1 = { uwtable noinline }
; CGSCC: attributes #[[ATTR5]] = { noreturn }
; CGSCC: attributes #[[ATTR6]] = { noinline noreturn nounwind uwtable }
; CGSCC: attributes #[[ATTR7]] = { noinline nounwind uwtable }
-; CGSCC: attributes #[[ATTR8:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CGSCC: attributes #[[ATTR8:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CGSCC: attributes #[[ATTR9:[0-9]+]] = { norecurse willreturn }
; CGSCC: attributes #[[ATTR10]] = { mustprogress noinline nounwind willreturn uwtable }
; CGSCC: attributes #[[ATTR11:[0-9]+]] = { noinline willreturn uwtable }
diff --git a/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll b/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll
index b8854e3061dbf4..d0442fa7d8e2e7 100644
--- a/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll
+++ b/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll
@@ -266,7 +266,7 @@ define i32 @commutative_intrinsic_intersection_failure(i32 %arg, i32 %arg1) stri
;.
; CHECK: attributes #[[ATTR0]] = { strictfp }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK: attributes #[[ATTR2]] = { memory(none) }
; CHECK: attributes #[[ATTR3]] = { memory(read) }
; CHECK: attributes #[[ATTR4]] = { alwaysinline memory(none) }
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/expand-fp-convert-small.ll b/llvm/test/Transforms/ExpandIRInsts/X86/expand-fp-convert-small.ll
index c468b9465ad631..9fe11c9d90dfd3 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/expand-fp-convert-small.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/expand-fp-convert-small.ll
@@ -336,7 +336,7 @@ define i8 @fptoui_sat_f16_i8(half %x) {
ret i8 %res
}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"branch_weights", i32 1, i32 1048575}
;.
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/expand-int-convert-small.ll b/llvm/test/Transforms/ExpandIRInsts/X86/expand-int-convert-small.ll
index a2eda588ac8dc3..5db0006f1a49f6 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/expand-int-convert-small.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/expand-int-convert-small.ll
@@ -374,7 +374,7 @@ define half @sitofp_i16_f16(i16 %x) {
ret half %res
}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"branch_weights", i32 1, i32 1048575}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1048575, i32 1}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-si129tofp.ll b/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-si129tofp.ll
index 4dccf9585e0a21..b2ba287237161e 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-si129tofp.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-si129tofp.ll
@@ -615,7 +615,7 @@ define <2 x float> @si129tofloatv2(<2 x i129> %a) {
!0 = !{!"function_entry_count", i64 1000}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-ui129tofp.ll b/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-ui129tofp.ll
index 262ed255a96959..39a7f196050c52 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-ui129tofp.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-convert-ui129tofp.ll
@@ -607,7 +607,7 @@ define <2 x float> @ui129tofloatv2(<2 x i129> %a) {
!0 = !{!"function_entry_count", i64 1000}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-optnone.ll b/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-optnone.ll
index 9d21c7ff6dbf63..63ec04eaf3ba09 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-optnone.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/expand-large-fp-optnone.ll
@@ -102,7 +102,7 @@ attributes #0 = { noinline optnone }
!0 = !{!"function_entry_count", i64 1000}
;.
; CHECK: attributes #[[ATTR0]] = { noinline optnone }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/sdiv129.ll b/llvm/test/Transforms/ExpandIRInsts/X86/sdiv129.ll
index 751bdbade15d96..5c82979d230e1e 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/sdiv129.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/sdiv129.ll
@@ -79,7 +79,7 @@ define void @sdiv129(ptr %ptr, ptr %out) nounwind !prof !0 {
!0 = !{!"function_entry_count", i64 1000}
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/srem129.ll b/llvm/test/Transforms/ExpandIRInsts/X86/srem129.ll
index 45491ccda2b191..5e9749bd3946af 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/srem129.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/srem129.ll
@@ -82,7 +82,7 @@ define void @test(ptr %ptr, ptr %out) nounwind !prof !0 {
!0 = !{!"function_entry_count", i64 1000}
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/udiv129.ll b/llvm/test/Transforms/ExpandIRInsts/X86/udiv129.ll
index 6ad696ae446fd5..5c65140b2db467 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/udiv129.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/udiv129.ll
@@ -68,7 +68,7 @@ define void @test(ptr %ptr, ptr %out) nounwind !prof !0 {
!0 = !{!"function_entry_count", i64 1000}
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/urem129.ll b/llvm/test/Transforms/ExpandIRInsts/X86/urem129.ll
index a4c4ac2cba3294..f6bf845fef788c 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/urem129.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/urem129.ll
@@ -72,7 +72,7 @@ define void @test(ptr %ptr, ptr %out) nounwind !prof !0 {
!0 = !{!"function_entry_count", i64 1000}
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandIRInsts/X86/vector.ll b/llvm/test/Transforms/ExpandIRInsts/X86/vector.ll
index 727e59b5bdef0a..ea4f26ef110ff7 100644
--- a/llvm/test/Transforms/ExpandIRInsts/X86/vector.ll
+++ b/llvm/test/Transforms/ExpandIRInsts/X86/vector.ll
@@ -538,7 +538,7 @@ define <vscale x 2 x i129> @sdiv129_scalable(<vscale x 2 x i129> %a, <vscale x 2
!0 = !{!"function_entry_count", i64 1000}
;.
; CHECK: attributes #[[ATTR0]] = { nounwind }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 1, i32 1048575}
diff --git a/llvm/test/Transforms/ExpandMemCmp/AArch64/memcmp.ll b/llvm/test/Transforms/ExpandMemCmp/AArch64/memcmp.ll
index b05c878d1f80aa..90a4d31ccedc2b 100644
--- a/llvm/test/Transforms/ExpandMemCmp/AArch64/memcmp.ll
+++ b/llvm/test/Transforms/ExpandMemCmp/AArch64/memcmp.ll
@@ -862,7 +862,7 @@ define i32 @cmp_eq16(ptr nocapture readonly %x, ptr nocapture readonly %y) {
!0 = !{!"function_entry_count", i64 1000}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"expand-memcmp"}
diff --git a/llvm/test/Transforms/ExpandMemCmp/X86/memcmp-x32.ll b/llvm/test/Transforms/ExpandMemCmp/X86/memcmp-x32.ll
index b7ab20ab6e779a..e88d0c5dc0d4b4 100644
--- a/llvm/test/Transforms/ExpandMemCmp/X86/memcmp-x32.ll
+++ b/llvm/test/Transforms/ExpandMemCmp/X86/memcmp-x32.ll
@@ -565,7 +565,7 @@ define i32 @cmp_eq16(ptr nocapture readonly %x, ptr nocapture readonly %y) {
!0 = !{!"function_entry_count", i64 1000}
;.
-; X32: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X32: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X32: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; X32: [[PROF1]] = !{!"unknown", !"expand-memcmp"}
diff --git a/llvm/test/Transforms/ExpandMemCmp/X86/memcmp.ll b/llvm/test/Transforms/ExpandMemCmp/X86/memcmp.ll
index c7a276ea7b60e2..9fb3b1289ce215 100644
--- a/llvm/test/Transforms/ExpandMemCmp/X86/memcmp.ll
+++ b/llvm/test/Transforms/ExpandMemCmp/X86/memcmp.ll
@@ -1160,9 +1160,9 @@ define i32 @cmp_eq16(ptr nocapture readonly %x, ptr nocapture readonly %y) {
!0 = !{!"function_entry_count", i64 1000}
;.
-; X64_1LD: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X64_1LD: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
-; X64_2LD: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X64_2LD: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X64_1LD: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; X64_1LD: [[PROF1]] = !{!"unknown", !"expand-memcmp"}
diff --git a/llvm/test/Transforms/InstCombine/icmp-select.ll b/llvm/test/Transforms/InstCombine/icmp-select.ll
index c1a7e4dfe1feab..a9c5d70effec85 100644
--- a/llvm/test/Transforms/InstCombine/icmp-select.ll
+++ b/llvm/test/Transforms/InstCombine/icmp-select.ll
@@ -960,7 +960,7 @@ end:
!1 = !{!"branch_weights", i32 2, i32 3}
!2 = !{!"branch_weights", i32 5, i32 7}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 5, i32 7}
diff --git a/llvm/test/Transforms/InstCombine/mul.ll b/llvm/test/Transforms/InstCombine/mul.ll
index 4bb0129f3ddfee..bed57fe816dd8a 100644
--- a/llvm/test/Transforms/InstCombine/mul.ll
+++ b/llvm/test/Transforms/InstCombine/mul.ll
@@ -2522,8 +2522,8 @@ define i1 @neg_mul_add_one_i1(i1 %x, i1 %y) {
!0 = !{!"function_entry_count", i64 1000}
!1 = !{!"branch_weights", i32 2, i32 3}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
diff --git a/llvm/test/Transforms/InstCombine/rem.ll b/llvm/test/Transforms/InstCombine/rem.ll
index d0730cc4ecce7f..ad43b94975a77b 100644
--- a/llvm/test/Transforms/InstCombine/rem.ll
+++ b/llvm/test/Transforms/InstCombine/rem.ll
@@ -1185,7 +1185,7 @@ bb2:
!0 = !{!"function_entry_count", i64 1000}
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"instcombine"}
diff --git a/llvm/test/Transforms/InstCombine/sdiv-canonicalize.ll b/llvm/test/Transforms/InstCombine/sdiv-canonicalize.ll
index ba8712bbb79027..dfec40b5103088 100644
--- a/llvm/test/Transforms/InstCombine/sdiv-canonicalize.ll
+++ b/llvm/test/Transforms/InstCombine/sdiv-canonicalize.ll
@@ -229,7 +229,7 @@ define i32 @sdiv_neg_neg_missing_nsw(i32 %x, i32 %b) {
!0 = !{!"function_entry_count", i64 1000}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"instcombine"}
diff --git a/llvm/test/Transforms/InstCombine/select-and-or.ll b/llvm/test/Transforms/InstCombine/select-and-or.ll
index 28d06a10ff5b63..75d1b0b7cddc24 100644
--- a/llvm/test/Transforms/InstCombine/select-and-or.ll
+++ b/llvm/test/Transforms/InstCombine/select-and-or.ll
@@ -1469,8 +1469,8 @@ define i8 @test_logical_commuted_and_ne_a_b(i1 %other_cond, i8 %a, i8 %b) {
!1 = !{!"branch_weights", i32 2, i32 3}
!2 = !{!"branch_weights", i32 3, i32 2}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 3, i32 2}
diff --git a/llvm/test/Transforms/InstCombine/select-select.ll b/llvm/test/Transforms/InstCombine/select-select.ll
index be6aed456fbf37..967db07c844cbe 100644
--- a/llvm/test/Transforms/InstCombine/select-select.ll
+++ b/llvm/test/Transforms/InstCombine/select-select.ll
@@ -695,7 +695,7 @@ declare void @use32(i32)
!2 = !{!"branch_weights", i32 5, i32 3}
!3 = !{!"branch_weights", i32 7, i32 11}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"instcombine"}
diff --git a/llvm/test/Transforms/InstCombine/select_meta.ll b/llvm/test/Transforms/InstCombine/select_meta.ll
index 41f912482b90e4..72b113c476d8b5 100644
--- a/llvm/test/Transforms/InstCombine/select_meta.ll
+++ b/llvm/test/Transforms/InstCombine/select_meta.ll
@@ -400,8 +400,8 @@ define i1 @fold_select_fcmp_fpmath(i1 %cond, float %a) {
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"branch_weights", i32 2, i32 10}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 10, i32 2}
diff --git a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-exp.ll b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-exp.ll
index e12074404ca646..54fc751fd26f08 100644
--- a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-exp.ll
+++ b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-exp.ll
@@ -621,7 +621,7 @@ define nofpclass(snan) float @qnan_result_demands_snan_src(i1 %cond, float %unkn
!0 = !{!"function_entry_count", i64 1000}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"instcombine"}
diff --git a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-fdiv.ll b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-fdiv.ll
index 117a3a1c6115ef..b2f717165f7c66 100644
--- a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-fdiv.ll
+++ b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass-fdiv.ll
@@ -2483,7 +2483,7 @@ attributes #3 = { denormal_fpenv(ieee|dynamic) }
; CHECK: attributes #[[ATTR1]] = { denormal_fpenv(preservesign) }
; CHECK: attributes #[[ATTR2]] = { denormal_fpenv(dynamic) }
; CHECK: attributes #[[ATTR3]] = { denormal_fpenv(ieee|dynamic) }
-; CHECK: attributes #[[ATTR4:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR4:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"instcombine"}
diff --git a/llvm/test/Transforms/InstCombine/sub-ashr-or-to-icmp-select.ll b/llvm/test/Transforms/InstCombine/sub-ashr-or-to-icmp-select.ll
index fffad1de08601c..3b04dc4d0989b7 100644
--- a/llvm/test/Transforms/InstCombine/sub-ashr-or-to-icmp-select.ll
+++ b/llvm/test/Transforms/InstCombine/sub-ashr-or-to-icmp-select.ll
@@ -321,7 +321,7 @@ define i32 @sub_ashr_or_i32_shift_wrong_bit(i32 %x, i32 %y) {
!0 = !{!"function_entry_count", i64 1000}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i64 1000}
; CHECK: [[PROF1]] = !{!"unknown", !"instcombine"}
diff --git a/llvm/test/Transforms/LoopIdiom/basic.ll b/llvm/test/Transforms/LoopIdiom/basic.ll
index 98a1f2bb40e234..d382ab8f2ffa59 100644
--- a/llvm/test/Transforms/LoopIdiom/basic.ll
+++ b/llvm/test/Transforms/LoopIdiom/basic.ll
@@ -1619,5 +1619,5 @@ define noalias ptr @_ZN8CMSPULog9beginImplEja(ptr nocapture writeonly %0) local_
; CHECK: attributes #[[ATTR1:[0-9]+]] = { nounwind }
; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
; CHECK: attributes #[[ATTR3:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: write) }
-; CHECK: attributes #[[ATTR4:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR4:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
diff --git a/llvm/test/Transforms/LoopIdiom/cyclic-redundancy-check.ll b/llvm/test/Transforms/LoopIdiom/cyclic-redundancy-check.ll
index 29b5875b1e7c68..79d737d38567f2 100644
--- a/llvm/test/Transforms/LoopIdiom/cyclic-redundancy-check.ll
+++ b/llvm/test/Transforms/LoopIdiom/cyclic-redundancy-check.ll
@@ -1412,5 +1412,5 @@ exit:
; TABLE: attributes #[[ATTR0]] = { optsize }
;.
; CLMUL: attributes #[[ATTR0]] = { optsize }
-; CLMUL: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CLMUL: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
diff --git a/llvm/test/Transforms/LoopVectorize/RISCV/veclib-function-calls.ll b/llvm/test/Transforms/LoopVectorize/RISCV/veclib-function-calls.ll
index 4667ed0ae1a766..c91c06d44e687e 100644
--- a/llvm/test/Transforms/LoopVectorize/RISCV/veclib-function-calls.ll
+++ b/llvm/test/Transforms/LoopVectorize/RISCV/veclib-function-calls.ll
@@ -2288,7 +2288,7 @@ define void @tgamma_f32(ptr noalias %in.ptr, ptr noalias %out.ptr) {
}
;.
; CHECK: attributes #[[ATTR0]] = { "target-features"="+v" }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK: attributes #[[ATTR2]] = { "vector-function-abi-variant"="_ZGVrNxv_acos(Sleef_acosdx_u10rvvm2)" }
; CHECK: attributes #[[ATTR3]] = { "vector-function-abi-variant"="_ZGVrNxv_acosf(Sleef_acosfx_u10rvvm2)" }
; CHECK: attributes #[[ATTR4]] = { "vector-function-abi-variant"="_ZGVrNxv_acosh(Sleef_acoshdx_u10rvvm2)" }
diff --git a/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable-dbg.ll b/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable-dbg.ll
index d23643b1f61035..2241025a75c3fd 100644
--- a/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable-dbg.ll
+++ b/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable-dbg.ll
@@ -93,11 +93,11 @@ define i1 @foo(ptr %p) {
;.
; AARCH64: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; AARCH64: attributes #[[ATTR1:[0-9]+]] = { naked noinline }
-; AARCH64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; AARCH64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; AARCH64-OFF: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; AARCH64-OFF: attributes #[[ATTR1:[0-9]+]] = { naked noinline }
-; AARCH64-OFF: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; AARCH64-OFF: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; AARCH64: [[META0:![0-9]+]] = !{i32 4, !"branch-target-enforcement", i32 1}
; AARCH64: [[META1:![0-9]+]] = !{i32 7, !"Dwarf Version", i32 5}
diff --git a/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable.ll b/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable.ll
index 272c704116359c..113c71b4902a0b 100644
--- a/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable.ll
+++ b/llvm/test/Transforms/LowerTypeTests/aarch64-jumptable.ll
@@ -62,7 +62,7 @@ define i1 @foo(ptr %p) {
;.
; AARCH64: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; AARCH64: attributes #[[ATTR1:[0-9]+]] = { naked noinline }
-; AARCH64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; AARCH64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; AARCH64: [[META0:![0-9]+]] = !{i32 4, !"branch-target-enforcement", i32 1}
; AARCH64: [[META1:![0-9]+]] = !{i32 0, !"typeid1"}
diff --git a/llvm/test/Transforms/LowerTypeTests/import-32bit.ll b/llvm/test/Transforms/LowerTypeTests/import-32bit.ll
index 7fa063ed9ec306..2f48876e773bbc 100644
--- a/llvm/test/Transforms/LowerTypeTests/import-32bit.ll
+++ b/llvm/test/Transforms/LowerTypeTests/import-32bit.ll
@@ -39,7 +39,7 @@ define i1 @allones32(ptr %p) {
}
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0]] = !{i32 0, i32 256}
; CHECK: [[META1]] = !{i32 -1, i32 -1}
diff --git a/llvm/test/Transforms/LowerTypeTests/import.ll b/llvm/test/Transforms/LowerTypeTests/import.ll
index 1583dda58cddc5..4018bccb365ba3 100644
--- a/llvm/test/Transforms/LowerTypeTests/import.ll
+++ b/llvm/test/Transforms/LowerTypeTests/import.ll
@@ -255,10 +255,10 @@ define i1 @single(ptr %p) {
}
;.
; X86: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; X86: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X86: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; ARM: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; ARM: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; ARM: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X86: [[META0]] = !{i64 0, i64 256}
; X86: [[META1]] = !{i64 0, i64 64}
diff --git a/llvm/test/Transforms/LowerTypeTests/x86-jumptable-dbg.ll b/llvm/test/Transforms/LowerTypeTests/x86-jumptable-dbg.ll
index 84f124a2db1274..ad2354edcb3f78 100644
--- a/llvm/test/Transforms/LowerTypeTests/x86-jumptable-dbg.ll
+++ b/llvm/test/Transforms/LowerTypeTests/x86-jumptable-dbg.ll
@@ -115,15 +115,15 @@ define i1 @foo(ptr %p) {
;.
; X86_32: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; X86_32: attributes #[[ATTR1:[0-9]+]] = { naked nocf_check noinline }
-; X86_32: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X86_32: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X86_64: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; X86_64: attributes #[[ATTR1:[0-9]+]] = { naked nocf_check noinline }
-; X86_64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X86_64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X86_64-OFF: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; X86_64-OFF: attributes #[[ATTR1:[0-9]+]] = { naked nocf_check noinline }
-; X86_64-OFF: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X86_64-OFF: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X86_32: [[META0:![0-9]+]] = !{i32 8, !"cf-protection-branch", i32 1}
; X86_32: [[META1:![0-9]+]] = !{i32 7, !"Dwarf Version", i32 5}
diff --git a/llvm/test/Transforms/LowerTypeTests/x86-jumptable.ll b/llvm/test/Transforms/LowerTypeTests/x86-jumptable.ll
index 4b3b2d812b81c6..c4c55a67887c21 100644
--- a/llvm/test/Transforms/LowerTypeTests/x86-jumptable.ll
+++ b/llvm/test/Transforms/LowerTypeTests/x86-jumptable.ll
@@ -99,11 +99,11 @@ define i1 @foo(ptr %p) {
;.
; X86_32: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; X86_32: attributes #[[ATTR1:[0-9]+]] = { naked nocf_check noinline }
-; X86_32: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X86_32: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X86_64: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; X86_64: attributes #[[ATTR1:[0-9]+]] = { naked nocf_check noinline }
-; X86_64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; X86_64: attributes #[[ATTR2:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; X86_32: [[META0:![0-9]+]] = !{i32 8, !"cf-protection-branch", i32 1}
; X86_32: [[META1:![0-9]+]] = !{i32 0, !"typeid1"}
diff --git a/llvm/test/Transforms/SimplifyCFG/X86/switch-of-powers-of-two.ll b/llvm/test/Transforms/SimplifyCFG/X86/switch-of-powers-of-two.ll
index ad95af47a286b5..762d8ad9b5ae7e 100644
--- a/llvm/test/Transforms/SimplifyCFG/X86/switch-of-powers-of-two.ll
+++ b/llvm/test/Transforms/SimplifyCFG/X86/switch-of-powers-of-two.ll
@@ -140,8 +140,8 @@ return:
!0 = !{!"function_entry_count", i32 10}
!1 = !{!"branch_weights", i32 10, i32 5, i32 7, i32 11, i32 13, i32 17}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[PROF0]] = !{!"function_entry_count", i32 10}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 58, i32 5}
diff --git a/llvm/test/Transforms/SimplifyCFG/rangereduce.ll b/llvm/test/Transforms/SimplifyCFG/rangereduce.ll
index fa9da7b9a28bf9..e5e4bf6264b702 100644
--- a/llvm/test/Transforms/SimplifyCFG/rangereduce.ll
+++ b/llvm/test/Transforms/SimplifyCFG/rangereduce.ll
@@ -426,7 +426,7 @@ declare void @side_effect(i32)
!1 = !{!"branch_weights", i32 5, i32 7, i32 11, i32 13, i32 17}
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { optsize }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
;.
; CHECK: [[META0:![0-9]+]] = !{!"function_entry_count", i32 100}
; CHECK: [[PROF1]] = !{!"branch_weights", i32 48, i32 5}
diff --git a/llvm/test/tools/llvm-reduce/remove-attributes-from-intrinsics.ll b/llvm/test/tools/llvm-reduce/remove-attributes-from-intrinsics.ll
index b800f9aa97c8fe..b48bbeed3d2066 100644
--- a/llvm/test/tools/llvm-reduce/remove-attributes-from-intrinsics.ll
+++ b/llvm/test/tools/llvm-reduce/remove-attributes-from-intrinsics.ll
@@ -26,7 +26,7 @@ define i32 @t(i32 %a) {
; CHECK-ALL: declare i32 @llvm.uadd.sat.i32(i32, i32) #0
declare i32 @llvm.uadd.sat.i32(i32, i32) #0
-; CHECK-ALL: attributes #0 = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK-ALL: attributes #0 = { nocallback nocreateundeforpoison nodivergencesource nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK-INTERESTINGNESS: attributes #1 = {
; CHECK-INTERESTINGNESS-SAME: "arg4"
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index b497c4dcd7806a..9bc6a858f9bf60 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -517,6 +517,8 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
isNoReturn = true;
else if (R->getName() == "IntrNoCallback")
isNoCallback = true;
+ else if (R->getName() == "IntrNoDivergenceSource")
+ isNoDivergenceSource = true;
else if (R->getName() == "IntrNoSync")
isNoSync = true;
else if (R->getName() == "IntrNoFree")
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index bd8a35dfc4de80..ecb9e8725a5dbe 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -90,6 +90,9 @@ struct CodeGenIntrinsic {
/// True if the intrinsic is no-callback.
bool isNoCallback = false;
+ /// True if the intrinsic cannot introduce divergence.
+ bool isNoDivergenceSource = false;
+
/// True if the intrinsic is no-sync.
bool isNoSync = false;
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index 61a4aa31aa2222..64433068e28592 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -507,9 +507,9 @@ static bool compareFnAttributes(const CodeGenIntrinsic *L,
auto TieBoolAttributes = [](const CodeGenIntrinsic *I) -> auto {
// Sort throwing intrinsics after non-throwing intrinsics.
return std::tie(I->canThrow, I->isNoDuplicate, I->isNoMerge, I->isNoReturn,
- I->isNoCallback, I->isNoSync, I->isNoFree, I->isWillReturn,
- I->isCold, I->isConvergent, I->isSpeculatable,
- I->hasSideEffects, I->isStrictFP,
+ I->isNoCallback, I->isNoDivergenceSource, I->isNoSync,
+ I->isNoFree, I->isWillReturn, I->isCold, I->isConvergent,
+ I->isSpeculatable, I->hasSideEffects, I->isStrictFP,
I->isNoCreateUndefOrPoison);
};
@@ -532,10 +532,11 @@ static bool compareFnAttributes(const CodeGenIntrinsic *L,
/// NoUnwind = !canThrow, so we need to negate it's sense to test if the
// intrinsic has NoUnwind attribute.
static bool hasFnAttributes(const CodeGenIntrinsic &Int) {
- return !Int.canThrow || Int.isNoReturn || Int.isNoCallback || Int.isNoSync ||
- Int.isNoFree || Int.isWillReturn || Int.isCold || Int.isNoDuplicate ||
- Int.isNoMerge || Int.isConvergent || Int.isSpeculatable ||
- Int.isStrictFP || Int.isNoCreateUndefOrPoison ||
+ return !Int.canThrow || Int.isNoReturn || Int.isNoCallback ||
+ Int.isNoDivergenceSource || Int.isNoSync || Int.isNoFree ||
+ Int.isWillReturn || Int.isCold || Int.isNoDuplicate || Int.isNoMerge ||
+ Int.isConvergent || Int.isSpeculatable || Int.isStrictFP ||
+ Int.isNoCreateUndefOrPoison ||
getEffectiveME(Int) != MemoryEffects::unknown();
}
@@ -679,6 +680,8 @@ static AttributeSet getIntrinsicFnAttributeSet(LLVMContext &C, unsigned ID) {
addAttribute("NoReturn");
if (Int.isNoCallback)
addAttribute("NoCallback");
+ if (Int.isNoDivergenceSource)
+ addAttribute("NoDivergenceSource");
if (Int.isNoSync)
addAttribute("NoSync");
if (Int.isNoFree)
>From 6f5d013d7acbba1b0ec0f5e38236c53e6a5483ef Mon Sep 17 00:00:00 2001
From: lezcano <lezcano-93 at hotmail.com>
Date: Thu, 15 Oct 2026 11:58:31 +0200
Subject: [PATCH 3/4] [NVPTX] Improve value uniformity analysis
Teach NVPTX TTI about warp-uniform values and conservative divergence
sources. Use getUniformityOperands to identify the defined, uniform
controls that make collective results uniform, and add analysis tests.
---
.../Target/NVPTX/NVPTXTargetTransformInfo.cpp | 265 +++++++++++----
.../Target/NVPTX/NVPTXTargetTransformInfo.h | 6 +-
.../UniformityAnalysis/NVPTX/intrinsics.ll | 310 ++++++++++++++++++
3 files changed, 522 insertions(+), 59 deletions(-)
create mode 100644 llvm/test/Analysis/UniformityAnalysis/NVPTX/intrinsics.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
index 9071c2a3671d2b..b49901944169f4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
@@ -9,7 +9,9 @@
#include "NVPTXTargetTransformInfo.h"
#include "NVVMProperties.h"
#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallBitVector.h"
#include "llvm/Analysis/LoopInfo.h"
+#include "llvm/Analysis/SimplifyQuery.h"
#include "llvm/Analysis/TargetTransformInfo.h"
#include "llvm/Analysis/ValueTracking.h"
#include "llvm/CodeGen/BasicTTIImpl.h"
@@ -18,6 +20,7 @@
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
+#include "llvm/IR/PatternMatch.h"
#include "llvm/IR/Value.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/ErrorHandling.h"
@@ -28,60 +31,6 @@ using namespace llvm;
#define DEBUG_TYPE "NVPTXtti"
-// Whether the given intrinsic reads threadIdx.x/y/z.
-static bool readsThreadIndex(const IntrinsicInst *II) {
- switch (II->getIntrinsicID()) {
- default: return false;
- case Intrinsic::nvvm_read_ptx_sreg_tid_x:
- case Intrinsic::nvvm_read_ptx_sreg_tid_y:
- case Intrinsic::nvvm_read_ptx_sreg_tid_z:
- return true;
- }
-}
-
-static bool readsLaneId(const IntrinsicInst *II) {
- return II->getIntrinsicID() == Intrinsic::nvvm_read_ptx_sreg_laneid;
-}
-
-bool NVPTXTTIImpl::isSourceOfDivergence(const Value *V) const {
- // Without inter-procedural analysis, we conservatively assume that arguments
- // to __device__ functions are divergent.
- if (const Argument *Arg = dyn_cast<Argument>(V))
- return !isKernelFunction(*Arg->getParent());
-
- if (const Instruction *I = dyn_cast<Instruction>(V)) {
- // Without pointer analysis, we conservatively assume values loaded from
- // generic or local address space are divergent.
- if (const LoadInst *LI = dyn_cast<LoadInst>(I)) {
- unsigned AS = LI->getPointerAddressSpace();
- return AS == ADDRESS_SPACE_GENERIC || AS == ADDRESS_SPACE_LOCAL;
- }
- // Atomic instructions may cause divergence. Atomic instructions are
- // executed sequentially across all threads in a warp. Therefore, an earlier
- // executed thread may see different memory inputs than a later executed
- // thread. For example, suppose *a = 0 initially.
- //
- // atom.global.add.s32 d, [a], 1
- //
- // returns 0 for the first thread that enters the critical region, and 1 for
- // the second thread.
- if (I->isAtomic())
- return true;
- if (const IntrinsicInst *II = dyn_cast<IntrinsicInst>(I)) {
- // Instructions that read threadIdx are obviously divergent.
- if (readsThreadIndex(II) || readsLaneId(II))
- return true;
- }
- // Conservatively consider the return value of function calls as divergent.
- // We could analyze callees with bodies more precisely using
- // inter-procedural analysis.
- if (isa<CallBase>(I))
- return true;
- }
-
- return false;
-}
-
// Convert NVVM intrinsics to target-generic LLVM code where possible.
static Instruction *convertNvvmIntrinsicToLlvm(InstCombiner &IC,
IntrinsicInst *II) {
@@ -691,9 +640,215 @@ bool NVPTXTTIImpl::shouldBuildLookupTablesForConstant(Constant *C) const {
return true;
}
+// Operands whose defined, uniform values make the collective result uniform.
+// Other data operands may vary between lanes.
+static SmallVector<unsigned, 2> getUniformityOperands(const IntrinsicInst &II) {
+ using namespace PatternMatch;
+ uint64_t Clamp, Index;
+ switch (II.getIntrinsicID()) {
+ case Intrinsic::nvvm_shfl_sync_idx_f32:
+ case Intrinsic::nvvm_shfl_sync_idx_f32p:
+ case Intrinsic::nvvm_shfl_sync_idx_i32:
+ case Intrinsic::nvvm_shfl_sync_idx_i32p:
+ // Require no segmentation or per-lane fallback to the input value.
+ if (match(II.getArgOperand(3), m_ConstantInt(Clamp)) &&
+ (Clamp & 0x1f00) == 0 &&
+ ((Clamp & 31) == 31 ||
+ (match(II.getArgOperand(2), m_ConstantInt(Index)) &&
+ (Index & 31) <= (Clamp & 31))))
+ return {0, 2};
+ return {};
+ case Intrinsic::nvvm_barrier_cta_red_and_aligned_all:
+ case Intrinsic::nvvm_barrier_cta_red_and_all:
+ case Intrinsic::nvvm_barrier_cta_red_or_aligned_all:
+ case Intrinsic::nvvm_barrier_cta_red_or_all:
+ case Intrinsic::nvvm_barrier_cta_red_popc_aligned_all:
+ case Intrinsic::nvvm_barrier_cta_red_popc_all:
+ case Intrinsic::nvvm_match_all_sync_i32p:
+ case Intrinsic::nvvm_match_all_sync_i64p:
+ case Intrinsic::nvvm_vote_all_sync:
+ case Intrinsic::nvvm_vote_any_sync:
+ case Intrinsic::nvvm_vote_ballot_sync:
+ case Intrinsic::nvvm_vote_uni_sync:
+ return {0};
+ case Intrinsic::nvvm_match_any_sync_i32:
+ case Intrinsic::nvvm_match_any_sync_i64:
+ case Intrinsic::nvvm_barrier_cta_red_and_aligned_count:
+ case Intrinsic::nvvm_barrier_cta_red_and_count:
+ case Intrinsic::nvvm_barrier_cta_red_or_aligned_count:
+ case Intrinsic::nvvm_barrier_cta_red_or_count:
+ case Intrinsic::nvvm_barrier_cta_red_popc_aligned_count:
+ case Intrinsic::nvvm_barrier_cta_red_popc_count:
+ return {0, 1};
+ case Intrinsic::nvvm_redux_sync_add:
+ case Intrinsic::nvvm_redux_sync_and:
+ case Intrinsic::nvvm_redux_sync_fmax:
+ case Intrinsic::nvvm_redux_sync_fmax_NaN:
+ case Intrinsic::nvvm_redux_sync_fmax_abs:
+ case Intrinsic::nvvm_redux_sync_fmax_abs_NaN:
+ case Intrinsic::nvvm_redux_sync_fmin:
+ case Intrinsic::nvvm_redux_sync_fmin_NaN:
+ case Intrinsic::nvvm_redux_sync_fmin_abs:
+ case Intrinsic::nvvm_redux_sync_fmin_abs_NaN:
+ case Intrinsic::nvvm_redux_sync_max:
+ case Intrinsic::nvvm_redux_sync_min:
+ case Intrinsic::nvvm_redux_sync_or:
+ case Intrinsic::nvvm_redux_sync_umax:
+ case Intrinsic::nvvm_redux_sync_umin:
+ case Intrinsic::nvvm_redux_sync_xor:
+ return {1};
+ default:
+ return {};
+ }
+}
+
ValueUniformity NVPTXTTIImpl::getValueUniformity(const Value *V) const {
- if (isSourceOfDivergence(V))
+ using namespace PatternMatch;
+
+ if (const auto *Arg = dyn_cast<Argument>(V))
+ return isKernelFunction(*Arg->getParent()) ? ValueUniformity::Default
+ : ValueUniformity::NeverUniform;
+
+ const auto *I = dyn_cast<Instruction>(V);
+ if (!I)
+ return ValueUniformity::Default;
+
+ // Even a uniform address can observe different writes between lanes.
+ if (const auto *LI = dyn_cast<LoadInst>(I))
+ return LI->isSimple() && LI->getPointerAddressSpace() == ADDRESS_SPACE_CONST
+ ? ValueUniformity::Default
+ : ValueUniformity::NeverUniform;
+ if (I->isAtomic() || isa<AllocaInst>(I))
return ValueUniformity::NeverUniform;
+ if (const auto *Cast = dyn_cast<AddrSpaceCastInst>(I))
+ if (Cast->getSrcAddressSpace() == ADDRESS_SPACE_LOCAL &&
+ Cast->getDestAddressSpace() == ADDRESS_SPACE_GENERIC)
+ return ValueUniformity::NeverUniform;
+ if (const auto *Freeze = dyn_cast<FreezeInst>(I))
+ if (!isGuaranteedNotToBeUndefOrPoison(Freeze->getOperand(0)))
+ return ValueUniformity::NeverUniform;
+
+ if (const auto *II = dyn_cast<IntrinsicInst>(I)) {
+ auto UniformityOperands = getUniformityOperands(*II);
+ if (!UniformityOperands.empty()) {
+ // Undef controls can select different, individually valid groups/lanes.
+ return llvm::all_of(UniformityOperands,
+ [&](unsigned Op) {
+ return isGuaranteedNotToBeUndefOrPoison(
+ II->getArgOperand(Op));
+ })
+ ? ValueUniformity::Custom
+ : ValueUniformity::NeverUniform;
+ }
+ Intrinsic::ID IID = II->getIntrinsicID();
+ if (const auto *CFP = dyn_cast<ConstrainedFPIntrinsic>(II)) {
+ switch (IID) {
+ case Intrinsic::experimental_constrained_lround:
+ case Intrinsic::experimental_constrained_llround:
+ case Intrinsic::experimental_constrained_lrint:
+ case Intrinsic::experimental_constrained_llrint:
+ // An out-of-range result can be an arbitrary defined integer.
+ return ValueUniformity::NeverUniform;
+ default:
+ return CFP->getExceptionBehavior() == fp::ebIgnore &&
+ CFP->getRoundingMode() != RoundingMode::Dynamic
+ ? ValueUniformity::Default
+ : ValueUniformity::NeverUniform;
+ }
+ }
+
+ switch (IID) {
+ case Intrinsic::frexp:
+ // The exponent is unspecified for NaN and infinity.
+ return isKnownNeverInfOrNaN(II->getArgOperand(0),
+ SimplifyQuery(I->getDataLayout()))
+ ? ValueUniformity::Default
+ : ValueUniformity::NeverUniform;
+ case Intrinsic::nvvm_read_ptx_sreg_tid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_z: {
+ auto ReqNTID = getReqNTID(*I->getFunction());
+ if (!isKernelFunction(*I->getFunction()) || ReqNTID.empty())
+ return ValueUniformity::NeverUniform;
+ ReqNTID.resize(3, 1);
+ unsigned Dim = IID == Intrinsic::nvvm_read_ptx_sreg_tid_x ? 0
+ : IID == Intrinsic::nvvm_read_ptx_sreg_tid_y ? 1
+ : 2;
+ unsigned Stride = llvm::product_of(ArrayRef(ReqNTID).take_front(Dim));
+ return ReqNTID[Dim] == 1 || Stride % 32 == 0
+ ? ValueUniformity::AlwaysUniform
+ : ValueUniformity::NeverUniform;
+ }
+ case Intrinsic::nvvm_is_explicit_cluster:
+ case Intrinsic::nvvm_read_ptx_sreg_aggr_smem_size:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_ctaid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_ctaid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_ctaid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_ctaid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_ctarank:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_nctaid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_nctaid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_nctaid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_nctaid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_cluster_nctarank:
+ case Intrinsic::nvvm_read_ptx_sreg_clusterid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_clusterid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_clusterid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_clusterid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_ctaid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_ctaid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_ctaid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_ctaid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_dynamic_smem_size:
+ case Intrinsic::nvvm_read_ptx_sreg_gridid:
+ case Intrinsic::nvvm_read_ptx_sreg_nclusterid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_nclusterid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_nclusterid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_nclusterid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_nctaid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_nctaid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_nctaid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_nctaid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_nsmid:
+ case Intrinsic::nvvm_read_ptx_sreg_ntid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_ntid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_ntid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_ntid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_nwarpid:
+ case Intrinsic::nvvm_read_ptx_sreg_reserved_smem_offset_begin:
+ case Intrinsic::nvvm_read_ptx_sreg_reserved_smem_offset_cap:
+ case Intrinsic::nvvm_read_ptx_sreg_reserved_smem_offset_end:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_w:
+ case Intrinsic::nvvm_read_ptx_sreg_total_smem_size:
+ case Intrinsic::nvvm_read_ptx_sreg_warpsize:
+ return ValueUniformity::AlwaysUniform;
+ default:
+ return ValueUniformity::NeverUniform;
+ }
+ }
+ if (isa<CallBase>(I))
+ return ValueUniformity::NeverUniform;
+
+ uint64_t C;
+ if ((match(V, m_Shr(m_Intrinsic<Intrinsic::nvvm_read_ptx_sreg_tid_x>(),
+ m_ConstantInt(C))) &&
+ C >= 5) ||
+ (match(V, m_c_And(m_Intrinsic<Intrinsic::nvvm_read_ptx_sreg_tid_x>(),
+ m_ConstantInt(C))) &&
+ (C & 31) == 0)) {
+ // maxntid limits only the total thread count, not individual dimensions.
+ const auto ReqNTID = getReqNTID(*I->getFunction());
+ if (isKernelFunction(*I->getFunction()) && !ReqNTID.empty() &&
+ (ReqNTID[0] <= 32 || ReqNTID[0] % 32 == 0 ||
+ llvm::all_of(drop_begin(ReqNTID), [](unsigned N) { return N == 1; })))
+ return ValueUniformity::AlwaysUniform;
+ }
return ValueUniformity::Default;
}
+
+bool NVPTXTTIImpl::isUniform(const Instruction *I,
+ const SmallBitVector &UniformArgs) const {
+ return llvm::all_of(getUniformityOperands(*cast<IntrinsicInst>(I)),
+ [&](unsigned Op) { return UniformArgs[Op]; });
+}
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h
index a9c8964c55eac3..c738d31a1e325d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h
@@ -37,10 +37,6 @@ class NVPTXTTIImpl final : public BasicTTIImplBase<NVPTXTTIImpl> {
const NVPTXSubtarget *getST() const { return ST; };
const NVPTXTargetLowering *getTLI() const { return TLI; };
- /// \returns true if the result of the value could potentially be
- /// different across threads in a warp.
- bool isSourceOfDivergence(const Value *V) const;
-
public:
explicit NVPTXTTIImpl(const NVPTXTargetMachine *TM, const Function &F)
: BaseT(TM, F.getDataLayout()), ST(TM->getSubtargetImpl()),
@@ -242,6 +238,8 @@ class NVPTXTTIImpl final : public BasicTTIImplBase<NVPTXTTIImpl> {
}
ValueUniformity getValueUniformity(const Value *V) const override;
+ bool isUniform(const Instruction *I,
+ const SmallBitVector &UniformArgs) const override;
};
} // end namespace llvm
diff --git a/llvm/test/Analysis/UniformityAnalysis/NVPTX/intrinsics.ll b/llvm/test/Analysis/UniformityAnalysis/NVPTX/intrinsics.ll
new file mode 100644
index 00000000000000..5a0aae29ee5ce5
--- /dev/null
+++ b/llvm/test/Analysis/UniformityAnalysis/NVPTX/intrinsics.ll
@@ -0,0 +1,310 @@
+; RUN: opt %s -passes='print<uniformity>' -disable-output > %t 2>&1
+; RUN: FileCheck %s < %t
+; RUN: FileCheck %s --check-prefix=UNIFORM < %t
+; UNIFORM-NOT: DIVERGENT: %uniform.
+
+target triple = "nvptx64-nvidia-cuda"
+
+; Values named uniform.* must remain uniform; mutable memory and unknown calls
+; can return different values to threads with identical operands.
+define ptx_kernel void @sources(i32 noundef %arg, ptr addrspace(1) %global, ptr addrspace(4) %constant, ptr addrspace(5) %local) {
+; CHECK-LABEL: for function 'sources'
+; CHECK-DAG: DIVERGENT: %mutable =
+; CHECK-DAG: DIVERGENT: %atomic =
+; CHECK-DAG: DIVERGENT: %volatile =
+; CHECK-DAG: DIVERGENT: %allocation =
+; CHECK-DAG: DIVERGENT: %local.generic =
+; CHECK-DAG: DIVERGENT: %frozen =
+; CHECK-DAG: DIVERGENT: %thread =
+; CHECK-DAG: DIVERGENT: %clock =
+; CHECK-DAG: DIVERGENT: %unknown =
+ %uniform.block = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+ %uniform.size = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
+ %uniform.min = call i32 @llvm.umin.i32(i32 %arg, i32 %uniform.block)
+ %uniform.constant = load i32, ptr addrspace(4) %constant
+ %uniform.frozen = freeze i32 %arg
+ %mutable = load i32, ptr addrspace(1) %global
+ %atomic = load atomic i32, ptr addrspace(1) %global monotonic, align 4
+ %volatile = load volatile i32, ptr addrspace(4) %constant
+ %allocation = alloca i32
+ %local.generic = addrspacecast ptr addrspace(5) %local to ptr
+ %frozen = freeze i32 poison
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %clock = call i32 @llvm.nvvm.read.ptx.sreg.clock()
+ %unknown = call i32 @unknown(i32 %arg)
+ ret void
+}
+
+define ptx_kernel void @collectives(i32 noundef %mask, i32 noundef %index, i32 %maybe.undef) {
+; CHECK-LABEL: for function 'collectives'
+; CHECK-DAG: DIVERGENT: %thread =
+; CHECK-DAG: DIVERGENT: %frozen.broadcast =
+; CHECK-DAG: DIVERGENT: %frozen.extracted =
+; CHECK-DAG: DIVERGENT: %segmented =
+; CHECK-DAG: DIVERGENT: %out.of.range =
+; CHECK-DAG: DIVERGENT: %undefined.match.data =
+; CHECK-DAG: DIVERGENT: %varying.index =
+; CHECK-DAG: DIVERGENT: %varying.mask =
+; CHECK-DAG: DIVERGENT: %undefined.mask =
+; CHECK-DAG: DIVERGENT: %unknown.mask =
+; CHECK-DAG: DIVERGENT: %match.any =
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %uniform.broadcast = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 -1, i32 %thread, i32 7, i32 31)
+ ; An inactive source can yield undef; freezing it can differ per lane.
+ %frozen.broadcast = freeze i32 %uniform.broadcast
+ %uniform.pair = call {i32, i1} @llvm.nvvm.shfl.sync.idx.i32p(i32 -1, i32 %thread, i32 7, i32 31)
+ %uniform.extracted = extractvalue {i32, i1} %uniform.pair, 0
+ %frozen.extracted = freeze i32 %uniform.extracted
+ %predicate = icmp ne i32 %thread, 0
+ %uniform.vote = call i32 @llvm.nvvm.vote.ballot.sync(i32 %mask, i1 %predicate)
+ %uniform.redux = call i32 @llvm.nvvm.redux.sync.add(i32 %thread, i32 %mask)
+ %fp = uitofp i32 %thread to float
+ %uniform.fp.redux = call float @llvm.nvvm.redux.sync.fmax(float %fp, i32 %mask)
+ %uniform.match = call {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32 %mask, i32 %thread)
+ %uniform.partial = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 %mask, i32 %thread, i32 %index, i32 31)
+ %uniform.barrier = call i32 @llvm.nvvm.barrier.cta.red.popc.count(i32 0, i32 32, i1 %predicate)
+ %uniform.clamped = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 %mask, i32 %thread, i32 7, i32 15)
+ %uniform.match.any = call i32 @llvm.nvvm.match.any.sync.i32(i32 %mask, i32 %index)
+ %segmented = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 -1, i32 %thread, i32 0, i32 4111)
+ %out.of.range = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 -1, i32 %thread, i32 16, i32 15)
+ %undefined.match.data = call i32 @llvm.nvvm.match.any.sync.i32(i32 -1, i32 undef)
+ %varying.index = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 -1, i32 %thread, i32 %thread, i32 31)
+ %varying.mask = call i32 @llvm.nvvm.vote.ballot.sync(i32 %thread, i1 true)
+ %undefined.mask = call i32 @llvm.nvvm.vote.ballot.sync(i32 undef, i1 true)
+ %unknown.mask = call i32 @llvm.nvvm.vote.ballot.sync(i32 %maybe.undef, i1 true)
+ %match.any = call i32 @llvm.nvvm.match.any.sync.i32(i32 %mask, i32 %thread)
+ ret void
+}
+
+; Data-only shuffles preserve uniform inputs, but their range predicates can vary.
+define ptx_kernel void @uniform_shuffle_data(i32 %data) {
+; CHECK-LABEL: for function 'uniform_shuffle_data'
+; CHECK-DAG: DIVERGENT: %varying.data =
+; CHECK-DAG: DIVERGENT: %range.predicate =
+; CHECK-DAG: DIVERGENT: %frozen.data =
+; CHECK-DAG: DIVERGENT: %legacy.varying =
+; CHECK-DAG: DIVERGENT: %legacy.range =
+; CHECK-DAG: DIVERGENT: %legacy.vote =
+; CHECK-DAG: DIVERGENT: %legacy.ballot =
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %uniform.up = call i32 @llvm.nvvm.shfl.sync.up.i32(i32 -1, i32 %data, i32 1, i32 0)
+ %uniform.down = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 -1, i32 %data, i32 1, i32 31)
+ %uniform.bfly = call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %data, i32 1, i32 31)
+ %uniform.segmented = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 -1, i32 %data, i32 0, i32 4111)
+ %varying.data = call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, i32 %thread, i32 1, i32 31)
+ %pair = call {i32, i1} @llvm.nvvm.shfl.sync.up.i32p(i32 -1, i32 %data, i32 1, i32 0)
+ %range.predicate = extractvalue {i32, i1} %pair, 1
+ %frozen.data = freeze i32 %uniform.up
+ %uniform.legacy = call i32 @llvm.nvvm.shfl.up.i32(i32 %data, i32 1, i32 0)
+ %legacy.varying = call i32 @llvm.nvvm.shfl.up.i32(i32 %thread, i32 1, i32 0)
+ %legacy.pair = call {i32, i1} @llvm.nvvm.shfl.up.i32p(i32 %data, i32 1, i32 0)
+ %legacy.range = extractvalue {i32, i1} %legacy.pair, 1
+ %uniform.legacy.vote = call i1 @llvm.nvvm.vote.all(i1 true)
+ %predicate = icmp ne i32 %thread, 0
+ %legacy.vote = call i1 @llvm.nvvm.vote.all(i1 %predicate)
+ %legacy.ballot = call i32 @llvm.nvvm.vote.ballot(i1 true)
+ ret void
+}
+
+; Separate groups can each return a uniform result that differs at the join.
+define ptx_kernel void @partial_groups() {
+; CHECK-LABEL: for function 'partial_groups'
+; CHECK: DIVERGENT: %joined =
+ %lane = call i32 @llvm.nvvm.read.ptx.sreg.laneid()
+ %first.half = icmp ult i32 %lane, 16
+ br i1 %first.half, label %low, label %high
+low:
+ %uniform.low = call i32 @llvm.nvvm.vote.ballot.sync(i32 65535, i1 true)
+ br label %join
+high:
+ %uniform.high = call i32 @llvm.nvvm.vote.ballot.sync(i32 -65536, i1 true)
+ br label %join
+join:
+ %joined = phi i32 [%uniform.low, %low], [%uniform.high, %high]
+ ret void
+}
+
+define ptx_kernel void @warp_index() "nvvm.reqntid"="64,2,1" {
+; CHECK-LABEL: for function 'warp_index'
+; CHECK: DIVERGENT: %thread =
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %uniform.warp = lshr i32 %thread, 5
+ %uniform.masked = and i32 %thread, -32
+ ret void
+}
+
+; A warp may straddle rows when the x dimension is not a multiple of 32.
+define ptx_kernel void @straddling_rows() "nvvm.reqntid"="48,2,1" {
+; CHECK-LABEL: for function 'straddling_rows'
+; CHECK: DIVERGENT: %warp =
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %warp = lshr i32 %thread, 5
+ ret void
+}
+
+; A maximum dimension does not establish the actual row boundaries.
+define ptx_kernel void @maximum_only() "nvvm.maxntid"="64,2,1" {
+; CHECK-LABEL: for function 'maximum_only'
+; CHECK: DIVERGENT: %warp =
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %warp = lshr i32 %thread, 5
+ ret void
+}
+
+declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+declare i32 @llvm.nvvm.read.ptx.sreg.clock()
+declare i32 @llvm.umin.i32(i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
+declare {i32, i1} @llvm.nvvm.shfl.sync.up.i32p(i32, i32, i32, i32)
+declare i32 @llvm.nvvm.vote.ballot.sync(i32, i1)
+declare i32 @llvm.nvvm.redux.sync.add(i32, i32)
+declare i32 @unknown(i32)
+
+; Ordinary arithmetic propagates divergence; warp collectives can remove it.
+define ptx_kernel void @arithmetic(float %x, i32 %n, <4 x i32> %vector) {
+; CHECK-LABEL: for function 'arithmetic'
+; CHECK-DAG: DIVERGENT: %varying.sqrt =
+; CHECK-DAG: DIVERGENT: %varying.random =
+; CHECK-DAG: DIVERGENT: %rounded =
+; CHECK-DAG: DIVERGENT: %exponent =
+; CHECK-DAG: DIVERGENT: %dynamic.rounding =
+; CHECK-DAG: DIVERGENT: %matrix =
+; CHECK-DAG: DIVERGENT: %varying.repeat =
+; CHECK-DAG: DIVERGENT: %varying.mask =
+; CHECK-DAG: DIVERGENT: %varying.ue5m3 =
+; CHECK-DAG: DIVERGENT: %varying.rounding =
+; CHECK-DAG: DIVERGENT: %strict.rounding =
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %fp = uitofp i32 %thread to float
+ %uniform.sqrt = call float @llvm.sqrt.f32(float %x)
+ %varying.sqrt = call float @llvm.sqrt.f32(float %fp)
+ %uniform.overflow = call {i32, i1} @llvm.uadd.with.overflow.i32(i32 %n, i32 1)
+ %uniform.bits = call i32 @llvm.ctpop.i32(i32 %n)
+ %uniform.vector = call <4 x i32> @llvm.vp.ctpop.v4i32(<4 x i32> %vector, <4 x i1> <i1 true, i1 true, i1 true, i1 true>, i32 4)
+ %uniform.reverse = call <4 x i32> @llvm.experimental.vp.reverse.v4i32(<4 x i32> %vector, <4 x i1> <i1 true, i1 true, i1 true, i1 true>, i32 4)
+ %uniform.vector.sum = call i32 @llvm.vector.reduce.add.v4i32(<4 x i32> %vector)
+ %uniform.repeat = call <vscale x 4 x i32> @llvm.vector.repeat.nxv4i32.v4i32(<4 x i32> %vector)
+ %varying.vector = insertelement <4 x i32> %vector, i32 %thread, i32 0
+ %varying.repeat = call <vscale x 4 x i32> @llvm.vector.repeat.nxv4i32.v4i32(<4 x i32> %varying.vector)
+ %uniform.mask.input = icmp ne <4 x i32> %vector, zeroinitializer
+ %uniform.mask = call <4 x i1> @llvm.mask.beforefirst.v4i1(<4 x i1> %uniform.mask.input)
+ %varying.mask.input = icmp ne <4 x i32> %varying.vector, zeroinitializer
+ %varying.mask = call <4 x i1> @llvm.mask.beforefirst.v4i1(<4 x i1> %varying.mask.input)
+ %uniform.ex2 = call float @llvm.nvvm.ex2.approx.f32(float %x)
+ %uniform.permute = call i32 @llvm.nvvm.prmt(i32 %n, i32 0, i32 0)
+ %uniform.random = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %x, float %x, i32 %n)
+ %varying.random = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %x, float %x, i32 %thread)
+ %uniform.ue5m3 = call i16 @llvm.nvvm.ff.to.ue5m3x2.rn(float %x, float %x)
+ %uniform.ue5m3.unpack = call <2 x half> @llvm.nvvm.ue5m3x2.to.f16x2.rn(i16 %uniform.ue5m3)
+ %varying.ue5m3 = call i16 @llvm.nvvm.ff.to.ue5m3x2.rn(float %x, float %fp)
+ %rounded = call i32 @llvm.lround.i32.f32(float %x)
+ %exponent = call {float, i32} @llvm.frexp.f32.i32(float %x)
+ %uniform.exponent = call {float, i32} @llvm.frexp.f32.i32(float 2.0)
+ %uniform.constrained = call float @llvm.experimental.constrained.fadd.f32(float %x, float %x, metadata !"round.tonearest", metadata !"fpexcept.ignore")
+ %uniform.rounding.down = call float @llvm.experimental.constrained.fadd.f32(float %x, float %x, metadata !"round.downward", metadata !"fpexcept.ignore")
+ %varying.rounding = call float @llvm.experimental.constrained.fadd.f32(float %x, float %fp, metadata !"round.downward", metadata !"fpexcept.ignore")
+ %strict.rounding = call float @llvm.experimental.constrained.fadd.f32(float %x, float %x, metadata !"round.downward", metadata !"fpexcept.strict")
+ %dynamic.rounding = call float @llvm.experimental.constrained.fadd.f32(float %x, float %x, metadata !"round.dynamic", metadata !"fpexcept.ignore")
+ ; A transpose redistributes packed components even with identical inputs.
+ %matrix = call i32 @llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16(i32 65538)
+ ret void
+}
+
+define ptx_kernel void @state_queries(ptr %p, ptr addrspace(1) %global, i64 %state) {
+; CHECK-LABEL: for function 'state_queries'
+; CHECK-DAG: DIVERGENT: %physical.warp =
+; CHECK-DAG: DIVERGENT: %physical.sm =
+; CHECK-DAG: DIVERGENT: %active =
+; CHECK-DAG: DIVERGENT: %barrier.layout =
+ %uniform.smem = call i32 @llvm.nvvm.read.ptx.sreg.dynamic_smem_size()
+ %uniform.space = call i1 @llvm.nvvm.isspacep.global(ptr %p)
+ %uniform.pending = call i32 @llvm.nvvm.mbarrier.pending.count(i64 %state)
+ %barrier.layout = call i1 @llvm.nvvm.mbarrier.check_layout.p0(ptr %p, i32 0)
+ %uniform.ldu = call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %global, i32 4)
+ %physical.warp = call i32 @llvm.nvvm.read.ptx.sreg.warpid()
+ %physical.sm = call i32 @llvm.nvvm.read.ptx.sreg.smid()
+ %active = call i32 @llvm.nvvm.activemask()
+ ret void
+}
+
+define ptx_kernel void @aligned_rows() "nvvm.reqntid"="32,2,2" {
+; CHECK-LABEL: for function 'aligned_rows'
+ %uniform.y = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+ %uniform.z = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
+ ret void
+}
+
+define ptx_kernel void @straddling_planes() "nvvm.reqntid"="16,3,2" {
+; CHECK-LABEL: for function 'straddling_planes'
+; CHECK-DAG: DIVERGENT: %y =
+; CHECK-DAG: DIVERGENT: %z =
+ %y = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+ %z = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
+ ret void
+}
+
+declare i32 @llvm.nvvm.read.ptx.sreg.laneid()
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
+declare {i32, i1} @llvm.nvvm.shfl.sync.idx.i32p(i32, i32, i32, i32)
+declare float @llvm.nvvm.redux.sync.fmax(float, i32)
+declare {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32, i32)
+declare i32 @llvm.nvvm.match.any.sync.i32(i32, i32)
+declare i32 @llvm.nvvm.barrier.cta.red.popc.count(i32, i32, i1)
+declare float @llvm.sqrt.f32(float)
+declare {i32, i1} @llvm.uadd.with.overflow.i32(i32, i32)
+declare i32 @llvm.ctpop.i32(i32)
+declare <4 x i32> @llvm.vp.ctpop.v4i32(<4 x i32>, <4 x i1>, i32)
+declare i32 @llvm.vector.reduce.add.v4i32(<4 x i32>)
+declare float @llvm.nvvm.ex2.approx.f32(float)
+declare i32 @llvm.nvvm.prmt(i32, i32, i32)
+declare <2 x half> @llvm.nvvm.ff2f16x2.rs(float, float, i32)
+declare i32 @llvm.lround.i32.f32(float)
+declare {float, i32} @llvm.frexp.f32.i32(float)
+declare float @llvm.experimental.constrained.fadd.f32(float, float, metadata, metadata)
+declare i32 @llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16(i32)
+declare i32 @llvm.nvvm.read.ptx.sreg.dynamic_smem_size()
+declare i1 @llvm.nvvm.isspacep.global(ptr)
+declare i32 @llvm.nvvm.mbarrier.pending.count(i64)
+declare i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1), i32)
+declare i32 @llvm.nvvm.read.ptx.sreg.warpid()
+declare i32 @llvm.nvvm.read.ptx.sreg.smid()
+declare i32 @llvm.nvvm.activemask()
+
+; Threads terminate at different iterations; later ballots exclude them.
+; A uniform result within the loop is divergent when used on different exits.
+define ptx_kernel void @collective_loop_exit(ptr %out) {
+; CHECK-LABEL: for function 'collective_loop_exit'
+; CHECK: DIVERGENT: %escaped =
+entry:
+ %lane = call i32 @llvm.nvvm.read.ptx.sreg.laneid()
+ br label %loop
+loop:
+ %iteration = phi i32 [0, %entry], [%next, %loop]
+ %uniform.ballot = call i32 @llvm.nvvm.vote.ballot.sync(i32 -1, i1 true)
+ %next = add i32 %iteration, 1
+ %done = icmp eq i32 %iteration, %lane
+ br i1 %done, label %exit, label %loop
+exit:
+ %escaped = add i32 %uniform.ballot, 0
+ %address = getelementptr i32, ptr %out, i32 %lane
+ store i32 %escaped, ptr %address
+ ret void
+}
+
+declare <4 x i32> @llvm.experimental.vp.reverse.v4i32(<4 x i32>, <4 x i1>, i32)
+
+declare i32 @llvm.nvvm.shfl.up.i32(i32, i32, i32)
+declare {i32, i1} @llvm.nvvm.shfl.up.i32p(i32, i32, i32)
+declare i1 @llvm.nvvm.vote.all(i1)
+declare i32 @llvm.nvvm.vote.ballot(i1)
+declare <vscale x 4 x i32> @llvm.vector.repeat.nxv4i32.v4i32(<4 x i32>)
+declare <4 x i1> @llvm.mask.beforefirst.v4i1(<4 x i1>)
+declare i16 @llvm.nvvm.ff.to.ue5m3x2.rn(float, float)
+declare <2 x half> @llvm.nvvm.ue5m3x2.to.f16x2.rn(i16)
+declare i1 @llvm.nvvm.mbarrier.check_layout.p0(ptr, i32)
>From bd4efd402995674cb153bc0402e3e125d3fa9d8e Mon Sep 17 00:00:00 2001
From: lezcano <lezcano-93 at hotmail.com>
Date: Thu, 15 Oct 2026 11:58:31 +0200
Subject: [PATCH 4/4] [NVPTX] Emit bra.uni for proven uniform branches
Use UniformityAnalysis during instruction selection to select uniform
conditional branches. Preserve the branch opcode through folding and
inversion, and drop uniformity when tail merging can combine different
predicates.
---
llvm/lib/Target/NVPTX/NVPTX.h | 2 +
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 22 ++-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp | 50 ++---
llvm/lib/Target/NVPTX/NVPTXInstrInfo.h | 2 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 19 +-
llvm/test/CodeGen/NVPTX/address-folder.ll | 8 +-
.../NVPTX/load-with-non-coherent-cache.ll | 4 +-
llvm/test/CodeGen/NVPTX/lower-byval-args.ll | 6 +-
.../NVPTX/uniform-branch-tail-merge.ll | 81 ++++++++
llvm/test/CodeGen/NVPTX/uniform-branches.ll | 175 ++++++++++++++++++
llvm/test/CodeGen/NVPTX/uniform-branches.mir | 95 ++++++++++
11 files changed, 427 insertions(+), 37 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/uniform-branch-tail-merge.ll
create mode 100644 llvm/test/CodeGen/NVPTX/uniform-branches.ll
create mode 100644 llvm/test/CodeGen/NVPTX/uniform-branches.mir
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 6a58b55acd29d4..af7882bec47742 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -222,6 +222,8 @@ class NVPTXForwardParamsPass
class NVPTXISelDAGToDAGPass : public SelectionDAGISelPass {
public:
NVPTXISelDAGToDAGPass(NVPTXTargetMachine &TM, CodeGenOptLevel OptLevel);
+ PreservedAnalyses run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM);
};
class NVPTXPeepholePass : public OptionalPassInfoMixin<NVPTXPeepholePass> {
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 848f819c5503e9..d0bafa90f100bb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -20,6 +20,7 @@
#include "llvm/ADT/MapVector.h"
#include "llvm/ADT/StringSwitch.h"
#include "llvm/ADT/Twine.h"
+#include "llvm/Analysis/UniformityAnalysis.h"
#include "llvm/Analysis/ValueTracking.h"
#include "llvm/CodeGen/ISDOpcodes.h"
#include "llvm/CodeGen/SelectionDAG.h"
@@ -33,6 +34,7 @@
#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/Metadata.h"
+#include "llvm/InitializePasses.h"
#include "llvm/Support/AtomicOrdering.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
@@ -173,6 +175,11 @@ class NVPTXDAGToDAGISelLegacy : public SelectionDAGISelLegacy {
static char ID;
explicit NVPTXDAGToDAGISelLegacy(NVPTXTargetMachine &tm,
CodeGenOptLevel OptLevel);
+
+ void getAnalysisUsage(AnalysisUsage &AU) const override {
+ AU.addRequired<UniformityInfoWrapperPass>();
+ SelectionDAGISelLegacy::getAnalysisUsage(AU);
+ }
};
} // end anonymous namespace
@@ -191,12 +198,25 @@ NVPTXDAGToDAGISelLegacy::NVPTXDAGToDAGISelLegacy(NVPTXTargetMachine &tm,
char NVPTXDAGToDAGISelLegacy::ID = 0;
-INITIALIZE_PASS(NVPTXDAGToDAGISelLegacy, DEBUG_TYPE, PASS_NAME, false, false)
+INITIALIZE_PASS_BEGIN(NVPTXDAGToDAGISelLegacy, DEBUG_TYPE, PASS_NAME, false,
+ false)
+INITIALIZE_PASS_DEPENDENCY(UniformityInfoWrapperPass)
+INITIALIZE_PASS_END(NVPTXDAGToDAGISelLegacy, DEBUG_TYPE, PASS_NAME, false,
+ false)
NVPTXISelDAGToDAGPass::NVPTXISelDAGToDAGPass(NVPTXTargetMachine &TM,
CodeGenOptLevel OptLevel)
: SelectionDAGISelPass(std::make_unique<NVPTXDAGToDAGISel>(TM, OptLevel)) {}
+PreservedAnalyses
+NVPTXISelDAGToDAGPass::run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM) {
+ auto &FAM = MFAM.getResult<FunctionAnalysisManagerMachineFunctionProxy>(MF)
+ .getManager();
+ FAM.getResult<UniformityInfoAnalysis>(MF.getFunction());
+ return SelectionDAGISelPass::run(MF, MFAM);
+}
+
NVPTXDAGToDAGISel::NVPTXDAGToDAGISel(NVPTXTargetMachine &tm,
CodeGenOptLevel OptLevel)
: SelectionDAGISel(tm, OptLevel), TM(tm) {}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
index e661f05ecd3351..6131ef0eabd0dd 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
@@ -100,11 +100,11 @@ bool NVPTXInstrInfo::analyzeBranch(MachineBasicBlock &MBB,
if (LastInst.getOpcode() == NVPTX::GOTO) {
TBB = LastInst.getOperand(0).getMBB();
return false;
- } else if (LastInst.getOpcode() == NVPTX::CBranch) {
+ } else if (LastInst.isConditionalBranch()) {
// Block ends with fall-through condbranch.
TBB = LastInst.getOperand(1).getMBB();
- Cond.push_back(LastInst.getOperand(0));
- Cond.push_back(LastInst.getOperand(2));
+ Cond.append({LastInst.getOperand(0), LastInst.getOperand(2),
+ MachineOperand::CreateImm(LastInst.getOpcode())});
return false;
}
// Otherwise, don't know what this is.
@@ -119,11 +119,11 @@ bool NVPTXInstrInfo::analyzeBranch(MachineBasicBlock &MBB,
return true;
// If the block ends with NVPTX::GOTO and NVPTX:CBranch, handle it.
- if (SecondLastInst.getOpcode() == NVPTX::CBranch &&
+ if (SecondLastInst.isConditionalBranch() &&
LastInst.getOpcode() == NVPTX::GOTO) {
TBB = SecondLastInst.getOperand(1).getMBB();
- Cond.push_back(SecondLastInst.getOperand(0));
- Cond.push_back(SecondLastInst.getOperand(2));
+ Cond.append({SecondLastInst.getOperand(0), SecondLastInst.getOperand(2),
+ MachineOperand::CreateImm(SecondLastInst.getOpcode())});
FBB = LastInst.getOperand(0).getMBB();
return false;
}
@@ -150,7 +150,7 @@ unsigned NVPTXInstrInfo::removeBranch(MachineBasicBlock &MBB,
if (I == MBB.begin())
return 0;
--I;
- if (I->getOpcode() != NVPTX::GOTO && I->getOpcode() != NVPTX::CBranch)
+ if (I->getOpcode() != NVPTX::GOTO && !I->isConditionalBranch())
return 0;
// Remove the branch.
@@ -161,7 +161,7 @@ unsigned NVPTXInstrInfo::removeBranch(MachineBasicBlock &MBB,
if (I == MBB.begin())
return 1;
--I;
- if (I->getOpcode() != NVPTX::CBranch)
+ if (!I->isConditionalBranch())
return 1;
// Remove the branch.
@@ -179,30 +179,34 @@ unsigned NVPTXInstrInfo::insertBranch(MachineBasicBlock &MBB,
// Shouldn't be a fall through.
assert(TBB && "insertBranch must not be told to insert a fallthrough");
- assert((Cond.size() == 2 || Cond.size() == 0) &&
- "NVPTX branch conditions have two components!");
-
- // One-way branch.
- if (!FBB) {
- if (Cond.empty()) // Unconditional branch
- BuildMI(&MBB, DL, get(NVPTX::GOTO)).addMBB(TBB);
- else // Conditional branch
- BuildMI(&MBB, DL, get(NVPTX::CBranch))
- .add(Cond[0])
- .addMBB(TBB)
- .add(Cond[1]);
+ assert((Cond.size() == 3 || Cond.empty()) &&
+ "NVPTX branch conditions have three components!");
+
+ if (Cond.empty()) {
+ BuildMI(&MBB, DL, get(NVPTX::GOTO)).addMBB(TBB);
return 1;
}
- // Two-way Conditional Branch.
- BuildMI(&MBB, DL, get(NVPTX::CBranch)).add(Cond[0]).addMBB(TBB).add(Cond[1]);
+ unsigned Opcode = Cond[2].getImm();
+ BuildMI(&MBB, DL, get(Opcode)).add(Cond[0]).addMBB(TBB).add(Cond[1]);
+ if (!FBB)
+ return 1;
BuildMI(&MBB, DL, get(NVPTX::GOTO)).addMBB(FBB);
return 2;
}
+void NVPTXInstrInfo::ReplaceTailWithBranchTo(MachineBasicBlock::iterator Tail,
+ MachineBasicBlock *NewDest) const {
+ TargetInstrInfo::ReplaceTailWithBranchTo(Tail, NewDest);
+ // Coalescing can give different uniform predicates the same register.
+ for (MachineInstr &MI : NewDest->terminators())
+ if (MI.getOpcode() == NVPTX::CBranchUni)
+ MI.setDesc(get(NVPTX::CBranch));
+}
+
bool NVPTXInstrInfo::reverseBranchCondition(
SmallVectorImpl<MachineOperand> &Cond) const {
- assert(Cond.size() == 2 && "Invalid NVPTX branch condition!");
+ assert(Cond.size() == 3 && "Invalid NVPTX branch condition!");
Cond[1].setImm(!Cond[1].getImm());
return false;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
index 957fcf5cae69a6..9d06bd0ce60448 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
@@ -69,6 +69,8 @@ class NVPTXInstrInfo : public NVPTXGenInstrInfo {
int *BytesAdded = nullptr) const override;
bool
reverseBranchCondition(SmallVectorImpl<MachineOperand> &Cond) const override;
+ void ReplaceTailWithBranchTo(MachineBasicBlock::iterator Tail,
+ MachineBasicBlock *NewDest) const override;
bool findCommutedOpIndices(const MachineInstr &MI, unsigned &SrcOpIdx1,
unsigned &SrcOpIdx2) const override;
MachineInstr *commuteInstructionImpl(MachineInstr &MI, bool NewMI,
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 8f1e48f29cad3c..998706374eb58e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2551,15 +2551,26 @@ defm : CVT_ROUND<frint, CvtRNI, CvtRNI_FTZ>;
// Control-flow
//-----------------------------------
+def uniform_brcond : PatFrag<(ops node:$cond, node:$target),
+ (brcond node:$cond, node:$target), [{
+ const BasicBlock *BB = CurDAG->getBasicBlock();
+ const Instruction *Term = BB ? BB->getTerminator() : nullptr;
+ // Exclude synthesized branches, such as stack-protector checks at returns.
+ return isa_and_nonnull<CondBrInst, SwitchInst>(Term) &&
+ CurDAG->getUniformityInfo()->isUniformTerminator(Term);
+}]>;
+
let isTerminator=1 in {
let isReturn=1, isBarrier=1 in
def Return : BasicNVPTXInst<(outs), (ins), "ret", [(retglue)]>;
let isBranch=1 in {
- def CBranch : NVPTXInst<(outs),
- (ins B1:$a, brtarget:$target, BranchFlag:$negate),
- "@${negate}$a bra \t$target;",
- [(brcond i1:$a, bb:$target)]>;
+ class ConditionalBranch<string mnemonic, list<dag> pattern>
+ : NVPTXInst<(outs), (ins B1:$a, brtarget:$target, BranchFlag:$negate),
+ "@${negate}$a " # mnemonic # " \t$target;", pattern>;
+ def CBranch : ConditionalBranch<"bra", [(brcond i1:$a, bb:$target)]>;
+ def CBranchUni : ConditionalBranch<"bra.uni",
+ [(uniform_brcond i1:$a, bb:$target)]>;
let isBarrier=1 in
def GOTO : BasicNVPTXInst<(outs), (ins brtarget:$target),
diff --git a/llvm/test/CodeGen/NVPTX/address-folder.ll b/llvm/test/CodeGen/NVPTX/address-folder.ll
index 813fd2c4396518..4ddccfcccb404f 100644
--- a/llvm/test/CodeGen/NVPTX/address-folder.ll
+++ b/llvm/test/CodeGen/NVPTX/address-folder.ll
@@ -16,7 +16,7 @@ define ptx_kernel void @test_kernel_cross_block(ptr byval([2 x i64]) align 8 %r,
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd2, [test_kernel_cross_block_param_1];
; CHECK-NEXT: setp.lt.s64 %p1, %rd2, 1;
-; CHECK-NEXT: @%p1 bra $L__BB0_2;
+; CHECK-NEXT: @%p1 bra.uni $L__BB0_2;
; CHECK-NEXT: // %bb.1: // %body
; CHECK-NEXT: ld.param.b64 %rd1, [test_kernel_cross_block_param_2];
; CHECK-NEXT: ld.param.b64 %rd3, [test_kernel_cross_block_param_0];
@@ -36,7 +36,7 @@ define ptx_kernel void @test_kernel_cross_block(ptr byval([2 x i64]) align 8 %r,
; O0-NEXT: ld.param.b64 %rd2, [test_kernel_cross_block_param_1];
; O0-NEXT: mov.b64 %rd1, test_kernel_cross_block_param_0;
; O0-NEXT: setp.lt.s64 %p1, %rd2, 1;
-; O0-NEXT: @%p1 bra $L__BB0_2;
+; O0-NEXT: @%p1 bra.uni $L__BB0_2;
; O0-NEXT: bra.uni $L__BB0_1;
; O0-NEXT: $L__BB0_1: // %body
; O0-NEXT: ld.param.b64 %rd4, [%rd1];
@@ -70,7 +70,7 @@ define ptx_kernel void @test_kernel_var_offset(ptr byval([8 x i64]) align 8 %r,
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd4, [test_kernel_var_offset_param_1];
; CHECK-NEXT: setp.lt.s64 %p1, %rd4, 1;
-; CHECK-NEXT: @%p1 bra $L__BB1_2;
+; CHECK-NEXT: @%p1 bra.uni $L__BB1_2;
; CHECK-NEXT: // %bb.1: // %body
; CHECK-NEXT: ld.param.b64 %rd2, [test_kernel_var_offset_param_2];
; CHECK-NEXT: mov.b64 %rd3, test_kernel_var_offset_param_0;
@@ -90,7 +90,7 @@ define ptx_kernel void @test_kernel_var_offset(ptr byval([8 x i64]) align 8 %r,
; O0-NEXT: ld.param.b64 %rd2, [test_kernel_var_offset_param_1];
; O0-NEXT: mov.b64 %rd1, test_kernel_var_offset_param_0;
; O0-NEXT: setp.lt.s64 %p1, %rd2, 1;
-; O0-NEXT: @%p1 bra $L__BB1_2;
+; O0-NEXT: @%p1 bra.uni $L__BB1_2;
; O0-NEXT: bra.uni $L__BB1_1;
; O0-NEXT: $L__BB1_1: // %body
; O0-NEXT: add.s64 %rd4, %rd1, %rd2;
diff --git a/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll b/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
index 7553c727b09c57..f926898157fd95 100644
--- a/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
+++ b/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
@@ -614,7 +614,7 @@ define ptx_kernel void @foo19(ptr noalias readonly %from, ptr %to, i32 %n) {
; SM20-NEXT: add.s64 %rd4, %rd4, 4;
; SM20-NEXT: add.s32 %r2, %r2, -1;
; SM20-NEXT: setp.ne.b32 %p1, %r2, 0;
-; SM20-NEXT: @%p1 bra $L__BB18_1;
+; SM20-NEXT: @%p1 bra.uni $L__BB18_1;
; SM20-NEXT: // %bb.2: // %exit
; SM20-NEXT: st.global.b32 [%rd1], %r3;
; SM20-NEXT: ret;
@@ -639,7 +639,7 @@ define ptx_kernel void @foo19(ptr noalias readonly %from, ptr %to, i32 %n) {
; SM35-NEXT: add.s64 %rd4, %rd4, 4;
; SM35-NEXT: add.s32 %r2, %r2, -1;
; SM35-NEXT: setp.ne.b32 %p1, %r2, 0;
-; SM35-NEXT: @%p1 bra $L__BB18_1;
+; SM35-NEXT: @%p1 bra.uni $L__BB18_1;
; SM35-NEXT: // %bb.2: // %exit
; SM35-NEXT: st.global.b32 [%rd1], %r3;
; SM35-NEXT: ret;
diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
index 8e44a5c5c99311..91feaee86e909a 100644
--- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
@@ -748,7 +748,7 @@ define ptx_kernel void @test_phi(ptr byval(%struct.S) align 4 %input1, ptr byval
; PTX_60-NEXT: ld.param.b64 %rd2, [test_phi_param_2];
; PTX_60-NEXT: cvta.to.global.u64 %rd1, %rd2;
; PTX_60-NEXT: ld.param.b32 %r1, [test_phi_param_0];
-; PTX_60-NEXT: @%p1 bra $L__BB13_2;
+; PTX_60-NEXT: @%p1 bra.uni $L__BB13_2;
; PTX_60-NEXT: // %bb.1: // %second
; PTX_60-NEXT: ld.param.b32 %r1, [test_phi_param_1+4];
; PTX_60-NEXT: $L__BB13_2: // %merge
@@ -769,7 +769,7 @@ define ptx_kernel void @test_phi(ptr byval(%struct.S) align 4 %input1, ptr byval
; PTX_70-NEXT: mov.b64 %rd4, test_phi_param_0;
; PTX_70-NEXT: ld.param.b64 %rd3, [test_phi_param_2];
; PTX_70-NEXT: cvta.to.global.u64 %rd1, %rd3;
-; PTX_70-NEXT: @%p1 bra $L__BB13_2;
+; PTX_70-NEXT: @%p1 bra.uni $L__BB13_2;
; PTX_70-NEXT: // %bb.1: // %second
; PTX_70-NEXT: mov.b64 %rd2, test_phi_param_1;
; PTX_70-NEXT: add.s64 %rd4, %rd2, 4;
@@ -854,7 +854,7 @@ define ptx_kernel void @test_phi_write(ptr byval(%struct.S) align 4 %input1, ptr
; PTX-NEXT: ld.param.b32 %r2, [test_phi_write_param_0];
; PTX-NEXT: st.local.b32 [%SPL+4], %r2;
; PTX-NEXT: add.u64 %rd2, %SPL, 4;
-; PTX-NEXT: @%p1 bra $L__BB14_2;
+; PTX-NEXT: @%p1 bra.uni $L__BB14_2;
; PTX-NEXT: // %bb.1: // %second
; PTX-NEXT: add.u64 %rd2, %SPL, 0;
; PTX-NEXT: $L__BB14_2: // %merge
diff --git a/llvm/test/CodeGen/NVPTX/uniform-branch-tail-merge.ll b/llvm/test/CodeGen/NVPTX/uniform-branch-tail-merge.ll
new file mode 100644
index 00000000000000..abaa9d2f3906c0
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/uniform-branch-tail-merge.ll
@@ -0,0 +1,81 @@
+; RUN: llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=CHECK,DEFAULT
+; RUN: llc -mtriple=nvptx64 -mcpu=sm_90 -O2 -enable-tail-merge -tail-merge-size=1 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=CHECK,MERGE
+
+; The side-exit PHI coalesces two different uniform predicates. With test=true,
+; a=0 and b=1, they differ between the half-warps after their branches merge.
+; CHECK-LABEL: .visible .entry separate_uniform_predicates(
+; DEFAULT: // left [[P:%p[0-9]+]]
+; DEFAULT: @{{!?}}[[P]] bra.uni
+; DEFAULT: // right [[P]]
+; DEFAULT: @{{!?}}[[P]] bra.uni
+; MERGE: // left [[P:%p[0-9]+]]
+; MERGE: // right [[P]]
+; MERGE: @{{!?}}[[P]] bra{{[ \t]}}
+
+define ptx_kernel void @separate_uniform_predicates(i1 %test, i32 %a, i32 %b) {
+entry:
+ %lane = call i32 @llvm.nvvm.read.ptx.sreg.laneid()
+ %split = icmp ult i32 %lane, 16
+ br i1 %split, label %left, label %right
+left:
+ %p = icmp eq i32 %a, 0
+ call void asm sideeffect "// left $0", "b"(i1 %p)
+ br i1 %test, label %lefttest, label %sideexit
+right:
+ %q = icmp eq i32 %b, 0
+ call void asm sideeffect "// right $0", "b"(i1 %q)
+ br i1 %test, label %righttest, label %sideexit
+lefttest:
+ br i1 %p, label %true, label %false
+righttest:
+ br i1 %q, label %true, label %false
+sideexit:
+ %r = phi i1 [%p, %left], [%q, %right]
+ call void asm sideeffect "// result $0", "b"(i1 %r)
+ ret void
+true:
+ call void asm sideeffect "// true", ""()
+ ret void
+false:
+ call void asm sideeffect "// false", ""()
+ ret void
+}
+
+; The adjacent fold keeps an existing path to the retained branch. Its predicate
+; stays uniform within each loop iteration.
+; CHECK-LABEL: .visible .entry adjacent_loop(
+; CHECK: setp.eq.b32 [[ITER_PRED:%p[0-9]+]], {{%r[0-9]+}}, 0;
+; CHECK: @{{!?}}[[ITER_PRED]] bra.uni
+define ptx_kernel void @adjacent_loop(ptr addrspace(1) %out, i1 %flip, i32 %limit) {
+entry:
+ %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %side = icmp ult i32 %tid, 16
+ %ptr = getelementptr i32, ptr addrspace(1) %out, i32 %tid
+ br label %header
+header:
+ %i = phi i32 [ 0, %entry ], [ %next, %latch ]
+ %p = icmp eq i32 %i, 0
+ %q = xor i1 %p, %flip
+ br i1 %side, label %a, label %c
+a:
+ br i1 %p, label %true, label %b
+b:
+ br i1 %p, label %true, label %false
+false:
+ store volatile i32 %i, ptr addrspace(1) %ptr
+ %stop = icmp sge i32 %i, %limit
+ br i1 %stop, label %exit, label %latch
+c:
+ br i1 %q, label %latch, label %b
+latch:
+ %next = add i32 %i, 1
+ br label %header
+true:
+ store volatile i32 99, ptr addrspace(1) %ptr
+ ret void
+exit:
+ ret void
+}
+
+declare i32 @llvm.nvvm.read.ptx.sreg.laneid()
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
diff --git a/llvm/test/CodeGen/NVPTX/uniform-branches.ll b/llvm/test/CodeGen/NVPTX/uniform-branches.ll
new file mode 100644
index 00000000000000..f8f5b4030ef5f0
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/uniform-branches.ll
@@ -0,0 +1,175 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -O0 -verify-machineinstrs | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -O2 -verify-machineinstrs | FileCheck %s --check-prefixes=CHECK,O2
+; RUN: llc -enable-new-pm < %s -mtriple=nvptx64 -mcpu=sm_90 -O0 -verify-machineinstrs | FileCheck %s
+; RUN: llc -enable-new-pm < %s -mtriple=nvptx64 -mcpu=sm_90 -O2 -verify-machineinstrs | FileCheck %s --check-prefixes=CHECK,O2
+
+target triple = "nvptx64-nvidia-cuda"
+
+define ptx_kernel void @uniform_scalars(i32 %a, i32 %b) {
+; CHECK-LABEL: .entry uniform_scalars(
+; CHECK: @{{!?}}%p{{[0-9]+}} bra.uni
+ %block = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+ %min = call i32 @llvm.umin.i32(i32 %a, i32 %block)
+ %c0 = icmp ult i32 %min, 7
+ %c1 = icmp ult i32 %b, 3
+ %both = and i1 %c0, %c1
+ %c2 = icmp eq i32 %a, %b
+ %cond = or i1 %both, %c2
+ br i1 %cond, label %then, label %exit
+then:
+ call void asm sideeffect "// taken", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+define void @device_argument(i32 %a) {
+; CHECK-LABEL: .func device_argument(
+; CHECK: @{{!?}}%p{{[0-9]+}} bra $L
+ %cond = icmp eq i32 %a, 0
+ br i1 %cond, label %then, label %exit
+then:
+ call void asm sideeffect "// taken", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+define ptx_kernel void @warp_dispatch() "nvvm.reqntid"="128" {
+; CHECK-LABEL: .entry warp_dispatch(
+; CHECK: @{{!?}}%p{{[0-9]+}} bra.uni
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %warp = lshr i32 %thread, 5
+ %cond = icmp ult i32 %warp, 2
+ br i1 %cond, label %then, label %exit
+then:
+ call void asm sideeffect "// default partition", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+define ptx_kernel void @uniform_loop(i32 %n) {
+; CHECK-LABEL: .entry uniform_loop(
+; CHECK: @{{!?}}%p{{[0-9]+}} bra.uni
+entry:
+ br label %loop
+loop:
+ %i = phi i32 [ 0, %entry ], [ %next, %loop ]
+ call void asm sideeffect "// iteration $0", "r"(i32 %i)
+ %next = add i32 %i, 1
+ %cond = icmp ult i32 %next, %n
+ br i1 %cond, label %loop, label %exit
+exit:
+ ret void
+}
+
+define ptx_kernel void @broadcast_condition() {
+; CHECK-LABEL: .entry broadcast_condition(
+; CHECK: shfl.sync.idx.b32
+; CHECK: @{{!?}}%p{{[0-9]+}} bra.uni
+ %thread = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ %broadcast = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 -1, i32 %thread, i32 0, i32 31)
+ %cond = icmp ult i32 %broadcast, 64
+ br i1 %cond, label %then, label %exit
+then:
+ call void asm sideeffect "// taken", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+; The induction variable is uniform inside the loop, but threads can observe
+; different iterations at the exit. Both conditional branches are divergent.
+define ptx_kernel void @temporal_divergence() {
+; CHECK-LABEL: .entry temporal_divergence(
+; CHECK: @{{!?}}%p{{[0-9]+}} bra $L
+; CHECK: @{{!?}}%p{{[0-9]+}} bra $L
+entry:
+ %lane = call i32 @llvm.nvvm.read.ptx.sreg.laneid()
+ br label %loop
+loop:
+ %i = phi i32 [ 0, %entry ], [ %next, %loop ]
+ call void asm sideeffect "// iteration $0", "r"(i32 %i)
+ %next = add i32 %i, 1
+ %done = icmp uge i32 %next, %lane
+ br i1 %done, label %test, label %loop
+test:
+ %cond = icmp eq i32 %i, 8
+ br i1 %cond, label %then, label %exit
+then:
+ call void asm sideeffect "// taken", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+define ptx_kernel void @uniform_switch(i32 %a) {
+; CHECK-LABEL: .entry uniform_switch(
+; CHECK-COUNT-2: @{{!?}}%p{{[0-9]+}} bra.uni
+ switch i32 %a, label %exit [ i32 4, label %four
+ i32 32, label %thirtytwo ]
+four:
+ call void asm sideeffect "// four", ""()
+ br label %exit
+thirtytwo:
+ call void asm sideeffect "// thirtytwo", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+; Opaque asm results are divergent, but independent branches can be uniform.
+define ptx_kernel void @opaque_control_flow(i32 %a) {
+; CHECK-LABEL: .entry opaque_control_flow(
+; CHECK: @{{!?}}%p{{[0-9]+}} bra.uni
+; CHECK: @{{!?}}%p{{[0-9]+}} bra $L
+ %uniform = icmp eq i32 %a, 0
+ br i1 %uniform, label %uniform_then, label %call
+uniform_then:
+ call void asm sideeffect "// uniform", ""()
+ br label %call
+call:
+ %value = callbr i32 asm sideeffect "mov.u32 $0, %laneid; // $1", "=r,!i"() to label %test [label %exit]
+test:
+ %cond = icmp eq i32 %a, %value
+ br i1 %cond, label %then, label %exit
+then:
+ call void asm sideeffect "// taken", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+; Wide comparison legalization must retain the branch's IR proof.
+define ptx_kernel void @wide_comparison(i128 %a, i128 %b) {
+; CHECK-LABEL: .entry wide_comparison(
+; CHECK: @{{!?}}%p{{[0-9]+}} bra.uni
+ %cond = icmp ult i128 %a, %b
+ br i1 %cond, label %then, label %exit
+then:
+ call void asm sideeffect "// taken", ""()
+ br label %exit
+exit:
+ ret void
+}
+
+; A uniform branch has the same memory effects as an ordinary branch.
+define ptx_kernel void @sink_load(ptr addrspace(3) %input, ptr addrspace(1) %output, i1 %cond) {
+; O2-LABEL: .entry sink_load(
+; O2: @{{!?}}%p{{[0-9]+}} bra.uni
+; O2: ld.shared.b32
+ %value = load i32, ptr addrspace(3) %input
+ br i1 %cond, label %store, label %exit
+store:
+ store i32 %value, ptr addrspace(1) %output
+ br label %exit
+exit:
+ ret void
+}
+
+declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+declare i32 @llvm.nvvm.read.ptx.sreg.laneid()
+declare i32 @llvm.umin.i32(i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
diff --git a/llvm/test/CodeGen/NVPTX/uniform-branches.mir b/llvm/test/CodeGen/NVPTX/uniform-branches.mir
new file mode 100644
index 00000000000000..9d593f66201b7a
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/uniform-branches.mir
@@ -0,0 +1,95 @@
+# RUN: llc -mtriple=nvptx64 -mcpu=sm_90 -run-pass=branch-folder -verify-machineinstrs -o - %s | FileCheck %s --check-prefix=FOLD
+# RUN: llc -mtriple=nvptx64 -mcpu=sm_90 -run-pass=machine-cse -verify-machineinstrs -o - %s | FileCheck %s --check-prefix=CSE
+# RUN: llc -mtriple=nvptx64 -mcpu=sm_90 -run-pass=branch-folder -enable-tail-merge -tail-merge-size=1 -verify-machineinstrs -o - %s | FileCheck %s --check-prefix=MERGE
+
+---
+name: uniform_loop
+tracksRegLiveness: true
+body: |
+ ; FOLD-LABEL: name: uniform_loop
+ ; FOLD: CBranchUni {{%[0-9]+}}, %bb.{{[0-9]+}}, 1
+ bb.0:
+ successors: %bb.1
+ GOTO %bb.1
+
+ bb.1:
+ successors: %bb.1, %bb.2
+ %0:b1 = IMPLICIT_DEF
+ CBranchUni %0, %bb.2, 0
+ GOTO %bb.1
+
+ bb.2:
+ Return
+...
+
+---
+name: compare_inversion
+tracksRegLiveness: true
+body: |
+ ; CSE-LABEL: name: compare_inversion
+ ; CSE: [[CMP:%[0-9]+]]:b1 = SETP_i32
+ ; CSE: CBranch [[CMP]], %bb.3, 0
+ ; CSE-NOT: SETP_i32
+ ; CSE: CBranchUni [[CMP]], %bb.3, 1
+ bb.0:
+ successors: %bb.1, %bb.3
+ %0:b32 = IMPLICIT_DEF
+ %1:b32 = IMPLICIT_DEF
+ %2:b1 = SETP_i32 %0, %1, 0
+ CBranch %2, %bb.3, 0
+ GOTO %bb.1
+
+ bb.1:
+ successors: %bb.2, %bb.3
+ %3:b1 = SETP_i32 %0, %1, 1
+ CBranchUni %3, %bb.3, 0
+ GOTO %bb.2
+
+ bb.2:
+ Return
+
+ bb.3:
+ Return
+...
+
+# Each branch can be uniform within its own incoming group of threads without
+# being uniform after the groups join. Merge the tails and drop uniformity.
+---
+name: separate_groups
+tracksRegLiveness: true
+body: |
+ ; MERGE-LABEL: name: separate_groups
+ ; MERGE: [[SPLIT:%[0-9]+]]:b1 = IMPLICIT_DEF
+ ; MERGE-NEXT: [[COND:%[0-9]+]]:b1 = IMPLICIT_DEF
+ ; MERGE-NOT: CBranchUni
+ ; MERGE: CBranch [[COND]],
+ ; MERGE-NOT: CBranchUni
+ ; MERGE-NOT: CBranch [[COND]],
+ ; MERGE: ...
+ bb.0:
+ successors: %bb.1, %bb.2
+ %0:b1 = IMPLICIT_DEF
+ %1:b1 = IMPLICIT_DEF
+ CBranch %0, %bb.2, 0
+ GOTO %bb.1
+
+ bb.1:
+ successors: %bb.3, %bb.4
+ INLINEASM &"// left", 1
+ CBranchUni %1, %bb.3, 0
+ GOTO %bb.4
+
+ bb.2:
+ successors: %bb.3, %bb.4
+ INLINEASM &"// right", 1
+ CBranchUni %1, %bb.3, 0
+ GOTO %bb.4
+
+ bb.3:
+ INLINEASM &"// true", 1
+ Return
+
+ bb.4:
+ INLINEASM &"// false", 1
+ Return
+...
More information about the cfe-commits
mailing list