46#include "llvm/IR/IntrinsicsNVPTX.h"
69#define DEBUG_TYPE "nvptx-lower"
81 cl::desc(
"NVPTX Specific: FMA contraction (0: don't do it"
82 " 1: do it 2: do it aggressively"),
87 cl::desc(
"NVPTX Specifies: 0 use div.approx, 1 use div.full, 2 use"
88 " IEEE Compliant F32 div.rnd if available."),
93 cl::desc(
"NVPTX Specific: 0 use sqrt.approx, 1 use sqrt.rn."),
97 "nvptx-force-min-byval-param-align",
cl::Hidden,
98 cl::desc(
"NVPTX Specific: force 4-byte minimal alignment for byval"
99 " params of device functions."),
181 Offsets->push_back(StartingOffset + 0);
182 Offsets->push_back(StartingOffset + 8);
189 if (
StructType *STy = dyn_cast<StructType>(Ty)) {
190 auto const *SL =
DL.getStructLayout(STy);
192 for(
auto *EI : STy->elements()) {
194 StartingOffset + SL->getElementOffset(ElementNum));
201 for (
unsigned i = 0, e = TempVTs.
size(); i != e; ++i) {
228 (NumElts % 4 == 0 || NumElts == 3)) {
231 NumElts = (NumElts + 3) / 4;
233 for (
unsigned j = 0; j != NumElts; ++j) {
241 Offsets->push_back(Off);
256 "Promotion is not suitable for scalars of size larger than 64-bits");
258 *PromotedVT = MVT::i1;
263 *PromotedVT = MVT::i8;
266 *PromotedVT = MVT::i16;
269 *PromotedVT = MVT::i32;
272 *PromotedVT = MVT::i64;
275 return EVT(*PromotedVT) != VT;
295 if (ParamAlignment < AccessSize)
298 if (Offsets[
Idx] & (AccessSize - 1))
301 EVT EltVT = ValueVTs[
Idx];
305 if (EltSize >= AccessSize)
308 unsigned NumElts = AccessSize / EltSize;
310 if (AccessSize != EltSize * NumElts)
314 if (
Idx + NumElts > ValueVTs.
size())
318 if (NumElts != 4 && NumElts != 2)
321 for (
unsigned j =
Idx + 1; j <
Idx + NumElts; ++j) {
323 if (ValueVTs[j] != EltVT)
327 if (Offsets[j] - Offsets[j - 1] != EltSize)
355 Align ParamAlignment,
bool IsVAArg =
false) {
365 for (
int I = 0, E = ValueVTs.
size();
I != E; ++
I) {
368 for (
unsigned AccessSize : {16, 8, 4, 2}) {
370 I, AccessSize, ValueVTs, Offsets, ParamAlignment);
379 assert(
I + 1 < E &&
"Not enough elements.");
385 assert(
I + 3 < E &&
"Not enough elements.");
456 Op, VT, IsOpSupported ? Action : NoBF16Action);
461 bool IsOpSupported =
false;
543 for (
MVT VT : {MVT::bf16, MVT::f16, MVT::v2bf16, MVT::v2f16, MVT::f32,
544 MVT::f64, MVT::i1, MVT::i8, MVT::i16, MVT::v2i16, MVT::v4i8,
545 MVT::i32, MVT::i64}) {
670 for (
const auto& Ty : {MVT::i16, MVT::i32, MVT::i64}) {
752 const bool IsFP16FP16x2NegAvailable = STI.
getSmVersion() >= 53 &&
755 for (
const auto &VT : {MVT::f16, MVT::v2f16})
780 for (
MVT VT : {MVT::bf16, MVT::f32, MVT::f64}) {
789 for (
MVT VT : {MVT::i1, MVT::i16, MVT::i32, MVT::i64}) {
818 for (
const auto &
Op :
842 return IsAtLeastSm80 ?
Legal : NotSm80Action;
848 setFP16OperationAction(
Op, MVT::v2f16, GetMinMaxAction(
Expand),
Expand);
855 setFP16OperationAction(
Op, MVT::f16, GetMinMaxAction(
Expand),
Expand);
858 setFP16OperationAction(
Op, MVT::v2f16, GetMinMaxAction(
Expand),
Expand);
875#define MAKE_CASE(V) \
1316 bool Reciprocal)
const {
1337 if (Reciprocal || ExtraSteps > 0) {
1339 return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_rsqrt_approx_ftz_f
1340 : Intrinsic::nvvm_rsqrt_approx_f);
1341 else if (VT == MVT::f64)
1342 return MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d);
1347 return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_sqrt_approx_ftz_f
1348 : Intrinsic::nvvm_sqrt_approx_f);
1356 DAG.
getConstant(Intrinsic::nvvm_rcp_approx_ftz_d,
DL, MVT::i32),
1357 MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d));
1379 std::optional<std::pair<unsigned, const APInt &>> VAInfo,
1380 const CallBase &CB,
unsigned UniqueCallSite)
const {
1384 assert(isABI &&
"Non-ABI compilation is not supported");
1388 std::string Prototype;
1390 O <<
"prototype_" << UniqueCallSite <<
" : .callprototype ";
1399 if (
auto *ITy = dyn_cast<IntegerType>(retTy)) {
1400 size = ITy->getBitWidth();
1403 "Floating point type expected here");
1411 O <<
".param .b" <<
size <<
" _";
1412 }
else if (isa<PointerType>(retTy)) {
1413 O <<
".param .b" << PtrVT.getSizeInBits() <<
" _";
1415 O <<
".param .align " << (retAlignment ? retAlignment->value() : 0)
1416 <<
" .b8 _[" <<
DL.getTypeAllocSize(retTy) <<
"]";
1427 unsigned NumArgs = VAInfo ? VAInfo->first : Args.size();
1428 for (
unsigned i = 0, OIdx = 0; i != NumArgs; ++i, ++OIdx) {
1429 Type *Ty = Args[i].Ty;
1435 if (!Outs[OIdx].Flags.isByVal()) {
1437 unsigned ParamAlign = 0;
1438 const CallInst *CallI = cast<CallInst>(&CB);
1440 if (!
getAlign(*CallI, i + 1, ParamAlign))
1442 O <<
".param .align " << ParamAlign <<
" .b8 ";
1444 O <<
"[" <<
DL.getTypeAllocSize(Ty) <<
"]";
1448 if (
unsigned len = vtparts.
size())
1454 (
getValueType(
DL, Ty) == MVT::i8 && Outs[OIdx].VT == MVT::i16)) &&
1455 "type mismatch between callee prototype and arguments");
1458 if (isa<IntegerType>(Ty)) {
1459 sz = cast<IntegerType>(Ty)->getBitWidth();
1461 }
else if (isa<PointerType>(Ty)) {
1462 sz = PtrVT.getSizeInBits();
1466 O <<
".param .b" << sz <<
" ";
1471 Type *ETy = Args[i].IndirectType;
1472 Align InitialAlign = Outs[OIdx].Flags.getNonZeroByValAlign();
1473 Align ParamByValAlign =
1476 O <<
".param .align " << ParamByValAlign.
value() <<
" .b8 ";
1478 O <<
"[" << Outs[OIdx].Flags.getByValSize() <<
"]";
1482 O << (first ?
"" :
",") <<
" .param .align " << VAInfo->second
1497 return DL.getABITypeAlign(Ty);
1500 unsigned Alignment = 0;
1503 if (!DirectCallee) {
1508 if (
const auto *CI = dyn_cast<CallInst>(CB)) {
1511 return Align(Alignment);
1520 return Align(Alignment);
1527 return DL.getABITypeAlign(Ty);
1531 switch (ElementType.getSimpleVT().SimpleTy) {
1536 ElementType = MVT::i16;
1541 ElementType = MVT::i32;
1544 ElementType = MVT::i64;
1556 unsigned ArgID,
const SDLoc &dl) {
1563 for (
unsigned i = 0, n = ElementType.getSizeInBits() / 8; i < n; i++) {
1589 EVT MergedType = ElementType;
1596 for (
unsigned i = 0, n = ElementType.getSizeInBits() / 8; i < n; i++) {
1623 if (ElementType != MergedType)
1634 "Support for variadic functions (unsized array parameter) introduced "
1635 "in PTX ISA version 6.0 and requires target sm_30.");
1651 assert(isABI &&
"Non-ABI compilation is not supported");
1673 unsigned VAOffset = 0;
1680 unsigned ParamCount = 0;
1693 for (
unsigned i = 0, e = Args.size(); i != e; ++i, ++OIdx) {
1694 EVT VT = Outs[OIdx].VT;
1695 Type *Ty = Args[i].Ty;
1697 bool IsByVal = Outs[OIdx].Flags.isByVal();
1702 assert((!IsByVal || Args[i].IndirectType) &&
1703 "byval arg must have indirect type");
1704 Type *ETy = (IsByVal ? Args[i].IndirectType : Ty);
1712 Align InitialAlign = Outs[OIdx].Flags.getNonZeroByValAlign();
1716 VAOffset =
alignTo(VAOffset, ArgAlign);
1718 ArgAlign = getArgumentAlignment(CB, Ty, ParamCount + 1,
DL);
1722 (IsByVal ? Outs[OIdx].Flags.getByValSize() :
DL.getTypeAllocSize(Ty));
1728 if (ParamCount == FirstVAArg) {
1734 DeclareParamVTs, DeclareParamOps);
1736 NeedAlign = PassAsArray;
1737 }
else if (PassAsArray) {
1754 SDValue DeclareScalarParamOps[] = {
1759 DeclareScalarParamOps);
1768 bool ExtendIntegerParam =
1773 for (
unsigned j = 0, je = VTs.
size(); j != je; ++j) {
1775 int CurOffset = Offsets[j];
1780 SDValue StVal = OutVals[OIdx];
1784 EltVT =
EVT(PromotedVT);
1789 StVal = DAG.
getNode(Ext, dl, PromotedVT, StVal);
1798 }
else if (ExtendIntegerParam) {
1799 assert(VTs.
size() == 1 &&
"Scalar can't have multiple parts.");
1803 dl, MVT::i32, StVal);
1814 if (VectorInfo[j] ==
PVF_SCALAR && !IsVAArg && PartAlign.has_value() &&
1817 assert(StoreOperands.
empty() &&
"Unfinished preceeding store.");
1819 DAG, Chain, IsByVal ? CurOffset + VAOffset : CurOffset, EltVT,
1820 StVal, InGlue, ParamCount, dl);
1831 assert(StoreOperands.
empty() &&
"Unfinished preceding store.");
1834 DAG.
getConstant(IsVAArg ? FirstVAArg : ParamCount, dl, MVT::i32));
1837 IsByVal ? CurOffset + VAOffset : (IsVAArg ? VAOffset : CurOffset),
1845 unsigned NumElts = StoreOperands.
size() - 3;
1865 EVT TheStoreType = ExtendIntegerParam ? MVT::i32 : EltVT;
1868 Op, dl, DAG.
getVTList(MVT::Other, MVT::Glue), StoreOperands,
1874 StoreOperands.
clear();
1878 if (!IsByVal && IsVAArg) {
1880 "Vectorization is expected to be disabled for variadics.");
1881 VAOffset +=
DL.getTypeAllocSize(
1888 assert(StoreOperands.
empty() &&
"Unfinished parameter store.");
1889 if (!IsByVal && VTs.
size() > 0)
1892 if (IsByVal && IsVAArg)
1900 if (Ins.size() > 0) {
1907 unsigned resultsz =
DL.getTypeAllocSizeInBits(
RetTy);
1918 retAlignment = getArgumentAlignment(CB,
RetTy, 0,
DL);
1919 assert(retAlignment &&
"retAlignment is guaranteed to be set");
1922 Chain, DAG.
getConstant(retAlignment->value(), dl, MVT::i32),
1940 VADeclareParam->
getVTList(), DeclareParamOps);
1948 if (isa<ExternalSymbolSDNode>(Callee)) {
1953 assert(CalleeFunc !=
nullptr &&
"Libcall callee must be set.");
1957 CalleeFunc->
addFnAttr(
"nvptx-libcall-callee",
"true");
1970 DL,
RetTy, Args, Outs, retAlignment,
1972 ? std::optional<std::pair<unsigned, const APInt &>>(std::make_pair(
1975 *CB, UniqueCallSite);
1988 Chain, DAG.
getConstant((Ins.size() == 0) ? 0 : 1, dl, MVT::i32), InGlue
1995 Chain = DAG.
getNode(Opcode, dl, PrintCallVTs, PrintCallOps);
2000 SDValue CallVoidOps[] = { Chain, Callee, InGlue };
2006 SDValue CallArgBeginOps[] = { Chain, InGlue };
2011 for (
unsigned i = 0, e = std::min(CLI.
NumFixedArgs + 1, ParamCount); i != e;
2021 Chain = DAG.
getNode(opcode, dl, CallArgVTs, CallArgOps);
2025 SDValue CallArgEndOps[] = { Chain,
2034 Chain, DAG.
getConstant(UniqueCallSite, dl, MVT::i32), InGlue};
2051 if (Ins.size() > 0) {
2055 assert(VTs.
size() == Ins.size() &&
"Bad value decomposition");
2057 Align RetAlign = getArgumentAlignment(CB,
RetTy, 0,
DL);
2066 bool ExtendIntegerRetVal =
2067 RetTy->isIntegerTy() &&
DL.getTypeAllocSizeInBits(
RetTy) < 32;
2069 for (
unsigned i = 0, e = VTs.
size(); i != e; ++i) {
2070 bool needTruncate =
false;
2071 EVT TheLoadType = VTs[i];
2072 EVT EltType = Ins[i].VT;
2077 TheLoadType =
EVT(PromotedVT);
2078 EltType =
EVT(PromotedVT);
2079 needTruncate =
true;
2082 if (ExtendIntegerRetVal) {
2083 TheLoadType = MVT::i32;
2085 needTruncate =
true;
2087 if (VTs[i].isInteger())
2088 needTruncate =
true;
2095 EltAlign <
DL.getABITypeAlign(
2097 assert(VecIdx == -1 && LoadVTs.
empty() &&
"Orphaned operand list.");
2099 DAG, Chain, Offsets[i], TheLoadType, InGlue, TempProxyRegOps, dl);
2101 ProxyRegTruncates.
push_back(std::optional<MVT>());
2110 assert(VecIdx == -1 && LoadVTs.
empty() &&
"Orphaned operand list.");
2117 unsigned NumElts = LoadVTs.
size();
2137 DAG.
getConstant(Offsets[VecIdx], dl, MVT::i32), InGlue};
2139 Op, dl, DAG.
getVTList(LoadVTs), LoadOperands, TheLoadType,
2143 for (
unsigned j = 0; j < NumElts; ++j) {
2147 ProxyRegTruncates.
push_back(std::optional<MVT>(Ins[VecIdx + j].VT));
2149 ProxyRegTruncates.
push_back(std::optional<MVT>());
2153 InGlue = RetVal.
getValue(NumElts + 1);
2163 DAG.
getCALLSEQ_END(Chain, UniqueCallSite, UniqueCallSite + 1, InGlue, dl);
2169 for (
unsigned i = 0; i < ProxyRegOps.
size(); ++i) {
2170 if (i < RetElts.
size() && RetElts[i]) {
2177 DAG.
getVTList(ProxyRegOps[i].getSimpleValueType(), MVT::Other, MVT::Glue),
2178 { Chain, ProxyRegOps[i], InGlue }
2181 Chain = Ret.getValue(1);
2182 InGlue = Ret.getValue(2);
2184 if (ProxyRegTruncates[i]) {
2191 for (
SDValue &
T : TempProxyRegOps) {
2194 DAG.
getVTList(
T.getSimpleValueType(), MVT::Other, MVT::Glue),
2195 {Chain, T.getOperand(0), InGlue});
2217 "Support for dynamic alloca introduced in PTX ISA version 7.3 and "
2218 "requires target sm_52.",
2228 uint64_t Align = cast<ConstantSDNode>(
Op.getOperand(2))->getZExtValue();
2242 SDValue MergeOps[] = {Alloca, Chain};
2254 unsigned NumOperands = Node->getNumOperands();
2255 for (
unsigned i = 0; i < NumOperands; ++i) {
2256 SDValue SubOp = Node->getOperand(i);
2260 for (
unsigned j = 0; j < NumSubElem; ++j) {
2274 EVT VT =
Op->getValueType(0);
2275 if (!(
Isv2x16VT(VT) || VT == MVT::v4i8))
2281 return Operand->isUndef() || isa<ConstantSDNode>(Operand) ||
2282 isa<ConstantFPSDNode>(Operand);
2286 if (VT == MVT::v4i8) {
2308 EVT VT =
Op->getValueType(0);
2310 return APInt(32, 0);
2312 if (VT == MVT::v2f16 || VT == MVT::v2bf16)
2313 Value = cast<ConstantFPSDNode>(Operand)->getValueAPF().bitcastToAPInt();
2314 else if (VT == MVT::v2i16 || VT == MVT::v4i8)
2320 if (VT == MVT::v4i8)
2322 return Value.zext(32);
2326 Value = GetOperand(
Op, 0) | GetOperand(
Op, 1).shl(16);
2327 }
else if (VT == MVT::v4i8) {
2328 Value = GetOperand(
Op, 0) | GetOperand(
Op, 1).shl(8) |
2329 GetOperand(
Op, 2).shl(16) | GetOperand(
Op, 3).shl(24);
2344 if (VectorVT == MVT::v4i8) {
2356 if (isa<ConstantSDNode>(
Index.getNode()))
2377 if (VectorVT != MVT::v4i8)
2381 if (
Value->isUndef())
2400 if (VectorVT != MVT::v4i8 ||
Op.getValueType() != MVT::v4i8)
2408 if (
I.value() != -1)
2409 Selector |= (
I.value() << (
I.index() * 4));
2427 EVT VT =
Op.getValueType();
2488 EVT VT =
Op.getValueType();
2539 EVT VT =
Op.getValueType();
2542 return LowerFROUND32(
Op, DAG);
2545 return LowerFROUND64(
Op, DAG);
2561 EVT VT =
Op.getValueType();
2567 const int SignBitMask = 0x80000000;
2570 const int PointFiveInBits = 0x3F000000;
2571 SDValue PointFiveWithSignRaw =
2602 EVT VT =
Op.getValueType();
2634 if (
Op.getValueType() == MVT::bf16) {
2638 DAG.
getNode(
Op.getOpcode(), Loc, MVT::f32,
Op.getOperand(0)),
2650 if (
Op.getOperand(0).getValueType() == MVT::bf16) {
2653 Op.getOpcode(), Loc,
Op.getValueType(),
2663 EVT NarrowVT =
Op.getValueType();
2700 EVT WideVT =
Op.getValueType();
2727 if (
Op.getValueType() != MVT::v2i16)
2729 EVT EltVT =
Op.getValueType().getVectorElementType();
2731 for (
int I = 0, E =
Op.getValueType().getVectorNumElements();
I < E;
I++) {
2734 [&](
const SDUse &O) {
2735 return DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, EltVT,
2736 O.get(), DAG.getIntPtrConstant(I, DL));
2747 switch (
Op.getOpcode()) {
2757 return LowerBUILD_VECTOR(
Op, DAG);
2761 return LowerEXTRACT_VECTOR_ELT(
Op, DAG);
2763 return LowerINSERT_VECTOR_ELT(
Op, DAG);
2765 return LowerVECTOR_SHUFFLE(
Op, DAG);
2767 return LowerCONCAT_VECTORS(
Op, DAG);
2769 return LowerSTORE(
Op, DAG);
2771 return LowerLOAD(
Op, DAG);
2773 return LowerShiftLeftParts(
Op, DAG);
2776 return LowerShiftRightParts(
Op, DAG);
2778 return LowerSelect(
Op, DAG);
2780 return LowerFROUND(
Op, DAG);
2783 return LowerINT_TO_FP(
Op, DAG);
2786 return LowerFP_TO_INT(
Op, DAG);
2788 return LowerFP_ROUND(
Op, DAG);
2790 return LowerFP_EXTEND(
Op, DAG);
2792 return LowerVAARG(
Op, DAG);
2794 return LowerVASTART(
Op, DAG);
2821 const Value *V = cast<SrcValueSDNode>(Node->getOperand(2))->getValue();
2822 EVT VT = Node->getValueType(0);
2824 SDValue Tmp1 = Node->getOperand(0);
2825 SDValue Tmp2 = Node->getOperand(1);
2826 const MaybeAlign MA(Node->getConstantOperandVal(3));
2864 SDValue Arg = getParamSymbol(DAG, -1, PtrVT);
2867 const Value *SV = cast<SrcValueSDNode>(
Op.getOperand(2))->getValue();
2868 return DAG.
getStore(
Op.getOperand(0),
DL, VAReg,
Op.getOperand(1),
2878 assert(
Op.getValueType() == MVT::i1 &&
"Custom lowering enabled only for i1");
2889 if (
Op.getValueType() == MVT::i1)
2890 return LowerLOADi1(
Op, DAG);
2894 EVT VT =
Op.getValueType();
2897 EVT MemVT =
Load->getMemoryVT();
2899 MemVT, *
Load->getMemOperand())) {
2919 "Custom lowering for i1 load only");
2921 LD->getPointerInfo(),
LD->getAlign(),
2922 LD->getMemOperand()->getFlags());
2927 SDValue Ops[] = { result,
LD->getChain() };
2936 return LowerSTOREi1(
Op, DAG);
2940 if ((
Isv2x16VT(VT) || VT == MVT::v4i8) &&
2942 VT, *
Store->getMemOperand()))
2950 return LowerSTOREVector(
Op, DAG);
2998 if (Alignment < PrefAlign) {
3007 unsigned Opcode = 0;
3014 bool NeedExt =
false;
3018 bool StoreF16x2 =
false;
3046 for (
unsigned i = 0; i < NumElts; ++i) {
3057 for (
unsigned i = 0; i < NumElts; ++i) {
3067 Ops.
append(
N->op_begin() + 2,
N->op_end());
3094 DAG.
getTruncStore(Tmp1, dl, Tmp3, Tmp2,
ST->getPointerInfo(), MVT::i8,
3095 ST->getAlign(),
ST->getMemOperand()->getFlags());
3123 std::vector<SDValue> OutChains;
3126 assert(isABI &&
"Non-ABI compilation is not supported");
3130 std::vector<Type *> argTypes;
3131 std::vector<const Argument *> theArgs;
3133 theArgs.push_back(&
I);
3134 argTypes.push_back(
I.getType());
3145 unsigned InsIdx = 0;
3147 for (
unsigned i = 0, e = theArgs.size(); i != e; ++i, ++InsIdx) {
3148 Type *Ty = argTypes[i];
3150 if (theArgs[i]->use_empty()) {
3156 if (vtparts.
empty())
3159 for (
unsigned parti = 0, parte = vtparts.
size(); parti != parte;
3164 if (vtparts.
size() > 0)
3171 for (
unsigned parti = 0; parti < NumRegs; ++parti) {
3188 bool aggregateIsPacked =
false;
3189 if (
StructType *STy = dyn_cast<StructType>(Ty))
3190 aggregateIsPacked = STy->isPacked();
3201 SDValue Arg = getParamSymbol(DAG, i, PtrVT);
3203 for (
unsigned parti = 0, parte = VTs.
size(); parti != parte; ++parti) {
3205 assert(VecIdx == -1 &&
"Orphaned vector.");
3210 if (VectorInfo[parti] &
PVF_LAST) {
3211 unsigned NumElts = parti - VecIdx + 1;
3212 EVT EltVT = VTs[parti];
3215 if (EltVT == MVT::i1)
3217 else if (
Isv2x16VT(EltVT) || EltVT == MVT::v4i8)
3231 if (aggregateIsPacked)
3234 return std::nullopt;
3246 P.getNode()->setIROrder(i + 1);
3247 for (
unsigned j = 0; j < NumElts; ++j) {
3251 if (EltVT == MVT::i1)
3254 else if (EltVT != LoadVT)
3266 Ins[InsIdx].VT.getFixedSizeInBits() >
3270 Elt = DAG.
getNode(Extend, dl, Ins[InsIdx].VT, Elt);
3293 assert(ObjectVT == Ins[InsIdx].VT &&
3294 "Ins type did not match function type");
3295 SDValue Arg = getParamSymbol(DAG, i, PtrVT);
3298 p.getNode()->setIROrder(i + 1);
3302 if (!OutChains.empty())
3318 for (
unsigned i = 0, n = ElementType.getSizeInBits() / 8; i < n; i++) {
3328 DAG.
getVTList(MVT::Other), StoreOperands,
3346 assert(isABI &&
"Non-ABI compilation is not supported");
3355 assert(VTs.
size() == OutVals.
size() &&
"Bad return value decomposition");
3357 for (
unsigned i = 0, e = VTs.
size(); i != e; ++i) {
3358 SDValue PromotedOutVal = OutVals[i];
3361 VTs[i] =
EVT(PromotedVT);
3366 PromotedOutVal = DAG.
getNode(Ext, dl, PromotedVT, PromotedOutVal);
3368 PromotedOutVals.
push_back(PromotedOutVal);
3379 bool ExtendIntegerRetVal =
3380 RetTy->isIntegerTy() &&
DL.getTypeAllocSizeInBits(
RetTy) < 32;
3383 for (
unsigned i = 0, e = VTs.
size(); i != e; ++i) {
3385 SDValue RetVal = PromotedOutVals[i];
3387 if (ExtendIntegerRetVal) {
3390 dl, MVT::i32, RetVal);
3400 EVT ElementType = ExtendIntegerRetVal ? MVT::i32 : VTs[i];
3401 Align ElementTypeAlign =
3402 DL.getABITypeAlign(ElementType.getTypeForEVT(
RetTy->getContext()));
3403 Align ElementAlign =
3405 if (ElementAlign < ElementTypeAlign) {
3406 assert(StoreOperands.
empty() &&
"Orphaned operand list.");
3418 assert(StoreOperands.
empty() &&
"Orphaned operand list.");
3429 unsigned NumElts = StoreOperands.
size() - 2;
3446 EVT TheStoreType = ExtendIntegerRetVal ? MVT::i32 : VTs[i];
3448 Op, dl, DAG.
getVTList(MVT::Other), StoreOperands, TheStoreType,
3451 StoreOperands.
clear();
3461 if (Constraint.
size() > 1)
3467 switch (Intrinsic) {
3471 case Intrinsic::nvvm_tex_1d_v4f32_s32:
3473 case Intrinsic::nvvm_tex_1d_v4f32_f32:
3475 case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
3477 case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
3479 case Intrinsic::nvvm_tex_1d_v4s32_s32:
3481 case Intrinsic::nvvm_tex_1d_v4s32_f32:
3483 case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
3485 case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
3487 case Intrinsic::nvvm_tex_1d_v4u32_s32:
3489 case Intrinsic::nvvm_tex_1d_v4u32_f32:
3491 case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
3493 case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
3496 case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
3498 case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
3500 case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
3502 case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
3504 case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
3506 case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
3508 case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
3510 case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
3512 case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
3514 case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
3516 case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
3518 case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
3521 case Intrinsic::nvvm_tex_2d_v4f32_s32:
3523 case Intrinsic::nvvm_tex_2d_v4f32_f32:
3525 case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
3527 case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
3529 case Intrinsic::nvvm_tex_2d_v4s32_s32:
3531 case Intrinsic::nvvm_tex_2d_v4s32_f32:
3533 case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
3535 case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
3537 case Intrinsic::nvvm_tex_2d_v4u32_s32:
3539 case Intrinsic::nvvm_tex_2d_v4u32_f32:
3541 case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
3543 case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
3546 case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
3548 case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
3550 case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
3552 case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
3554 case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
3556 case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
3558 case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
3560 case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
3562 case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
3564 case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
3566 case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
3568 case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
3571 case Intrinsic::nvvm_tex_3d_v4f32_s32:
3573 case Intrinsic::nvvm_tex_3d_v4f32_f32:
3575 case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
3577 case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
3579 case Intrinsic::nvvm_tex_3d_v4s32_s32:
3581 case Intrinsic::nvvm_tex_3d_v4s32_f32:
3583 case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
3585 case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
3587 case Intrinsic::nvvm_tex_3d_v4u32_s32:
3589 case Intrinsic::nvvm_tex_3d_v4u32_f32:
3591 case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
3593 case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
3596 case Intrinsic::nvvm_tex_cube_v4f32_f32:
3598 case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
3600 case Intrinsic::nvvm_tex_cube_v4s32_f32:
3602 case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
3604 case Intrinsic::nvvm_tex_cube_v4u32_f32:
3606 case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
3609 case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
3611 case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
3613 case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
3615 case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
3617 case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
3619 case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
3622 case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
3624 case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
3626 case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
3628 case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
3630 case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
3632 case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
3634 case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
3636 case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
3638 case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
3640 case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
3642 case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
3644 case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
3647 case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
3649 case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
3651 case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
3653 case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
3655 case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
3657 case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
3659 case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
3661 case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
3663 case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
3665 case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
3667 case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
3669 case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
3672 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
3674 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
3676 case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
3678 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
3680 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
3682 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
3684 case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
3686 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
3688 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
3690 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
3692 case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
3694 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
3697 case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
3699 case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
3701 case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
3703 case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
3705 case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
3707 case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
3709 case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
3711 case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
3713 case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
3715 case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
3717 case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
3719 case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
3722 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
3724 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
3726 case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
3728 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
3730 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
3732 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
3734 case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
3736 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
3738 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
3740 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
3742 case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
3744 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
3747 case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
3749 case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
3751 case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
3753 case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
3755 case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
3757 case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
3759 case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
3761 case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
3763 case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
3765 case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
3767 case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
3769 case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
3772 case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
3774 case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
3776 case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
3778 case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
3780 case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
3782 case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
3785 case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
3787 case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
3789 case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
3791 case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
3793 case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
3795 case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
3798 case Intrinsic::nvvm_tex_unified_cube_grad_v4f32_f32:
3800 case Intrinsic::nvvm_tex_unified_cube_grad_v4s32_f32:
3802 case Intrinsic::nvvm_tex_unified_cube_grad_v4u32_f32:
3804 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4f32_f32:
3806 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4s32_f32:
3808 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4u32_f32:
3811 case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
3813 case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
3815 case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
3817 case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
3819 case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
3821 case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
3823 case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
3825 case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
3827 case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
3829 case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
3831 case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
3833 case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
3839 switch (Intrinsic) {
3842 case Intrinsic::nvvm_suld_1d_i8_clamp:
3844 case Intrinsic::nvvm_suld_1d_i16_clamp:
3846 case Intrinsic::nvvm_suld_1d_i32_clamp:
3848 case Intrinsic::nvvm_suld_1d_i64_clamp:
3850 case Intrinsic::nvvm_suld_1d_v2i8_clamp:
3852 case Intrinsic::nvvm_suld_1d_v2i16_clamp:
3854 case Intrinsic::nvvm_suld_1d_v2i32_clamp:
3856 case Intrinsic::nvvm_suld_1d_v2i64_clamp:
3858 case Intrinsic::nvvm_suld_1d_v4i8_clamp:
3860 case Intrinsic::nvvm_suld_1d_v4i16_clamp:
3862 case Intrinsic::nvvm_suld_1d_v4i32_clamp:
3864 case Intrinsic::nvvm_suld_1d_array_i8_clamp:
3866 case Intrinsic::nvvm_suld_1d_array_i16_clamp:
3868 case Intrinsic::nvvm_suld_1d_array_i32_clamp:
3870 case Intrinsic::nvvm_suld_1d_array_i64_clamp:
3872 case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
3874 case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
3876 case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
3878 case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
3880 case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
3882 case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
3884 case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
3886 case Intrinsic::nvvm_suld_2d_i8_clamp:
3888 case Intrinsic::nvvm_suld_2d_i16_clamp:
3890 case Intrinsic::nvvm_suld_2d_i32_clamp:
3892 case Intrinsic::nvvm_suld_2d_i64_clamp:
3894 case Intrinsic::nvvm_suld_2d_v2i8_clamp:
3896 case Intrinsic::nvvm_suld_2d_v2i16_clamp:
3898 case Intrinsic::nvvm_suld_2d_v2i32_clamp:
3900 case Intrinsic::nvvm_suld_2d_v2i64_clamp:
3902 case Intrinsic::nvvm_suld_2d_v4i8_clamp:
3904 case Intrinsic::nvvm_suld_2d_v4i16_clamp:
3906 case Intrinsic::nvvm_suld_2d_v4i32_clamp:
3908 case Intrinsic::nvvm_suld_2d_array_i8_clamp:
3910 case Intrinsic::nvvm_suld_2d_array_i16_clamp:
3912 case Intrinsic::nvvm_suld_2d_array_i32_clamp:
3914 case Intrinsic::nvvm_suld_2d_array_i64_clamp:
3916 case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
3918 case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
3920 case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
3922 case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
3924 case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
3926 case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
3928 case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
3930 case Intrinsic::nvvm_suld_3d_i8_clamp:
3932 case Intrinsic::nvvm_suld_3d_i16_clamp:
3934 case Intrinsic::nvvm_suld_3d_i32_clamp:
3936 case Intrinsic::nvvm_suld_3d_i64_clamp:
3938 case Intrinsic::nvvm_suld_3d_v2i8_clamp:
3940 case Intrinsic::nvvm_suld_3d_v2i16_clamp:
3942 case Intrinsic::nvvm_suld_3d_v2i32_clamp:
3944 case Intrinsic::nvvm_suld_3d_v2i64_clamp:
3946 case Intrinsic::nvvm_suld_3d_v4i8_clamp:
3948 case Intrinsic::nvvm_suld_3d_v4i16_clamp:
3950 case Intrinsic::nvvm_suld_3d_v4i32_clamp:
3952 case Intrinsic::nvvm_suld_1d_i8_trap:
3954 case Intrinsic::nvvm_suld_1d_i16_trap:
3956 case Intrinsic::nvvm_suld_1d_i32_trap:
3958 case Intrinsic::nvvm_suld_1d_i64_trap:
3960 case Intrinsic::nvvm_suld_1d_v2i8_trap:
3962 case Intrinsic::nvvm_suld_1d_v2i16_trap:
3964 case Intrinsic::nvvm_suld_1d_v2i32_trap:
3966 case Intrinsic::nvvm_suld_1d_v2i64_trap:
3968 case Intrinsic::nvvm_suld_1d_v4i8_trap:
3970 case Intrinsic::nvvm_suld_1d_v4i16_trap:
3972 case Intrinsic::nvvm_suld_1d_v4i32_trap:
3974 case Intrinsic::nvvm_suld_1d_array_i8_trap:
3976 case Intrinsic::nvvm_suld_1d_array_i16_trap:
3978 case Intrinsic::nvvm_suld_1d_array_i32_trap:
3980 case Intrinsic::nvvm_suld_1d_array_i64_trap:
3982 case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
3984 case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
3986 case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
3988 case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
3990 case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
3992 case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
3994 case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
3996 case Intrinsic::nvvm_suld_2d_i8_trap:
3998 case Intrinsic::nvvm_suld_2d_i16_trap:
4000 case Intrinsic::nvvm_suld_2d_i32_trap:
4002 case Intrinsic::nvvm_suld_2d_i64_trap:
4004 case Intrinsic::nvvm_suld_2d_v2i8_trap:
4006 case Intrinsic::nvvm_suld_2d_v2i16_trap:
4008 case Intrinsic::nvvm_suld_2d_v2i32_trap:
4010 case Intrinsic::nvvm_suld_2d_v2i64_trap:
4012 case Intrinsic::nvvm_suld_2d_v4i8_trap:
4014 case Intrinsic::nvvm_suld_2d_v4i16_trap:
4016 case Intrinsic::nvvm_suld_2d_v4i32_trap:
4018 case Intrinsic::nvvm_suld_2d_array_i8_trap:
4020 case Intrinsic::nvvm_suld_2d_array_i16_trap:
4022 case Intrinsic::nvvm_suld_2d_array_i32_trap:
4024 case Intrinsic::nvvm_suld_2d_array_i64_trap:
4026 case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
4028 case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
4030 case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
4032 case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
4034 case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
4036 case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
4038 case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
4040 case Intrinsic::nvvm_suld_3d_i8_trap:
4042 case Intrinsic::nvvm_suld_3d_i16_trap:
4044 case Intrinsic::nvvm_suld_3d_i32_trap:
4046 case Intrinsic::nvvm_suld_3d_i64_trap:
4048 case Intrinsic::nvvm_suld_3d_v2i8_trap:
4050 case Intrinsic::nvvm_suld_3d_v2i16_trap:
4052 case Intrinsic::nvvm_suld_3d_v2i32_trap:
4054 case Intrinsic::nvvm_suld_3d_v2i64_trap:
4056 case Intrinsic::nvvm_suld_3d_v4i8_trap:
4058 case Intrinsic::nvvm_suld_3d_v4i16_trap:
4060 case Intrinsic::nvvm_suld_3d_v4i32_trap:
4062 case Intrinsic::nvvm_suld_1d_i8_zero:
4064 case Intrinsic::nvvm_suld_1d_i16_zero:
4066 case Intrinsic::nvvm_suld_1d_i32_zero:
4068 case Intrinsic::nvvm_suld_1d_i64_zero:
4070 case Intrinsic::nvvm_suld_1d_v2i8_zero:
4072 case Intrinsic::nvvm_suld_1d_v2i16_zero:
4074 case Intrinsic::nvvm_suld_1d_v2i32_zero:
4076 case Intrinsic::nvvm_suld_1d_v2i64_zero:
4078 case Intrinsic::nvvm_suld_1d_v4i8_zero:
4080 case Intrinsic::nvvm_suld_1d_v4i16_zero:
4082 case Intrinsic::nvvm_suld_1d_v4i32_zero:
4084 case Intrinsic::nvvm_suld_1d_array_i8_zero:
4086 case Intrinsic::nvvm_suld_1d_array_i16_zero:
4088 case Intrinsic::nvvm_suld_1d_array_i32_zero:
4090 case Intrinsic::nvvm_suld_1d_array_i64_zero:
4092 case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
4094 case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
4096 case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
4098 case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
4100 case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
4102 case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
4104 case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
4106 case Intrinsic::nvvm_suld_2d_i8_zero:
4108 case Intrinsic::nvvm_suld_2d_i16_zero:
4110 case Intrinsic::nvvm_suld_2d_i32_zero:
4112 case Intrinsic::nvvm_suld_2d_i64_zero:
4114 case Intrinsic::nvvm_suld_2d_v2i8_zero:
4116 case Intrinsic::nvvm_suld_2d_v2i16_zero:
4118 case Intrinsic::nvvm_suld_2d_v2i32_zero:
4120 case Intrinsic::nvvm_suld_2d_v2i64_zero:
4122 case Intrinsic::nvvm_suld_2d_v4i8_zero:
4124 case Intrinsic::nvvm_suld_2d_v4i16_zero:
4126 case Intrinsic::nvvm_suld_2d_v4i32_zero:
4128 case Intrinsic::nvvm_suld_2d_array_i8_zero:
4130 case Intrinsic::nvvm_suld_2d_array_i16_zero:
4132 case Intrinsic::nvvm_suld_2d_array_i32_zero:
4134 case Intrinsic::nvvm_suld_2d_array_i64_zero:
4136 case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
4138 case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
4140 case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
4142 case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
4144 case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
4146 case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
4148 case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
4150 case Intrinsic::nvvm_suld_3d_i8_zero:
4152 case Intrinsic::nvvm_suld_3d_i16_zero:
4154 case Intrinsic::nvvm_suld_3d_i32_zero:
4156 case Intrinsic::nvvm_suld_3d_i64_zero:
4158 case Intrinsic::nvvm_suld_3d_v2i8_zero:
4160 case Intrinsic::nvvm_suld_3d_v2i16_zero:
4162 case Intrinsic::nvvm_suld_3d_v2i32_zero:
4164 case Intrinsic::nvvm_suld_3d_v2i64_zero:
4166 case Intrinsic::nvvm_suld_3d_v4i8_zero:
4168 case Intrinsic::nvvm_suld_3d_v4i16_zero:
4170 case Intrinsic::nvvm_suld_3d_v4i32_zero:
4183 switch (Intrinsic) {
4186 case Intrinsic::nvvm_match_all_sync_i32p:
4187 case Intrinsic::nvvm_match_all_sync_i64p:
4192 Info.memVT = MVT::i1;
4197 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col:
4198 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row:
4199 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col_stride:
4200 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row_stride:
4201 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col:
4202 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row:
4203 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col_stride:
4204 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row_stride:
4205 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col:
4206 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row:
4207 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col_stride:
4208 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row_stride:
4209 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col:
4210 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row:
4211 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col_stride:
4212 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row_stride:
4213 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col:
4214 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row:
4215 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col_stride:
4216 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row_stride:
4217 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col:
4218 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row:
4219 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col_stride:
4220 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row_stride: {
4222 Info.memVT = MVT::v8f16;
4223 Info.ptrVal =
I.getArgOperand(0);
4229 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col:
4230 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col_stride:
4231 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col_stride:
4232 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col:
4233 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row:
4234 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row_stride:
4235 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row_stride:
4236 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row:
4237 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col:
4238 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col_stride:
4239 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row:
4240 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row_stride:
4241 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col:
4242 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col_stride:
4243 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col_stride:
4244 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col:
4245 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row:
4246 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row_stride:
4247 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row_stride:
4248 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row:
4249 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col:
4250 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col_stride:
4251 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row:
4252 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row_stride: {
4254 Info.memVT = MVT::v2i32;
4255 Info.ptrVal =
I.getArgOperand(0);
4262 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col:
4263 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col_stride:
4264 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col_stride:
4265 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col:
4266 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row:
4267 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row_stride:
4268 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row_stride:
4269 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row:
4270 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col:
4271 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col_stride:
4272 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row:
4273 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row_stride:
4274 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col:
4275 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col_stride:
4276 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row:
4277 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row_stride:
4279 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col:
4280 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col_stride:
4281 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col_stride:
4282 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col:
4283 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row:
4284 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row_stride:
4285 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row_stride:
4286 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row:
4287 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col:
4288 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col_stride:
4289 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row:
4290 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row_stride:
4291 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col:
4292 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col_stride:
4293 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row:
4294 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row_stride:
4295 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16:
4296 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_trans_b16: {
4298 Info.memVT = MVT::v4i32;
4299 Info.ptrVal =
I.getArgOperand(0);
4306 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col:
4307 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col_stride:
4308 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col_stride:
4309 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col:
4310 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row:
4311 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row_stride:
4312 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row_stride:
4313 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row:
4315 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col:
4316 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col_stride:
4317 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col_stride:
4318 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col:
4319 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row:
4320 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row_stride:
4321 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row_stride:
4322 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row:
4323 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row:
4324 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row_stride:
4325 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col:
4326 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col_stride:
4327 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row:
4328 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row_stride:
4329 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row_stride:
4330 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row:
4331 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col:
4332 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col_stride:
4333 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col_stride:
4334 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col:
4335 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16:
4336 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_trans_b16: {
4338 Info.memVT = MVT::i32;
4339 Info.ptrVal =
I.getArgOperand(0);
4346 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col:
4347 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row:
4348 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col_stride:
4349 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row_stride:
4350 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col:
4351 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row:
4352 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col_stride:
4353 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row_stride:
4354 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col:
4355 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row:
4356 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col_stride:
4357 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row_stride: {
4359 Info.memVT = MVT::v4f16;
4360 Info.ptrVal =
I.getArgOperand(0);
4367 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col:
4368 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row:
4369 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col_stride:
4370 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row_stride:
4371 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col:
4372 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row:
4373 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col_stride:
4374 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row_stride:
4375 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col:
4376 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row:
4377 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col_stride:
4378 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row_stride:
4379 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col:
4380 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row:
4381 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col_stride:
4382 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row_stride: {
4384 Info.memVT = MVT::v8f32;
4385 Info.ptrVal =
I.getArgOperand(0);
4392 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col:
4393 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col_stride:
4394 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row:
4395 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row_stride:
4397 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col:
4398 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col_stride:
4399 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row:
4400 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row_stride:
4402 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col:
4403 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col_stride:
4404 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row:
4405 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row_stride:
4406 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col:
4407 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col_stride:
4408 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row:
4409 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row_stride:
4410 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col:
4411 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col_stride:
4412 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row:
4413 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row_stride: {
4415 Info.memVT = MVT::v8i32;
4416 Info.ptrVal =
I.getArgOperand(0);
4423 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col:
4424 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col_stride:
4425 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row:
4426 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row_stride:
4427 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col:
4428 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col_stride:
4429 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row:
4430 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row_stride:
4431 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16:
4432 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_trans_b16: {
4434 Info.memVT = MVT::v2i32;
4435 Info.ptrVal =
I.getArgOperand(0);
4442 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col:
4443 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col_stride:
4444 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row:
4445 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row_stride:
4447 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col:
4448 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col_stride:
4449 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row:
4450 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row_stride: {
4452 Info.memVT = MVT::f64;
4453 Info.ptrVal =
I.getArgOperand(0);
4460 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col:
4461 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col_stride:
4462 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row:
4463 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row_stride: {
4465 Info.memVT = MVT::v2f64;
4466 Info.ptrVal =
I.getArgOperand(0);
4473 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col:
4474 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row:
4475 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col_stride:
4476 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row_stride:
4477 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col:
4478 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row:
4479 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col_stride:
4480 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row_stride:
4481 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col:
4482 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row:
4483 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col_stride:
4484 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row_stride: {
4486 Info.memVT = MVT::v4f16;
4487 Info.ptrVal =
I.getArgOperand(0);
4494 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col:
4495 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row:
4496 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col_stride:
4497 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row_stride:
4498 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col:
4499 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row:
4500 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col_stride:
4501 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row_stride:
4502 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col:
4503 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row:
4504 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col_stride:
4505 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row_stride:
4506 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col:
4507 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row:
4508 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col_stride:
4509 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row_stride: {
4511 Info.memVT = MVT::v8f32;
4512 Info.ptrVal =
I.getArgOperand(0);
4519 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col:
4520 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col_stride:
4521 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row:
4522 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row_stride:
4523 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col:
4524 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col_stride:
4525 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row:
4526 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row_stride:
4527 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col:
4528 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col_stride:
4529 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row:
4530 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row_stride: {
4532 Info.memVT = MVT::v8i32;
4533 Info.ptrVal =
I.getArgOperand(0);
4540 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col:
4541 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col_stride:
4542 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row:
4543 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row_stride:
4544 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col:
4545 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col_stride:
4546 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row:
4547 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row_stride: {
4549 Info.memVT = MVT::v2i32;
4550 Info.ptrVal =
I.getArgOperand(0);
4557 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col:
4558 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col_stride:
4559 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row:
4560 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row_stride: {
4562 Info.memVT = MVT::v2f64;
4563 Info.ptrVal =
I.getArgOperand(0);
4570 case Intrinsic::nvvm_atomic_load_inc_32:
4571 case Intrinsic::nvvm_atomic_load_dec_32:
4573 case Intrinsic::nvvm_atomic_add_gen_f_cta:
4574 case Intrinsic::nvvm_atomic_add_gen_f_sys:
4575 case Intrinsic::nvvm_atomic_add_gen_i_cta:
4576 case Intrinsic::nvvm_atomic_add_gen_i_sys:
4577 case Intrinsic::nvvm_atomic_and_gen_i_cta:
4578 case Intrinsic::nvvm_atomic_and_gen_i_sys:
4579 case Intrinsic::nvvm_atomic_cas_gen_i_cta:
4580 case Intrinsic::nvvm_atomic_cas_gen_i_sys:
4581 case Intrinsic::nvvm_atomic_dec_gen_i_cta:
4582 case Intrinsic::nvvm_atomic_dec_gen_i_sys:
4583 case Intrinsic::nvvm_atomic_inc_gen_i_cta:
4584 case Intrinsic::nvvm_atomic_inc_gen_i_sys:
4585 case Intrinsic::nvvm_atomic_max_gen_i_cta:
4586 case Intrinsic::nvvm_atomic_max_gen_i_sys:
4587 case Intrinsic::nvvm_atomic_min_gen_i_cta:
4588 case Intrinsic::nvvm_atomic_min_gen_i_sys:
4589 case Intrinsic::nvvm_atomic_or_gen_i_cta:
4590 case Intrinsic::nvvm_atomic_or_gen_i_sys:
4591 case Intrinsic::nvvm_atomic_exch_gen_i_cta:
4592 case Intrinsic::nvvm_atomic_exch_gen_i_sys:
4593 case Intrinsic::nvvm_atomic_xor_gen_i_cta:
4594 case Intrinsic::nvvm_atomic_xor_gen_i_sys: {
4595 auto &
DL =
I.getModule()->getDataLayout();
4598 Info.ptrVal =
I.getArgOperand(0);
4605 case Intrinsic::nvvm_ldu_global_i:
4606 case Intrinsic::nvvm_ldu_global_f:
4607 case Intrinsic::nvvm_ldu_global_p: {
4608 auto &
DL =
I.getModule()->getDataLayout();
4610 if (Intrinsic == Intrinsic::nvvm_ldu_global_i)
4612 else if(Intrinsic == Intrinsic::nvvm_ldu_global_p)
4616 Info.ptrVal =
I.getArgOperand(0);
4619 Info.align = cast<ConstantInt>(
I.getArgOperand(1))->getMaybeAlignValue();
4623 case Intrinsic::nvvm_ldg_global_i:
4624 case Intrinsic::nvvm_ldg_global_f:
4625 case Intrinsic::nvvm_ldg_global_p: {
4626 auto &
DL =
I.getModule()->getDataLayout();
4629 if (Intrinsic == Intrinsic::nvvm_ldg_global_i)
4631 else if(Intrinsic == Intrinsic::nvvm_ldg_global_p)
4635 Info.ptrVal =
I.getArgOperand(0);
4638 Info.align = cast<ConstantInt>(
I.getArgOperand(1))->getMaybeAlignValue();
4643 case Intrinsic::nvvm_tex_1d_v4f32_s32:
4644 case Intrinsic::nvvm_tex_1d_v4f32_f32:
4645 case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
4646 case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
4647 case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
4648 case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
4649 case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
4650 case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
4651 case Intrinsic::nvvm_tex_2d_v4f32_s32:
4652 case Intrinsic::nvvm_tex_2d_v4f32_f32:
4653 case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
4654 case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
4655 case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
4656 case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
4657 case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
4658 case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
4659 case Intrinsic::nvvm_tex_3d_v4f32_s32:
4660 case Intrinsic::nvvm_tex_3d_v4f32_f32:
4661 case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
4662 case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
4663 case Intrinsic::nvvm_tex_cube_v4f32_f32:
4664 case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
4665 case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
4666 case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
4667 case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
4668 case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
4669 case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
4670 case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
4671 case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
4672 case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
4673 case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
4674 case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
4675 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
4676 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
4677 case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
4678 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
4679 case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
4680 case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
4681 case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
4682 case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
4683 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
4684 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
4685 case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
4686 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
4687 case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
4688 case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
4689 case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
4690 case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
4691 case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
4692 case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
4693 case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
4694 case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
4695 case Intrinsic::nvvm_tex_unified_cube_grad_v4f32_f32:
4696 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4f32_f32:
4697 case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
4698 case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
4699 case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
4700 case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
4702 Info.memVT = MVT::v4f32;
4703 Info.ptrVal =
nullptr;
4709 case Intrinsic::nvvm_tex_1d_v4s32_s32:
4710 case Intrinsic::nvvm_tex_1d_v4s32_f32:
4711 case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
4712 case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
4713 case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
4714 case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
4715 case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
4716 case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
4717 case Intrinsic::nvvm_tex_2d_v4s32_s32:
4718 case Intrinsic::nvvm_tex_2d_v4s32_f32:
4719 case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
4720 case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
4721 case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
4722 case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
4723 case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
4724 case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
4725 case Intrinsic::nvvm_tex_3d_v4s32_s32:
4726 case Intrinsic::nvvm_tex_3d_v4s32_f32:
4727 case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
4728 case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
4729 case Intrinsic::nvvm_tex_cube_v4s32_f32:
4730 case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
4731 case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
4732 case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
4733 case Intrinsic::nvvm_tex_cube_v4u32_f32:
4734 case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
4735 case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
4736 case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
4737 case Intrinsic::nvvm_tex_1d_v4u32_s32:
4738 case Intrinsic::nvvm_tex_1d_v4u32_f32:
4739 case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
4740 case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
4741 case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
4742 case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
4743 case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
4744 case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
4745 case Intrinsic::nvvm_tex_2d_v4u32_s32:
4746 case Intrinsic::nvvm_tex_2d_v4u32_f32:
4747 case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
4748 case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
4749 case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
4750 case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
4751 case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
4752 case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
4753 case Intrinsic::nvvm_tex_3d_v4u32_s32:
4754 case Intrinsic::nvvm_tex_3d_v4u32_f32:
4755 case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
4756 case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
4757 case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
4758 case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
4759 case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
4760 case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
4761 case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
4762 case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
4763 case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
4764 case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
4765 case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
4766 case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
4767 case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
4768 case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
4769 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
4770 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
4771 case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
4772 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
4773 case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
4774 case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
4775 case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
4776 case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
4777 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
4778 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
4779 case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
4780 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
4781 case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
4782 case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
4783 case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
4784 case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
4785 case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
4786 case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
4787 case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
4788 case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
4789 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
4790 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
4791 case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
4792 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
4793 case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
4794 case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
4795 case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
4796 case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
4797 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
4798 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
4799 case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
4800 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
4801 case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
4802 case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
4803 case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
4804 case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
4805 case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
4806 case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
4807 case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
4808 case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
4809 case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
4810 case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
4811 case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
4812 case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
4813 case Intrinsic::nvvm_tex_unified_cube_grad_v4s32_f32:
4814 case Intrinsic::nvvm_tex_unified_cube_grad_v4u32_f32:
4815 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4s32_f32:
4816 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4u32_f32:
4817 case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
4818 case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
4819 case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
4820 case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
4821 case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
4822 case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
4823 case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
4824 case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
4826 Info.memVT = MVT::v4i32;
4827 Info.ptrVal =
nullptr;
4833 case Intrinsic::nvvm_suld_1d_i8_clamp:
4834 case Intrinsic::nvvm_suld_1d_v2i8_clamp:
4835 case Intrinsic::nvvm_suld_1d_v4i8_clamp:
4836 case Intrinsic::nvvm_suld_1d_array_i8_clamp:
4837 case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
4838 case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
4839 case Intrinsic::nvvm_suld_2d_i8_clamp:
4840 case Intrinsic::nvvm_suld_2d_v2i8_clamp:
4841 case Intrinsic::nvvm_suld_2d_v4i8_clamp:
4842 case Intrinsic::nvvm_suld_2d_array_i8_clamp:
4843 case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
4844 case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
4845 case Intrinsic::nvvm_suld_3d_i8_clamp:
4846 case Intrinsic::nvvm_suld_3d_v2i8_clamp:
4847 case Intrinsic::nvvm_suld_3d_v4i8_clamp:
4848 case Intrinsic::nvvm_suld_1d_i8_trap:
4849 case Intrinsic::nvvm_suld_1d_v2i8_trap:
4850 case Intrinsic::nvvm_suld_1d_v4i8_trap:
4851 case Intrinsic::nvvm_suld_1d_array_i8_trap:
4852 case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
4853 case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
4854 case Intrinsic::nvvm_suld_2d_i8_trap:
4855 case Intrinsic::nvvm_suld_2d_v2i8_trap:
4856 case Intrinsic::nvvm_suld_2d_v4i8_trap:
4857 case Intrinsic::nvvm_suld_2d_array_i8_trap:
4858 case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
4859 case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
4860 case Intrinsic::nvvm_suld_3d_i8_trap:
4861 case Intrinsic::nvvm_suld_3d_v2i8_trap:
4862 case Intrinsic::nvvm_suld_3d_v4i8_trap:
4863 case Intrinsic::nvvm_suld_1d_i8_zero:
4864 case Intrinsic::nvvm_suld_1d_v2i8_zero:
4865 case Intrinsic::nvvm_suld_1d_v4i8_zero:
4866 case Intrinsic::nvvm_suld_1d_array_i8_zero:
4867 case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
4868 case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
4869 case Intrinsic::nvvm_suld_2d_i8_zero:
4870 case Intrinsic::nvvm_suld_2d_v2i8_zero:
4871 case Intrinsic::nvvm_suld_2d_v4i8_zero:
4872 case Intrinsic::nvvm_suld_2d_array_i8_zero:
4873 case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
4874 case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
4875 case Intrinsic::nvvm_suld_3d_i8_zero:
4876 case Intrinsic::nvvm_suld_3d_v2i8_zero:
4877 case Intrinsic::nvvm_suld_3d_v4i8_zero:
4879 Info.memVT = MVT::i8;
4880 Info.ptrVal =
nullptr;
4886 case Intrinsic::nvvm_suld_1d_i16_clamp:
4887 case Intrinsic::nvvm_suld_1d_v2i16_clamp:
4888 case Intrinsic::nvvm_suld_1d_v4i16_clamp:
4889 case Intrinsic::nvvm_suld_1d_array_i16_clamp:
4890 case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
4891 case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
4892 case Intrinsic::nvvm_suld_2d_i16_clamp:
4893 case Intrinsic::nvvm_suld_2d_v2i16_clamp:
4894 case Intrinsic::nvvm_suld_2d_v4i16_clamp:
4895 case Intrinsic::nvvm_suld_2d_array_i16_clamp:
4896 case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
4897 case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
4898 case Intrinsic::nvvm_suld_3d_i16_clamp:
4899 case Intrinsic::nvvm_suld_3d_v2i16_clamp:
4900 case Intrinsic::nvvm_suld_3d_v4i16_clamp:
4901 case Intrinsic::nvvm_suld_1d_i16_trap:
4902 case Intrinsic::nvvm_suld_1d_v2i16_trap:
4903 case Intrinsic::nvvm_suld_1d_v4i16_trap:
4904 case Intrinsic::nvvm_suld_1d_array_i16_trap:
4905 case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
4906 case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
4907 case Intrinsic::nvvm_suld_2d_i16_trap:
4908 case Intrinsic::nvvm_suld_2d_v2i16_trap:
4909 case Intrinsic::nvvm_suld_2d_v4i16_trap:
4910 case Intrinsic::nvvm_suld_2d_array_i16_trap:
4911 case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
4912 case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
4913 case Intrinsic::nvvm_suld_3d_i16_trap:
4914 case Intrinsic::nvvm_suld_3d_v2i16_trap:
4915 case Intrinsic::nvvm_suld_3d_v4i16_trap:
4916 case Intrinsic::nvvm_suld_1d_i16_zero:
4917 case Intrinsic::nvvm_suld_1d_v2i16_zero:
4918 case Intrinsic::nvvm_suld_1d_v4i16_zero:
4919 case Intrinsic::nvvm_suld_1d_array_i16_zero:
4920 case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
4921 case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
4922 case Intrinsic::nvvm_suld_2d_i16_zero:
4923 case Intrinsic::nvvm_suld_2d_v2i16_zero:
4924 case Intrinsic::nvvm_suld_2d_v4i16_zero:
4925 case Intrinsic::nvvm_suld_2d_array_i16_zero:
4926 case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
4927 case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
4928 case Intrinsic::nvvm_suld_3d_i16_zero:
4929 case Intrinsic::nvvm_suld_3d_v2i16_zero:
4930 case Intrinsic::nvvm_suld_3d_v4i16_zero:
4932 Info.memVT = MVT::i16;
4933 Info.ptrVal =
nullptr;
4939 case Intrinsic::nvvm_suld_1d_i32_clamp:
4940 case Intrinsic::nvvm_suld_1d_v2i32_clamp:
4941 case Intrinsic::nvvm_suld_1d_v4i32_clamp:
4942 case Intrinsic::nvvm_suld_1d_array_i32_clamp:
4943 case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
4944 case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
4945 case Intrinsic::nvvm_suld_2d_i32_clamp:
4946 case Intrinsic::nvvm_suld_2d_v2i32_clamp:
4947 case Intrinsic::nvvm_suld_2d_v4i32_clamp:
4948 case Intrinsic::nvvm_suld_2d_array_i32_clamp:
4949 case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
4950 case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
4951 case Intrinsic::nvvm_suld_3d_i32_clamp:
4952 case Intrinsic::nvvm_suld_3d_v2i32_clamp:
4953 case Intrinsic::nvvm_suld_3d_v4i32_clamp:
4954 case Intrinsic::nvvm_suld_1d_i32_trap:
4955 case Intrinsic::nvvm_suld_1d_v2i32_trap:
4956 case Intrinsic::nvvm_suld_1d_v4i32_trap:
4957 case Intrinsic::nvvm_suld_1d_array_i32_trap:
4958 case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
4959 case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
4960 case Intrinsic::nvvm_suld_2d_i32_trap:
4961 case Intrinsic::nvvm_suld_2d_v2i32_trap:
4962 case Intrinsic::nvvm_suld_2d_v4i32_trap:
4963 case Intrinsic::nvvm_suld_2d_array_i32_trap:
4964 case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
4965 case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
4966 case Intrinsic::nvvm_suld_3d_i32_trap:
4967 case Intrinsic::nvvm_suld_3d_v2i32_trap:
4968 case Intrinsic::nvvm_suld_3d_v4i32_trap:
4969 case Intrinsic::nvvm_suld_1d_i32_zero:
4970 case Intrinsic::nvvm_suld_1d_v2i32_zero:
4971 case Intrinsic::nvvm_suld_1d_v4i32_zero:
4972 case Intrinsic::nvvm_suld_1d_array_i32_zero:
4973 case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
4974 case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
4975 case Intrinsic::nvvm_suld_2d_i32_zero:
4976 case Intrinsic::nvvm_suld_2d_v2i32_zero:
4977 case Intrinsic::nvvm_suld_2d_v4i32_zero:
4978 case Intrinsic::nvvm_suld_2d_array_i32_zero:
4979 case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
4980 case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
4981 case Intrinsic::nvvm_suld_3d_i32_zero:
4982 case Intrinsic::nvvm_suld_3d_v2i32_zero:
4983 case Intrinsic::nvvm_suld_3d_v4i32_zero:
4985 Info.memVT = MVT::i32;
4986 Info.ptrVal =
nullptr;
4992 case Intrinsic::nvvm_suld_1d_i64_clamp:
4993 case Intrinsic::nvvm_suld_1d_v2i64_clamp:
4994 case Intrinsic::nvvm_suld_1d_array_i64_clamp:
4995 case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
4996 case Intrinsic::nvvm_suld_2d_i64_clamp:
4997 case Intrinsic::nvvm_suld_2d_v2i64_clamp:
4998 case Intrinsic::nvvm_suld_2d_array_i64_clamp:
4999 case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
5000 case Intrinsic::nvvm_suld_3d_i64_clamp:
5001 case Intrinsic::nvvm_suld_3d_v2i64_clamp:
5002 case Intrinsic::nvvm_suld_1d_i64_trap:
5003 case Intrinsic::nvvm_suld_1d_v2i64_trap:
5004 case Intrinsic::nvvm_suld_1d_array_i64_trap:
5005 case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
5006 case Intrinsic::nvvm_suld_2d_i64_trap:
5007 case Intrinsic::nvvm_suld_2d_v2i64_trap:
5008 case Intrinsic::nvvm_suld_2d_array_i64_trap:
5009 case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
5010 case Intrinsic::nvvm_suld_3d_i64_trap:
5011 case Intrinsic::nvvm_suld_3d_v2i64_trap:
5012 case Intrinsic::nvvm_suld_1d_i64_zero:
5013 case Intrinsic::nvvm_suld_1d_v2i64_zero:
5014 case Intrinsic::nvvm_suld_1d_array_i64_zero:
5015 case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
5016 case Intrinsic::nvvm_suld_2d_i64_zero:
5017 case Intrinsic::nvvm_suld_2d_v2i64_zero:
5018 case Intrinsic::nvvm_suld_2d_array_i64_zero:
5019 case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
5020 case Intrinsic::nvvm_suld_3d_i64_zero:
5021 case Intrinsic::nvvm_suld_3d_v2i64_zero:
5023 Info.memVT = MVT::i64;
5024 Info.ptrVal =
nullptr;
5042 const uint64_t ABITypeAlign =
DL.getABITypeAlign(ArgTy).value();
5047 if (!
F || !
F->hasLocalLinkage() ||
5048 F->hasAddressTaken(
nullptr,
5052 return Align(ABITypeAlign);
5062 Align ArgAlign = InitialAlign;
5077 ArgAlign = std::max(ArgAlign,
Align(4));
5087 std::string ParamName;
5092 ParamStr <<
"_vararg";
5094 ParamStr <<
"_param_" <<
Idx;
5143 if (Constraint.
size() == 1) {
5144 switch (Constraint[0]) {
5162std::pair<unsigned, const TargetRegisterClass *>
5166 if (Constraint.
size() == 1) {
5167 switch (Constraint[0]) {
5169 return std::make_pair(0U, &NVPTX::Int1RegsRegClass);
5171 return std::make_pair(0U, &NVPTX::Int16RegsRegClass);
5173 return std::make_pair(0U, &NVPTX::Int16RegsRegClass);
5175 return std::make_pair(0U, &NVPTX::Int32RegsRegClass);
5178 return std::make_pair(0U, &NVPTX::Int64RegsRegClass);
5180 return std::make_pair(0U, &NVPTX::Float32RegsRegClass);
5182 return std::make_pair(0U, &NVPTX::Float64RegsRegClass);
5216 return F.getFnAttribute(
"unsafe-fp-math").getValueAsBool();
5249 if (VT == MVT::f32 || VT == MVT::f64) {
5264 int nonAddCount = 0;
5273 int orderNo =
N->getIROrder();
5279 if (orderNo - orderNo2 < 500)
5284 bool opIsLive =
false;
5288 if (isa<ConstantSDNode>(left) || isa<ConstantSDNode>(right))
5293 int orderNo3 =
User->getIROrder();
5294 if (orderNo3 > orderNo) {
5302 int orderNo3 =
User->getIROrder();
5303 if (orderNo3 > orderNo) {
5323 for (std::size_t
I = 2, OpsCount =
N->ops().size();
I != OpsCount; ++
I)
5324 if (!
N->getOperand(
I).isUndef())
5329 return N->getOperand(0);
5360 if (isa<ConstantSDNode>(Val)) {
5374 ConstantSDNode *BFEBits = dyn_cast<ConstantSDNode>(BFE.getOperand(0));
5386 if (MaskVal != (
uint64_t(1) << BFEBitsVal) - 1)
5410 if (MaskVal != 0xff) {
5415 MemSDNode *Mem = dyn_cast<MemSDNode>(Val);
5422 if (MemVT != MVT::v2i8 && MemVT != MVT::v4i8) {
5435 if (AExt.
getNode() !=
nullptr) {
5460 EVT VT =
N->getValueType(0);
5464 const SDValue &Num =
N->getOperand(0);
5465 const SDValue &Den =
N->getOperand(1);
5468 if (U->getOpcode() == DivOpc && U->getOperand(0) == Num &&
5469 U->getOperand(1) == Den) {
5496 EVT OrigVT =
Op.getOperand(0).getValueType();
5502 EVT OrigVT =
Op.getOperand(0).getValueType();
5529 IsSigned = (LHSSign ==
Signed);
5533 const APInt &Val = CI->getAPIntValue();
5535 return Val.
isIntN(OptSize);
5544 return LHSSign == RHSSign;
5554 EVT MulType =
N->getValueType(0);
5555 if (MulType != MVT::i32 && MulType != MVT::i64) {
5566 if (isa<ConstantSDNode>(
LHS)) {
5595 if (MulType == MVT::i32) {
5596 DemotedVT = MVT::i16;
5598 DemotedVT = MVT::i32;
5615 return DCI.
DAG.
getNode(Opc,
DL, MulType, TruncLHS, TruncRHS);
5647 EVT CCType =
N->getValueType(0);
5651 EVT AType =
A.getValueType();
5652 if (!(CCType == MVT::v2i1 && (AType == MVT::v2f16 || AType == MVT::v2bf16)))
5655 if (
A.getValueType() == MVT::v2bf16 &&
SmVersion < 90)
5666 DL, DCI.
DAG.
getVTList(MVT::i1, MVT::i1), {A, B, N->getOperand(2)});
5683 VectorVT == MVT::v4i8 || VectorVT == MVT::v8i8)
5688 if (!(VectorBits == 16 || VectorBits == 32 || VectorBits == 64))
5708 if (EltVT != EltIVT)
5711 if (EltVT !=
N->getValueType(0))
5721 if (VectorVT != MVT::v4i8)
5732 for (
int I = 0;
I < 4; ++
I) {
5759 EVT VT =
N->getValueType(0);
5760 if (VT != MVT::v16i8)
5767 EVT NewVT = MVT::v4i32;
5770 EVT RetVTs[] = {EltVT, EltVT, EltVT, EltVT, MVT::Other};
5775 LD->getMemOperand());
5780 for (
unsigned i = 0; i < NumElts; i++)
5789 DAGCombinerInfo &DCI)
const {
5791 switch (
N->getOpcode()) {
5824 EVT ResVT =
N->getValueType(0);
5857 Align Alignment = LD->getAlign();
5861 if (Alignment < PrefAlign) {
5876 bool NeedTrunc =
false;
5882 unsigned Opcode = 0;
5884 bool Load16x2 =
false;
5891 LdResVTs = DAG.
getVTList(EltVT, EltVT, MVT::Other);
5895 EVT ListVTs[] = { EltVT, EltVT, EltVT, EltVT, MVT::Other };
5920 EVT ListVTs[] = {VVT, VVT, VVT, VVT, MVT::Other};
5935 LD->getMemOperand());
5941 for (
unsigned i = 0; i < NumElts; ++i) {
5951 for (
unsigned i = 0; i < NumElts; ++i) {
5978 case Intrinsic::nvvm_ldg_global_i:
5979 case Intrinsic::nvvm_ldg_global_f:
5980 case Intrinsic::nvvm_ldg_global_p:
5981 case Intrinsic::nvvm_ldu_global_i:
5982 case Intrinsic::nvvm_ldu_global_f:
5983 case Intrinsic::nvvm_ldu_global_p: {
5984 EVT ResVT =
N->getValueType(0);
5996 bool NeedTrunc =
false;
6002 unsigned Opcode = 0;
6012 case Intrinsic::nvvm_ldg_global_i:
6013 case Intrinsic::nvvm_ldg_global_f:
6014 case Intrinsic::nvvm_ldg_global_p:
6017 case Intrinsic::nvvm_ldu_global_i:
6018 case Intrinsic::nvvm_ldu_global_f:
6019 case Intrinsic::nvvm_ldu_global_p:
6023 LdResVTs = DAG.
getVTList(EltVT, EltVT, MVT::Other);
6029 case Intrinsic::nvvm_ldg_global_i:
6030 case Intrinsic::nvvm_ldg_global_f:
6031 case Intrinsic::nvvm_ldg_global_p:
6034 case Intrinsic::nvvm_ldu_global_i:
6035 case Intrinsic::nvvm_ldu_global_f:
6036 case Intrinsic::nvvm_ldu_global_p:
6040 EVT ListVTs[] = { EltVT, EltVT, EltVT, EltVT, MVT::Other };
6053 OtherOps.
append(
N->op_begin() + 2,
N->op_end());
6063 for (
unsigned i = 0; i < NumElts; ++i) {
6081 "Custom handling of non-i8 ldu/ldg?");
6105void NVPTXTargetLowering::ReplaceNodeResults(
6107 switch (
N->getOpcode()) {
6137 auto ITy = cast<llvm::IntegerType>(Ty);
6146 switch (ITy->getBitWidth()) {
6165 switch (ITy->getBitWidth()) {
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
amdgpu AMDGPU Register Bank Select
This file implements a class to represent arbitrary precision integral constant values and operations...
static SDValue PerformLOADCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const ARMSubtarget *Subtarget)
static SDValue PerformADDCombineWithOperands(SDNode *N, SDValue N0, SDValue N1, TargetLowering::DAGCombinerInfo &DCI, const ARMSubtarget *Subtarget)
PerformADDCombineWithOperands - Try DAG combinations for an ADD with operands N0 and N1.
static SDValue PerformADDCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const ARMSubtarget *Subtarget)
PerformADDCombine - Target-specific dag combine xforms for ISD::ADD.
static SDValue PerformVSELECTCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const ARMSubtarget *Subtarget)
static SDValue PerformMULCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const ARMSubtarget *Subtarget)
static SDValue PerformANDCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const ARMSubtarget *Subtarget)
Function Alias Analysis Results
This file contains the simple types necessary to represent the attributes associated with functions a...
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
Analysis containing CSE Info
This file contains the declarations for the subclasses of Constant, which represent the different fla...
Returns the sub type a function will return at a given Idx Should correspond to the result type of an ExtractValue instruction executed with just that one unsigned Idx
This file contains the declarations of entities that describe floating point environment and related ...
static DebugLoc getDebugLoc(MachineBasicBlock::instr_iterator FirstMI, MachineBasicBlock::instr_iterator LastMI)
Return the first found DebugLoc that has a DILocation, given a range of instructions.
unsigned const TargetRegisterInfo * TRI
Module.h This file contains the declarations for the Module class.
static cl::opt< bool > sched4reg("nvptx-sched4reg", cl::desc("NVPTX Specific: schedule for register pressue"), cl::init(false))
static SDValue PerformEXTRACTCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
static cl::opt< unsigned > FMAContractLevelOpt("nvptx-fma-level", cl::Hidden, cl::desc("NVPTX Specific: FMA contraction (0: don't do it" " 1: do it 2: do it aggressively"), cl::init(2))
static bool IsPTXVectorType(MVT VT)
static cl::opt< int > UsePrecDivF32("nvptx-prec-divf32", cl::Hidden, cl::desc("NVPTX Specifies: 0 use div.approx, 1 use div.full, 2 use" " IEEE Compliant F32 div.rnd if available."), cl::init(2))
static void ReplaceLoadVector(SDNode *N, SelectionDAG &DAG, SmallVectorImpl< SDValue > &Results)
ReplaceVectorLoad - Convert vector loads into multi-output scalar loads.
static unsigned getOpcForSurfaceInstr(unsigned Intrinsic)
static bool Is16bitsType(MVT VT)
static bool IsTypePassedAsArray(const Type *Ty)
static SmallVector< ParamVectorizationFlags, 16 > VectorizePTXValueVTs(const SmallVectorImpl< EVT > &ValueVTs, const SmallVectorImpl< uint64_t > &Offsets, Align ParamAlignment, bool IsVAArg=false)
static unsigned CanMergeParamLoadStoresStartingAt(unsigned Idx, uint32_t AccessSize, const SmallVectorImpl< EVT > &ValueVTs, const SmallVectorImpl< uint64_t > &Offsets, Align ParamAlignment)
static void ReplaceINTRINSIC_W_CHAIN(SDNode *N, SelectionDAG &DAG, SmallVectorImpl< SDValue > &Results)
static unsigned getOpcForTextureInstr(unsigned Intrinsic)
static SDValue LowerVectorArith(SDValue Op, SelectionDAG &DAG)
static void ComputePTXValueVTs(const TargetLowering &TLI, const DataLayout &DL, Type *Ty, SmallVectorImpl< EVT > &ValueVTs, SmallVectorImpl< uint64_t > *Offsets=nullptr, uint64_t StartingOffset=0)
ComputePTXValueVTs - For the given Type Ty, returns the set of primitive EVTs that compose it.
static bool IsMulWideOperandDemotable(SDValue Op, unsigned OptSize, OperandSignedness &S)
IsMulWideOperandDemotable - Checks if the provided DAG node is an operand that can be demoted to OptS...
static SDValue LowerUnalignedStoreParam(SelectionDAG &DAG, SDValue Chain, uint64_t Offset, EVT ElementType, SDValue StVal, SDValue &InGlue, unsigned ArgID, const SDLoc &dl)
static SDValue PerformREMCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
static SDValue PerformStoreRetvalCombine(SDNode *N)
static bool AreMulWideOperandsDemotable(SDValue LHS, SDValue RHS, unsigned OptSize, bool &IsSigned)
AreMulWideOperandsDemotable - Checks if the given LHS and RHS operands can be demoted to OptSize bits...
static bool adjustElementType(EVT &ElementType)
static SDValue TryMULWIDECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
TryMULWIDECombine - Attempt to replace a multiply of M bits with a multiply of M/2 bits that produces...
static cl::opt< bool > UsePrecSqrtF32("nvptx-prec-sqrtf32", cl::Hidden, cl::desc("NVPTX Specific: 0 use sqrt.approx, 1 use sqrt.rn."), cl::init(true))
static SDValue LowerUnalignedStoreRet(SelectionDAG &DAG, SDValue Chain, uint64_t Offset, EVT ElementType, SDValue RetVal, const SDLoc &dl)
static bool PromoteScalarIntegerPTX(const EVT &VT, MVT *PromotedVT)
PromoteScalarIntegerPTX Used to make sure the arguments/returns are suitable for passing and promote ...
static SDValue PerformSETCCCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, unsigned int SmVersion)
static SDValue LowerUnalignedLoadRetParam(SelectionDAG &DAG, SDValue &Chain, uint64_t Offset, EVT ElementType, SDValue &InGlue, SmallVectorImpl< SDValue > &TempProxyRegOps, const SDLoc &dl)
static std::atomic< unsigned > GlobalUniqueCallSite
static cl::opt< bool > ForceMinByValParamAlign("nvptx-force-min-byval-param-align", cl::Hidden, cl::desc("NVPTX Specific: force 4-byte minimal alignment for byval" " params of device functions."), cl::init(false))
static SDValue PerformSHLCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
PerformSHLCombine - Runs PTX-specific DAG combine patterns on SHL nodes.
const char LLVMTargetMachineRef TM
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
This file defines the SmallVector class.
This file describes how to lower LLVM code to machine code.
Class for arbitrary precision integers.
bool isSignedIntN(unsigned N) const
Check if this APInt has an N-bits signed integer value.
bool slt(const APInt &RHS) const
Signed less than comparison.
bool isIntN(unsigned N) const
Check if this APInt has an N-bits unsigned integer value.
bool sge(const APInt &RHS) const
Signed greater or equal comparison.
This class represents an incoming formal argument to a Function.
an instruction that atomically reads a memory location, combines it with another value,...
@ Min
*p = old <signed v ? old : v
@ Max
*p = old >signed v ? old : v
@ UMin
*p = old <unsigned v ? old : v
@ UMax
*p = old >unsigned v ? old : v
bool isFloatingPointOperation() const
BinOp getOperation() const
MaybeAlign getParamAlignment(unsigned ArgNo) const
Return the alignment for the specified function parameter.
bool hasParamAttr(unsigned ArgNo, Attribute::AttrKind Kind) const
Return true if the attribute exists for the given argument.
Base class for all callable instructions (InvokeInst and CallInst) Holds everything related to callin...
Function * getCalledFunction() const
Returns the function called, or null if this is an indirect function invocation or the function signa...
This class represents a function call, abstracting a target machine's calling convention.
uint64_t getZExtValue() const
const APInt & getAPIntValue() const
static Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
This class represents an Operation in the Expression.
uint64_t getNumOperands() const
A parsed version of the target data layout string in and methods for querying it.
TypeSize getTypeAllocSize(Type *Ty) const
Returns the offset in bytes between successive objects of the specified type, including alignment pad...
Align getPrefTypeAlign(Type *Ty) const
Returns the preferred stack/global alignment for the specified type.
Diagnostic information for unsupported feature in backend.
void addFnAttr(Attribute::AttrKind Kind)
Add function attributes to this function.
Type * getReturnType() const
Returns the type of the ret val.
unsigned getAddressSpace() const
const GlobalValue * getGlobal() const
const Function * getFunction() const
Return the function this instruction belongs to.
void diagnose(const DiagnosticInfo &DI)
Report a message to the currently installed diagnostic handler.
This class is used to represent ISD::LOAD nodes.
MCSection * getDataSection() const
Instances of this class represent a uniqued identifier for a section in the current translation unit.
StringRef getName() const
getName - Get the symbol name.
unsigned getVectorNumElements() const
bool isScalableVector() const
Return true if this is a vector value type where the runtime length is machine dependent.
static auto integer_valuetypes()
static auto fixedlen_vector_valuetypes()
static MVT getIntegerVT(unsigned BitWidth)
MVT getScalarType() const
If this is a vector, return the element type, otherwise return this.
DenormalMode getDenormalMode(const fltSemantics &FPType) const
Returns the denormal handling type for the default rounding mode of the function.
Function & getFunction()
Return the LLVM function that this machine code represents.
const LLVMTargetMachine & getTarget() const
getTarget - Return the target machine this machine code is compiled with
@ MODereferenceable
The memory access is dereferenceable (i.e., doesn't trap).
@ MOLoad
The memory access reads data.
@ MOInvariant
The memory access always returns the same value (or traps).
@ MOStore
The memory access writes data.
This SDNode is used for target intrinsics that touch memory and need an associated MachineMemOperand.
This is an abstract virtual class for memory operations.
MachineMemOperand * getMemOperand() const
Return a MachineMemOperand object describing the memory reference performed by operation.
EVT getMemoryVT() const
Return the type of the in-memory value.
unsigned getMaxRequiredAlignment() const
bool hasAtomMinMax64() const
bool hasAtomAddF64() const
const NVPTXTargetLowering * getTargetLowering() const override
unsigned getPTXVersion() const
const NVPTXRegisterInfo * getRegisterInfo() const override
unsigned int getSmVersion() const
bool hasAtomBitwise64() const
bool allowFP16Math() const
ConstraintType getConstraintType(StringRef Constraint) const override
getConstraintType - Given a constraint letter, return the type of constraint it is for this target.
SDValue LowerOperation(SDValue Op, SelectionDAG &DAG) const override
This callback is invoked for operations that are unsupported by the target, which are registered to u...
const NVPTXTargetMachine * nvTM
SDValue LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const
NVPTXTargetLowering(const NVPTXTargetMachine &TM, const NVPTXSubtarget &STI)
bool useF32FTZ(const MachineFunction &MF) const
SDValue getSqrtEstimate(SDValue Operand, SelectionDAG &DAG, int Enabled, int &ExtraSteps, bool &UseOneConst, bool Reciprocal) const override
Hooks for building estimates in place of slower divisions and square roots.
SDValue LowerReturn(SDValue Chain, CallingConv::ID CallConv, bool isVarArg, const SmallVectorImpl< ISD::OutputArg > &Outs, const SmallVectorImpl< SDValue > &OutVals, const SDLoc &dl, SelectionDAG &DAG) const override
This hook must be implemented to lower outgoing return values, described by the Outs array,...
SDValue LowerFormalArguments(SDValue Chain, CallingConv::ID CallConv, bool isVarArg, const SmallVectorImpl< ISD::InputArg > &Ins, const SDLoc &dl, SelectionDAG &DAG, SmallVectorImpl< SDValue > &InVals) const override
This hook must be implemented to lower the incoming (formal) arguments, described by the Ins array,...
void LowerAsmOperandForConstraint(SDValue Op, StringRef Constraint, std::vector< SDValue > &Ops, SelectionDAG &DAG) const override
Lower the specified operand into the Ops vector.
std::string getParamName(const Function *F, int Idx) const
TargetLoweringBase::LegalizeTypeAction getPreferredVectorAction(MVT VT) const override
Return the preferred vector type legalization action.
std::string getPrototype(const DataLayout &DL, Type *, const ArgListTy &, const SmallVectorImpl< ISD::OutputArg > &, MaybeAlign retAlignment, std::optional< std::pair< unsigned, const APInt & > > VAInfo, const CallBase &CB, unsigned UniqueCallSite) const
Align getFunctionParamOptimizedAlign(const Function *F, Type *ArgTy, const DataLayout &DL) const
getFunctionParamOptimizedAlign - since function arguments are passed via .param space,...
SDValue LowerDYNAMIC_STACKALLOC(SDValue Op, SelectionDAG &DAG) const
EVT getSetCCResultType(const DataLayout &DL, LLVMContext &Ctx, EVT VT) const override
Return the ValueType of the result of SETCC operations.
std::pair< unsigned, const TargetRegisterClass * > getRegForInlineAsmConstraint(const TargetRegisterInfo *TRI, StringRef Constraint, MVT VT) const override
Given a physical register constraint (e.g.
bool isLegalAddressingMode(const DataLayout &DL, const AddrMode &AM, Type *Ty, unsigned AS, Instruction *I=nullptr) const override
isLegalAddressingMode - Return true if the addressing mode represented by AM is legal for this target...
AtomicExpansionKind shouldExpandAtomicRMWInIR(AtomicRMWInst *AI) const override
Returns how the IR-level AtomicExpand pass should expand the given AtomicRMW, if at all.
Align getFunctionByValParamAlign(const Function *F, Type *ArgTy, Align InitialAlign, const DataLayout &DL) const
Helper for computing alignment of a device function byval parameter.
bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallInst &I, MachineFunction &MF, unsigned Intrinsic) const override
Given an intrinsic, checks if on the target the intrinsic will need to map to a MemIntrinsicNode (tou...
const char * getTargetNodeName(unsigned Opcode) const override
This method returns the name of a target specific DAG node.
bool allowFMA(MachineFunction &MF, CodeGenOptLevel OptLevel) const
bool usePrecSqrtF32() const
bool allowUnsafeFPMath(MachineFunction &MF) const
int getDivF32Level() const
SDValue LowerCall(CallLoweringInfo &CLI, SmallVectorImpl< SDValue > &InVals) const override
This hook must be implemented to lower calls into the specified DAG.
UniqueStringSaver & getStrPool() const
MCSection * SelectSectionForGlobal(const GlobalObject *GO, SectionKind Kind, const TargetMachine &TM) const override
~NVPTXTargetObjectFile() override
static PointerType * get(Type *ElementType, unsigned AddressSpace)
This constructs a pointer to an object of the specified type in a numbered address space.
Wrapper class for IR location info (IR ordering and DebugLoc) to be passed into SDNode creation funct...
Represents one node in the SelectionDAG.
const APInt & getAsAPIntVal() const
Helper method returns the APInt value of a ConstantSDNode.
bool isMachineOpcode() const
Test if this node has a post-isel opcode, directly corresponding to a MachineInstr opcode.
unsigned getOpcode() const
Return the SelectionDAG opcode value for this node.
bool hasOneUse() const
Return true if there is exactly one use of this node.
unsigned getIROrder() const
Return the node ordering.
iterator_range< use_iterator > uses()
uint64_t getAsZExtVal() const
Helper method returns the zero-extended integer value of a ConstantSDNode.
unsigned getNumOperands() const
Return the number of values used by this operation.
unsigned getMachineOpcode() const
This may only be called if isMachineOpcode returns true.
SDVTList getVTList() const
const SDValue & getOperand(unsigned Num) const
uint64_t getConstantOperandVal(unsigned Num) const
Helper method returns the integer value of a ConstantSDNode operand.
const APInt & getConstantOperandAPInt(unsigned Num) const
Helper method returns the APInt of a ConstantSDNode operand.
EVT getValueType(unsigned ResNo) const
Return the type of a specified result.
bool isUndef() const
Return true if the type of the node type undefined.
Represents a use of a SDNode.
Unlike LLVM values, Selection DAG nodes may return multiple values as the result of a computation.
SDNode * getNode() const
get the SDNode which holds the desired result
SDValue getValue(unsigned R) const
EVT getValueType() const
Return the ValueType of the referenced return value.
TypeSize getValueSizeInBits() const
Returns the size of the value in bits.
const SDValue & getOperand(unsigned i) const
MVT getSimpleValueType() const
Return the simple ValueType of the referenced return value.
unsigned getOpcode() const
SectionKind - This is a simple POD value that classifies the properties of a section.
This is used to represent a portion of an LLVM function in a low-level Data Dependence DAG representa...
SDValue getTargetGlobalAddress(const GlobalValue *GV, const SDLoc &DL, EVT VT, int64_t offset=0, unsigned TargetFlags=0)
const SDValue & getRoot() const
Return the root tag of the SelectionDAG.
SDValue getMergeValues(ArrayRef< SDValue > Ops, const SDLoc &dl)
Create a MERGE_VALUES node from the given operands.
SDVTList getVTList(EVT VT)
Return an SDVTList that represents the list of values specified.
SDValue getSetCC(const SDLoc &DL, EVT VT, SDValue LHS, SDValue RHS, ISD::CondCode Cond, SDValue Chain=SDValue(), bool IsSignaling=false)
Helper function to make it easier to build SetCC's if you just have an ISD::CondCode instead of an SD...
SDValue getSymbolFunctionGlobalAddress(SDValue Op, Function **TargetFunction=nullptr)
Return a GlobalAddress of the function from the current module with name matching the given ExternalS...
SDValue getConstantFP(double Val, const SDLoc &DL, EVT VT, bool isTarget=false)
Create a ConstantFPSDNode wrapping a constant value.
SDValue getLoad(EVT VT, const SDLoc &dl, SDValue Chain, SDValue Ptr, MachinePointerInfo PtrInfo, MaybeAlign Alignment=MaybeAlign(), MachineMemOperand::Flags MMOFlags=MachineMemOperand::MONone, const AAMDNodes &AAInfo=AAMDNodes(), const MDNode *Ranges=nullptr)
Loads are not normal binary operators: their result type is not determined by their operands,...
const TargetLowering & getTargetLoweringInfo() const
SDNode * MorphNodeTo(SDNode *N, unsigned Opc, SDVTList VTs, ArrayRef< SDValue > Ops)
This mutates the specified node to have the specified return type, opcode, and operands.
SDValue getCALLSEQ_END(SDValue Chain, SDValue Op1, SDValue Op2, SDValue InGlue, const SDLoc &DL)
Return a new CALLSEQ_END node, which always must have a glue result (to ensure it's not CSE'd).
SDValue getBuildVector(EVT VT, const SDLoc &DL, ArrayRef< SDValue > Ops)
Return an ISD::BUILD_VECTOR node.
SDValue getBitcast(EVT VT, SDValue V)
Return a bitcast using the SDLoc of the value operand, and casting to the provided type.
const DataLayout & getDataLayout() const
SDValue getConstant(uint64_t Val, const SDLoc &DL, EVT VT, bool isTarget=false, bool isOpaque=false)
Create a ConstantSDNode wrapping a constant value.
SDValue getTruncStore(SDValue Chain, const SDLoc &dl, SDValue Val, SDValue Ptr, MachinePointerInfo PtrInfo, EVT SVT, Align Alignment, MachineMemOperand::Flags MMOFlags=MachineMemOperand::MONone, const AAMDNodes &AAInfo=AAMDNodes())
void ReplaceAllUsesWith(SDValue From, SDValue To)
Modify anything using 'From' to use 'To' instead.
SDValue getStore(SDValue Chain, const SDLoc &dl, SDValue Val, SDValue Ptr, MachinePointerInfo PtrInfo, Align Alignment, MachineMemOperand::Flags MMOFlags=MachineMemOperand::MONone, const AAMDNodes &AAInfo=AAMDNodes())
Helper function to build ISD::STORE nodes.
SDValue getCALLSEQ_START(SDValue Chain, uint64_t InSize, uint64_t OutSize, const SDLoc &DL)
Return a new CALLSEQ_START node, that starts new call frame, in which InSize bytes are set up inside ...
void RemoveDeadNode(SDNode *N)
Remove the specified node from the system.
SDValue getAnyExtOrTrunc(SDValue Op, const SDLoc &DL, EVT VT)
Convert Op, which must be of integer type, to the integer type VT, by either any-extending or truncat...
SDValue getSelectCC(const SDLoc &DL, SDValue LHS, SDValue RHS, SDValue True, SDValue False, ISD::CondCode Cond)
Helper function to make it easier to build SelectCC's if you just have an ISD::CondCode instead of an...
SDValue getIntPtrConstant(uint64_t Val, const SDLoc &DL, bool isTarget=false)
SDValue getNode(unsigned Opcode, const SDLoc &DL, EVT VT, ArrayRef< SDUse > Ops)
Gets or creates the specified node.
SDValue getFPExtendOrRound(SDValue Op, const SDLoc &DL, EVT VT)
Convert Op, which must be of float type, to the float type VT, by either extending or rounding (by tr...
SDValue getTargetConstant(uint64_t Val, const SDLoc &DL, EVT VT, bool isOpaque=false)
MachineFunction & getMachineFunction() const
SDValue getZExtOrTrunc(SDValue Op, const SDLoc &DL, EVT VT)
Convert Op, which must be of integer type, to the integer type VT, by either zero-extending or trunca...
LLVMContext * getContext() const
const SDValue & setRoot(SDValue N)
Set the current root tag of the SelectionDAG.
SDValue getMemIntrinsicNode(unsigned Opcode, const SDLoc &dl, SDVTList VTList, ArrayRef< SDValue > Ops, EVT MemVT, MachinePointerInfo PtrInfo, Align Alignment, MachineMemOperand::Flags Flags=MachineMemOperand::MOLoad|MachineMemOperand::MOStore, LocationSize Size=0, const AAMDNodes &AAInfo=AAMDNodes())
Creates a MemIntrinsicNode that may produce a result and takes a list of operands.
SDValue getTargetExternalSymbol(const char *Sym, EVT VT, unsigned TargetFlags=0)
This SDNode is used to implement the code generator support for the llvm IR shufflevector instruction...
ArrayRef< int > getMask() const
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
void assign(size_type NumElts, ValueParamT Elt)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
This class is used to represent ISD::STORE nodes.
StringRef - Represent a constant reference to a string, i.e.
constexpr size_t size() const
size - Get the string size.
constexpr const char * data() const
data - Get a pointer to the start of the string (which may not be null terminated).
Class to represent struct types.
void setBooleanVectorContents(BooleanContent Ty)
Specify how the target extends the result of a vector boolean value from a vector of i1 to a wider ty...
void setOperationAction(unsigned Op, MVT VT, LegalizeAction Action)
Indicate that the specified operation does not work with the specified type and indicate what to do a...
EVT getValueType(const DataLayout &DL, Type *Ty, bool AllowUnknown=false) const
Return the EVT corresponding to this LLVM type.
LegalizeAction
This enum indicates whether operations are valid for a target, and if not, what action should be used...
unsigned MaxStoresPerMemcpyOptSize
Likewise for functions with the OptSize attribute.
const TargetMachine & getTargetMachine() const
LegalizeTypeAction
This enum indicates whether a types are legal for a target, and if not, what action should be used to...
void addBypassSlowDiv(unsigned int SlowBitWidth, unsigned int FastBitWidth)
Tells the code generator which bitwidths to bypass.
virtual unsigned getNumRegisters(LLVMContext &Context, EVT VT, std::optional< MVT > RegisterVT=std::nullopt) const
Return the number of registers that this ValueType will eventually require.
void setMaxAtomicSizeInBitsSupported(unsigned SizeInBits)
Set the maximum atomic operation size supported by the backend.
virtual TargetLoweringBase::LegalizeTypeAction getPreferredVectorAction(MVT VT) const
Return the preferred vector type legalization action.
unsigned MaxStoresPerMemsetOptSize
Likewise for functions with the OptSize attribute.
void setBooleanContents(BooleanContent Ty)
Specify how the target extends the result of integer and floating point boolean values from i1 to a w...
unsigned MaxStoresPerMemmove
Specify maximum number of store instructions per memmove call.
void computeRegisterProperties(const TargetRegisterInfo *TRI)
Once all of the register classes are added, this allows us to compute derived properties we expose.
unsigned MaxStoresPerMemmoveOptSize
Likewise for functions with the OptSize attribute.
void addRegisterClass(MVT VT, const TargetRegisterClass *RC)
Add the specified register class as an available regclass for the specified value type.
virtual MVT getPointerTy(const DataLayout &DL, uint32_t AS=0) const
Return the pointer type for the given address space, defaults to the pointer type from the data layou...
unsigned MaxStoresPerMemset
Specify maximum number of store instructions per memset call.
void setTruncStoreAction(MVT ValVT, MVT MemVT, LegalizeAction Action)
Indicate that the specified truncating store does not work with the specified type and indicate what ...
@ ZeroOrNegativeOneBooleanContent
void setMinCmpXchgSizeInBits(unsigned SizeInBits)
Sets the minimum cmpxchg or ll/sc size supported by the backend.
void AddPromotedToType(unsigned Opc, MVT OrigVT, MVT DestVT)
If Opc/OrigVT is specified as being promoted, the promotion code defaults to trying a larger integer/...
AtomicExpansionKind
Enum that specifies what an atomic load/AtomicRMWInst is expanded to, if at all.
void setTargetDAGCombine(ArrayRef< ISD::NodeType > NTs)
Targets should invoke this method for each target independent node that they want to provide a custom...
Align getMinStackArgumentAlignment() const
Return the minimum stack alignment of an argument.
void setLoadExtAction(unsigned ExtType, MVT ValVT, MVT MemVT, LegalizeAction Action)
Indicate that the specified load with extension does not work with the specified type and indicate wh...
std::vector< ArgListEntry > ArgListTy
bool allowsMemoryAccessForAlignment(LLVMContext &Context, const DataLayout &DL, EVT VT, unsigned AddrSpace=0, Align Alignment=Align(1), MachineMemOperand::Flags Flags=MachineMemOperand::MONone, unsigned *Fast=nullptr) const
This function returns true if the memory access is aligned or if the target allows this specific unal...
unsigned MaxStoresPerMemcpy
Specify maximum number of store instructions per memcpy call.
void setSchedulingPreference(Sched::Preference Pref)
Specify the target scheduling preference.
void setJumpIsExpensive(bool isExpensive=true)
Tells the code generator not to expand logic operations on comparison predicates into separate sequen...
LegalizeAction getOperationAction(unsigned Op, EVT VT) const
Return how this operation should be treated: either it is legal, needs to be promoted to a larger siz...
This class defines information used to lower LLVM code to legal SelectionDAG operators that the targe...
SDValue expandUnalignedStore(StoreSDNode *ST, SelectionDAG &DAG) const
Expands an unaligned store to 2 half-size stores for integer values, and possibly more for vectors.
virtual ConstraintType getConstraintType(StringRef Constraint) const
Given a constraint, return the type of constraint it is for this target.
std::pair< SDValue, SDValue > expandUnalignedLoad(LoadSDNode *LD, SelectionDAG &DAG) const
Expands an unaligned load to 2 half-size loads for an integer, and possibly more for vectors.
virtual std::pair< unsigned, const TargetRegisterClass * > getRegForInlineAsmConstraint(const TargetRegisterInfo *TRI, StringRef Constraint, MVT VT) const
Given a physical register constraint (e.g.
SDValue expandRoundInexactToOdd(EVT ResultVT, SDValue Op, const SDLoc &DL, SelectionDAG &DAG) const
Truncate Op to ResultVT.
SDValue expandFP_ROUND(SDNode *Node, SelectionDAG &DAG) const
Expand round(fp) to fp conversion.
virtual void LowerAsmOperandForConstraint(SDValue Op, StringRef Constraint, std::vector< SDValue > &Ops, SelectionDAG &DAG) const
Lower the specified operand into the Ops vector.
Primary interface to the complete machine description for the target machine.
MCSymbol * getSymbol(const GlobalValue *GV) const
CodeGenOptLevel getOptLevel() const
Returns the optimization level: None, Less, Default, or Aggressive.
unsigned UnsafeFPMath
UnsafeFPMath - This flag is enabled when the -enable-unsafe-fp-math flag is specified on the command ...
FPOpFusion::FPOpFusionMode AllowFPOpFusion
AllowFPOpFusion - This flag is set by the -fp-contract=xxx option.
TargetRegisterInfo base class - We assume that the target defines a static array of TargetRegisterDes...
The instances of the Type class are immutable: once they are created, they are never changed.
bool isVectorTy() const
True if this is an instance of VectorType.
bool isFloatTy() const
Return true if this is 'float', a 32-bit IEEE fp type.
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
@ VoidTyID
type with no size
bool isAggregateType() const
Return true if the type is an aggregate type.
bool isHalfTy() const
Return true if this is 'half', a 16-bit IEEE fp type.
bool isDoubleTy() const
Return true if this is 'double', a 64-bit IEEE fp type.
bool isFloatingPointTy() const
Return true if this is one of the floating-point types.
bool isIntegerTy() const
True if this is an instance of IntegerType.
TypeID getTypeID() const
Return the type id for the type.
TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
StringRef save(const char *S)
LLVM Value Representation.
Type * getType() const
All values are typed, get the type of this value.
A raw_ostream that writes to an std::string.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ C
The default llvm calling convention, compatible with C.
NodeType
ISD::NodeType enum - This enum defines the target-independent operators for a SelectionDAG.
@ SETCC
SetCC operator - This evaluates to a true value iff the condition is true.
@ SMUL_LOHI
SMUL_LOHI/UMUL_LOHI - Multiply two integers of type iN, producing a signed/unsigned value of type i[2...
@ BSWAP
Byte Swap and Counting operators.
@ VAEND
VAEND, VASTART - VAEND and VASTART have three operands: an input chain, pointer, and a SRCVALUE.
@ ADDC
Carry-setting nodes for multiple precision addition and subtraction.
@ ADD
Simple integer binary arithmetic operators.
@ LOAD
LOAD and STORE have token chains as their first operand, then the same operands as an LLVM load/store...
@ ANY_EXTEND
ANY_EXTEND - Used for integer types. The high bits are undefined.
@ FMA
FMA - Perform a * b + c with no intermediate rounding step.
@ INTRINSIC_VOID
OUTCHAIN = INTRINSIC_VOID(INCHAIN, INTRINSICID, arg1, arg2, ...) This node represents a target intrin...
@ SINT_TO_FP
[SU]INT_TO_FP - These operators convert integers (whose interpreted sign depends on the first letter)...
@ CONCAT_VECTORS
CONCAT_VECTORS(VECTOR0, VECTOR1, ...) - Given a number of values of vector type with the same length ...
@ FADD
Simple binary floating point operators.
@ ABS
ABS - Determine the unsigned absolute value of a signed integer value of the same bitwidth.
@ SDIVREM
SDIVREM/UDIVREM - Divide two integers and produce both a quotient and remainder result.
@ BITCAST
BITCAST - This operator converts between integer, vector and FP values, as if the value was stored to...
@ SIGN_EXTEND
Conversion operators.
@ READSTEADYCOUNTER
READSTEADYCOUNTER - This corresponds to the readfixedcounter intrinsic.
@ FNEG
Perform various unary floating-point operations inspired by libm.
@ BR_CC
BR_CC - Conditional branch.
@ SSUBO
Same for subtraction.
@ BRIND
BRIND - Indirect branch.
@ BR_JT
BR_JT - Jumptable branch.
@ SSUBSAT
RESULT = [US]SUBSAT(LHS, RHS) - Perform saturation subtraction on 2 integers with the same bit width ...
@ SELECT
Select(COND, TRUEVAL, FALSEVAL).
@ UNDEF
UNDEF - An undefined node.
@ VACOPY
VACOPY - VACOPY has 5 operands: an input chain, a destination pointer, a source pointer,...
@ SADDO
RESULT, BOOL = [SU]ADDO(LHS, RHS) - Overflow-aware nodes for addition.
@ MULHU
MULHU/MULHS - Multiply high - Multiply two integers of type iN, producing an unsigned/signed value of...
@ SHL
Shift and rotation operations.
@ VECTOR_SHUFFLE
VECTOR_SHUFFLE(VEC1, VEC2) - Returns a vector, of the same type as VEC1/VEC2.
@ EXTRACT_SUBVECTOR
EXTRACT_SUBVECTOR(VECTOR, IDX) - Returns a subvector from VECTOR.
@ EXTRACT_VECTOR_ELT
EXTRACT_VECTOR_ELT(VECTOR, IDX) - Returns a single element from VECTOR identified by the (potentially...
@ ZERO_EXTEND
ZERO_EXTEND - Used for integer types, zeroing the new bits.
@ SELECT_CC
Select with condition operator - This selects between a true value and a false value (ops #2 and #3) ...
@ FMINNUM
FMINNUM/FMAXNUM - Perform floating-point minimum or maximum on two values.
@ SSHLSAT
RESULT = [US]SHLSAT(LHS, RHS) - Perform saturation left shift.
@ SMULO
Same for multiplication.
@ DYNAMIC_STACKALLOC
DYNAMIC_STACKALLOC - Allocate some number of bytes on the stack aligned to a specified boundary.
@ SIGN_EXTEND_INREG
SIGN_EXTEND_INREG - This operator atomically performs a SHL/SRA pair to sign extend a small value in ...
@ SMIN
[US]{MIN/MAX} - Binary minimum or maximum of signed or unsigned integers.
@ FP_EXTEND
X = FP_EXTEND(Y) - Extend a smaller FP type into a larger FP type.
@ VSELECT
Select with a vector condition (op #0) and two vector operands (ops #1 and #2), returning a vector re...
@ UADDO_CARRY
Carry-using nodes for multiple precision addition and subtraction.
@ BF16_TO_FP
BF16_TO_FP, FP_TO_BF16 - These operators are used to perform promotions and truncation for bfloat16.
@ FRAMEADDR
FRAMEADDR, RETURNADDR - These nodes represent llvm.frameaddress and llvm.returnaddress on the DAG.
@ FMINIMUM
FMINIMUM/FMAXIMUM - NaN-propagating minimum/maximum that also treat -0.0 as less than 0....
@ FP_TO_SINT
FP_TO_[US]INT - Convert a floating point value to a signed or unsigned integer.
@ READCYCLECOUNTER
READCYCLECOUNTER - This corresponds to the readcyclecounter intrinsic.
@ AND
Bitwise operators - logical and, logical or, logical xor.
@ TRAP
TRAP - Trapping instruction.
@ INTRINSIC_WO_CHAIN
RESULT = INTRINSIC_WO_CHAIN(INTRINSICID, arg1, arg2, ...) This node represents a target intrinsic fun...
@ ADDE
Carry-using nodes for multiple precision addition and subtraction.
@ INSERT_VECTOR_ELT
INSERT_VECTOR_ELT(VECTOR, VAL, IDX) - Returns VECTOR with the element at IDX replaced with VAL.
@ TokenFactor
TokenFactor - This node takes multiple tokens as input and produces a single token result.
@ FP_ROUND
X = FP_ROUND(Y, TRUNC) - Rounding 'Y' from a larger floating point type down to the precision of the ...
@ TRUNCATE
TRUNCATE - Completely drop the high bits.
@ VAARG
VAARG - VAARG has four operands: an input chain, a pointer, a SRCVALUE, and the alignment.
@ SHL_PARTS
SHL_PARTS/SRA_PARTS/SRL_PARTS - These operators are used for expanded integer shift operations.
@ FCOPYSIGN
FCOPYSIGN(X, Y) - Return the value of X with the sign of Y.
@ SADDSAT
RESULT = [US]ADDSAT(LHS, RHS) - Perform saturation addition on 2 integers with the same bit width (W)...
@ SADDO_CARRY
Carry-using overflow-aware nodes for multiple precision addition and subtraction.
@ INTRINSIC_W_CHAIN
RESULT,OUTCHAIN = INTRINSIC_W_CHAIN(INCHAIN, INTRINSICID, arg1, ...) This node represents a target in...
@ BUILD_VECTOR
BUILD_VECTOR(ELT0, ELT1, ELT2, ELT3,...) - Return a fixed-width vector with the specified,...
@ Bitcast
Perform the operation on a different, but equivalently sized type.
@ TexUnified1DS32FloatLevel
@ Tex1DArrayFloatFloatLevel
@ TexUnified2DU32FloatGrad
@ Tld4UnifiedG2DFloatFloat
@ TexUnifiedCubeArrayFloatFloatLevel
@ Tld4UnifiedR2DFloatFloat
@ Tex2DArrayS32FloatLevel
@ TexUnified1DArrayFloatFloatLevel
@ TexUnified2DFloatFloatLevel
@ TexUnified3DFloatFloatLevel
@ TexUnified1DFloatFloatLevel
@ TexUnified2DArrayU32Float
@ TexUnified1DArrayFloatFloat
@ Tex1DArrayFloatFloatGrad
@ TexUnifiedCubeArrayU32FloatGrad
@ TexUnified1DFloatFloatGrad
@ TexUnifiedCubeFloatFloatGrad
@ TexUnified2DArrayFloatFloat
@ TexUnified3DU32FloatLevel
@ TexUnified1DArrayU32Float
@ TexUnified2DArrayFloatFloatLevel
@ TexUnified2DFloatFloatGrad
@ TexUnified2DArrayU32S32
@ TexUnifiedCubeArrayS32FloatLevel
@ TexUnified1DArrayS32Float
@ TexUnified1DArrayS32FloatLevel
@ TexUnified2DS32FloatLevel
@ TexUnified3DU32FloatGrad
@ TexUnifiedCubeU32FloatLevel
@ TexUnified2DArrayU32FloatGrad
@ TexUnifiedCubeFloatFloatLevel
@ TexUnified1DArrayFloatS32
@ TexUnifiedCubeS32FloatLevel
@ TexUnified1DS32FloatGrad
@ Tex2DArrayFloatFloatLevel
@ TexUnifiedCubeArrayFloatFloat
@ TexUnifiedCubeArrayFloatFloatGrad
@ TexUnifiedCubeFloatFloat
@ TexUnified1DArrayU32S32
@ TexUnified3DFloatFloatGrad
@ Tld4UnifiedA2DFloatFloat
@ TexUnified3DS32FloatGrad
@ TexUnified2DU32FloatLevel
@ TexUnified1DArrayS32S32
@ TexCubeArrayFloatFloatLevel
@ TexUnified1DU32FloatGrad
@ TexCubeArrayS32FloatLevel
@ Tex2DArrayU32FloatLevel
@ Tex1DArrayU32FloatLevel
@ TexUnified2DArrayU32FloatLevel
@ TexUnified1DArrayFloatFloatGrad
@ TexUnifiedCubeS32FloatGrad
@ TexCubeArrayU32FloatLevel
@ TexUnified3DS32FloatLevel
@ TexUnifiedCubeArrayS32FloatGrad
@ TexUnified2DArrayS32Float
@ Tex2DArrayFloatFloatGrad
@ TexUnifiedCubeArrayS32Float
@ TexUnified2DArrayS32FloatLevel
@ Tex1DArrayS32FloatLevel
@ TexUnifiedCubeArrayU32FloatLevel
@ TexUnified2DArrayS32S32
@ TexUnified2DArrayFloatFloatGrad
@ TexUnifiedCubeU32FloatGrad
@ Tld4UnifiedB2DFloatFloat
@ TexUnified1DArrayU32FloatLevel
@ TexUnified1DArrayS32FloatGrad
@ TexUnified2DS32FloatGrad
@ TexUnified2DArrayS32FloatGrad
@ TexUnified1DU32FloatLevel
@ TexUnifiedCubeArrayU32Float
@ TexUnified2DArrayFloatS32
@ TexUnified1DArrayU32FloatGrad
initializer< Ty > init(const Ty &Val)
This is an optimization pass for GlobalISel generic memory operations.
static bool isIndirectCall(const MachineInstr &MI)
bool shouldEmitPTXNoReturn(const Value *V, const TargetMachine &TM)
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
auto enumerate(FirstRange &&First, RestRanges &&...Rest)
Given two or more input ranges, returns a new range whose values are are tuples (A,...
bool getAlign(const Function &F, unsigned index, unsigned &align)
uint64_t PowerOf2Ceil(uint64_t A)
Returns the power of two which is greater than or equal to the given value.
OutputIt transform(R &&Range, OutputIt d_first, UnaryFunction F)
Wrapper function around std::transform to apply a function to a range and store the result elsewhere.
unsigned promoteScalarArgumentSize(unsigned size)
void report_fatal_error(Error Err, bool gen_crash_diag=true)
Report a serious error, calling any installed error handler.
CodeGenOptLevel
Code generation optimization level.
uint64_t alignTo(uint64_t Size, Align A)
Returns a multiple of A needed to store Size bytes.
DWARFExpression::Operation Op
void ComputeValueVTs(const TargetLowering &TLI, const DataLayout &DL, Type *Ty, SmallVectorImpl< EVT > &ValueVTs, SmallVectorImpl< EVT > *MemVTs, SmallVectorImpl< TypeSize > *Offsets=nullptr, TypeSize StartingOffset=TypeSize::getZero())
ComputeValueVTs - Given an LLVM IR type, compute a sequence of EVTs that represent all the individual...
constexpr unsigned BitWidth
bool isKernelFunction(const Function &F)
Function * getMaybeBitcastedCallee(const CallBase *CB)
Align commonAlignment(Align A, uint64_t Offset)
Returns the alignment that satisfies both alignments.
void swap(llvm::BitVector &LHS, llvm::BitVector &RHS)
Implement std::swap in terms of BitVector swap.
static const fltSemantics & IEEEsingle() LLVM_READNONE
This struct is a compact representation of a valid (non-zero power of two) alignment.
uint64_t value() const
This is a hole in the type system and should not be abused.
@ PreserveSign
The sign of a flushed-to-zero number is preserved in the sign of 0.
DenormalModeKind Output
Denormal flushing mode for floating point instruction results in the default floating point environme...
TypeSize getStoreSize() const
Return the number of bytes overwritten by a store of the specified value type.
bool isSimple() const
Test if the given EVT is simple (as opposed to being extended).
static EVT getVectorVT(LLVMContext &Context, EVT VT, unsigned NumElements, bool IsScalable=false)
Returns the EVT that represents a vector NumElements in length, where each element is of type VT.
EVT changeTypeToInteger() const
Return the type converted to an equivalently sized integer or vector with integer element type.
bool isFloatingPoint() const
Return true if this is a FP or a vector FP type.
TypeSize getSizeInBits() const
Return the size of the specified value type in bits.
uint64_t getScalarSizeInBits() const
MVT getSimpleVT() const
Return the SimpleValueType held in the specified simple EVT.
uint64_t getFixedSizeInBits() const
Return the size of the specified fixed width value type in bits.
bool isVector() const
Return true if this is a vector value type.
EVT getScalarType() const
If this is a vector type, return the element type, otherwise return this.
Type * getTypeForEVT(LLVMContext &Context) const
This method returns an LLVM type corresponding to the specified EVT.
EVT getVectorElementType() const
Given a vector type, return the type of each element.
bool isScalarInteger() const
Return true if this is an integer, but not a vector.
EVT changeVectorElementType(EVT EltVT) const
Return a VT for a vector type whose attributes match ourselves with the exception of the element type...
unsigned getVectorNumElements() const
Given a vector type, return the number of elements it contains.
bool isInteger() const
Return true if this is an integer or a vector integer type.
This class contains a discriminated union of information about pointers in memory operands,...
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
This represents a list of ValueType's that has been intern'd by a SelectionDAG.
This represents an addressing mode of: BaseGV + BaseOffs + BaseReg + Scale*ScaleReg + ScalableOffset*...
This structure contains all information that is necessary for lowering calls.
SmallVector< ISD::InputArg, 32 > Ins
SmallVector< ISD::OutputArg, 32 > Outs
SmallVector< SDValue, 32 > OutVals
SDValue CombineTo(SDNode *N, ArrayRef< SDValue > To, bool AddTo=true)