From 6ec96839594467cd2a9282c277cf94af70fc66b2 Mon Sep 17 00:00:00 2001 From: Eric Christopher Date: Thu, 19 Feb 2015 00:08:27 +0000 Subject: [PATCH] Remove all use of is64bit off of NVPTXSubtarget and clean up code accordingly. This changes the constructors of a number of classes that don't need to know the subtarget's 64-bitness. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@229787 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 4 +- lib/Target/NVPTX/NVPTXFrameLowering.cpp | 7 +- lib/Target/NVPTX/NVPTXFrameLowering.h | 10 ++- lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 95 +++++++++++-------------- lib/Target/NVPTX/NVPTXISelDAGToDAG.h | 1 + lib/Target/NVPTX/NVPTXInstrInfo.cpp | 4 +- lib/Target/NVPTX/NVPTXInstrInfo.h | 2 +- lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 3 +- lib/Target/NVPTX/NVPTXRegisterInfo.h | 8 +-- lib/Target/NVPTX/NVPTXSubtarget.cpp | 9 ++- lib/Target/NVPTX/NVPTXSubtarget.h | 5 +- lib/Target/NVPTX/NVPTXTargetMachine.cpp | 2 +- 12 files changed, 61 insertions(+), 89 deletions(-) diff --git a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index b87aec35ccf..dd3dc7ef1cb 100644 --- a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -819,7 +819,7 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { StringRef CPU = TM.getTargetCPU(); StringRef FS = TM.getTargetFeatureString(); const NVPTXTargetMachine &NTM = static_cast(TM); - const NVPTXSubtarget STI(TT, CPU, FS, NTM, NTM.is64Bit()); + const NVPTXSubtarget STI(TT, CPU, FS, NTM); SmallString<128> Str1; raw_svector_ostream OS1(Str1); @@ -1625,7 +1625,7 @@ void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters( if (NumBytes) { O << "\t.local .align " << MFI->getMaxAlignment() << " .b8 \t" << DEPOTNAME << getFunctionNumber() << "[" << NumBytes << "];\n"; - if (nvptxSubtarget->is64Bit()) { + if (static_cast(MF.getTarget()).is64Bit()) { O << "\t.reg .b64 \t%SP;\n"; O << "\t.reg .b64 \t%SPL;\n"; } else { diff --git a/lib/Target/NVPTX/NVPTXFrameLowering.cpp b/lib/Target/NVPTX/NVPTXFrameLowering.cpp index 314df3828b8..34d3a66adce 100644 --- a/lib/Target/NVPTX/NVPTXFrameLowering.cpp +++ b/lib/Target/NVPTX/NVPTXFrameLowering.cpp @@ -26,9 +26,8 @@ using namespace llvm; -NVPTXFrameLowering::NVPTXFrameLowering(NVPTXSubtarget &STI) - : TargetFrameLowering(TargetFrameLowering::StackGrowsUp, 8, 0), - is64bit(STI.is64Bit()) {} +NVPTXFrameLowering::NVPTXFrameLowering() + : TargetFrameLowering(TargetFrameLowering::StackGrowsUp, 8, 0) {} bool NVPTXFrameLowering::hasFP(const MachineFunction &MF) const { return true; } @@ -45,7 +44,7 @@ void NVPTXFrameLowering::emitPrologue(MachineFunction &MF) const { // mov %SPL, %depot; // cvta.local %SP, %SPL; - if (is64bit) { + if (static_cast(MF.getTarget()).is64Bit()) { unsigned LocalReg = MRI.createVirtualRegister(&NVPTX::Int64RegsRegClass); MachineInstr *MI = BuildMI(MBB, MBBI, dl, MF.getSubtarget().getInstrInfo()->get( diff --git a/lib/Target/NVPTX/NVPTXFrameLowering.h b/lib/Target/NVPTX/NVPTXFrameLowering.h index 0846b78d58e..d1e0a5ceb30 100644 --- a/lib/Target/NVPTX/NVPTXFrameLowering.h +++ b/lib/Target/NVPTX/NVPTXFrameLowering.h @@ -19,18 +19,16 @@ namespace llvm { class NVPTXSubtarget; class NVPTXFrameLowering : public TargetFrameLowering { - bool is64bit; - public: - explicit NVPTXFrameLowering(NVPTXSubtarget &STI); + explicit NVPTXFrameLowering(); bool hasFP(const MachineFunction &MF) const override; void emitPrologue(MachineFunction &MF) const override; void emitEpilogue(MachineFunction &MF, MachineBasicBlock &MBB) const override; - void eliminateCallFramePseudoInstr(MachineFunction &MF, - MachineBasicBlock &MBB, - MachineBasicBlock::iterator I) const override; + void + eliminateCallFramePseudoInstr(MachineFunction &MF, MachineBasicBlock &MBB, + MachineBasicBlock::iterator I) const override; }; } // End llvm namespace diff --git a/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp index 22992b47f40..e01c7801507 100644 --- a/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp +++ b/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp @@ -50,7 +50,7 @@ FunctionPass *llvm::createNVPTXISelDag(NVPTXTargetMachine &TM, NVPTXDAGToDAGISel::NVPTXDAGToDAGISel(NVPTXTargetMachine &tm, CodeGenOpt::Level OptLevel) - : SelectionDAGISel(tm, OptLevel) { + : SelectionDAGISel(tm, OptLevel), TM(tm) { doMulWide = (OptLevel > 0); } @@ -580,20 +580,16 @@ SDNode *NVPTXDAGToDAGISel::SelectAddrSpaceCast(SDNode *N) { switch (SrcAddrSpace) { default: report_fatal_error("Bad address space in addrspacecast"); case ADDRESS_SPACE_GLOBAL: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_global_yes_64 - : NVPTX::cvta_global_yes; + Opc = TM.is64Bit() ? NVPTX::cvta_global_yes_64 : NVPTX::cvta_global_yes; break; case ADDRESS_SPACE_SHARED: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_shared_yes_64 - : NVPTX::cvta_shared_yes; + Opc = TM.is64Bit() ? NVPTX::cvta_shared_yes_64 : NVPTX::cvta_shared_yes; break; case ADDRESS_SPACE_CONST: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_const_yes_64 - : NVPTX::cvta_const_yes; + Opc = TM.is64Bit() ? NVPTX::cvta_const_yes_64 : NVPTX::cvta_const_yes; break; case ADDRESS_SPACE_LOCAL: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_local_yes_64 - : NVPTX::cvta_local_yes; + Opc = TM.is64Bit() ? NVPTX::cvta_local_yes_64 : NVPTX::cvta_local_yes; break; } return CurDAG->getMachineNode(Opc, SDLoc(N), N->getValueType(0), Src); @@ -605,20 +601,20 @@ SDNode *NVPTXDAGToDAGISel::SelectAddrSpaceCast(SDNode *N) { switch (DstAddrSpace) { default: report_fatal_error("Bad address space in addrspacecast"); case ADDRESS_SPACE_GLOBAL: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_to_global_yes_64 - : NVPTX::cvta_to_global_yes; + Opc = TM.is64Bit() ? NVPTX::cvta_to_global_yes_64 + : NVPTX::cvta_to_global_yes; break; case ADDRESS_SPACE_SHARED: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_to_shared_yes_64 - : NVPTX::cvta_to_shared_yes; + Opc = TM.is64Bit() ? NVPTX::cvta_to_shared_yes_64 + : NVPTX::cvta_to_shared_yes; break; case ADDRESS_SPACE_CONST: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_to_const_yes_64 - : NVPTX::cvta_to_const_yes; + Opc = + TM.is64Bit() ? NVPTX::cvta_to_const_yes_64 : NVPTX::cvta_to_const_yes; break; case ADDRESS_SPACE_LOCAL: - Opc = Subtarget->is64Bit() ? NVPTX::cvta_to_local_yes_64 - : NVPTX::cvta_to_local_yes; + Opc = + TM.is64Bit() ? NVPTX::cvta_to_local_yes_64 : NVPTX::cvta_to_local_yes; break; } return CurDAG->getMachineNode(Opc, SDLoc(N), N->getValueType(0), Src); @@ -714,9 +710,8 @@ SDNode *NVPTXDAGToDAGISel::SelectLoad(SDNode *N) { getI32Imm(vecType), getI32Imm(fromType), getI32Imm(fromTypeWidth), Addr, Chain }; NVPTXLD = CurDAG->getMachineNode(Opcode, dl, TargetVT, MVT::Other, Ops); - } else if (Subtarget->is64Bit() - ? SelectADDRsi64(N1.getNode(), N1, Base, Offset) - : SelectADDRsi(N1.getNode(), N1, Base, Offset)) { + } else if (TM.is64Bit() ? SelectADDRsi64(N1.getNode(), N1, Base, Offset) + : SelectADDRsi(N1.getNode(), N1, Base, Offset)) { switch (TargetVT) { case MVT::i8: Opcode = NVPTX::LD_i8_asi; @@ -743,10 +738,9 @@ SDNode *NVPTXDAGToDAGISel::SelectLoad(SDNode *N) { getI32Imm(vecType), getI32Imm(fromType), getI32Imm(fromTypeWidth), Base, Offset, Chain }; NVPTXLD = CurDAG->getMachineNode(Opcode, dl, TargetVT, MVT::Other, Ops); - } else if (Subtarget->is64Bit() - ? SelectADDRri64(N1.getNode(), N1, Base, Offset) - : SelectADDRri(N1.getNode(), N1, Base, Offset)) { - if (Subtarget->is64Bit()) { + } else if (TM.is64Bit() ? SelectADDRri64(N1.getNode(), N1, Base, Offset) + : SelectADDRri(N1.getNode(), N1, Base, Offset)) { + if (TM.is64Bit()) { switch (TargetVT) { case MVT::i8: Opcode = NVPTX::LD_i8_ari_64; @@ -798,7 +792,7 @@ SDNode *NVPTXDAGToDAGISel::SelectLoad(SDNode *N) { getI32Imm(fromTypeWidth), Base, Offset, Chain }; NVPTXLD = CurDAG->getMachineNode(Opcode, dl, TargetVT, MVT::Other, Ops); } else { - if (Subtarget->is64Bit()) { + if (TM.is64Bit()) { switch (TargetVT) { case MVT::i8: Opcode = NVPTX::LD_i8_areg_64; @@ -975,9 +969,8 @@ SDNode *NVPTXDAGToDAGISel::SelectLoadVector(SDNode *N) { getI32Imm(VecType), getI32Imm(FromType), getI32Imm(FromTypeWidth), Addr, Chain }; LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops); - } else if (Subtarget->is64Bit() - ? SelectADDRsi64(Op1.getNode(), Op1, Base, Offset) - : SelectADDRsi(Op1.getNode(), Op1, Base, Offset)) { + } else if (TM.is64Bit() ? SelectADDRsi64(Op1.getNode(), Op1, Base, Offset) + : SelectADDRsi(Op1.getNode(), Op1, Base, Offset)) { switch (N->getOpcode()) { default: return nullptr; @@ -1029,10 +1022,9 @@ SDNode *NVPTXDAGToDAGISel::SelectLoadVector(SDNode *N) { getI32Imm(VecType), getI32Imm(FromType), getI32Imm(FromTypeWidth), Base, Offset, Chain }; LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops); - } else if (Subtarget->is64Bit() - ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset) - : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) { - if (Subtarget->is64Bit()) { + } else if (TM.is64Bit() ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset) + : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) { + if (TM.is64Bit()) { switch (N->getOpcode()) { default: return nullptr; @@ -1134,7 +1126,7 @@ SDNode *NVPTXDAGToDAGISel::SelectLoadVector(SDNode *N) { LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops); } else { - if (Subtarget->is64Bit()) { + if (TM.is64Bit()) { switch (N->getOpcode()) { default: return nullptr; @@ -1426,10 +1418,9 @@ SDNode *NVPTXDAGToDAGISel::SelectLDGLDU(SDNode *N) { SDValue Ops[] = { Addr, Chain }; LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops); - } else if (Subtarget->is64Bit() - ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset) - : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) { - if (Subtarget->is64Bit()) { + } else if (TM.is64Bit() ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset) + : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) { + if (TM.is64Bit()) { switch (N->getOpcode()) { default: return nullptr; @@ -1711,7 +1702,7 @@ SDNode *NVPTXDAGToDAGISel::SelectLDGLDU(SDNode *N) { LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops); } else { - if (Subtarget->is64Bit()) { + if (TM.is64Bit()) { switch (N->getOpcode()) { default: return nullptr; @@ -2084,9 +2075,8 @@ SDNode *NVPTXDAGToDAGISel::SelectStore(SDNode *N) { getI32Imm(vecType), getI32Imm(toType), getI32Imm(toTypeWidth), Addr, Chain }; NVPTXST = CurDAG->getMachineNode(Opcode, dl, MVT::Other, Ops); - } else if (Subtarget->is64Bit() - ? SelectADDRsi64(N2.getNode(), N2, Base, Offset) - : SelectADDRsi(N2.getNode(), N2, Base, Offset)) { + } else if (TM.is64Bit() ? SelectADDRsi64(N2.getNode(), N2, Base, Offset) + : SelectADDRsi(N2.getNode(), N2, Base, Offset)) { switch (SourceVT) { case MVT::i8: Opcode = NVPTX::ST_i8_asi; @@ -2113,10 +2103,9 @@ SDNode *NVPTXDAGToDAGISel::SelectStore(SDNode *N) { getI32Imm(vecType), getI32Imm(toType), getI32Imm(toTypeWidth), Base, Offset, Chain }; NVPTXST = CurDAG->getMachineNode(Opcode, dl, MVT::Other, Ops); - } else if (Subtarget->is64Bit() - ? SelectADDRri64(N2.getNode(), N2, Base, Offset) - : SelectADDRri(N2.getNode(), N2, Base, Offset)) { - if (Subtarget->is64Bit()) { + } else if (TM.is64Bit() ? SelectADDRri64(N2.getNode(), N2, Base, Offset) + : SelectADDRri(N2.getNode(), N2, Base, Offset)) { + if (TM.is64Bit()) { switch (SourceVT) { case MVT::i8: Opcode = NVPTX::ST_i8_ari_64; @@ -2168,7 +2157,7 @@ SDNode *NVPTXDAGToDAGISel::SelectStore(SDNode *N) { getI32Imm(toTypeWidth), Base, Offset, Chain }; NVPTXST = CurDAG->getMachineNode(Opcode, dl, MVT::Other, Ops); } else { - if (Subtarget->is64Bit()) { + if (TM.is64Bit()) { switch (SourceVT) { case MVT::i8: Opcode = NVPTX::ST_i8_areg_64; @@ -2345,9 +2334,8 @@ SDNode *NVPTXDAGToDAGISel::SelectStoreVector(SDNode *N) { break; } StOps.push_back(Addr); - } else if (Subtarget->is64Bit() - ? SelectADDRsi64(N2.getNode(), N2, Base, Offset) - : SelectADDRsi(N2.getNode(), N2, Base, Offset)) { + } else if (TM.is64Bit() ? SelectADDRsi64(N2.getNode(), N2, Base, Offset) + : SelectADDRsi(N2.getNode(), N2, Base, Offset)) { switch (N->getOpcode()) { default: return nullptr; @@ -2396,10 +2384,9 @@ SDNode *NVPTXDAGToDAGISel::SelectStoreVector(SDNode *N) { } StOps.push_back(Base); StOps.push_back(Offset); - } else if (Subtarget->is64Bit() - ? SelectADDRri64(N2.getNode(), N2, Base, Offset) - : SelectADDRri(N2.getNode(), N2, Base, Offset)) { - if (Subtarget->is64Bit()) { + } else if (TM.is64Bit() ? SelectADDRri64(N2.getNode(), N2, Base, Offset) + : SelectADDRri(N2.getNode(), N2, Base, Offset)) { + if (TM.is64Bit()) { switch (N->getOpcode()) { default: return nullptr; @@ -2497,7 +2484,7 @@ SDNode *NVPTXDAGToDAGISel::SelectStoreVector(SDNode *N) { StOps.push_back(Base); StOps.push_back(Offset); } else { - if (Subtarget->is64Bit()) { + if (TM.is64Bit()) { switch (N->getOpcode()) { default: return nullptr; diff --git a/lib/Target/NVPTX/NVPTXISelDAGToDAG.h b/lib/Target/NVPTX/NVPTXISelDAGToDAG.h index 1548c920c0e..ca432b53be8 100644 --- a/lib/Target/NVPTX/NVPTXISelDAGToDAG.h +++ b/lib/Target/NVPTX/NVPTXISelDAGToDAG.h @@ -26,6 +26,7 @@ using namespace llvm; namespace { class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel { + const NVPTXTargetMachine &TM; // If true, generate mul.wide from sext and mul bool doMulWide; diff --git a/lib/Target/NVPTX/NVPTXInstrInfo.cpp b/lib/Target/NVPTX/NVPTXInstrInfo.cpp index 740ca0328ef..dabc3be43a3 100644 --- a/lib/Target/NVPTX/NVPTXInstrInfo.cpp +++ b/lib/Target/NVPTX/NVPTXInstrInfo.cpp @@ -28,9 +28,7 @@ using namespace llvm; // Pin the vtable to this file. void NVPTXInstrInfo::anchor() {} -// FIXME: Add the subtarget support on this constructor. -NVPTXInstrInfo::NVPTXInstrInfo(NVPTXSubtarget &STI) - : NVPTXGenInstrInfo(), RegInfo(STI) {} +NVPTXInstrInfo::NVPTXInstrInfo() : NVPTXGenInstrInfo(), RegInfo() {} void NVPTXInstrInfo::copyPhysReg( MachineBasicBlock &MBB, MachineBasicBlock::iterator I, DebugLoc DL, diff --git a/lib/Target/NVPTX/NVPTXInstrInfo.h b/lib/Target/NVPTX/NVPTXInstrInfo.h index 6de75364a82..9b5d491dfeb 100644 --- a/lib/Target/NVPTX/NVPTXInstrInfo.h +++ b/lib/Target/NVPTX/NVPTXInstrInfo.h @@ -27,7 +27,7 @@ class NVPTXInstrInfo : public NVPTXGenInstrInfo { const NVPTXRegisterInfo RegInfo; virtual void anchor(); public: - explicit NVPTXInstrInfo(NVPTXSubtarget &STI); + explicit NVPTXInstrInfo(); const NVPTXRegisterInfo &getRegisterInfo() const { return RegInfo; } diff --git a/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/lib/Target/NVPTX/NVPTXRegisterInfo.cpp index 358ccce3981..5ca96e4efb2 100644 --- a/lib/Target/NVPTX/NVPTXRegisterInfo.cpp +++ b/lib/Target/NVPTX/NVPTXRegisterInfo.cpp @@ -71,8 +71,7 @@ std::string getNVPTXRegClassStr(TargetRegisterClass const *RC) { } } -NVPTXRegisterInfo::NVPTXRegisterInfo(const NVPTXSubtarget &st) - : NVPTXGenRegisterInfo(0), Is64Bit(st.is64Bit()) {} +NVPTXRegisterInfo::NVPTXRegisterInfo() : NVPTXGenRegisterInfo(0) {} #define GET_REGINFO_TARGET_DESC #include "NVPTXGenRegisterInfo.inc" diff --git a/lib/Target/NVPTX/NVPTXRegisterInfo.h b/lib/Target/NVPTX/NVPTXRegisterInfo.h index d2e67331f78..75b8f153fa6 100644 --- a/lib/Target/NVPTX/NVPTXRegisterInfo.h +++ b/lib/Target/NVPTX/NVPTXRegisterInfo.h @@ -22,19 +22,13 @@ #include "NVPTXGenRegisterInfo.inc" namespace llvm { - -// Forward Declarations. -class TargetInstrInfo; -class NVPTXSubtarget; - class NVPTXRegisterInfo : public NVPTXGenRegisterInfo { private: - bool Is64Bit; // Hold Strings that can be free'd all together with NVPTXRegisterInfo ManagedStringPool ManagedStrPool; public: - NVPTXRegisterInfo(const NVPTXSubtarget &st); + NVPTXRegisterInfo(); //------------------------------------------------------ // Pure virtual functions from TargetRegisterInfo diff --git a/lib/Target/NVPTX/NVPTXSubtarget.cpp b/lib/Target/NVPTX/NVPTXSubtarget.cpp index 85738565528..069d6e179dd 100644 --- a/lib/Target/NVPTX/NVPTXSubtarget.cpp +++ b/lib/Target/NVPTX/NVPTXSubtarget.cpp @@ -45,11 +45,10 @@ NVPTXSubtarget &NVPTXSubtarget::initializeSubtargetDependencies(StringRef CPU, NVPTXSubtarget::NVPTXSubtarget(const std::string &TT, const std::string &CPU, const std::string &FS, - const NVPTXTargetMachine &TM, bool is64Bit) - : NVPTXGenSubtargetInfo(TT, CPU, FS), Is64Bit(is64Bit), PTXVersion(0), - SmVersion(20), TM(TM), - InstrInfo(initializeSubtargetDependencies(CPU, FS)), TLInfo(TM, *this), - TSInfo(TM.getDataLayout()), FrameLowering(*this) {} + const NVPTXTargetMachine &TM) + : NVPTXGenSubtargetInfo(TT, CPU, FS), PTXVersion(0), SmVersion(20), TM(TM), + InstrInfo(), TLInfo(TM, initializeSubtargetDependencies(CPU, FS)), + TSInfo(TM.getDataLayout()), FrameLowering() {} bool NVPTXSubtarget::hasImageHandles() const { // Enable handles for Kepler+, where CUDA supports indirect surfaces and diff --git a/lib/Target/NVPTX/NVPTXSubtarget.h b/lib/Target/NVPTX/NVPTXSubtarget.h index 0a31fb03302..e9833e5823c 100644 --- a/lib/Target/NVPTX/NVPTXSubtarget.h +++ b/lib/Target/NVPTX/NVPTXSubtarget.h @@ -32,7 +32,6 @@ namespace llvm { class NVPTXSubtarget : public NVPTXGenSubtargetInfo { virtual void anchor(); std::string TargetName; - bool Is64Bit; // PTX version x.y is represented as 10*x+y, e.g. 3.1 == 31 unsigned PTXVersion; @@ -54,8 +53,7 @@ public: /// of the specified module. /// NVPTXSubtarget(const std::string &TT, const std::string &CPU, - const std::string &FS, const NVPTXTargetMachine &TM, - bool is64Bit); + const std::string &FS, const NVPTXTargetMachine &TM); const TargetFrameLowering *getFrameLowering() const override { return &FrameLowering; @@ -95,7 +93,6 @@ public: inline bool hasROT32() const { return hasHWROT32() || hasSWROT32(); } inline bool hasROT64() const { return SmVersion >= 20; } bool hasImageHandles() const; - bool is64Bit() const { return Is64Bit; } unsigned int getSmVersion() const { return SmVersion; } std::string getTargetName() const { return TargetName; } diff --git a/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/lib/Target/NVPTX/NVPTXTargetMachine.cpp index 7c6ecc387be..1a267a6a114 100644 --- a/lib/Target/NVPTX/NVPTXTargetMachine.cpp +++ b/lib/Target/NVPTX/NVPTXTargetMachine.cpp @@ -88,7 +88,7 @@ NVPTXTargetMachine::NVPTXTargetMachine(const Target &T, StringRef TT, CodeGenOpt::Level OL, bool is64bit) : LLVMTargetMachine(T, TT, CPU, FS, Options, RM, CM, OL), is64bit(is64bit), TLOF(make_unique()), - DL(computeDataLayout(is64bit)), Subtarget(TT, CPU, FS, *this, is64bit) { + DL(computeDataLayout(is64bit)), Subtarget(TT, CPU, FS, *this) { if (Triple(TT).getOS() == Triple::NVCL) drvInterface = NVPTX::NVCL; else -- 2.34.1