Skip to content

[IR] Add new function attribute nocreateundeforpoison #164809

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Nov 4, 2025
Merged

[IR] Add new function attribute nocreateundeforpoison #164809

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
ac95b57
[IR] Add new function attribute nocreateundeforpoison
jayfoad Oct 23, 2025
40a3c93
Update standard intrinsics
jayfoad Oct 23, 2025
bf32374
Add comments
jayfoad Oct 23, 2025
8652177
Adopt nikic's wording
jayfoad Oct 31, 2025
ab58489
Merge branch 'main' into nocreateundeforpoison
jayfoad Nov 3, 2025
5657809
Update MLIR tests
jayfoad Nov 3, 2025
2477738
Update Clang tests
jayfoad Nov 3, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions llvm/docs/LangRef.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Copy link
Contributor

@arsenm arsenm Oct 24, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does this need to clarify the interaction with fast math flags?

Copy link
Contributor Author

@jayfoad jayfoad Oct 24, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 nocreateundeforpoison attribute. In that sense, poison-generating fast math flags take priority over this attribute.

?

Copy link
Contributor

@arsenm arsenm Oct 24, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, something like that

Copy link
Contributor

@arsenm arsenm Oct 24, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But same applies for other poison-generating metadata and attributes like range and !range

Copy link
Contributor Author

@jayfoad jayfoad Oct 24, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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:

  1. Having a range attribute on the return value, and returning a value outside that range
  2. Having arguments that are poison
  3. Not having the nocreateundeforpoison attribute
  4. ... 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?

Copy link
Contributor

@arsenm arsenm Oct 25, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think phrasing it in terms of what it doesn't mean is worse. Can we phrase it in terms of... base value behavior?

Copy link
Contributor Author

@jayfoad jayfoad Oct 29, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

@nikic nikic Oct 29, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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."

Copy link
Contributor Author

@jayfoad jayfoad Oct 31, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks, I used that.


Call Site Attributes
----------------------
Expand Down
1 change: 1 addition & 0 deletions llvm/include/llvm/Bitcode/LLVMBitCodes.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
Expand Down
5 changes: 5 additions & 0 deletions llvm/include/llvm/IR/Attributes.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]>;

Expand Down
4 changes: 4 additions & 0 deletions llvm/include/llvm/IR/Intrinsics.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
//===----------------------------------------------------------------------===//
Expand Down
2 changes: 1 addition & 1 deletion llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]
Copy link
Contributor

@nikic nikic Oct 23, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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?

Copy link
Contributor Author

@jayfoad jayfoad Oct 23, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sure, done.

>;

// Fused single-precision multiply-add with legacy behaviour for the multiply,
Expand Down
9 changes: 6 additions & 3 deletions llvm/lib/Analysis/ValueTracking.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down Expand Up @@ -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: {
Expand Down
2 changes: 2 additions & 0 deletions llvm/lib/Bitcode/Reader/BitcodeReader.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
}

Expand Down
2 changes: 2 additions & 0 deletions llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down
1 change: 1 addition & 0 deletions llvm/lib/Transforms/Utils/CodeExtractor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down
6 changes: 6 additions & 0 deletions llvm/test/Bitcode/attributes.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 }

Expand Down Expand Up @@ -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 }
Expand Down
12 changes: 12 additions & 0 deletions llvm/test/Transforms/InstCombine/AMDGPU/fmul_legacy.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
2 changes: 2 additions & 0 deletions llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
3 changes: 3 additions & 0 deletions llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
8 changes: 6 additions & 2 deletions llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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 {
Expand Down Expand Up @@ -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()) {
Expand Down