diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp index 99a7fdb9d1e21..4f3ee26990cef 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp @@ -612,10 +612,10 @@ bool NVPTXDAGToDAGISel::SelectSETP_F16X2(SDNode *N) { bool NVPTXDAGToDAGISel::tryEXTRACT_VECTOR_ELEMENT(SDNode *N) { SDValue Vector = N->getOperand(0); - // We only care about f16x2 as it's the only real vector type we + // We only care about 16x2 as it's the only real vector type we // need to deal with. MVT VT = Vector.getSimpleValueType(); - if (!(VT == MVT::v2f16 || VT == MVT::v2bf16)) + if (!Isv2x16VT(VT)) return false; // Find and record all uses of this vector that extract element 0 or 1. SmallVector E0, E1; @@ -828,6 +828,7 @@ pickOpcodeForVT(MVT::SimpleValueType VT, unsigned Opcode_i8, return Opcode_i16; case MVT::v2f16: case MVT::v2bf16: + case MVT::v2i16: return Opcode_i32; case MVT::f32: return Opcode_f32; @@ -909,9 +910,8 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) { // Vector Setting unsigned vecType = NVPTX::PTXLdStInstCode::Scalar; if (SimpleVT.isVector()) { - assert((LoadedVT == MVT::v2f16 || LoadedVT == MVT::v2bf16) && - "Unexpected vector type"); - // v2f16/v2bf16 is loaded using ld.b32 + assert(Isv2x16VT(LoadedVT) && "Unexpected vector type"); + // v2f16/v2bf16/v2i16 is loaded using ld.b32 fromTypeWidth = 32; } @@ -1061,10 +1061,10 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) { EVT EltVT = N->getValueType(0); - // v8f16 is a special case. PTX doesn't have ld.v8.f16 - // instruction. Instead, we split the vector into v2f16 chunks and + // v8x16 is a special case. PTX doesn't have ld.v8.16 + // instruction. Instead, we split the vector into v2x16 chunks and // load them with ld.v4.b32. - if (EltVT == MVT::v2f16 || EltVT == MVT::v2bf16) { + if (Isv2x16VT(EltVT)) { assert(N->getOpcode() == NVPTXISD::LoadV4 && "Unexpected load opcode."); EltVT = MVT::i32; FromType = NVPTX::PTXLdStInstCode::Untyped; @@ -1260,12 +1260,13 @@ bool NVPTXDAGToDAGISel::tryLDGLDU(SDNode *N) { if (EltVT.isVector()) { NumElts = EltVT.getVectorNumElements(); EltVT = EltVT.getVectorElementType(); - // vectors of f16 are loaded/stored as multiples of v2f16 elements. + // vectors of 16bits type are loaded/stored as multiples of v2x16 elements. if ((EltVT == MVT::f16 && N->getValueType(0) == MVT::v2f16) || - (EltVT == MVT::bf16 && N->getValueType(0) == MVT::v2bf16)) { - assert(NumElts % 2 == 0 && "Vector must have even number of elements"); - EltVT = N->getValueType(0); - NumElts /= 2; + (EltVT == MVT::bf16 && N->getValueType(0) == MVT::v2bf16) || + (EltVT == MVT::i16 && N->getValueType(0) == MVT::v2i16)) { + assert(NumElts % 2 == 0 && "Vector must have even number of elements"); + EltVT = N->getValueType(0); + NumElts /= 2; } } @@ -1678,9 +1679,8 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode *N) { MVT ScalarVT = SimpleVT.getScalarType(); unsigned toTypeWidth = ScalarVT.getSizeInBits(); if (SimpleVT.isVector()) { - assert((StoreVT == MVT::v2f16 || StoreVT == MVT::v2bf16) && - "Unexpected vector type"); - // v2f16 is stored using st.b32 + assert(Isv2x16VT(StoreVT) && "Unexpected vector type"); + // v2x16 is stored using st.b32 toTypeWidth = 32; } @@ -1844,10 +1844,10 @@ bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) { return false; } - // v8f16 is a special case. PTX doesn't have st.v8.f16 - // instruction. Instead, we split the vector into v2f16 chunks and + // v8x16 is a special case. PTX doesn't have st.v8.x16 + // instruction. Instead, we split the vector into v2x16 chunks and // store them with st.v4.b32. - if (EltVT == MVT::v2f16 || EltVT == MVT::v2bf16) { + if (Isv2x16VT(EltVT)) { assert(N->getOpcode() == NVPTXISD::StoreV4 && "Unexpected load opcode."); EltVT = MVT::i32; ToType = NVPTX::PTXLdStInstCode::Untyped; diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 5a4c8749b8053..b11c381158a11 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -133,6 +133,7 @@ static bool IsPTXVectorType(MVT VT) { case MVT::v4i8: case MVT::v2i16: case MVT::v4i16: + case MVT::v8i16: // <4 x i16x2> case MVT::v2i32: case MVT::v4i32: case MVT::v2i64: @@ -149,12 +150,9 @@ static bool IsPTXVectorType(MVT VT) { } } -static bool Isv2f16Orv2bf16Type(EVT VT) { - return (VT == MVT::v2f16 || VT == MVT::v2bf16); -} - -static bool Isf16Orbf16Type(MVT VT) { - return (VT.SimpleTy == MVT::f16 || VT.SimpleTy == MVT::bf16); +static bool Is16bitsType(MVT VT) { + return (VT.SimpleTy == MVT::f16 || VT.SimpleTy == MVT::bf16 || + VT.SimpleTy == MVT::i16); } /// ComputePTXValueVTs - For the given Type \p Ty, returns the set of primitive @@ -207,8 +205,20 @@ static void ComputePTXValueVTs(const TargetLowering &TLI, const DataLayout &DL, // Vectors with an even number of f16 elements will be passed to // us as an array of v2f16/v2bf16 elements. We must match this so we // stay in sync with Ins/Outs. - if ((Isf16Orbf16Type(EltVT.getSimpleVT())) && NumElts % 2 == 0) { - EltVT = EltVT == MVT::f16 ? MVT::v2f16 : MVT::v2bf16; + if ((Is16bitsType(EltVT.getSimpleVT())) && NumElts % 2 == 0) { + switch (EltVT.getSimpleVT().SimpleTy) { + case MVT::f16: + EltVT = MVT::v2f16; + break; + case MVT::bf16: + EltVT = MVT::v2bf16; + break; + case MVT::i16: + EltVT = MVT::v2i16; + break; + default: + llvm_unreachable("Unexpected type"); + } NumElts /= 2; } for (unsigned j = 0; j != NumElts; ++j) { @@ -427,8 +437,26 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, Op, VT, IsOpSupported ? Action : NoBF16Action); }; + auto setI16x2OperationAction = [&](unsigned Op, MVT VT, LegalizeAction Action, + LegalizeAction NoI16x2Action) { + bool IsOpSupported = false; + // instructions are available on sm_90 only + switch (Op) { + case ISD::ADD: + case ISD::SMAX: + case ISD::SMIN: + case ISD::UMIN: + case ISD::UMAX: + case ISD::SUB: + IsOpSupported = STI.getSmVersion() >= 90 && STI.getPTXVersion() >= 80; + break; + } + setOperationAction(Op, VT, IsOpSupported ? Action : NoI16x2Action); + }; + addRegisterClass(MVT::i1, &NVPTX::Int1RegsRegClass); addRegisterClass(MVT::i16, &NVPTX::Int16RegsRegClass); + addRegisterClass(MVT::v2i16, &NVPTX::Int32RegsRegClass); addRegisterClass(MVT::i32, &NVPTX::Int32RegsRegClass); addRegisterClass(MVT::i64, &NVPTX::Int64RegsRegClass); addRegisterClass(MVT::f32, &NVPTX::Float32RegsRegClass); @@ -459,9 +487,17 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, setBF16OperationAction(ISD::SETCC, MVT::bf16, Legal, Promote); setBF16OperationAction(ISD::SETCC, MVT::v2bf16, Legal, Expand); + + // Conversion to/from i16/i16x2 is always legal. + setOperationAction(ISD::BUILD_VECTOR, MVT::v2i16, Custom); + setOperationAction(ISD::EXTRACT_VECTOR_ELT, MVT::v2i16, Custom); + setOperationAction(ISD::INSERT_VECTOR_ELT, MVT::v2i16, Expand); + setOperationAction(ISD::VECTOR_SHUFFLE, MVT::v2i16, Expand); + // Operations not directly supported by NVPTX. - for (MVT VT : {MVT::bf16, MVT::f16, MVT::v2bf16, MVT::v2f16, MVT::f32, - MVT::f64, MVT::i1, MVT::i8, MVT::i16, MVT::i32, MVT::i64}) { + for (MVT VT : + {MVT::bf16, MVT::f16, MVT::v2bf16, MVT::v2f16, MVT::f32, MVT::f64, + MVT::i1, MVT::i8, MVT::i16, MVT::v2i16, MVT::i32, MVT::i64}) { setOperationAction(ISD::SELECT_CC, VT, Expand); setOperationAction(ISD::BR_CC, VT, Expand); } @@ -473,6 +509,7 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::i16, Legal); setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::i8 , Legal); setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::i1, Expand); + setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::v2i16, Expand); setOperationAction(ISD::SHL_PARTS, MVT::i32 , Custom); setOperationAction(ISD::SRA_PARTS, MVT::i32 , Custom); @@ -493,10 +530,13 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, setOperationAction(ISD::ROTR, MVT::i32, Legal); setOperationAction(ISD::ROTL, MVT::i16, Expand); + setOperationAction(ISD::ROTL, MVT::v2i16, Expand); setOperationAction(ISD::ROTR, MVT::i16, Expand); + setOperationAction(ISD::ROTR, MVT::v2i16, Expand); setOperationAction(ISD::ROTL, MVT::i8, Expand); setOperationAction(ISD::ROTR, MVT::i8, Expand); setOperationAction(ISD::BSWAP, MVT::i16, Expand); + setOperationAction(ISD::BSWAP, MVT::v2i16, Expand); setOperationAction(ISD::BSWAP, MVT::i32, Expand); setOperationAction(ISD::BSWAP, MVT::i64, Expand); @@ -589,6 +629,21 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, setOperationAction(ISD::CTLZ, Ty, Legal); } + setI16x2OperationAction(ISD::ABS, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::SMIN, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::SMAX, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::UMIN, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::UMAX, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::CTPOP, MVT::v2i16, Legal, Expand); + setI16x2OperationAction(ISD::CTLZ, MVT::v2i16, Legal, Expand); + + setI16x2OperationAction(ISD::ADD, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::SUB, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::MUL, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::SHL, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::SREM, MVT::v2i16, Legal, Custom); + setI16x2OperationAction(ISD::UREM, MVT::v2i16, Legal, Custom); + setOperationAction(ISD::ADDC, MVT::i32, Legal); setOperationAction(ISD::ADDE, MVT::i32, Legal); setOperationAction(ISD::SUBC, MVT::i32, Legal); @@ -601,6 +656,7 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, } setOperationAction(ISD::CTTZ, MVT::i16, Expand); + setOperationAction(ISD::CTTZ, MVT::v2i16, Expand); setOperationAction(ISD::CTTZ, MVT::i32, Expand); setOperationAction(ISD::CTTZ, MVT::i64, Expand); @@ -1323,7 +1379,7 @@ NVPTXTargetLowering::getPreferredVectorAction(MVT VT) const { if (!VT.isScalableVector() && VT.getVectorNumElements() != 1 && VT.getScalarType() == MVT::i1) return TypeSplitVector; - if (Isv2f16Orv2bf16Type(VT)) + if (Isv2x16VT(VT)) return TypeLegal; return TargetLoweringBase::getPreferredVectorAction(VT); } @@ -2103,15 +2159,31 @@ NVPTXTargetLowering::LowerCONCAT_VECTORS(SDValue Op, SelectionDAG &DAG) const { // generates good SASS in both cases. SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, SelectionDAG &DAG) const { - if (!(Isv2f16Orv2bf16Type(Op->getValueType(0)) && - isa(Op->getOperand(0)) && - isa(Op->getOperand(1)))) + EVT VT = Op->getValueType(0); + if (!(Isv2x16VT(VT))) return Op; + APInt E0; + APInt E1; + if (VT == MVT::v2f16 || VT == MVT::v2bf16) { + if (!(isa(Op->getOperand(0)) && + isa(Op->getOperand(1)))) + return Op; + + E0 = cast(Op->getOperand(0)) + ->getValueAPF() + .bitcastToAPInt(); + E1 = cast(Op->getOperand(1)) + ->getValueAPF() + .bitcastToAPInt(); + } else { + assert(VT == MVT::v2i16); + if (!(isa(Op->getOperand(0)) && + isa(Op->getOperand(1)))) + return Op; - APInt E0 = - cast(Op->getOperand(0))->getValueAPF().bitcastToAPInt(); - APInt E1 = - cast(Op->getOperand(1))->getValueAPF().bitcastToAPInt(); + E0 = cast(Op->getOperand(0))->getAPIntValue(); + E1 = cast(Op->getOperand(1))->getAPIntValue(); + } SDValue Const = DAG.getConstant(E1.zext(32).shl(16) | E0.zext(32), SDLoc(Op), MVT::i32); return DAG.getNode(ISD::BITCAST, SDLoc(Op), Op->getValueType(0), Const); @@ -2127,7 +2199,7 @@ SDValue NVPTXTargetLowering::LowerEXTRACT_VECTOR_ELT(SDValue Op, // Extract individual elements and select one of them. SDValue Vector = Op->getOperand(0); EVT VectorVT = Vector.getValueType(); - assert(VectorVT == MVT::v2f16 && "Unexpected vector type."); + assert(Isv2x16VT(VectorVT) && "Unexpected vector type."); EVT EltVT = VectorVT.getVectorElementType(); SDLoc dl(Op.getNode()); @@ -2352,7 +2424,25 @@ SDValue NVPTXTargetLowering::LowerFROUND64(SDValue Op, return DAG.getNode(ISD::SELECT, SL, VT, IsLarge, A, RoundedA); } - +static SDValue LowerVectorArith(SDValue Op, SelectionDAG &DAG) { + SDLoc DL(Op); + if (Op.getValueType() != MVT::v2i16) + return Op; + EVT EltVT = Op.getValueType().getVectorElementType(); + SmallVector VecElements; + for (int I = 0, E = Op.getValueType().getVectorNumElements(); I < E; I++) { + SmallVector ScalarArgs; + llvm::transform(Op->ops(), std::back_inserter(ScalarArgs), + [&](const SDUse &O) { + return DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, EltVT, + O.get(), DAG.getIntPtrConstant(I, DL)); + }); + VecElements.push_back(DAG.getNode(Op.getOpcode(), DL, EltVT, ScalarArgs)); + } + SDValue V = + DAG.getNode(ISD::BUILD_VECTOR, DL, Op.getValueType(), VecElements); + return V; +} SDValue NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const { @@ -2390,6 +2480,18 @@ NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const { return LowerVAARG(Op, DAG); case ISD::VASTART: return LowerVASTART(Op, DAG); + case ISD::ABS: + case ISD::SMIN: + case ISD::SMAX: + case ISD::UMIN: + case ISD::UMAX: + case ISD::ADD: + case ISD::SUB: + case ISD::MUL: + case ISD::SHL: + case ISD::SREM: + case ISD::UREM: + return LowerVectorArith(Op, DAG); default: llvm_unreachable("Custom lowering not defined for operation"); } @@ -2473,9 +2575,9 @@ SDValue NVPTXTargetLowering::LowerLOAD(SDValue Op, SelectionDAG &DAG) const { if (Op.getValueType() == MVT::i1) return LowerLOADi1(Op, DAG); - // v2f16 is legal, so we can't rely on legalizer to handle unaligned - // loads and have to handle it here. - if (Isv2f16Orv2bf16Type(Op.getValueType())) { + // v2f16/v2bf16/v2i16 are legal, so we can't rely on legalizer to handle + // unaligned loads and have to handle it here. + if (Isv2x16VT(Op.getValueType())) { LoadSDNode *Load = cast(Op); EVT MemVT = Load->getMemoryVT(); if (!allowsMemoryAccessForAlignment(*DAG.getContext(), DAG.getDataLayout(), @@ -2520,13 +2622,13 @@ SDValue NVPTXTargetLowering::LowerSTORE(SDValue Op, SelectionDAG &DAG) const { // v2f16 is legal, so we can't rely on legalizer to handle unaligned // stores and have to handle it here. - if (Isv2f16Orv2bf16Type(VT) && + if (Isv2x16VT(VT) && !allowsMemoryAccessForAlignment(*DAG.getContext(), DAG.getDataLayout(), VT, *Store->getMemOperand())) return expandUnalignedStore(Store, DAG); - // v2f16 and v2bf16 don't need special handling. - if (VT == MVT::v2f16 || VT == MVT::v2bf16) + // v2f16, v2bf16 and v2i16 don't need special handling. + if (Isv2x16VT(VT)) return SDValue(); if (VT.isVector()) @@ -2567,6 +2669,7 @@ NVPTXTargetLowering::LowerSTOREVector(SDValue Op, SelectionDAG &DAG) const { case MVT::v4f32: case MVT::v8f16: // <4 x f16x2> case MVT::v8bf16: // <4 x bf16x2> + case MVT::v8i16: // <4 x i16x2> // This is a "native" vector type break; } @@ -2611,8 +2714,7 @@ NVPTXTargetLowering::LowerSTOREVector(SDValue Op, SelectionDAG &DAG) const { // v8f16 is a special case. PTX doesn't have st.v8.f16 // instruction. Instead, we split the vector into v2f16 chunks and // store them with st.v4.b32. - assert(Isf16Orbf16Type(EltVT.getSimpleVT()) && - "Wrong type for the vector."); + assert(Is16bitsType(EltVT.getSimpleVT()) && "Wrong type for the vector."); Opcode = NVPTXISD::StoreV4; StoreF16x2 = true; break; @@ -2798,7 +2900,7 @@ SDValue NVPTXTargetLowering::LowerFormalArguments( EVT LoadVT = EltVT; if (EltVT == MVT::i1) LoadVT = MVT::i8; - else if (Isv2f16Orv2bf16Type(EltVT)) + else if (Isv2x16VT(EltVT)) // getLoad needs a vector type, but it can't handle // vectors which contain v2f16 or v2bf16 elements. So we must load // using i32 here and then bitcast back. @@ -2824,7 +2926,7 @@ SDValue NVPTXTargetLowering::LowerFormalArguments( if (EltVT == MVT::i1) Elt = DAG.getNode(ISD::TRUNCATE, dl, MVT::i1, Elt); // v2f16 was loaded as an i32. Now we must bitcast it back. - else if (Isv2f16Orv2bf16Type(EltVT)) + else if (Isv2x16VT(EltVT)) Elt = DAG.getNode(ISD::BITCAST, dl, EltVT, Elt); // If a promoted integer type is used, truncate down to the original @@ -5202,7 +5304,9 @@ static void ReplaceLoadVector(SDNode *N, SelectionDAG &DAG, case MVT::v4i32: case MVT::v4f16: case MVT::v4f32: - case MVT::v8f16: // <4 x f16x2> + case MVT::v8f16: // <4 x f16x2> + case MVT::v8bf16: // <4 x bf16x2> + case MVT::v8i16: // <4 x i16x2> // This is a "native" vector type break; } @@ -5236,7 +5340,7 @@ static void ReplaceLoadVector(SDNode *N, SelectionDAG &DAG, unsigned Opcode = 0; SDVTList LdResVTs; - bool LoadF16x2 = false; + bool Load16x2 = false; switch (NumElts) { default: @@ -5255,11 +5359,23 @@ static void ReplaceLoadVector(SDNode *N, SelectionDAG &DAG, // v8f16 is a special case. PTX doesn't have ld.v8.f16 // instruction. Instead, we split the vector into v2f16 chunks and // load them with ld.v4.b32. - assert(Isf16Orbf16Type(EltVT.getSimpleVT()) && - "Unsupported v8 vector type."); - LoadF16x2 = true; + assert(Is16bitsType(EltVT.getSimpleVT()) && "Unsupported v8 vector type."); + Load16x2 = true; Opcode = NVPTXISD::LoadV4; - EVT VVT = (EltVT == MVT::f16) ? MVT::v2f16 : MVT::v2bf16; + EVT VVT; + switch (EltVT.getSimpleVT().SimpleTy) { + case MVT::f16: + VVT = MVT::v2f16; + break; + case MVT::bf16: + VVT = MVT::v2bf16; + break; + case MVT::i16: + VVT = MVT::v2i16; + break; + default: + llvm_unreachable("Unsupported v8 vector type."); + } EVT ListVTs[] = {VVT, VVT, VVT, VVT, MVT::Other}; LdResVTs = DAG.getVTList(ListVTs); break; @@ -5278,7 +5394,7 @@ static void ReplaceLoadVector(SDNode *N, SelectionDAG &DAG, LD->getMemOperand()); SmallVector ScalarRes; - if (LoadF16x2) { + if (Load16x2) { // Split v2f16 subvectors back into individual elements. NumElts /= 2; for (unsigned i = 0; i < NumElts; ++i) { diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h index ccd80359bf80b..87cd52c954a7b 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h @@ -617,6 +617,7 @@ class NVPTXTargetLowering : public TargetLowering { Align getArgumentAlignment(SDValue Callee, const CallBase *CB, Type *Ty, unsigned Idx, const DataLayout &DL) const; }; + } // namespace llvm #endif diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td index 4d4dcca2f53e6..d1c75f8bea89b 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -165,6 +165,7 @@ class ValueToRegClass { NVPTXRegClass ret = !cond( !eq(name, "i1"): Int1Regs, !eq(name, "i16"): Int16Regs, + !eq(name, "v2i16"): Int32Regs, !eq(name, "i32"): Int32Regs, !eq(name, "i64"): Int64Regs, !eq(name, "f16"): Int16Regs, @@ -214,6 +215,12 @@ multiclass I3 { [(set Int16Regs:$dst, (OpNode Int16Regs:$a, (imm):$b))]>; } +class I16x2 : + NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$a, Int32Regs:$b), + !strconcat(OpcStr, "16x2 \t$dst, $a, $b;"), + [(set Int32Regs:$dst, (OpNode (v2i16 Int32Regs:$a), (v2i16 Int32Regs:$b)))]>, + Requires<[hasPTX<80>, hasSM<90>]>; + // Template for instructions which take 3 int args. The instructions are // named ".s32" (e.g. "addc.cc.s32"). multiclass ADD_SUB_INT_CARRY { @@ -740,12 +747,10 @@ defm SELP_f64 : SELP_PATTERN<"f64", f64, Float64Regs, f64imm, fpimm>; // def v2f16imm : Operand; // defm SELP_f16x2 : SELP_PATTERN<"b32", v2f16, Int32Regs, v2f16imm, imm>; -def SELP_f16x2rr : - NVPTXInst<(outs Int32Regs:$dst), - (ins Int32Regs:$a, Int32Regs:$b, Int1Regs:$p), - "selp.b32 \t$dst, $a, $b, $p;", - [(set Int32Regs:$dst, - (select Int1Regs:$p, (v2f16 Int32Regs:$a), (v2f16 Int32Regs:$b)))]>; +foreach vt = [v2f16, v2bf16, v2i16] in { +def : Pat<(vt (select Int1Regs:$p, (vt Int32Regs:$a), (vt Int32Regs:$b))), + (SELP_b32rr Int32Regs:$a, Int32Regs:$b, Int1Regs:$p)>; +} //----------------------------------- // Test Instructions @@ -787,6 +792,9 @@ defm SUB_i1 : ADD_SUB_i1; defm ADD : I3<"add.s", add>; defm SUB : I3<"sub.s", sub>; +def ADD16x2 : I16x2<"add.s", add>; +def SUB16x2 : I16x2<"sub.s", sub>; + // in32 and int64 addition and subtraction with carry-out. defm ADDCC : ADD_SUB_INT_CARRY<"add.cc", addc>; defm SUBCC : ADD_SUB_INT_CARRY<"sub.cc", subc>; @@ -826,6 +834,12 @@ defm UMAX : I3<"max.u", umax>; defm SMIN : I3<"min.s", smin>; defm UMIN : I3<"min.u", umin>; +def SMAX16x2 : I16x2<"max.s", smax>; +def UMAX16x2 : I16x2<"max.u", umax>; +def SMIN16x2 : I16x2<"min.s", smin>; +def UMIN16x2 : I16x2<"min.u", umin>; + + // // Wide multiplication // @@ -2633,7 +2647,7 @@ foreach vt = [f16, bf16] in { def: Pat<(vt (ProxyReg vt:$src)), (ProxyRegI16 Int16Regs:$src)>; } -foreach vt = [v2f16, v2bf16] in { +foreach vt = [v2f16, v2bf16, v2i16] in { def: Pat<(vt (ProxyReg vt:$src)), (ProxyRegI32 Int32Regs:$src)>; } @@ -2929,15 +2943,13 @@ def BITCONVERT_32_F2I : F_BITCONVERT<"32", f32, i32>; def BITCONVERT_64_I2F : F_BITCONVERT<"64", i64, f64>; def BITCONVERT_64_F2I : F_BITCONVERT<"64", f64, i64>; -foreach vt = [v2f16, v2bf16] in { +foreach vt = [v2f16, v2bf16, v2i16] in { def: Pat<(vt (bitconvert (i32 UInt32Const:$a))), (IMOVB32ri UInt32Const:$a)>; -def: Pat<(vt (bitconvert (i32 Int32Regs:$a))), - (ProxyRegI32 Int32Regs:$a)>; -def: Pat<(i32 (bitconvert (vt Int32Regs:$a))), - (ProxyRegI32 Int32Regs:$a)>; def: Pat<(vt (bitconvert (f32 Float32Regs:$a))), (BITCONVERT_32_F2I Float32Regs:$a)>; +def: Pat<(f32 (bitconvert (vt Int32Regs:$a))), + (BITCONVERT_32_I2F Int32Regs:$a)>; } foreach vt = [f16, bf16] in { def: Pat<(vt (bitconvert (i16 UInt16Const:$a))), @@ -2948,6 +2960,15 @@ def: Pat<(i16 (bitconvert (vt Int16Regs:$a))), (ProxyRegI16 Int16Regs:$a)>; } +foreach ta = [v2f16, v2bf16, v2i16, i32] in { + foreach tb = [v2f16, v2bf16, v2i16, i32] in { + if !ne(ta, tb) then { + def: Pat<(ta (bitconvert (tb Int32Regs:$a))), + (ProxyRegI32 Int32Regs:$a)>; + } + } +} + // NOTE: pred->fp are currently sub-optimal due to an issue in TableGen where // we cannot specify floating-point literals in isel patterns. Therefore, we // use an integer selp to select either 1 or 0 and then cvt to floating-point. @@ -3286,19 +3307,18 @@ def : Pat<(i32 (trunc (srl Int64Regs:$s, (i32 32)))), def : Pat<(i32 (trunc (sra Int64Regs:$s, (i32 32)))), (I64toI32H Int64Regs:$s)>; -def : Pat<(f16 (extractelt (v2f16 Int32Regs:$src), 0)), +foreach vt = [v2f16, v2bf16, v2i16] in { +def : Pat<(extractelt (vt Int32Regs:$src), 0), (I32toI16L Int32Regs:$src)>; -def : Pat<(f16 (extractelt (v2f16 Int32Regs:$src), 1)), +def : Pat<(extractelt (vt Int32Regs:$src), 1), (I32toI16H Int32Regs:$src)>; +} def : Pat<(v2f16 (build_vector (f16 Int16Regs:$a), (f16 Int16Regs:$b))), (V2I16toI32 Int16Regs:$a, Int16Regs:$b)>; - -def : Pat<(bf16 (extractelt (v2bf16 Int32Regs:$src), 0)), - (I32toI16L Int32Regs:$src)>; -def : Pat<(bf16 (extractelt (v2bf16 Int32Regs:$src), 1)), - (I32toI16H Int32Regs:$src)>; def : Pat<(v2bf16 (build_vector (bf16 Int16Regs:$a), (bf16 Int16Regs:$b))), (V2I16toI32 Int16Regs:$a, Int16Regs:$b)>; +def : Pat<(v2i16 (build_vector (i16 Int16Regs:$a), (i16 Int16Regs:$b))), + (V2I16toI32 Int16Regs:$a, Int16Regs:$b)>; // Count leading zeros let hasSideEffects = false in { diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td index b62460e8cd31f..ed9dabf39dd7a 100644 --- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td @@ -58,7 +58,7 @@ foreach i = 0...31 in { //===----------------------------------------------------------------------===// def Int1Regs : NVPTXRegClass<[i1], 8, (add (sequence "P%u", 0, 4))>; def Int16Regs : NVPTXRegClass<[i16, f16, bf16], 16, (add (sequence "RS%u", 0, 4))>; -def Int32Regs : NVPTXRegClass<[i32, v2f16, v2bf16], 32, +def Int32Regs : NVPTXRegClass<[i32, v2f16, v2bf16, v2i16], 32, (add (sequence "R%u", 0, 4), VRFrame32, VRFrameLocal32)>; def Int64Regs : NVPTXRegClass<[i64], 64, (add (sequence "RL%u", 0, 4), VRFrame64, VRFrameLocal64)>; diff --git a/llvm/lib/Target/NVPTX/NVPTXUtilities.cpp b/llvm/lib/Target/NVPTX/NVPTXUtilities.cpp index 988910810da65..c3737f9fcca82 100644 --- a/llvm/lib/Target/NVPTX/NVPTXUtilities.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXUtilities.cpp @@ -348,4 +348,8 @@ bool shouldEmitPTXNoReturn(const Value *V, const TargetMachine &TM) { !isKernelFunction(*F); } +bool Isv2x16VT(EVT VT) { + return (VT == MVT::v2f16 || VT == MVT::v2bf16 || VT == MVT::v2i16); +} + } // namespace llvm diff --git a/llvm/lib/Target/NVPTX/NVPTXUtilities.h b/llvm/lib/Target/NVPTX/NVPTXUtilities.h index f980ea3dec0b8..521f8198911f2 100644 --- a/llvm/lib/Target/NVPTX/NVPTXUtilities.h +++ b/llvm/lib/Target/NVPTX/NVPTXUtilities.h @@ -13,6 +13,7 @@ #ifndef LLVM_LIB_TARGET_NVPTX_NVPTXUTILITIES_H #define LLVM_LIB_TARGET_NVPTX_NVPTXUTILITIES_H +#include "llvm/CodeGen/ValueTypes.h" #include "llvm/IR/Function.h" #include "llvm/IR/GlobalVariable.h" #include "llvm/IR/IntrinsicInst.h" @@ -74,6 +75,8 @@ inline unsigned promoteScalarArgumentSize(unsigned size) { } bool shouldEmitPTXNoReturn(const Value *V, const TargetMachine &TM); + +bool Isv2x16VT(EVT VT); } #endif diff --git a/llvm/test/CodeGen/NVPTX/dag-cse.ll b/llvm/test/CodeGen/NVPTX/dag-cse.ll index 0b21cdebd87cd..bbfedf42ad548 100644 --- a/llvm/test/CodeGen/NVPTX/dag-cse.ll +++ b/llvm/test/CodeGen/NVPTX/dag-cse.ll @@ -12,8 +12,8 @@ ; CHECK: ld.global.v2.u8 {%[[B1:rs[0-9]+]], %[[B2:rs[0-9]+]]}, [a]; ; CHECK: st.global.v2.u8 [b], {%[[B1]], %[[B2]]}; ; -; CHECK: ld.global.v2.u16 {%[[C1:rs[0-9]+]], %[[C2:rs[0-9]+]]}, [a]; -; CHECK: st.global.v2.u16 [c], {%[[C1]], %[[C2]]}; +; CHECK: ld.global.u32 %[[C:r[0-9]+]], [a]; +; CHECK: st.global.u32 [c], %[[C]]; define void @test1() #0 { %1 = load <2 x i8>, ptr addrspace(1) @a, align 8 diff --git a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll index dba6691e30aa2..4eb538628a8de 100644 --- a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll @@ -993,9 +993,7 @@ define <2 x double> @test_fpext_2xdouble(<2 x half> %a) #0 { ; CHECK-LABEL: test_bitcast_2xhalf_to_2xi16( ; CHECK: ld.param.u32 [[A:%r[0-9]+]], [test_bitcast_2xhalf_to_2xi16_param_0]; -; CHECK-DAG: cvt.u16.u32 [[R0:%rs[0-9]+]], [[A]] -; CHECK-DAG: mov.b32 {tmp, [[R1:%rs[0-9]+]]}, [[A]]; -; CHECK: st.param.v2.b16 [func_retval0+0], {[[R0]], [[R1]]} +; CHECK: st.param.b32 [func_retval0+0], [[A]] ; CHECK: ret; define <2 x i16> @test_bitcast_2xhalf_to_2xi16(<2 x half> %a) #0 { %r = bitcast <2 x half> %a to <2 x i16> @@ -1003,11 +1001,7 @@ define <2 x i16> @test_bitcast_2xhalf_to_2xi16(<2 x half> %a) #0 { } ; CHECK-LABEL: test_bitcast_2xi16_to_2xhalf( -; CHECK: ld.param.v2.u16 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [test_bitcast_2xi16_to_2xhalf_param_0]; -; CHECK-DAG: cvt.u32.u16 [[R0:%r[0-9]+]], [[RS0]]; -; CHECK-DAG: cvt.u32.u16 [[R1:%r[0-9]+]], [[RS1]]; -; CHECK-DAG: shl.b32 [[R1H:%r[0-9]+]], [[R1]], 16; -; CHECK-DAG: or.b32 [[R:%r[0-9]+]], [[R0]], [[R1H]]; +; CHECK: ld.param.u32 [[R:%r[0-9]+]], [test_bitcast_2xi16_to_2xhalf_param_0]; ; CHECK: st.param.b32 [func_retval0+0], [[R]]; ; CHECK: ret; define <2 x half> @test_bitcast_2xi16_to_2xhalf(<2 x i16> %a) #0 { diff --git a/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll new file mode 100644 index 0000000000000..6fa16bf393f95 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll @@ -0,0 +1,537 @@ +; ## Support i16x2 instructions +; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx80 -asm-verbose=false \ +; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ +; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,I16x2 %s +; RUN: %if ptxas %{ \ +; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -asm-verbose=false \ +; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ +; RUN: | %ptxas-verify -arch=sm_53 \ +; RUN: %} +; ## No support for i16x2 instructions +; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 -asm-verbose=false \ +; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ +; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,NO-I16x2 %s +; RUN: %if ptxas %{ \ +; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 -asm-verbose=false \ +; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ +; RUN: | %ptxas-verify -arch=sm_53 \ +; RUN: %} + +target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128" + +; COMMON-LABEL: test_ret_const( +; COMMON: mov.b32 [[R:%r[0-9+]]], 131073; +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_ret_const() #0 { + ret <2 x i16> +} + +; COMMON-LABEL: test_extract_0( +; COMMON: ld.param.u32 [[A:%r[0-9]+]], [test_extract_0_param_0]; +; COMMON: mov.b32 {[[RS:%rs[0-9]+]], tmp}, [[A]]; +; COMMON: cvt.u32.u16 [[R:%r[0-9]+]], [[RS]]; +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define i16 @test_extract_0(<2 x i16> %a) #0 { + %e = extractelement <2 x i16> %a, i32 0 + ret i16 %e +} + +; COMMON-LABEL: test_extract_1( +; COMMON: ld.param.u32 [[A:%r[0-9]+]], [test_extract_1_param_0]; +; COMMON: mov.b32 {tmp, [[RS:%rs[0-9]+]]}, [[A]]; +; COMMON: cvt.u32.u16 [[R:%r[0-9]+]], [[RS]]; +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define i16 @test_extract_1(<2 x i16> %a) #0 { + %e = extractelement <2 x i16> %a, i32 1 + ret i16 %e +} + +; COMMON-LABEL: test_extract_i( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_extract_i_param_0]; +; COMMON-DAG: ld.param.u64 [[IDX:%rd[0-9]+]], [test_extract_i_param_1]; +; COMMON-DAG: setp.eq.s64 [[PRED:%p[0-9]+]], [[IDX]], 0; +; COMMON-DAG: mov.b32 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]]}, [[A]]; +; COMMON: selp.b16 [[RS:%rs[0-9]+]], [[E0]], [[E1]], [[PRED]]; +; COMMON: cvt.u32.u16 [[R:%r[0-9]+]], [[RS]]; +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define i16 @test_extract_i(<2 x i16> %a, i64 %idx) #0 { + %e = extractelement <2 x i16> %a, i64 %idx + ret i16 %e +} + +; COMMON-LABEL: test_add( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_add_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_add_param_1]; +; +; I16x2-NEXT: add.s16x2 [[R:%r[0-9]+]], [[A]], [[B]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: mov.b32 {[[RS2:%rs[0-9]+]], [[RS3:%rs[0-9]+]]}, [[B]]; +; NO-I16x2-DAG: add.s16 [[RS4:%rs[0-9]+]], [[RS0]], [[RS2]]; +; NO-I16x2-DAG: add.s16 [[RS5:%rs[0-9]+]], [[RS1]], [[RS3]]; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS4]], [[RS5]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_add(<2 x i16> %a, <2 x i16> %b) #0 { + %r = add <2 x i16> %a, %b + ret <2 x i16> %r +} + +; Check that we can lower add with immediate arguments. +; COMMON-LABEL: test_add_imm_0( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_add_imm_0_param_0]; +; +; I16x2: mov.b32 [[I:%r[0-9+]]], 131073; +; I16x2: add.s16x2 [[R:%r[0-9]+]], [[A]], [[I]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: add.s16 [[RS2:%rs[0-9]+]], [[RS0]], 1; +; NO-I16x2-DAG: add.s16 [[RS3:%rs[0-9]+]], [[RS1]], 2; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS2]], [[RS3]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_add_imm_0(<2 x i16> %a) #0 { + %r = add <2 x i16> , %a + ret <2 x i16> %r +} + +; COMMON-LABEL: test_add_imm_1( +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_add_imm_1_param_0]; +; +; I16x2: mov.b32 [[I:%r[0-9+]]], 131073; +; I16x2: add.s16x2 [[R:%r[0-9]+]], [[A]], [[I]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: add.s16 [[RS2:%rs[0-9]+]], [[RS0]], 1; +; NO-I16x2-DAG: add.s16 [[RS3:%rs[0-9]+]], [[RS1]], 2; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS2]], [[RS3]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_add_imm_1(<2 x i16> %a) #0 { + %r = add <2 x i16> %a, + ret <2 x i16> %r +} + +; COMMON-LABEL: test_sub( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_sub_param_0]; +; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_sub_param_1]; +; I16x2: sub.s16x2 [[R:%r[0-9]+]], [[A]], [[B]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: mov.b32 {[[RS2:%rs[0-9]+]], [[RS3:%rs[0-9]+]]}, [[B]]; +; NO-I16x2-DAG: sub.s16 [[RS4:%rs[0-9]+]], [[RS0]], [[RS2]]; +; NO-I16x2-DAG: sub.s16 [[RS5:%rs[0-9]+]], [[RS1]], [[RS3]]; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS4]], [[RS5]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_sub(<2 x i16> %a, <2 x i16> %b) #0 { + %r = sub <2 x i16> %a, %b + ret <2 x i16> %r +} + +; COMMON-LABEL: test_smax( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_smax_param_0]; +; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_smax_param_1]; +; I16x2: max.s16x2 [[R:%r[0-9]+]], [[A]], [[B]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: mov.b32 {[[RS2:%rs[0-9]+]], [[RS3:%rs[0-9]+]]}, [[B]]; +; NO-I16x2-DAG: max.s16 [[RS4:%rs[0-9]+]], [[RS0]], [[RS2]]; +; NO-I16x2-DAG: max.s16 [[RS5:%rs[0-9]+]], [[RS1]], [[RS3]]; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS4]], [[RS5]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_smax(<2 x i16> %a, <2 x i16> %b) #0 { + %cmp = icmp sgt <2 x i16> %a, %b + %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b + ret <2 x i16> %r +} + +; COMMON-LABEL: test_umax( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_umax_param_0]; +; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_umax_param_1]; +; I16x2: max.u16x2 [[R:%r[0-9]+]], [[A]], [[B]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: mov.b32 {[[RS2:%rs[0-9]+]], [[RS3:%rs[0-9]+]]}, [[B]]; +; NO-I16x2-DAG: max.u16 [[RS4:%rs[0-9]+]], [[RS0]], [[RS2]]; +; NO-I16x2-DAG: max.u16 [[RS5:%rs[0-9]+]], [[RS1]], [[RS3]]; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS4]], [[RS5]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_umax(<2 x i16> %a, <2 x i16> %b) #0 { + %cmp = icmp ugt <2 x i16> %a, %b + %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b + ret <2 x i16> %r +} + +; COMMON-LABEL: test_smin( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_smin_param_0]; +; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_smin_param_1]; +; I16x2: min.s16x2 [[R:%r[0-9]+]], [[A]], [[B]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: mov.b32 {[[RS2:%rs[0-9]+]], [[RS3:%rs[0-9]+]]}, [[B]]; +; NO-I16x2-DAG: min.s16 [[RS4:%rs[0-9]+]], [[RS0]], [[RS2]]; +; NO-I16x2-DAG: min.s16 [[RS5:%rs[0-9]+]], [[RS1]], [[RS3]]; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS4]], [[RS5]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_smin(<2 x i16> %a, <2 x i16> %b) #0 { + %cmp = icmp sle <2 x i16> %a, %b + %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b + ret <2 x i16> %r +} + +; COMMON-LABEL: test_umin( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_umin_param_0]; +; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_umin_param_1]; +; I16x2: min.u16x2 [[R:%r[0-9]+]], [[A]], [[B]]; +; +; NO-I16x2-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; NO-I16x2-DAG: mov.b32 {[[RS2:%rs[0-9]+]], [[RS3:%rs[0-9]+]]}, [[B]]; +; NO-I16x2-DAG: min.u16 [[RS4:%rs[0-9]+]], [[RS0]], [[RS2]]; +; NO-I16x2-DAG: min.u16 [[RS5:%rs[0-9]+]], [[RS1]], [[RS3]]; +; NO-I16x2-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS4]], [[RS5]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_umin(<2 x i16> %a, <2 x i16> %b) #0 { + %cmp = icmp ule <2 x i16> %a, %b + %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b + ret <2 x i16> %r +} + +; COMMON-LABEL: test_mul( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_mul_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_mul_param_1]; +; +; COMMON-DAG: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[A]]; +; COMMON-DAG: mov.b32 {[[RS2:%rs[0-9]+]], [[RS3:%rs[0-9]+]]}, [[B]]; +; COMMON-DAG: mul.lo.s16 [[RS4:%rs[0-9]+]], [[RS0]], [[RS2]]; +; COMMON-DAG: mul.lo.s16 [[RS5:%rs[0-9]+]], [[RS1]], [[RS3]]; +; COMMON-DAG: mov.b32 [[R:%r[0-9]+]], {[[RS4]], [[RS5]]}; +; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_mul(<2 x i16> %a, <2 x i16> %b) #0 { + %r = mul <2 x i16> %a, %b + ret <2 x i16> %r +} + + +; COMMON-LABEL: .func test_ldst_v2i16( +; COMMON-DAG: ld.param.u64 [[A:%rd[0-9]+]], [test_ldst_v2i16_param_0]; +; COMMON-DAG: ld.param.u64 [[B:%rd[0-9]+]], [test_ldst_v2i16_param_1]; +; COMMON-DAG: ld.u32 [[E:%r[0-9]+]], [[[A]]]; +; COMMON-DAG: st.u32 [[[B]]], [[E]]; +; COMMON: ret; +define void @test_ldst_v2i16(ptr %a, ptr %b) { + %t1 = load <2 x i16>, ptr %a + store <2 x i16> %t1, ptr %b, align 16 + ret void +} + +; COMMON-LABEL: .func test_ldst_v3i16( +; COMMON-DAG: ld.param.u64 %[[A:rd[0-9]+]], [test_ldst_v3i16_param_0]; +; COMMON-DAG: ld.param.u64 %[[B:rd[0-9]+]], [test_ldst_v3i16_param_1]; +; -- v3 is inconvenient to capture as it's lowered as ld.b64 + fair +; number of bitshifting instructions that may change at llvm's whim. +; So we only verify that we only issue correct number of writes using +; correct offset, but not the values we write. +; COMMON-DAG: ld.u64 +; COMMON-DAG: st.u32 [%[[B]]], +; COMMON-DAG: st.u16 [%[[B]]+4], +; COMMON: ret; +define void @test_ldst_v3i16(ptr %a, ptr %b) { + %t1 = load <3 x i16>, ptr %a + store <3 x i16> %t1, ptr %b, align 16 + ret void +} + +; COMMON-LABEL: .func test_ldst_v4i16( +; COMMON-DAG: ld.param.u64 %[[A:rd[0-9]+]], [test_ldst_v4i16_param_0]; +; COMMON-DAG: ld.param.u64 %[[B:rd[0-9]+]], [test_ldst_v4i16_param_1]; +; COMMON-DAG: ld.v4.u16 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]], [[E2:%rs[0-9]+]], [[E3:%rs[0-9]+]]}, [%[[A]]]; +; COMMON-DAG: st.v4.u16 [%[[B]]], {[[E0]], [[E1]], [[E2]], [[E3]]}; +; COMMON: ret; +define void @test_ldst_v4i16(ptr %a, ptr %b) { + %t1 = load <4 x i16>, ptr %a + store <4 x i16> %t1, ptr %b, align 16 + ret void +} + +; COMMON-LABEL: .func test_ldst_v8i16( +; COMMON-DAG: ld.param.u64 %[[A:rd[0-9]+]], [test_ldst_v8i16_param_0]; +; COMMON-DAG: ld.param.u64 %[[B:rd[0-9]+]], [test_ldst_v8i16_param_1]; +; COMMON-DAG: ld.v4.b32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]], [[E2:%r[0-9]+]], [[E3:%r[0-9]+]]}, [%[[A]]]; +; COMMON-DAG: st.v4.b32 [%[[B]]], {[[E0]], [[E1]], [[E2]], [[E3]]}; +; COMMON: ret; +define void @test_ldst_v8i16(ptr %a, ptr %b) { + %t1 = load <8 x i16>, ptr %a + store <8 x i16> %t1, ptr %b, align 16 + ret void +} + +declare <2 x i16> @test_callee(<2 x i16> %a, <2 x i16> %b) #0 + +; COMMON-LABEL: test_call( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_call_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_call_param_1]; +; COMMON: { +; COMMON-DAG: .param .align 4 .b8 param0[4]; +; COMMON-DAG: .param .align 4 .b8 param1[4]; +; COMMON-DAG: st.param.b32 [param0+0], [[A]]; +; COMMON-DAG: st.param.b32 [param1+0], [[B]]; +; COMMON-DAG: .param .align 4 .b8 retval0[4]; +; COMMON: call.uni (retval0), +; COMMON-NEXT: test_callee, +; COMMON: ); +; COMMON-NEXT: ld.param.b32 [[R:%r[0-9]+]], [retval0+0]; +; COMMON-NEXT: } +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_call(<2 x i16> %a, <2 x i16> %b) #0 { + %r = call <2 x i16> @test_callee(<2 x i16> %a, <2 x i16> %b) + ret <2 x i16> %r +} + +; COMMON-LABEL: test_call_flipped( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_call_flipped_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_call_flipped_param_1]; +; COMMON: { +; COMMON-DAG: .param .align 4 .b8 param0[4]; +; COMMON-DAG: .param .align 4 .b8 param1[4]; +; COMMON-DAG: st.param.b32 [param0+0], [[B]]; +; COMMON-DAG: st.param.b32 [param1+0], [[A]]; +; COMMON-DAG: .param .align 4 .b8 retval0[4]; +; COMMON: call.uni (retval0), +; COMMON-NEXT: test_callee, +; COMMON: ); +; COMMON-NEXT: ld.param.b32 [[R:%r[0-9]+]], [retval0+0]; +; COMMON-NEXT: } +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_call_flipped(<2 x i16> %a, <2 x i16> %b) #0 { + %r = call <2 x i16> @test_callee(<2 x i16> %b, <2 x i16> %a) + ret <2 x i16> %r +} + +; COMMON-LABEL: test_tailcall_flipped( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_tailcall_flipped_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_tailcall_flipped_param_1]; +; COMMON: { +; COMMON-DAG: .param .align 4 .b8 param0[4]; +; COMMON-DAG: .param .align 4 .b8 param1[4]; +; COMMON-DAG: st.param.b32 [param0+0], [[B]]; +; COMMON-DAG: st.param.b32 [param1+0], [[A]]; +; COMMON-DAG: .param .align 4 .b8 retval0[4]; +; COMMON: call.uni (retval0), +; COMMON-NEXT: test_callee, +; COMMON: ); +; COMMON-NEXT: ld.param.b32 [[R:%r[0-9]+]], [retval0+0]; +; COMMON-NEXT: } +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_tailcall_flipped(<2 x i16> %a, <2 x i16> %b) #0 { + %r = tail call <2 x i16> @test_callee(<2 x i16> %b, <2 x i16> %a) + ret <2 x i16> %r +} + +; COMMON-LABEL: test_select( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_select_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_select_param_1]; +; COMMON-DAG: ld.param.u8 [[C:%rs[0-9]+]], [test_select_param_2] +; COMMON-DAG: setp.eq.b16 [[PRED:%p[0-9]+]], %rs{{.*}}, 1; +; COMMON-NEXT: selp.b32 [[R:%r[0-9]+]], [[A]], [[B]], [[PRED]]; +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_select(<2 x i16> %a, <2 x i16> %b, i1 zeroext %c) #0 { + %r = select i1 %c, <2 x i16> %a, <2 x i16> %b + ret <2 x i16> %r +} + +; COMMON-LABEL: test_select_cc( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_select_cc_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_select_cc_param_1]; +; COMMON-DAG: ld.param.u32 [[C:%r[0-9]+]], [test_select_cc_param_2]; +; COMMON-DAG: ld.param.u32 [[D:%r[0-9]+]], [test_select_cc_param_3]; +; COMMON-DAG: mov.b32 {[[C0:%rs[0-9]+]], [[C1:%rs[0-9]+]]}, [[C]] +; COMMON-DAG: mov.b32 {[[D0:%rs[0-9]+]], [[D1:%rs[0-9]+]]}, [[D]] +; COMMON-DAG: setp.ne.s16 [[P0:%p[0-9]+]], [[C0]], [[D0]] +; COMMON-DAG: setp.ne.s16 [[P1:%p[0-9]+]], [[C1]], [[D1]] +; COMMON-DAG: mov.b32 {[[A0:%rs[0-9]+]], [[A1:%rs[0-9]+]]}, [[A]] +; COMMON-DAG: mov.b32 {[[B0:%rs[0-9]+]], [[B1:%rs[0-9]+]]}, [[B]] +; COMMON-DAG: selp.b16 [[R0:%rs[0-9]+]], [[A0]], [[B0]], [[P0]]; +; COMMON-DAG: selp.b16 [[R1:%rs[0-9]+]], [[A1]], [[B1]], [[P1]]; +; COMMON: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]} +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_select_cc(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> %d) #0 { + %cc = icmp ne <2 x i16> %c, %d + %r = select <2 x i1> %cc, <2 x i16> %a, <2 x i16> %b + ret <2 x i16> %r +} + +; COMMON-LABEL: test_select_cc_i32_i16( +; COMMON-DAG: ld.param.v2.u32 {[[A0:%r[0-9]+]], [[A1:%r[0-9]+]]}, [test_select_cc_i32_i16_param_0]; +; COMMON-DAG: ld.param.v2.u32 {[[B0:%r[0-9]+]], [[B1:%r[0-9]+]]}, [test_select_cc_i32_i16_param_1]; +; COMMON-DAG: ld.param.u32 [[C:%r[0-9]+]], [test_select_cc_i32_i16_param_2]; +; COMMON-DAG: ld.param.u32 [[D:%r[0-9]+]], [test_select_cc_i32_i16_param_3]; +; COMMON-DAG: mov.b32 {[[C0:%rs[0-9]+]], [[C1:%rs[0-9]+]]}, [[C]] +; COMMON-DAG: mov.b32 {[[D0:%rs[0-9]+]], [[D1:%rs[0-9]+]]}, [[D]] +; COMMON-DAG: setp.ne.s16 [[P0:%p[0-9]+]], [[C0]], [[D0]] +; COMMON-DAG: setp.ne.s16 [[P1:%p[0-9]+]], [[C1]], [[D1]] +; COMMON-DAG: selp.b32 [[R0:%r[0-9]+]], [[A0]], [[B0]], [[P0]]; +; COMMON-DAG: selp.b32 [[R1:%r[0-9]+]], [[A1]], [[B1]], [[P1]]; +; COMMON-NEXT: st.param.v2.b32 [func_retval0+0], {[[R0]], [[R1]]}; +; COMMON-NEXT: ret; +define <2 x i32> @test_select_cc_i32_i16(<2 x i32> %a, <2 x i32> %b, + <2 x i16> %c, <2 x i16> %d) #0 { + %cc = icmp ne <2 x i16> %c, %d + %r = select <2 x i1> %cc, <2 x i32> %a, <2 x i32> %b + ret <2 x i32> %r +} + +; COMMON-LABEL: test_select_cc_i16_i32( +; COMMON-DAG: ld.param.u32 [[A:%r[0-9]+]], [test_select_cc_i16_i32_param_0]; +; COMMON-DAG: ld.param.u32 [[B:%r[0-9]+]], [test_select_cc_i16_i32_param_1]; +; COMMON-DAG: ld.param.v2.u32 {[[C0:%r[0-9]+]], [[C1:%r[0-9]+]]}, [test_select_cc_i16_i32_param_2]; +; COMMON-DAG: ld.param.v2.u32 {[[D0:%r[0-9]+]], [[D1:%r[0-9]+]]}, [test_select_cc_i16_i32_param_3]; +; COMMON-DAG: setp.ne.s32 [[P0:%p[0-9]+]], [[C0]], [[D0]] +; COMMON-DAG: setp.ne.s32 [[P1:%p[0-9]+]], [[C1]], [[D1]] +; COMMON-DAG: mov.b32 {[[A0:%rs[0-9]+]], [[A1:%rs[0-9]+]]}, [[A]] +; COMMON-DAG: mov.b32 {[[B0:%rs[0-9]+]], [[B1:%rs[0-9]+]]}, [[B]] +; COMMON-DAG: selp.b16 [[R0:%rs[0-9]+]], [[A0]], [[B0]], [[P0]]; +; COMMON-DAG: selp.b16 [[R1:%rs[0-9]+]], [[A1]], [[B1]], [[P1]]; +; COMMON: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]} +; COMMON-NEXT: st.param.b32 [func_retval0+0], [[R]]; +; COMMON-NEXT: ret; +define <2 x i16> @test_select_cc_i16_i32(<2 x i16> %a, <2 x i16> %b, + <2 x i32> %c, <2 x i32> %d) #0 { + %cc = icmp ne <2 x i32> %c, %d + %r = select <2 x i1> %cc, <2 x i16> %a, <2 x i16> %b + ret <2 x i16> %r +} + + +; COMMON-LABEL: test_trunc_2xi32( +; COMMON: ld.param.v2.u32 {[[A0:%r[0-9]+]], [[A1:%r[0-9]+]]}, [test_trunc_2xi32_param_0]; +; COMMON-DAG: cvt.u16.u32 [[R0:%rs[0-9]+]], [[A0]]; +; COMMON-DAG: cvt.u16.u32 [[R1:%rs[0-9]+]], [[A1]]; +; COMMON: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]} +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define <2 x i16> @test_trunc_2xi32(<2 x i32> %a) #0 { + %r = trunc <2 x i32> %a to <2 x i16> + ret <2 x i16> %r +} + +; COMMON-LABEL: test_trunc_2xi64( +; COMMON: ld.param.v2.u64 {[[A0:%rd[0-9]+]], [[A1:%rd[0-9]+]]}, [test_trunc_2xi64_param_0]; +; COMMON-DAG: cvt.u16.u64 [[R0:%rs[0-9]+]], [[A0]]; +; COMMON-DAG: cvt.u16.u64 [[R1:%rs[0-9]+]], [[A1]]; +; COMMON: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]} +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define <2 x i16> @test_trunc_2xi64(<2 x i64> %a) #0 { + %r = trunc <2 x i64> %a to <2 x i16> + ret <2 x i16> %r +} + +; COMMON-LABEL: test_zext_2xi32( +; COMMON: ld.param.u32 [[A:%r[0-9]+]], [test_zext_2xi32_param_0]; +; COMMON: mov.b32 {[[A0:%rs[0-9]+]], [[A1:%rs[0-9]+]]}, [[A]] +; COMMON-DAG: cvt.u32.u16 [[R0:%r[0-9]+]], [[A0]]; +; COMMON-DAG: cvt.u32.u16 [[R1:%r[0-9]+]], [[A1]]; +; COMMON-NEXT: st.param.v2.b32 [func_retval0+0], {[[R0]], [[R1]]}; +; COMMON: ret; +define <2 x i32> @test_zext_2xi32(<2 x i16> %a) #0 { + %r = zext <2 x i16> %a to <2 x i32> + ret <2 x i32> %r +} + +; COMMON-LABEL: test_zext_2xi64( +; COMMON: ld.param.u32 [[A:%r[0-9]+]], [test_zext_2xi64_param_0]; +; COMMON: mov.b32 {[[A0:%rs[0-9]+]], [[A1:%rs[0-9]+]]}, [[A]] +; COMMON-DAG: cvt.u64.u16 [[R0:%rd[0-9]+]], [[A0]]; +; COMMON-DAG: cvt.u64.u16 [[R1:%rd[0-9]+]], [[A1]]; +; COMMON-NEXT: st.param.v2.b64 [func_retval0+0], {[[R0]], [[R1]]}; +; COMMON: ret; +define <2 x i64> @test_zext_2xi64(<2 x i16> %a) #0 { + %r = zext <2 x i16> %a to <2 x i64> + ret <2 x i64> %r +} + +; COMMON-LABEL: test_bitcast_i32_to_2xi16( +; COMMON: ld.param.u32 [[R:%r[0-9]+]], [test_bitcast_i32_to_2xi16_param_0]; +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define <2 x i16> @test_bitcast_i32_to_2xi16(i32 %a) #0 { + %r = bitcast i32 %a to <2 x i16> + ret <2 x i16> %r +} + +; COMMON-LABEL: test_bitcast_2xi16_to_i32( +; COMMON: ld.param.u32 [[R:%r[0-9]+]], [test_bitcast_2xi16_to_i32_param_0]; +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define i32 @test_bitcast_2xi16_to_i32(<2 x i16> %a) #0 { + %r = bitcast <2 x i16> %a to i32 + ret i32 %r +} + +; COMMON-LABEL: test_bitcast_2xi16_to_2xhalf( +; COMMON: ld.param.u16 [[RS1:%rs[0-9]+]], [test_bitcast_2xi16_to_2xhalf_param_0]; +; COMMON: mov.u16 [[RS2:%rs[0-9]+]], 5; +; COMMON: mov.b32 [[R:%r[0-9]+]], {[[RS1]], [[RS2]]}; +; COMMON: st.param.b32 [func_retval0+0], [[R]]; +; COMMON: ret; +define <2 x half> @test_bitcast_2xi16_to_2xhalf(i16 %a) #0 { + %ins.0 = insertelement <2 x i16> undef, i16 %a, i32 0 + %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1 + %r = bitcast <2 x i16> %ins.1 to <2 x half> + ret <2 x half> %r +} + + +; COMMON-LABEL: test_shufflevector( +; COMMON: ld.param.u32 [[R:%r[0-9]+]], [test_shufflevector_param_0]; +; COMMON: mov.b32 {[[RS0:%rs[0-9]+]], [[RS1:%rs[0-9]+]]}, [[R]]; +; COMMON: mov.b32 [[R1:%r[0-9]+]], {[[RS1]], [[RS0]]}; +; COMMON: st.param.b32 [func_retval0+0], [[R1]]; +; COMMON: ret; +define <2 x i16> @test_shufflevector(<2 x i16> %a) #0 { + %s = shufflevector <2 x i16> %a, <2 x i16> undef, <2 x i32> + ret <2 x i16> %s +} + +; COMMON-LABEL: test_insertelement( +; COMMON: ld.param.u16 [[B:%rs[0-9]+]], [test_insertelement_param_1]; +; COMMON: ld.param.u32 [[A:%r[0-9]+]], [test_insertelement_param_0]; +; COMMON: { .reg .b16 tmp; mov.b32 {[[R0:%rs[0-9]+]], tmp}, [[A]]; } +; COMMON: mov.b32 [[R1:%r[0-9]+]], {[[R0]], [[B]]}; +; COMMON: st.param.b32 [func_retval0+0], [[R1]]; +; COMMON: ret; +define <2 x i16> @test_insertelement(<2 x i16> %a, i16 %x) #0 { + %i = insertelement <2 x i16> %a, i16 %x, i64 1 + ret <2 x i16> %i +} + +attributes #0 = { nounwind } diff --git a/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll b/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll index aa9a1280abcd6..9012339fb6b1e 100644 --- a/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll +++ b/llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll @@ -80,9 +80,9 @@ define void @foo7(ptr noalias readonly %from, ptr %to) { } ; SM20-LABEL: .visible .entry foo8( -; SM20: ld.global.v2.u16 +; SM20: ld.global.u32 ; SM35-LABEL: .visible .entry foo8( -; SM35: ld.global.nc.v2.u16 +; SM35: ld.global.nc.u32 define void @foo8(ptr noalias readonly %from, ptr %to) { %1 = load <2 x i16>, ptr %from store <2 x i16> %1, ptr %to diff --git a/llvm/test/CodeGen/NVPTX/param-load-store.ll b/llvm/test/CodeGen/NVPTX/param-load-store.ll index 313a0915d2030..2d87271e30ae0 100644 --- a/llvm/test/CodeGen/NVPTX/param-load-store.ll +++ b/llvm/test/CodeGen/NVPTX/param-load-store.ll @@ -326,7 +326,8 @@ define signext i16 @test_i16s(i16 signext %a) { ; CHECK-LABEL: test_v3i16( ; CHECK-NEXT: .param .align 8 .b8 test_v3i16_param_0[8] ; CHECK-DAG: ld.param.u16 [[E2:%rs[0-9]+]], [test_v3i16_param_0+4]; -; CHECK-DAG: ld.param.v2.u16 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]]}, [test_v3i16_param_0]; +; CHECK-DAG: ld.param.u32 [[R:%r[0-9]+]], [test_v3i16_param_0]; +; CHECK-DAG: mov.b32 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]]}, [[R]]; ; CHECK: .param .align 8 .b8 param0[8]; ; CHECK: st.param.v2.b16 [param0+0], {[[E0]], [[E1]]}; ; CHECK: st.param.b16 [param0+4], [[E2]]; @@ -346,14 +347,14 @@ define <3 x i16> @test_v3i16(<3 x i16> %a) { ; CHECK: .func (.param .align 8 .b8 func_retval0[8]) ; CHECK-LABEL: test_v4i16( ; CHECK-NEXT: .param .align 8 .b8 test_v4i16_param_0[8] -; CHECK: ld.param.v4.u16 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]], [[E2:%rs[0-9]+]], [[E3:%rs[0-9]+]]}, [test_v4i16_param_0] +; CHECK: ld.param.v2.u32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]]}, [test_v4i16_param_0] ; CHECK: .param .align 8 .b8 param0[8]; -; CHECK: st.param.v4.b16 [param0+0], {[[E0]], [[E1]], [[E2]], [[E3]]}; +; CHECK: st.param.v2.b32 [param0+0], {[[E0]], [[E1]]}; ; CHECK: .param .align 8 .b8 retval0[8]; ; CHECK: call.uni (retval0), ; CHECK-NEXT: test_v4i16, -; CHECK: ld.param.v4.b16 {[[RE0:%rs[0-9]+]], [[RE1:%rs[0-9]+]], [[RE2:%rs[0-9]+]], [[RE3:%rs[0-9]+]]}, [retval0+0]; -; CHECK: st.param.v4.b16 [func_retval0+0], {[[RE0]], [[RE1]], [[RE2]], [[RE3]]} +; CHECK: ld.param.v2.b32 {[[RE0:%r[0-9]+]], [[RE1:%r[0-9]+]]}, [retval0+0]; +; CHECK: st.param.v2.b32 [func_retval0+0], {[[RE0]], [[RE1]]} ; CHECK-NEXT: ret; define <4 x i16> @test_v4i16(<4 x i16> %a) { %r = tail call <4 x i16> @test_v4i16(<4 x i16> %a); @@ -365,6 +366,10 @@ define <4 x i16> @test_v4i16(<4 x i16> %a) { ; CHECK-NEXT: .param .align 16 .b8 test_v5i16_param_0[16] ; CHECK-DAG: ld.param.u16 [[E4:%rs[0-9]+]], [test_v5i16_param_0+8]; ; CHECK-DAG: ld.param.v4.u16 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]], [[E2:%rs[0-9]+]], [[E3:%rs[0-9]+]]}, [test_v5i16_param_0] +; CHECK-DAG: mov.b32 [[R0:%r[0-9]+]], {[[E0]], [[E1]]}; +; CHECK-DAG: mov.b32 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]]}, [[R0]]; +; CHECK-DAG: mov.b32 [[R1:%r[0-9]+]], {[[E2]], [[E3]]}; +; CHECK-DAG: mov.b32 {[[E2:%rs[0-9]+]], [[E3:%rs[0-9]+]]}, [[R1]]; ; CHECK: .param .align 16 .b8 param0[16]; ; CHECK-DAG: st.param.v4.b16 [param0+0], {[[E0]], [[E1]], [[E2]], [[E3]]}; ; CHECK-DAG: st.param.b16 [param0+8], [[E4]]; diff --git a/llvm/test/CodeGen/NVPTX/vec-param-load.ll b/llvm/test/CodeGen/NVPTX/vec-param-load.ll index f8b16c1feab5e..f4f5c26be3474 100644 --- a/llvm/test/CodeGen/NVPTX/vec-param-load.ll +++ b/llvm/test/CodeGen/NVPTX/vec-param-load.ll @@ -70,14 +70,10 @@ define <8 x i64> @test_v8i64(<8 x i64> %a) { define <16 x i16> @test_v16i16(<16 x i16> %a) { ; CHECK-LABEL: test_v16i16( -; CHECK-DAG: ld.param.v4.u16 {[[V_12_15:(%rs[0-9]+[, ]*){4}]]}, [test_v16i16_param_0+24]; -; CHECK-DAG: ld.param.v4.u16 {[[V_8_11:(%rs[0-9]+[, ]*){4}]]}, [test_v16i16_param_0+16]; -; CHECK-DAG: ld.param.v4.u16 {[[V_4_7:(%rs[0-9]+[, ]*){4}]]}, [test_v16i16_param_0+8]; -; CHECK-DAG: ld.param.v4.u16 {[[V_0_3:(%rs[0-9]+[, ]*){4}]]}, [test_v16i16_param_0]; -; CHECK-DAG: st.param.v4.b16 [func_retval0+0], {[[V_0_3]]} -; CHECK-DAG: st.param.v4.b16 [func_retval0+8], {[[V_4_7]]} -; CHECK-DAG: st.param.v4.b16 [func_retval0+16], {[[V_8_11]]} -; CHECK-DAG: st.param.v4.b16 [func_retval0+24], {[[V_12_15]]} +; CHECK-DAG: ld.param.v4.u32 {[[V_8_15:(%r[0-9]+[, ]*){4}]]}, [test_v16i16_param_0+16]; +; CHECK-DAG: ld.param.v4.u32 {[[V_0_7:(%r[0-9]+[, ]*){4}]]}, [test_v16i16_param_0]; +; CHECK-DAG: st.param.v4.b32 [func_retval0+0], {[[V_0_7]]} +; CHECK-DAG: st.param.v4.b32 [func_retval0+16], {[[V_8_15]]} ; CHECK: ret; ret <16 x i16> %a } diff --git a/llvm/test/Transforms/SLPVectorizer/NVPTX/non-vectorizable-intrinsic-inseltpoison.ll b/llvm/test/Transforms/SLPVectorizer/NVPTX/non-vectorizable-intrinsic-inseltpoison.ll deleted file mode 100644 index aeb4aa8078aa2..0000000000000 --- a/llvm/test/Transforms/SLPVectorizer/NVPTX/non-vectorizable-intrinsic-inseltpoison.ll +++ /dev/null @@ -1,57 +0,0 @@ -; NOTE: Assertions have been autogenerated by utils/update_test_checks.py -; RUN: opt < %s -passes=slp-vectorizer -o - -S -slp-threshold=-1000 | FileCheck %s - -target datalayout = "e-p:32:32-i64:64-v16:16-v32:32-n16:32:64" -target triple = "nvptx--nvidiacl" - -; CTLZ cannot be vectorized currently because the second argument is a scalar -; for both the scalar and vector forms of the intrinsic. In the future it -; should be possible to vectorize such functions. -; Test causes an assert if LLVM tries to vectorize CTLZ. - -define <2 x i8> @cltz_test(<2 x i8> %x) #0 { -; CHECK-LABEL: @cltz_test( -; CHECK-NEXT: entry: -; CHECK-NEXT: [[TMP0:%.*]] = extractelement <2 x i8> [[X:%.*]], i32 0 -; CHECK-NEXT: [[CALL_I:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP0]], i1 false) -; CHECK-NEXT: [[VECINIT:%.*]] = insertelement <2 x i8> poison, i8 [[CALL_I]], i32 0 -; CHECK-NEXT: [[TMP1:%.*]] = extractelement <2 x i8> [[X]], i32 1 -; CHECK-NEXT: [[CALL_I4:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP1]], i1 false) -; CHECK-NEXT: [[VECINIT2:%.*]] = insertelement <2 x i8> [[VECINIT]], i8 [[CALL_I4]], i32 1 -; CHECK-NEXT: ret <2 x i8> [[VECINIT2]] -; -entry: - %0 = extractelement <2 x i8> %x, i32 0 - %call.i = call i8 @llvm.ctlz.i8(i8 %0, i1 false) - %vecinit = insertelement <2 x i8> poison, i8 %call.i, i32 0 - %1 = extractelement <2 x i8> %x, i32 1 - %call.i4 = call i8 @llvm.ctlz.i8(i8 %1, i1 false) - %vecinit2 = insertelement <2 x i8> %vecinit, i8 %call.i4, i32 1 - ret <2 x i8> %vecinit2 -} - -define <2 x i8> @cltz_test2(<2 x i8> %x) #1 { -; CHECK-LABEL: @cltz_test2( -; CHECK-NEXT: entry: -; CHECK-NEXT: [[TMP0:%.*]] = extractelement <2 x i8> [[X:%.*]], i32 0 -; CHECK-NEXT: [[TMP1:%.*]] = extractelement <2 x i8> [[X]], i32 1 -; CHECK-NEXT: [[CALL_I:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP0]], i1 false) -; CHECK-NEXT: [[CALL_I4:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP1]], i1 false) -; CHECK-NEXT: [[VECINIT:%.*]] = insertelement <2 x i8> poison, i8 [[CALL_I]], i32 0 -; CHECK-NEXT: [[VECINIT2:%.*]] = insertelement <2 x i8> [[VECINIT]], i8 [[CALL_I4]], i32 1 -; CHECK-NEXT: ret <2 x i8> [[VECINIT2]] -; -entry: - %0 = extractelement <2 x i8> %x, i32 0 - %1 = extractelement <2 x i8> %x, i32 1 - %call.i = call i8 @llvm.ctlz.i8(i8 %0, i1 false) - %call.i4 = call i8 @llvm.ctlz.i8(i8 %1, i1 false) - %vecinit = insertelement <2 x i8> poison, i8 %call.i, i32 0 - %vecinit2 = insertelement <2 x i8> %vecinit, i8 %call.i4, i32 1 - ret <2 x i8> %vecinit2 -} - -declare i8 @llvm.ctlz.i8(i8, i1) #3 - -attributes #0 = { alwaysinline nounwind "less-precise-fpmad"="false" "frame-pointer"="all" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -attributes #1 = { nounwind readnone } diff --git a/llvm/test/Transforms/SLPVectorizer/NVPTX/non-vectorizable-intrinsic.ll b/llvm/test/Transforms/SLPVectorizer/NVPTX/non-vectorizable-intrinsic.ll deleted file mode 100644 index 4f92fa7c3d0f4..0000000000000 --- a/llvm/test/Transforms/SLPVectorizer/NVPTX/non-vectorizable-intrinsic.ll +++ /dev/null @@ -1,57 +0,0 @@ -; NOTE: Assertions have been autogenerated by utils/update_test_checks.py -; RUN: opt < %s -passes=slp-vectorizer -o - -S -slp-threshold=-1000 | FileCheck %s - -target datalayout = "e-p:32:32-i64:64-v16:16-v32:32-n16:32:64" -target triple = "nvptx--nvidiacl" - -; CTLZ cannot be vectorized currently because the second argument is a scalar -; for both the scalar and vector forms of the intrinsic. In the future it -; should be possible to vectorize such functions. -; Test causes an assert if LLVM tries to vectorize CTLZ. - -define <2 x i8> @cltz_test(<2 x i8> %x) #0 { -; CHECK-LABEL: @cltz_test( -; CHECK-NEXT: entry: -; CHECK-NEXT: [[TMP0:%.*]] = extractelement <2 x i8> [[X:%.*]], i32 0 -; CHECK-NEXT: [[CALL_I:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP0]], i1 false) -; CHECK-NEXT: [[VECINIT:%.*]] = insertelement <2 x i8> undef, i8 [[CALL_I]], i32 0 -; CHECK-NEXT: [[TMP1:%.*]] = extractelement <2 x i8> [[X]], i32 1 -; CHECK-NEXT: [[CALL_I4:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP1]], i1 false) -; CHECK-NEXT: [[VECINIT2:%.*]] = insertelement <2 x i8> [[VECINIT]], i8 [[CALL_I4]], i32 1 -; CHECK-NEXT: ret <2 x i8> [[VECINIT2]] -; -entry: - %0 = extractelement <2 x i8> %x, i32 0 - %call.i = call i8 @llvm.ctlz.i8(i8 %0, i1 false) - %vecinit = insertelement <2 x i8> undef, i8 %call.i, i32 0 - %1 = extractelement <2 x i8> %x, i32 1 - %call.i4 = call i8 @llvm.ctlz.i8(i8 %1, i1 false) - %vecinit2 = insertelement <2 x i8> %vecinit, i8 %call.i4, i32 1 - ret <2 x i8> %vecinit2 -} - -define <2 x i8> @cltz_test2(<2 x i8> %x) #1 { -; CHECK-LABEL: @cltz_test2( -; CHECK-NEXT: entry: -; CHECK-NEXT: [[TMP0:%.*]] = extractelement <2 x i8> [[X:%.*]], i32 0 -; CHECK-NEXT: [[TMP1:%.*]] = extractelement <2 x i8> [[X]], i32 1 -; CHECK-NEXT: [[CALL_I:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP0]], i1 false) -; CHECK-NEXT: [[CALL_I4:%.*]] = call i8 @llvm.ctlz.i8(i8 [[TMP1]], i1 false) -; CHECK-NEXT: [[VECINIT:%.*]] = insertelement <2 x i8> undef, i8 [[CALL_I]], i32 0 -; CHECK-NEXT: [[VECINIT2:%.*]] = insertelement <2 x i8> [[VECINIT]], i8 [[CALL_I4]], i32 1 -; CHECK-NEXT: ret <2 x i8> [[VECINIT2]] -; -entry: - %0 = extractelement <2 x i8> %x, i32 0 - %1 = extractelement <2 x i8> %x, i32 1 - %call.i = call i8 @llvm.ctlz.i8(i8 %0, i1 false) - %call.i4 = call i8 @llvm.ctlz.i8(i8 %1, i1 false) - %vecinit = insertelement <2 x i8> undef, i8 %call.i, i32 0 - %vecinit2 = insertelement <2 x i8> %vecinit, i8 %call.i4, i32 1 - ret <2 x i8> %vecinit2 -} - -declare i8 @llvm.ctlz.i8(i8, i1) #3 - -attributes #0 = { alwaysinline nounwind "less-precise-fpmad"="false" "frame-pointer"="all" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -attributes #1 = { nounwind readnone } diff --git a/llvm/test/Transforms/SLPVectorizer/NVPTX/vectorizable-intrinsic.ll b/llvm/test/Transforms/SLPVectorizer/NVPTX/vectorizable-intrinsic.ll new file mode 100644 index 0000000000000..e6ad2adba1759 --- /dev/null +++ b/llvm/test/Transforms/SLPVectorizer/NVPTX/vectorizable-intrinsic.ll @@ -0,0 +1,42 @@ +; RUN: opt < %s -passes=slp-vectorizer -o - -S -slp-threshold=-1000 | FileCheck %s + +target datalayout = "e-p:32:32-i64:64-v16:16-v32:32-n16:32:64" +target triple = "nvptx--nvidiacl" + +; Test that CTLZ can be vectorized currently even though the second argument is a scalar + +define <2 x i8> @cltz_test(<2 x i8> %x) #0 { +; CHECK-LABEL: @cltz_test( +; CHECK: [[VEC:%.*]] = call <2 x i8> @llvm.ctlz.v2i8(<2 x i8> %{{.*}}, i1 false) +; CHECK-NEXT: ret <2 x i8> [[VEC]] +; +entry: + %0 = extractelement <2 x i8> %x, i32 0 + %call.i = call i8 @llvm.ctlz.i8(i8 %0, i1 false) + %vecinit = insertelement <2 x i8> undef, i8 %call.i, i32 0 + %1 = extractelement <2 x i8> %x, i32 1 + %call.i4 = call i8 @llvm.ctlz.i8(i8 %1, i1 false) + %vecinit2 = insertelement <2 x i8> %vecinit, i8 %call.i4, i32 1 + ret <2 x i8> %vecinit2 +} + + +define <2 x i8> @cltz_test_poison(<2 x i8> %x) #0 { +; CHECK-LABEL: @cltz_test_poison( +; CHECK: [[VEC:%.*]] = call <2 x i8> @llvm.ctlz.v2i8(<2 x i8> %{{.*}}, i1 false) +; CHECK-NEXT: ret <2 x i8> [[VEC]] +; +entry: + %0 = extractelement <2 x i8> %x, i32 0 + %call.i = call i8 @llvm.ctlz.i8(i8 %0, i1 false) + %vecinit = insertelement <2 x i8> poison, i8 %call.i, i32 0 + %1 = extractelement <2 x i8> %x, i32 1 + %call.i4 = call i8 @llvm.ctlz.i8(i8 %1, i1 false) + %vecinit2 = insertelement <2 x i8> %vecinit, i8 %call.i4, i32 1 + ret <2 x i8> %vecinit2 +} + +declare i8 @llvm.ctlz.i8(i8, i1) #3 + +attributes #0 = { alwaysinline nounwind "less-precise-fpmad"="false" "frame-pointer"="all" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { nounwind readnone }