[llvm] r229787 - Remove all use of is64bit off of NVPTXSubtarget and clean up code
Eric Christopher
echristo at gmail.com
Wed Feb 18 16:08:27 PST 2015
Author: echristo
Date: Wed Feb 18 18:08:27 2015
New Revision: 229787
URL: http://llvm.org/viewvc/llvm-project?rev=229787&view=rev
Log:
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.
Modified:
llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.h
llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.h
llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.h
llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h
llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp Wed Feb 18 18:08:27 2015
@@ -819,7 +819,7 @@ bool NVPTXAsmPrinter::doInitialization(M
StringRef CPU = TM.getTargetCPU();
StringRef FS = TM.getTargetFeatureString();
const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(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::setAndEmitFunction
if (NumBytes) {
O << "\t.local .align " << MFI->getMaxAlignment() << " .b8 \t" << DEPOTNAME
<< getFunctionNumber() << "[" << NumBytes << "];\n";
- if (nvptxSubtarget->is64Bit()) {
+ if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
O << "\t.reg .b64 \t%SP;\n";
O << "\t.reg .b64 \t%SPL;\n";
} else {
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.cpp?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.cpp Wed Feb 18 18:08:27 2015
@@ -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(Ma
// mov %SPL, %depot;
// cvta.local %SP, %SPL;
- if (is64bit) {
+ if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
unsigned LocalReg = MRI.createVirtualRegister(&NVPTX::Int64RegsRegClass);
MachineInstr *MI =
BuildMI(MBB, MBBI, dl, MF.getSubtarget().getInstrInfo()->get(
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.h?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXFrameLowering.h Wed Feb 18 18:08:27 2015
@@ -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
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp Wed Feb 18 18:08:27 2015
@@ -50,7 +50,7 @@ FunctionPass *llvm::createNVPTXISelDag(N
NVPTXDAGToDAGISel::NVPTXDAGToDAGISel(NVPTXTargetMachine &tm,
CodeGenOpt::Level OptLevel)
- : SelectionDAGISel(tm, OptLevel) {
+ : SelectionDAGISel(tm, OptLevel), TM(tm) {
doMulWide = (OptLevel > 0);
}
@@ -580,20 +580,16 @@ SDNode *NVPTXDAGToDAGISel::SelectAddrSpa
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::SelectAddrSpa
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(SD
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(SD
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(SD
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::SelectLoadVec
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::SelectLoadVec
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::SelectLoadVec
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(
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(
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(S
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(S
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(S
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::SelectStoreVe
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::SelectStoreVe
}
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::SelectStoreVe
StOps.push_back(Base);
StOps.push_back(Offset);
} else {
- if (Subtarget->is64Bit()) {
+ if (TM.is64Bit()) {
switch (N->getOpcode()) {
default:
return nullptr;
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h Wed Feb 18 18:08:27 2015
@@ -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;
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp Wed Feb 18 18:08:27 2015
@@ -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,
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.h?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.h Wed Feb 18 18:08:27 2015
@@ -27,7 +27,7 @@ class NVPTXInstrInfo : public NVPTXGenIn
const NVPTXRegisterInfo RegInfo;
virtual void anchor();
public:
- explicit NVPTXInstrInfo(NVPTXSubtarget &STI);
+ explicit NVPTXInstrInfo();
const NVPTXRegisterInfo &getRegisterInfo() const { return RegInfo; }
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.cpp?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.cpp Wed Feb 18 18:08:27 2015
@@ -71,8 +71,7 @@ std::string getNVPTXRegClassStr(TargetRe
}
}
-NVPTXRegisterInfo::NVPTXRegisterInfo(const NVPTXSubtarget &st)
- : NVPTXGenRegisterInfo(0), Is64Bit(st.is64Bit()) {}
+NVPTXRegisterInfo::NVPTXRegisterInfo() : NVPTXGenRegisterInfo(0) {}
#define GET_REGINFO_TARGET_DESC
#include "NVPTXGenRegisterInfo.inc"
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.h?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.h Wed Feb 18 18:08:27 2015
@@ -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
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.cpp?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.cpp Wed Feb 18 18:08:27 2015
@@ -45,11 +45,10 @@ NVPTXSubtarget &NVPTXSubtarget::initiali
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
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h Wed Feb 18 18:08:27 2015
@@ -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; }
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp?rev=229787&r1=229786&r2=229787&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp Wed Feb 18 18:08:27 2015
@@ -88,7 +88,7 @@ NVPTXTargetMachine::NVPTXTargetMachine(c
CodeGenOpt::Level OL, bool is64bit)
: LLVMTargetMachine(T, TT, CPU, FS, Options, RM, CM, OL), is64bit(is64bit),
TLOF(make_unique<NVPTXTargetObjectFile>()),
- 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
More information about the llvm-commits
mailing list