[IR] Add new function attribute nocreateundeforpoison #164809
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
llvmbot
added
backend:AMDGPU
tablegen
llvm:instcombine
|
There is some prior work in https://reviews.llvm.org/D152798 |
|
@llvm/pr-subscribers-backend-risc-v @llvm/pr-subscribers-llvm-ir Author: Jay Foad (jayfoad) ChangesAlso add a corresponding intrinsic property that can be used to mark As a smoke test this patch adds the new property to Full diff: https://github.com/llvm/llvm-project/pull/164809.diff 14 Files Affected:
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()) {
|
|
@llvm/pr-subscribers-backend-amdgpu Author: Jay Foad (jayfoad) ChangesAlso add a corresponding intrinsic property that can be used to mark As a smoke test this patch adds the new property to Full diff: https://github.com/llvm/llvm-project/pull/164809.diff 14 Files Affected:
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()) {
|
nikic
reviewed
Oct 23, 2025
| 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] |
There was a problem hiding this comment.
Instead of adding this to a new intrinsic, can we migrate the existing intrinsics (that don't need special handling) to use the attribute instead of being hardcoded in ValueTracking?
llvmbot
added
backend:AArch64
backend:RISC-V
LTO
nikic
approved these changes
Oct 23, 2025
| // 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>]>; |
There was a problem hiding this comment.
Add a TODO to move these into the IntrNoCreateUndefOrPoison case above? I don't think any of these produce poison, they were just missing from the switch...
arsenm
reviewed
Oct 24, 2025
llvm/docs/LangRef.rst
Outdated
| ``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. |
There was a problem hiding this comment.
Does this need to clarify the interaction with fast math flags?
There was a problem hiding this comment.
Good point. I'm not sure exactly how to word it. How about:
Note that a call instruction with poison-generating fast math flags may still generate poison, even if the call site or called function have the
nocreateundeforpoisonattribute. In that sense, poison-generating fast math flags take priority over this attribute.
?
There was a problem hiding this comment.
But same applies for other poison-generating metadata and attributes like range and !range
There was a problem hiding this comment.
For the interaction with other attributes: I do not want to come up with an order of precedence of attributes, saying which ones take priority over other ones, and I do not think it is necessary. The way I understand it is there are a bunch of reasons why a call might return poison, including:
- Having a
rangeattribute on the return value, and returning a value outside that range - Having arguments that are poison
- Not having the
nocreateundeforpoisonattribute - ... and so on.
I'm not sure how to explain this succinctly in the LangRef. Maybe I need to invert the description:
``nocreateundeforpoison``
Lack of this attribute indicates that the result of the function may be undef
or poison even if all arguments are not undef and not poison. Returning
undef or poison from a function or call site with this attribute triggers
undefined behavior.
Does that help?
There was a problem hiding this comment.
I think phrasing it in terms of what it doesn't mean is worse. Can we phrase it in terms of... base value behavior?
There was a problem hiding this comment.
Not sure what you mean by "base value" and I don't see that term used elsewhere. Are you suggesting we separate out the callee's internal view of the return value and the caller's external view of it? (Like "intensional" and "extensional" definitions in logic/semantics.) nocreateundeforpoison says something about the internal view. Poison-generating attributes and fast math flags only affect the external view, e.g. nnan says that if the internal is NaN then the external is poison.
There was a problem hiding this comment.
I think something like this would be fine? "This attribute indicates that the result of the function (prior to application of return attributes/metadata) will not be undef or poison if all arguments are not undef and not poison."
There was a problem hiding this comment.
Thanks, I used that.
Does |
No. Just like ValueTracking's canCreateUndefOrPoison, it tells you whether the result might be undef/poison when the arguments are not undef/poison. It says nothing about the result when the arguments are undef/poison. That would need a separate attribute. |
|
Overall, +1 on this attribute (and I assume a bunch of AMDGPU intrinsics will get it in a followup patch?) |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
6 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.