diff options
Diffstat (limited to 'llvm/lib/Target/NVPTX')
22 files changed, 68 insertions, 40 deletions
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp index b6eefe206268a..82d332ab3f08d 100644 --- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp +++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp @@ -72,9 +72,10 @@ void NVPTXInstPrinter::printRegName(raw_ostream &OS, unsigned RegNo) const { OS << VReg; } -void NVPTXInstPrinter::printInst(const MCInst *MI, raw_ostream &OS, - StringRef Annot, const MCSubtargetInfo &STI) { - printInstruction(MI, OS); +void NVPTXInstPrinter::printInst(const MCInst *MI, uint64_t Address, + StringRef Annot, const MCSubtargetInfo &STI, + raw_ostream &OS) { + printInstruction(MI, Address, OS); // Next always print the annotation. printAnnotation(OS, Annot); diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h index c38472925a29d..cee0e7eec54af 100644 --- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h +++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h @@ -25,11 +25,11 @@ public: const MCRegisterInfo &MRI); void printRegName(raw_ostream &OS, unsigned RegNo) const override; - void printInst(const MCInst *MI, raw_ostream &OS, StringRef Annot, - const MCSubtargetInfo &STI) override; + void printInst(const MCInst *MI, uint64_t Address, StringRef Annot, + const MCSubtargetInfo &STI, raw_ostream &OS) override; // Autogenerated by tblgen. - void printInstruction(const MCInst *MI, raw_ostream &O); + void printInstruction(const MCInst *MI, uint64_t Address, raw_ostream &O); static const char *getRegisterName(unsigned RegNo); // End diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp index 556745825a154..7e1da9b7a94b9 100644 --- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp +++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp @@ -17,7 +17,8 @@ using namespace llvm; void NVPTXMCAsmInfo::anchor() {} -NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple) { +NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple, + const MCTargetOptions &Options) { if (TheTriple.getArch() == Triple::nvptx64) { CodePointerSize = CalleeSaveStackSlotSize = 8; } diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h index e888526da8987..ce5ca99c53970 100644 --- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h +++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h @@ -23,7 +23,8 @@ class NVPTXMCAsmInfo : public MCAsmInfo { virtual void anchor(); public: - explicit NVPTXMCAsmInfo(const Triple &TheTriple); + explicit NVPTXMCAsmInfo(const Triple &TheTriple, + const MCTargetOptions &Options); /// Return true if the .section directive should be omitted when /// emitting \p SectionName. For example: diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp index c8b85b2718a65..d758c2c869591 100644 --- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp +++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp @@ -66,7 +66,7 @@ static MCTargetStreamer *createTargetAsmStreamer(MCStreamer &S, } // Force static initialization. -extern "C" void LLVMInitializeNVPTXTargetMC() { +extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXTargetMC() { for (Target *T : {&getTheNVPTXTarget32(), &getTheNVPTXTarget64()}) { // Register the MC asm info. RegisterMCAsmInfo<NVPTXMCAsmInfo> X(*T); diff --git a/llvm/lib/Target/NVPTX/ManagedStringPool.h b/llvm/lib/Target/NVPTX/ManagedStringPool.h index bbcbb45980404..fb838d1fcd7fa 100644 --- a/llvm/lib/Target/NVPTX/ManagedStringPool.h +++ b/llvm/lib/Target/NVPTX/ManagedStringPool.h @@ -32,7 +32,7 @@ public: SmallVectorImpl<std::string *>::iterator Current = Pool.begin(); while (Current != Pool.end()) { delete *Current; - Current++; + ++Current; } } diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index 307f4d58c3ab2..7117438dc503f 100644 --- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -2258,7 +2258,7 @@ void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum, } // Force static initialization. -extern "C" void LLVMInitializeNVPTXAsmPrinter() { +extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXAsmPrinter() { RegisterAsmPrinter<NVPTXAsmPrinter> X(getTheNVPTXTarget32()); RegisterAsmPrinter<NVPTXAsmPrinter> Y(getTheNVPTXTarget64()); } diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp index 3d2447d75c775..4296eca6a8dfd 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp @@ -11,11 +11,12 @@ //===----------------------------------------------------------------------===// #include "NVPTXISelDAGToDAG.h" -#include "NVPTXUtilities.h" #include "MCTargetDesc/NVPTXBaseInfo.h" +#include "NVPTXUtilities.h" #include "llvm/Analysis/ValueTracking.h" #include "llvm/IR/GlobalValue.h" #include "llvm/IR/Instructions.h" +#include "llvm/IR/IntrinsicsNVPTX.h" #include "llvm/Support/AtomicOrdering.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" @@ -611,7 +612,7 @@ bool NVPTXDAGToDAGISel::tryEXTRACT_VECTOR_ELEMENT(SDNode *N) { // Find and record all uses of this vector that extract element 0 or 1. SmallVector<SDNode *, 4> E0, E1; - for (const auto &U : Vector.getNode()->uses()) { + for (auto U : Vector.getNode()->uses()) { if (U->getOpcode() != ISD::EXTRACT_VECTOR_ELT) continue; if (U->getOperand(0) != Vector) @@ -885,7 +886,7 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) { MVT SimpleVT = LoadedVT.getSimpleVT(); MVT ScalarVT = SimpleVT.getScalarType(); // Read at least 8 bits (predicates are stored as 8-bit values) - unsigned fromTypeWidth = std::max(8U, ScalarVT.getSizeInBits()); + unsigned fromTypeWidth = std::max(8U, (unsigned)ScalarVT.getSizeInBits()); unsigned int fromType; // Vector Setting @@ -1030,7 +1031,7 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) { // Float : ISD::NON_EXTLOAD or ISD::EXTLOAD and the type is float MVT ScalarVT = SimpleVT.getScalarType(); // Read at least 8 bits (predicates are stored as 8-bit values) - unsigned FromTypeWidth = std::max(8U, ScalarVT.getSizeInBits()); + unsigned FromTypeWidth = std::max(8U, (unsigned)ScalarVT.getSizeInBits()); unsigned int FromType; // The last operand holds the original LoadSDNode::getExtensionType() value unsigned ExtensionType = cast<ConstantSDNode>( diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h index e4e5069b7a807..ce5136016f962 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h +++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h @@ -46,7 +46,7 @@ public: return "NVPTX DAG->DAG Pattern Instruction Selection"; } bool runOnMachineFunction(MachineFunction &MF) override; - const NVPTXSubtarget *Subtarget; + const NVPTXSubtarget *Subtarget = nullptr; bool SelectInlineAsmMemoryOperand(const SDValue &Op, unsigned ConstraintID, diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 9acd0bea66fdf..92f71c687c461 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -39,6 +39,7 @@ #include "llvm/IR/GlobalValue.h" #include "llvm/IR/Instruction.h" #include "llvm/IR/Instructions.h" +#include "llvm/IR/IntrinsicsNVPTX.h" #include "llvm/IR/Module.h" #include "llvm/IR/Type.h" #include "llvm/IR/Value.h" diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h index ef645fc1e541f..546fe49808e2d 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h @@ -538,7 +538,10 @@ public: bool allowFMA(MachineFunction &MF, CodeGenOpt::Level OptLevel) const; bool allowUnsafeFPMath(MachineFunction &MF) const; - bool isFMAFasterThanFMulAndFAdd(EVT) const override { return true; } + bool isFMAFasterThanFMulAndFAdd(const MachineFunction &MF, + EVT) const override { + return true; + } bool enableAggressiveFMAFusion(EVT VT) const override { return true; } diff --git a/llvm/lib/Target/NVPTX/NVPTXImageOptimizer.cpp b/llvm/lib/Target/NVPTX/NVPTXImageOptimizer.cpp index 74ab2f7b8453d..a846c2fada269 100644 --- a/llvm/lib/Target/NVPTX/NVPTXImageOptimizer.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXImageOptimizer.cpp @@ -18,6 +18,7 @@ #include "llvm/Analysis/ConstantFolding.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/Intrinsics.h" +#include "llvm/IR/IntrinsicsNVPTX.h" #include "llvm/IR/Module.h" #include "llvm/Pass.h" diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp index f928b44c91e0b..cff230289e600 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp @@ -31,8 +31,8 @@ NVPTXInstrInfo::NVPTXInstrInfo() : NVPTXGenInstrInfo(), RegInfo() {} void NVPTXInstrInfo::copyPhysReg(MachineBasicBlock &MBB, MachineBasicBlock::iterator I, - const DebugLoc &DL, unsigned DestReg, - unsigned SrcReg, bool KillSrc) const { + const DebugLoc &DL, MCRegister DestReg, + MCRegister SrcReg, bool KillSrc) const { const MachineRegisterInfo &MRI = MBB.getParent()->getRegInfo(); const TargetRegisterClass *DestRC = MRI.getRegClass(DestReg); const TargetRegisterClass *SrcRC = MRI.getRegClass(SrcReg); diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h index 7c0912808f7bd..40f6adc964d7e 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h @@ -49,7 +49,7 @@ public: */ void copyPhysReg(MachineBasicBlock &MBB, MachineBasicBlock::iterator I, - const DebugLoc &DL, unsigned DestReg, unsigned SrcReg, + const DebugLoc &DL, MCRegister DestReg, MCRegister SrcReg, bool KillSrc) const override; // Branch analysis. diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index c52195fb04495..76a4a1d4030a9 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -7400,7 +7400,9 @@ class WMMA_REGINFO<WMMA_REGS r> // u4/s4/b1 -> s32 @ m8n8k32 (u4/s4), m8n8k128(b1) !or(!eq(geom,"m8n8k128"), - !eq(geom,"m8n8k32")) : [hasSM75, hasPTX63]); + !eq(geom,"m8n8k32")) : [hasSM75, hasPTX63], + + !eq(geom, "m8n8k4") : [hasSM70, hasPTX64]); // template DAGs for instruction inputs/output. dag Outs = !dag(outs, ptx_regs, reg_names); @@ -7546,25 +7548,37 @@ class WMMA_MMA<WMMA_REGINFO FragA, WMMA_REGINFO FragB, let OutOperandList = FragD.Outs; let InOperandList = !con(Args, (ins MmaCode:$ptx)); string TypeList = !cond( + !eq(FragD.geom, "m8n8k4") : "." # FragD.ptx_elt_type + # ".f16.f16." + # FragC.ptx_elt_type, !eq(FragD.ptx_elt_type, "s32") : ".s32" # "." # FragA.ptx_elt_type # "." # FragB.ptx_elt_type # ".s32", 1: "." # FragD.ptx_elt_type # "." # FragC.ptx_elt_type, ); - let AsmString = "wmma.mma" - # !if(!eq(FragA.ptx_elt_type, "b1"), ".xor.popc", "") - # ".sync" - # "${ptx:aligned}" - # "." # ALayout - # "." # BLayout - # "." # FragA.geom - # TypeList - # !if(Satfinite, ".satfinite", "") # "\n\t\t" - # FragD.regstring # ",\n\t\t" - # FragA.regstring # ",\n\t\t" - # FragB.regstring # ",\n\t\t" - # FragC.regstring # ";"; + let AsmString = !if(!eq(FragA.geom, "m8n8k4"), + "mma.sync.aligned.m8n8k4" + # "." # ALayout + # "." # BLayout + # TypeList # "\n\t\t" + # FragD.regstring # ",\n\t\t" + # FragA.regstring # ",\n\t\t" + # FragB.regstring # ",\n\t\t" + # FragC.regstring # ";", + "wmma.mma" + # !if(!eq(FragA.ptx_elt_type, "b1"), ".xor.popc", "") + # ".sync" + # "${ptx:aligned}" + # "." # ALayout + # "." # BLayout + # "." # FragA.geom + # TypeList + # !if(Satfinite, ".satfinite", "") # "\n\t\t" + # FragD.regstring # ",\n\t\t" + # FragA.regstring # ",\n\t\t" + # FragB.regstring # ",\n\t\t" + # FragC.regstring # ";"); } defset list<WMMA_INSTR> MMAs = { diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp index 945b7286b03cc..cf488b06f06c7 100644 --- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp @@ -70,8 +70,7 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) { for (auto &I : BB) { if (auto allocaInst = dyn_cast<AllocaInst>(&I)) { Changed = true; - auto PTy = dyn_cast<PointerType>(allocaInst->getType()); - auto ETy = PTy->getElementType(); + auto ETy = cast<PointerType>(allocaInst->getType())->getElementType(); auto LocalAddrTy = PointerType::get(ETy, ADDRESS_SPACE_LOCAL); auto NewASCToLocal = new AddrSpaceCastInst(allocaInst, LocalAddrTy, ""); auto GenericAddrTy = PointerType::get(ETy, ADDRESS_SPACE_GENERIC); diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp index f58fb5717773d..0778706d936a7 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp @@ -72,7 +72,7 @@ void initializeNVPTXProxyRegErasurePass(PassRegistry &); } // end namespace llvm -extern "C" void LLVMInitializeNVPTXTarget() { +extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXTarget() { // Register the target. RegisterTargetMachine<NVPTXTargetMachine32> X(getTheNVPTXTarget32()); RegisterTargetMachine<NVPTXTargetMachine64> Y(getTheNVPTXTarget64()); diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp index be0416f90fca3..afc40a7abed08 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp @@ -14,6 +14,7 @@ #include "llvm/CodeGen/BasicTTIImpl.h" #include "llvm/CodeGen/CostTable.h" #include "llvm/CodeGen/TargetLowering.h" +#include "llvm/IR/IntrinsicsNVPTX.h" #include "llvm/Support/Debug.h" using namespace llvm; @@ -113,7 +114,8 @@ bool NVPTXTTIImpl::isSourceOfDivergence(const Value *V) { int NVPTXTTIImpl::getArithmeticInstrCost( unsigned Opcode, Type *Ty, TTI::OperandValueKind Opd1Info, TTI::OperandValueKind Opd2Info, TTI::OperandValueProperties Opd1PropInfo, - TTI::OperandValueProperties Opd2PropInfo, ArrayRef<const Value *> Args) { + TTI::OperandValueProperties Opd2PropInfo, ArrayRef<const Value *> Args, + const Instruction *CxtI) { // Legalize the type. std::pair<int, MVT> LT = TLI->getTypeLegalizationCost(DL, Ty); diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h index b179a28fa713b..864d8b91a89a5 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h +++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h @@ -91,7 +91,8 @@ public: TTI::OperandValueKind Opd2Info = TTI::OK_AnyValue, TTI::OperandValueProperties Opd1PropInfo = TTI::OP_None, TTI::OperandValueProperties Opd2PropInfo = TTI::OP_None, - ArrayRef<const Value *> Args = ArrayRef<const Value *>()); + ArrayRef<const Value *> Args = ArrayRef<const Value *>(), + const Instruction *CxtI = nullptr); void getUnrollingPreferences(Loop *L, ScalarEvolution &SE, TTI::UnrollingPreferences &UP); diff --git a/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp b/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp index 5cf7b6691e638..baaedc7ac87cc 100644 --- a/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp +++ b/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp @@ -16,6 +16,8 @@ #include "llvm/IR/InstIterator.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/Intrinsics.h" +#include "llvm/IR/IntrinsicsNVPTX.h" +#include "llvm/Support/CommandLine.h" using namespace llvm; diff --git a/llvm/lib/Target/NVPTX/NVVMReflect.cpp b/llvm/lib/Target/NVPTX/NVVMReflect.cpp index 634a052e2ee72..ae166dc5a8d57 100644 --- a/llvm/lib/Target/NVPTX/NVVMReflect.cpp +++ b/llvm/lib/Target/NVPTX/NVVMReflect.cpp @@ -27,6 +27,7 @@ #include "llvm/IR/InstIterator.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/Intrinsics.h" +#include "llvm/IR/IntrinsicsNVPTX.h" #include "llvm/IR/Module.h" #include "llvm/IR/Type.h" #include "llvm/Pass.h" diff --git a/llvm/lib/Target/NVPTX/TargetInfo/NVPTXTargetInfo.cpp b/llvm/lib/Target/NVPTX/TargetInfo/NVPTXTargetInfo.cpp index 2c71ec58ec421..235be9c0dbbbc 100644 --- a/llvm/lib/Target/NVPTX/TargetInfo/NVPTXTargetInfo.cpp +++ b/llvm/lib/Target/NVPTX/TargetInfo/NVPTXTargetInfo.cpp @@ -19,7 +19,7 @@ Target &llvm::getTheNVPTXTarget64() { return TheNVPTXTarget64; } -extern "C" void LLVMInitializeNVPTXTargetInfo() { +extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXTargetInfo() { RegisterTarget<Triple::nvptx> X(getTheNVPTXTarget32(), "nvptx", "NVIDIA PTX 32-bit", "NVPTX"); RegisterTarget<Triple::nvptx64> Y(getTheNVPTXTarget64(), "nvptx64", |