[llvm] [IR] Add new function attribute nocreateundeforpoison (PR #164809)
Jay Foad via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 23 07:10:46 PDT 2025
https://github.com/jayfoad updated https://github.com/llvm/llvm-project/pull/164809
>From ac95b57f3ae2abee2af43d198c972993dcfdcfaa Mon Sep 17 00:00:00 2001
From: Jay Foad <jay.foad at amd.com>
Date: Thu, 23 Oct 2025 12:44:54 +0100
Subject: [PATCH 1/2] [IR] Add new function attribute nocreateundeforpoison
Also add a corresponding intrinsic property that can be used to mark
intrinsics that do not introduce poison, for example simple arithmetic
intrinsics that propagate poison just like a simple arithmetic
instruction.
As a smoke test this patch adds the new property to
llvm.amdgcn.fmul.legacy.
---
llvm/docs/LangRef.rst | 4 ++++
llvm/include/llvm/Bitcode/LLVMBitCodes.h | 1 +
llvm/include/llvm/IR/Attributes.td | 5 +++++
llvm/include/llvm/IR/Intrinsics.td | 4 ++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 2 +-
llvm/lib/Analysis/ValueTracking.cpp | 9 ++++++---
llvm/lib/Bitcode/Reader/BitcodeReader.cpp | 2 ++
llvm/lib/Bitcode/Writer/BitcodeWriter.cpp | 2 ++
llvm/lib/Transforms/Utils/CodeExtractor.cpp | 1 +
llvm/test/Bitcode/attributes.ll | 6 ++++++
.../Transforms/InstCombine/AMDGPU/fmul_legacy.ll | 12 ++++++++++++
llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp | 2 ++
llvm/utils/TableGen/Basic/CodeGenIntrinsics.h | 3 +++
llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp | 8 ++++++--
14 files changed, 55 insertions(+), 6 deletions(-)
diff --git a/llvm/docs/LangRef.rst b/llvm/docs/LangRef.rst
index 1c6823be44dcb..779597ab61e5a 100644
--- a/llvm/docs/LangRef.rst
+++ b/llvm/docs/LangRef.rst
@@ -2741,6 +2741,10 @@ For example:
``"nooutline"``
This attribute indicates that outlining passes should not modify the
function.
+``nocreateundeforpoison``
+ This attribute indicates that the result of the function will not be undef
+ or poison if all arguments are not undef and not poison. Otherwise, it is
+ undefined behavior.
Call Site Attributes
----------------------
diff --git a/llvm/include/llvm/Bitcode/LLVMBitCodes.h b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
index 464f475098ec5..b0c5beae631ce 100644
--- a/llvm/include/llvm/Bitcode/LLVMBitCodes.h
+++ b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
@@ -801,6 +801,7 @@ enum AttributeKindCodes {
ATTR_KIND_CAPTURES = 102,
ATTR_KIND_DEAD_ON_RETURN = 103,
ATTR_KIND_SANITIZE_ALLOC_TOKEN = 104,
+ ATTR_KIND_NO_CREATE_UNDEF_OR_POISON = 105,
};
enum ComdatSelectionKindCodes {
diff --git a/llvm/include/llvm/IR/Attributes.td b/llvm/include/llvm/IR/Attributes.td
index 8ce2b1bea8fac..c086a39616249 100644
--- a/llvm/include/llvm/IR/Attributes.td
+++ b/llvm/include/llvm/IR/Attributes.td
@@ -183,6 +183,11 @@ def NoCallback : EnumAttr<"nocallback", IntersectAnd, [FnAttr]>;
/// Specify how the pointer may be captured.
def Captures : IntAttr<"captures", IntersectCustom, [ParamAttr]>;
+/// Result will not be undef or poison if all arguments are not undef and not
+/// poison.
+def NoCreateUndefOrPoison
+ : EnumAttr<"nocreateundeforpoison", IntersectAnd, [FnAttr]>;
+
/// Function is not a source of divergence.
def NoDivergenceSource : EnumAttr<"nodivergencesource", IntersectAnd, [FnAttr]>;
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index e6cce9a4eea1d..e5fb0a9afad3c 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -186,6 +186,10 @@ def IntrSpeculatable : IntrinsicProperty;
// defined by the hasSideEffects property of the TableGen Instruction class.
def IntrHasSideEffects : IntrinsicProperty;
+// Result will not be undef or poison if all arguments are not undef and not
+// poison.
+def IntrNoCreateUndefOrPoison : IntrinsicProperty;
+
//===----------------------------------------------------------------------===//
// IIT constants and utils
//===----------------------------------------------------------------------===//
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 8e35109061792..cc375a35ca822 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -449,7 +449,7 @@ def int_amdgcn_log_clamp : DefaultAttrsIntrinsic<
def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">,
DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
- [IntrNoMem, IntrSpeculatable, Commutative]
+ [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
// Fused single-precision multiply-add with legacy behaviour for the multiply,
diff --git a/llvm/lib/Analysis/ValueTracking.cpp b/llvm/lib/Analysis/ValueTracking.cpp
index 0a72076f51824..c293b93f528b5 100644
--- a/llvm/lib/Analysis/ValueTracking.cpp
+++ b/llvm/lib/Analysis/ValueTracking.cpp
@@ -7446,8 +7446,10 @@ static bool canCreateUndefOrPoison(const Operator *Op, UndefPoisonKind Kind,
return false;
case Intrinsic::sshl_sat:
case Intrinsic::ushl_sat:
- return includesPoison(Kind) &&
- !shiftAmountKnownInRange(II->getArgOperand(1));
+ if (!includesPoison(Kind) ||
+ shiftAmountKnownInRange(II->getArgOperand(1)))
+ return false;
+ break;
case Intrinsic::fma:
case Intrinsic::fmuladd:
case Intrinsic::sqrt:
@@ -7496,7 +7498,8 @@ static bool canCreateUndefOrPoison(const Operator *Op, UndefPoisonKind Kind,
case Instruction::CallBr:
case Instruction::Invoke: {
const auto *CB = cast<CallBase>(Op);
- return !CB->hasRetAttr(Attribute::NoUndef);
+ return !CB->hasRetAttr(Attribute::NoUndef) &&
+ !CB->hasFnAttr(Attribute::NoCreateUndefOrPoison);
}
case Instruction::InsertElement:
case Instruction::ExtractElement: {
diff --git a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
index 466dcb02696f4..5a5fce6315161 100644
--- a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
+++ b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
@@ -2257,6 +2257,8 @@ static Attribute::AttrKind getAttrFromCode(uint64_t Code) {
return Attribute::Captures;
case bitc::ATTR_KIND_DEAD_ON_RETURN:
return Attribute::DeadOnReturn;
+ case bitc::ATTR_KIND_NO_CREATE_UNDEF_OR_POISON:
+ return Attribute::NoCreateUndefOrPoison;
}
}
diff --git a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
index 61aa7c2f5af53..3141c5a49592e 100644
--- a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
+++ b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
@@ -956,6 +956,8 @@ static uint64_t getAttrKindEncoding(Attribute::AttrKind Kind) {
return bitc::ATTR_KIND_CAPTURES;
case Attribute::DeadOnReturn:
return bitc::ATTR_KIND_DEAD_ON_RETURN;
+ case Attribute::NoCreateUndefOrPoison:
+ return bitc::ATTR_KIND_NO_CREATE_UNDEF_OR_POISON;
case Attribute::EndAttrKinds:
llvm_unreachable("Can not encode end-attribute kinds marker.");
case Attribute::None:
diff --git a/llvm/lib/Transforms/Utils/CodeExtractor.cpp b/llvm/lib/Transforms/Utils/CodeExtractor.cpp
index 5ba6f95f5fae8..608661583c3db 100644
--- a/llvm/lib/Transforms/Utils/CodeExtractor.cpp
+++ b/llvm/lib/Transforms/Utils/CodeExtractor.cpp
@@ -933,6 +933,7 @@ Function *CodeExtractor::constructFunctionDeclaration(
case Attribute::CoroDestroyOnlyWhenComplete:
case Attribute::CoroElideSafe:
case Attribute::NoDivergenceSource:
+ case Attribute::NoCreateUndefOrPoison:
continue;
// Those attributes should be safe to propagate to the extracted function.
case Attribute::AlwaysInline:
diff --git a/llvm/test/Bitcode/attributes.ll b/llvm/test/Bitcode/attributes.ll
index aef7810fe2c3b..107a98aebeeb8 100644
--- a/llvm/test/Bitcode/attributes.ll
+++ b/llvm/test/Bitcode/attributes.ll
@@ -521,6 +521,11 @@ define void @f_sanitize_alloc_token() sanitize_alloc_token {
ret void;
}
+; CHECK: define void @f_no_create_undef_or_poison() #56
+define void @f_no_create_undef_or_poison() nocreateundeforpoison {
+ ret void;
+}
+
; CHECK: define void @f87() [[FNRETTHUNKEXTERN:#[0-9]+]]
define void @f87() fn_ret_thunk_extern { ret void }
@@ -633,6 +638,7 @@ define void @dead_on_return(ptr dead_on_return %p) {
; CHECK: attributes #53 = { sanitize_realtime }
; CHECK: attributes #54 = { sanitize_realtime_blocking }
; CHECK: attributes #55 = { sanitize_alloc_token }
+; CHECK: attributes #56 = { nocreateundeforpoison }
; CHECK: attributes [[FNRETTHUNKEXTERN]] = { fn_ret_thunk_extern }
; CHECK: attributes [[SKIPPROFILE]] = { skipprofile }
; CHECK: attributes [[OPTDEBUG]] = { optdebug }
diff --git a/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll b/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll
index 899b611b929d1..dd5f6cc9a81bd 100644
--- a/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll
+++ b/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll
@@ -91,6 +91,18 @@ define float @test_var_poison(float %x) {
ret float %call
}
+define float @test_freeze(float %x, float %y) {
+; CHECK-LABEL: @test_freeze(
+; CHECK-NEXT: [[FR:%.*]] = freeze float [[CALL:%.*]]
+; CHECK-NEXT: [[Y_FR:%.*]] = freeze float [[Y:%.*]]
+; CHECK-NEXT: [[CALL1:%.*]] = call float @llvm.amdgcn.fmul.legacy(float [[FR]], float [[Y_FR]])
+; CHECK-NEXT: ret float [[CALL1]]
+;
+ %call = call float @llvm.amdgcn.fmul.legacy(float %x, float %y)
+ %fr = freeze float %call
+ ret float %fr
+}
+
declare float @llvm.amdgcn.fmul.legacy(float, float)
declare float @llvm.fabs.f32(float)
declare void @llvm.assume(i1 noundef)
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index be7537c83da3a..cace48950518d 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -401,6 +401,8 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
hasSideEffects = true;
else if (R->getName() == "IntrStrictFP")
isStrictFP = true;
+ else if (R->getName() == "IntrNoCreateUndefOrPoison")
+ isNoCreateUndefOrPoison = true;
else if (R->isSubClassOf("NoCapture")) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
addArgAttribute(ArgNo, NoCapture);
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index 2e86149514f46..15e803c4feba1 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -114,6 +114,9 @@ struct CodeGenIntrinsic {
// True if the intrinsic is marked as strictfp.
bool isStrictFP = false;
+ // True if the intrinsic is marked as IntrNoCreateUndefOrPoison.
+ bool isNoCreateUndefOrPoison = false;
+
enum ArgAttrKind {
NoCapture,
NoAlias,
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index 75dffb18fca5a..452d2b08f25c3 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -421,7 +421,8 @@ static bool compareFnAttributes(const CodeGenIntrinsic *L,
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->hasSideEffects, I->isStrictFP,
+ I->isNoCreateUndefOrPoison);
};
auto TieL = TieBoolAttributes(L);
@@ -446,7 +447,8 @@ 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 || getEffectiveME(Int) != MemoryEffects::unknown();
+ Int.isStrictFP || Int.isNoCreateUndefOrPoison ||
+ getEffectiveME(Int) != MemoryEffects::unknown();
}
namespace {
@@ -605,6 +607,8 @@ static AttributeSet getIntrinsicFnAttributeSet(LLVMContext &C, unsigned ID) {
addAttribute("Speculatable");
if (Int.isStrictFP)
addAttribute("StrictFP");
+ if (Int.isNoCreateUndefOrPoison)
+ addAttribute("NoCreateUndefOrPoison");
const MemoryEffects ME = getEffectiveME(Int);
if (ME != MemoryEffects::unknown()) {
>From 40a3c93fdacae87a0705669e5ddcdda97080957b Mon Sep 17 00:00:00 2001
From: Jay Foad <jay.foad at amd.com>
Date: Thu, 23 Oct 2025 14:49:35 +0100
Subject: [PATCH 2/2] Update standard intrinsics
---
llvm/include/llvm/IR/Intrinsics.td | 80 ++++++++++---------
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 2 +-
llvm/lib/Analysis/ValueTracking.cpp | 67 ----------------
.../AArch64/replace-with-veclib-armpl.ll | 3 +-
.../AMDGPU/amdgpu-simplify-libcall-rootn.ll | 2 +-
.../replace-with-veclib-sleef-scalable.ll | 2 +-
llvm/test/Feature/intrinsics.ll | 5 +-
llvm/test/Linker/drop-attribute.ll | 2 +-
llvm/test/Transforms/Attributor/nofree.ll | 6 +-
llvm/test/Transforms/Attributor/nosync.ll | 2 +-
llvm/test/Transforms/Attributor/willreturn.ll | 16 ++--
.../EarlyCSE/replace-calls-def-attrs.ll | 2 +-
.../InstCombine/AMDGPU/fmul_legacy.ll | 12 ---
llvm/test/Transforms/LoopIdiom/basic.ll | 2 +-
.../AArch64/expand-exp.ll | 2 +-
.../Transforms/SimplifyCFG/rangereduce.ll | 2 +-
.../remove-attributes-from-intrinsics.ll | 2 +-
17 files changed, 71 insertions(+), 138 deletions(-)
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index e5fb0a9afad3c..55d03eee7e9c9 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -1043,7 +1043,7 @@ def int_experimental_memset_pattern
// FIXME: Add version of these floating point intrinsics which allow non-default
// rounding modes and FP exception handling.
-let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison] in {
def int_fma : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>,
LLVMMatchType<0>]>;
@@ -1056,16 +1056,8 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
// environment so they can be treated as readnone.
def int_sqrt : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
def int_powi : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_anyint_ty]>;
- def int_asin : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_acos : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_atan : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_atan2 : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, LLVMMatchType<0>]>;
def int_sin : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
def int_cos : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_tan : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_sinh : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_cosh : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_tanh : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
def int_pow : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>]>;
def int_log : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
@@ -1084,12 +1076,6 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
def int_nearbyint : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
def int_round : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
def int_roundeven : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
- def int_sincos : DefaultAttrsIntrinsic<[LLVMMatchType<0>, LLVMMatchType<0>],
- [llvm_anyfloat_ty]>;
- def int_sincospi : DefaultAttrsIntrinsic<[LLVMMatchType<0>, LLVMMatchType<0>],
- [llvm_anyfloat_ty]>;
- def int_modf : DefaultAttrsIntrinsic<[LLVMMatchType<0>, LLVMMatchType<0>],
- [llvm_anyfloat_ty]>;
// Truncate a floating point number with a specific rounding mode
def int_fptrunc_round : DefaultAttrsIntrinsic<[ llvm_anyfloat_ty ],
@@ -1101,6 +1087,8 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
def int_arithmetic_fence : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>],
[IntrNoMem]>;
+ // If the value doesn't fit an unspecified value is returned (but this
+ // is not poison).
def int_lround : DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty]>;
def int_llround : DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty]>;
def int_lrint : DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty]>;
@@ -1114,29 +1102,49 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
def int_frexp : DefaultAttrsIntrinsic<[llvm_anyfloat_ty, llvm_anyint_ty], [LLVMMatchType<0>]>;
}
+let IntrProperties = [IntrNoMem, IntrSpeculatable] 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.
+ def int_asin : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
+ def int_acos : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
+ def int_atan : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
+ def int_atan2 : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, LLVMMatchType<0>]>;
+ def int_tan : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
+ def int_sinh : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
+ def int_cosh : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
+ def int_tanh : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;
+ def int_sincos : DefaultAttrsIntrinsic<[LLVMMatchType<0>, LLVMMatchType<0>],
+ [llvm_anyfloat_ty]>;
+ def int_sincospi : DefaultAttrsIntrinsic<[LLVMMatchType<0>, LLVMMatchType<0>],
+ [llvm_anyfloat_ty]>;
+ def int_modf : DefaultAttrsIntrinsic<[LLVMMatchType<0>, LLVMMatchType<0>],
+ [llvm_anyfloat_ty]>;
+}
+
def int_minnum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]
+ [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
def int_maxnum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]
+ [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
def int_minimum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]
+ [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
def int_maximum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]
+ [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
def int_minimumnum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]
+ [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
def int_maximumnum : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]
+ [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
>;
// Internal interface for object size checking
@@ -1168,7 +1176,7 @@ let IntrProperties = [IntrInaccessibleMemOnly] in {
def int_is_fpclass
: DefaultAttrsIntrinsic<[LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>],
[llvm_anyfloat_ty, llvm_i32_ty],
- [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<1>>]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison, ImmArg<ArgIndex<1>>]>;
//===--------------- Constrained Floating Point Intrinsics ----------------===//
//
@@ -1410,7 +1418,7 @@ def int_expect_with_probability : DefaultAttrsIntrinsic<[llvm_anyint_ty],
//
// None of these intrinsics accesses memory at all.
-let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison] in {
def int_bswap: DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>]>;
def int_ctpop: DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>]>;
def int_bitreverse : DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>]>;
@@ -1526,7 +1534,7 @@ def int_adjust_trampoline : DefaultAttrsIntrinsic<
//
// Expose the carry flag from add operations on two integrals.
-let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison] in {
def int_sadd_with_overflow : DefaultAttrsIntrinsic<[llvm_anyint_ty,
LLVMScalarOrSameVectorWidth<0, llvm_i1_ty>],
[LLVMMatchType<0>, LLVMMatchType<0>]>;
@@ -1552,16 +1560,16 @@ let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
//
def int_sadd_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison, Commutative]>;
def int_uadd_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable, Commutative]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison, Commutative]>;
def int_ssub_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]>;
def int_usub_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]>;
def int_sshl_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem, IntrSpeculatable]>;
@@ -1616,22 +1624,22 @@ def int_abs : DefaultAttrsIntrinsic<
def int_smax : DefaultAttrsIntrinsic<
[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]>;
def int_smin : DefaultAttrsIntrinsic<
[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]>;
def int_umax : DefaultAttrsIntrinsic<
[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]>;
def int_umin : DefaultAttrsIntrinsic<
[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]>;
def int_scmp : DefaultAttrsIntrinsic<
[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>],
- [IntrNoMem, IntrSpeculatable, Range<RetIndex, -1, 2>]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison, Range<RetIndex, -1, 2>]>;
def int_ucmp : DefaultAttrsIntrinsic<
[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>],
- [IntrNoMem, IntrSpeculatable, Range<RetIndex, -1, 2>]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison, Range<RetIndex, -1, 2>]>;
//===------------------------- Memory Use Markers -------------------------===//
//
@@ -1873,7 +1881,7 @@ def int_convert_from_fp16 : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [llvm_i16_
}
// Saturating floating point to integer intrinsics
-let IntrProperties = [IntrNoMem, IntrSpeculatable] in {
+let IntrProperties = [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison] in {
def int_fptoui_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty]>;
def int_fptosi_sat : DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty]>;
}
@@ -1896,7 +1904,7 @@ def int_fake_use : DefaultAttrsIntrinsic<[], [llvm_vararg_ty],
// First argument must be pointer or vector of pointer. This is checked by the
// verifier.
def int_ptrmask: DefaultAttrsIntrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_anyint_ty],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, IntrNoCreateUndefOrPoison]>;
// Intrinsic to wrap a thread local variable.
def int_threadlocal_address : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [LLVMMatchType<0>],
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index cc375a35ca822..8e35109061792 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -449,7 +449,7 @@ def int_amdgcn_log_clamp : DefaultAttrsIntrinsic<
def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">,
DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
- [IntrNoMem, IntrSpeculatable, Commutative, IntrNoCreateUndefOrPoison]
+ [IntrNoMem, IntrSpeculatable, Commutative]
>;
// Fused single-precision multiply-add with legacy behaviour for the multiply,
diff --git a/llvm/lib/Analysis/ValueTracking.cpp b/llvm/lib/Analysis/ValueTracking.cpp
index c293b93f528b5..523374bdc472f 100644
--- a/llvm/lib/Analysis/ValueTracking.cpp
+++ b/llvm/lib/Analysis/ValueTracking.cpp
@@ -7419,79 +7419,12 @@ static bool canCreateUndefOrPoison(const Operator *Op, UndefPoisonKind Kind,
if (cast<ConstantInt>(II->getArgOperand(1))->isNullValue())
return false;
break;
- case Intrinsic::ctpop:
- case Intrinsic::bswap:
- case Intrinsic::bitreverse:
- case Intrinsic::fshl:
- case Intrinsic::fshr:
- case Intrinsic::smax:
- case Intrinsic::smin:
- case Intrinsic::scmp:
- case Intrinsic::umax:
- case Intrinsic::umin:
- case Intrinsic::ucmp:
- case Intrinsic::ptrmask:
- case Intrinsic::fptoui_sat:
- case Intrinsic::fptosi_sat:
- case Intrinsic::sadd_with_overflow:
- case Intrinsic::ssub_with_overflow:
- case Intrinsic::smul_with_overflow:
- case Intrinsic::uadd_with_overflow:
- case Intrinsic::usub_with_overflow:
- case Intrinsic::umul_with_overflow:
- case Intrinsic::sadd_sat:
- case Intrinsic::uadd_sat:
- case Intrinsic::ssub_sat:
- case Intrinsic::usub_sat:
- return false;
case Intrinsic::sshl_sat:
case Intrinsic::ushl_sat:
if (!includesPoison(Kind) ||
shiftAmountKnownInRange(II->getArgOperand(1)))
return false;
break;
- case Intrinsic::fma:
- case Intrinsic::fmuladd:
- case Intrinsic::sqrt:
- case Intrinsic::powi:
- case Intrinsic::sin:
- case Intrinsic::cos:
- case Intrinsic::pow:
- case Intrinsic::log:
- case Intrinsic::log10:
- case Intrinsic::log2:
- case Intrinsic::exp:
- case Intrinsic::exp2:
- case Intrinsic::exp10:
- case Intrinsic::fabs:
- case Intrinsic::copysign:
- case Intrinsic::floor:
- case Intrinsic::ceil:
- case Intrinsic::trunc:
- case Intrinsic::rint:
- case Intrinsic::nearbyint:
- case Intrinsic::round:
- case Intrinsic::roundeven:
- case Intrinsic::fptrunc_round:
- case Intrinsic::canonicalize:
- case Intrinsic::arithmetic_fence:
- case Intrinsic::minnum:
- case Intrinsic::maxnum:
- case Intrinsic::minimum:
- case Intrinsic::maximum:
- case Intrinsic::minimumnum:
- case Intrinsic::maximumnum:
- case Intrinsic::is_fpclass:
- case Intrinsic::ldexp:
- case Intrinsic::frexp:
- return false;
- case Intrinsic::lround:
- case Intrinsic::llround:
- case Intrinsic::lrint:
- case Intrinsic::llrint:
- // If the value doesn't fit an unspecified value is returned (but this
- // is not poison).
- return false;
}
}
[[fallthrough]];
diff --git a/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll b/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll
index 71c6380177b3a..8a0ac6d4ace7a 100644
--- a/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll
+++ b/llvm/test/CodeGen/AArch64/replace-with-veclib-armpl.ll
@@ -780,6 +780,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 nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison 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) }
;.
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll
index f6ae516faf2f7..89d039406f375 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-rootn.ll
@@ -1489,7 +1489,7 @@ attributes #2 = { noinline }
!0 = !{float 3.0}
;.
; CHECK: attributes #[[ATTR0]] = { strictfp }
-; CHECK: attributes #[[ATTR1:[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 #[[ATTR2:[0-9]+]] = { nounwind memory(read) }
; CHECK: attributes #[[ATTR3]] = { noinline }
; CHECK: 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 c489bc3681876..aa63552eb4b63 100644
--- a/llvm/test/CodeGen/RISCV/replace-with-veclib-sleef-scalable.ll
+++ b/llvm/test/CodeGen/RISCV/replace-with-veclib-sleef-scalable.ll
@@ -488,5 +488,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]] = { "target-features"="+v" }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+v" }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) "target-features"="+v" }
;.
diff --git a/llvm/test/Feature/intrinsics.ll b/llvm/test/Feature/intrinsics.ll
index b6abc0fff6db7..a2da8f29116e1 100644
--- a/llvm/test/Feature/intrinsics.ll
+++ b/llvm/test/Feature/intrinsics.ll
@@ -64,7 +64,7 @@ define void @libm() {
; FIXME: test ALL the intrinsics in this file.
-; CHECK: declare void @llvm.trap() #1
+; CHECK: declare void @llvm.trap() #2
declare void @llvm.trap()
define void @trap() {
@@ -72,5 +72,4 @@ define void @trap() {
ret void
}
-; CHECK: attributes #0 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #1 = { cold noreturn nounwind memory(inaccessiblemem: write) }
+; CHECK: attributes #2 = { cold noreturn nounwind memory(inaccessiblemem: write) }
diff --git a/llvm/test/Linker/drop-attribute.ll b/llvm/test/Linker/drop-attribute.ll
index 9be95a89109b4..3d4c13c2ffc75 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 nofree nosync nounwind speculatable willreturn
+; CHECK: ; Function Attrs: nocallback nocreateundeforpoison 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/Transforms/Attributor/nofree.ll b/llvm/test/Transforms/Attributor/nofree.ll
index 2a9d5d91ae053..94aa79aa327f4 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 nofree nosync nounwind speculatable willreturn memory(none)
+; CHECK: Function Attrs: nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)
declare float @llvm.floor.f32(float)
define void @call_floor(float %a) #0 {
@@ -489,7 +489,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 nofree nosync nounwind speculatable willreturn memory(none) }
+; TUNIT: attributes #[[ATTR6:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
; TUNIT: attributes #[[ATTR7]] = { nofree nounwind }
; TUNIT: attributes #[[ATTR8]] = { nobuiltin nofree nounwind }
; TUNIT: attributes #[[ATTR9]] = { nosync memory(none) }
@@ -506,7 +506,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 nofree nosync nounwind speculatable willreturn memory(none) }
+; CGSCC: attributes #[[ATTR6:[0-9]+]] = { nocallback nocreateundeforpoison 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 7ef46e8e94c9e..c15bd775ddb4d 100644
--- a/llvm/test/Transforms/Attributor/nosync.ll
+++ b/llvm/test/Transforms/Attributor/nosync.ll
@@ -454,7 +454,7 @@ define void @nosync_convergent_callee_test() {
; CHECK: attributes #[[ATTR14:[0-9]+]] = { convergent memory(none) }
; CHECK: attributes #[[ATTR15]] = { memory(none) }
; CHECK: attributes #[[ATTR16]] = { nounwind }
-; CHECK: attributes #[[ATTR17:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR17:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK: attributes #[[ATTR18]] = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) }
; CHECK: attributes #[[ATTR19]] = { nosync memory(none) }
; CHECK: attributes #[[ATTR20]] = { nofree nounwind }
diff --git a/llvm/test/Transforms/Attributor/willreturn.ll b/llvm/test/Transforms/Attributor/willreturn.ll
index b7ac7fc2970b0..d65480b05759a 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 nofree nosync nounwind speculatable willreturn memory(none)
+; CHECK: Function Attrs: nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)
declare float @llvm.floor.f32(float)
define void @call_floor(float %a) #0 {
@@ -425,7 +425,7 @@ define i32 @loop_constant_trip_count(ptr nocapture readonly %0) #0 {
; CHECK-NEXT: [[TMP4:%.*]] = phi i64 [ 0, [[TMP1:%.*]] ], [ [[TMP9:%.*]], [[TMP3]] ]
; CHECK-NEXT: [[TMP5:%.*]] = phi i32 [ 0, [[TMP1]] ], [ [[TMP8]], [[TMP3]] ]
; CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 [[TMP4]]
-; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4
+; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4, !invariant.load [[META0:![0-9]+]]
; CHECK-NEXT: [[TMP8]] = add nsw i32 [[TMP7]], [[TMP5]]
; CHECK-NEXT: [[TMP9]] = add nuw nsw i64 [[TMP4]], 1
; CHECK-NEXT: [[TMP10:%.*]] = icmp eq i64 [[TMP9]], 10
@@ -472,7 +472,7 @@ define i32 @loop_trip_count_unbound(i32 %0, i32 %1, ptr nocapture readonly %2, i
; CHECK-NEXT: [[TMP10:%.*]] = phi i32 [ [[TMP14]], [[TMP8]] ], [ 0, [[TMP4]] ]
; CHECK-NEXT: [[TMP11:%.*]] = zext i32 [[TMP9]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 [[TMP11]]
-; CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP12]], align 4
+; CHECK-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP12]], align 4, !invariant.load [[META0]]
; CHECK-NEXT: [[TMP14]] = add nsw i32 [[TMP13]], [[TMP10]]
; CHECK-NEXT: [[TMP15]] = add i32 [[TMP9]], [[TMP3]]
; CHECK-NEXT: [[TMP16:%.*]] = icmp eq i32 [[TMP15]], [[TMP1]]
@@ -522,7 +522,7 @@ define i32 @loop_trip_dec(i32 %0, ptr nocapture readonly %1) local_unnamed_addr
; CHECK-NEXT: [[TMP7:%.*]] = phi i64 [ [[TMP5]], [[TMP4]] ], [ [[TMP12:%.*]], [[TMP6]] ]
; CHECK-NEXT: [[TMP8:%.*]] = phi i32 [ 0, [[TMP4]] ], [ [[TMP11:%.*]], [[TMP6]] ]
; CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[TMP7]]
-; CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 4
+; CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr [[TMP9]], align 4, !invariant.load [[META0]]
; CHECK-NEXT: [[TMP11]] = add nsw i32 [[TMP10]], [[TMP8]]
; CHECK-NEXT: [[TMP12]] = add nsw i64 [[TMP7]], -1
; CHECK-NEXT: [[TMP13:%.*]] = icmp sgt i64 [[TMP7]], 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 nofree nosync nounwind speculatable willreturn memory(none) }
+; TUNIT: attributes #[[ATTR8:[0-9]+]] = { nocallback nocreateundeforpoison 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 nofree nosync nounwind speculatable willreturn memory(none) }
+; CGSCC: attributes #[[ATTR8:[0-9]+]] = { nocallback nocreateundeforpoison 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 }
@@ -1364,3 +1364,7 @@ attributes #1 = { uwtable noinline }
; CGSCC: attributes #[[ATTR37]] = { nosync willreturn memory(read) }
; CGSCC: attributes #[[ATTR38]] = { willreturn memory(read) }
;.
+; TUNIT: [[META0]] = !{}
+;.
+; CGSCC: [[META0]] = !{}
+;.
diff --git a/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll b/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll
index cf871e5714bf5..1dbffd962a638 100644
--- a/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll
+++ b/llvm/test/Transforms/EarlyCSE/replace-calls-def-attrs.ll
@@ -262,7 +262,7 @@ define i32 @commutative_intrinsic_intersection_failure(i32 %arg, i32 %arg1) {
}
;.
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK: attributes #[[ATTR1]] = { memory(none) }
; CHECK: attributes #[[ATTR2]] = { memory(read) }
; CHECK: attributes #[[ATTR3]] = { alwaysinline memory(none) }
diff --git a/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll b/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll
index dd5f6cc9a81bd..899b611b929d1 100644
--- a/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll
+++ b/llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll
@@ -91,18 +91,6 @@ define float @test_var_poison(float %x) {
ret float %call
}
-define float @test_freeze(float %x, float %y) {
-; CHECK-LABEL: @test_freeze(
-; CHECK-NEXT: [[FR:%.*]] = freeze float [[CALL:%.*]]
-; CHECK-NEXT: [[Y_FR:%.*]] = freeze float [[Y:%.*]]
-; CHECK-NEXT: [[CALL1:%.*]] = call float @llvm.amdgcn.fmul.legacy(float [[FR]], float [[Y_FR]])
-; CHECK-NEXT: ret float [[CALL1]]
-;
- %call = call float @llvm.amdgcn.fmul.legacy(float %x, float %y)
- %fr = freeze float %call
- ret float %fr
-}
-
declare float @llvm.amdgcn.fmul.legacy(float, float)
declare float @llvm.fabs.f32(float)
declare void @llvm.assume(i1 noundef)
diff --git a/llvm/test/Transforms/LoopIdiom/basic.ll b/llvm/test/Transforms/LoopIdiom/basic.ll
index e8ea912246728..9deccf5352ea8 100644
--- a/llvm/test/Transforms/LoopIdiom/basic.ll
+++ b/llvm/test/Transforms/LoopIdiom/basic.ll
@@ -1620,5 +1620,5 @@ define noalias ptr @_ZN8CMSPULog9beginImplEja(ptr nocapture writeonly %0) local_
; CHECK: attributes #[[ATTR1:[0-9]+]] = { nounwind }
; CHECK: attributes #[[ATTR2:[0-9]+]] = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }
; CHECK: attributes #[[ATTR3:[0-9]+]] = { nocallback nofree nounwind willreturn memory(argmem: write) }
-; CHECK: attributes #[[ATTR4:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR4:[0-9]+]] = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
;.
diff --git a/llvm/test/Transforms/PreISelIntrinsicLowering/AArch64/expand-exp.ll b/llvm/test/Transforms/PreISelIntrinsicLowering/AArch64/expand-exp.ll
index 09f583f9242d5..3416584729317 100644
--- a/llvm/test/Transforms/PreISelIntrinsicLowering/AArch64/expand-exp.ll
+++ b/llvm/test/Transforms/PreISelIntrinsicLowering/AArch64/expand-exp.ll
@@ -38,5 +38,5 @@ define <4 x float> @fixed_vec_exp(<4 x float> %input) {
declare <4 x float> @llvm.exp.v4f32(<4 x float>) #0
declare <vscale x 4 x float> @llvm.exp.nxv4f32(<vscale x 4 x float>) #0
-; CHECK: attributes #0 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #0 = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
attributes #0 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
diff --git a/llvm/test/Transforms/SimplifyCFG/rangereduce.ll b/llvm/test/Transforms/SimplifyCFG/rangereduce.ll
index d1fba91d1e505..169803f7aa012 100644
--- a/llvm/test/Transforms/SimplifyCFG/rangereduce.ll
+++ b/llvm/test/Transforms/SimplifyCFG/rangereduce.ll
@@ -321,7 +321,7 @@ three:
!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 nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nocreateundeforpoison 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 a6ace2246fa8e..b800f9aa97c8f 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 nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK-ALL: attributes #0 = { nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK-INTERESTINGNESS: attributes #1 = {
; CHECK-INTERESTINGNESS-SAME: "arg4"
More information about the llvm-commits
mailing list