[llvm-commits] [llvm] r145947 - in /llvm/trunk: lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h lib/Target/PTX/PTXAsmPrinter.cpp lib/Target/PTX/PTXISelLowering.cpp lib/Target/PTX/PTXInstrInfo.td lib/Target/PTX/PTXMFInfoExtract.cpp lib/Target/PTX/PTXMachineFunctionInfo.h test/CodeGen/PTX/mov.ll test/CodeGen/PTX/parameter-order.ll
Justin Holewinski
justin.holewinski at gmail.com
Tue Dec 6 09:39:48 PST 2011
Author: jholewinski
Date: Tue Dec 6 11:39:48 2011
New Revision: 145947
URL: http://llvm.org/viewvc/llvm-project?rev=145947&view=rev
Log:
PTX: Continue to fix up the register mess.
Modified:
llvm/trunk/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp
llvm/trunk/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h
llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp
llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp
llvm/trunk/lib/Target/PTX/PTXInstrInfo.td
llvm/trunk/lib/Target/PTX/PTXMFInfoExtract.cpp
llvm/trunk/lib/Target/PTX/PTXMachineFunctionInfo.h
llvm/trunk/test/CodeGen/PTX/mov.ll
llvm/trunk/test/CodeGen/PTX/parameter-order.ll
Modified: llvm/trunk/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp (original)
+++ llvm/trunk/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp Tue Dec 6 11:39:48 2011
@@ -39,32 +39,45 @@
void PTXInstPrinter::printRegName(raw_ostream &OS, unsigned RegNo) const {
// Decode the register number into type and offset
- unsigned RegType = RegNo & 0xF;
- unsigned RegOffset = RegNo >> 4;
+ unsigned RegSpace = RegNo & 0x7;
+ unsigned RegType = (RegNo >> 3) & 0x7;
+ unsigned RegOffset = RegNo >> 6;
// Print the register
OS << "%";
- switch (RegType) {
+ switch (RegSpace) {
default:
- llvm_unreachable("Unknown register type!");
- case PTXRegisterType::Pred:
- OS << "p";
- break;
- case PTXRegisterType::B16:
- OS << "rh";
- break;
- case PTXRegisterType::B32:
- OS << "r";
- break;
- case PTXRegisterType::B64:
- OS << "rd";
+ llvm_unreachable("Unknown register space!");
+ case PTXRegisterSpace::Reg:
+ switch (RegType) {
+ default:
+ llvm_unreachable("Unknown register type!");
+ case PTXRegisterType::Pred:
+ OS << "p";
+ break;
+ case PTXRegisterType::B16:
+ OS << "rh";
+ break;
+ case PTXRegisterType::B32:
+ OS << "r";
+ break;
+ case PTXRegisterType::B64:
+ OS << "rd";
+ break;
+ case PTXRegisterType::F32:
+ OS << "f";
+ break;
+ case PTXRegisterType::F64:
+ OS << "fd";
+ break;
+ }
break;
- case PTXRegisterType::F32:
- OS << "f";
+ case PTXRegisterSpace::Return:
+ OS << "ret";
break;
- case PTXRegisterType::F64:
- OS << "fd";
+ case PTXRegisterSpace::Argument:
+ OS << "arg";
break;
}
Modified: llvm/trunk/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h (original)
+++ llvm/trunk/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h Tue Dec 6 11:39:48 2011
@@ -17,6 +17,8 @@
#ifndef PTXBASEINFO_H
#define PTXBASEINFO_H
+#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/raw_ostream.h"
#include "PTXMCTargetDesc.h"
namespace llvm {
@@ -69,6 +71,63 @@
F64
};
} // namespace PTXRegisterType
+
+ namespace PTXRegisterSpace {
+ // Register space encoded in MCOperands
+ enum {
+ Reg = 0,
+ Local,
+ Param,
+ Argument,
+ Return
+ };
+ }
+
+ inline static void decodeRegisterName(raw_ostream &OS,
+ unsigned EncodedReg) {
+ OS << "%";
+
+ unsigned RegSpace = EncodedReg & 0x7;
+ unsigned RegType = (EncodedReg >> 3) & 0x7;
+ unsigned RegOffset = EncodedReg >> 6;
+
+ switch (RegSpace) {
+ default:
+ llvm_unreachable("Unknown register space!");
+ case PTXRegisterSpace::Reg:
+ switch (RegType) {
+ default:
+ llvm_unreachable("Unknown register type!");
+ case PTXRegisterType::Pred:
+ OS << "p";
+ break;
+ case PTXRegisterType::B16:
+ OS << "rh";
+ break;
+ case PTXRegisterType::B32:
+ OS << "r";
+ break;
+ case PTXRegisterType::B64:
+ OS << "rd";
+ break;
+ case PTXRegisterType::F32:
+ OS << "f";
+ break;
+ case PTXRegisterType::F64:
+ OS << "fd";
+ break;
+ }
+ break;
+ case PTXRegisterSpace::Return:
+ OS << "ret";
+ break;
+ case PTXRegisterSpace::Argument:
+ OS << "arg";
+ break;
+ }
+
+ OS << RegOffset;
+ }
} // namespace llvm
#endif
Modified: llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp Tue Dec 6 11:39:48 2011
@@ -51,23 +51,23 @@
static const char PARAM_PREFIX[] = "__param_";
static const char RETURN_PREFIX[] = "__ret_";
-static const char *getRegisterTypeName(unsigned RegNo,
- const MachineRegisterInfo& MRI) {
- const TargetRegisterClass *TRC = MRI.getRegClass(RegNo);
-
-#define TEST_REGCLS(cls, clsstr) \
- if (PTX::cls ## RegisterClass == TRC) return # clsstr;
-
- TEST_REGCLS(RegPred, pred);
- TEST_REGCLS(RegI16, b16);
- TEST_REGCLS(RegI32, b32);
- TEST_REGCLS(RegI64, b64);
- TEST_REGCLS(RegF32, b32);
- TEST_REGCLS(RegF64, b64);
-#undef TEST_REGCLS
-
- llvm_unreachable("Not in any register class!");
- return NULL;
+static const char *getRegisterTypeName(unsigned RegType) {
+ switch (RegType) {
+ default:
+ llvm_unreachable("Unknown register type");
+ case PTXRegisterType::Pred:
+ return ".pred";
+ case PTXRegisterType::B16:
+ return ".b16";
+ case PTXRegisterType::B32:
+ return ".b32";
+ case PTXRegisterType::B64:
+ return ".b64";
+ case PTXRegisterType::F32:
+ return ".f32";
+ case PTXRegisterType::F64:
+ return ".f64";
+ }
}
static const char *getStateSpaceName(unsigned addressSpace) {
@@ -188,32 +188,32 @@
unsigned numRegs;
// pred
- numRegs = MFI->getNumRegistersForClass(PTX::RegPredRegisterClass);
+ numRegs = MFI->countRegisters(PTXRegisterType::Pred, PTXRegisterSpace::Reg);
if(numRegs > 0)
os << "\t.reg .pred %p<" << numRegs << ">;\n";
// i16
- numRegs = MFI->getNumRegistersForClass(PTX::RegI16RegisterClass);
+ numRegs = MFI->countRegisters(PTXRegisterType::B16, PTXRegisterSpace::Reg);
if(numRegs > 0)
os << "\t.reg .b16 %rh<" << numRegs << ">;\n";
// i32
- numRegs = MFI->getNumRegistersForClass(PTX::RegI32RegisterClass);
+ numRegs = MFI->countRegisters(PTXRegisterType::B32, PTXRegisterSpace::Reg);
if(numRegs > 0)
os << "\t.reg .b32 %r<" << numRegs << ">;\n";
// i64
- numRegs = MFI->getNumRegistersForClass(PTX::RegI64RegisterClass);
+ numRegs = MFI->countRegisters(PTXRegisterType::B64, PTXRegisterSpace::Reg);
if(numRegs > 0)
os << "\t.reg .b64 %rd<" << numRegs << ">;\n";
// f32
- numRegs = MFI->getNumRegistersForClass(PTX::RegF32RegisterClass);
+ numRegs = MFI->countRegisters(PTXRegisterType::F32, PTXRegisterSpace::Reg);
if(numRegs > 0)
os << "\t.reg .f32 %f<" << numRegs << ">;\n";
// f64
- numRegs = MFI->getNumRegistersForClass(PTX::RegF64RegisterClass);
+ numRegs = MFI->countRegisters(PTXRegisterType::F64, PTXRegisterSpace::Reg);
if(numRegs > 0)
os << "\t.reg .f64 %fd<" << numRegs << ">;\n";
@@ -368,7 +368,6 @@
const PTXParamManager &PM = MFI->getParamManager();
const bool isKernel = MFI->isKernel();
const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>();
- const MachineRegisterInfo& MRI = MF->getRegInfo();
SmallString<128> decl;
raw_svector_ostream os(decl);
@@ -391,7 +390,7 @@
if (i != b)
os << ", ";
- os << ".reg ." << getRegisterTypeName(*i, MRI) << ' '
+ os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' '
<< MFI->getRegisterName(*i);
}
}
@@ -450,7 +449,7 @@
if (i != b)
os << ", ";
- os << ".reg ." << getRegisterTypeName(*i, MRI) << ' '
+ os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' '
<< MFI->getRegisterName(*i);
}
}
@@ -521,34 +520,14 @@
MCOperand PTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
MCOperand MCOp;
const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>();
- const MachineRegisterInfo& MRI = MF->getRegInfo();
- const TargetRegisterClass* TRC;
- unsigned RegType;
- unsigned RegOffset;
unsigned EncodedReg;
switch (MO.getType()) {
default:
llvm_unreachable("Unknown operand type");
case MachineOperand::MO_Register:
if (MO.getReg() > 0) {
- TRC = MRI.getRegClass(MO.getReg());
- // Determine which PTX register type to use
- if (TRC == PTX::RegPredRegisterClass)
- RegType = PTXRegisterType::Pred;
- else if (TRC == PTX::RegI16RegisterClass)
- RegType = PTXRegisterType::B16;
- else if (TRC == PTX::RegI32RegisterClass)
- RegType = PTXRegisterType::B32;
- else if (TRC == PTX::RegI64RegisterClass)
- RegType = PTXRegisterType::B64;
- else if (TRC == PTX::RegF32RegisterClass)
- RegType = PTXRegisterType::F32;
- else if (TRC == PTX::RegF64RegisterClass)
- RegType = PTXRegisterType::F64;
- // Determine our virtual register offset
- RegOffset = MFI->getOffsetForRegister(TRC, MO.getReg());
// Encode the register
- EncodedReg = (RegOffset << 4) | RegType;
+ EncodedReg = MFI->getEncodedRegister(MO.getReg());
} else {
EncodedReg = 0;
}
Modified: llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp (original)
+++ llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp Tue Dec 6 11:39:48 2011
@@ -243,6 +243,30 @@
for (unsigned i = 0, e = Ins.size(); i != e; ++i) {
EVT RegVT = Ins[i].VT;
TargetRegisterClass* TRC = getRegClassFor(RegVT);
+ unsigned RegType;
+
+ // Determine which register class we need
+ if (RegVT == MVT::i1) {
+ RegType = PTXRegisterType::Pred;
+ }
+ else if (RegVT == MVT::i16) {
+ RegType = PTXRegisterType::B16;
+ }
+ else if (RegVT == MVT::i32) {
+ RegType = PTXRegisterType::B32;
+ }
+ else if (RegVT == MVT::i64) {
+ RegType = PTXRegisterType::B64;
+ }
+ else if (RegVT == MVT::f32) {
+ RegType = PTXRegisterType::F32;
+ }
+ else if (RegVT == MVT::f64) {
+ RegType = PTXRegisterType::F64;
+ }
+ else {
+ llvm_unreachable("Unknown parameter type");
+ }
// Use a unique index in the instruction to prevent instruction folding.
// Yes, this is a hack.
@@ -253,7 +277,7 @@
InVals.push_back(ArgValue);
- MFI->addArgReg(Reg);
+ MFI->addRegister(Reg, RegType, PTXRegisterSpace::Argument);
}
}
@@ -304,25 +328,32 @@
for (unsigned i = 0, e = Outs.size(); i != e; ++i) {
EVT RegVT = Outs[i].VT;
TargetRegisterClass* TRC = 0;
+ unsigned RegType;
// Determine which register class we need
if (RegVT == MVT::i1) {
TRC = PTX::RegPredRegisterClass;
+ RegType = PTXRegisterType::Pred;
}
else if (RegVT == MVT::i16) {
TRC = PTX::RegI16RegisterClass;
+ RegType = PTXRegisterType::B16;
}
else if (RegVT == MVT::i32) {
TRC = PTX::RegI32RegisterClass;
+ RegType = PTXRegisterType::B32;
}
else if (RegVT == MVT::i64) {
TRC = PTX::RegI64RegisterClass;
+ RegType = PTXRegisterType::B64;
}
else if (RegVT == MVT::f32) {
TRC = PTX::RegF32RegisterClass;
+ RegType = PTXRegisterType::F32;
}
else if (RegVT == MVT::f64) {
TRC = PTX::RegF64RegisterClass;
+ RegType = PTXRegisterType::F64;
}
else {
llvm_unreachable("Unknown parameter type");
@@ -335,7 +366,7 @@
Chain = DAG.getNode(PTXISD::WRITE_PARAM, dl, MVT::Other, Copy, OutReg);
- MFI->addRetReg(Reg);
+ MFI->addRegister(Reg, RegType, PTXRegisterSpace::Return);
}
}
Modified: llvm/trunk/lib/Target/PTX/PTXInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXInstrInfo.td?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXInstrInfo.td (original)
+++ llvm/trunk/lib/Target/PTX/PTXInstrInfo.td Tue Dec 6 11:39:48 2011
@@ -825,17 +825,17 @@
///===- Parameter Passing Pseudo-Instructions -----------------------------===//
def READPARAMPRED : InstPTX<(outs RegPred:$a), (ins i32imm:$b),
- "mov.pred\t$a, %param$b", []>;
+ "mov.pred\t$a, %arg$b", []>;
def READPARAMI16 : InstPTX<(outs RegI16:$a), (ins i32imm:$b),
- "mov.b16\t$a, %param$b", []>;
+ "mov.b16\t$a, %arg$b", []>;
def READPARAMI32 : InstPTX<(outs RegI32:$a), (ins i32imm:$b),
- "mov.b32\t$a, %param$b", []>;
+ "mov.b32\t$a, %arg$b", []>;
def READPARAMI64 : InstPTX<(outs RegI64:$a), (ins i32imm:$b),
- "mov.b64\t$a, %param$b", []>;
+ "mov.b64\t$a, %arg$b", []>;
def READPARAMF32 : InstPTX<(outs RegF32:$a), (ins i32imm:$b),
- "mov.f32\t$a, %param$b", []>;
+ "mov.f32\t$a, %arg$b", []>;
def READPARAMF64 : InstPTX<(outs RegF64:$a), (ins i32imm:$b),
- "mov.f64\t$a, %param$b", []>;
+ "mov.f64\t$a, %arg$b", []>;
def WRITEPARAMPRED : InstPTX<(outs), (ins RegPred:$a), "//w", []>;
def WRITEPARAMI16 : InstPTX<(outs), (ins RegI16:$a), "//w", []>;
Modified: llvm/trunk/lib/Target/PTX/PTXMFInfoExtract.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXMFInfoExtract.cpp?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXMFInfoExtract.cpp (original)
+++ llvm/trunk/lib/Target/PTX/PTXMFInfoExtract.cpp Tue Dec 6 11:39:48 2011
@@ -58,7 +58,20 @@
for (unsigned i = 0; i < MRI.getNumVirtRegs(); ++i) {
unsigned Reg = TargetRegisterInfo::index2VirtReg(i);
const TargetRegisterClass *TRC = MRI.getRegClass(Reg);
- MFI->addVirtualRegister(TRC, Reg);
+ unsigned RegType;
+ if (TRC == PTX::RegPredRegisterClass)
+ RegType = PTXRegisterType::Pred;
+ else if (TRC == PTX::RegI16RegisterClass)
+ RegType = PTXRegisterType::B16;
+ else if (TRC == PTX::RegI32RegisterClass)
+ RegType = PTXRegisterType::B32;
+ else if (TRC == PTX::RegI64RegisterClass)
+ RegType = PTXRegisterType::B64;
+ else if (TRC == PTX::RegF32RegisterClass)
+ RegType = PTXRegisterType::F32;
+ else if (TRC == PTX::RegF64RegisterClass)
+ RegType = PTXRegisterType::F64;
+ MFI->addRegister(Reg, RegType, PTXRegisterSpace::Reg);
}
return false;
Modified: llvm/trunk/lib/Target/PTX/PTXMachineFunctionInfo.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXMachineFunctionInfo.h?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXMachineFunctionInfo.h (original)
+++ llvm/trunk/lib/Target/PTX/PTXMachineFunctionInfo.h Tue Dec 6 11:39:48 2011
@@ -35,15 +35,22 @@
DenseSet<unsigned> RegArgs;
DenseSet<unsigned> RegRets;
- typedef std::vector<unsigned> RegisterList;
- typedef DenseMap<const TargetRegisterClass*, RegisterList> RegisterMap;
- typedef DenseMap<unsigned, std::string> RegisterNameMap;
typedef DenseMap<int, std::string> FrameMap;
- RegisterMap UsedRegs;
- RegisterNameMap RegNames;
FrameMap FrameSymbols;
+ struct RegisterInfo {
+ unsigned Reg;
+ unsigned Type;
+ unsigned Space;
+ unsigned Offset;
+ unsigned Encoded;
+ };
+
+ typedef DenseMap<unsigned, RegisterInfo> RegisterInfoMap;
+
+ RegisterInfoMap RegInfo;
+
PTXParamManager ParamManager;
public:
@@ -51,13 +58,7 @@
PTXMachineFunctionInfo(MachineFunction &MF)
: IsKernel(false) {
- UsedRegs[PTX::RegPredRegisterClass] = RegisterList();
- UsedRegs[PTX::RegI16RegisterClass] = RegisterList();
- UsedRegs[PTX::RegI32RegisterClass] = RegisterList();
- UsedRegs[PTX::RegI64RegisterClass] = RegisterList();
- UsedRegs[PTX::RegF32RegisterClass] = RegisterList();
- UsedRegs[PTX::RegF64RegisterClass] = RegisterList();
- }
+ }
/// getParamManager - Returns the PTXParamManager instance for this function.
PTXParamManager& getParamManager() { return ParamManager; }
@@ -78,81 +79,106 @@
reg_iterator retreg_begin() const { return RegRets.begin(); }
reg_iterator retreg_end() const { return RegRets.end(); }
+ /// addRegister - Adds a virtual register to the set of all used registers
+ void addRegister(unsigned Reg, unsigned RegType, unsigned RegSpace) {
+ if (!RegInfo.count(Reg)) {
+ RegisterInfo Info;
+ Info.Reg = Reg;
+ Info.Type = RegType;
+ Info.Space = RegSpace;
+
+ // Determine register offset
+ Info.Offset = 0;
+ for(RegisterInfoMap::const_iterator i = RegInfo.begin(),
+ e = RegInfo.end(); i != e; ++i) {
+ const RegisterInfo& RI = i->second;
+ if (RI.Space == RegSpace)
+ if (RI.Space != PTXRegisterSpace::Reg || RI.Type == Info.Type)
+ Info.Offset++;
+ }
+
+ // Encode the register data into a single register number
+ Info.Encoded = (Info.Offset << 6) | (Info.Type << 3) | Info.Space;
+
+ RegInfo[Reg] = Info;
+
+ if (RegSpace == PTXRegisterSpace::Argument)
+ RegArgs.insert(Reg);
+ else if (RegSpace == PTXRegisterSpace::Return)
+ RegRets.insert(Reg);
+ }
+ }
+
+ /// countRegisters - Returns the number of registers of the given type and
+ /// space.
+ unsigned countRegisters(unsigned RegType, unsigned RegSpace) const {
+ unsigned Count = 0;
+ for(RegisterInfoMap::const_iterator i = RegInfo.begin(), e = RegInfo.end();
+ i != e; ++i) {
+ const RegisterInfo& RI = i->second;
+ if (RI.Type == RegType && RI.Space == RegSpace)
+ Count++;
+ }
+ return Count;
+ }
+
+ /// getEncodedRegister - Returns the encoded value of the register.
+ unsigned getEncodedRegister(unsigned Reg) const {
+ return RegInfo.lookup(Reg).Encoded;
+ }
+
/// addRetReg - Adds a register to the set of return-value registers.
void addRetReg(unsigned Reg) {
if (!RegRets.count(Reg)) {
RegRets.insert(Reg);
- std::string name;
- name = "%ret";
- name += utostr(RegRets.size() - 1);
- RegNames[Reg] = name;
}
}
/// addArgReg - Adds a register to the set of function argument registers.
void addArgReg(unsigned Reg) {
RegArgs.insert(Reg);
- std::string name;
- name = "%param";
- name += utostr(RegArgs.size() - 1);
- RegNames[Reg] = name;
- }
-
- /// addVirtualRegister - Adds a virtual register to the set of all used
- /// registers in the function.
- void addVirtualRegister(const TargetRegisterClass *TRC, unsigned Reg) {
- std::string name;
-
- // Do not count registers that are argument/return registers.
- if (!RegRets.count(Reg) && !RegArgs.count(Reg)) {
- UsedRegs[TRC].push_back(Reg);
- if (TRC == PTX::RegPredRegisterClass)
- name = "%p";
- else if (TRC == PTX::RegI16RegisterClass)
- name = "%rh";
- else if (TRC == PTX::RegI32RegisterClass)
- name = "%r";
- else if (TRC == PTX::RegI64RegisterClass)
- name = "%rd";
- else if (TRC == PTX::RegF32RegisterClass)
- name = "%f";
- else if (TRC == PTX::RegF64RegisterClass)
- name = "%fd";
- else
- llvm_unreachable("Invalid register class");
-
- name += utostr(UsedRegs[TRC].size() - 1);
- RegNames[Reg] = name;
- }
}
/// getRegisterName - Returns the name of the specified virtual register. This
/// name is used during PTX emission.
- const char *getRegisterName(unsigned Reg) const {
- if (RegNames.count(Reg))
- return RegNames.find(Reg)->second.c_str();
+ std::string getRegisterName(unsigned Reg) const {
+ if (RegInfo.count(Reg)) {
+ const RegisterInfo& RI = RegInfo.lookup(Reg);
+ std::string Name;
+ raw_string_ostream NameStr(Name);
+ decodeRegisterName(NameStr, RI.Encoded);
+ NameStr.flush();
+ return Name;
+ }
else if (Reg == PTX::NoRegister)
return "%noreg";
else
llvm_unreachable("Register not in register name map");
}
- /// getNumRegistersForClass - Returns the number of virtual registers that are
- /// used for the specified register class.
- unsigned getNumRegistersForClass(const TargetRegisterClass *TRC) const {
- return UsedRegs.lookup(TRC).size();
+ /// getEncodedRegisterName - Returns the name of the encoded register.
+ std::string getEncodedRegisterName(unsigned EncodedReg) const {
+ std::string Name;
+ raw_string_ostream NameStr(Name);
+ decodeRegisterName(NameStr, EncodedReg);
+ NameStr.flush();
+ return Name;
+ }
+
+ /// getRegisterType - Returns the type of the specified virtual register.
+ unsigned getRegisterType(unsigned Reg) const {
+ if (RegInfo.count(Reg))
+ return RegInfo.lookup(Reg).Type;
+ else
+ llvm_unreachable("Unknown register");
}
/// getOffsetForRegister - Returns the offset of the virtual register
- unsigned getOffsetForRegister(const TargetRegisterClass *TRC,
- unsigned Reg) const {
- const RegisterList &RegList = UsedRegs.lookup(TRC);
- for (unsigned i = 0, e = RegList.size(); i != e; ++i) {
- if (RegList[i] == Reg)
- return i;
- }
- //llvm_unreachable("Unknown virtual register");
- return 0;
+ unsigned getOffsetForRegister(unsigned Reg) const {
+ if (RegInfo.count(Reg))
+ return RegInfo.lookup(Reg).Offset;
+ else
+ return 0;
}
/// getFrameSymbol - Returns the symbol name for the given FrameIndex.
Modified: llvm/trunk/test/CodeGen/PTX/mov.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/mov.ll?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/mov.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/mov.ll Tue Dec 6 11:39:48 2011
@@ -31,31 +31,31 @@
}
define ptx_device i16 @t2_u16(i16 %x) {
-; CHECK: mov.b16 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.b16 %ret{{[0-9]+}}, %arg{{[0-9]+}};
; CHECK: ret;
ret i16 %x
}
define ptx_device i32 @t2_u32(i32 %x) {
-; CHECK: mov.b32 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.b32 %ret{{[0-9]+}}, %arg{{[0-9]+}};
; CHECK: ret;
ret i32 %x
}
define ptx_device i64 @t2_u64(i64 %x) {
-; CHECK: mov.b64 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.b64 %ret{{[0-9]+}}, %arg{{[0-9]+}};
; CHECK: ret;
ret i64 %x
}
define ptx_device float @t3_f32(float %x) {
-; CHECK: mov.f32 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.f32 %ret{{[0-9]+}}, %arg{{[0-9]+}};
; CHECK: ret;
ret float %x
}
define ptx_device double @t3_f64(double %x) {
-; CHECK: mov.f64 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.f64 %ret{{[0-9]+}}, %arg{{[0-9]+}};
; CHECK: ret;
ret double %x
}
Modified: llvm/trunk/test/CodeGen/PTX/parameter-order.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/parameter-order.ll?rev=145947&r1=145946&r2=145947&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/parameter-order.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/parameter-order.ll Tue Dec 6 11:39:48 2011
@@ -1,6 +1,6 @@
; RUN: llc < %s -march=ptx32 | FileCheck %s
-; CHECK: .func (.reg .b32 %ret{{[0-9]+}}) test_parameter_order (.reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}})
+; CHECK: .func (.reg .b32 %ret{{[0-9]+}}) test_parameter_order (.reg .f32 %arg{{[0-9]+}}, .reg .b32 %arg{{[0-9]+}}, .reg .b32 %arg{{[0-9]+}}, .reg .f32 %arg{{[0-9]+}})
define ptx_device i32 @test_parameter_order(float %a, i32 %b, i32 %c, float %d) {
; CHECK: sub.u32 %ret{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}
%result = sub i32 %b, %c
More information about the llvm-commits
mailing list