[llvm] [IR] Add new function attribute nocreateundeforpoison (PR #164809)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 23 04:57:17 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-ir
Author: Jay Foad (jayfoad)
<details>
<summary>Changes</summary>
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.
---
Full diff: https://github.com/llvm/llvm-project/pull/164809.diff
14 Files Affected:
- (modified) llvm/docs/LangRef.rst (+4)
- (modified) llvm/include/llvm/Bitcode/LLVMBitCodes.h (+1)
- (modified) llvm/include/llvm/IR/Attributes.td (+5)
- (modified) llvm/include/llvm/IR/Intrinsics.td (+4)
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+1-1)
- (modified) llvm/lib/Analysis/ValueTracking.cpp (+6-3)
- (modified) llvm/lib/Bitcode/Reader/BitcodeReader.cpp (+2)
- (modified) llvm/lib/Bitcode/Writer/BitcodeWriter.cpp (+2)
- (modified) llvm/lib/Transforms/Utils/CodeExtractor.cpp (+1)
- (modified) llvm/test/Bitcode/attributes.ll (+6)
- (modified) llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll (+12)
- (modified) llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp (+2)
- (modified) llvm/utils/TableGen/Basic/CodeGenIntrinsics.h (+3)
- (modified) llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp (+6-2)
``````````diff
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()) {
``````````
</details>
https://github.com/llvm/llvm-project/pull/164809
More information about the llvm-commits
mailing list