diff --git a/llvm/include/llvm/CodeGen/GlobalISel/IRTranslator.h b/llvm/include/llvm/CodeGen/GlobalISel/IRTranslator.h index 3828d859212cb..2f3f55a58a517 100644 --- a/llvm/include/llvm/CodeGen/GlobalISel/IRTranslator.h +++ b/llvm/include/llvm/CodeGen/GlobalISel/IRTranslator.h @@ -299,7 +299,7 @@ class IRTranslator : public MachineFunctionPass { bool translateIntrinsic( const CallBase &CB, Intrinsic::ID ID, MachineIRBuilder &MIRBuilder, - const TargetLowering::IntrinsicInfo *TgtMemIntrinsicInfo = nullptr); + ArrayRef TgtMemIntrinsicInfos = {}); /// When an invoke or a cleanupret unwinds to the next EH pad, there are /// many places it could ultimately go. In the IR, we have a single unwind diff --git a/llvm/include/llvm/CodeGen/SelectionDAG.h b/llvm/include/llvm/CodeGen/SelectionDAG.h index 89619da7c9f50..ed695dc990bae 100644 --- a/llvm/include/llvm/CodeGen/SelectionDAG.h +++ b/llvm/include/llvm/CodeGen/SelectionDAG.h @@ -438,10 +438,18 @@ class SelectionDAG { template static uint16_t getSyntheticNodeSubclassData(unsigned Opc, unsigned Order, - SDVTList VTs, EVT MemoryVT, - MachineMemOperand *MMO) { + SDVTList VTs, EVT MemoryVT, + MachineMemOperand *MMO) { return SDNodeTy(Opc, Order, DebugLoc(), VTs, MemoryVT, MMO) - .getRawSubclassData(); + .getRawSubclassData(); + } + + template + static uint16_t getSyntheticNodeSubclassData( + unsigned Opc, unsigned Order, SDVTList VTs, EVT MemoryVT, + PointerUnion MemRefs) { + return SDNodeTy(Opc, Order, DebugLoc(), VTs, MemoryVT, MemRefs) + .getRawSubclassData(); } void createOperands(SDNode *Node, ArrayRef Vals); @@ -1481,6 +1489,12 @@ class SelectionDAG { SDVTList VTList, ArrayRef Ops, EVT MemVT, MachineMemOperand *MMO); + /// getMemIntrinsicNode - Creates a MemIntrinsicNode with multiple MMOs. + LLVM_ABI SDValue getMemIntrinsicNode(unsigned Opcode, const SDLoc &dl, + SDVTList VTList, ArrayRef Ops, + EVT MemVT, + ArrayRef MMOs); + /// Creates a LifetimeSDNode that starts (`IsStart==true`) or ends /// (`IsStart==false`) the lifetime of the `FrameIndex`. LLVM_ABI SDValue getLifetimeNode(bool IsStart, const SDLoc &dl, SDValue Chain, diff --git a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h index 536dca4602c03..a50bd8dab5407 100644 --- a/llvm/include/llvm/CodeGen/SelectionDAGNodes.h +++ b/llvm/include/llvm/CodeGen/SelectionDAGNodes.h @@ -1411,19 +1411,26 @@ class MemSDNode : public SDNode { EVT MemoryVT; protected: - /// Memory reference information. - MachineMemOperand *MMO; + /// Memory reference information. Must always have at least one MMO. + /// - MachineMemOperand*: exactly 1 MMO (common case) + /// - MachineMemOperand**: pointer to array, size at offset -1 + PointerUnion MemRefs; public: - LLVM_ABI MemSDNode(unsigned Opc, unsigned Order, const DebugLoc &dl, - SDVTList VTs, EVT memvt, MachineMemOperand *MMO); + /// Constructor that supports single or multiple MMOs. For single MMO, pass + /// the MMO pointer directly. For multiple MMOs, pre-allocate storage with + /// count at offset -1 and pass pointer to array. + LLVM_ABI + MemSDNode(unsigned Opc, unsigned Order, const DebugLoc &dl, SDVTList VTs, + EVT memvt, + PointerUnion memrefs); - bool readMem() const { return MMO->isLoad(); } - bool writeMem() const { return MMO->isStore(); } + bool readMem() const { return getMemOperand()->isLoad(); } + bool writeMem() const { return getMemOperand()->isStore(); } /// Returns alignment and volatility of the memory access - Align getBaseAlign() const { return MMO->getBaseAlign(); } - Align getAlign() const { return MMO->getAlign(); } + Align getBaseAlign() const { return getMemOperand()->getBaseAlign(); } + Align getAlign() const { return getMemOperand()->getAlign(); } /// Return the SubclassData value, without HasDebugValue. This contains an /// encoding of the volatile flag, as well as bits used by subclasses. This @@ -1450,36 +1457,40 @@ class MemSDNode : public SDNode { bool isInvariant() const { return MemSDNodeBits.IsInvariant; } // Returns the offset from the location of the access. - int64_t getSrcValueOffset() const { return MMO->getOffset(); } + int64_t getSrcValueOffset() const { return getMemOperand()->getOffset(); } /// Returns the AA info that describes the dereference. - AAMDNodes getAAInfo() const { return MMO->getAAInfo(); } + AAMDNodes getAAInfo() const { return getMemOperand()->getAAInfo(); } /// Returns the Ranges that describes the dereference. - const MDNode *getRanges() const { return MMO->getRanges(); } + const MDNode *getRanges() const { return getMemOperand()->getRanges(); } /// Returns the synchronization scope ID for this memory operation. - SyncScope::ID getSyncScopeID() const { return MMO->getSyncScopeID(); } + SyncScope::ID getSyncScopeID() const { + return getMemOperand()->getSyncScopeID(); + } /// Return the atomic ordering requirements for this memory operation. For /// cmpxchg atomic operations, return the atomic ordering requirements when /// store occurs. AtomicOrdering getSuccessOrdering() const { - return MMO->getSuccessOrdering(); + return getMemOperand()->getSuccessOrdering(); } /// Return a single atomic ordering that is at least as strong as both the /// success and failure orderings for an atomic operation. (For operations /// other than cmpxchg, this is equivalent to getSuccessOrdering().) - AtomicOrdering getMergedOrdering() const { return MMO->getMergedOrdering(); } + AtomicOrdering getMergedOrdering() const { + return getMemOperand()->getMergedOrdering(); + } /// Return true if the memory operation ordering is Unordered or higher. - bool isAtomic() const { return MMO->isAtomic(); } + bool isAtomic() const { return getMemOperand()->isAtomic(); } /// Returns true if the memory operation doesn't imply any ordering /// constraints on surrounding memory operations beyond the normal memory /// aliasing rules. - bool isUnordered() const { return MMO->isUnordered(); } + bool isUnordered() const { return getMemOperand()->isUnordered(); } /// Returns true if the memory operation is neither atomic or volatile. bool isSimple() const { return !isAtomic() && !isVolatile(); } @@ -1487,12 +1498,37 @@ class MemSDNode : public SDNode { /// Return the type of the in-memory value. EVT getMemoryVT() const { return MemoryVT; } - /// Return a MachineMemOperand object describing the memory + /// Return the unique MachineMemOperand object describing the memory /// reference performed by operation. - MachineMemOperand *getMemOperand() const { return MMO; } + /// Asserts if multiple MMOs are present - use memoperands() instead. + MachineMemOperand *getMemOperand() const { + assert(!isa(MemRefs) && + "Use memoperands() for nodes with multiple memory operands"); + return cast(MemRefs); + } + + /// Return the number of memory operands. + size_t getNumMemOperands() const { + if (isa(MemRefs)) + return 1; + MachineMemOperand **Array = cast(MemRefs); + return reinterpret_cast(Array)[-1]; + } + + /// Return true if this node has exactly one memory operand. + bool hasUniqueMemOperand() const { return isa(MemRefs); } + + /// Return the memory operands for this node. + ArrayRef memoperands() const { + if (isa(MemRefs)) + return ArrayRef(MemRefs.getAddrOfPtr1(), 1); + MachineMemOperand **Array = cast(MemRefs); + size_t Count = reinterpret_cast(Array)[-1]; + return ArrayRef(Array, Count); + } const MachinePointerInfo &getPointerInfo() const { - return MMO->getPointerInfo(); + return getMemOperand()->getPointerInfo(); } /// Return the address space for the associated pointer @@ -1501,19 +1537,35 @@ class MemSDNode : public SDNode { } /// Update this MemSDNode's MachineMemOperand information - /// to reflect the alignment of NewMMO, if it has a greater alignment. + /// to reflect the alignment of NewMMOs, if they have greater alignment. /// This must only be used when the new alignment applies to all users of - /// this MachineMemOperand. - void refineAlignment(const MachineMemOperand *NewMMO) { - MMO->refineAlignment(NewMMO); + /// these MachineMemOperands. The NewMMOs array must parallel memoperands(). + void refineAlignment(ArrayRef NewMMOs) { + ArrayRef MMOs = memoperands(); + assert(NewMMOs.size() == MMOs.size() && "MMO count mismatch"); + for (auto [MMO, NewMMO] : zip(MMOs, NewMMOs)) + MMO->refineAlignment(NewMMO); + } + + void refineAlignment(MachineMemOperand *NewMMO) { + refineAlignment(ArrayRef(NewMMO)); } - void refineRanges(const MachineMemOperand *NewMMO) { - // If this node has range metadata that is different than NewMMO, clear the - // range metadata. + /// Refine range metadata for all MMOs. The NewMMOs array must parallel + /// memoperands(). For each pair, if ranges differ, the stored range is + /// cleared. + void refineRanges(ArrayRef NewMMOs) { + ArrayRef MMOs = memoperands(); + assert(NewMMOs.size() == MMOs.size() && "MMO count mismatch"); // FIXME: Union the ranges instead? - if (getRanges() && getRanges() != NewMMO->getRanges()) - MMO->clearRanges(); + for (auto [MMO, NewMMO] : zip(MMOs, NewMMOs)) { + if (MMO->getRanges() && MMO->getRanges() != NewMMO->getRanges()) + MMO->clearRanges(); + } + } + + void refineRanges(MachineMemOperand *NewMMO) { + refineRanges(ArrayRef(NewMMO)); } const SDValue &getChain() const { return getOperand(0); } @@ -1626,7 +1678,7 @@ class AtomicSDNode : public MemSDNode { /// when store does not occur. AtomicOrdering getFailureOrdering() const { assert(isCompareAndSwap() && "Must be cmpxchg operation"); - return MMO->getFailureOrdering(); + return getMemOperand()->getFailureOrdering(); } // Methods to support isa and dyn_cast @@ -1666,9 +1718,11 @@ class AtomicSDNode : public MemSDNode { /// opcode (see `SelectionDAGTargetInfo::isTargetMemoryOpcode`). class MemIntrinsicSDNode : public MemSDNode { public: - MemIntrinsicSDNode(unsigned Opc, unsigned Order, const DebugLoc &dl, - SDVTList VTs, EVT MemoryVT, MachineMemOperand *MMO) - : MemSDNode(Opc, Order, dl, VTs, MemoryVT, MMO) { + MemIntrinsicSDNode( + unsigned Opc, unsigned Order, const DebugLoc &dl, SDVTList VTs, + EVT MemoryVT, + PointerUnion MemRefs) + : MemSDNode(Opc, Order, dl, VTs, MemoryVT, MemRefs) { SDNodeBits.IsMemIntrinsic = true; } diff --git a/llvm/include/llvm/CodeGen/TargetLowering.h b/llvm/include/llvm/CodeGen/TargetLowering.h index 442225bdec01f..ada4ffd3bcc89 100644 --- a/llvm/include/llvm/CodeGen/TargetLowering.h +++ b/llvm/include/llvm/CodeGen/TargetLowering.h @@ -1244,15 +1244,32 @@ class LLVM_ABI TargetLoweringBase { }; /// Given an intrinsic, checks if on the target the intrinsic will need to map - /// to a MemIntrinsicNode (touches memory). If this is the case, it returns - /// true and store the intrinsic information into the IntrinsicInfo that was - /// passed to the function. + /// to a MemIntrinsicNode (touches memory). If this is the case, it stores + /// the intrinsic information into the IntrinsicInfo vector passed to the + /// function. The vector may contain multiple entries for intrinsics that + /// access multiple memory locations. + virtual void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, + unsigned Intrinsic) const { + // The default implementation forwards to the legacy single-info overload + // for compatibility. + IntrinsicInfo Info; + if (getTgtMemIntrinsic(Info, I, MF, Intrinsic)) + Infos.push_back(Info); + } + +protected: + /// This is a legacy single-info overload. New code should override the + /// SmallVectorImpl overload instead to support multiple memory operands. + /// + /// TODO: Remove this once the refactoring is complete. virtual bool getTgtMemIntrinsic(IntrinsicInfo &, const CallBase &, MachineFunction &, unsigned /*Intrinsic*/) const { return false; } +public: /// Returns true if the target can instruction select the specified FP /// immediate natively. If false, the legalizer will materialize the FP /// immediate as a load from a constant pool. diff --git a/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp b/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp index a0fe900778cca..126199849b033 100644 --- a/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp +++ b/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp @@ -2819,20 +2819,16 @@ bool IRTranslator::translateCall(const User &U, MachineIRBuilder &MIRBuilder) { if (translateKnownIntrinsic(CI, ID, MIRBuilder)) return true; - TargetLowering::IntrinsicInfo Info; - bool IsTgtMemIntrinsic = TLI->getTgtMemIntrinsic(Info, CI, *MF, ID); + SmallVector Infos; + TLI->getTgtMemIntrinsic(Infos, CI, *MF, ID); - return translateIntrinsic(CI, ID, MIRBuilder, - IsTgtMemIntrinsic ? &Info : nullptr); + return translateIntrinsic(CI, ID, MIRBuilder, Infos); } /// Translate a call or callbr to an intrinsic. -/// Depending on whether TLI->getTgtMemIntrinsic() is true, TgtMemIntrinsicInfo -/// is a pointer to the correspondingly populated IntrinsicInfo object. -/// Otherwise, this pointer is null. bool IRTranslator::translateIntrinsic( const CallBase &CB, Intrinsic::ID ID, MachineIRBuilder &MIRBuilder, - const TargetLowering::IntrinsicInfo *TgtMemIntrinsicInfo) { + ArrayRef TgtMemIntrinsicInfos) { ArrayRef ResultRegs; if (!CB.getType()->isVoidTy()) ResultRegs = getOrCreateVRegs(CB); @@ -2874,30 +2870,25 @@ bool IRTranslator::translateIntrinsic( } } - // Add a MachineMemOperand if it is a target mem intrinsic. - if (TgtMemIntrinsicInfo) { - const Function *F = CB.getCalledFunction(); + // Add MachineMemOperands for each memory access described by the target. + for (const auto &Info : TgtMemIntrinsicInfos) { + Align Alignment = Info.align.value_or( + DL->getABITypeAlign(Info.memVT.getTypeForEVT(CB.getContext()))); + LLT MemTy = Info.memVT.isSimple() + ? getLLTForMVT(Info.memVT.getSimpleVT()) + : LLT::scalar(Info.memVT.getStoreSizeInBits()); - Align Alignment = TgtMemIntrinsicInfo->align.value_or(DL->getABITypeAlign( - TgtMemIntrinsicInfo->memVT.getTypeForEVT(F->getContext()))); - LLT MemTy = - TgtMemIntrinsicInfo->memVT.isSimple() - ? getLLTForMVT(TgtMemIntrinsicInfo->memVT.getSimpleVT()) - : LLT::scalar(TgtMemIntrinsicInfo->memVT.getStoreSizeInBits()); - - // TODO: We currently just fallback to address space 0 if getTgtMemIntrinsic - // didn't yield anything useful. + // TODO: We currently just fallback to address space 0 if + // getTgtMemIntrinsic didn't yield anything useful. MachinePointerInfo MPI; - if (TgtMemIntrinsicInfo->ptrVal) { - MPI = MachinePointerInfo(TgtMemIntrinsicInfo->ptrVal, - TgtMemIntrinsicInfo->offset); - } else if (TgtMemIntrinsicInfo->fallbackAddressSpace) { - MPI = MachinePointerInfo(*TgtMemIntrinsicInfo->fallbackAddressSpace); + if (Info.ptrVal) { + MPI = MachinePointerInfo(Info.ptrVal, Info.offset); + } else if (Info.fallbackAddressSpace) { + MPI = MachinePointerInfo(*Info.fallbackAddressSpace); } MIB.addMemOperand(MF->getMachineMemOperand( - MPI, TgtMemIntrinsicInfo->flags, MemTy, Alignment, CB.getAAMetadata(), - /*Ranges=*/nullptr, TgtMemIntrinsicInfo->ssid, - TgtMemIntrinsicInfo->order, TgtMemIntrinsicInfo->failureOrder)); + MPI, Info.flags, MemTy, Alignment, CB.getAAMetadata(), + /*Ranges=*/nullptr, Info.ssid, Info.order, Info.failureOrder)); } if (CB.isConvergent()) { diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp index eb15aa8ce2261..df69f0870d27a 100644 --- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp @@ -1208,7 +1208,7 @@ bool DAGCombiner::reassociationCanBreakAddressingModePattern(unsigned Opc, for (SDNode *Node : N->users()) { auto *LoadStore = dyn_cast(Node); - if (!LoadStore) + if (!LoadStore || !LoadStore->hasUniqueMemOperand()) return false; // Is x[offset2] a legal addressing mode? If so then diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp index c49e056dba5ac..302b8059e4df0 100644 --- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp @@ -989,9 +989,11 @@ static void AddNodeIDCustom(FoldingSetNodeID &ID, const SDNode *N) { // to check. if (auto *MN = dyn_cast(N)) { ID.AddInteger(MN->getRawSubclassData()); - ID.AddInteger(MN->getPointerInfo().getAddrSpace()); - ID.AddInteger(MN->getMemOperand()->getFlags()); ID.AddInteger(MN->getMemoryVT().getRawBits()); + for (const MachineMemOperand *MMO : MN->memoperands()) { + ID.AddInteger(MMO->getPointerInfo().getAddrSpace()); + ID.AddInteger(MMO->getFlags()); + } } } @@ -1304,7 +1306,7 @@ SelectionDAG::AddModifiedNodeToCSEMaps(SDNode *N) { // recursive merging of other unrelated nodes down the line. Existing->intersectFlagsWith(N->getFlags()); if (auto *MemNode = dyn_cast(Existing)) - MemNode->refineRanges(cast(N)->getMemOperand()); + MemNode->refineRanges(cast(N)->memoperands()); ReplaceAllUsesWith(N, Existing); // N is now dead. Inform the listeners and delete it. @@ -9831,6 +9833,14 @@ SDValue SelectionDAG::getMemIntrinsicNode(unsigned Opcode, const SDLoc &dl, SDVTList VTList, ArrayRef Ops, EVT MemVT, MachineMemOperand *MMO) { + return getMemIntrinsicNode(Opcode, dl, VTList, Ops, MemVT, ArrayRef(MMO)); +} + +SDValue SelectionDAG::getMemIntrinsicNode(unsigned Opcode, const SDLoc &dl, + SDVTList VTList, + ArrayRef Ops, EVT MemVT, + ArrayRef MMOs) { + assert(!MMOs.empty() && "Must have at least one MMO"); assert( (Opcode == ISD::INTRINSIC_VOID || Opcode == ISD::INTRINSIC_W_CHAIN || Opcode == ISD::PREFETCH || @@ -9838,30 +9848,47 @@ SDValue SelectionDAG::getMemIntrinsicNode(unsigned Opcode, const SDLoc &dl, Opcode >= ISD::BUILTIN_OP_END && TSI->isTargetMemoryOpcode(Opcode))) && "Opcode is not a memory-accessing opcode!"); + PointerUnion MemRefs; + if (MMOs.size() == 1) { + MemRefs = MMOs[0]; + } else { + // Allocate: [size_t count][MMO*][MMO*]... + size_t AllocSize = + sizeof(size_t) + MMOs.size() * sizeof(MachineMemOperand *); + void *Buffer = Allocator.Allocate(AllocSize, alignof(size_t)); + size_t *CountPtr = static_cast(Buffer); + *CountPtr = MMOs.size(); + MachineMemOperand **Array = + reinterpret_cast(CountPtr + 1); + llvm::copy(MMOs, Array); + MemRefs = Array; + } + // Memoize the node unless it returns a glue result. MemIntrinsicSDNode *N; if (VTList.VTs[VTList.NumVTs-1] != MVT::Glue) { FoldingSetNodeID ID; AddNodeIDNode(ID, Opcode, VTList, Ops); ID.AddInteger(getSyntheticNodeSubclassData( - Opcode, dl.getIROrder(), VTList, MemVT, MMO)); - ID.AddInteger(MMO->getPointerInfo().getAddrSpace()); - ID.AddInteger(MMO->getFlags()); + Opcode, dl.getIROrder(), VTList, MemVT, MemRefs)); ID.AddInteger(MemVT.getRawBits()); + for (const MachineMemOperand *MMO : MMOs) { + ID.AddInteger(MMO->getPointerInfo().getAddrSpace()); + ID.AddInteger(MMO->getFlags()); + } void *IP = nullptr; if (SDNode *E = FindNodeOrInsertPos(ID, dl, IP)) { - cast(E)->refineAlignment(MMO); + cast(E)->refineAlignment(MMOs); return SDValue(E, 0); } N = newSDNode(Opcode, dl.getIROrder(), dl.getDebugLoc(), - VTList, MemVT, MMO); + VTList, MemVT, MemRefs); createOperands(N, Ops); - - CSEMap.InsertNode(N, IP); + CSEMap.InsertNode(N, IP); } else { N = newSDNode(Opcode, dl.getIROrder(), dl.getDebugLoc(), - VTList, MemVT, MMO); + VTList, MemVT, MemRefs); createOperands(N, Ops); } InsertNode(N); @@ -13285,21 +13312,33 @@ HandleSDNode::~HandleSDNode() { DropOperands(); } -MemSDNode::MemSDNode(unsigned Opc, unsigned Order, const DebugLoc &dl, - SDVTList VTs, EVT memvt, MachineMemOperand *mmo) - : SDNode(Opc, Order, dl, VTs), MemoryVT(memvt), MMO(mmo) { - MemSDNodeBits.IsVolatile = MMO->isVolatile(); - MemSDNodeBits.IsNonTemporal = MMO->isNonTemporal(); - MemSDNodeBits.IsDereferenceable = MMO->isDereferenceable(); - MemSDNodeBits.IsInvariant = MMO->isInvariant(); - - // We check here that the size of the memory operand fits within the size of - // the MMO. This is because the MMO might indicate only a possible address - // range instead of specifying the affected memory addresses precisely. - assert( - (!MMO->getType().isValid() || - TypeSize::isKnownLE(memvt.getStoreSize(), MMO->getSize().getValue())) && - "Size mismatch!"); +MemSDNode::MemSDNode( + unsigned Opc, unsigned Order, const DebugLoc &dl, SDVTList VTs, EVT memvt, + PointerUnion memrefs) + : SDNode(Opc, Order, dl, VTs), MemoryVT(memvt), MemRefs(memrefs) { + bool IsVolatile = false; + bool IsNonTemporal = false; + bool IsDereferenceable = true; + bool IsInvariant = true; + for (const MachineMemOperand *MMO : memoperands()) { + IsVolatile |= MMO->isVolatile(); + IsNonTemporal |= MMO->isNonTemporal(); + IsDereferenceable &= MMO->isDereferenceable(); + IsInvariant &= MMO->isInvariant(); + } + MemSDNodeBits.IsVolatile = IsVolatile; + MemSDNodeBits.IsNonTemporal = IsNonTemporal; + MemSDNodeBits.IsDereferenceable = IsDereferenceable; + MemSDNodeBits.IsInvariant = IsInvariant; + + // For the single-MMO case, we check here that the size of the memory operand + // fits within the size of the MMO. This is because the MMO might indicate + // only a possible address range instead of specifying the affected memory + // addresses precisely. + assert((getNumMemOperands() != 1 || !getMemOperand()->getType().isValid() || + TypeSize::isKnownLE(memvt.getStoreSize(), + getMemOperand()->getSize().getValue())) && + "Size mismatch!"); } /// Profile - Gather unique data for the node. diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp index 18cb69a47d85f..6045b55130925 100644 --- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp @@ -3514,10 +3514,12 @@ void SelectionDAGBuilder::visitInvoke(const InvokeInst &I) { /// - they do not need custom argument handling (no /// TLI.CollectTargetIntrinsicOperands()) void SelectionDAGBuilder::visitCallBrIntrinsic(const CallBrInst &I) { - TargetLowering::IntrinsicInfo Info; - assert(!DAG.getTargetLoweringInfo().getTgtMemIntrinsic( - Info, I, DAG.getMachineFunction(), I.getIntrinsicID()) && - "Intrinsic touches memory"); +#ifndef NDEBUG + SmallVector Infos; + DAG.getTargetLoweringInfo().getTgtMemIntrinsic( + Infos, I, DAG.getMachineFunction(), I.getIntrinsicID()); + assert(Infos.empty() && "Intrinsic touches memory"); +#endif auto [HasChain, OnlyLoad] = getTargetIntrinsicCallProperties(I); @@ -5485,14 +5487,15 @@ void SelectionDAGBuilder::visitTargetIntrinsic(const CallInst &I, unsigned Intrinsic) { auto [HasChain, OnlyLoad] = getTargetIntrinsicCallProperties(I); - // Info is set by getTgtMemIntrinsic - TargetLowering::IntrinsicInfo Info; + // Infos is set by getTgtMemIntrinsic. + SmallVector Infos; const TargetLowering &TLI = DAG.getTargetLoweringInfo(); - bool IsTgtMemIntrinsic = - TLI.getTgtMemIntrinsic(Info, I, DAG.getMachineFunction(), Intrinsic); + TLI.getTgtMemIntrinsic(Infos, I, DAG.getMachineFunction(), Intrinsic); + // Use the first (primary) info determines the node opcode. + TargetLowering::IntrinsicInfo *Info = !Infos.empty() ? &Infos[0] : nullptr; - SmallVector Ops = getTargetIntrinsicOperands( - I, HasChain, OnlyLoad, IsTgtMemIntrinsic ? &Info : nullptr); + SmallVector Ops = + getTargetIntrinsicOperands(I, HasChain, OnlyLoad, Info); SDVTList VTs = getTargetIntrinsicVTList(I, HasChain); // Propagate fast-math-flags from IR to node(s). @@ -5506,26 +5509,32 @@ void SelectionDAGBuilder::visitTargetIntrinsic(const CallInst &I, // In some cases, custom collection of operands from CallInst I may be needed. TLI.CollectTargetIntrinsicOperands(I, Ops, DAG); - if (IsTgtMemIntrinsic) { + if (!Infos.empty()) { // This is target intrinsic that touches memory - // - // TODO: We currently just fallback to address space 0 if getTgtMemIntrinsic - // didn't yield anything useful. - MachinePointerInfo MPI; - if (Info.ptrVal) - MPI = MachinePointerInfo(Info.ptrVal, Info.offset); - else if (Info.fallbackAddressSpace) - MPI = MachinePointerInfo(*Info.fallbackAddressSpace); - EVT MemVT = Info.memVT; - LocationSize Size = LocationSize::precise(Info.size); - if (Size.hasValue() && !Size.getValue()) - Size = LocationSize::precise(MemVT.getStoreSize()); - Align Alignment = Info.align.value_or(DAG.getEVTAlign(MemVT)); - MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand( - MPI, Info.flags, Size, Alignment, I.getAAMetadata(), /*Ranges=*/nullptr, - Info.ssid, Info.order, Info.failureOrder); - Result = - DAG.getMemIntrinsicNode(Info.opc, getCurSDLoc(), VTs, Ops, MemVT, MMO); + // Create MachineMemOperands for each memory access described by the target. + MachineFunction &MF = DAG.getMachineFunction(); + SmallVector MMOs; + for (const auto &Info : Infos) { + // TODO: We currently just fallback to address space 0 if + // getTgtMemIntrinsic didn't yield anything useful. + MachinePointerInfo MPI; + if (Info.ptrVal) + MPI = MachinePointerInfo(Info.ptrVal, Info.offset); + else if (Info.fallbackAddressSpace) + MPI = MachinePointerInfo(*Info.fallbackAddressSpace); + EVT MemVT = Info.memVT; + LocationSize Size = LocationSize::precise(Info.size); + if (Size.hasValue() && !Size.getValue()) + Size = LocationSize::precise(MemVT.getStoreSize()); + Align Alignment = Info.align.value_or(DAG.getEVTAlign(MemVT)); + MachineMemOperand *MMO = MF.getMachineMemOperand( + MPI, Info.flags, Size, Alignment, I.getAAMetadata(), + /*Ranges=*/nullptr, Info.ssid, Info.order, Info.failureOrder); + MMOs.push_back(MMO); + } + + Result = DAG.getMemIntrinsicNode(Info->opc, getCurSDLoc(), VTs, Ops, + Info->memVT, MMOs); } else { Result = getTargetNonMemIntrinsicNode(*I.getType(), HasChain, Ops, VTs); } diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp index 66ecb40e48954..a213396f3df90 100644 --- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp @@ -933,7 +933,9 @@ void SDNode::print_details(raw_ostream &OS, const SelectionDAG *G) const { OS << ">"; } else if (const MemSDNode *M = dyn_cast(this)) { OS << "<"; - printMemOperand(OS, *M->getMemOperand(), G); + interleaveComma(M->memoperands(), OS, [&](const MachineMemOperand *MMO) { + printMemOperand(OS, *MMO, G); + }); if (auto *A = dyn_cast(M)) if (A->getOpcode() == ISD::ATOMIC_LOAD) { bool doExt = true; diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp index bd21c95c0ff93..e7cb0a3574b4a 100644 --- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp @@ -3597,7 +3597,7 @@ void SelectionDAGISel::SelectCodeCommon(SDNode *NodeToMatch, } case OPC_RecordMemRef: if (auto *MN = dyn_cast(N)) - MatchedMemRefs.push_back(MN->getMemOperand()); + llvm::append_range(MatchedMemRefs, MN->memoperands()); else { LLVM_DEBUG(dbgs() << "Expected MemSDNode "; N->dump(CurDAG); dbgs() << '\n'); diff --git a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp index 003ec37b86f56..8c815ac940511 100644 --- a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp +++ b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp @@ -17453,7 +17453,7 @@ SDValue AArch64TargetLowering::LowerVSCALE(SDValue Op, /// Set the IntrinsicInfo for the `aarch64_sve_st` intrinsics. template -static bool +static void setInfoSVEStN(const AArch64TargetLowering &TLI, const DataLayout &DL, AArch64TargetLowering::IntrinsicInfo &Info, const CallBase &CI) { Info.opc = ISD::INTRINSIC_VOID; @@ -17473,24 +17473,29 @@ setInfoSVEStN(const AArch64TargetLowering &TLI, const DataLayout &DL, Info.offset = 0; Info.align.reset(); Info.flags = MachineMemOperand::MOStore; - return true; } /// getTgtMemIntrinsic - Represent NEON load and store intrinsics as /// MemIntrinsicNodes. The associated MachineMemOperands record the alignment /// specified in the intrinsic calls. -bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void AArch64TargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; auto &DL = I.getDataLayout(); switch (Intrinsic) { case Intrinsic::aarch64_sve_st2: - return setInfoSVEStN<2>(*this, DL, Info, I); + setInfoSVEStN<2>(*this, DL, Info, I); + Infos.push_back(Info); + return; case Intrinsic::aarch64_sve_st3: - return setInfoSVEStN<3>(*this, DL, Info, I); + setInfoSVEStN<3>(*this, DL, Info, I); + Infos.push_back(Info); + return; case Intrinsic::aarch64_sve_st4: - return setInfoSVEStN<4>(*this, DL, Info, I); + setInfoSVEStN<4>(*this, DL, Info, I); + Infos.push_back(Info); + return; case Intrinsic::aarch64_neon_ld2: case Intrinsic::aarch64_neon_ld3: case Intrinsic::aarch64_neon_ld4: @@ -17505,7 +17510,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align.reset(); // volatile loads with NEON intrinsics not supported Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_neon_ld2lane: case Intrinsic::aarch64_neon_ld3lane: @@ -17526,7 +17532,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align.reset(); // volatile loads with NEON intrinsics not supported Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_neon_st2: case Intrinsic::aarch64_neon_st3: @@ -17548,7 +17555,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align.reset(); // volatile stores with NEON intrinsics not supported Info.flags = MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_neon_st2lane: case Intrinsic::aarch64_neon_st3lane: @@ -17572,7 +17580,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align.reset(); // volatile stores with NEON intrinsics not supported Info.flags = MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_ldaxr: case Intrinsic::aarch64_ldxr: { @@ -17583,7 +17592,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = DL.getABITypeAlign(ValTy); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_stlxr: case Intrinsic::aarch64_stxr: { @@ -17594,7 +17604,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = DL.getABITypeAlign(ValTy); Info.flags = MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_ldaxp: case Intrinsic::aarch64_ldxp: @@ -17604,7 +17615,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(16); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; case Intrinsic::aarch64_stlxp: case Intrinsic::aarch64_stxp: Info.opc = ISD::INTRINSIC_W_CHAIN; @@ -17613,7 +17625,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(16); Info.flags = MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; case Intrinsic::aarch64_sve_ldnt1: { Type *ElTy = cast(I.getType())->getElementType(); Info.opc = ISD::INTRINSIC_W_CHAIN; @@ -17622,7 +17635,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = DL.getABITypeAlign(ElTy); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MONonTemporal; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_sve_stnt1: { Type *ElTy = @@ -17633,7 +17647,8 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = DL.getABITypeAlign(ElTy); Info.flags = MachineMemOperand::MOStore | MachineMemOperand::MONonTemporal; - return true; + Infos.push_back(Info); + return; } case Intrinsic::aarch64_mops_memset_tag: { Value *Dst = I.getArgOperand(0); @@ -17646,13 +17661,12 @@ bool AArch64TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.flags = MachineMemOperand::MOStore; // The size of the memory being operated on is unknown at this point Info.size = MemoryLocation::UnknownSize; - return true; + Infos.push_back(Info); + return; } default: break; } - - return false; } bool AArch64TargetLowering::shouldReduceLoadWidth( diff --git a/llvm/lib/Target/AArch64/AArch64ISelLowering.h b/llvm/lib/Target/AArch64/AArch64ISelLowering.h index aa6110a4ce39d..89a8858550ca2 100644 --- a/llvm/lib/Target/AArch64/AArch64ISelLowering.h +++ b/llvm/lib/Target/AArch64/AArch64ISelLowering.h @@ -208,8 +208,8 @@ class AArch64TargetLowering : public TargetLowering { EmitInstrWithCustomInserter(MachineInstr &MI, MachineBasicBlock *MBB) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; bool shouldReduceLoadWidth(SDNode *Load, ISD::LoadExtType ExtTy, EVT NewVT, diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index b043d5354042d..46b07100b8fe0 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -1340,10 +1340,11 @@ static void getCoopAtomicOperandsInfo(const CallBase &CI, bool IsLoad, Info.ssid = CI.getContext().getOrInsertSyncScopeID(Scope); } -bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, +void SITargetLowering::getTgtMemIntrinsic(SmallVectorImpl &Infos, const CallBase &CI, MachineFunction &MF, unsigned IntrID) const { + IntrinsicInfo Info; Info.flags = MachineMemOperand::MONone; if (CI.hasMetadata(LLVMContext::MD_invariant_load)) Info.flags |= MachineMemOperand::MOInvariant; @@ -1357,7 +1358,7 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Intrinsic::getFnAttributes(CI.getContext(), (Intrinsic::ID)IntrID); MemoryEffects ME = Attr.getMemoryEffects(); if (ME.doesNotAccessMemory()) - return false; + return; // TODO: Should images get their own address space? Info.fallbackAddressSpace = AMDGPUAS::BUFFER_RESOURCE; @@ -1453,7 +1454,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, unsigned Width = cast(CI.getArgOperand(2))->getZExtValue(); Info.memVT = EVT::getIntegerVT(CI.getContext(), Width * 8); Info.ptrVal = CI.getArgOperand(1); - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_raw_atomic_buffer_load: case Intrinsic::amdgcn_raw_ptr_atomic_buffer_load: @@ -1463,11 +1465,13 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, memVTFromLoadIntrReturn(*this, MF.getDataLayout(), CI.getType(), std::numeric_limits::max()); Info.flags &= ~MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } } } - return true; + Infos.push_back(Info); + return; } switch (IntrID) { @@ -1483,7 +1487,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, if (!Vol->isZero()) Info.flags |= MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_ds_add_gs_reg_rtn: case Intrinsic::amdgcn_ds_sub_gs_reg_rtn: { @@ -1492,7 +1497,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.ptrVal = nullptr; Info.fallbackAddressSpace = AMDGPUAS::STREAMOUT_REGISTER; Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_ds_append: case Intrinsic::amdgcn_ds_consume: { @@ -1506,7 +1512,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, if (!Vol->isZero()) Info.flags |= MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_ds_atomic_async_barrier_arrive_b64: case Intrinsic::amdgcn_ds_atomic_barrier_arrive_rtn_b64: { @@ -1519,7 +1526,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.size = 8; Info.align.reset(); Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_image_bvh_dual_intersect_ray: case Intrinsic::amdgcn_image_bvh_intersect_ray: @@ -1535,7 +1543,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align.reset(); Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_global_atomic_fmin_num: case Intrinsic::amdgcn_global_atomic_fmax_num: @@ -1549,7 +1558,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MODereferenceable | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_flat_load_monitor_b32: case Intrinsic::amdgcn_flat_load_monitor_b64: @@ -1577,7 +1587,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.ptrVal = CI.getOperand(0); Info.align.reset(); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_cooperative_atomic_load_32x4B: case Intrinsic::amdgcn_cooperative_atomic_load_16x8B: @@ -1587,7 +1598,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.ptrVal = CI.getOperand(0); Info.align.reset(); getCoopAtomicOperandsInfo(CI, /*IsLoad=*/true, Info); - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_cooperative_atomic_store_32x4B: case Intrinsic::amdgcn_cooperative_atomic_store_16x8B: @@ -1597,7 +1609,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.ptrVal = CI.getArgOperand(0); Info.align.reset(); getCoopAtomicOperandsInfo(CI, /*IsLoad=*/false, Info); - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_ds_gws_init: case Intrinsic::amdgcn_ds_gws_barrier: @@ -1622,7 +1635,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.flags |= MachineMemOperand::MOLoad; else Info.flags |= MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_global_load_async_to_lds_b8: case Intrinsic::amdgcn_global_load_async_to_lds_b32: @@ -1636,7 +1650,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = EVT::getIntegerVT(CI.getContext(), getIntrMemWidth(IntrID)); Info.ptrVal = CI.getArgOperand(1); Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_global_store_async_from_lds_b8: case Intrinsic::amdgcn_global_store_async_from_lds_b32: @@ -1646,7 +1661,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = EVT::getIntegerVT(CI.getContext(), getIntrMemWidth(IntrID)); Info.ptrVal = CI.getArgOperand(0); Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_load_to_lds: case Intrinsic::amdgcn_global_load_lds: { @@ -1658,7 +1674,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, auto *Aux = cast(CI.getArgOperand(CI.arg_size() - 1)); if (Aux->getZExtValue() & AMDGPU::CPol::VOLATILE) Info.flags |= MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_ds_bvh_stack_rtn: case Intrinsic::amdgcn_ds_bvh_stack_push4_pop1_rtn: @@ -1678,7 +1695,8 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(4); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::amdgcn_s_prefetch_data: case Intrinsic::amdgcn_flat_prefetch: @@ -1687,10 +1705,11 @@ bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = EVT::getIntegerVT(CI.getContext(), 8); Info.ptrVal = CI.getArgOperand(0); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } default: - return false; + return; } } diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.h b/llvm/lib/Target/AMDGPU/SIISelLowering.h index d56e5ea1f9685..87d3b5bfee150 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.h +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.h @@ -336,7 +336,7 @@ class SITargetLowering final : public AMDGPUTargetLowering { MVT getPointerTy(const DataLayout &DL, unsigned AS) const override; MVT getPointerMemTy(const DataLayout &DL, unsigned AS) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &, const CallBase &, + void getTgtMemIntrinsic(SmallVectorImpl &, const CallBase &, MachineFunction &MF, unsigned IntrinsicID) const override; diff --git a/llvm/lib/Target/ARM/ARMISelLowering.cpp b/llvm/lib/Target/ARM/ARMISelLowering.cpp index a25c59bff80af..522da3374f3e3 100644 --- a/llvm/lib/Target/ARM/ARMISelLowering.cpp +++ b/llvm/lib/Target/ARM/ARMISelLowering.cpp @@ -20813,10 +20813,10 @@ bool ARMTargetLowering::isFPImmLegal(const APFloat &Imm, EVT VT, /// getTgtMemIntrinsic - Represent NEON load and store intrinsics as /// MemIntrinsicNodes. The associated MachineMemOperands record the alignment /// specified in the intrinsic calls. -bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void ARMTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; switch (Intrinsic) { case Intrinsic::arm_neon_vld1: case Intrinsic::arm_neon_vld2: @@ -20839,7 +20839,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = cast(AlignArg)->getMaybeAlignValue(); // volatile loads with NEON intrinsics not supported Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_neon_vld1x2: case Intrinsic::arm_neon_vld1x3: @@ -20854,7 +20855,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = I.getParamAlign(I.arg_size() - 1).valueOrOne(); // volatile loads with NEON intrinsics not supported Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_neon_vst1: case Intrinsic::arm_neon_vst2: @@ -20880,7 +20882,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = cast(AlignArg)->getMaybeAlignValue(); // volatile stores with NEON intrinsics not supported Info.flags = MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_neon_vst1x2: case Intrinsic::arm_neon_vst1x3: @@ -20901,7 +20904,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = I.getParamAlign(0).valueOrOne(); // volatile stores with NEON intrinsics not supported Info.flags = MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vld2q: case Intrinsic::arm_mve_vld4q: { @@ -20915,7 +20919,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(VecTy->getScalarSizeInBits() / 8); // volatile loads with MVE intrinsics not supported Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vst2q: case Intrinsic::arm_mve_vst4q: { @@ -20929,7 +20934,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(VecTy->getScalarSizeInBits() / 8); // volatile stores with MVE intrinsics not supported Info.flags = MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vldr_gather_base: case Intrinsic::arm_mve_vldr_gather_base_predicated: { @@ -20938,7 +20944,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::getVT(I.getType()); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vldr_gather_base_wb: case Intrinsic::arm_mve_vldr_gather_base_wb_predicated: { @@ -20947,7 +20954,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::getVT(I.getType()->getContainedType(0)); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vldr_gather_offset: case Intrinsic::arm_mve_vldr_gather_offset_predicated: { @@ -20959,7 +20967,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, DataVT.getVectorNumElements()); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vstr_scatter_base: case Intrinsic::arm_mve_vstr_scatter_base_predicated: { @@ -20968,7 +20977,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::getVT(I.getArgOperand(2)->getType()); Info.align = Align(1); Info.flags |= MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vstr_scatter_base_wb: case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated: { @@ -20977,7 +20987,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::getVT(I.getArgOperand(2)->getType()); Info.align = Align(1); Info.flags |= MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_mve_vstr_scatter_offset: case Intrinsic::arm_mve_vstr_scatter_offset_predicated: { @@ -20989,7 +21000,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, DataVT.getVectorNumElements()); Info.align = Align(1); Info.flags |= MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_ldaex: case Intrinsic::arm_ldrex: { @@ -21001,7 +21013,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = DL.getABITypeAlign(ValTy); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_stlex: case Intrinsic::arm_strex: { @@ -21013,7 +21026,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = DL.getABITypeAlign(ValTy); Info.flags = MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::arm_stlexd: case Intrinsic::arm_strexd: @@ -21023,7 +21037,8 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(8); Info.flags = MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; case Intrinsic::arm_ldaexd: case Intrinsic::arm_ldrexd: @@ -21033,13 +21048,12 @@ bool ARMTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(8); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; default: break; } - - return false; } /// Returns true if it is beneficial to convert a load of a constant diff --git a/llvm/lib/Target/ARM/ARMISelLowering.h b/llvm/lib/Target/ARM/ARMISelLowering.h index aec56cb69e01f..6e7d6d3acfe9f 100644 --- a/llvm/lib/Target/ARM/ARMISelLowering.h +++ b/llvm/lib/Target/ARM/ARMISelLowering.h @@ -317,8 +317,8 @@ class VectorType; bool isFPImmLegal(const APFloat &Imm, EVT VT, bool ForCodeSize = false) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; /// Returns true if it is beneficial to convert a load of a constant diff --git a/llvm/lib/Target/Hexagon/HexagonISelLowering.cpp b/llvm/lib/Target/Hexagon/HexagonISelLowering.cpp index 9557f31957ded..0ceb095cdc653 100644 --- a/llvm/lib/Target/Hexagon/HexagonISelLowering.cpp +++ b/llvm/lib/Target/Hexagon/HexagonISelLowering.cpp @@ -2026,13 +2026,12 @@ static Value *getUnderLyingObjectForBrevLdIntr(Value *V) { } /// Given an intrinsic, checks if on the target the intrinsic will need to map -/// to a MemIntrinsicNode (touches memory). If this is the case, it returns -/// true and store the intrinsic information into the IntrinsicInfo that was -/// passed to the function. -bool HexagonTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +/// to a MemIntrinsicNode (touches memory). If this is the case, it stores +/// the intrinsic information into the Infos vector. +void HexagonTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; switch (Intrinsic) { case Intrinsic::hexagon_L2_loadrd_pbr: case Intrinsic::hexagon_L2_loadri_pbr: @@ -2055,7 +2054,8 @@ bool HexagonTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = DL.getABITypeAlign(Info.memVT.getTypeForEVT(Cont)); Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::hexagon_V6_vgathermw: case Intrinsic::hexagon_V6_vgathermw_128B: @@ -2079,15 +2079,14 @@ bool HexagonTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = MaybeAlign(M.getDataLayout().getTypeAllocSizeInBits(VecTy) / 8); - Info.flags = MachineMemOperand::MOLoad | - MachineMemOperand::MOStore | + Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } default: break; } - return false; } bool HexagonTargetLowering::hasBitTest(SDValue X, SDValue Y) const { diff --git a/llvm/lib/Target/Hexagon/HexagonISelLowering.h b/llvm/lib/Target/Hexagon/HexagonISelLowering.h index d576de4049e6b..f882fe03d465a 100644 --- a/llvm/lib/Target/Hexagon/HexagonISelLowering.h +++ b/llvm/lib/Target/Hexagon/HexagonISelLowering.h @@ -49,8 +49,8 @@ class HexagonTargetLowering : public TargetLowering { const SmallVectorImpl &OutVals, const SmallVectorImpl &Ins, SelectionDAG& DAG) const; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; bool isTruncateFree(Type *Ty1, Type *Ty2) const override; diff --git a/llvm/lib/Target/LoongArch/LoongArchISelLowering.cpp b/llvm/lib/Target/LoongArch/LoongArchISelLowering.cpp index 486500e5f3f2d..4f603baab573f 100644 --- a/llvm/lib/Target/LoongArch/LoongArchISelLowering.cpp +++ b/llvm/lib/Target/LoongArch/LoongArchISelLowering.cpp @@ -9037,17 +9037,17 @@ bool LoongArchTargetLowering::hasAndNot(SDValue Y) const { return VT.isScalarInteger() && !isa(Y); } -bool LoongArchTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void LoongArchTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { switch (Intrinsic) { default: - return false; + return; case Intrinsic::loongarch_masked_atomicrmw_xchg_i32: case Intrinsic::loongarch_masked_atomicrmw_add_i32: case Intrinsic::loongarch_masked_atomicrmw_sub_i32: - case Intrinsic::loongarch_masked_atomicrmw_nand_i32: + case Intrinsic::loongarch_masked_atomicrmw_nand_i32: { + IntrinsicInfo Info; Info.opc = ISD::INTRINSIC_W_CHAIN; Info.memVT = MVT::i32; Info.ptrVal = I.getArgOperand(0); @@ -9055,9 +9055,11 @@ bool LoongArchTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(4); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; // TODO: Add more Intrinsics later. } + } } // When -mlamcas is enabled, MinCmpXchgSizeInBits will be set to 8, diff --git a/llvm/lib/Target/LoongArch/LoongArchISelLowering.h b/llvm/lib/Target/LoongArch/LoongArchISelLowering.h index 126ea055829eb..1fd9b6b237fe5 100644 --- a/llvm/lib/Target/LoongArch/LoongArchISelLowering.h +++ b/llvm/lib/Target/LoongArch/LoongArchISelLowering.h @@ -78,8 +78,8 @@ class LoongArchTargetLowering : public TargetLowering { Value *NewVal, Value *Mask, AtomicOrdering Ord) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; bool isFMAFasterThanFMulAndFAdd(const MachineFunction &MF, diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 1be35a1c67457..02b2b217aff51 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -4247,13 +4247,13 @@ void NVPTXTargetLowering::LowerAsmOperandForConstraint( // because we need the information that is only available in the "Value" type // of destination // pointer. In particular, the address space information. -bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void NVPTXTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; switch (Intrinsic) { default: - return false; + return; case Intrinsic::nvvm_match_all_sync_i32p: case Intrinsic::nvvm_match_all_sync_i64p: Info.opc = ISD::INTRINSIC_W_CHAIN; @@ -4264,7 +4264,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, // Our result depends on both our and other thread's arguments. Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col: case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row: case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col_stride: @@ -4295,7 +4296,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col: case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col_stride: @@ -4327,7 +4329,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(8); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col: @@ -4376,7 +4379,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col: @@ -4418,7 +4422,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(4); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col: @@ -4439,7 +4444,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col: @@ -4464,7 +4470,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col: @@ -4495,7 +4502,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col: @@ -4519,7 +4527,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(8); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col: @@ -4537,7 +4546,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(8); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col: @@ -4550,7 +4560,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col: @@ -4571,7 +4582,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col: @@ -4596,7 +4608,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col: @@ -4617,7 +4630,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col: @@ -4637,7 +4651,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align = Align(8); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col: @@ -4650,7 +4665,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x1_b16: @@ -4662,7 +4678,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align = Align(4); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x4_b16: @@ -4674,7 +4691,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_atomic_add_gen_f_cta: @@ -4706,7 +4724,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_prefetch_tensormap: { @@ -4718,7 +4737,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tensormap_replace_global_address: @@ -4729,7 +4749,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tensormap_replace_rank: @@ -4747,7 +4768,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_ldu_global_i: @@ -4760,7 +4782,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.flags = MachineMemOperand::MOLoad; Info.align = cast(I.getArgOperand(1))->getMaybeAlignValue(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tex_1d_v4f32_s32: case Intrinsic::nvvm_tex_1d_v4f32_f32: @@ -4826,7 +4849,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; case Intrinsic::nvvm_tex_1d_v4s32_s32: case Intrinsic::nvvm_tex_1d_v4s32_f32: @@ -4950,7 +4974,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; case Intrinsic::nvvm_suld_1d_i8_clamp: case Intrinsic::nvvm_suld_1d_v2i8_clamp: @@ -5003,7 +5028,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; case Intrinsic::nvvm_suld_1d_i16_clamp: case Intrinsic::nvvm_suld_1d_v2i16_clamp: @@ -5056,7 +5082,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; case Intrinsic::nvvm_suld_1d_i32_clamp: case Intrinsic::nvvm_suld_1d_v2i32_clamp: @@ -5109,7 +5136,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; case Intrinsic::nvvm_suld_1d_i64_clamp: case Intrinsic::nvvm_suld_1d_v2i64_clamp: @@ -5147,7 +5175,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; case Intrinsic::nvvm_tcgen05_ld_16x64b_x1: case Intrinsic::nvvm_tcgen05_ld_32x32b_x1: @@ -5158,7 +5187,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_16x64b_x2: @@ -5173,7 +5203,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x2_f32: @@ -5184,7 +5215,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_16x64b_x4: @@ -5200,7 +5232,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x4_f32: @@ -5211,7 +5244,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_16x64b_x8: @@ -5227,7 +5261,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x8_f32: @@ -5238,7 +5273,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_16x64b_x16: @@ -5254,7 +5290,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x16_f32: @@ -5265,7 +5302,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_16x64b_x32: @@ -5281,7 +5319,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x32_f32: @@ -5292,7 +5331,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_16x64b_x64: @@ -5308,7 +5348,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x64_f32: @@ -5319,7 +5360,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_16x64b_x128: @@ -5335,7 +5377,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x128_f32: @@ -5346,7 +5389,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x1: @@ -5358,7 +5402,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x2: @@ -5371,7 +5416,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x4: @@ -5385,7 +5431,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x8: @@ -5399,7 +5446,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x16: @@ -5413,7 +5461,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x32: @@ -5427,7 +5476,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x64: @@ -5441,7 +5491,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_st_16x64b_x128: @@ -5455,7 +5506,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOStore; Info.align.reset(); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1: case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1: @@ -5478,7 +5530,8 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2: @@ -5502,10 +5555,10 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore; Info.align = Align(16); - return true; + Infos.push_back(Info); + return; } } - return false; } /// getFunctionParamOptimizedAlign - since function arguments are passed via diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h index 20d49f7a6b252..9f35fe1e866fa 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h @@ -32,8 +32,8 @@ class NVPTXTargetLowering : public TargetLowering { const NVPTXSubtarget &STI); SDValue LowerOperation(SDValue Op, SelectionDAG &DAG) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; Align getFunctionArgumentAlignment(const Function *F, Type *Ty, unsigned Idx, diff --git a/llvm/lib/Target/PowerPC/PPCISelLowering.cpp b/llvm/lib/Target/PowerPC/PPCISelLowering.cpp index 5451b5fe9e36a..6c87d1cc9e920 100644 --- a/llvm/lib/Target/PowerPC/PPCISelLowering.cpp +++ b/llvm/lib/Target/PowerPC/PPCISelLowering.cpp @@ -18631,10 +18631,10 @@ PPCTargetLowering::isOffsetFoldingLegal(const GlobalAddressSDNode *GA) const { return false; } -bool PPCTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void PPCTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; switch (Intrinsic) { case Intrinsic::ppc_atomicrmw_xchg_i128: case Intrinsic::ppc_atomicrmw_add_i128: @@ -18651,7 +18651,8 @@ bool PPCTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(16); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; case Intrinsic::ppc_atomic_load_i128: Info.opc = ISD::INTRINSIC_W_CHAIN; Info.memVT = MVT::i128; @@ -18659,7 +18660,8 @@ bool PPCTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(16); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; case Intrinsic::ppc_atomic_store_i128: Info.opc = ISD::INTRINSIC_VOID; Info.memVT = MVT::i128; @@ -18667,7 +18669,8 @@ bool PPCTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(16); Info.flags = MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; case Intrinsic::ppc_altivec_lvx: case Intrinsic::ppc_altivec_lvxl: case Intrinsic::ppc_altivec_lvebx: @@ -18706,7 +18709,8 @@ bool PPCTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.size = 2*VT.getStoreSize()-1; Info.align = Align(1); Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; } case Intrinsic::ppc_altivec_stvx: case Intrinsic::ppc_altivec_stvxl: @@ -18746,7 +18750,8 @@ bool PPCTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.size = 2*VT.getStoreSize()-1; Info.align = Align(1); Info.flags = MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } case Intrinsic::ppc_stdcx: case Intrinsic::ppc_stwcx: @@ -18777,13 +18782,12 @@ bool PPCTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Alignment; Info.flags = MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } default: break; } - - return false; } /// It returns EVT::Other if the type should be determined using generic diff --git a/llvm/lib/Target/PowerPC/PPCISelLowering.h b/llvm/lib/Target/PowerPC/PPCISelLowering.h index 52e79469c78da..c74f6a6db8a3c 100644 --- a/llvm/lib/Target/PowerPC/PPCISelLowering.h +++ b/llvm/lib/Target/PowerPC/PPCISelLowering.h @@ -492,8 +492,8 @@ namespace llvm { bool isOffsetFoldingLegal(const GlobalAddressSDNode *GA) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; /// It returns EVT::Other if the type should be determined using generic diff --git a/llvm/lib/Target/RISCV/RISCVISelLowering.cpp b/llvm/lib/Target/RISCV/RISCVISelLowering.cpp index 90809828e9653..4e46a334bd5c9 100644 --- a/llvm/lib/Target/RISCV/RISCVISelLowering.cpp +++ b/llvm/lib/Target/RISCV/RISCVISelLowering.cpp @@ -1953,10 +1953,10 @@ bool RISCVTargetLowering::shouldExpandCttzElements(EVT VT) const { VT.getVectorElementType() != MVT::i1 || !isTypeLegal(VT); } -bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void RISCVTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; auto &DL = I.getDataLayout(); auto SetRVVLoadStoreInfo = [&](unsigned PtrOp, bool IsStore, @@ -1997,7 +1997,7 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.size = MemoryLocation::UnknownSize; Info.flags |= IsStore ? MachineMemOperand::MOStore : MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); }; if (I.hasMetadata(LLVMContext::MD_nontemporal)) @@ -2006,7 +2006,7 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.flags |= RISCVTargetLowering::getTargetMMOFlags(I); switch (Intrinsic) { default: - return false; + return; case Intrinsic::riscv_masked_atomicrmw_xchg: case Intrinsic::riscv_masked_atomicrmw_add: case Intrinsic::riscv_masked_atomicrmw_sub: @@ -2028,7 +2028,8 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(4); Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; case Intrinsic::riscv_seg2_load_mask: case Intrinsic::riscv_seg3_load_mask: case Intrinsic::riscv_seg4_load_mask: @@ -2043,8 +2044,9 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_sseg6_load_mask: case Intrinsic::riscv_sseg7_load_mask: case Intrinsic::riscv_sseg8_load_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ 0, /*IsStore*/ false, - /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ 0, /*IsStore*/ false, + /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + return; case Intrinsic::riscv_seg2_store_mask: case Intrinsic::riscv_seg3_store_mask: case Intrinsic::riscv_seg4_store_mask: @@ -2053,9 +2055,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_seg7_store_mask: case Intrinsic::riscv_seg8_store_mask: // Operands are (vec, ..., vec, ptr, mask, vl) - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 3, - /*IsStore*/ true, - /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 3, + /*IsStore*/ true, + /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + return; case Intrinsic::riscv_sseg2_store_mask: case Intrinsic::riscv_sseg3_store_mask: case Intrinsic::riscv_sseg4_store_mask: @@ -2064,47 +2067,53 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_sseg7_store_mask: case Intrinsic::riscv_sseg8_store_mask: // Operands are (vec, ..., vec, ptr, offset, mask, vl) - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, - /*IsStore*/ true, - /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, + /*IsStore*/ true, + /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + return; case Intrinsic::riscv_vlm: - return SetRVVLoadStoreInfo(/*PtrOp*/ 0, - /*IsStore*/ false, - /*IsUnitStrided*/ true, - /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ 0, + /*IsStore*/ false, + /*IsUnitStrided*/ true, + /*UsePtrVal*/ true); + return; case Intrinsic::riscv_vle: case Intrinsic::riscv_vle_mask: case Intrinsic::riscv_vleff: case Intrinsic::riscv_vleff_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ 1, - /*IsStore*/ false, - /*IsUnitStrided*/ true, - /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ 1, + /*IsStore*/ false, + /*IsUnitStrided*/ true, + /*UsePtrVal*/ true); + return; case Intrinsic::riscv_vsm: case Intrinsic::riscv_vse: case Intrinsic::riscv_vse_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ 1, - /*IsStore*/ true, - /*IsUnitStrided*/ true, - /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ 1, + /*IsStore*/ true, + /*IsUnitStrided*/ true, + /*UsePtrVal*/ true); + return; case Intrinsic::riscv_vlse: case Intrinsic::riscv_vlse_mask: case Intrinsic::riscv_vloxei: case Intrinsic::riscv_vloxei_mask: case Intrinsic::riscv_vluxei: case Intrinsic::riscv_vluxei_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ 1, - /*IsStore*/ false, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ 1, + /*IsStore*/ false, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_vsse: case Intrinsic::riscv_vsse_mask: case Intrinsic::riscv_vsoxei: case Intrinsic::riscv_vsoxei_mask: case Intrinsic::riscv_vsuxei: case Intrinsic::riscv_vsuxei_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ 1, - /*IsStore*/ true, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ 1, + /*IsStore*/ true, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_vlseg2: case Intrinsic::riscv_vlseg3: case Intrinsic::riscv_vlseg4: @@ -2119,9 +2128,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vlseg6ff: case Intrinsic::riscv_vlseg7ff: case Intrinsic::riscv_vlseg8ff: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 3, - /*IsStore*/ false, - /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 3, + /*IsStore*/ false, + /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + return; case Intrinsic::riscv_vlseg2_mask: case Intrinsic::riscv_vlseg3_mask: case Intrinsic::riscv_vlseg4_mask: @@ -2136,9 +2146,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vlseg6ff_mask: case Intrinsic::riscv_vlseg7ff_mask: case Intrinsic::riscv_vlseg8ff_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 5, - /*IsStore*/ false, - /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 5, + /*IsStore*/ false, + /*IsUnitStrided*/ false, /*UsePtrVal*/ true); + return; case Intrinsic::riscv_vlsseg2: case Intrinsic::riscv_vlsseg3: case Intrinsic::riscv_vlsseg4: @@ -2160,9 +2171,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vluxseg6: case Intrinsic::riscv_vluxseg7: case Intrinsic::riscv_vluxseg8: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, - /*IsStore*/ false, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, + /*IsStore*/ false, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_vlsseg2_mask: case Intrinsic::riscv_vlsseg3_mask: case Intrinsic::riscv_vlsseg4_mask: @@ -2184,9 +2196,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vluxseg6_mask: case Intrinsic::riscv_vluxseg7_mask: case Intrinsic::riscv_vluxseg8_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 6, - /*IsStore*/ false, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 6, + /*IsStore*/ false, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_vsseg2: case Intrinsic::riscv_vsseg3: case Intrinsic::riscv_vsseg4: @@ -2194,9 +2207,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vsseg6: case Intrinsic::riscv_vsseg7: case Intrinsic::riscv_vsseg8: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 3, - /*IsStore*/ true, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 3, + /*IsStore*/ true, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_vsseg2_mask: case Intrinsic::riscv_vsseg3_mask: case Intrinsic::riscv_vsseg4_mask: @@ -2204,9 +2218,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vsseg6_mask: case Intrinsic::riscv_vsseg7_mask: case Intrinsic::riscv_vsseg8_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, - /*IsStore*/ true, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, + /*IsStore*/ true, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_vssseg2: case Intrinsic::riscv_vssseg3: case Intrinsic::riscv_vssseg4: @@ -2228,9 +2243,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vsuxseg6: case Intrinsic::riscv_vsuxseg7: case Intrinsic::riscv_vsuxseg8: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, - /*IsStore*/ true, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 4, + /*IsStore*/ true, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_vssseg2_mask: case Intrinsic::riscv_vssseg3_mask: case Intrinsic::riscv_vssseg4_mask: @@ -2252,9 +2268,10 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, case Intrinsic::riscv_vsuxseg6_mask: case Intrinsic::riscv_vsuxseg7_mask: case Intrinsic::riscv_vsuxseg8_mask: - return SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 5, - /*IsStore*/ true, - /*IsUnitStrided*/ false); + SetRVVLoadStoreInfo(/*PtrOp*/ I.arg_size() - 5, + /*IsStore*/ true, + /*IsUnitStrided*/ false); + return; case Intrinsic::riscv_sf_vlte8: case Intrinsic::riscv_sf_vlte16: case Intrinsic::riscv_sf_vlte32: @@ -2281,7 +2298,8 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, } Info.size = MemoryLocation::UnknownSize; Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::riscv_sf_vste8: case Intrinsic::riscv_sf_vste16: case Intrinsic::riscv_sf_vste32: @@ -2308,7 +2326,8 @@ bool RISCVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, } Info.size = MemoryLocation::UnknownSize; Info.flags |= MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; } } diff --git a/llvm/lib/Target/RISCV/RISCVISelLowering.h b/llvm/lib/Target/RISCV/RISCVISelLowering.h index c203a6460992c..4a27bef2013ac 100644 --- a/llvm/lib/Target/RISCV/RISCVISelLowering.h +++ b/llvm/lib/Target/RISCV/RISCVISelLowering.h @@ -35,8 +35,8 @@ class RISCVTargetLowering : public TargetLowering { const RISCVSubtarget &getSubtarget() const { return Subtarget; } - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; bool isLegalAddressingMode(const DataLayout &DL, const AddrMode &AM, Type *Ty, unsigned AS, diff --git a/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp b/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp index 36fa5fa9a70cb..3e5ce4b90ea4a 100644 --- a/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp @@ -93,10 +93,10 @@ MVT SPIRVTargetLowering::getRegisterTypeForCallingConv(LLVMContext &Context, return getRegisterType(Context, VT); } -bool SPIRVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void SPIRVTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; unsigned AlignIdx = 3; switch (Intrinsic) { case Intrinsic::spv_load: @@ -112,13 +112,12 @@ bool SPIRVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::i64; // TODO: take into account opaque pointers (don't use getElementType). // MVT::getVT(PtrTy->getElementType()); - return true; - break; + Infos.push_back(Info); + return; } default: break; } - return false; } std::pair diff --git a/llvm/lib/Target/SPIRV/SPIRVISelLowering.h b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h index 5746832c8fd95..462605ab6fe36 100644 --- a/llvm/lib/Target/SPIRV/SPIRVISelLowering.h +++ b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h @@ -48,8 +48,8 @@ class SPIRVTargetLowering : public TargetLowering { EVT VT) const override; MVT getRegisterTypeForCallingConv(LLVMContext &Context, CallingConv::ID CC, EVT VT) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; std::pair diff --git a/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.cpp b/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.cpp index 5abf0e8f59d2a..abd5b0c0ad9d6 100644 --- a/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.cpp +++ b/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.cpp @@ -1065,10 +1065,10 @@ EVT WebAssemblyTargetLowering::getSetCCResultType(const DataLayout &DL, return EVT::getIntegerVT(C, 32); } -bool WebAssemblyTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void WebAssemblyTargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; switch (Intrinsic) { case Intrinsic::wasm_memory_atomic_notify: Info.opc = ISD::INTRINSIC_W_CHAIN; @@ -1083,7 +1083,8 @@ bool WebAssemblyTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, // instructions are treated as volatiles in the backend, so we should be // consistent. The same applies for wasm_atomic_wait intrinsics too. Info.flags = MachineMemOperand::MOVolatile | MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::wasm_memory_atomic_wait32: Info.opc = ISD::INTRINSIC_W_CHAIN; Info.memVT = MVT::i32; @@ -1091,7 +1092,8 @@ bool WebAssemblyTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(4); Info.flags = MachineMemOperand::MOVolatile | MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::wasm_memory_atomic_wait64: Info.opc = ISD::INTRINSIC_W_CHAIN; Info.memVT = MVT::i64; @@ -1099,7 +1101,8 @@ bool WebAssemblyTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(8); Info.flags = MachineMemOperand::MOVolatile | MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::wasm_loadf16_f32: Info.opc = ISD::INTRINSIC_W_CHAIN; Info.memVT = MVT::f16; @@ -1107,7 +1110,8 @@ bool WebAssemblyTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(2); Info.flags = MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::wasm_storef16_f32: Info.opc = ISD::INTRINSIC_VOID; Info.memVT = MVT::f16; @@ -1115,9 +1119,10 @@ bool WebAssemblyTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.offset = 0; Info.align = Align(2); Info.flags = MachineMemOperand::MOStore; - return true; + Infos.push_back(Info); + return; default: - return false; + return; } } diff --git a/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.h b/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.h index 204384f06ab25..b1c9db37a2b18 100644 --- a/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.h +++ b/llvm/lib/Target/WebAssembly/WebAssemblyISelLowering.h @@ -61,8 +61,8 @@ class WebAssemblyTargetLowering final : public TargetLowering { bool isOffsetFoldingLegal(const GlobalAddressSDNode *GA) const override; EVT getSetCCResultType(const DataLayout &DL, LLVMContext &Context, EVT VT) const override; - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; void computeKnownBitsForTargetNode(const SDValue Op, KnownBits &Known, diff --git a/llvm/lib/Target/X86/X86ISelLowering.cpp b/llvm/lib/Target/X86/X86ISelLowering.cpp index 15f0815fe4e06..1837c8bbedf0e 100644 --- a/llvm/lib/Target/X86/X86ISelLowering.cpp +++ b/llvm/lib/Target/X86/X86ISelLowering.cpp @@ -3164,10 +3164,10 @@ static bool useVPTERNLOG(const X86Subtarget &Subtarget, MVT VT) { VT.is512BitVector(); } -bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, - const CallBase &I, - MachineFunction &MF, - unsigned Intrinsic) const { +void X86TargetLowering::getTgtMemIntrinsic( + SmallVectorImpl &Infos, const CallBase &I, + MachineFunction &MF, unsigned Intrinsic) const { + IntrinsicInfo Info; Info.flags = MachineMemOperand::MONone; Info.offset = 0; @@ -3181,7 +3181,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = EVT::getIntegerVT(I.getType()->getContext(), 48); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::x86_aesenc256kl: case Intrinsic::x86_aesdec256kl: Info.opc = ISD::INTRINSIC_W_CHAIN; @@ -3189,7 +3190,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = EVT::getIntegerVT(I.getType()->getContext(), 64); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::x86_aesencwide128kl: case Intrinsic::x86_aesdecwide128kl: Info.opc = ISD::INTRINSIC_W_CHAIN; @@ -3197,7 +3199,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = EVT::getIntegerVT(I.getType()->getContext(), 48); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::x86_aesencwide256kl: case Intrinsic::x86_aesdecwide256kl: Info.opc = ISD::INTRINSIC_W_CHAIN; @@ -3205,7 +3208,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = EVT::getIntegerVT(I.getType()->getContext(), 64); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - return true; + Infos.push_back(Info); + return; case Intrinsic::x86_cmpccxadd32: case Intrinsic::x86_cmpccxadd64: case Intrinsic::x86_atomic_bts: @@ -3218,7 +3222,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(Size); Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::x86_atomic_bts_rm: case Intrinsic::x86_atomic_btc_rm: @@ -3230,7 +3235,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(Size); Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } case Intrinsic::x86_aadd32: case Intrinsic::x86_aadd64: @@ -3252,10 +3258,11 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.align = Align(Size); Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore | MachineMemOperand::MOVolatile; - return true; + Infos.push_back(Info); + return; } } - return false; + return; } switch (IntrData->Type) { @@ -3276,7 +3283,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::getVectorVT(ScalarVT, VT.getVectorNumElements()); Info.align = Align(1); Info.flags |= MachineMemOperand::MOStore; - break; + Infos.push_back(Info); + return; } case GATHER: case GATHER_AVX2: { @@ -3289,7 +3297,8 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::getVectorVT(DataVT.getVectorElementType(), NumElts); Info.align = Align(1); Info.flags |= MachineMemOperand::MOLoad; - break; + Infos.push_back(Info); + return; } case SCATTER: { Info.opc = ISD::INTRINSIC_VOID; @@ -3301,13 +3310,12 @@ bool X86TargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info, Info.memVT = MVT::getVectorVT(DataVT.getVectorElementType(), NumElts); Info.align = Align(1); Info.flags |= MachineMemOperand::MOStore; - break; + Infos.push_back(Info); + return; } default: - return false; + return; } - - return true; } /// Returns true if the target can instruction select the diff --git a/llvm/lib/Target/X86/X86ISelLowering.h b/llvm/lib/Target/X86/X86ISelLowering.h index 30faca54e13f6..1e8010c1a339a 100644 --- a/llvm/lib/Target/X86/X86ISelLowering.h +++ b/llvm/lib/Target/X86/X86ISelLowering.h @@ -1503,12 +1503,12 @@ namespace llvm { unsigned SelectOpcode, SDValue X, SDValue Y) const override; - /// Given an intrinsic, checks if on the target the intrinsic will need to map - /// to a MemIntrinsicNode (touches memory). If this is the case, it returns - /// true and stores the intrinsic information into the IntrinsicInfo that was - /// passed to the function. - bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallBase &I, - MachineFunction &MF, + /// Given an intrinsic, checks if on the target the intrinsic will need to + /// map to a MemIntrinsicNode (touches memory). If this is the case, it + /// returns true and stores the intrinsic information into the IntrinsicInfo + /// that was passed to the function. + void getTgtMemIntrinsic(SmallVectorImpl &Infos, + const CallBase &I, MachineFunction &MF, unsigned Intrinsic) const override; /// Returns true if the target can instruction select the