6502/llvm-6502
1
0
Fork 0
mirror of https://github.com/c64scene-ar/llvm-6502.git synced 2025-01-23 02:32:11 +00:00
Code Issues Projects Releases Wiki Activity

Remove all use of is64bit off of NVPTXSubtarget and clean up code

Browse Source 
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
This commit is contained in:
Eric Christopher 2015-02-19 00:08:27 +00:00
parent c72978539d
commit 6ec9683959
12 changed files with 61 additions and 89 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
lib/Target/NVPTX/NVPTXAsmPrinter.cpp
View File

@ -819,7 +819,7 @@ bool NVPTXAsmPrinter::doInitialization(Module &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::setAndEmitFunctionVirtualRegisters(
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 {

7
lib/Target/NVPTX/NVPTXFrameLowering.cpp
View File

@ -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<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
unsigned LocalReg = MRI.createVirtualRegister(&NVPTX::Int64RegsRegClass);
MachineInstr *MI =
BuildMI(MBB, MBBI, dl, MF.getSubtarget().getInstrInfo()->get(

10
lib/Target/NVPTX/NVPTXFrameLowering.h
View File

@ -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

95
lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
View File

@ -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;

1
lib/Target/NVPTX/NVPTXISelDAGToDAG.h
View File

@ -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;

4
lib/Target/NVPTX/NVPTXInstrInfo.cpp
View File

@ -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,

2
lib/Target/NVPTX/NVPTXInstrInfo.h
View File

@ -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; }

3
lib/Target/NVPTX/NVPTXRegisterInfo.cpp
View File

@ -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"

8
lib/Target/NVPTX/NVPTXRegisterInfo.h
View File

@ -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

9
lib/Target/NVPTX/NVPTXSubtarget.cpp
View File

@ -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

5
lib/Target/NVPTX/NVPTXSubtarget.h
View File

@ -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; }

2
lib/Target/NVPTX/NVPTXTargetMachine.cpp
View File

@ -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<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