[llvm] [IR][TableGen] Add RangeSet TableGen intrinsic property (PR #203623)
Kirill Vedernikov via llvm-commits
llvm-commits at lists.llvm.org
Mon Jul 13 03:09:45 PDT 2026
https://github.com/kvederni updated https://github.com/llvm/llvm-project/pull/203623
>From 553b5ce799315807665f5d393c2687348e2de39d Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Fri, 12 Jun 2026 21:41:21 +0200
Subject: [PATCH 1/3] [IR][TableGen] Add rangeset IR attribute and RangeSet
TableGen intrinsic property
This change introduces a rangeset IR attribute and a TableGen intrinsic
property for the same. Existing IR range(<ty> lo, hi) syntax and
half-open semantics are unchanged. Existing Range<idx, lo, hi> keeps its
current half-open semantics and continues to emit the existing range
attribute.
IR syntax: rangeset(<ty> (<lo0>, <hi0>), (<lo1>, <hi1>), ...) where
(<lo>, <hi>) are closed intervals and must be ordered and
non-overlapped.
TableGen syntax: RangeSet<idx, [[lo0, hi0], [lo1, hi1], ...]> where
[lo, hi] are closed ordered and non-overlapped intervals.
---
llvm/include/llvm/AsmParser/LLParser.h | 2 +
llvm/include/llvm/Bitcode/LLVMBitCodes.h | 1 +
llvm/include/llvm/IR/Attributes.h | 10 ++
llvm/include/llvm/IR/Attributes.td | 4 +
llvm/include/llvm/IR/Intrinsics.td | 23 +++-
llvm/lib/AsmParser/LLParser.cpp | 85 ++++++++++--
llvm/lib/Bitcode/Reader/BitcodeReader.cpp | 6 +-
llvm/lib/Bitcode/Writer/BitcodeWriter.cpp | 2 +
.../CodeGen/SelectionDAG/TargetLowering.cpp | 9 +-
llvm/lib/IR/Attributes.cpp | 129 +++++++++++++++++-
llvm/lib/IR/Verifier.cpp | 30 +++-
llvm/test/Assembler/invalid-immarg.ll | 22 +--
.../Assembler/rangeset-attribute-invalid.ll | 36 +++++
llvm/test/Assembler/rangeset-attribute.ll | 19 +++
.../TableGen/intrinsic-rangeset-errors.td | 27 ++++
llvm/test/TableGen/intrinsic-rangeset.td | 28 ++++
llvm/unittests/IR/AttributesTest.cpp | 9 ++
.../TableGen/Basic/CodeGenIntrinsics.cpp | 67 +++++++++
llvm/utils/TableGen/Basic/CodeGenIntrinsics.h | 13 +-
.../utils/TableGen/Basic/IntrinsicEmitter.cpp | 16 ++-
20 files changed, 500 insertions(+), 38 deletions(-)
create mode 100644 llvm/test/Assembler/rangeset-attribute-invalid.ll
create mode 100644 llvm/test/Assembler/rangeset-attribute.ll
create mode 100644 llvm/test/TableGen/intrinsic-rangeset-errors.td
create mode 100644 llvm/test/TableGen/intrinsic-rangeset.td
diff --git a/llvm/include/llvm/AsmParser/LLParser.h b/llvm/include/llvm/AsmParser/LLParser.h
index 788b56cb78f08..22df21290c015 100644
--- a/llvm/include/llvm/AsmParser/LLParser.h
+++ b/llvm/include/llvm/AsmParser/LLParser.h
@@ -309,6 +309,7 @@ namespace llvm {
Loc = Lex.getLoc();
return parseUInt64(Val);
}
+ bool parseAPInt(unsigned BitWidth, APInt &Val);
bool parseFlag(unsigned &Val);
bool parseStringAttribute(AttrBuilder &B);
@@ -412,6 +413,7 @@ namespace llvm {
std::vector<unsigned> &FwdRefAttrGrps,
bool inAttrGrp, LocTy &BuiltinLoc);
bool parseRangeAttr(AttrBuilder &B);
+ bool parseRangeSetAttr(AttrBuilder &B);
bool parseInitializesAttr(AttrBuilder &B);
bool parseCapturesAttr(AttrBuilder &B);
bool parseRequiredTypeAttr(AttrBuilder &B, lltok::Kind AttrToken,
diff --git a/llvm/include/llvm/Bitcode/LLVMBitCodes.h b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
index 358f9a65a80af..2d3efb653683a 100644
--- a/llvm/include/llvm/Bitcode/LLVMBitCodes.h
+++ b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
@@ -826,6 +826,7 @@ enum AttributeKindCodes {
ATTR_KIND_NOOUTLINE = 107,
ATTR_KIND_FLATTEN = 108,
ATTR_KIND_NOIPA = 109,
+ ATTR_KIND_RANGE_SET = 110,
};
enum ComdatSelectionKindCodes {
diff --git a/llvm/include/llvm/IR/Attributes.h b/llvm/include/llvm/IR/Attributes.h
index 049f699b2980b..2c157aab122ac 100644
--- a/llvm/include/llvm/IR/Attributes.h
+++ b/llvm/include/llvm/IR/Attributes.h
@@ -355,6 +355,9 @@ class Attribute {
/// Returns the value of the range attribute.
LLVM_ABI const ConstantRange &getRange() const;
+ /// Returns the value of the rangeset attribute.
+ LLVM_ABI ArrayRef<ConstantRange> getRangeSet() const;
+
/// Returns the value of the initializes attribute.
LLVM_ABI ArrayRef<ConstantRange> getInitializes() const;
@@ -1339,6 +1342,9 @@ class AttrBuilder {
/// Add range attribute.
LLVM_ABI AttrBuilder &addRangeAttr(const ConstantRange &CR);
+ /// Add rangeset attribute.
+ LLVM_ABI AttrBuilder &addRangeSetAttr(ArrayRef<ConstantRange> CRs);
+
/// Add a ConstantRangeList attribute with the given ranges.
LLVM_ABI AttrBuilder &addConstantRangeListAttr(Attribute::AttrKind Kind,
ArrayRef<ConstantRange> Val);
@@ -1369,6 +1375,10 @@ enum AttributeSafetyKind : uint8_t {
/// follows the same type rules as FPMathOperator.
LLVM_ABI bool isNoFPClassCompatibleType(Type *Ty);
+/// Returns true if the ranges are ordered, non-overlapping, and canonical for
+/// the 'rangeset' attribute.
+LLVM_ABI bool isOrderedRangeSet(ArrayRef<ConstantRange> Ranges);
+
/// Which attributes cannot be applied to a type. The argument \p AS
/// is used as a hint for the attributes whose compatibility is being
/// checked against \p Ty. This does not mean the return will be a
diff --git a/llvm/include/llvm/IR/Attributes.td b/llvm/include/llvm/IR/Attributes.td
index 4e45100b54d38..5dbe593a7ba17 100644
--- a/llvm/include/llvm/IR/Attributes.td
+++ b/llvm/include/llvm/IR/Attributes.td
@@ -282,6 +282,10 @@ def Preallocated : TypeAttr<"preallocated", IntersectPreserve, [FnAttr, ParamAtt
/// Parameter or return value is within the specified range.
def Range : ConstantRangeAttr<"range", IntersectCustom, [ParamAttr, RetAttr]>;
+/// Parameter or return value is within one of the specified ranges.
+def RangeSet : ConstantRangeListAttr<"rangeset", IntersectCustom,
+ [ParamAttr, RetAttr]>;
+
/// Function does not access memory.
def ReadNone : EnumAttr<"readnone", IntersectAnd, [ParamAttr]>;
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 37c9c783465d6..1e0b9e17f121f 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -173,10 +173,29 @@ class ReadNone<ArgIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
}
+// The return value or argument is in the union of the specified closed ranges.
+// Range entries must be ordered and non-overlapping.
+class RangeSet<AttrIndex idx, list<list<int>> ranges> : IntrinsicProperty {
+ list<list<int>> ValidRanges = !filter(r, ranges, !eq(!size(r), 2));
+
+ assert !not(!empty(ranges)), "RangeSet requires at least one range";
+ assert !eq(!size(ValidRanges), !size(ranges)),
+ "RangeSet requires each range to have exactly two bounds";
+ assert !empty(!filter(r, ValidRanges, !lt(r[1], r[0]))),
+ "RangeSet requires lower <= upper for each range";
+ assert !empty(!filter(i, !range(1, !size(ValidRanges)),
+ !le(ValidRanges[i][0],
+ ValidRanges[!sub(i, 1)][1]))),
+ "RangeSet requires ordered and non-overlapping ranges";
+
+ int ArgNo = idx.Value;
+ list<list<int>> Ranges = ranges;
+}
+
// The return value or argument is in the range [lower, upper),
// where lower and upper are interpreted as signed integers.
-class Range<AttrIndex idx, int lower, int upper> : IntrinsicProperty {
- int ArgNo = idx.Value;
+class Range<AttrIndex idx, int lower, int upper>
+ : RangeSet<idx, [[lower, !sub(upper, 1)]]> {
int Lower = lower;
int Upper = upper;
}
diff --git a/llvm/lib/AsmParser/LLParser.cpp b/llvm/lib/AsmParser/LLParser.cpp
index 38d10587b104e..2a24c2c1a916d 100644
--- a/llvm/lib/AsmParser/LLParser.cpp
+++ b/llvm/lib/AsmParser/LLParser.cpp
@@ -1733,6 +1733,8 @@ bool LLParser::parseEnumAttribute(Attribute::AttrKind Attr, AttrBuilder &B,
}
case Attribute::Range:
return parseRangeAttr(B);
+ case Attribute::RangeSet:
+ return parseRangeSetAttr(B);
case Attribute::Initializes:
return parseInitializesAttr(B);
case Attribute::Captures:
@@ -1989,6 +1991,18 @@ bool LLParser::parseUInt64(uint64_t &Val) {
return false;
}
+/// parseAPInt
+/// ::= apint
+bool LLParser::parseAPInt(unsigned BitWidth, APInt &Val) {
+ if (Lex.getKind() != lltok::APSInt)
+ return tokError("expected integer");
+ if (Lex.getAPSIntVal().getBitWidth() > BitWidth)
+ return tokError("integer is too large for the bit width of specified type");
+ Val = Lex.getAPSIntVal().extend(BitWidth);
+ Lex.Lex();
+ return false;
+}
+
/// parseTLSModel
/// := 'localdynamic'
/// := 'initialexec'
@@ -3396,17 +3410,6 @@ bool LLParser::parseRangeAttr(AttrBuilder &B) {
Type *Ty = nullptr;
LocTy TyLoc;
- auto ParseAPSInt = [&](unsigned BitWidth, APInt &Val) {
- if (Lex.getKind() != lltok::APSInt)
- return tokError("expected integer");
- if (Lex.getAPSIntVal().getBitWidth() > BitWidth)
- return tokError(
- "integer is too large for the bit width of specified type");
- Val = Lex.getAPSIntVal().extend(BitWidth);
- Lex.Lex();
- return false;
- };
-
if (parseToken(lltok::lparen, "expected '('") || parseType(Ty, TyLoc))
return true;
if (!Ty->isIntegerTy())
@@ -3414,8 +3417,8 @@ bool LLParser::parseRangeAttr(AttrBuilder &B) {
unsigned BitWidth = Ty->getPrimitiveSizeInBits();
- if (ParseAPSInt(BitWidth, Lower) ||
- parseToken(lltok::comma, "expected ','") || ParseAPSInt(BitWidth, Upper))
+ if (parseAPInt(BitWidth, Lower) || parseToken(lltok::comma, "expected ','") ||
+ parseAPInt(BitWidth, Upper))
return true;
if (Lower == Upper && !Lower.isZero())
return tokError("the range represent the empty set but limits aren't 0!");
@@ -3427,6 +3430,62 @@ bool LLParser::parseRangeAttr(AttrBuilder &B) {
return false;
}
+/// parseRangeSetAttr
+/// ::= rangeset(<ty> (<n>,<n>)[,(<n>,<n>)]*)
+bool LLParser::parseRangeSetAttr(AttrBuilder &B) {
+ Lex.Lex();
+
+ Type *Ty = nullptr;
+ LocTy TyLoc;
+
+ if (parseToken(lltok::lparen, "expected '('"))
+ return true;
+
+ TyLoc = Lex.getLoc();
+ if (Lex.getKind() != lltok::Type)
+ return tokError("expected type");
+ Ty = Lex.getTyVal();
+ Lex.Lex();
+
+ if (!Ty->isIntegerTy())
+ return error(TyLoc, "the rangeset must have integer type!");
+
+ unsigned BitWidth = Ty->getPrimitiveSizeInBits();
+
+ SmallVector<ConstantRange, 2> RangeList;
+ do {
+ APInt Lower, Upper;
+ if (parseToken(lltok::lparen, "expected '('") ||
+ parseAPInt(BitWidth, Lower) ||
+ parseToken(lltok::comma, "expected ','") ||
+ parseAPInt(BitWidth, Upper) ||
+ parseToken(lltok::rparen, "expected ')'"))
+ return true;
+
+ if (Upper.slt(Lower))
+ return tokError("rangeset requires lower <= upper");
+
+ ConstantRange Range = ConstantRange::getNonEmpty(Lower, Upper + 1);
+ if (!RangeList.empty() && !RangeList.back().isFullSet() &&
+ !RangeList.back().getUpper().isMinSignedValue() &&
+ Lower == RangeList.back().getUpper()) {
+ RangeList.back() = ConstantRange::getNonEmpty(RangeList.back().getLower(),
+ Range.getUpper());
+ continue;
+ }
+ RangeList.push_back(Range);
+ } while (EatIfPresent(lltok::comma));
+
+ if (parseToken(lltok::rparen, "expected ')'"))
+ return true;
+
+ if (!AttributeFuncs::isOrderedRangeSet(RangeList))
+ return tokError("Invalid (unordered or overlapping) range set");
+
+ B.addRangeSetAttr(RangeList);
+ return false;
+}
+
/// parseInitializesAttr
/// ::= initializes((Lo1,Hi1),(Lo2,Hi2),...)
bool LLParser::parseInitializesAttr(AttrBuilder &B) {
diff --git a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
index a5a11ab6221d4..cec11303ca40b 100644
--- a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
+++ b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
@@ -2304,6 +2304,8 @@ static Attribute::AttrKind getAttrFromCode(uint64_t Code) {
return Attribute::DeadOnUnwind;
case bitc::ATTR_KIND_RANGE:
return Attribute::Range;
+ case bitc::ATTR_KIND_RANGE_SET:
+ return Attribute::RangeSet;
case bitc::ATTR_KIND_INITIALIZES:
return Attribute::Initializes;
case bitc::ATTR_KIND_CORO_ELIDE_SAFE:
@@ -2582,7 +2584,9 @@ Error BitcodeReader::parseAttributeGroupBlock() {
}
i--;
- if (!ConstantRangeList::isOrderedRanges(Val))
+ if (Kind == Attribute::RangeSet
+ ? !AttributeFuncs::isOrderedRangeSet(Val)
+ : !ConstantRangeList::isOrderedRanges(Val))
return error("Invalid (unordered or overlapping) range list");
B.addConstantRangeListAttr(Kind, Val);
} else {
diff --git a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
index a6e7941a06699..210ba1191526a 100644
--- a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
+++ b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
@@ -998,6 +998,8 @@ static uint64_t getAttrKindEncoding(Attribute::AttrKind Kind) {
return bitc::ATTR_KIND_DEAD_ON_UNWIND;
case Attribute::Range:
return bitc::ATTR_KIND_RANGE;
+ case Attribute::RangeSet:
+ return bitc::ATTR_KIND_RANGE_SET;
case Attribute::Initializes:
return bitc::ATTR_KIND_INITIALIZES;
case Attribute::NoExt:
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index bca34c5c347ee..8fc28ea2b8b0e 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -70,10 +70,11 @@ bool TargetLowering::isInTailCallPosition(SelectionDAG &DAG, SDNode *Node,
// the return. Ignore following attributes because they don't affect the
// call sequence.
AttrBuilder CallerAttrs(F.getContext(), F.getAttributes().getRetAttrs());
- for (const auto &Attr : {Attribute::Alignment, Attribute::Dereferenceable,
- Attribute::DereferenceableOrNull, Attribute::NoAlias,
- Attribute::NonNull, Attribute::NoUndef,
- Attribute::Range, Attribute::NoFPClass})
+ for (const auto &Attr :
+ {Attribute::Alignment, Attribute::Dereferenceable,
+ Attribute::DereferenceableOrNull, Attribute::NoAlias,
+ Attribute::NonNull, Attribute::NoUndef, Attribute::Range,
+ Attribute::RangeSet, Attribute::NoFPClass})
CallerAttrs.removeAttribute(Attr);
if (CallerAttrs.hasAttributes())
diff --git a/llvm/lib/IR/Attributes.cpp b/llvm/lib/IR/Attributes.cpp
index 4087b25951a1c..dec94c313c884 100644
--- a/llvm/lib/IR/Attributes.cpp
+++ b/llvm/lib/IR/Attributes.cpp
@@ -92,6 +92,64 @@ unpackVScaleRangeArgs(uint64_t Value) {
MaxValue > 0 ? MaxValue : std::optional<unsigned>());
}
+static void appendRangeSetRange(SmallVectorImpl<ConstantRange> &Ranges,
+ const ConstantRange &CR) {
+ if (CR.isEmptySet())
+ return;
+ if (CR.isFullSet()) {
+ Ranges.assign(1, CR);
+ return;
+ }
+ if (Ranges.empty()) {
+ Ranges.push_back(CR);
+ return;
+ }
+ if (Ranges.front().isFullSet())
+ return;
+
+ APInt PreviousUpper = Ranges.back().getUpper();
+ if (PreviousUpper.isMinSignedValue())
+ return;
+ if (CR.getLower().sle(PreviousUpper)) {
+ APInt NewUpper = CR.getUpper().isMinSignedValue()
+ ? APInt::getSignedMinValue(CR.getBitWidth())
+ : APIntOps::smax(PreviousUpper, CR.getUpper());
+ Ranges.back() = ConstantRange(Ranges.back().getLower(), NewUpper);
+ if (Ranges.back().isFullSet())
+ Ranges.truncate(1);
+ return;
+ }
+
+ Ranges.push_back(CR);
+}
+
+static SmallVector<ConstantRange, 2>
+unionRangeSetLists(ArrayRef<ConstantRange> LHS, ArrayRef<ConstantRange> RHS) {
+ assert(AttributeFuncs::isOrderedRangeSet(LHS) &&
+ AttributeFuncs::isOrderedRangeSet(RHS) &&
+ "rangeset attribute must be ordered");
+ if (LHS.empty())
+ return SmallVector<ConstantRange, 2>(RHS);
+ if (RHS.empty())
+ return SmallVector<ConstantRange, 2>(LHS);
+
+ SmallVector<ConstantRange, 2> Result;
+ size_t LHSIdx = 0, RHSIdx = 0;
+ while (LHSIdx < LHS.size() || RHSIdx < RHS.size()) {
+ const ConstantRange *NextRange;
+ if (RHSIdx == RHS.size() ||
+ (LHSIdx < LHS.size() &&
+ LHS[LHSIdx].getLower().slt(RHS[RHSIdx].getLower())))
+ NextRange = &LHS[LHSIdx++];
+ else
+ NextRange = &RHS[RHSIdx++];
+ appendRangeSetRange(Result, *NextRange);
+ if (Result.size() == 1 && Result.front().isFullSet())
+ break;
+ }
+ return Result;
+}
+
Attribute Attribute::get(LLVMContext &Context, Attribute::AttrKind Kind,
uint64_t Val) {
bool IsIntAttr = Attribute::isIntAttrKind(Kind);
@@ -529,6 +587,12 @@ const ConstantRange &Attribute::getRange() const {
return pImpl->getValueAsConstantRange();
}
+ArrayRef<ConstantRange> Attribute::getRangeSet() const {
+ assert(hasAttribute(Attribute::RangeSet) &&
+ "Trying to get rangeset attr from non-ConstantRangeList attribute");
+ return pImpl->getValueAsConstantRangeList();
+}
+
ArrayRef<ConstantRange> Attribute::getInitializes() const {
assert(hasAttribute(Attribute::Initializes) &&
"Trying to get initializes attr from non-ConstantRangeList attribute");
@@ -739,6 +803,21 @@ std::string Attribute::getAsString(bool InAttrGrp) const {
return Result;
}
+ if (hasAttribute(Attribute::RangeSet)) {
+ std::string Result;
+ raw_string_ostream OS(Result);
+ ArrayRef<ConstantRange> Ranges = getRangeSet();
+ assert(!Ranges.empty() &&
+ "rangeset attribute must have at least one range");
+ OS << "rangeset(";
+ OS << "i" << Ranges.front().getBitWidth() << " ";
+ interleaveComma(Ranges, OS, [&](const ConstantRange &CR) {
+ OS << "(" << CR.getLower() << ", " << CR.getUpper() - 1 << ")";
+ });
+ OS << ")";
+ return Result;
+ }
+
if (hasAttribute(Attribute::Initializes)) {
std::string Result;
raw_string_ostream OS(Result);
@@ -1150,6 +1229,11 @@ AttributeSet::intersectWith(LLVMContext &C, AttributeSet Other) const {
if (!NewRange.isFullSet())
Intersected.addRangeAttr(NewRange);
} break;
+ case Attribute::RangeSet: {
+ SmallVector<ConstantRange> NewRangeSet =
+ unionRangeSetLists(Attr0.getRangeSet(), Attr1.getRangeSet());
+ Intersected.addRangeSetAttr(NewRangeSet);
+ } break;
default:
llvm_unreachable("Unknown attribute with custom intersection rule");
}
@@ -2371,9 +2455,17 @@ AttrBuilder &AttrBuilder::addRangeAttr(const ConstantRange &CR) {
return addConstantRangeAttr(Attribute::Range, CR);
}
+AttrBuilder &AttrBuilder::addRangeSetAttr(ArrayRef<ConstantRange> CRs) {
+ if (CRs.empty() || (CRs.size() == 1 && CRs.front().isFullSet()))
+ return *this;
+ return addAttribute(Attribute::get(Ctx, Attribute::RangeSet, CRs));
+}
+
AttrBuilder &
AttrBuilder::addConstantRangeListAttr(Attribute::AttrKind Kind,
ArrayRef<ConstantRange> Val) {
+ if (Kind == Attribute::RangeSet)
+ return addRangeSetAttr(Val);
return addAttribute(Attribute::get(Ctx, Kind, Val));
}
@@ -2479,6 +2571,33 @@ bool AttributeFuncs::isNoFPClassCompatibleType(Type *Ty) {
return FPMathOperator::isSupportedFloatingPointType(Ty);
}
+bool AttributeFuncs::isOrderedRangeSet(ArrayRef<ConstantRange> Ranges) {
+ if (Ranges.empty())
+ return true;
+
+ for (unsigned I = 0; I < Ranges.size(); ++I) {
+ const ConstantRange &Range = Ranges[I];
+ if (Range.isEmptySet())
+ return false;
+ if (I != 0 && Range.getBitWidth() != Ranges.front().getBitWidth())
+ return false;
+ if (Range.isFullSet())
+ return Ranges.size() == 1;
+ if (!Range.getLower().slt(Range.getUpper()) &&
+ !Range.getUpper().isMinSignedValue())
+ return false;
+ if (I == 0)
+ continue;
+
+ const ConstantRange &PrevRange = Ranges[I - 1];
+ if (PrevRange.isFullSet() || PrevRange.getUpper().isMinSignedValue())
+ return false;
+ if (Range.getLower().sle(PrevRange.getUpper()))
+ return false;
+ }
+ return true;
+}
+
/// Which attributes cannot be applied to a type.
AttributeMask AttributeFuncs::typeIncompatible(Type *Ty, AttributeSet AS,
AttributeSafetyKind ASK) {
@@ -2499,12 +2618,20 @@ AttributeMask AttributeFuncs::typeIncompatible(Type *Ty, AttributeSet AS,
if (!Ty->isIntOrIntVectorTy()) {
// Attributes that only apply to integers or vector of integers.
if (ASK & ASK_SAFE_TO_DROP)
- Incompatible.addAttribute(Attribute::Range);
+ Incompatible.addAttribute(Attribute::Range)
+ .addAttribute(Attribute::RangeSet);
} else {
Attribute RangeAttr = AS.getAttribute(Attribute::Range);
if (RangeAttr.isValid() &&
RangeAttr.getRange().getBitWidth() != Ty->getScalarSizeInBits())
Incompatible.addAttribute(Attribute::Range);
+ Attribute RangeSetAttr = AS.getAttribute(Attribute::RangeSet);
+ if (RangeSetAttr.isValid()) {
+ ArrayRef<ConstantRange> Ranges = RangeSetAttr.getRangeSet();
+ if (!Ranges.empty() &&
+ Ranges.front().getBitWidth() != Ty->getScalarSizeInBits())
+ Incompatible.addAttribute(Attribute::RangeSet);
+ }
}
if (!Ty->isPointerTy()) {
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index dd5140c61dc67..b2c0e42b0b54e 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -2036,11 +2036,12 @@ void Verifier::verifyParameterAttrs(AttributeSet Attrs, Type *Ty,
V);
if (Attrs.hasAttribute(Attribute::ImmArg)) {
- unsigned AttrCount =
- Attrs.getNumAttributes() - Attrs.hasAttribute(Attribute::Range);
+ unsigned AttrCount = Attrs.getNumAttributes() -
+ Attrs.hasAttribute(Attribute::Range) -
+ Attrs.hasAttribute(Attribute::RangeSet);
Check(AttrCount == 1,
"Attribute 'immarg' is incompatible with other attributes except the "
- "'range' attribute",
+ "'range' and 'rangeset' attributes",
V);
}
@@ -2185,6 +2186,16 @@ void Verifier::verifyParameterAttrs(AttributeSet Attrs, Type *Ty,
Check(Ty->isIntOrIntVectorTy(CR.getBitWidth()),
"Range bit width must match type bit width!", V);
}
+ if (Attrs.hasAttribute(Attribute::RangeSet)) {
+ ArrayRef<ConstantRange> Ranges =
+ Attrs.getAttribute(Attribute::RangeSet).getRangeSet();
+ Check(!Ranges.empty(), "Attribute 'rangeset' does not support empty list",
+ V);
+ Check(AttributeFuncs::isOrderedRangeSet(Ranges),
+ "Attribute 'rangeset' does not support unordered ranges", V);
+ Check(Ty->isIntOrIntVectorTy(Ranges.front().getBitWidth()),
+ "RangeSet bit width must match type bit width!", V);
+ }
}
void Verifier::checkUnsignedBaseTenFuncAttr(AttributeList Attrs, StringRef Attr,
@@ -3908,6 +3919,19 @@ void Verifier::visitCallBase(CallBase &Call) {
Call);
}
}
+ if (Call.paramHasAttr(i, Attribute::RangeSet)) {
+ if (auto *CI = dyn_cast<ConstantInt>(ArgVal)) {
+ ArrayRef<ConstantRange> Ranges =
+ Call.getParamAttr(i, Attribute::RangeSet).getRangeSet();
+ bool IsContained = llvm::any_of(Ranges, [&](const ConstantRange &CR) {
+ return CR.contains(CI->getValue());
+ });
+ Check(IsContained,
+ "immarg value " + Twine(CI->getValue().getSExtValue()) +
+ " out of rangeset",
+ Call);
+ }
+ }
}
if (Call.paramHasAttr(i, Attribute::Preallocated)) {
diff --git a/llvm/test/Assembler/invalid-immarg.ll b/llvm/test/Assembler/invalid-immarg.ll
index df69a66c67332..c04f1d74b882c 100644
--- a/llvm/test/Assembler/invalid-immarg.ll
+++ b/llvm/test/Assembler/invalid-immarg.ll
@@ -1,34 +1,34 @@
; RUN: not llvm-as < %s -o /dev/null 2>&1 | FileCheck %s
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.byval(ptr byval(i32) immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.inalloca(ptr inalloca(i32) immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.inreg(i32 inreg immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.nest(ptr nest immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.sret(ptr sret(i32) immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.zeroext(i32 zeroext immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.signext(i32 signext immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.returned(i32 returned immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.noalias(ptr noalias immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.readnone(ptr readnone immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
declare void @llvm.immarg.readonly(ptr readonly immarg)
diff --git a/llvm/test/Assembler/rangeset-attribute-invalid.ll b/llvm/test/Assembler/rangeset-attribute-invalid.ll
new file mode 100644
index 0000000000000..98605a1ac4b3a
--- /dev/null
+++ b/llvm/test/Assembler/rangeset-attribute-invalid.ll
@@ -0,0 +1,36 @@
+; RUN: split-file %s %t
+; RUN: not llvm-as < %t/non_integer_type.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=NON-INTEGER-TYPE
+; RUN: not llvm-as < %t/missing_type.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=MISSING-TYPE
+; RUN: not llvm-as < %t/empty.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=EMPTY
+; RUN: not llvm-as < %t/bad_bounds.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=BAD-BOUNDS
+; RUN: not llvm-as < %t/overlap.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=OVERLAP
+; RUN: not llvm-as < %t/unordered.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=UNORDERED
+; RUN: not llvm-as < %t/missing_comma.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=MISSING-COMMA
+
+;--- non_integer_type.ll
+; NON-INTEGER-TYPE: the rangeset must have integer type!
+declare rangeset(float (0, 2)) float @non_integer_type()
+
+;--- missing_type.ll
+; MISSING-TYPE: expected type
+declare rangeset((0, 2)) i32 @missing_type()
+
+;--- empty.ll
+; EMPTY: expected '('
+declare rangeset(i32) i32 @empty()
+
+;--- bad_bounds.ll
+; BAD-BOUNDS: rangeset requires lower <= upper
+declare rangeset(i32 (7, 5)) i32 @bad_bounds()
+
+;--- overlap.ll
+; OVERLAP: Invalid (unordered or overlapping) range set
+declare rangeset(i32 (0, 2), (2, 4)) i32 @overlap()
+
+;--- unordered.ll
+; UNORDERED: Invalid (unordered or overlapping) range set
+declare rangeset(i32 (5, 8), (0, 2)) i32 @unordered()
+
+;--- missing_comma.ll
+; MISSING-COMMA: expected ','
+declare rangeset(i32 (0 2)) i32 @missing_comma()
diff --git a/llvm/test/Assembler/rangeset-attribute.ll b/llvm/test/Assembler/rangeset-attribute.ll
new file mode 100644
index 0000000000000..a47ab51040fb2
--- /dev/null
+++ b/llvm/test/Assembler/rangeset-attribute.ll
@@ -0,0 +1,19 @@
+; RUN: llvm-as < %s | llvm-dis | FileCheck %s
+
+; CHECK: declare rangeset(i32 (0, 2), (5, 8)) i32 @ret_rangeset()
+declare rangeset(i32 (0, 2), (5, 8)) i32 @ret_rangeset()
+
+; CHECK: declare void @param_rangeset(i32 rangeset(i32 (0, 2), (5, 8)))
+declare void @param_rangeset(i32 rangeset(i32 (0, 2), (5, 8)))
+
+; CHECK: declare rangeset(i32 (0, 2), (5, 8)) <4 x i32> @ret_rangeset_vec()
+declare rangeset(i32 (0, 2), (5, 8)) <4 x i32> @ret_rangeset_vec()
+
+; CHECK: declare void @param_rangeset_vec(<4 x i32> rangeset(i32 (0, 2), (5, 8)))
+declare void @param_rangeset_vec(<4 x i32> rangeset(i32 (0, 2), (5, 8)))
+
+; CHECK: declare rangeset(i32 (1, 2147483647)) i32 @ret_rangeset_terminal()
+declare rangeset(i32 (1, 2147483647)) i32 @ret_rangeset_terminal()
+
+; CHECK: declare rangeset(i32 (0, 4), (6, 8)) i32 @ret_rangeset_coalesced()
+declare rangeset(i32 (0, 2), (3, 4), (6, 8)) i32 @ret_rangeset_coalesced()
diff --git a/llvm/test/TableGen/intrinsic-rangeset-errors.td b/llvm/test/TableGen/intrinsic-rangeset-errors.td
new file mode 100644
index 0000000000000..795220cb3d00b
--- /dev/null
+++ b/llvm/test/TableGen/intrinsic-rangeset-errors.td
@@ -0,0 +1,27 @@
+// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DEMPTY %s 2>&1 | FileCheck --check-prefix=EMPTY %s -DFILE=%s
+// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DBAD_BOUNDS %s 2>&1 | FileCheck --check-prefix=BAD-BOUNDS %s -DFILE=%s
+// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DOVERLAP %s 2>&1 | FileCheck --check-prefix=OVERLAP %s -DFILE=%s
+
+#define TEST_INTRINSICS_SUPPRESS_DEFS
+include "llvm/IR/Intrinsics.td"
+
+#ifdef EMPTY
+// EMPTY: error: assertion failed: RangeSet requires at least one range
+def int_rangeset_empty : Intrinsic<
+ [llvm_i32_ty], [llvm_i16_ty],
+ [ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, []>]>;
+#endif
+
+#ifdef BAD_BOUNDS
+// BAD-BOUNDS: error: assertion failed: RangeSet requires lower <= upper for each range
+def int_rangeset_bad_bounds : Intrinsic<
+ [llvm_i32_ty], [llvm_i16_ty],
+ [ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, [[4, 3]]>]>;
+#endif
+
+#ifdef OVERLAP
+// OVERLAP: error: assertion failed: RangeSet requires ordered and non-overlapping ranges
+def int_rangeset_overlap : Intrinsic<
+ [llvm_i32_ty], [llvm_i16_ty],
+ [ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, [[0, 2], [2, 4]]>]>;
+#endif
diff --git a/llvm/test/TableGen/intrinsic-rangeset.td b/llvm/test/TableGen/intrinsic-rangeset.td
new file mode 100644
index 0000000000000..a5a3101ce26bb
--- /dev/null
+++ b/llvm/test/TableGen/intrinsic-rangeset.td
@@ -0,0 +1,28 @@
+// RUN: llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s | FileCheck %s
+
+#define TEST_INTRINSICS_SUPPRESS_DEFS
+include "llvm/IR/Intrinsics.td"
+
+def int_range_alias_tst : Intrinsic<
+ [llvm_i32_ty], [llvm_i16_ty],
+ [ImmArg<ArgIndex<0>>, Range<ArgIndex<0>, 0, 9>]>;
+
+def int_rangeset_tst : Intrinsic<
+ [llvm_i32_ty], [llvm_i16_ty],
+ [ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, [[0, 2], [5, 8]]>]>;
+
+def int_rangeset_coalesce_tst : Intrinsic<
+ [llvm_i32_ty], [llvm_i16_ty],
+ [ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, [[0, 2], [3, 4], [6, 8]]>]>;
+
+def int_rangeset_terminal_tst : Intrinsic<
+ [llvm_i32_ty], [],
+ [RangeSet<RetIndex, [[1, 2147483647]]>]>;
+
+// CHECK-DAG: Attribute::get(C, Attribute::Range, ConstantRange(APInt(BitWidth, 0, {{.*}} APInt(BitWidth, 9,
+
+// CHECK-DAG: Attribute::get(C, Attribute::RangeSet, ArrayRef<ConstantRange>{ConstantRange(APInt(BitWidth, 0, {{.*}} APInt(BitWidth, 3, {{.*}} ConstantRange(APInt(BitWidth, 5, {{.*}} APInt(BitWidth, 9,
+
+// CHECK-DAG: Attribute::get(C, Attribute::RangeSet, ArrayRef<ConstantRange>{ConstantRange(APInt(BitWidth, 0, {{.*}} APInt(BitWidth, 5, {{.*}} ConstantRange(APInt(BitWidth, 6, {{.*}} APInt(BitWidth, 9,
+
+// CHECK-DAG: Attribute::get(C, Attribute::RangeSet, ArrayRef<ConstantRange>{ConstantRange(APInt(BitWidth, 1, {{.*}} APInt(BitWidth, 2147483648,
diff --git a/llvm/unittests/IR/AttributesTest.cpp b/llvm/unittests/IR/AttributesTest.cpp
index 5e795ef771713..6ecea71f7d310 100644
--- a/llvm/unittests/IR/AttributesTest.cpp
+++ b/llvm/unittests/IR/AttributesTest.cpp
@@ -466,6 +466,7 @@ TEST(Attributes, SetIntersect) {
V1 = FPClassTest::fcNan;
break;
case Attribute::Range:
+ case Attribute::RangeSet:
break;
case Attribute::Captures:
V0 = CaptureInfo(CaptureComponents::AddressIsNull,
@@ -554,6 +555,14 @@ TEST(Attributes, SetIntersect) {
ASSERT_EQ(Res->getAttribute(Kind).getRange(),
ConstantRange(APInt(32, 0), APInt(32, 20)));
break;
+ case Attribute::RangeSet: {
+ ArrayRef<ConstantRange> RangeSet =
+ Res->getAttribute(Kind).getRangeSet();
+ ASSERT_EQ(RangeSet.size(), 2u);
+ ASSERT_EQ(RangeSet[0], CR0);
+ ASSERT_EQ(RangeSet[1], CR1);
+ break;
+ }
case Attribute::Captures:
ASSERT_EQ(Res->getCaptureInfo(),
CaptureInfo(CaptureComponents::AddressIsNull,
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 76fc2a4f38811..8df3d47400736 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -22,6 +22,7 @@
#include "llvm/TableGen/Record.h"
#include <algorithm>
#include <cassert>
+#include <limits>
using namespace llvm;
// As the type of more than one return values is represented as an anonymous
@@ -31,6 +32,44 @@ using namespace llvm;
// return is 257.
static constexpr unsigned MaxNumReturn = 257;
+static std::pair<int64_t, int64_t>
+getHalfOpenRange(const Record *R, const ListInit *Range, StringRef AttrName) {
+ if (!Range || Range->size() != 2)
+ PrintFatalError(R->getLoc(),
+ Twine(AttrName) + " range must have exactly two bounds");
+
+ auto GetBound = [&](const Init *Bound) -> int64_t {
+ if (const auto *II = dyn_cast<IntInit>(Bound))
+ return II->getValue();
+ PrintFatalError(R->getLoc(),
+ Twine(AttrName) +
+ " bound is not an integer: " + Bound->getAsString());
+ };
+
+ int64_t Lower = GetBound(Range->getElement(0));
+ int64_t Upper = GetBound(Range->getElement(1));
+ if (Upper == std::numeric_limits<int64_t>::max())
+ PrintFatalError(R->getLoc(),
+ Twine(AttrName) + " closed upper bound is too large");
+ return std::pair<int64_t, int64_t>(Lower, Upper + 1);
+}
+
+static void
+appendHalfOpenRange(const Record *R,
+ SmallVectorImpl<std::pair<int64_t, int64_t>> &Ranges,
+ std::pair<int64_t, int64_t> Range) {
+ if (!Ranges.empty()) {
+ if (Range.first < Ranges.back().second)
+ PrintFatalError(R->getLoc(),
+ "RangeSet ranges must be ordered and non-overlapping");
+ if (Range.first == Ranges.back().second) {
+ Ranges.back().second = Range.second;
+ return;
+ }
+ }
+ Ranges.push_back(Range);
+}
+
//===----------------------------------------------------------------------===//
// CodeGenIntrinsic Implementation
//===----------------------------------------------------------------------===//
@@ -566,6 +605,25 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
int64_t Lower = R->getValueAsInt("Lower");
int64_t Upper = R->getValueAsInt("Upper");
addArgAttribute(ArgNo, Range, Lower, Upper);
+ } else if (R->isSubClassOf("RangeSet")) {
+ unsigned ArgNo = R->getValueAsInt("ArgNo");
+ const ListInit *RangeList = R->getValueAsListInit("Ranges");
+ SmallVector<std::pair<int64_t, int64_t>, 4> Ranges;
+
+ if (ArgNo < ArgumentAttributes.size()) {
+ for (const auto &Entry : ArgumentAttributes[ArgNo]) {
+ if (Entry.Kind != RangeSet)
+ continue;
+ PrintFatalError(R->getLoc(),
+ "duplicate RangeSet for the same argument index");
+ }
+ }
+
+ for (const Init *Range : RangeList->getElements())
+ appendHalfOpenRange(
+ R, Ranges,
+ getHalfOpenRange(R, dyn_cast<ListInit>(Range), "RangeSet"));
+ addRangeSetAttribute(ArgNo, std::move(Ranges));
} else if (R->isSubClassOf("ArgInfo")) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
if (ArgNo < 1)
@@ -634,6 +692,15 @@ void CodeGenIntrinsic::addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V,
ArgumentAttributes[Idx].emplace_back(AK, V, V2);
}
+void CodeGenIntrinsic::addRangeSetAttribute(
+ unsigned Idx, SmallVector<std::pair<int64_t, int64_t>, 4> Ranges) {
+ if (Idx >= ArgumentAttributes.size())
+ ArgumentAttributes.resize(Idx + 1);
+ ArgAttribute Attr(RangeSet, 0, 0);
+ Attr.Ranges = std::move(Ranges);
+ ArgumentAttributes[Idx].push_back(std::move(Attr));
+}
+
void CodeGenIntrinsic::addPrettyPrintFunction(unsigned ArgIdx,
StringRef ArgName,
StringRef FuncName) {
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index f17149dea6b07..866f32718d9bb 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -134,19 +134,26 @@ struct CodeGenIntrinsic {
Alignment,
Dereferenceable,
Range,
+ RangeSet,
};
struct ArgAttribute {
ArgAttrKind Kind;
uint64_t Value;
uint64_t Value2;
+ SmallVector<std::pair<int64_t, int64_t>, 4> Ranges;
ArgAttribute(ArgAttrKind K, uint64_t V, uint64_t V2)
: Kind(K), Value(V), Value2(V2) {}
bool operator<(const ArgAttribute &Other) const {
- return std::tie(Kind, Value, Value2) <
- std::tie(Other.Kind, Other.Value, Other.Value2);
+ if (std::tie(Kind, Value, Value2) !=
+ std::tie(Other.Kind, Other.Value, Other.Value2))
+ return std::tie(Kind, Value, Value2) <
+ std::tie(Other.Kind, Other.Value, Other.Value2);
+ return std::lexicographical_compare(Ranges.begin(), Ranges.end(),
+ Other.Ranges.begin(),
+ Other.Ranges.end());
}
};
@@ -155,6 +162,8 @@ struct CodeGenIntrinsic {
void addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V = 0,
uint64_t V2 = 0);
+ void addRangeSetAttribute(unsigned Idx,
+ SmallVector<std::pair<int64_t, int64_t>, 4> Ranges);
/// Structure to store pretty print and argument information.
struct PrettyPrintArgInfo {
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index 23086c422ec1c..b462ba69139b2 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -577,6 +577,8 @@ static StringRef getArgAttrEnumName(CodeGenIntrinsic::ArgAttrKind Kind) {
return "Dereferenceable";
case CodeGenIntrinsic::Range:
return "Range";
+ case CodeGenIntrinsic::RangeSet:
+ return "RangeSet";
}
llvm_unreachable("Unknown CodeGenIntrinsic::ArgAttrKind enum");
}
@@ -632,7 +634,19 @@ static AttributeSet getIntrinsicArgAttributeSet(LLVMContext &C, unsigned ID,
"/*implicitTrunc=*/true), APInt(BitWidth, {}, "
"/*isSigned=*/true, /*implicitTrunc=*/true))),\n",
AttrName, (int64_t)Attr.Value, (int64_t)Attr.Value2);
- else
+ else if (Attr.Kind == CodeGenIntrinsic::RangeSet) {
+ OS << formatv(" Attribute::get(C, Attribute::{}, "
+ "ArrayRef<ConstantRange>{{",
+ AttrName);
+ interleaveComma(Attr.Ranges, OS, [&](auto Range) {
+ OS << formatv("ConstantRange(APInt(BitWidth, {}, "
+ "/*isSigned=*/true, /*implicitTrunc=*/true), "
+ "APInt(BitWidth, {}, /*isSigned=*/true, "
+ "/*implicitTrunc=*/true))",
+ Range.first, Range.second);
+ });
+ OS << "}),\n";
+ } else
OS << formatv(" Attribute::get(C, Attribute::{}),\n", AttrName);
}
OS << " });";
>From 4f2271447aaa5fe119b7d8b819e66fa4ebe0481a Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Fri, 26 Jun 2026 12:48:03 +0200
Subject: [PATCH 2/3] [IR][TableGen] Restrict RangeSet to immarg validation
Keep RangeSet local to intrinsic TableGen definitions instead of
modelling it as a general IR attribute.
---
llvm/include/llvm/AsmParser/LLParser.h | 2 -
llvm/include/llvm/Bitcode/LLVMBitCodes.h | 1 -
llvm/include/llvm/IR/Attributes.h | 10 --
llvm/include/llvm/IR/Attributes.td | 4 -
llvm/include/llvm/IR/Intrinsics.h | 16 ++-
llvm/include/llvm/IR/Intrinsics.td | 8 +-
llvm/lib/AsmParser/LLParser.cpp | 85 ++----------
llvm/lib/Bitcode/Reader/BitcodeReader.cpp | 6 +-
llvm/lib/Bitcode/Writer/BitcodeWriter.cpp | 2 -
.../CodeGen/SelectionDAG/TargetLowering.cpp | 9 +-
llvm/lib/IR/Attributes.cpp | 129 +-----------------
llvm/lib/IR/Intrinsics.cpp | 3 +
llvm/lib/IR/Verifier.cpp | 36 ++---
llvm/test/Assembler/invalid-immarg.ll | 22 +--
.../Assembler/rangeset-attribute-invalid.ll | 36 -----
llvm/test/Assembler/rangeset-attribute.ll | 19 ---
.../TableGen/intrinsic-rangeset-errors.td | 8 ++
llvm/test/TableGen/intrinsic-rangeset.td | 22 ++-
llvm/unittests/IR/AttributesTest.cpp | 9 --
.../TableGen/Basic/CodeGenIntrinsics.cpp | 29 ++--
llvm/utils/TableGen/Basic/CodeGenIntrinsics.h | 23 ++--
.../utils/TableGen/Basic/IntrinsicEmitter.cpp | 67 +++++++--
22 files changed, 154 insertions(+), 392 deletions(-)
delete mode 100644 llvm/test/Assembler/rangeset-attribute-invalid.ll
delete mode 100644 llvm/test/Assembler/rangeset-attribute.ll
diff --git a/llvm/include/llvm/AsmParser/LLParser.h b/llvm/include/llvm/AsmParser/LLParser.h
index 22df21290c015..788b56cb78f08 100644
--- a/llvm/include/llvm/AsmParser/LLParser.h
+++ b/llvm/include/llvm/AsmParser/LLParser.h
@@ -309,7 +309,6 @@ namespace llvm {
Loc = Lex.getLoc();
return parseUInt64(Val);
}
- bool parseAPInt(unsigned BitWidth, APInt &Val);
bool parseFlag(unsigned &Val);
bool parseStringAttribute(AttrBuilder &B);
@@ -413,7 +412,6 @@ namespace llvm {
std::vector<unsigned> &FwdRefAttrGrps,
bool inAttrGrp, LocTy &BuiltinLoc);
bool parseRangeAttr(AttrBuilder &B);
- bool parseRangeSetAttr(AttrBuilder &B);
bool parseInitializesAttr(AttrBuilder &B);
bool parseCapturesAttr(AttrBuilder &B);
bool parseRequiredTypeAttr(AttrBuilder &B, lltok::Kind AttrToken,
diff --git a/llvm/include/llvm/Bitcode/LLVMBitCodes.h b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
index 2d3efb653683a..358f9a65a80af 100644
--- a/llvm/include/llvm/Bitcode/LLVMBitCodes.h
+++ b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
@@ -826,7 +826,6 @@ enum AttributeKindCodes {
ATTR_KIND_NOOUTLINE = 107,
ATTR_KIND_FLATTEN = 108,
ATTR_KIND_NOIPA = 109,
- ATTR_KIND_RANGE_SET = 110,
};
enum ComdatSelectionKindCodes {
diff --git a/llvm/include/llvm/IR/Attributes.h b/llvm/include/llvm/IR/Attributes.h
index 2c157aab122ac..049f699b2980b 100644
--- a/llvm/include/llvm/IR/Attributes.h
+++ b/llvm/include/llvm/IR/Attributes.h
@@ -355,9 +355,6 @@ class Attribute {
/// Returns the value of the range attribute.
LLVM_ABI const ConstantRange &getRange() const;
- /// Returns the value of the rangeset attribute.
- LLVM_ABI ArrayRef<ConstantRange> getRangeSet() const;
-
/// Returns the value of the initializes attribute.
LLVM_ABI ArrayRef<ConstantRange> getInitializes() const;
@@ -1342,9 +1339,6 @@ class AttrBuilder {
/// Add range attribute.
LLVM_ABI AttrBuilder &addRangeAttr(const ConstantRange &CR);
- /// Add rangeset attribute.
- LLVM_ABI AttrBuilder &addRangeSetAttr(ArrayRef<ConstantRange> CRs);
-
/// Add a ConstantRangeList attribute with the given ranges.
LLVM_ABI AttrBuilder &addConstantRangeListAttr(Attribute::AttrKind Kind,
ArrayRef<ConstantRange> Val);
@@ -1375,10 +1369,6 @@ enum AttributeSafetyKind : uint8_t {
/// follows the same type rules as FPMathOperator.
LLVM_ABI bool isNoFPClassCompatibleType(Type *Ty);
-/// Returns true if the ranges are ordered, non-overlapping, and canonical for
-/// the 'rangeset' attribute.
-LLVM_ABI bool isOrderedRangeSet(ArrayRef<ConstantRange> Ranges);
-
/// Which attributes cannot be applied to a type. The argument \p AS
/// is used as a hint for the attributes whose compatibility is being
/// checked against \p Ty. This does not mean the return will be a
diff --git a/llvm/include/llvm/IR/Attributes.td b/llvm/include/llvm/IR/Attributes.td
index 5dbe593a7ba17..4e45100b54d38 100644
--- a/llvm/include/llvm/IR/Attributes.td
+++ b/llvm/include/llvm/IR/Attributes.td
@@ -282,10 +282,6 @@ def Preallocated : TypeAttr<"preallocated", IntersectPreserve, [FnAttr, ParamAtt
/// Parameter or return value is within the specified range.
def Range : ConstantRangeAttr<"range", IntersectCustom, [ParamAttr, RetAttr]>;
-/// Parameter or return value is within one of the specified ranges.
-def RangeSet : ConstantRangeListAttr<"rangeset", IntersectCustom,
- [ParamAttr, RetAttr]>;
-
/// Function does not access memory.
def ReadNone : EnumAttr<"readnone", IntersectAnd, [ParamAttr]>;
diff --git a/llvm/include/llvm/IR/Intrinsics.h b/llvm/include/llvm/IR/Intrinsics.h
index a4799751832bf..1210407e11fad 100644
--- a/llvm/include/llvm/IR/Intrinsics.h
+++ b/llvm/include/llvm/IR/Intrinsics.h
@@ -25,15 +25,16 @@
namespace llvm {
-class Type;
-class FunctionType;
+class APInt;
+class AttributeList;
+class AttributeSet;
+class Constant;
class Function;
+class FunctionType;
class LLVMContext;
class Module;
-class AttributeList;
-class AttributeSet;
class raw_ostream;
-class Constant;
+class Type;
/// This namespace contains an enum with a value for every intrinsic/builtin
/// function known by LLVM. The enum values are returned by
@@ -327,6 +328,11 @@ LLVM_ABI void printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS,
LLVM_ABI void printFPClassMask(raw_ostream &OS, const Constant *ImmArgVal);
+/// Return true if the immediate argument value is accepted by any RangeSet
+/// constraint specified for this intrinsic argument.
+LLVM_ABI bool isImmArgValueInRangeSet(ID IID, unsigned ArgIdx,
+ const APInt &Value);
+
} // namespace Intrinsic
} // namespace llvm
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 1e0b9e17f121f..25add373994e2 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -173,9 +173,9 @@ class ReadNone<ArgIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
}
-// The return value or argument is in the union of the specified closed ranges.
+// The immediate argument is in the union of the specified closed ranges.
// Range entries must be ordered and non-overlapping.
-class RangeSet<AttrIndex idx, list<list<int>> ranges> : IntrinsicProperty {
+class RangeSet<ArgIndex idx, list<list<int>> ranges> : IntrinsicProperty {
list<list<int>> ValidRanges = !filter(r, ranges, !eq(!size(r), 2));
assert !not(!empty(ranges)), "RangeSet requires at least one range";
@@ -194,8 +194,8 @@ class RangeSet<AttrIndex idx, list<list<int>> ranges> : IntrinsicProperty {
// The return value or argument is in the range [lower, upper),
// where lower and upper are interpreted as signed integers.
-class Range<AttrIndex idx, int lower, int upper>
- : RangeSet<idx, [[lower, !sub(upper, 1)]]> {
+class Range<AttrIndex idx, int lower, int upper> : IntrinsicProperty {
+ int ArgNo = idx.Value;
int Lower = lower;
int Upper = upper;
}
diff --git a/llvm/lib/AsmParser/LLParser.cpp b/llvm/lib/AsmParser/LLParser.cpp
index 2a24c2c1a916d..38d10587b104e 100644
--- a/llvm/lib/AsmParser/LLParser.cpp
+++ b/llvm/lib/AsmParser/LLParser.cpp
@@ -1733,8 +1733,6 @@ bool LLParser::parseEnumAttribute(Attribute::AttrKind Attr, AttrBuilder &B,
}
case Attribute::Range:
return parseRangeAttr(B);
- case Attribute::RangeSet:
- return parseRangeSetAttr(B);
case Attribute::Initializes:
return parseInitializesAttr(B);
case Attribute::Captures:
@@ -1991,18 +1989,6 @@ bool LLParser::parseUInt64(uint64_t &Val) {
return false;
}
-/// parseAPInt
-/// ::= apint
-bool LLParser::parseAPInt(unsigned BitWidth, APInt &Val) {
- if (Lex.getKind() != lltok::APSInt)
- return tokError("expected integer");
- if (Lex.getAPSIntVal().getBitWidth() > BitWidth)
- return tokError("integer is too large for the bit width of specified type");
- Val = Lex.getAPSIntVal().extend(BitWidth);
- Lex.Lex();
- return false;
-}
-
/// parseTLSModel
/// := 'localdynamic'
/// := 'initialexec'
@@ -3410,6 +3396,17 @@ bool LLParser::parseRangeAttr(AttrBuilder &B) {
Type *Ty = nullptr;
LocTy TyLoc;
+ auto ParseAPSInt = [&](unsigned BitWidth, APInt &Val) {
+ if (Lex.getKind() != lltok::APSInt)
+ return tokError("expected integer");
+ if (Lex.getAPSIntVal().getBitWidth() > BitWidth)
+ return tokError(
+ "integer is too large for the bit width of specified type");
+ Val = Lex.getAPSIntVal().extend(BitWidth);
+ Lex.Lex();
+ return false;
+ };
+
if (parseToken(lltok::lparen, "expected '('") || parseType(Ty, TyLoc))
return true;
if (!Ty->isIntegerTy())
@@ -3417,8 +3414,8 @@ bool LLParser::parseRangeAttr(AttrBuilder &B) {
unsigned BitWidth = Ty->getPrimitiveSizeInBits();
- if (parseAPInt(BitWidth, Lower) || parseToken(lltok::comma, "expected ','") ||
- parseAPInt(BitWidth, Upper))
+ if (ParseAPSInt(BitWidth, Lower) ||
+ parseToken(lltok::comma, "expected ','") || ParseAPSInt(BitWidth, Upper))
return true;
if (Lower == Upper && !Lower.isZero())
return tokError("the range represent the empty set but limits aren't 0!");
@@ -3430,62 +3427,6 @@ bool LLParser::parseRangeAttr(AttrBuilder &B) {
return false;
}
-/// parseRangeSetAttr
-/// ::= rangeset(<ty> (<n>,<n>)[,(<n>,<n>)]*)
-bool LLParser::parseRangeSetAttr(AttrBuilder &B) {
- Lex.Lex();
-
- Type *Ty = nullptr;
- LocTy TyLoc;
-
- if (parseToken(lltok::lparen, "expected '('"))
- return true;
-
- TyLoc = Lex.getLoc();
- if (Lex.getKind() != lltok::Type)
- return tokError("expected type");
- Ty = Lex.getTyVal();
- Lex.Lex();
-
- if (!Ty->isIntegerTy())
- return error(TyLoc, "the rangeset must have integer type!");
-
- unsigned BitWidth = Ty->getPrimitiveSizeInBits();
-
- SmallVector<ConstantRange, 2> RangeList;
- do {
- APInt Lower, Upper;
- if (parseToken(lltok::lparen, "expected '('") ||
- parseAPInt(BitWidth, Lower) ||
- parseToken(lltok::comma, "expected ','") ||
- parseAPInt(BitWidth, Upper) ||
- parseToken(lltok::rparen, "expected ')'"))
- return true;
-
- if (Upper.slt(Lower))
- return tokError("rangeset requires lower <= upper");
-
- ConstantRange Range = ConstantRange::getNonEmpty(Lower, Upper + 1);
- if (!RangeList.empty() && !RangeList.back().isFullSet() &&
- !RangeList.back().getUpper().isMinSignedValue() &&
- Lower == RangeList.back().getUpper()) {
- RangeList.back() = ConstantRange::getNonEmpty(RangeList.back().getLower(),
- Range.getUpper());
- continue;
- }
- RangeList.push_back(Range);
- } while (EatIfPresent(lltok::comma));
-
- if (parseToken(lltok::rparen, "expected ')'"))
- return true;
-
- if (!AttributeFuncs::isOrderedRangeSet(RangeList))
- return tokError("Invalid (unordered or overlapping) range set");
-
- B.addRangeSetAttr(RangeList);
- return false;
-}
-
/// parseInitializesAttr
/// ::= initializes((Lo1,Hi1),(Lo2,Hi2),...)
bool LLParser::parseInitializesAttr(AttrBuilder &B) {
diff --git a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
index cec11303ca40b..a5a11ab6221d4 100644
--- a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
+++ b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
@@ -2304,8 +2304,6 @@ static Attribute::AttrKind getAttrFromCode(uint64_t Code) {
return Attribute::DeadOnUnwind;
case bitc::ATTR_KIND_RANGE:
return Attribute::Range;
- case bitc::ATTR_KIND_RANGE_SET:
- return Attribute::RangeSet;
case bitc::ATTR_KIND_INITIALIZES:
return Attribute::Initializes;
case bitc::ATTR_KIND_CORO_ELIDE_SAFE:
@@ -2584,9 +2582,7 @@ Error BitcodeReader::parseAttributeGroupBlock() {
}
i--;
- if (Kind == Attribute::RangeSet
- ? !AttributeFuncs::isOrderedRangeSet(Val)
- : !ConstantRangeList::isOrderedRanges(Val))
+ if (!ConstantRangeList::isOrderedRanges(Val))
return error("Invalid (unordered or overlapping) range list");
B.addConstantRangeListAttr(Kind, Val);
} else {
diff --git a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
index 210ba1191526a..a6e7941a06699 100644
--- a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
+++ b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
@@ -998,8 +998,6 @@ static uint64_t getAttrKindEncoding(Attribute::AttrKind Kind) {
return bitc::ATTR_KIND_DEAD_ON_UNWIND;
case Attribute::Range:
return bitc::ATTR_KIND_RANGE;
- case Attribute::RangeSet:
- return bitc::ATTR_KIND_RANGE_SET;
case Attribute::Initializes:
return bitc::ATTR_KIND_INITIALIZES;
case Attribute::NoExt:
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index 8fc28ea2b8b0e..bca34c5c347ee 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -70,11 +70,10 @@ bool TargetLowering::isInTailCallPosition(SelectionDAG &DAG, SDNode *Node,
// the return. Ignore following attributes because they don't affect the
// call sequence.
AttrBuilder CallerAttrs(F.getContext(), F.getAttributes().getRetAttrs());
- for (const auto &Attr :
- {Attribute::Alignment, Attribute::Dereferenceable,
- Attribute::DereferenceableOrNull, Attribute::NoAlias,
- Attribute::NonNull, Attribute::NoUndef, Attribute::Range,
- Attribute::RangeSet, Attribute::NoFPClass})
+ for (const auto &Attr : {Attribute::Alignment, Attribute::Dereferenceable,
+ Attribute::DereferenceableOrNull, Attribute::NoAlias,
+ Attribute::NonNull, Attribute::NoUndef,
+ Attribute::Range, Attribute::NoFPClass})
CallerAttrs.removeAttribute(Attr);
if (CallerAttrs.hasAttributes())
diff --git a/llvm/lib/IR/Attributes.cpp b/llvm/lib/IR/Attributes.cpp
index dec94c313c884..4087b25951a1c 100644
--- a/llvm/lib/IR/Attributes.cpp
+++ b/llvm/lib/IR/Attributes.cpp
@@ -92,64 +92,6 @@ unpackVScaleRangeArgs(uint64_t Value) {
MaxValue > 0 ? MaxValue : std::optional<unsigned>());
}
-static void appendRangeSetRange(SmallVectorImpl<ConstantRange> &Ranges,
- const ConstantRange &CR) {
- if (CR.isEmptySet())
- return;
- if (CR.isFullSet()) {
- Ranges.assign(1, CR);
- return;
- }
- if (Ranges.empty()) {
- Ranges.push_back(CR);
- return;
- }
- if (Ranges.front().isFullSet())
- return;
-
- APInt PreviousUpper = Ranges.back().getUpper();
- if (PreviousUpper.isMinSignedValue())
- return;
- if (CR.getLower().sle(PreviousUpper)) {
- APInt NewUpper = CR.getUpper().isMinSignedValue()
- ? APInt::getSignedMinValue(CR.getBitWidth())
- : APIntOps::smax(PreviousUpper, CR.getUpper());
- Ranges.back() = ConstantRange(Ranges.back().getLower(), NewUpper);
- if (Ranges.back().isFullSet())
- Ranges.truncate(1);
- return;
- }
-
- Ranges.push_back(CR);
-}
-
-static SmallVector<ConstantRange, 2>
-unionRangeSetLists(ArrayRef<ConstantRange> LHS, ArrayRef<ConstantRange> RHS) {
- assert(AttributeFuncs::isOrderedRangeSet(LHS) &&
- AttributeFuncs::isOrderedRangeSet(RHS) &&
- "rangeset attribute must be ordered");
- if (LHS.empty())
- return SmallVector<ConstantRange, 2>(RHS);
- if (RHS.empty())
- return SmallVector<ConstantRange, 2>(LHS);
-
- SmallVector<ConstantRange, 2> Result;
- size_t LHSIdx = 0, RHSIdx = 0;
- while (LHSIdx < LHS.size() || RHSIdx < RHS.size()) {
- const ConstantRange *NextRange;
- if (RHSIdx == RHS.size() ||
- (LHSIdx < LHS.size() &&
- LHS[LHSIdx].getLower().slt(RHS[RHSIdx].getLower())))
- NextRange = &LHS[LHSIdx++];
- else
- NextRange = &RHS[RHSIdx++];
- appendRangeSetRange(Result, *NextRange);
- if (Result.size() == 1 && Result.front().isFullSet())
- break;
- }
- return Result;
-}
-
Attribute Attribute::get(LLVMContext &Context, Attribute::AttrKind Kind,
uint64_t Val) {
bool IsIntAttr = Attribute::isIntAttrKind(Kind);
@@ -587,12 +529,6 @@ const ConstantRange &Attribute::getRange() const {
return pImpl->getValueAsConstantRange();
}
-ArrayRef<ConstantRange> Attribute::getRangeSet() const {
- assert(hasAttribute(Attribute::RangeSet) &&
- "Trying to get rangeset attr from non-ConstantRangeList attribute");
- return pImpl->getValueAsConstantRangeList();
-}
-
ArrayRef<ConstantRange> Attribute::getInitializes() const {
assert(hasAttribute(Attribute::Initializes) &&
"Trying to get initializes attr from non-ConstantRangeList attribute");
@@ -803,21 +739,6 @@ std::string Attribute::getAsString(bool InAttrGrp) const {
return Result;
}
- if (hasAttribute(Attribute::RangeSet)) {
- std::string Result;
- raw_string_ostream OS(Result);
- ArrayRef<ConstantRange> Ranges = getRangeSet();
- assert(!Ranges.empty() &&
- "rangeset attribute must have at least one range");
- OS << "rangeset(";
- OS << "i" << Ranges.front().getBitWidth() << " ";
- interleaveComma(Ranges, OS, [&](const ConstantRange &CR) {
- OS << "(" << CR.getLower() << ", " << CR.getUpper() - 1 << ")";
- });
- OS << ")";
- return Result;
- }
-
if (hasAttribute(Attribute::Initializes)) {
std::string Result;
raw_string_ostream OS(Result);
@@ -1229,11 +1150,6 @@ AttributeSet::intersectWith(LLVMContext &C, AttributeSet Other) const {
if (!NewRange.isFullSet())
Intersected.addRangeAttr(NewRange);
} break;
- case Attribute::RangeSet: {
- SmallVector<ConstantRange> NewRangeSet =
- unionRangeSetLists(Attr0.getRangeSet(), Attr1.getRangeSet());
- Intersected.addRangeSetAttr(NewRangeSet);
- } break;
default:
llvm_unreachable("Unknown attribute with custom intersection rule");
}
@@ -2455,17 +2371,9 @@ AttrBuilder &AttrBuilder::addRangeAttr(const ConstantRange &CR) {
return addConstantRangeAttr(Attribute::Range, CR);
}
-AttrBuilder &AttrBuilder::addRangeSetAttr(ArrayRef<ConstantRange> CRs) {
- if (CRs.empty() || (CRs.size() == 1 && CRs.front().isFullSet()))
- return *this;
- return addAttribute(Attribute::get(Ctx, Attribute::RangeSet, CRs));
-}
-
AttrBuilder &
AttrBuilder::addConstantRangeListAttr(Attribute::AttrKind Kind,
ArrayRef<ConstantRange> Val) {
- if (Kind == Attribute::RangeSet)
- return addRangeSetAttr(Val);
return addAttribute(Attribute::get(Ctx, Kind, Val));
}
@@ -2571,33 +2479,6 @@ bool AttributeFuncs::isNoFPClassCompatibleType(Type *Ty) {
return FPMathOperator::isSupportedFloatingPointType(Ty);
}
-bool AttributeFuncs::isOrderedRangeSet(ArrayRef<ConstantRange> Ranges) {
- if (Ranges.empty())
- return true;
-
- for (unsigned I = 0; I < Ranges.size(); ++I) {
- const ConstantRange &Range = Ranges[I];
- if (Range.isEmptySet())
- return false;
- if (I != 0 && Range.getBitWidth() != Ranges.front().getBitWidth())
- return false;
- if (Range.isFullSet())
- return Ranges.size() == 1;
- if (!Range.getLower().slt(Range.getUpper()) &&
- !Range.getUpper().isMinSignedValue())
- return false;
- if (I == 0)
- continue;
-
- const ConstantRange &PrevRange = Ranges[I - 1];
- if (PrevRange.isFullSet() || PrevRange.getUpper().isMinSignedValue())
- return false;
- if (Range.getLower().sle(PrevRange.getUpper()))
- return false;
- }
- return true;
-}
-
/// Which attributes cannot be applied to a type.
AttributeMask AttributeFuncs::typeIncompatible(Type *Ty, AttributeSet AS,
AttributeSafetyKind ASK) {
@@ -2618,20 +2499,12 @@ AttributeMask AttributeFuncs::typeIncompatible(Type *Ty, AttributeSet AS,
if (!Ty->isIntOrIntVectorTy()) {
// Attributes that only apply to integers or vector of integers.
if (ASK & ASK_SAFE_TO_DROP)
- Incompatible.addAttribute(Attribute::Range)
- .addAttribute(Attribute::RangeSet);
+ Incompatible.addAttribute(Attribute::Range);
} else {
Attribute RangeAttr = AS.getAttribute(Attribute::Range);
if (RangeAttr.isValid() &&
RangeAttr.getRange().getBitWidth() != Ty->getScalarSizeInBits())
Incompatible.addAttribute(Attribute::Range);
- Attribute RangeSetAttr = AS.getAttribute(Attribute::RangeSet);
- if (RangeSetAttr.isValid()) {
- ArrayRef<ConstantRange> Ranges = RangeSetAttr.getRangeSet();
- if (!Ranges.empty() &&
- Ranges.front().getBitWidth() != Ty->getScalarSizeInBits())
- Incompatible.addAttribute(Attribute::RangeSet);
- }
}
if (!Ty->isPointerTy()) {
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index 266af8e06a230..fefd407a5556b 100644
--- a/llvm/lib/IR/Intrinsics.cpp
+++ b/llvm/lib/IR/Intrinsics.cpp
@@ -1478,6 +1478,9 @@ LLVM_ABI void Intrinsic::printFPClassMask(raw_ostream &OS,
OS << static_cast<FPClassTest>(Val);
}
+#define GET_INTRINSIC_IMMARG_RANGE_SET_CHECKS
+#include "llvm/IR/IntrinsicImpl.inc"
+
#define GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
#include "llvm/IR/IntrinsicImpl.inc"
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index b2c0e42b0b54e..39e8b1a18cdc4 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -2036,12 +2036,11 @@ void Verifier::verifyParameterAttrs(AttributeSet Attrs, Type *Ty,
V);
if (Attrs.hasAttribute(Attribute::ImmArg)) {
- unsigned AttrCount = Attrs.getNumAttributes() -
- Attrs.hasAttribute(Attribute::Range) -
- Attrs.hasAttribute(Attribute::RangeSet);
+ unsigned AttrCount =
+ Attrs.getNumAttributes() - Attrs.hasAttribute(Attribute::Range);
Check(AttrCount == 1,
"Attribute 'immarg' is incompatible with other attributes except the "
- "'range' and 'rangeset' attributes",
+ "'range' attribute",
V);
}
@@ -2186,16 +2185,6 @@ void Verifier::verifyParameterAttrs(AttributeSet Attrs, Type *Ty,
Check(Ty->isIntOrIntVectorTy(CR.getBitWidth()),
"Range bit width must match type bit width!", V);
}
- if (Attrs.hasAttribute(Attribute::RangeSet)) {
- ArrayRef<ConstantRange> Ranges =
- Attrs.getAttribute(Attribute::RangeSet).getRangeSet();
- Check(!Ranges.empty(), "Attribute 'rangeset' does not support empty list",
- V);
- Check(AttributeFuncs::isOrderedRangeSet(Ranges),
- "Attribute 'rangeset' does not support unordered ranges", V);
- Check(Ty->isIntOrIntVectorTy(Ranges.front().getBitWidth()),
- "RangeSet bit width must match type bit width!", V);
- }
}
void Verifier::checkUnsignedBaseTenFuncAttr(AttributeList Attrs, StringRef Attr,
@@ -3919,19 +3908,12 @@ void Verifier::visitCallBase(CallBase &Call) {
Call);
}
}
- if (Call.paramHasAttr(i, Attribute::RangeSet)) {
- if (auto *CI = dyn_cast<ConstantInt>(ArgVal)) {
- ArrayRef<ConstantRange> Ranges =
- Call.getParamAttr(i, Attribute::RangeSet).getRangeSet();
- bool IsContained = llvm::any_of(Ranges, [&](const ConstantRange &CR) {
- return CR.contains(CI->getValue());
- });
- Check(IsContained,
- "immarg value " + Twine(CI->getValue().getSExtValue()) +
- " out of rangeset",
- Call);
- }
- }
+ if (auto *CI = dyn_cast<ConstantInt>(ArgVal))
+ Check(Intrinsic::isImmArgValueInRangeSet(Call.getIntrinsicID(), i,
+ CI->getValue()),
+ "immarg value " + Twine(CI->getValue().getSExtValue()) +
+ " out of rangeset",
+ Call);
}
if (Call.paramHasAttr(i, Attribute::Preallocated)) {
diff --git a/llvm/test/Assembler/invalid-immarg.ll b/llvm/test/Assembler/invalid-immarg.ll
index c04f1d74b882c..df69a66c67332 100644
--- a/llvm/test/Assembler/invalid-immarg.ll
+++ b/llvm/test/Assembler/invalid-immarg.ll
@@ -1,34 +1,34 @@
; RUN: not llvm-as < %s -o /dev/null 2>&1 | FileCheck %s
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.byval(ptr byval(i32) immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.inalloca(ptr inalloca(i32) immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.inreg(i32 inreg immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.nest(ptr nest immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.sret(ptr sret(i32) immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.zeroext(i32 zeroext immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.signext(i32 signext immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.returned(i32 returned immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.noalias(ptr noalias immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.readnone(ptr readnone immarg)
-; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' and 'rangeset' attributes
+; CHECK: Attribute 'immarg' is incompatible with other attributes except the 'range' attribute
declare void @llvm.immarg.readonly(ptr readonly immarg)
diff --git a/llvm/test/Assembler/rangeset-attribute-invalid.ll b/llvm/test/Assembler/rangeset-attribute-invalid.ll
deleted file mode 100644
index 98605a1ac4b3a..0000000000000
--- a/llvm/test/Assembler/rangeset-attribute-invalid.ll
+++ /dev/null
@@ -1,36 +0,0 @@
-; RUN: split-file %s %t
-; RUN: not llvm-as < %t/non_integer_type.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=NON-INTEGER-TYPE
-; RUN: not llvm-as < %t/missing_type.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=MISSING-TYPE
-; RUN: not llvm-as < %t/empty.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=EMPTY
-; RUN: not llvm-as < %t/bad_bounds.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=BAD-BOUNDS
-; RUN: not llvm-as < %t/overlap.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=OVERLAP
-; RUN: not llvm-as < %t/unordered.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=UNORDERED
-; RUN: not llvm-as < %t/missing_comma.ll -o /dev/null 2>&1 | FileCheck %s --check-prefix=MISSING-COMMA
-
-;--- non_integer_type.ll
-; NON-INTEGER-TYPE: the rangeset must have integer type!
-declare rangeset(float (0, 2)) float @non_integer_type()
-
-;--- missing_type.ll
-; MISSING-TYPE: expected type
-declare rangeset((0, 2)) i32 @missing_type()
-
-;--- empty.ll
-; EMPTY: expected '('
-declare rangeset(i32) i32 @empty()
-
-;--- bad_bounds.ll
-; BAD-BOUNDS: rangeset requires lower <= upper
-declare rangeset(i32 (7, 5)) i32 @bad_bounds()
-
-;--- overlap.ll
-; OVERLAP: Invalid (unordered or overlapping) range set
-declare rangeset(i32 (0, 2), (2, 4)) i32 @overlap()
-
-;--- unordered.ll
-; UNORDERED: Invalid (unordered or overlapping) range set
-declare rangeset(i32 (5, 8), (0, 2)) i32 @unordered()
-
-;--- missing_comma.ll
-; MISSING-COMMA: expected ','
-declare rangeset(i32 (0 2)) i32 @missing_comma()
diff --git a/llvm/test/Assembler/rangeset-attribute.ll b/llvm/test/Assembler/rangeset-attribute.ll
deleted file mode 100644
index a47ab51040fb2..0000000000000
--- a/llvm/test/Assembler/rangeset-attribute.ll
+++ /dev/null
@@ -1,19 +0,0 @@
-; RUN: llvm-as < %s | llvm-dis | FileCheck %s
-
-; CHECK: declare rangeset(i32 (0, 2), (5, 8)) i32 @ret_rangeset()
-declare rangeset(i32 (0, 2), (5, 8)) i32 @ret_rangeset()
-
-; CHECK: declare void @param_rangeset(i32 rangeset(i32 (0, 2), (5, 8)))
-declare void @param_rangeset(i32 rangeset(i32 (0, 2), (5, 8)))
-
-; CHECK: declare rangeset(i32 (0, 2), (5, 8)) <4 x i32> @ret_rangeset_vec()
-declare rangeset(i32 (0, 2), (5, 8)) <4 x i32> @ret_rangeset_vec()
-
-; CHECK: declare void @param_rangeset_vec(<4 x i32> rangeset(i32 (0, 2), (5, 8)))
-declare void @param_rangeset_vec(<4 x i32> rangeset(i32 (0, 2), (5, 8)))
-
-; CHECK: declare rangeset(i32 (1, 2147483647)) i32 @ret_rangeset_terminal()
-declare rangeset(i32 (1, 2147483647)) i32 @ret_rangeset_terminal()
-
-; CHECK: declare rangeset(i32 (0, 4), (6, 8)) i32 @ret_rangeset_coalesced()
-declare rangeset(i32 (0, 2), (3, 4), (6, 8)) i32 @ret_rangeset_coalesced()
diff --git a/llvm/test/TableGen/intrinsic-rangeset-errors.td b/llvm/test/TableGen/intrinsic-rangeset-errors.td
index 795220cb3d00b..7a6629bde2485 100644
--- a/llvm/test/TableGen/intrinsic-rangeset-errors.td
+++ b/llvm/test/TableGen/intrinsic-rangeset-errors.td
@@ -1,6 +1,7 @@
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DEMPTY %s 2>&1 | FileCheck --check-prefix=EMPTY %s -DFILE=%s
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DBAD_BOUNDS %s 2>&1 | FileCheck --check-prefix=BAD-BOUNDS %s -DFILE=%s
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DOVERLAP %s 2>&1 | FileCheck --check-prefix=OVERLAP %s -DFILE=%s
+// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include -DMISSING_IMMARG %s 2>&1 | FileCheck --check-prefix=MISSING-IMMARG %s -DFILE=%s
#define TEST_INTRINSICS_SUPPRESS_DEFS
include "llvm/IR/Intrinsics.td"
@@ -25,3 +26,10 @@ def int_rangeset_overlap : Intrinsic<
[llvm_i32_ty], [llvm_i16_ty],
[ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, [[0, 2], [2, 4]]>]>;
#endif
+
+#ifdef MISSING_IMMARG
+// MISSING-IMMARG: error: RangeSet requires ImmArg for the same argument index
+def int_rangeset_missing_immarg : Intrinsic<
+ [llvm_i32_ty], [llvm_i16_ty],
+ [RangeSet<ArgIndex<0>, [[0, 2]]>]>;
+#endif
diff --git a/llvm/test/TableGen/intrinsic-rangeset.td b/llvm/test/TableGen/intrinsic-rangeset.td
index a5a3101ce26bb..b14d70e26b062 100644
--- a/llvm/test/TableGen/intrinsic-rangeset.td
+++ b/llvm/test/TableGen/intrinsic-rangeset.td
@@ -3,10 +3,6 @@
#define TEST_INTRINSICS_SUPPRESS_DEFS
include "llvm/IR/Intrinsics.td"
-def int_range_alias_tst : Intrinsic<
- [llvm_i32_ty], [llvm_i16_ty],
- [ImmArg<ArgIndex<0>>, Range<ArgIndex<0>, 0, 9>]>;
-
def int_rangeset_tst : Intrinsic<
[llvm_i32_ty], [llvm_i16_ty],
[ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, [[0, 2], [5, 8]]>]>;
@@ -15,14 +11,16 @@ def int_rangeset_coalesce_tst : Intrinsic<
[llvm_i32_ty], [llvm_i16_ty],
[ImmArg<ArgIndex<0>>, RangeSet<ArgIndex<0>, [[0, 2], [3, 4], [6, 8]]>]>;
-def int_rangeset_terminal_tst : Intrinsic<
- [llvm_i32_ty], [],
- [RangeSet<RetIndex, [[1, 2147483647]]>]>;
-
-// CHECK-DAG: Attribute::get(C, Attribute::Range, ConstantRange(APInt(BitWidth, 0, {{.*}} APInt(BitWidth, 9,
+// CHECK-NOT: Attribute::RangeSet
-// CHECK-DAG: Attribute::get(C, Attribute::RangeSet, ArrayRef<ConstantRange>{ConstantRange(APInt(BitWidth, 0, {{.*}} APInt(BitWidth, 3, {{.*}} ConstantRange(APInt(BitWidth, 5, {{.*}} APInt(BitWidth, 9,
+// CHECK: bool Intrinsic::isImmArgValueInRangeSet
-// CHECK-DAG: Attribute::get(C, Attribute::RangeSet, ArrayRef<ConstantRange>{ConstantRange(APInt(BitWidth, 0, {{.*}} APInt(BitWidth, 5, {{.*}} ConstantRange(APInt(BitWidth, 6, {{.*}} APInt(BitWidth, 9,
+// CHECK: case rangeset_coalesce_tst:
+// CHECK: case 0:
+// CHECK: APInt(Value.getBitWidth(), 0, {{.*}}APInt(Value.getBitWidth(), 5,
+// CHECK: APInt(Value.getBitWidth(), 6, {{.*}}APInt(Value.getBitWidth(), 9,
-// CHECK-DAG: Attribute::get(C, Attribute::RangeSet, ArrayRef<ConstantRange>{ConstantRange(APInt(BitWidth, 1, {{.*}} APInt(BitWidth, 2147483648,
+// CHECK: case rangeset_tst:
+// CHECK: case 0:
+// CHECK: APInt(Value.getBitWidth(), 0, {{.*}}APInt(Value.getBitWidth(), 3,
+// CHECK: APInt(Value.getBitWidth(), 5, {{.*}}APInt(Value.getBitWidth(), 9,
diff --git a/llvm/unittests/IR/AttributesTest.cpp b/llvm/unittests/IR/AttributesTest.cpp
index 6ecea71f7d310..5e795ef771713 100644
--- a/llvm/unittests/IR/AttributesTest.cpp
+++ b/llvm/unittests/IR/AttributesTest.cpp
@@ -466,7 +466,6 @@ TEST(Attributes, SetIntersect) {
V1 = FPClassTest::fcNan;
break;
case Attribute::Range:
- case Attribute::RangeSet:
break;
case Attribute::Captures:
V0 = CaptureInfo(CaptureComponents::AddressIsNull,
@@ -555,14 +554,6 @@ TEST(Attributes, SetIntersect) {
ASSERT_EQ(Res->getAttribute(Kind).getRange(),
ConstantRange(APInt(32, 0), APInt(32, 20)));
break;
- case Attribute::RangeSet: {
- ArrayRef<ConstantRange> RangeSet =
- Res->getAttribute(Kind).getRangeSet();
- ASSERT_EQ(RangeSet.size(), 2u);
- ASSERT_EQ(RangeSet[0], CR0);
- ASSERT_EQ(RangeSet[1], CR1);
- break;
- }
case Attribute::Captures:
ASSERT_EQ(Res->getCaptureInfo(),
CaptureInfo(CaptureComponents::AddressIsNull,
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 8df3d47400736..a09fae9663b86 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -427,6 +427,11 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
for (auto &Attrs : ArgumentAttributes)
llvm::sort(Attrs);
+ for (const ImmArgRangeSet &RangeSet : ImmArgRangeSets)
+ if (!isParamImmArg(RangeSet.ArgNo))
+ PrintFatalError(TheDef->getLoc(),
+ "RangeSet requires ImmArg for the same argument index");
+
// Default values are not yet supported for overloaded intrinsics
// (overloaded support will come in a follow-up).
if (isOverloaded &&
@@ -610,20 +615,20 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
const ListInit *RangeList = R->getValueAsListInit("Ranges");
SmallVector<std::pair<int64_t, int64_t>, 4> Ranges;
- if (ArgNo < ArgumentAttributes.size()) {
- for (const auto &Entry : ArgumentAttributes[ArgNo]) {
- if (Entry.Kind != RangeSet)
- continue;
+ if (ArgNo < 1)
+ PrintFatalError(R->getLoc(), "RangeSet should only apply to arguments");
+ unsigned ParamNo = ArgNo - 1;
+
+ for (const ImmArgRangeSet &RangeSet : ImmArgRangeSets)
+ if (RangeSet.ArgNo == ParamNo)
PrintFatalError(R->getLoc(),
"duplicate RangeSet for the same argument index");
- }
- }
for (const Init *Range : RangeList->getElements())
appendHalfOpenRange(
R, Ranges,
getHalfOpenRange(R, dyn_cast<ListInit>(Range), "RangeSet"));
- addRangeSetAttribute(ArgNo, std::move(Ranges));
+ addImmArgRangeSet(ParamNo, std::move(Ranges));
} else if (R->isSubClassOf("ArgInfo")) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
if (ArgNo < 1)
@@ -692,13 +697,9 @@ void CodeGenIntrinsic::addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V,
ArgumentAttributes[Idx].emplace_back(AK, V, V2);
}
-void CodeGenIntrinsic::addRangeSetAttribute(
- unsigned Idx, SmallVector<std::pair<int64_t, int64_t>, 4> Ranges) {
- if (Idx >= ArgumentAttributes.size())
- ArgumentAttributes.resize(Idx + 1);
- ArgAttribute Attr(RangeSet, 0, 0);
- Attr.Ranges = std::move(Ranges);
- ArgumentAttributes[Idx].push_back(std::move(Attr));
+void CodeGenIntrinsic::addImmArgRangeSet(
+ unsigned ArgNo, SmallVector<std::pair<int64_t, int64_t>, 4> Ranges) {
+ ImmArgRangeSets.push_back({ArgNo, std::move(Ranges)});
}
void CodeGenIntrinsic::addPrettyPrintFunction(unsigned ArgIdx,
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index 866f32718d9bb..ede0811232eb3 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -134,36 +134,37 @@ struct CodeGenIntrinsic {
Alignment,
Dereferenceable,
Range,
- RangeSet,
};
struct ArgAttribute {
ArgAttrKind Kind;
uint64_t Value;
uint64_t Value2;
- SmallVector<std::pair<int64_t, int64_t>, 4> Ranges;
ArgAttribute(ArgAttrKind K, uint64_t V, uint64_t V2)
: Kind(K), Value(V), Value2(V2) {}
bool operator<(const ArgAttribute &Other) const {
- if (std::tie(Kind, Value, Value2) !=
- std::tie(Other.Kind, Other.Value, Other.Value2))
- return std::tie(Kind, Value, Value2) <
- std::tie(Other.Kind, Other.Value, Other.Value2);
- return std::lexicographical_compare(Ranges.begin(), Ranges.end(),
- Other.Ranges.begin(),
- Other.Ranges.end());
+ return std::tie(Kind, Value, Value2) <
+ std::tie(Other.Kind, Other.Value, Other.Value2);
}
};
+ struct ImmArgRangeSet {
+ unsigned ArgNo;
+ SmallVector<std::pair<int64_t, int64_t>, 4> Ranges;
+ };
+
/// Vector of attributes for each argument.
SmallVector<SmallVector<ArgAttribute, 0>> ArgumentAttributes;
+ /// Range-set constraints for immediate arguments, indexed by argument number.
+ SmallVector<ImmArgRangeSet> ImmArgRangeSets;
+
void addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V = 0,
uint64_t V2 = 0);
- void addRangeSetAttribute(unsigned Idx,
- SmallVector<std::pair<int64_t, int64_t>, 4> Ranges);
+ void addImmArgRangeSet(unsigned ArgNo,
+ SmallVector<std::pair<int64_t, int64_t>, 4> Ranges);
/// Structure to store pretty print and argument information.
struct PrettyPrintArgInfo {
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index b462ba69139b2..12616975308a9 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -73,6 +73,8 @@ class IntrinsicEmitter {
function_ref<bool(const CodeGenIntrinsic &Int)> GetProperty);
void EmitGenerator(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
void EmitAttributes(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
+ void EmitImmArgRangeSetChecks(const CodeGenIntrinsicTable &Ints,
+ raw_ostream &OS);
void EmitPrettyPrintArguments(const CodeGenIntrinsicTable &Ints,
raw_ostream &OS);
void EmitDefaultArgValuesTable(const CodeGenIntrinsicTable &Ints,
@@ -130,6 +132,9 @@ void IntrinsicEmitter::run(raw_ostream &OS, bool Enums) {
// Emit the intrinsic parameter attributes.
EmitAttributes(Ints, OS);
+ // Emit immediate argument range-set checks.
+ EmitImmArgRangeSetChecks(Ints, OS);
+
// Emit the intrinsic ID -> pretty print table.
EmitIntrinsicToPrettyPrintTable(Ints, OS);
@@ -577,8 +582,6 @@ static StringRef getArgAttrEnumName(CodeGenIntrinsic::ArgAttrKind Kind) {
return "Dereferenceable";
case CodeGenIntrinsic::Range:
return "Range";
- case CodeGenIntrinsic::RangeSet:
- return "RangeSet";
}
llvm_unreachable("Unknown CodeGenIntrinsic::ArgAttrKind enum");
}
@@ -634,19 +637,7 @@ static AttributeSet getIntrinsicArgAttributeSet(LLVMContext &C, unsigned ID,
"/*implicitTrunc=*/true), APInt(BitWidth, {}, "
"/*isSigned=*/true, /*implicitTrunc=*/true))),\n",
AttrName, (int64_t)Attr.Value, (int64_t)Attr.Value2);
- else if (Attr.Kind == CodeGenIntrinsic::RangeSet) {
- OS << formatv(" Attribute::get(C, Attribute::{}, "
- "ArrayRef<ConstantRange>{{",
- AttrName);
- interleaveComma(Attr.Ranges, OS, [&](auto Range) {
- OS << formatv("ConstantRange(APInt(BitWidth, {}, "
- "/*isSigned=*/true, /*implicitTrunc=*/true), "
- "APInt(BitWidth, {}, /*isSigned=*/true, "
- "/*implicitTrunc=*/true))",
- Range.first, Range.second);
- });
- OS << "}),\n";
- } else
+ else
OS << formatv(" Attribute::get(C, Attribute::{}),\n", AttrName);
}
OS << " });";
@@ -900,6 +891,52 @@ AttributeSet Intrinsic::getFnAttributes(LLVMContext &C, ID id) {{
NoFunctionAttrsID);
}
+void IntrinsicEmitter::EmitImmArgRangeSetChecks(
+ const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
+ IfDefEmitter IfDef(OS, "GET_INTRINSIC_IMMARG_RANGE_SET_CHECKS");
+ OS << R"(
+bool Intrinsic::isImmArgValueInRangeSet(ID IID, unsigned ArgIdx,
+ const APInt &Value) {
+ using namespace Intrinsic;
+ switch (IID) {
+)";
+
+ for (const CodeGenIntrinsic &Int : Ints) {
+ if (Int.ImmArgRangeSets.empty())
+ continue;
+
+ OS << " case " << Int.EnumName << ":\n";
+ OS << " switch (ArgIdx) {\n";
+ for (const CodeGenIntrinsic::ImmArgRangeSet &RangeSet :
+ Int.ImmArgRangeSets) {
+ OS << " case " << RangeSet.ArgNo << ":\n";
+ OS << " return ";
+ interleave(
+ RangeSet.Ranges,
+ [&](const std::pair<int64_t, int64_t> &Range) {
+ OS << formatv("ConstantRange(APInt(Value.getBitWidth(), {}, "
+ "/*isSigned=*/true, /*implicitTrunc=*/true), "
+ "APInt(Value.getBitWidth(), {}, "
+ "/*isSigned=*/true, "
+ "/*implicitTrunc=*/true)).contains(Value)",
+ Range.first, Range.second);
+ },
+ [&] { OS << " ||\n "; });
+ OS << ";\n";
+ }
+ OS << R"( default:
+ return true;
+ }
+)";
+ }
+
+ OS << R"( default:
+ return true;
+ }
+}
+)";
+}
+
void IntrinsicEmitter::EmitIntrinsicToPrettyPrintTable(
const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
EmitIntrinsicBitTable(Ints, OS, "GET_INTRINSIC_PRETTY_PRINT_TABLE", "PPTable",
>From ad5b31e8370ac356e9391efba4ea9fd4e6fc3a0d Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Fri, 10 Jul 2026 11:00:57 +0200
Subject: [PATCH 3/3] [NVVM][NVPTX] Add cp.async.bulk.reduce.s2g.<type>
intrinsics
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 36 ++++++
llvm/include/llvm/IR/NVVMIntrinsicUtils.h | 30 +++--
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp | 104 +++++++++++++----
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.h | 5 +-
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 110 +++++++++++++++---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h | 1 +
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 44 +++++--
.../cp-async-bulk-reduce-s2g-ptx92-invalid.ll | 11 ++
.../NVPTX/cp-async-bulk-reduce-s2g-ptx92.ll | 27 +++++
.../CodeGen/NVPTX/cp-async-bulk-reduce-s2g.ll | 65 +++++++++++
.../NVPTX/cp-async-bulk-reduce-s2g-invalid.ll | 31 +++++
11 files changed, 401 insertions(+), 63 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92-invalid.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92.ll
create mode 100644 llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g.ll
create mode 100644 llvm/test/Verifier/NVPTX/cp-async-bulk-reduce-s2g-invalid.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 40f3c40a76bc6..dd47c57691c9f 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -3031,6 +3031,42 @@ def int_nvvm_cp_async_bulk_shared_cta_to_global
[IntrConvergent, IntrArgMemOnly,
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>]>;
+// From Shared CTA to Global memory with reduction
+class CPAsyncBulkReduceS2GIntrinsic<list<list<int>> red_ops>
+ : DefaultAttrsIntrinsicFlags<[],
+ [llvm_global_ptr_ty, // dst_gmem_ptr
+ llvm_shared_ptr_ty, // src_smem_ptr
+ llvm_i32_ty, // copy_size
+ llvm_i32_ty, // reduction Op enum - nvvm::ReductionOp in NVVMIntrinsicUtils.h
+ llvm_i32_ty, // scope Op enum - nvvm::MMScope in NVVMIntrinsicUtils.h
+ llvm_i64_ty], // cache_hint
+ [llvm_i1_ty], // Flag for cache_hint
+ [IntrConvergent, IntrArgMemOnly,
+ WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>,
+ ImmArg<ArgIndex<3>>, RangeSet<ArgIndex<3>, red_ops>,
+ ImmArg<ArgIndex<4>>, Range<ArgIndex<4>, 0, 4>]>;
+
+// RangeSet bounds are closed.
+foreach type = ["b32", "b64"] in
+ def int_nvvm_cp_async_bulk_reduce_s2g_ # type
+ : CPAsyncBulkReduceS2GIntrinsic<[[0, 2]]>;
+
+foreach type = ["u32"] in
+ def int_nvvm_cp_async_bulk_reduce_s2g_ # type
+ : CPAsyncBulkReduceS2GIntrinsic<[[3, 7]]>;
+
+foreach type = ["s32", "u64", "f16", "bf16"] in
+ def int_nvvm_cp_async_bulk_reduce_s2g_ # type
+ : CPAsyncBulkReduceS2GIntrinsic<[[3, 3], [6, 7]]>;
+
+foreach type = ["s64"] in
+ def int_nvvm_cp_async_bulk_reduce_s2g_ # type
+ : CPAsyncBulkReduceS2GIntrinsic<[[6, 7]]>;
+
+foreach type = ["f32", "f64"] in
+ def int_nvvm_cp_async_bulk_reduce_s2g_ # type
+ : CPAsyncBulkReduceS2GIntrinsic<[[3, 3]]>;
+
// From Shared CTA to Global memory with bytemask
def int_nvvm_cp_async_bulk_shared_cta_to_global_bytemask
: DefaultAttrsIntrinsic<[],
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index 083be6d4247c1..b671c9a4b9563 100644
--- a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
+++ b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
@@ -27,18 +27,24 @@
namespace llvm {
namespace nvvm {
-// Reduction Ops supported with TMA Copy from Shared
-// to Global Memory for the "cp.reduce.async.bulk.tensor.*"
-// family of PTX instructions.
-enum class TMAReductionOp : uint8_t {
- ADD = 0,
- MIN = 1,
- MAX = 2,
- INC = 3,
- DEC = 4,
- AND = 5,
- OR = 6,
- XOR = 7,
+// Reduction op
+enum class ReductionOp : uint8_t {
+ AND = 0,
+ OR = 1,
+ XOR = 2,
+ ADD = 3,
+ INC = 4,
+ DEC = 5,
+ MIN = 6,
+ MAX = 7,
+};
+
+// Memory scope
+enum class MMScope : uint8_t {
+ GPU = 0,
+ CTA = 1,
+ SYSTEM = 2,
+ CLUSTER = 3,
};
// Enum to represent the cta_group::1 and
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index 108f869effe7c..559aed9104d02 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -37,6 +37,16 @@ static bool hasParamSubqualifiers(const MCSubtargetInfo &STI) {
return STI.hasFeature(NVPTX::PTX83);
}
+static bool hasPTXFeatureAtLeast(const MCSubtargetInfo &STI,
+ unsigned FirstPTXFeature) {
+ // NVPTX.td emits PTX features in version order before SM features.
+ const FeatureBitset &Features = STI.getFeatureBits();
+ for (unsigned Feature = FirstPTXFeature; Feature < NVPTX::SM20; ++Feature)
+ if (Features[Feature])
+ return true;
+ return false;
+}
+
NVPTXInstPrinter::NVPTXInstPrinter(const MCAsmInfo &MAI, const MCInstrInfo &MII,
const MCRegisterInfo &MRI)
: MCInstPrinter(MAI, MII, MRI) {}
@@ -485,40 +495,88 @@ void NVPTXInstPrinter::printPrmtMode(const MCInst *MI, int OpNum,
}
}
-void NVPTXInstPrinter::printTmaReductionMode(const MCInst *MI, int OpNum,
- const MCSubtargetInfo &,
- raw_ostream &O) {
- const MCOperand &MO = MI->getOperand(OpNum);
- using RedTy = nvvm::TMAReductionOp;
-
- switch (static_cast<RedTy>(MO.getImm())) {
- case RedTy::ADD:
- O << ".add";
+static void printReductionOp(int64_t Imm, bool NoftzAdd, raw_ostream &O) {
+ switch (static_cast<nvvm::ReductionOp>(Imm)) {
+ case nvvm::ReductionOp::AND:
+ O << ".and";
return;
- case RedTy::MIN:
- O << ".min";
+ case nvvm::ReductionOp::OR:
+ O << ".or";
return;
- case RedTy::MAX:
- O << ".max";
+ case nvvm::ReductionOp::XOR:
+ O << ".xor";
+ return;
+ case nvvm::ReductionOp::ADD:
+ O << (NoftzAdd ? ".add.noftz" : ".add");
return;
- case RedTy::INC:
+ case nvvm::ReductionOp::INC:
O << ".inc";
return;
- case RedTy::DEC:
+ case nvvm::ReductionOp::DEC:
O << ".dec";
return;
- case RedTy::AND:
- O << ".and";
+ case nvvm::ReductionOp::MIN:
+ O << ".min";
return;
- case RedTy::OR:
- O << ".or";
+ case nvvm::ReductionOp::MAX:
+ O << ".max";
return;
- case RedTy::XOR:
- O << ".xor";
+ }
+ llvm_unreachable("Invalid reduction op");
+}
+
+static void printMMScope(nvvm::MMScope Scope, raw_ostream &O) {
+ switch (Scope) {
+ case nvvm::MMScope::GPU:
+ O << ".gpu";
+ return;
+ case nvvm::MMScope::CTA:
+ O << ".cta";
+ return;
+ case nvvm::MMScope::SYSTEM:
+ O << ".sys";
+ return;
+ case nvvm::MMScope::CLUSTER:
+ O << ".cluster";
+ return;
+ }
+ report_fatal_error("Invalid memory scope");
+}
+
+void NVPTXInstPrinter::printCpReduceAsyncBulk(const MCInst *MI, int OpNum,
+ const MCSubtargetInfo &STI,
+ raw_ostream &O,
+ StringRef Modifier) {
+ int64_t Imm = MI->getOperand(OpNum).getImm();
+
+ if (Modifier == "redop") {
+ printReductionOp(Imm, /*NoftzAdd=*/false, O);
return;
}
- llvm_unreachable(
- "Invalid Reduction Op in printCpAsyncBulkTensorReductionMode");
+
+ if (Modifier == "redop_noftz") {
+ printReductionOp(Imm, /*NoftzAdd=*/true, O);
+ return;
+ }
+
+ if (Modifier == "scope") {
+ nvvm::MMScope Scope = static_cast<nvvm::MMScope>(Imm);
+ bool HasSemScope = hasPTXFeatureAtLeast(STI, NVPTX::PTX93);
+ if (!HasSemScope) {
+ if (Scope == nvvm::MMScope::SYSTEM)
+ return;
+ report_fatal_error(
+ "cp.reduce.async.bulk sem.scope requires PTX ISA 9.3 or higher");
+ }
+
+ O << ".relaxed";
+ printMMScope(Scope, O);
+ return;
+ }
+
+ report_fatal_error(formatv(
+ "NVPTX cp.reduce.async.bulk modifier printer does not support \"{}\".",
+ Modifier));
}
void NVPTXInstPrinter::printCTAGroup(const MCInst *MI, int OpNum,
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
index acd7822f1fc5a..a9f71bf317dda 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
@@ -56,8 +56,9 @@ class NVPTXInstPrinter : public MCInstPrinter {
raw_ostream &O);
void printPrmtMode(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
raw_ostream &O);
- void printTmaReductionMode(const MCInst *MI, int OpNum,
- const MCSubtargetInfo &STI, raw_ostream &O);
+ void printCpReduceAsyncBulk(const MCInst *MI, int OpNum,
+ const MCSubtargetInfo &STI, raw_ostream &O,
+ StringRef Modifier = {});
void printCTAGroup(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
raw_ostream &O);
void printCallOperand(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index ede1deb5400b0..a290810d9e293 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -1986,6 +1986,68 @@ void NVPTXDAGToDAGISel::SelectCpAsyncBulkTensorReduceCommon(SDNode *N,
ReplaceNode(N, CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops));
}
+#define GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(TYPE, IS_CACHE_HINT) \
+ ((IS_CACHE_HINT) ? NVPTX::CP_ASYNC_BULK_REDUCE_S2G_##TYPE##_CH \
+ : NVPTX::CP_ASYNC_BULK_REDUCE_S2G_##TYPE)
+
+static unsigned getCpAsyncBulkReduceS2GOpcode(unsigned IID, bool IsCacheHint) {
+ switch (IID) {
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_f16:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(F16, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_bf16:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(BF16, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_b32:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(B32, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_u32:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(U32, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_s32:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(S32, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_b64:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(B64, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_u64:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(U64, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_s64:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(S64, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_f32:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(F32, IsCacheHint);
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_f64:
+ return GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE(F64, IsCacheHint);
+ }
+ llvm_unreachable("unhandled cp.async.bulk.reduce.s2g lowering");
+}
+
+#undef GET_CP_ASYNC_BULK_REDUCE_S2G_OPCODE
+
+void NVPTXDAGToDAGISel::SelectCpAsyncBulkReduceS2G(SDNode *N) {
+ // Operands are {Chain, Intrinsic-ID, dst, src, size, red_op, scope,
+ // cache_hint, cache_hint_flag}. The cache hint operand is only emitted when
+ // its flag is set.
+ bool IsCacheHint = N->getConstantOperandVal(8) == 1;
+ unsigned RedOp = N->getConstantOperandVal(5);
+ unsigned Scope = N->getConstantOperandVal(6);
+
+ if (Scope != static_cast<unsigned>(nvvm::MMScope::SYSTEM) &&
+ Subtarget->getPTXVersion() < 93)
+ report_fatal_error(
+ "cp.reduce.async.bulk sem.scope requires PTX ISA 9.3 or higher");
+
+ SDLoc DL(N);
+ const auto [DstBase, DstOffset] = selectADDR(N->getOperand(2), CurDAG);
+ const auto [SrcBase, SrcOffset] = selectADDR(N->getOperand(3), CurDAG);
+
+ SmallVector<SDValue, 10> Ops{DstBase, DstOffset, SrcBase, SrcOffset,
+ N->getOperand(4)};
+ if (IsCacheHint)
+ Ops.push_back(N->getOperand(7));
+ Ops.push_back(getI32Imm(Scope, DL));
+ Ops.push_back(getI32Imm(RedOp, DL));
+ Ops.push_back(N->getOperand(0)); // Chain operand
+
+ unsigned IID = N->getConstantOperandVal(1);
+ unsigned Opcode = getCpAsyncBulkReduceS2GOpcode(IID, IsCacheHint);
+ ReplaceNode(N, CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops));
+}
+
#define TCGEN05_ST_OPCODE(SHAPE, NUM) \
(enableUnpack ? NVPTX::TCGEN05_ST_##SHAPE##_##NUM##_UNPACK \
: NVPTX::TCGEN05_ST_##SHAPE##_##NUM)
@@ -2101,22 +2163,34 @@ void NVPTXDAGToDAGISel::SelectTcgen05St(SDNode *N, bool hasOffset) {
bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
unsigned IID = N->getConstantOperandVal(1);
- using TMARedTy = llvm::nvvm::TMAReductionOp;
- auto CastTy = [](TMARedTy Op) { return static_cast<unsigned>(Op); };
+ using RedTy = llvm::nvvm::ReductionOp;
+ auto CastTy = [](RedTy Op) { return static_cast<unsigned>(Op); };
switch (IID) {
default:
return false;
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_f16:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_bf16:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_b32:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_u32:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_s32:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_b64:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_u64:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_s64:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_f32:
+ case Intrinsic::nvvm_cp_async_bulk_reduce_s2g_f64:
+ SelectCpAsyncBulkReduceS2G(N);
+ return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_1d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_2d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::ADD));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::ADD));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::ADD),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::ADD),
/*IsIm2Col=*/true);
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_1d:
@@ -2124,12 +2198,12 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::MIN));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::MIN));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::MIN),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::MIN),
/*IsIm2Col=*/true);
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_1d:
@@ -2137,12 +2211,12 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::MAX));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::MAX));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::MAX),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::MAX),
/*IsIm2Col=*/true);
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_1d:
@@ -2150,12 +2224,12 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::INC));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::INC));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::INC),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::INC),
/*IsIm2Col=*/true);
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_1d:
@@ -2163,12 +2237,12 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::DEC));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::DEC));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::DEC),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::DEC),
/*IsIm2Col=*/true);
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_1d:
@@ -2176,12 +2250,12 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::AND));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::AND));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::AND),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::AND),
/*IsIm2Col=*/true);
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_1d:
@@ -2189,12 +2263,12 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::OR));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::OR));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::OR),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::OR),
/*IsIm2Col=*/true);
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_1d:
@@ -2202,12 +2276,12 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::XOR));
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::XOR));
return true;
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_im2col_3d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_im2col_4d:
case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_im2col_5d:
- SelectCpAsyncBulkTensorReduceCommon(N, CastTy(TMARedTy::XOR),
+ SelectCpAsyncBulkTensorReduceCommon(N, CastTy(RedTy::XOR),
/*IsIm2Col=*/true);
return true;
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
index fcb5700dcb6d4..48efd9a8078ea 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
@@ -90,6 +90,7 @@ class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel {
void SelectI128toV2I64(SDNode *N);
void SelectCpAsyncBulkTensorReduceCommon(SDNode *N, unsigned RedOp,
bool IsIm2Col = false);
+ void SelectCpAsyncBulkReduceS2G(SDNode *N);
void SelectTcgen05Ld(SDNode *N, bool hasOffset = false);
void SelectTcgen05St(SDNode *N, bool hasOffset = false);
void selectAtomicSwap128(SDNode *N);
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 5af5ae2fe9e52..65047be73dbd7 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -562,6 +562,38 @@ multiclass CP_ASYNC_BULK_S2G_INTR<bit has_ch> {
defm CP_ASYNC_BULK_S2G : CP_ASYNC_BULK_S2G_INTR<has_ch = 0>;
defm CP_ASYNC_BULK_S2G_CH : CP_ASYNC_BULK_S2G_INTR<has_ch = 1>;
+def CpReduceAsyncBulkFlags : Operand<i32> {
+ let PrintMethod = "printCpReduceAsyncBulk";
+}
+
+// TMA Async Bulk Reduction from Shared CTA to Global memory
+multiclass CP_ASYNC_BULK_REDUCE_S2G_INTR<string type, bit noftz_add> {
+ defvar RedOp = !if(noftz_add, "redop_noftz", "redop");
+ defvar Base =
+ "cp.reduce.async.bulk${scope:scope}.global.shared::cta.bulk_group";
+
+ def "" : NVPTXInst<(outs),
+ (ins ADDR:$dst, ADDR:$src, B32:$size,
+ CpReduceAsyncBulkFlags:$scope, CpReduceAsyncBulkFlags:$red_op),
+ Base # "${red_op:" # RedOp # "}." # type # " [$dst], [$src], $size;">,
+ Requires<[hasPTX<80>, hasSM<90>]>;
+
+ def _CH : NVPTXInst<(outs),
+ (ins ADDR:$dst, ADDR:$src, B32:$size, B64:$ch,
+ CpReduceAsyncBulkFlags:$scope, CpReduceAsyncBulkFlags:$red_op),
+ Base # ".L2::cache_hint${red_op:" # RedOp # "}." # type #
+ " [$dst], [$src], $size, $ch;">,
+ Requires<[hasPTX<80>, hasSM<90>]>;
+}
+
+foreach type = ["f16", "bf16"] in
+ defm CP_ASYNC_BULK_REDUCE_S2G_ # !toupper(type) :
+ CP_ASYNC_BULK_REDUCE_S2G_INTR<type, noftz_add = 1>;
+
+foreach type = ["b32", "u32", "s32", "b64", "u64", "s64", "f32", "f64"] in
+ defm CP_ASYNC_BULK_REDUCE_S2G_ # !toupper(type) :
+ CP_ASYNC_BULK_REDUCE_S2G_INTR<type, noftz_add = 0>;
+
multiclass CP_ASYNC_BULK_G2S_INTR<bit has_ch> {
defvar Intr = int_nvvm_cp_async_bulk_global_to_shared_cluster;
@@ -852,10 +884,6 @@ foreach dim = 1...5 in {
defm TMA_S2G_TILE_SCATTER4_2D : TMA_TENSOR_S2G_INTR<5, "tile_scatter4",
[callSubtarget<"hasTMABlackwellSupport">]>;
-def TMAReductionFlags : Operand<i32> {
- let PrintMethod = "printTmaReductionMode";
-}
-
// TMA Copy from Shared to Global memory with Reduction
multiclass CP_ASYNC_BULK_TENSOR_REDUCE_INTR<int dim, bit shared32, string mode> {
defvar dims_dag = TMA_DIMS_UTIL<dim>.ins_dag;
@@ -870,12 +898,12 @@ multiclass CP_ASYNC_BULK_TENSOR_REDUCE_INTR<int dim, bit shared32, string mode>
defvar suffix = "." # mode_asm_str # ".bulk_group";
def "" : NVPTXInst<(outs),
- !con((ins rc:$src, B64:$tmap), dims_dag, (ins TMAReductionFlags:$red_op)),
- !strconcat(prefix, "${red_op}", suffix, asm_str, ";")>,
+ !con((ins rc:$src, B64:$tmap), dims_dag, (ins CpReduceAsyncBulkFlags:$red_op)),
+ !strconcat(prefix, "${red_op:redop}", suffix, asm_str, ";")>,
Requires<[hasPTX<80>, hasSM<90>]>;
def _CH : NVPTXInst<(outs),
- !con((ins rc:$src, B64:$tmap), dims_dag, (ins B64:$ch, TMAReductionFlags:$red_op)),
- !strconcat(prefix, "${red_op}", suffix, ".L2::cache_hint", asm_str, ", $ch;")>,
+ !con((ins rc:$src, B64:$tmap), dims_dag, (ins B64:$ch, CpReduceAsyncBulkFlags:$red_op)),
+ !strconcat(prefix, "${red_op:redop}", suffix, ".L2::cache_hint", asm_str, ", $ch;")>,
Requires<[hasPTX<80>, hasSM<90>]>;
}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92-invalid.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92-invalid.ll
new file mode 100644
index 0000000000000..7ffc5834e7097
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92-invalid.ll
@@ -0,0 +1,11 @@
+; RUN: not --crash llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 2>&1 | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+
+define void @cp_async_bulk_reduce_s2g_ptx92_invalid_scope(ptr addrspace(3) %src, ptr addrspace(1) %dst, i32 %size, i64 %ch) {
+; CHECK: LLVM ERROR: cp.reduce.async.bulk sem.scope requires PTX ISA 9.3 or higher
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 1, i64 %ch, i1 0)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92.ll
new file mode 100644
index 0000000000000..7518af4439bff
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g-ptx92.ll
@@ -0,0 +1,27 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 | FileCheck %s
+; RUN: %if ptxas-sm_90 && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx92 | %ptxas-verify -arch=sm_90 %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.u64(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+
+define void @cp_async_bulk_reduce_s2g_ptx92(ptr addrspace(3) %src, ptr addrspace(1) %dst, i32 %size, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_reduce_s2g_ptx92(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_reduce_s2g_ptx92_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [cp_async_bulk_reduce_s2g_ptx92_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [cp_async_bulk_reduce_s2g_ptx92_param_2];
+; CHECK-NEXT: cp.reduce.async.bulk.global.shared::cta.bulk_group.add.u32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: ld.param.b64 %rd3, [cp_async_bulk_reduce_s2g_ptx92_param_3];
+; CHECK-NEXT: cp.reduce.async.bulk.global.shared::cta.bulk_group.L2::cache_hint.max.u64 [%rd2], [%rd1], %r1, %rd3;
+; CHECK-NEXT: ret;
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 2, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u64(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 7, i32 2, i64 %ch, i1 1)
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g.ll
new file mode 100644
index 0000000000000..8b9fdf2343472
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-reduce-s2g.ll
@@ -0,0 +1,65 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx93 | FileCheck %s
+; RUN: %if ptxas-sm_90 && ptxas-isa-9.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx93 | %ptxas-verify -arch=sm_90 %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.f16(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.bf16(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.b32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.s32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.b64(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.u64(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.s64(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.f32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.f64(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+
+define void @cp_async_bulk_reduce_s2g(ptr addrspace(3) %src, ptr addrspace(1) %dst, i32 %size, i64 %ch) {
+; CHECK-LABEL: cp_async_bulk_reduce_s2g(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_reduce_s2g_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [cp_async_bulk_reduce_s2g_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [cp_async_bulk_reduce_s2g_param_2];
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.sys.global.shared::cta.bulk_group.add.noftz.f16 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cta.global.shared::cta.bulk_group.min.bf16 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cluster.global.shared::cta.bulk_group.and.b32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: ld.param.b64 %rd3, [cp_async_bulk_reduce_s2g_param_3];
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.gpu.global.shared::cta.bulk_group.xor.b32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.sys.global.shared::cta.bulk_group.L2::cache_hint.or.b64 [%rd2], [%rd1], %r1, %rd3;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cta.global.shared::cta.bulk_group.add.u32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cluster.global.shared::cta.bulk_group.inc.u32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.gpu.global.shared::cta.bulk_group.dec.u32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.sys.global.shared::cta.bulk_group.min.s32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cta.global.shared::cta.bulk_group.max.u64 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cluster.global.shared::cta.bulk_group.min.s64 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.gpu.global.shared::cta.bulk_group.add.f32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.sys.global.shared::cta.bulk_group.L2::cache_hint.add.f64 [%rd2], [%rd1], %r1, %rd3;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cta.global.shared::cta.bulk_group.add.u32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.cluster.global.shared::cta.bulk_group.add.u32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.gpu.global.shared::cta.bulk_group.add.u32 [%rd2], [%rd1], %r1;
+; CHECK-NEXT: cp.reduce.async.bulk.relaxed.sys.global.shared::cta.bulk_group.L2::cache_hint.add.u32 [%rd2], [%rd1], %r1, %rd3;
+; CHECK-NEXT: ret;
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.f16(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 2, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.bf16(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 6, i32 1, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.b32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 0, i32 3, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.b32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 2, i32 0, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.b64(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 1, i32 2, i64 %ch, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 1, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 4, i32 3, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 5, i32 0, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.s32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 6, i32 2, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u64(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 7, i32 1, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.s64(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 6, i32 3, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.f32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 0, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.f64(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 2, i64 %ch, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 1, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 3, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 0, i64 %ch, i1 0)
+ tail call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 2, i64 %ch, i1 1)
+ ret void
+}
diff --git a/llvm/test/Verifier/NVPTX/cp-async-bulk-reduce-s2g-invalid.ll b/llvm/test/Verifier/NVPTX/cp-async-bulk-reduce-s2g-invalid.ll
new file mode 100644
index 0000000000000..99d1ea14a61ae
--- /dev/null
+++ b/llvm/test/Verifier/NVPTX/cp-async-bulk-reduce-s2g-invalid.ll
@@ -0,0 +1,31 @@
+; RUN: not llvm-as -disable-output < %s 2>&1 | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.b32(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+declare void @llvm.nvvm.cp.async.bulk.reduce.s2g.f16(ptr addrspace(1), ptr addrspace(3), i32, i32 immarg, i32 immarg, i64, i1 immarg)
+
+define void @invalid_u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch) {
+ ; CHECK: immarg value 0 out of rangeset
+ call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 0, i32 2, i64 %ch, i1 0)
+ ret void
+}
+
+define void @invalid_b32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch) {
+ ; CHECK: immarg value 3 out of rangeset
+ call void @llvm.nvvm.cp.async.bulk.reduce.s2g.b32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 2, i64 %ch, i1 0)
+ ret void
+}
+
+define void @invalid_f16(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch) {
+ ; CHECK: immarg value 4 out of rangeset
+ call void @llvm.nvvm.cp.async.bulk.reduce.s2g.f16(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 4, i32 2, i64 %ch, i1 0)
+ ret void
+}
+
+define void @invalid_scope(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch) {
+ ; CHECK: immarg value 4 for arg 4 out of range [0,4)
+ call void @llvm.nvvm.cp.async.bulk.reduce.s2g.u32(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i32 3, i32 4, i64 %ch, i1 0)
+ ret void
+}
More information about the llvm-commits
mailing list