Page MenuHomePhabricator
Differential D38978 Diff 129093 lib/Target/NVPTX/NVPTXAsmPrinter.cpp

Changeset View

Standalone View

View Options

lib/Target/NVPTX/NVPTXAsmPrinter.cpp

Show First 20 Lines • Show All 86 Lines • ▼ Show 20 Lines
#include <sstream> #include <sstream>
#include <string> #include <string>
#include <utility> #include <utility>
#include <vector> #include <vector>
using namespace llvm; using namespace llvm;
#define DEPOTNAME "__local_depot" #define DEPOTNAME "__local_depot"
#define SHARED_DEPOTNAME "__shared_depot"
static cl::opt<bool> static cl::opt<bool>
EmitLineNumbers("nvptx-emit-line-numbers", cl::Hidden, EmitLineNumbers("nvptx-emit-line-numbers", cl::Hidden,
cl::desc("NVPTX Specific: Emit Line numbers even without -G"), cl::desc("NVPTX Specific: Emit Line numbers even without -G"),
cl::init(true)); cl::init(true));
static cl::opt<bool> static cl::opt<bool>
InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, cl::Hidden, InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, cl::Hidden,
▲ Show 20 Lines • Show All 1,613 Lines • ▼ Show 20 Linesvoid NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
const MachineFunction &MF) { const MachineFunction &MF) {
SmallString<128> Str; SmallString<128> Str;
raw_svector_ostream O(Str); raw_svector_ostream O(Str);
// Map the global virtual register number to a register class specific // Map the global virtual register number to a register class specific
// virtual register number starting from 1 with that class. // virtual register number starting from 1 with that class.
const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo(); const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo();
//unsigned numRegClasses = TRI->getNumRegClasses(); //unsigned numRegClasses = TRI->getNumRegClasses();
bool IsKernelFunction = isKernelFunction(MF.getFunction());
bool GenerateSharedDepot =
MF.getFunction().hasFnAttribute("has-nvptx-shared-depot");
// Emit the Fake Stack Object // Emit the Fake Stack Object
const MachineFrameInfo &MFI = MF.getFrameInfo(); const MachineFrameInfo &MFI = MF.getFrameInfo();
int NumBytes = (int) MFI.getStackSize(); int NumBytes = (int) MFI.getStackSize();
if (NumBytes) { if (NumBytes) {
O << "\t.local .align " << MFI.getMaxAlignment() << " .b8 \t" << DEPOTNAME O << "\t.local .align " << MFI.getMaxAlignment() << " .b8 \t" << DEPOTNAME
<< getFunctionNumber() << "[" << NumBytes << "];\n"; << getFunctionNumber() << "[" << NumBytes << "];\n";
if (IsKernelFunction && GenerateSharedDepot) {
O << "\t.shared .align " << MFI.getMaxAlignment()
hfinkelUnsubmitted
Done
ReplyInline Actions

Line too long.

hfinkel: Line too long.
<< " .b8 \t" << SHARED_DEPOTNAME << getFunctionNumber()
<< "[" << NumBytes << "];\n";
}
if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) { if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
O << "\t.reg .b64 \t%SP;\n"; O << "\t.reg .b64 \t%SP;\n";
O << "\t.reg .b64 \t%SPL;\n"; O << "\t.reg .b64 \t%SPL;\n";
if (IsKernelFunction && GenerateSharedDepot) {
O << "\t.reg .b64 \t%SPS;\n";
O << "\t.reg .b64 \t%SPSH;\n";
}
} else { } else {
O << "\t.reg .b32 \t%SP;\n"; O << "\t.reg .b32 \t%SP;\n";
O << "\t.reg .b32 \t%SPL;\n"; O << "\t.reg .b32 \t%SPL;\n";
if (IsKernelFunction && GenerateSharedDepot) {
O << "\t.reg .b32 \t%SPS;\n";
O << "\t.reg .b32 \t%SPSH;\n";
traUnsubmitted
Done
ReplyInline Actions

Nit: the name should end with S as the L in SPL was for 'local' address space. which then gets converted to generic AS. In your case it will be in shared space, hence S would be more appropriate.

tra: Nit: the name should end with `S` as the L in `SPL` was for 'local' address space. which then…
}
} }
} }
// Go through all virtual registers to establish the mapping between the // Go through all virtual registers to establish the mapping between the
// global virtual // global virtual
// register number and the per class virtual register number. // register number and the per class virtual register number.
// We use the per class virtual register number in the ptx output. // We use the per class virtual register number in the ptx output.
unsigned int numVRs = MRI->getNumVirtRegs(); unsigned int numVRs = MRI->getNumVirtRegs();
▲ Show 20 Lines • Show All 608 Lines • ▼ Show 20 Lines
void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum, void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum,
raw_ostream &O, const char *Modifier) { raw_ostream &O, const char *Modifier) {
const MachineOperand &MO = MI->getOperand(opNum); const MachineOperand &MO = MI->getOperand(opNum);
switch (MO.getType()) { switch (MO.getType()) {
case MachineOperand::MO_Register: case MachineOperand::MO_Register:
if (TargetRegisterInfo::isPhysicalRegister(MO.getReg())) { if (TargetRegisterInfo::isPhysicalRegister(MO.getReg())) {
if (MO.getReg() == NVPTX::VRDepot) if (MO.getReg() == NVPTX::VRDepot)
O << DEPOTNAME << getFunctionNumber(); O << DEPOTNAME << getFunctionNumber();
else if (MO.getReg() == NVPTX::VRSharedDepot)
O << SHARED_DEPOTNAME << getFunctionNumber();
else else
O << NVPTXInstPrinter::getRegisterName(MO.getReg()); O << NVPTXInstPrinter::getRegisterName(MO.getReg());
} else { } else {
emitVirtualRegister(MO.getReg(), O); emitVirtualRegister(MO.getReg(), O);
} }
return; return;
case MachineOperand::MO_Immediate: case MachineOperand::MO_Immediate:
▲ Show 20 Lines • Show All 85 Lines • Show Last 20 Lines