X-Git-Url: http://plrg.eecs.uci.edu/git/?a=blobdiff_plain;f=lib%2FTarget%2FNVPTX%2FNVPTXAsmPrinter.cpp;h=3beee65b7ef09da840373af2023c38f52577245c;hb=579cebfb15c5f80cc8bbc7d51da9f7827424125a;hp=31ab68158ca6cc953ec0706ccc89c09f4025f4d2;hpb=2085d00d09f4f3678a6c67da46df3a04e31b499c;p=oota-llvm.git diff --git a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index 31ab68158ca..3beee65b7ef 100644 --- a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -13,70 +13,64 @@ //===----------------------------------------------------------------------===// #include "NVPTXAsmPrinter.h" +#include "InstPrinter/NVPTXInstPrinter.h" +#include "MCTargetDesc/NVPTXMCAsmInfo.h" #include "NVPTX.h" #include "NVPTXInstrInfo.h" -#include "NVPTXTargetMachine.h" +#include "NVPTXMCExpr.h" +#include "NVPTXMachineFunctionInfo.h" #include "NVPTXRegisterInfo.h" +#include "NVPTXTargetMachine.h" #include "NVPTXUtilities.h" -#include "MCTargetDesc/NVPTXMCAsmInfo.h" -#include "NVPTXNumRegisters.h" +#include "cl_common_defines.h" #include "llvm/ADT/StringExtras.h" -#include "llvm/DebugInfo.h" -#include "llvm/Function.h" -#include "llvm/GlobalVariable.h" -#include "llvm/Module.h" +#include "llvm/Analysis/ConstantFolding.h" #include "llvm/CodeGen/Analysis.h" -#include "llvm/CodeGen/MachineRegisterInfo.h" #include "llvm/CodeGen/MachineFrameInfo.h" +#include "llvm/CodeGen/MachineLoopInfo.h" #include "llvm/CodeGen/MachineModuleInfo.h" +#include "llvm/CodeGen/MachineRegisterInfo.h" +#include "llvm/IR/DebugInfo.h" +#include "llvm/IR/DerivedTypes.h" +#include "llvm/IR/Function.h" +#include "llvm/IR/GlobalVariable.h" +#include "llvm/IR/Mangler.h" +#include "llvm/IR/Module.h" +#include "llvm/IR/Operator.h" #include "llvm/MC/MCStreamer.h" #include "llvm/MC/MCSymbol.h" -#include "llvm/Target/Mangler.h" -#include "llvm/Target/TargetLoweringObjectFile.h" -#include "llvm/Support/TargetRegistry.h" +#include "llvm/Support/CommandLine.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/FormattedStream.h" -#include "llvm/DerivedTypes.h" -#include "llvm/Support/TimeValue.h" -#include "llvm/Support/CommandLine.h" -#include "llvm/Analysis/ConstantFolding.h" #include "llvm/Support/Path.h" -#include "llvm/Assembly/Writer.h" -#include "cl_common_defines.h" +#include "llvm/Support/TargetRegistry.h" +#include "llvm/Support/TimeValue.h" +#include "llvm/Target/TargetLoweringObjectFile.h" +#include "llvm/Transforms/Utils/UnrollLoop.h" #include using namespace llvm; - -#include "NVPTXGenAsmWriter.inc" - -bool RegAllocNilUsed = true; - #define DEPOTNAME "__local_depot" static cl::opt -EmitLineNumbers("nvptx-emit-line-numbers", +EmitLineNumbers("nvptx-emit-line-numbers", cl::Hidden, cl::desc("NVPTX Specific: Emit Line numbers even without -G"), cl::init(true)); -namespace llvm { -bool InterleaveSrcInPtx = false; -} - -static cl::optInterleaveSrc("nvptx-emit-src", - cl::ZeroOrMore, - cl::desc("NVPTX Specific: Emit source line in ptx file"), - cl::location(llvm::InterleaveSrcInPtx)); - +static cl::opt +InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, cl::Hidden, + cl::desc("NVPTX Specific: Emit source line in ptx file"), + cl::init(false)); namespace { /// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V /// depends. -void DiscoverDependentGlobals(Value *V, - DenseSet &Globals) { - if (GlobalVariable *GV = dyn_cast(V)) +void DiscoverDependentGlobals(const Value *V, + DenseSet &Globals) { + if (const GlobalVariable *GV = dyn_cast(V)) Globals.insert(GV); else { - if (User *U = dyn_cast(V)) { + if (const User *U = dyn_cast(V)) { for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) { DiscoverDependentGlobals(U->getOperand(i), Globals); } @@ -87,27 +81,26 @@ void DiscoverDependentGlobals(Value *V, /// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable /// instances to be emitted, but only after any dependents have been added /// first. -void VisitGlobalVariableForEmission(GlobalVariable *GV, - SmallVectorImpl &Order, - DenseSet &Visited, - DenseSet &Visiting) { +void VisitGlobalVariableForEmission( + const GlobalVariable *GV, SmallVectorImpl &Order, + DenseSet &Visited, + DenseSet &Visiting) { // Have we already visited this one? - if (Visited.count(GV)) return; + if (Visited.count(GV)) + return; // Do we have a circular dependency? - if (Visiting.count(GV)) + if (!Visiting.insert(GV).second) report_fatal_error("Circular dependency found in global variable set"); - // Start visiting this global - Visiting.insert(GV); - // Make sure we visit all dependents first - DenseSet Others; + DenseSet Others; for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i) DiscoverDependentGlobals(GV->getOperand(i), Others); - - for (DenseSet::iterator I = Others.begin(), - E = Others.end(); I != E; ++I) + + for (DenseSet::iterator I = Others.begin(), + E = Others.end(); + I != E; ++I) VisitGlobalVariableForEmission(*I, Order, Visited, Visiting); // Now we can visit ourself @@ -117,144 +110,7 @@ void VisitGlobalVariableForEmission(GlobalVariable *GV, } } -// @TODO: This is a copy from AsmPrinter.cpp. The function is static, so we -// cannot just link to the existing version. -/// LowerConstant - Lower the specified LLVM Constant to an MCExpr. -/// -using namespace nvptx; -const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { - MCContext &Ctx = AP.OutContext; - - if (CV->isNullValue() || isa(CV)) - return MCConstantExpr::Create(0, Ctx); - - if (const ConstantInt *CI = dyn_cast(CV)) - return MCConstantExpr::Create(CI->getZExtValue(), Ctx); - - if (const GlobalValue *GV = dyn_cast(CV)) - return MCSymbolRefExpr::Create(AP.Mang->getSymbol(GV), Ctx); - - if (const BlockAddress *BA = dyn_cast(CV)) - return MCSymbolRefExpr::Create(AP.GetBlockAddressSymbol(BA), Ctx); - - const ConstantExpr *CE = dyn_cast(CV); - if (CE == 0) - llvm_unreachable("Unknown constant value to lower!"); - - - switch (CE->getOpcode()) { - default: - // If the code isn't optimized, there may be outstanding folding - // opportunities. Attempt to fold the expression using DataLayout as a - // last resort before giving up. - if (Constant *C = - ConstantFoldConstantExpression(CE, AP.TM.getDataLayout())) - if (C != CE) - return LowerConstant(C, AP); - - // Otherwise report the problem to the user. - { - std::string S; - raw_string_ostream OS(S); - OS << "Unsupported expression in static initializer: "; - WriteAsOperand(OS, CE, /*PrintType=*/false, - !AP.MF ? 0 : AP.MF->getFunction()->getParent()); - report_fatal_error(OS.str()); - } - case Instruction::GetElementPtr: { - const DataLayout &TD = *AP.TM.getDataLayout(); - // Generate a symbolic expression for the byte address - const Constant *PtrVal = CE->getOperand(0); - SmallVector IdxVec(CE->op_begin()+1, CE->op_end()); - int64_t Offset = TD.getIndexedOffset(PtrVal->getType(), IdxVec); - - const MCExpr *Base = LowerConstant(CE->getOperand(0), AP); - if (Offset == 0) - return Base; - - // Truncate/sext the offset to the pointer size. - if (TD.getPointerSizeInBits() != 64) { - int SExtAmount = 64-TD.getPointerSizeInBits(); - Offset = (Offset << SExtAmount) >> SExtAmount; - } - - return MCBinaryExpr::CreateAdd(Base, MCConstantExpr::Create(Offset, Ctx), - Ctx); - } - - case Instruction::Trunc: - // We emit the value and depend on the assembler to truncate the generated - // expression properly. This is important for differences between - // blockaddress labels. Since the two labels are in the same function, it - // is reasonable to treat their delta as a 32-bit value. - // FALL THROUGH. - case Instruction::BitCast: - return LowerConstant(CE->getOperand(0), AP); - - case Instruction::IntToPtr: { - const DataLayout &TD = *AP.TM.getDataLayout(); - // Handle casts to pointers by changing them into casts to the appropriate - // integer type. This promotes constant folding and simplifies this code. - Constant *Op = CE->getOperand(0); - Op = ConstantExpr::getIntegerCast(Op, TD.getIntPtrType(CV->getContext()), - false/*ZExt*/); - return LowerConstant(Op, AP); - } - - case Instruction::PtrToInt: { - const DataLayout &TD = *AP.TM.getDataLayout(); - // Support only foldable casts to/from pointers that can be eliminated by - // changing the pointer to the appropriately sized integer type. - Constant *Op = CE->getOperand(0); - Type *Ty = CE->getType(); - - const MCExpr *OpExpr = LowerConstant(Op, AP); - - // We can emit the pointer value into this slot if the slot is an - // integer slot equal to the size of the pointer. - if (TD.getTypeAllocSize(Ty) == TD.getTypeAllocSize(Op->getType())) - return OpExpr; - - // Otherwise the pointer is smaller than the resultant integer, mask off - // the high bits so we are sure to get a proper truncation if the input is - // a constant expr. - unsigned InBits = TD.getTypeAllocSizeInBits(Op->getType()); - const MCExpr *MaskExpr = MCConstantExpr::Create(~0ULL >> (64-InBits), Ctx); - return MCBinaryExpr::CreateAnd(OpExpr, MaskExpr, Ctx); - } - - // The MC library also has a right-shift operator, but it isn't consistently - // signed or unsigned between different targets. - case Instruction::Add: - case Instruction::Sub: - case Instruction::Mul: - case Instruction::SDiv: - case Instruction::SRem: - case Instruction::Shl: - case Instruction::And: - case Instruction::Or: - case Instruction::Xor: { - const MCExpr *LHS = LowerConstant(CE->getOperand(0), AP); - const MCExpr *RHS = LowerConstant(CE->getOperand(1), AP); - switch (CE->getOpcode()) { - default: llvm_unreachable("Unknown binary operator constant cast expr"); - case Instruction::Add: return MCBinaryExpr::CreateAdd(LHS, RHS, Ctx); - case Instruction::Sub: return MCBinaryExpr::CreateSub(LHS, RHS, Ctx); - case Instruction::Mul: return MCBinaryExpr::CreateMul(LHS, RHS, Ctx); - case Instruction::SDiv: return MCBinaryExpr::CreateDiv(LHS, RHS, Ctx); - case Instruction::SRem: return MCBinaryExpr::CreateMod(LHS, RHS, Ctx); - case Instruction::Shl: return MCBinaryExpr::CreateShl(LHS, RHS, Ctx); - case Instruction::And: return MCBinaryExpr::CreateAnd(LHS, RHS, Ctx); - case Instruction::Or: return MCBinaryExpr::CreateOr (LHS, RHS, Ctx); - case Instruction::Xor: return MCBinaryExpr::CreateXor(LHS, RHS, Ctx); - } - } - } -} - - -void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) -{ +void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) { if (!EmitLineNumbers) return; if (ignoreLoc(MI)) @@ -262,7 +118,7 @@ void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) DebugLoc curLoc = MI.getDebugLoc(); - if (prevDebugLoc.isUnknown() && curLoc.isUnknown()) + if (!prevDebugLoc && !curLoc) return; if (prevDebugLoc == curLoc) @@ -270,59 +126,225 @@ void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) prevDebugLoc = curLoc; - if (curLoc.isUnknown()) + if (!curLoc) return; + auto *Scope = cast_or_null(curLoc.getScope()); + if (!Scope) + return; - const MachineFunction *MF = MI.getParent()->getParent(); - //const TargetMachine &TM = MF->getTarget(); - - const LLVMContext &ctx = MF->getFunction()->getContext(); - DIScope Scope(curLoc.getScope(ctx)); - - if (!Scope.Verify()) - return; - - StringRef fileName(Scope.getFilename()); - StringRef dirName(Scope.getDirectory()); + StringRef fileName(Scope->getFilename()); + StringRef dirName(Scope->getDirectory()); SmallString<128> FullPathName = dirName; if (!dirName.empty() && !sys::path::is_absolute(fileName)) { sys::path::append(FullPathName, fileName); - fileName = FullPathName.str(); + fileName = FullPathName; } - if (filenameMap.find(fileName.str()) == filenameMap.end()) + if (filenameMap.find(fileName) == filenameMap.end()) return; - // Emit the line from the source file. - if (llvm::InterleaveSrcInPtx) - this->emitSrcInText(fileName.str(), curLoc.getLine()); + if (InterleaveSrc) + this->emitSrcInText(fileName, curLoc.getLine()); std::stringstream temp; - temp << "\t.loc " << filenameMap[fileName.str()] - << " " << curLoc.getLine() << " " << curLoc.getCol(); - OutStreamer.EmitRawText(Twine(temp.str().c_str())); + temp << "\t.loc " << filenameMap[fileName] << " " << curLoc.getLine() + << " " << curLoc.getCol(); + OutStreamer->EmitRawText(temp.str()); } void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) { SmallString<128> Str; raw_svector_ostream OS(Str); - if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) + if (static_cast(TM).getDrvInterface() == NVPTX::CUDA) emitLineNumberAsDotLoc(*MI); - printInstruction(MI, OS); - OutStreamer.EmitRawText(OS.str()); + + MCInst Inst; + lowerToMCInst(MI, Inst); + EmitToStreamer(*OutStreamer, Inst); +} + +// Handle symbol backtracking for targets that do not support image handles +bool NVPTXAsmPrinter::lowerImageHandleOperand(const MachineInstr *MI, + unsigned OpNo, MCOperand &MCOp) { + const MachineOperand &MO = MI->getOperand(OpNo); + const MCInstrDesc &MCID = MI->getDesc(); + + if (MCID.TSFlags & NVPTXII::IsTexFlag) { + // This is a texture fetch, so operand 4 is a texref and operand 5 is + // a samplerref + if (OpNo == 4 && MO.isImm()) { + lowerImageHandleSymbol(MO.getImm(), MCOp); + return true; + } + if (OpNo == 5 && MO.isImm() && !(MCID.TSFlags & NVPTXII::IsTexModeUnifiedFlag)) { + lowerImageHandleSymbol(MO.getImm(), MCOp); + return true; + } + + return false; + } else if (MCID.TSFlags & NVPTXII::IsSuldMask) { + unsigned VecSize = + 1 << (((MCID.TSFlags & NVPTXII::IsSuldMask) >> NVPTXII::IsSuldShift) - 1); + + // For a surface load of vector size N, the Nth operand will be the surfref + if (OpNo == VecSize && MO.isImm()) { + lowerImageHandleSymbol(MO.getImm(), MCOp); + return true; + } + + return false; + } else if (MCID.TSFlags & NVPTXII::IsSustFlag) { + // This is a surface store, so operand 0 is a surfref + if (OpNo == 0 && MO.isImm()) { + lowerImageHandleSymbol(MO.getImm(), MCOp); + return true; + } + + return false; + } else if (MCID.TSFlags & NVPTXII::IsSurfTexQueryFlag) { + // This is a query, so operand 1 is a surfref/texref + if (OpNo == 1 && MO.isImm()) { + lowerImageHandleSymbol(MO.getImm(), MCOp); + return true; + } + + return false; + } + + return false; +} + +void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) { + // Ewwww + TargetMachine &TM = const_cast(MF->getTarget()); + NVPTXTargetMachine &nvTM = static_cast(TM); + const NVPTXMachineFunctionInfo *MFI = MF->getInfo(); + const char *Sym = MFI->getImageHandleSymbol(Index); + std::string *SymNamePtr = + nvTM.getManagedStrPool()->getManagedString(Sym); + MCOp = GetSymbolRef(OutContext.GetOrCreateSymbol( + StringRef(SymNamePtr->c_str()))); +} + +void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) { + OutMI.setOpcode(MI->getOpcode()); + // Special: Do not mangle symbol operand of CALL_PROTOTYPE + if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) { + const MachineOperand &MO = MI->getOperand(0); + OutMI.addOperand(GetSymbolRef( + OutContext.GetOrCreateSymbol(Twine(MO.getSymbolName())))); + return; + } + + for (unsigned i = 0, e = MI->getNumOperands(); i != e; ++i) { + const MachineOperand &MO = MI->getOperand(i); + + MCOperand MCOp; + if (!nvptxSubtarget->hasImageHandles()) { + if (lowerImageHandleOperand(MI, i, MCOp)) { + OutMI.addOperand(MCOp); + continue; + } + } + + if (lowerOperand(MO, MCOp)) + OutMI.addOperand(MCOp); + } +} + +bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO, + MCOperand &MCOp) { + switch (MO.getType()) { + default: llvm_unreachable("unknown operand type"); + case MachineOperand::MO_Register: + MCOp = MCOperand::CreateReg(encodeVirtualRegister(MO.getReg())); + break; + case MachineOperand::MO_Immediate: + MCOp = MCOperand::CreateImm(MO.getImm()); + break; + case MachineOperand::MO_MachineBasicBlock: + MCOp = MCOperand::CreateExpr(MCSymbolRefExpr::Create( + MO.getMBB()->getSymbol(), OutContext)); + break; + case MachineOperand::MO_ExternalSymbol: + MCOp = GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName())); + break; + case MachineOperand::MO_GlobalAddress: + MCOp = GetSymbolRef(getSymbol(MO.getGlobal())); + break; + case MachineOperand::MO_FPImmediate: { + const ConstantFP *Cnt = MO.getFPImm(); + APFloat Val = Cnt->getValueAPF(); + + switch (Cnt->getType()->getTypeID()) { + default: report_fatal_error("Unsupported FP type"); break; + case Type::FloatTyID: + MCOp = MCOperand::CreateExpr( + NVPTXFloatMCExpr::CreateConstantFPSingle(Val, OutContext)); + break; + case Type::DoubleTyID: + MCOp = MCOperand::CreateExpr( + NVPTXFloatMCExpr::CreateConstantFPDouble(Val, OutContext)); + break; + } + break; + } + } + return true; +} + +unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) { + if (TargetRegisterInfo::isVirtualRegister(Reg)) { + const TargetRegisterClass *RC = MRI->getRegClass(Reg); + + DenseMap &RegMap = VRegMapping[RC]; + unsigned RegNum = RegMap[Reg]; + + // Encode the register class in the upper 4 bits + // Must be kept in sync with NVPTXInstPrinter::printRegName + unsigned Ret = 0; + if (RC == &NVPTX::Int1RegsRegClass) { + Ret = (1 << 28); + } else if (RC == &NVPTX::Int16RegsRegClass) { + Ret = (2 << 28); + } else if (RC == &NVPTX::Int32RegsRegClass) { + Ret = (3 << 28); + } else if (RC == &NVPTX::Int64RegsRegClass) { + Ret = (4 << 28); + } else if (RC == &NVPTX::Float32RegsRegClass) { + Ret = (5 << 28); + } else if (RC == &NVPTX::Float64RegsRegClass) { + Ret = (6 << 28); + } else { + report_fatal_error("Bad register class"); + } + + // Insert the vreg number + Ret |= (RegNum & 0x0FFFFFFF); + return Ret; + } else { + // Some special-use registers are actually physical registers. + // Encode this as the register class ID of 0 and the real register ID. + return Reg & 0x0FFFFFFF; + } } -void NVPTXAsmPrinter::printReturnValStr(const Function *F, - raw_ostream &O) -{ +MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) { + const MCExpr *Expr; + Expr = MCSymbolRefExpr::Create(Symbol, MCSymbolRefExpr::VK_None, + OutContext); + return MCOperand::CreateExpr(Expr); +} + +void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) { const DataLayout *TD = TM.getDataLayout(); - const TargetLowering *TLI = TM.getTargetLowering(); + const TargetLowering *TLI = nvptxSubtarget->getTargetLowering(); Type *Ty = F->getReturnType(); - bool isABI = (nvptxSubtarget.getSmVersion() >= 20); + bool isABI = (nvptxSubtarget->getSmVersion() >= 20); if (Ty->getTypeID() == Type::VoidTyID) return; @@ -330,57 +352,35 @@ void NVPTXAsmPrinter::printReturnValStr(const Function *F, O << " ("; if (isABI) { - if (Ty->isPrimitiveType() || Ty->isIntegerTy()) { + if (Ty->isFloatingPointTy() || Ty->isIntegerTy()) { unsigned size = 0; if (const IntegerType *ITy = dyn_cast(Ty)) { size = ITy->getBitWidth(); - if (size < 32) size = 32; + if (size < 32) + size = 32; } else { - assert(Ty->isFloatingPointTy() && - "Floating point type expected here"); + assert(Ty->isFloatingPointTy() && "Floating point type expected here"); size = Ty->getPrimitiveSizeInBits(); } O << ".param .b" << size << " func_retval0"; - } - else if (isa(Ty)) { + } else if (isa(Ty)) { O << ".param .b" << TLI->getPointerTy().getSizeInBits() - << " func_retval0"; - } else { - if ((Ty->getTypeID() == Type::StructTyID) || - isa(Ty)) { - SmallVector vtparts; - ComputeValueVTs(*TLI, Ty, vtparts); - unsigned totalsz = 0; - for (unsigned i=0,e=vtparts.size(); i!=e; ++i) { - unsigned elems = 1; - EVT elemtype = vtparts[i]; - if (vtparts[i].isVector()) { - elems = vtparts[i].getVectorNumElements(); - elemtype = vtparts[i].getVectorElementType(); - } - for (unsigned j=0, je=elems; j!=je; ++j) { - unsigned sz = elemtype.getSizeInBits(); - if (elemtype.isInteger() && (sz < 8)) sz = 8; - totalsz += sz/8; - } - } - unsigned retAlignment = 0; - if (!llvm::getAlign(*F, 0, retAlignment)) - retAlignment = TD->getABITypeAlignment(Ty); - O << ".param .align " - << retAlignment - << " .b8 func_retval0[" - << totalsz << "]"; - } else - assert(false && - "Unknown return type"); - } + << " func_retval0"; + } else if ((Ty->getTypeID() == Type::StructTyID) || isa(Ty)) { + unsigned totalsz = TD->getTypeAllocSize(Ty); + unsigned retAlignment = 0; + if (!llvm::getAlign(*F, 0, retAlignment)) + retAlignment = TD->getABITypeAlignment(Ty); + O << ".param .align " << retAlignment << " .b8 func_retval0[" << totalsz + << "]"; + } else + llvm_unreachable("Unknown return type"); } else { SmallVector vtparts; ComputeValueVTs(*TLI, Ty, vtparts); unsigned idx = 0; - for (unsigned i=0,e=vtparts.size(); i!=e; ++i) { + for (unsigned i = 0, e = vtparts.size(); i != e; ++i) { unsigned elems = 1; EVT elemtype = vtparts[i]; if (vtparts[i].isVector()) { @@ -388,14 +388,16 @@ void NVPTXAsmPrinter::printReturnValStr(const Function *F, elemtype = vtparts[i].getVectorElementType(); } - for (unsigned j=0, je=elems; j!=je; ++j) { + for (unsigned j = 0, je = elems; j != je; ++j) { unsigned sz = elemtype.getSizeInBits(); - if (elemtype.isInteger() && (sz < 32)) sz = 32; + if (elemtype.isInteger() && (sz < 32)) + sz = 32; O << ".reg .b" << sz << " func_retval" << idx; - if (j(); + // TODO: isLoopHeader() should take "const MachineBasicBlock *". + // We insert .pragma "nounroll" only to the loop header. + if (!LI.isLoopHeader(const_cast(&MBB))) + return false; + + // llvm.loop.unroll.disable is marked on the back edges of a loop. Therefore, + // we iterate through each back edge of the loop with header MBB, and check + // whether its metadata contains llvm.loop.unroll.disable. + for (auto I = MBB.pred_begin(); I != MBB.pred_end(); ++I) { + const MachineBasicBlock *PMBB = *I; + if (LI.getLoopFor(PMBB) != LI.getLoopFor(&MBB)) { + // Edges from other loops to MBB are not back edges. + continue; + } + if (const BasicBlock *PBB = PMBB->getBasicBlock()) { + if (MDNode *LoopID = PBB->getTerminator()->getMetadata("llvm.loop")) { + if (GetUnrollMetadata(LoopID, "llvm.loop.unroll.disable")) + return true; + } + } + } + return false; +} + +void NVPTXAsmPrinter::EmitBasicBlockStart(const MachineBasicBlock &MBB) const { + AsmPrinter::EmitBasicBlockStart(MBB); + if (isLoopHeaderOfNoUnroll(MBB)) + OutStreamer->EmitRawText(StringRef("\t.pragma \"nounroll\";\n")); +} + void NVPTXAsmPrinter::EmitFunctionEntryLabel() { SmallString<128> Str; raw_svector_ostream O(Str); + if (!GlobalsEmitted) { + emitGlobals(*MF->getFunction()->getParent()); + GlobalsEmitted = true; + } + // Set up MRI = &MF->getRegInfo(); F = MF->getFunction(); - emitLinkageDirective(F,O); + emitLinkageDirective(F, O); if (llvm::isKernelFunction(*F)) O << ".entry "; else { @@ -431,314 +474,182 @@ void NVPTXAsmPrinter::EmitFunctionEntryLabel() { if (llvm::isKernelFunction(*F)) emitKernelFunctionDirectives(*F, O); - OutStreamer.EmitRawText(O.str()); + OutStreamer->EmitRawText(O.str()); prevDebugLoc = DebugLoc(); } void NVPTXAsmPrinter::EmitFunctionBodyStart() { - const TargetRegisterInfo &TRI = *TM.getRegisterInfo(); - unsigned numRegClasses = TRI.getNumRegClasses(); - VRidGlobal2LocalMap = new std::map[numRegClasses+1]; - OutStreamer.EmitRawText(StringRef("{\n")); + VRegMapping.clear(); + OutStreamer->EmitRawText(StringRef("{\n")); setAndEmitFunctionVirtualRegisters(*MF); SmallString<128> Str; raw_svector_ostream O(Str); emitDemotedVars(MF->getFunction(), O); - OutStreamer.EmitRawText(O.str()); + OutStreamer->EmitRawText(O.str()); } void NVPTXAsmPrinter::EmitFunctionBodyEnd() { - OutStreamer.EmitRawText(StringRef("}\n")); - delete []VRidGlobal2LocalMap; + OutStreamer->EmitRawText(StringRef("}\n")); + VRegMapping.clear(); } +void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const { + unsigned RegNo = MI->getOperand(0).getReg(); + if (TargetRegisterInfo::isVirtualRegister(RegNo)) { + OutStreamer->AddComment(Twine("implicit-def: ") + + getVirtualRegisterName(RegNo)); + } else { + OutStreamer->AddComment(Twine("implicit-def: ") + + nvptxSubtarget->getRegisterInfo()->getName(RegNo)); + } + OutStreamer->AddBlankLine(); +} -void -NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function& F, - raw_ostream &O) const { +void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, + raw_ostream &O) const { // If the NVVM IR has some of reqntid* specified, then output // the reqntid directive, and set the unspecified ones to 1. // If none of reqntid* is specified, don't output reqntid directive. unsigned reqntidx, reqntidy, reqntidz; bool specified = false; - if (llvm::getReqNTIDx(F, reqntidx) == false) reqntidx = 1; - else specified = true; - if (llvm::getReqNTIDy(F, reqntidy) == false) reqntidy = 1; - else specified = true; - if (llvm::getReqNTIDz(F, reqntidz) == false) reqntidz = 1; - else specified = true; + if (!llvm::getReqNTIDx(F, reqntidx)) + reqntidx = 1; + else + specified = true; + if (!llvm::getReqNTIDy(F, reqntidy)) + reqntidy = 1; + else + specified = true; + if (!llvm::getReqNTIDz(F, reqntidz)) + reqntidz = 1; + else + specified = true; if (specified) - O << ".reqntid " << reqntidx << ", " - << reqntidy << ", " << reqntidz << "\n"; + O << ".reqntid " << reqntidx << ", " << reqntidy << ", " << reqntidz + << "\n"; // If the NVVM IR has some of maxntid* specified, then output // the maxntid directive, and set the unspecified ones to 1. // If none of maxntid* is specified, don't output maxntid directive. unsigned maxntidx, maxntidy, maxntidz; specified = false; - if (llvm::getMaxNTIDx(F, maxntidx) == false) maxntidx = 1; - else specified = true; - if (llvm::getMaxNTIDy(F, maxntidy) == false) maxntidy = 1; - else specified = true; - if (llvm::getMaxNTIDz(F, maxntidz) == false) maxntidz = 1; - else specified = true; + if (!llvm::getMaxNTIDx(F, maxntidx)) + maxntidx = 1; + else + specified = true; + if (!llvm::getMaxNTIDy(F, maxntidy)) + maxntidy = 1; + else + specified = true; + if (!llvm::getMaxNTIDz(F, maxntidz)) + maxntidz = 1; + else + specified = true; if (specified) - O << ".maxntid " << maxntidx << ", " - << maxntidy << ", " << maxntidz << "\n"; + O << ".maxntid " << maxntidx << ", " << maxntidy << ", " << maxntidz + << "\n"; unsigned mincta; if (llvm::getMinCTASm(F, mincta)) O << ".minnctapersm " << mincta << "\n"; } -void -NVPTXAsmPrinter::getVirtualRegisterName(unsigned vr, bool isVec, - raw_ostream &O) { - const TargetRegisterClass * RC = MRI->getRegClass(vr); - unsigned id = RC->getID(); +std::string +NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const { + const TargetRegisterClass *RC = MRI->getRegClass(Reg); - std::map ®map = VRidGlobal2LocalMap[id]; - unsigned mapped_vr = regmap[vr]; + std::string Name; + raw_string_ostream NameStr(Name); - if (!isVec) { - O << getNVPTXRegClassStr(RC) << mapped_vr; - return; - } - // Vector virtual register - if (getNVPTXVectorSize(RC) == 4) - O << "{" - << getNVPTXRegClassStr(RC) << mapped_vr << "_0, " - << getNVPTXRegClassStr(RC) << mapped_vr << "_1, " - << getNVPTXRegClassStr(RC) << mapped_vr << "_2, " - << getNVPTXRegClassStr(RC) << mapped_vr << "_3" - << "}"; - else if (getNVPTXVectorSize(RC) == 2) - O << "{" - << getNVPTXRegClassStr(RC) << mapped_vr << "_0, " - << getNVPTXRegClassStr(RC) << mapped_vr << "_1" - << "}"; - else - llvm_unreachable("Unsupported vector size"); + VRegRCMap::const_iterator I = VRegMapping.find(RC); + assert(I != VRegMapping.end() && "Bad register class"); + const DenseMap &RegMap = I->second; + + VRegMap::const_iterator VI = RegMap.find(Reg); + assert(VI != RegMap.end() && "Bad virtual register"); + unsigned MappedVR = VI->second; + + NameStr << getNVPTXRegClassStr(RC) << MappedVR; + + NameStr.flush(); + return Name; } -void -NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr, bool isVec, - raw_ostream &O) { - getVirtualRegisterName(vr, isVec, O); +void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr, + raw_ostream &O) { + O << getVirtualRegisterName(vr); } -void NVPTXAsmPrinter::printVecModifiedImmediate(const MachineOperand &MO, - const char *Modifier, - raw_ostream &O) { - static const char vecelem[] = {'0', '1', '2', '3', '0', '1', '2', '3'}; - int Imm = (int)MO.getImm(); - if(0 == strcmp(Modifier, "vecelem")) +void NVPTXAsmPrinter::printVecModifiedImmediate( + const MachineOperand &MO, const char *Modifier, raw_ostream &O) { + static const char vecelem[] = { '0', '1', '2', '3', '0', '1', '2', '3' }; + int Imm = (int) MO.getImm(); + if (0 == strcmp(Modifier, "vecelem")) O << "_" << vecelem[Imm]; - else if(0 == strcmp(Modifier, "vecv4comm1")) { - if((Imm < 0) || (Imm > 3)) + else if (0 == strcmp(Modifier, "vecv4comm1")) { + if ((Imm < 0) || (Imm > 3)) O << "//"; - } - else if(0 == strcmp(Modifier, "vecv4comm2")) { - if((Imm < 4) || (Imm > 7)) + } else if (0 == strcmp(Modifier, "vecv4comm2")) { + if ((Imm < 4) || (Imm > 7)) O << "//"; - } - else if(0 == strcmp(Modifier, "vecv4pos")) { - if(Imm < 0) Imm = 0; - O << "_" << vecelem[Imm%4]; - } - else if(0 == strcmp(Modifier, "vecv2comm1")) { - if((Imm < 0) || (Imm > 1)) + } else if (0 == strcmp(Modifier, "vecv4pos")) { + if (Imm < 0) + Imm = 0; + O << "_" << vecelem[Imm % 4]; + } else if (0 == strcmp(Modifier, "vecv2comm1")) { + if ((Imm < 0) || (Imm > 1)) O << "//"; - } - else if(0 == strcmp(Modifier, "vecv2comm2")) { - if((Imm < 2) || (Imm > 3)) + } else if (0 == strcmp(Modifier, "vecv2comm2")) { + if ((Imm < 2) || (Imm > 3)) O << "//"; - } - else if(0 == strcmp(Modifier, "vecv2pos")) { - if(Imm < 0) Imm = 0; - O << "_" << vecelem[Imm%2]; - } - else + } else if (0 == strcmp(Modifier, "vecv2pos")) { + if (Imm < 0) + Imm = 0; + O << "_" << vecelem[Imm % 2]; + } else llvm_unreachable("Unknown Modifier on immediate operand"); } -void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum, - raw_ostream &O, const char *Modifier) { - const MachineOperand &MO = MI->getOperand(opNum); - switch (MO.getType()) { - case MachineOperand::MO_Register: - if (TargetRegisterInfo::isPhysicalRegister(MO.getReg())) { - if (MO.getReg() == NVPTX::VRDepot) - O << DEPOTNAME << getFunctionNumber(); - else - O << getRegisterName(MO.getReg()); - } else { - if (!Modifier) - emitVirtualRegister(MO.getReg(), false, O); - else { - if (strcmp(Modifier, "vecfull") == 0) - emitVirtualRegister(MO.getReg(), true, O); - else - llvm_unreachable( - "Don't know how to handle the modifier on virtual register."); - } - } - return; - - case MachineOperand::MO_Immediate: - if (!Modifier) - O << MO.getImm(); - else if (strstr(Modifier, "vec") == Modifier) - printVecModifiedImmediate(MO, Modifier, O); - else - llvm_unreachable("Don't know how to handle modifier on immediate operand"); - return; - - case MachineOperand::MO_FPImmediate: - printFPConstant(MO.getFPImm(), O); - break; - - case MachineOperand::MO_GlobalAddress: - O << *Mang->getSymbol(MO.getGlobal()); - break; - - case MachineOperand::MO_ExternalSymbol: { - const char * symbname = MO.getSymbolName(); - if (strstr(symbname, ".PARAM") == symbname) { - unsigned index; - sscanf(symbname+6, "%u[];", &index); - printParamName(index, O); - } - else if (strstr(symbname, ".HLPPARAM") == symbname) { - unsigned index; - sscanf(symbname+9, "%u[];", &index); - O << *CurrentFnSym << "_param_" << index << "_offset"; - } - else - O << symbname; - break; - } - - case MachineOperand::MO_MachineBasicBlock: - O << *MO.getMBB()->getSymbol(); - return; - - default: - llvm_unreachable("Operand type not supported."); - } -} - -void NVPTXAsmPrinter:: -printImplicitDef(const MachineInstr *MI, raw_ostream &O) const { -#ifndef __OPTIMIZE__ - O << "\t// Implicit def :"; - //printOperand(MI, 0); - O << "\n"; -#endif -} - -void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum, - raw_ostream &O, const char *Modifier) { - printOperand(MI, opNum, O); - - if (Modifier && !strcmp(Modifier, "add")) { - O << ", "; - printOperand(MI, opNum+1, O); - } else { - if (MI->getOperand(opNum+1).isImm() && - MI->getOperand(opNum+1).getImm() == 0) - return; // don't print ',0' or '+0' - O << "+"; - printOperand(MI, opNum+1, O); - } -} -void NVPTXAsmPrinter::printLdStCode(const MachineInstr *MI, int opNum, - raw_ostream &O, const char *Modifier) -{ - if (Modifier) { - const MachineOperand &MO = MI->getOperand(opNum); - int Imm = (int)MO.getImm(); - if (!strcmp(Modifier, "volatile")) { - if (Imm) - O << ".volatile"; - } else if (!strcmp(Modifier, "addsp")) { - switch (Imm) { - case NVPTX::PTXLdStInstCode::GLOBAL: O << ".global"; break; - case NVPTX::PTXLdStInstCode::SHARED: O << ".shared"; break; - case NVPTX::PTXLdStInstCode::LOCAL: O << ".local"; break; - case NVPTX::PTXLdStInstCode::PARAM: O << ".param"; break; - case NVPTX::PTXLdStInstCode::CONSTANT: O << ".const"; break; - case NVPTX::PTXLdStInstCode::GENERIC: - if (!nvptxSubtarget.hasGenericLdSt()) - O << ".global"; - break; - default: - llvm_unreachable("Wrong Address Space"); - } - } - else if (!strcmp(Modifier, "sign")) { - if (Imm==NVPTX::PTXLdStInstCode::Signed) - O << "s"; - else if (Imm==NVPTX::PTXLdStInstCode::Unsigned) - O << "u"; - else - O << "f"; - } - else if (!strcmp(Modifier, "vec")) { - if (Imm==NVPTX::PTXLdStInstCode::V2) - O << ".v2"; - else if (Imm==NVPTX::PTXLdStInstCode::V4) - O << ".v4"; - } - else - llvm_unreachable("Unknown Modifier"); - } - else - llvm_unreachable("Empty Modifier"); -} -void NVPTXAsmPrinter::emitDeclaration (const Function *F, raw_ostream &O) { +void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) { - emitLinkageDirective(F,O); + emitLinkageDirective(F, O); if (llvm::isKernelFunction(*F)) O << ".entry "; else O << ".func "; printReturnValStr(F, O); - O << *CurrentFnSym << "\n"; + O << *getSymbol(F) << "\n"; emitFunctionParamList(F, O); O << ";\n"; } -static bool usedInGlobalVarDef(const Constant *C) -{ +static bool usedInGlobalVarDef(const Constant *C) { if (!C) return false; if (const GlobalVariable *GV = dyn_cast(C)) { - if (GV->getName().str() == "llvm.used") + if (GV->getName() == "llvm.used") return false; return true; } - for (Value::const_use_iterator ui=C->use_begin(), ue=C->use_end(); - ui!=ue; ++ui) { - const Constant *C = dyn_cast(*ui); - if (usedInGlobalVarDef(C)) - return true; - } + for (const User *U : C->users()) + if (const Constant *C = dyn_cast(U)) + if (usedInGlobalVarDef(C)) + return true; + return false; } -static bool usedInOneFunc(const User *U, Function const *&oneFunc) -{ +static bool usedInOneFunc(const User *U, Function const *&oneFunc) { if (const GlobalVariable *othergv = dyn_cast(U)) { - if (othergv->getName().str() == "llvm.used") + if (othergv->getName() == "llvm.used") return true; } @@ -749,22 +660,14 @@ static bool usedInOneFunc(const User *U, Function const *&oneFunc) return false; oneFunc = curFunc; return true; - } - else + } else return false; } - if (const MDNode *md = dyn_cast(U)) - if (md->hasName() && ((md->getName().str() == "llvm.dbg.gv") || - (md->getName().str() == "llvm.dbg.sp"))) - return true; - - - for (User::const_use_iterator ui=U->use_begin(), ue=U->use_end(); - ui!=ue; ++ui) { - if (usedInOneFunc(*ui, oneFunc) == false) + for (const User *UU : U->users()) + if (!usedInOneFunc(UU, oneFunc)) return false; - } + return true; } @@ -776,16 +679,16 @@ static bool usedInOneFunc(const User *U, Function const *&oneFunc) * 3. Is the global variable referenced only in one function? */ static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) { - if (gv->hasInternalLinkage() == false) + if (!gv->hasInternalLinkage()) return false; const PointerType *Pty = gv->getType(); if (Pty->getAddressSpace() != llvm::ADDRESS_SPACE_SHARED) return false; - const Function *oneFunc = 0; + const Function *oneFunc = nullptr; bool flag = usedInOneFunc(gv, oneFunc); - if (flag == false) + if (!flag) return false; if (!oneFunc) return false; @@ -795,16 +698,17 @@ static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) { static bool useFuncSeen(const Constant *C, llvm::DenseMap &seenMap) { - for (Value::const_use_iterator ui=C->use_begin(), ue=C->use_end(); - ui!=ue; ++ui) { - if (const Constant *cu = dyn_cast(*ui)) { + for (const User *U : C->users()) { + if (const Constant *cu = dyn_cast(U)) { if (useFuncSeen(cu, seenMap)) return true; - } else if (const Instruction *I = dyn_cast(*ui)) { + } else if (const Instruction *I = dyn_cast(U)) { const BasicBlock *bb = I->getParent(); - if (!bb) continue; + if (!bb) + continue; const Function *caller = bb->getParent(); - if (!caller) continue; + if (!caller) + continue; if (seenMap.find(caller) != seenMap.end()) return true; } @@ -812,10 +716,9 @@ static bool useFuncSeen(const Constant *C, return false; } -void NVPTXAsmPrinter::emitDeclarations (Module &M, raw_ostream &O) { +void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) { llvm::DenseMap seenMap; - for (Module::const_iterator FI=M.begin(), FE=M.end(); - FI!=FE; ++FI) { + for (Module::const_iterator FI = M.begin(), FE = M.end(); FI != FE; ++FI) { const Function *F = FI; if (F->isDeclaration()) { @@ -823,42 +726,40 @@ void NVPTXAsmPrinter::emitDeclarations (Module &M, raw_ostream &O) { continue; if (F->getIntrinsicID()) continue; - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); continue; } - for (Value::const_use_iterator iter=F->use_begin(), - iterEnd=F->use_end(); iter!=iterEnd; ++iter) { - if (const Constant *C = dyn_cast(*iter)) { + for (const User *U : F->users()) { + if (const Constant *C = dyn_cast(U)) { if (usedInGlobalVarDef(C)) { // The use is in the initialization of a global variable // that is a function pointer, so print a declaration // for the original function - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); break; } // Emit a declaration of this function if the function that // uses this constant expr has already been seen. if (useFuncSeen(C, seenMap)) { - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); break; } } - if (!isa(*iter)) continue; - const Instruction *instr = cast(*iter); + if (!isa(U)) + continue; + const Instruction *instr = cast(U); const BasicBlock *bb = instr->getParent(); - if (!bb) continue; + if (!bb) + continue; const Function *caller = bb->getParent(); - if (!caller) continue; + if (!caller) + continue; // If a caller has already been seen, then the caller is // appearing in the module before the callee. so print out // a declaration for the callee. if (seenMap.find(caller) != seenMap.end()) { - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); break; } @@ -871,42 +772,46 @@ void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) { DebugInfoFinder DbgFinder; DbgFinder.processModule(M); - unsigned i=1; - for (DebugInfoFinder::iterator I = DbgFinder.compile_unit_begin(), - E = DbgFinder.compile_unit_end(); I != E; ++I) { - DICompileUnit DIUnit(*I); - StringRef Filename(DIUnit.getFilename()); - StringRef Dirname(DIUnit.getDirectory()); + unsigned i = 1; + for (const MDCompileUnit *DIUnit : DbgFinder.compile_units()) { + StringRef Filename = DIUnit->getFilename(); + StringRef Dirname = DIUnit->getDirectory(); SmallString<128> FullPathName = Dirname; if (!Dirname.empty() && !sys::path::is_absolute(Filename)) { sys::path::append(FullPathName, Filename); - Filename = FullPathName.str(); + Filename = FullPathName; } - if (filenameMap.find(Filename.str()) != filenameMap.end()) + if (filenameMap.find(Filename) != filenameMap.end()) continue; - filenameMap[Filename.str()] = i; - OutStreamer.EmitDwarfFileDirective(i, "", Filename.str()); + filenameMap[Filename] = i; + OutStreamer->EmitDwarfFileDirective(i, "", Filename); ++i; } - for (DebugInfoFinder::iterator I = DbgFinder.subprogram_begin(), - E = DbgFinder.subprogram_end(); I != E; ++I) { - DISubprogram SP(*I); - StringRef Filename(SP.getFilename()); - StringRef Dirname(SP.getDirectory()); + for (MDSubprogram *SP : DbgFinder.subprograms()) { + StringRef Filename = SP->getFilename(); + StringRef Dirname = SP->getDirectory(); SmallString<128> FullPathName = Dirname; if (!Dirname.empty() && !sys::path::is_absolute(Filename)) { sys::path::append(FullPathName, Filename); - Filename = FullPathName.str(); + Filename = FullPathName; } - if (filenameMap.find(Filename.str()) != filenameMap.end()) + if (filenameMap.find(Filename) != filenameMap.end()) continue; - filenameMap[Filename.str()] = i; + filenameMap[Filename] = i; ++i; } } -bool NVPTXAsmPrinter::doInitialization (Module &M) { +bool NVPTXAsmPrinter::doInitialization(Module &M) { + // Construct a default subtarget off of the TargetMachine defaults. The + // rest of NVPTX isn't friendly to change subtargets per function and + // so the default TargetMachine will have all of the options. + StringRef TT = TM.getTargetTriple(); + StringRef CPU = TM.getTargetCPU(); + StringRef FS = TM.getTargetFeatureString(); + const NVPTXTargetMachine &NTM = static_cast(TM); + const NVPTXSubtarget STI(TT, CPU, FS, NTM); SmallString<128> Str1; raw_svector_ostream OS1(Str1); @@ -918,23 +823,38 @@ bool NVPTXAsmPrinter::doInitialization (Module &M) { //bool Result = AsmPrinter::doInitialization(M); // Initialize TargetLoweringObjectFile. - const_cast(getObjFileLowering()) - .Initialize(OutContext, TM); + const_cast(getObjFileLowering()) + .Initialize(OutContext, TM); - Mang = new Mangler(OutContext, *TM.getDataLayout()); + Mang = new Mangler(TM.getDataLayout()); // Emit header before any dwarf directives are emitted below. - emitHeader(M, OS1); - OutStreamer.EmitRawText(OS1.str()); - + emitHeader(M, OS1, STI); + OutStreamer->EmitRawText(OS1.str()); // Already commented out //bool Result = AsmPrinter::doInitialization(M); + // Emit module-level inline asm if it exists. + if (!M.getModuleInlineAsm().empty()) { + OutStreamer->AddComment("Start of file scope inline assembly"); + OutStreamer->AddBlankLine(); + OutStreamer->EmitRawText(StringRef(M.getModuleInlineAsm())); + OutStreamer->AddBlankLine(); + OutStreamer->AddComment("End of file scope inline assembly"); + OutStreamer->AddBlankLine(); + } - if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) + // If we're not NVCL we're CUDA, go ahead and emit filenames. + if (Triple(TM.getTargetTriple()).getOS() != Triple::NVCL) recordAndEmitFilenames(M); + GlobalsEmitted = false; + + return false; // success +} + +void NVPTXAsmPrinter::emitGlobals(const Module &M) { SmallString<128> Str2; raw_svector_ostream OS2(Str2); @@ -945,16 +865,16 @@ bool NVPTXAsmPrinter::doInitialization (Module &M) { // global variable in order, and ensure that we emit it *after* its dependent // globals. We use a little extra memory maintaining both a set and a list to // have fast searches while maintaining a strict ordering. - SmallVector Globals; - DenseSet GVVisited; - DenseSet GVVisiting; + SmallVector Globals; + DenseSet GVVisited; + DenseSet GVVisiting; // Visit each global variable, in order - for (Module::global_iterator I = M.global_begin(), E = M.global_end(); + for (Module::const_global_iterator I = M.global_begin(), E = M.global_end(); I != E; ++I) VisitGlobalVariableForEmission(I, Globals, GVVisited, GVVisiting); - assert(GVVisited.size() == M.getGlobalList().size() && + assert(GVVisited.size() == M.getGlobalList().size() && "Missed a global variable"); assert(GVVisiting.size() == 0 && "Did not fully process a global variable"); @@ -964,26 +884,27 @@ bool NVPTXAsmPrinter::doInitialization (Module &M) { OS2 << '\n'; - OutStreamer.EmitRawText(OS2.str()); - return false; // success + OutStreamer->EmitRawText(OS2.str()); } -void NVPTXAsmPrinter::emitHeader (Module &M, raw_ostream &O) { +void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O, + const NVPTXSubtarget &STI) { O << "//\n"; O << "// Generated by LLVM NVPTX Back-End\n"; O << "//\n"; O << "\n"; - unsigned PTXVersion = nvptxSubtarget.getPTXVersion(); + unsigned PTXVersion = STI.getPTXVersion(); O << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n"; O << ".target "; - O << nvptxSubtarget.getTargetName(); + O << STI.getTargetName(); - if (nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) + const NVPTXTargetMachine &NTM = static_cast(TM); + if (NTM.getDrvInterface() == NVPTX::NVCL) O << ", texmode_independent"; - if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) { - if (!nvptxSubtarget.hasDouble()) + else { + if (!STI.hasDouble()) O << ", map_f64_to_f32"; } @@ -993,7 +914,7 @@ void NVPTXAsmPrinter::emitHeader (Module &M, raw_ostream &O) { O << "\n"; O << ".address_size "; - if (nvptxSubtarget.is64Bit()) + if (NTM.is64Bit()) O << "64"; else O << "32"; @@ -1003,17 +924,24 @@ void NVPTXAsmPrinter::emitHeader (Module &M, raw_ostream &O) { } bool NVPTXAsmPrinter::doFinalization(Module &M) { + // If we did not emit any functions, then the global declarations have not + // yet been emitted. + if (!GlobalsEmitted) { + emitGlobals(M); + GlobalsEmitted = true; + } + // XXX Temproarily remove global variables so that doFinalization() will not // emit them again (global variables are emitted at beginning). Module::GlobalListType &global_list = M.getGlobalList(); int i, n = global_list.size(); - GlobalVariable **gv_array = new GlobalVariable* [n]; + GlobalVariable **gv_array = new GlobalVariable *[n]; // first, back-up GlobalVariable in gv_array i = 0; for (Module::global_iterator I = global_list.begin(), E = global_list.end(); - I != E; ++I) + I != E; ++I) gv_array[i++] = &*I; // second, empty global_list @@ -1024,13 +952,14 @@ bool NVPTXAsmPrinter::doFinalization(Module &M) { bool ret = AsmPrinter::doFinalization(M); // now we restore global variables - for (i = 0; i < n; i ++) + for (i = 0; i < n; i++) global_list.insert(global_list.end(), gv_array[i]); + clearAnnotationCache(&M); + delete[] gv_array; return ret; - //bool Result = AsmPrinter::doFinalization(M); // Instead of calling the parents doFinalization, we may // clone parents doFinalization and customize here. @@ -1049,10 +978,14 @@ bool NVPTXAsmPrinter::doFinalization(Module &M) { // external global variable with init -> .visible // external without init -> .extern // appending -> not allowed, assert. - -void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue* V, raw_ostream &O) -{ - if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) { +// for any linkage other than +// internal, private, linker_private, +// linker_private_weak, linker_private_weak_def_auto, +// we emit -> .weak. + +void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V, + raw_ostream &O) { + if (static_cast(TM).getDrvInterface() == NVPTX::CUDA) { if (V->hasExternalLinkage()) { if (isa(V)) { const GlobalVariable *GVar = cast(V); @@ -1071,23 +1004,31 @@ void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue* V, raw_ostream &O) msg.append("Error: "); msg.append("Symbol "); if (V->hasName()) - msg.append(V->getName().str()); + msg.append(V->getName()); msg.append("has unsupported appending linkage type"); llvm_unreachable(msg.c_str()); + } else if (!V->hasInternalLinkage() && + !V->hasPrivateLinkage()) { + O << ".weak "; } } } - -void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O, +void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, + raw_ostream &O, bool processDemoted) { // Skip meta data if (GVar->hasSection()) { - if (GVar->getSection() == "llvm.metadata") + if (GVar->getSection() == StringRef("llvm.metadata")) return; } + // Skip LLVM intrinsic global variables + if (GVar->getName().startswith("llvm.") || + GVar->getName().startswith("nvvm.")) + return; + const DataLayout *TD = TM.getDataLayout(); // GlobalVariables are always constant pointers themselves. @@ -1099,6 +1040,10 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O, O << ".visible "; else O << ".extern "; + } else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() || + GVar->hasAvailableExternallyLinkage() || + GVar->hasCommonLinkage()) { + O << ".weak "; } if (llvm::isTexture(*GVar)) { @@ -1123,37 +1068,55 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O, if (llvm::isSampler(*GVar)) { O << ".global .samplerref " << llvm::getSamplerName(*GVar); - Constant *Initializer = NULL; + const Constant *Initializer = nullptr; if (GVar->hasInitializer()) Initializer = GVar->getInitializer(); - ConstantInt *CI = NULL; + const ConstantInt *CI = nullptr; if (Initializer) CI = dyn_cast(Initializer); if (CI) { - unsigned sample=CI->getZExtValue(); + unsigned sample = CI->getZExtValue(); O << " = { "; - for (int i =0, addr=((sample & __CLK_ADDRESS_MASK ) >> - __CLK_ADDRESS_BASE) ; i < 3 ; i++) { + for (int i = 0, + addr = ((sample & __CLK_ADDRESS_MASK) >> __CLK_ADDRESS_BASE); + i < 3; i++) { O << "addr_mode_" << i << " = "; switch (addr) { - case 0: O << "wrap"; break; - case 1: O << "clamp_to_border"; break; - case 2: O << "clamp_to_edge"; break; - case 3: O << "wrap"; break; - case 4: O << "mirror"; break; + case 0: + O << "wrap"; + break; + case 1: + O << "clamp_to_border"; + break; + case 2: + O << "clamp_to_edge"; + break; + case 3: + O << "wrap"; + break; + case 4: + O << "mirror"; + break; } - O <<", "; + O << ", "; } O << "filter_mode = "; - switch (( sample & __CLK_FILTER_MASK ) >> __CLK_FILTER_BASE ) { - case 0: O << "nearest"; break; - case 1: O << "linear"; break; - case 2: assert ( 0 && "Anisotropic filtering is not supported"); - default: O << "nearest"; break; + switch ((sample & __CLK_FILTER_MASK) >> __CLK_FILTER_BASE) { + case 0: + O << "nearest"; + break; + case 1: + O << "linear"; + break; + case 2: + llvm_unreachable("Anisotropic filtering is not supported"); + default: + O << "nearest"; + break; } - if (!(( sample &__CLK_NORMALIZED_MASK ) >> __CLK_NORMALIZED_BASE)) { + if (!((sample & __CLK_NORMALIZED_MASK) >> __CLK_NORMALIZED_BASE)) { O << ", force_unnormalized_coords = 1"; } O << " }"; @@ -1175,13 +1138,13 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O, return; } - const Function *demotedFunc = 0; + const Function *demotedFunc = nullptr; if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) { - O << "// " << GVar->getName().str() << " has been demoted\n"; + O << "// " << GVar->getName() << " has been demoted\n"; if (localDecls.find(demotedFunc) != localDecls.end()) localDecls[demotedFunc].push_back(GVar); else { - std::vector temp; + std::vector temp; temp.push_back(GVar); localDecls[demotedFunc] = temp; } @@ -1190,32 +1153,51 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O, O << "."; emitPTXAddressSpace(PTy->getAddressSpace(), O); + + if (isManaged(*GVar)) { + O << " .attribute(.managed)"; + } + if (GVar->getAlignment() == 0) O << " .align " << (int) TD->getPrefTypeAlignment(ETy); else O << " .align " << GVar->getAlignment(); - - if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa(ETy)) { + if (ETy->isFloatingPointTy() || ETy->isIntegerTy() || ETy->isPointerTy()) { O << " ."; - O << getPTXFundamentalTypeStr(ETy, false); + // Special case: ABI requires that we use .u8 for predicates + if (ETy->isIntegerTy(1)) + O << "u8"; + else + O << getPTXFundamentalTypeStr(ETy, false); O << " "; - O << *Mang->getSymbol(GVar); + O << *getSymbol(GVar); // Ptx allows variable initilization only for constant and global state // spaces. - if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) || - (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) || - (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) - && GVar->hasInitializer()) { - Constant *Initializer = GVar->getInitializer(); - if (!Initializer->isNullValue()) { - O << " = " ; - printScalarConstant(Initializer, O); + if (GVar->hasInitializer()) { + if ((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) || + (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) { + const Constant *Initializer = GVar->getInitializer(); + // 'undef' is treated as there is no value specified. + if (!Initializer->isNullValue() && !isa(Initializer)) { + O << " = "; + printScalarConstant(Initializer, O); + } + } else { + // The frontend adds zero-initializer to variables that don't have an + // initial value, so skip warning for this case. + if (!GVar->getInitializer()->isNullValue()) { + std::string warnMsg = + ("initial value of '" + GVar->getName() + + "' is not allowed in addrspace(" + + Twine(llvm::utostr_32(PTy->getAddressSpace())) + ")").str(); + report_fatal_error(warnMsg.c_str()); + } } } } else { - unsigned int ElementSize =0; + unsigned int ElementSize = 0; // Although PTX has direct support for struct type and array type and // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for @@ -1229,54 +1211,48 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O, // Ptx allows variable initilization only for constant and // global state spaces. if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) || - (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) || - (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) - && GVar->hasInitializer()) { - Constant *Initializer = GVar->getInitializer(); - if (!isa(Initializer) && - !Initializer->isNullValue()) { + (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) && + GVar->hasInitializer()) { + const Constant *Initializer = GVar->getInitializer(); + if (!isa(Initializer) && !Initializer->isNullValue()) { AggBuffer aggBuffer(ElementSize, O, *this); bufferAggregateConstant(Initializer, &aggBuffer); if (aggBuffer.numSymbols) { - if (nvptxSubtarget.is64Bit()) { - O << " .u64 " << *Mang->getSymbol(GVar) <<"[" ; - O << ElementSize/8; - } - else { - O << " .u32 " << *Mang->getSymbol(GVar) <<"[" ; - O << ElementSize/4; + if (static_cast(TM).is64Bit()) { + O << " .u64 " << *getSymbol(GVar) << "["; + O << ElementSize / 8; + } else { + O << " .u32 " << *getSymbol(GVar) << "["; + O << ElementSize / 4; } O << "]"; - } - else { - O << " .b8 " << *Mang->getSymbol(GVar) <<"[" ; + } else { + O << " .b8 " << *getSymbol(GVar) << "["; O << ElementSize; O << "]"; } - O << " = {" ; + O << " = {"; aggBuffer.print(); O << "}"; - } - else { - O << " .b8 " << *Mang->getSymbol(GVar) ; + } else { + O << " .b8 " << *getSymbol(GVar); if (ElementSize) { - O <<"[" ; + O << "["; O << ElementSize; O << "]"; } } - } - else { - O << " .b8 " << *Mang->getSymbol(GVar); + } else { + O << " .b8 " << *getSymbol(GVar); if (ElementSize) { - O <<"[" ; + O << "["; O << ElementSize; O << "]"; } } break; default: - assert( 0 && "type not supported yet"); + llvm_unreachable("type not supported yet"); } } @@ -1287,9 +1263,9 @@ void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) { if (localDecls.find(f) == localDecls.end()) return; - std::vector &gvars = localDecls[f]; + std::vector &gvars = localDecls[f]; - for (unsigned i=0, e=gvars.size(); i!=e; ++i) { + for (unsigned i = 0, e = gvars.size(); i != e; ++i) { O << "\t// demoted variable\n\t"; printModuleLevelGV(gvars[i], O, true); } @@ -1299,32 +1275,25 @@ void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace, raw_ostream &O) const { switch (AddressSpace) { case llvm::ADDRESS_SPACE_LOCAL: - O << "local" ; + O << "local"; break; case llvm::ADDRESS_SPACE_GLOBAL: - O << "global" ; + O << "global"; break; case llvm::ADDRESS_SPACE_CONST: - // This logic should be consistent with that in - // getCodeAddrSpace() (NVPTXISelDATToDAT.cpp) - if (nvptxSubtarget.hasGenericLdSt()) - O << "global" ; - else - O << "const" ; - break; - case llvm::ADDRESS_SPACE_CONST_NOT_GEN: - O << "const" ; + O << "const"; break; case llvm::ADDRESS_SPACE_SHARED: - O << "shared" ; + O << "shared"; break; default: - llvm_unreachable("unexpected address space"); + report_fatal_error("Bad address space found while emitting PTX"); + break; } } -std::string NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty, - bool useB4PTR) const { +std::string +NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty, bool useB4PTR) const { switch (Ty->getTypeID()) { default: llvm_unreachable("unexpected type"); @@ -1347,18 +1316,21 @@ std::string NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty, case Type::DoubleTyID: return "f64"; case Type::PointerTyID: - if (nvptxSubtarget.is64Bit()) - if (useB4PTR) return "b64"; - else return "u64"; + if (static_cast(TM).is64Bit()) + if (useB4PTR) + return "b64"; + else + return "u64"; + else if (useB4PTR) + return "b32"; else - if (useB4PTR) return "b32"; - else return "u32"; + return "u32"; } llvm_unreachable("unexpected type"); - return NULL; + return nullptr; } -void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable* GVar, +void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar, raw_ostream &O) { const DataLayout *TD = TM.getDataLayout(); @@ -1374,15 +1346,15 @@ void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable* GVar, else O << " .align " << GVar->getAlignment(); - if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa(ETy)) { + if (ETy->isFloatingPointTy() || ETy->isIntegerTy() || ETy->isPointerTy()) { O << " ."; O << getPTXFundamentalTypeStr(ETy); O << " "; - O << *Mang->getSymbol(GVar); + O << *getSymbol(GVar); return; } - int64_t ElementSize =0; + int64_t ElementSize = 0; // Although PTX has direct support for struct type and array type and LLVM IR // is very similar to PTX, the LLVM CodeGen does not support for targets that @@ -1393,46 +1365,32 @@ void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable* GVar, case Type::ArrayTyID: case Type::VectorTyID: ElementSize = TD->getTypeStoreSize(ETy); - O << " .b8 " << *Mang->getSymbol(GVar) <<"[" ; + O << " .b8 " << *getSymbol(GVar) << "["; if (ElementSize) { - O << itostr(ElementSize) ; + O << itostr(ElementSize); } O << "]"; break; default: - assert( 0 && "type not supported yet"); + llvm_unreachable("type not supported yet"); } - return ; + return; } - -static unsigned int -getOpenCLAlignment(const DataLayout *TD, - Type *Ty) { - if (Ty->isPrimitiveType() || Ty->isIntegerTy() || isa(Ty)) +static unsigned int getOpenCLAlignment(const DataLayout *TD, Type *Ty) { + if (Ty->isSingleValueType()) return TD->getPrefTypeAlignment(Ty); const ArrayType *ATy = dyn_cast(Ty); if (ATy) return getOpenCLAlignment(TD, ATy->getElementType()); - const VectorType *VTy = dyn_cast(Ty); - if (VTy) { - Type *ETy = VTy->getElementType(); - unsigned int numE = VTy->getNumElements(); - unsigned int alignE = TD->getPrefTypeAlignment(ETy); - if (numE == 3) - return 4*alignE; - else - return numE*alignE; - } - const StructType *STy = dyn_cast(Ty); if (STy) { unsigned int alignStruct = 1; // Go through each element of the struct and find the // largest alignment. - for (unsigned i=0, e=STy->getNumElements(); i != e; i++) { + for (unsigned i = 0, e = STy->getNumElements(); i != e; i++) { Type *ETy = STy->getElementType(i); unsigned int align = getOpenCLAlignment(TD, ETy); if (align > alignStruct) @@ -1449,57 +1407,28 @@ getOpenCLAlignment(const DataLayout *TD, void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I, int paramIndex, raw_ostream &O) { - if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) || - (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) - O << *CurrentFnSym << "_param_" << paramIndex; - else { - std::string argName = I->getName(); - const char *p = argName.c_str(); - while (*p) { - if (*p == '.') - O << "_"; - else - O << *p; - p++; - } - } + O << *getSymbol(I->getParent()) << "_param_" << paramIndex; } void NVPTXAsmPrinter::printParamName(int paramIndex, raw_ostream &O) { - Function::const_arg_iterator I, E; - int i = 0; - - if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) || - (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) { - O << *CurrentFnSym << "_param_" << paramIndex; - return; - } - - for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, i++) { - if (i==paramIndex) { - printParamName(I, paramIndex, O); - return; - } - } - llvm_unreachable("paramIndex out of bound"); + O << *CurrentFnSym << "_param_" << paramIndex; } -void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, - raw_ostream &O) { +void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { const DataLayout *TD = TM.getDataLayout(); - const AttrListPtr &PAL = F->getAttributes(); - const TargetLowering *TLI = TM.getTargetLowering(); + const AttributeSet &PAL = F->getAttributes(); + const TargetLowering *TLI = nvptxSubtarget->getTargetLowering(); Function::const_arg_iterator I, E; unsigned paramIndex = 0; bool first = true; bool isKernelFunc = llvm::isKernelFunction(*F); - bool isABI = (nvptxSubtarget.getSmVersion() >= 20); + bool isABI = (nvptxSubtarget->getSmVersion() >= 20); MVT thePointerTy = TLI->getPointerTy(); O << "(\n"; for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) { - const Type *Ty = I->getType(); + Type *Ty = I->getType(); if (!first) O << ",\n"; @@ -1507,22 +1436,51 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, first = false; // Handle image/sampler parameters - if (llvm::isSampler(*I) || llvm::isImage(*I)) { - if (llvm::isImage(*I)) { - std::string sname = I->getName(); - if (llvm::isImageWriteOnly(*I)) - O << "\t.param .surfref " << *CurrentFnSym << "_param_" << paramIndex; - else // Default image is read_only - O << "\t.param .texref " << *CurrentFnSym << "_param_" << paramIndex; + if (isKernelFunction(*F)) { + if (isSampler(*I) || isImage(*I)) { + if (isImage(*I)) { + std::string sname = I->getName(); + if (isImageWriteOnly(*I) || isImageReadWrite(*I)) { + if (nvptxSubtarget->hasImageHandles()) + O << "\t.param .u64 .ptr .surfref "; + else + O << "\t.param .surfref "; + O << *CurrentFnSym << "_param_" << paramIndex; + } + else { // Default image is read_only + if (nvptxSubtarget->hasImageHandles()) + O << "\t.param .u64 .ptr .texref "; + else + O << "\t.param .texref "; + O << *CurrentFnSym << "_param_" << paramIndex; + } + } else { + if (nvptxSubtarget->hasImageHandles()) + O << "\t.param .u64 .ptr .samplerref "; + else + O << "\t.param .samplerref "; + O << *CurrentFnSym << "_param_" << paramIndex; + } + continue; } - else // Should be llvm::isSampler(*I) - O << "\t.param .samplerref " << *CurrentFnSym << "_param_" - << paramIndex; - continue; } - if (PAL.getParamAttributes(paramIndex+1). - hasAttribute(Attributes::ByVal) == false) { + if (!PAL.hasAttribute(paramIndex + 1, Attribute::ByVal)) { + if (Ty->isAggregateType() || Ty->isVectorTy()) { + // Just print .param .align .b8 .param[size]; + // = PAL.getparamalignment + // size = typeallocsize of element type + unsigned align = PAL.getParamAlignment(paramIndex + 1); + if (align == 0) + align = TD->getABITypeAlignment(Ty); + + unsigned sz = TD->getTypeAllocSize(Ty); + O << "\t.param .align " << align << " .b8 "; + printParamName(I, paramIndex, O); + O << "[" << sz << "]"; + + continue; + } // Just a scalar const PointerType *PTy = dyn_cast(Ty); if (isKernelFunc) { @@ -1530,44 +1488,49 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, // Special handling for pointer arguments to kernel O << "\t.param .u" << thePointerTy.getSizeInBits() << " "; - if (nvptxSubtarget.getDrvInterface() != NVPTX::CUDA) { + if (static_cast(TM).getDrvInterface() != + NVPTX::CUDA) { Type *ETy = PTy->getElementType(); int addrSpace = PTy->getAddressSpace(); - switch(addrSpace) { + switch (addrSpace) { default: O << ".ptr "; break; - case llvm::ADDRESS_SPACE_CONST_NOT_GEN: + case llvm::ADDRESS_SPACE_CONST: O << ".ptr .const "; break; case llvm::ADDRESS_SPACE_SHARED: O << ".ptr .shared "; break; case llvm::ADDRESS_SPACE_GLOBAL: - case llvm::ADDRESS_SPACE_CONST: O << ".ptr .global "; break; } - O << ".align " << (int)getOpenCLAlignment(TD, ETy) << " "; + O << ".align " << (int) getOpenCLAlignment(TD, ETy) << " "; } printParamName(I, paramIndex, O); continue; } // non-pointer scalar to kernel func - O << "\t.param ." - << getPTXFundamentalTypeStr(Ty) << " "; + O << "\t.param ."; + // Special case: predicate operands become .u8 types + if (Ty->isIntegerTy(1)) + O << "u8"; + else + O << getPTXFundamentalTypeStr(Ty); + O << " "; printParamName(I, paramIndex, O); continue; } // Non-kernel function, just print .param .b for ABI - // and .reg .b for non ABY + // and .reg .b for non-ABI unsigned sz = 0; if (isa(Ty)) { sz = cast(Ty)->getBitWidth(); - if (sz < 32) sz = 32; - } - else if (isa(Ty)) + if (sz < 32) + sz = 32; + } else if (isa(Ty)) sz = thePointerTy.getSizeInBits(); else sz = Ty->getPrimitiveSizeInBits(); @@ -1581,21 +1544,19 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, // param has byVal attribute. So should be a pointer const PointerType *PTy = dyn_cast(Ty); - assert(PTy && - "Param with byval attribute should be a pointer type"); + assert(PTy && "Param with byval attribute should be a pointer type"); Type *ETy = PTy->getElementType(); if (isABI || isKernelFunc) { - // Just print .param .b8 .align .param[size]; + // Just print .param .align .b8 .param[size]; // = PAL.getparamalignment // size = typeallocsize of element type - unsigned align = PAL.getParamAlignment(paramIndex+1); + unsigned align = PAL.getParamAlignment(paramIndex + 1); if (align == 0) align = TD->getABITypeAlignment(ETy); unsigned sz = TD->getTypeAllocSize(ETy); - O << "\t.param .align " << align - << " .b8 "; + O << "\t.param .align " << align << " .b8 "; printParamName(I, paramIndex, O); O << "[" << sz << "]"; continue; @@ -1606,7 +1567,7 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, // each vector element. SmallVector vtparts; ComputeValueVTs(*TLI, ETy, vtparts); - for (unsigned i=0,e=vtparts.size(); i!=e; ++i) { + for (unsigned i = 0, e = vtparts.size(); i != e; ++i) { unsigned elems = 1; EVT elemtype = vtparts[i]; if (vtparts[i].isVector()) { @@ -1614,15 +1575,17 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, elemtype = vtparts[i].getVectorElementType(); } - for (unsigned j=0,je=elems; j!=je; ++j) { + for (unsigned j = 0, je = elems; j != je; ++j) { unsigned sz = elemtype.getSizeInBits(); - if (elemtype.isInteger() && (sz < 32)) sz = 32; + if (elemtype.isInteger() && (sz < 32)) + sz = 32; O << "\t.reg .b" << sz << " "; printParamName(I, paramIndex, O); - if (j Str; raw_svector_ostream O(Str); // Map the global virtual register number to a register class specific // virtual register number starting from 1 with that class. - const TargetRegisterInfo *TRI = MF.getTarget().getRegisterInfo(); + const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo(); //unsigned numRegClasses = TRI->getNumRegClasses(); // Emit the Fake Stack Object const MachineFrameInfo *MFI = MF.getFrameInfo(); int NumBytes = (int) MFI->getStackSize(); if (NumBytes) { - O << "\t.local .align " << MFI->getMaxAlignment() << " .b8 \t" - << DEPOTNAME - << getFunctionNumber() << "[" << NumBytes << "];\n"; - if (nvptxSubtarget.is64Bit()) { + O << "\t.local .align " << MFI->getMaxAlignment() << " .b8 \t" << DEPOTNAME + << getFunctionNumber() << "[" << NumBytes << "];\n"; + if (static_cast(MF.getTarget()).is64Bit()) { O << "\t.reg .b64 \t%SP;\n"; O << "\t.reg .b64 \t%SPL;\n"; - } - else { + } else { O << "\t.reg .b32 \t%SP;\n"; O << "\t.reg .b32 \t%SPL;\n"; } @@ -1672,72 +1632,57 @@ setAndEmitFunctionVirtualRegisters(const MachineFunction &MF) { // register number and the per class virtual register number. // We use the per class virtual register number in the ptx output. unsigned int numVRs = MRI->getNumVirtRegs(); - for (unsigned i=0; i< numVRs; i++) { + for (unsigned i = 0; i < numVRs; i++) { unsigned int vr = TRI->index2VirtReg(i); const TargetRegisterClass *RC = MRI->getRegClass(vr); - std::map ®map = VRidGlobal2LocalMap[RC->getID()]; + DenseMap ®map = VRegMapping[RC]; int n = regmap.size(); - regmap.insert(std::make_pair(vr, n+1)); + regmap.insert(std::make_pair(vr, n + 1)); } // Emit register declarations // @TODO: Extract out the real register usage - O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s64 %rl<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .f64 %fl<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s64 %rd<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .f64 %fd<" << NVPTXNumRegisters << ">;\n"; // Emit declaration of the virtual registers or 'physical' registers for // each register class - //for (unsigned i=0; i< numRegClasses; i++) { - // std::map ®map = VRidGlobal2LocalMap[i]; - // const TargetRegisterClass *RC = TRI->getRegClass(i); - // std::string rcname = getNVPTXRegClassName(RC); - // std::string rcStr = getNVPTXRegClassStr(RC); - // //int n = regmap.size(); - // if (!isNVPTXVectorRegClass(RC)) { - // O << "\t.reg " << rcname << " \t" << rcStr << "<" - // << NVPTXNumRegisters << ">;\n"; - // } - - // Only declare those registers that may be used. And do not emit vector - // registers as - // they are all elementized to scalar registers. - //if (n && !isNVPTXVectorRegClass(RC)) { - // if (RegAllocNilUsed) { - // O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1) - // << ">;\n"; - // } - // else { - // O << "\t.reg " << rcname << " \t" << StrToUpper(rcStr) - // << "<" << 32 << ">;\n"; - // } - //} - //} - - OutStreamer.EmitRawText(O.str()); -} + for (unsigned i=0; i< TRI->getNumRegClasses(); i++) { + const TargetRegisterClass *RC = TRI->getRegClass(i); + DenseMap ®map = VRegMapping[RC]; + std::string rcname = getNVPTXRegClassName(RC); + std::string rcStr = getNVPTXRegClassStr(RC); + int n = regmap.size(); + + // Only declare those registers that may be used. + if (n) { + O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1) + << ">;\n"; + } + } + OutStreamer->EmitRawText(O.str()); +} void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) { - APFloat APF = APFloat(Fp->getValueAPF()); // make a copy + APFloat APF = APFloat(Fp->getValueAPF()); // make a copy bool ignored; unsigned int numHex; const char *lead; - if (Fp->getType()->getTypeID()==Type::FloatTyID) { + if (Fp->getType()->getTypeID() == Type::FloatTyID) { numHex = 8; lead = "0f"; - APF.convert(APFloat::IEEEsingle, APFloat::rmNearestTiesToEven, - &ignored); + APF.convert(APFloat::IEEEsingle, APFloat::rmNearestTiesToEven, &ignored); } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) { numHex = 16; lead = "0d"; - APF.convert(APFloat::IEEEdouble, APFloat::rmNearestTiesToEven, - &ignored); + APF.convert(APFloat::IEEEdouble, APFloat::rmNearestTiesToEven, &ignored); } else llvm_unreachable("unsupported fp type"); @@ -1749,12 +1694,12 @@ void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) { O << utohexstr(API.getZExtValue()); } -void NVPTXAsmPrinter::printScalarConstant(Constant *CPV, raw_ostream &O) { - if (ConstantInt *CI = dyn_cast(CPV)) { +void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) { + if (const ConstantInt *CI = dyn_cast(CPV)) { O << CI->getValue(); return; } - if (ConstantFP *CFP = dyn_cast(CPV)) { + if (const ConstantFP *CFP = dyn_cast(CPV)) { printFPConstant(CFP, O); return; } @@ -1762,32 +1707,53 @@ void NVPTXAsmPrinter::printScalarConstant(Constant *CPV, raw_ostream &O) { O << "0"; return; } - if (GlobalValue *GVar = dyn_cast(CPV)) { - O << *Mang->getSymbol(GVar); + if (const GlobalValue *GVar = dyn_cast(CPV)) { + PointerType *PTy = dyn_cast(GVar->getType()); + bool IsNonGenericPointer = false; + if (PTy && PTy->getAddressSpace() != 0) { + IsNonGenericPointer = true; + } + if (EmitGeneric && !isa(CPV) && !IsNonGenericPointer) { + O << "generic("; + O << *getSymbol(GVar); + O << ")"; + } else { + O << *getSymbol(GVar); + } return; } - if (ConstantExpr *Cexpr = dyn_cast(CPV)) { - Value *v = Cexpr->stripPointerCasts(); - if (GlobalValue *GVar = dyn_cast(v)) { - O << *Mang->getSymbol(GVar); + if (const ConstantExpr *Cexpr = dyn_cast(CPV)) { + const Value *v = Cexpr->stripPointerCasts(); + PointerType *PTy = dyn_cast(Cexpr->getType()); + bool IsNonGenericPointer = false; + if (PTy && PTy->getAddressSpace() != 0) { + IsNonGenericPointer = true; + } + if (const GlobalValue *GVar = dyn_cast(v)) { + if (EmitGeneric && !isa(v) && !IsNonGenericPointer) { + O << "generic("; + O << *getSymbol(GVar); + O << ")"; + } else { + O << *getSymbol(GVar); + } return; } else { - O << *LowerConstant(CPV, *this); + O << *lowerConstant(CPV); return; } } llvm_unreachable("Not scalar type found in printScalarConstant()"); } - -void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, +void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, AggBuffer *aggBuffer) { const DataLayout *TD = TM.getDataLayout(); if (isa(CPV) || CPV->isNullValue()) { int s = TD->getTypeAllocSize(CPV->getType()); - if (saddZeros(s); return; @@ -1798,56 +1764,53 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, case Type::IntegerTyID: { const Type *ETy = CPV->getType(); - if ( ETy == Type::getInt8Ty(CPV->getContext()) ){ - unsigned char c = - (unsigned char)(dyn_cast(CPV))->getZExtValue(); + if (ETy == Type::getInt8Ty(CPV->getContext())) { + unsigned char c = (unsigned char)cast(CPV)->getZExtValue(); ptr = &c; aggBuffer->addBytes(ptr, 1, Bytes); - } else if ( ETy == Type::getInt16Ty(CPV->getContext()) ) { - short int16 = - (short)(dyn_cast(CPV))->getZExtValue(); - ptr = (unsigned char*)&int16; + } else if (ETy == Type::getInt16Ty(CPV->getContext())) { + short int16 = (short)cast(CPV)->getZExtValue(); + ptr = (unsigned char *)&int16; aggBuffer->addBytes(ptr, 2, Bytes); - } else if ( ETy == Type::getInt32Ty(CPV->getContext()) ) { - if (ConstantInt *constInt = dyn_cast(CPV)) { - int int32 =(int)(constInt->getZExtValue()); - ptr = (unsigned char*)&int32; + } else if (ETy == Type::getInt32Ty(CPV->getContext())) { + if (const ConstantInt *constInt = dyn_cast(CPV)) { + int int32 = (int)(constInt->getZExtValue()); + ptr = (unsigned char *)&int32; aggBuffer->addBytes(ptr, 4, Bytes); break; - } else if (ConstantExpr *Cexpr = dyn_cast(CPV)) { - if (ConstantInt *constInt = - dyn_cast(ConstantFoldConstantExpression( - Cexpr, TD))) { - int int32 =(int)(constInt->getZExtValue()); - ptr = (unsigned char*)&int32; + } else if (const ConstantExpr *Cexpr = dyn_cast(CPV)) { + if (const ConstantInt *constInt = dyn_cast( + ConstantFoldConstantExpression(Cexpr, *TD))) { + int int32 = (int)(constInt->getZExtValue()); + ptr = (unsigned char *)&int32; aggBuffer->addBytes(ptr, 4, Bytes); break; } if (Cexpr->getOpcode() == Instruction::PtrToInt) { Value *v = Cexpr->getOperand(0)->stripPointerCasts(); - aggBuffer->addSymbol(v); + aggBuffer->addSymbol(v, Cexpr->getOperand(0)); aggBuffer->addZeros(4); break; } } llvm_unreachable("unsupported integer const type"); - } else if (ETy == Type::getInt64Ty(CPV->getContext()) ) { - if (ConstantInt *constInt = dyn_cast(CPV)) { - long long int64 =(long long)(constInt->getZExtValue()); - ptr = (unsigned char*)&int64; + } else if (ETy == Type::getInt64Ty(CPV->getContext())) { + if (const ConstantInt *constInt = dyn_cast(CPV)) { + long long int64 = (long long)(constInt->getZExtValue()); + ptr = (unsigned char *)&int64; aggBuffer->addBytes(ptr, 8, Bytes); break; - } else if (ConstantExpr *Cexpr = dyn_cast(CPV)) { - if (ConstantInt *constInt = dyn_cast( - ConstantFoldConstantExpression(Cexpr, TD))) { - long long int64 =(long long)(constInt->getZExtValue()); - ptr = (unsigned char*)&int64; + } else if (const ConstantExpr *Cexpr = dyn_cast(CPV)) { + if (const ConstantInt *constInt = dyn_cast( + ConstantFoldConstantExpression(Cexpr, *TD))) { + long long int64 = (long long)(constInt->getZExtValue()); + ptr = (unsigned char *)&int64; aggBuffer->addBytes(ptr, 8, Bytes); break; } if (Cexpr->getOpcode() == Instruction::PtrToInt) { Value *v = Cexpr->getOperand(0)->stripPointerCasts(); - aggBuffer->addSymbol(v); + aggBuffer->addSymbol(v, Cexpr->getOperand(0)); aggBuffer->addZeros(8); break; } @@ -1859,29 +1822,27 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, } case Type::FloatTyID: case Type::DoubleTyID: { - ConstantFP *CFP = dyn_cast(CPV); - const Type* Ty = CFP->getType(); + const ConstantFP *CFP = dyn_cast(CPV); + const Type *Ty = CFP->getType(); if (Ty == Type::getFloatTy(CPV->getContext())) { - float float32 = (float)CFP->getValueAPF().convertToFloat(); - ptr = (unsigned char*)&float32; + float float32 = (float) CFP->getValueAPF().convertToFloat(); + ptr = (unsigned char *)&float32; aggBuffer->addBytes(ptr, 4, Bytes); } else if (Ty == Type::getDoubleTy(CPV->getContext())) { double float64 = CFP->getValueAPF().convertToDouble(); - ptr = (unsigned char*)&float64; + ptr = (unsigned char *)&float64; aggBuffer->addBytes(ptr, 8, Bytes); - } - else { + } else { llvm_unreachable("unsupported fp const type"); } break; } case Type::PointerTyID: { - if (GlobalValue *GVar = dyn_cast(CPV)) { - aggBuffer->addSymbol(GVar); - } - else if (ConstantExpr *Cexpr = dyn_cast(CPV)) { - Value *v = Cexpr->stripPointerCasts(); - aggBuffer->addSymbol(v); + if (const GlobalValue *GVar = dyn_cast(CPV)) { + aggBuffer->addSymbol(GVar, GVar); + } else if (const ConstantExpr *Cexpr = dyn_cast(CPV)) { + const Value *v = Cexpr->stripPointerCasts(); + aggBuffer->addSymbol(v, Cexpr); } unsigned int s = TD->getTypeAllocSize(CPV->getType()); aggBuffer->addZeros(s); @@ -1892,13 +1853,12 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, case Type::VectorTyID: case Type::StructTyID: { if (isa(CPV) || isa(CPV) || - isa(CPV)) { + isa(CPV) || isa(CPV)) { int ElementSize = TD->getTypeAllocSize(CPV->getType()); bufferAggregateConstant(CPV, aggBuffer); - if ( Bytes > ElementSize ) - aggBuffer->addZeros(Bytes-ElementSize); - } - else if (isa(CPV)) + if (Bytes > ElementSize) + aggBuffer->addZeros(Bytes - ElementSize); + } else if (isa(CPV)) aggBuffer->addZeros(Bytes); else llvm_unreachable("Unexpected Constant type"); @@ -1910,7 +1870,7 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, } } -void NVPTXAsmPrinter::bufferAggregateConstant(Constant *CPV, +void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV, AggBuffer *aggBuffer) { const DataLayout *TD = TM.getDataLayout(); int Bytes; @@ -1924,7 +1884,7 @@ void NVPTXAsmPrinter::bufferAggregateConstant(Constant *CPV, } if (const ConstantDataSequential *CDS = - dyn_cast(CPV)) { + dyn_cast(CPV)) { if (CDS->getNumElements()) for (unsigned i = 0; i < CDS->getNumElements(); ++i) bufferLEByte(cast(CDS->getElementAsConstant(i)), 0, @@ -1932,20 +1892,18 @@ void NVPTXAsmPrinter::bufferAggregateConstant(Constant *CPV, return; } - if (isa(CPV)) { if (CPV->getNumOperands()) { StructType *ST = cast(CPV->getType()); for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) { - if ( i == (e - 1)) + if (i == (e - 1)) Bytes = TD->getStructLayout(ST)->getElementOffset(0) + - TD->getTypeAllocSize(ST) - - TD->getStructLayout(ST)->getElementOffset(i); + TD->getTypeAllocSize(ST) - + TD->getStructLayout(ST)->getElementOffset(i); else - Bytes = TD->getStructLayout(ST)->getElementOffset(i+1) - - TD->getStructLayout(ST)->getElementOffset(i); - bufferLEByte(cast(CPV->getOperand(i)), Bytes, - aggBuffer); + Bytes = TD->getStructLayout(ST)->getElementOffset(i + 1) - + TD->getStructLayout(ST)->getElementOffset(i); + bufferLEByte(cast(CPV->getOperand(i)), Bytes, aggBuffer); } } return; @@ -1956,28 +1914,83 @@ void NVPTXAsmPrinter::bufferAggregateConstant(Constant *CPV, // buildTypeNameMap - Run through symbol table looking for type names. // - bool NVPTXAsmPrinter::isImageType(const Type *Ty) { std::map::iterator PI = TypeNameMap.find(Ty); - if (PI != TypeNameMap.end() && - (!PI->second.compare("struct._image1d_t") || - !PI->second.compare("struct._image2d_t") || - !PI->second.compare("struct._image3d_t"))) + if (PI != TypeNameMap.end() && (!PI->second.compare("struct._image1d_t") || + !PI->second.compare("struct._image2d_t") || + !PI->second.compare("struct._image3d_t"))) return true; return false; } + +bool NVPTXAsmPrinter::ignoreLoc(const MachineInstr &MI) { + switch (MI.getOpcode()) { + default: + return false; + case NVPTX::CallArgBeginInst: + case NVPTX::CallArgEndInst0: + case NVPTX::CallArgEndInst1: + case NVPTX::CallArgF32: + case NVPTX::CallArgF64: + case NVPTX::CallArgI16: + case NVPTX::CallArgI32: + case NVPTX::CallArgI32imm: + case NVPTX::CallArgI64: + case NVPTX::CallArgParam: + case NVPTX::CallVoidInst: + case NVPTX::CallVoidInstReg: + case NVPTX::Callseq_End: + case NVPTX::CallVoidInstReg64: + case NVPTX::DeclareParamInst: + case NVPTX::DeclareRetMemInst: + case NVPTX::DeclareRetRegInst: + case NVPTX::DeclareRetScalarInst: + case NVPTX::DeclareScalarParamInst: + case NVPTX::DeclareScalarRegInst: + case NVPTX::StoreParamF32: + case NVPTX::StoreParamF64: + case NVPTX::StoreParamI16: + case NVPTX::StoreParamI32: + case NVPTX::StoreParamI64: + case NVPTX::StoreParamI8: + case NVPTX::StoreRetvalF32: + case NVPTX::StoreRetvalF64: + case NVPTX::StoreRetvalI16: + case NVPTX::StoreRetvalI32: + case NVPTX::StoreRetvalI64: + case NVPTX::StoreRetvalI8: + case NVPTX::LastCallArgF32: + case NVPTX::LastCallArgF64: + case NVPTX::LastCallArgI16: + case NVPTX::LastCallArgI32: + case NVPTX::LastCallArgI32imm: + case NVPTX::LastCallArgI64: + case NVPTX::LastCallArgParam: + case NVPTX::LoadParamMemF32: + case NVPTX::LoadParamMemF64: + case NVPTX::LoadParamMemI16: + case NVPTX::LoadParamMemI32: + case NVPTX::LoadParamMemI64: + case NVPTX::LoadParamMemI8: + case NVPTX::PrototypeInst: + case NVPTX::DBG_VALUE: + return true; + } + return false; +} + /// PrintAsmOperand - Print out an operand for an inline asm expression. /// bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, unsigned AsmVariant, - const char *ExtraCode, - raw_ostream &O) { + const char *ExtraCode, raw_ostream &O) { if (ExtraCode && ExtraCode[0]) { - if (ExtraCode[1] != 0) return true; // Unknown modifier. + if (ExtraCode[1] != 0) + return true; // Unknown modifier. switch (ExtraCode[0]) { default: @@ -1993,13 +2006,11 @@ bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, return false; } -bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI, - unsigned OpNo, - unsigned AsmVariant, - const char *ExtraCode, - raw_ostream &O) { +bool NVPTXAsmPrinter::PrintAsmMemoryOperand( + const MachineInstr *MI, unsigned OpNo, unsigned AsmVariant, + const char *ExtraCode, raw_ostream &O) { if (ExtraCode && ExtraCode[0]) - return true; // Unknown modifier + return true; // Unknown modifier O << '['; printMemOperand(MI, OpNo, O); @@ -2008,86 +2019,67 @@ bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI, return false; } -bool NVPTXAsmPrinter::ignoreLoc(const MachineInstr &MI) -{ - switch(MI.getOpcode()) { +void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum, + raw_ostream &O, const char *Modifier) { + const MachineOperand &MO = MI->getOperand(opNum); + switch (MO.getType()) { + case MachineOperand::MO_Register: + if (TargetRegisterInfo::isPhysicalRegister(MO.getReg())) { + if (MO.getReg() == NVPTX::VRDepot) + O << DEPOTNAME << getFunctionNumber(); + else + O << NVPTXInstPrinter::getRegisterName(MO.getReg()); + } else { + emitVirtualRegister(MO.getReg(), O); + } + return; + + case MachineOperand::MO_Immediate: + if (!Modifier) + O << MO.getImm(); + else if (strstr(Modifier, "vec") == Modifier) + printVecModifiedImmediate(MO, Modifier, O); + else + llvm_unreachable( + "Don't know how to handle modifier on immediate operand"); + return; + + case MachineOperand::MO_FPImmediate: + printFPConstant(MO.getFPImm(), O); + break; + + case MachineOperand::MO_GlobalAddress: + O << *getSymbol(MO.getGlobal()); + break; + + case MachineOperand::MO_MachineBasicBlock: + O << *MO.getMBB()->getSymbol(); + return; + default: - return false; - case NVPTX::CallArgBeginInst: case NVPTX::CallArgEndInst0: - case NVPTX::CallArgEndInst1: case NVPTX::CallArgF32: - case NVPTX::CallArgF64: case NVPTX::CallArgI16: - case NVPTX::CallArgI32: case NVPTX::CallArgI32imm: - case NVPTX::CallArgI64: case NVPTX::CallArgI8: - case NVPTX::CallArgParam: case NVPTX::CallVoidInst: - case NVPTX::CallVoidInstReg: case NVPTX::Callseq_End: - case NVPTX::CallVoidInstReg64: - case NVPTX::DeclareParamInst: case NVPTX::DeclareRetMemInst: - case NVPTX::DeclareRetRegInst: case NVPTX::DeclareRetScalarInst: - case NVPTX::DeclareScalarParamInst: case NVPTX::DeclareScalarRegInst: - case NVPTX::StoreParamF32: case NVPTX::StoreParamF64: - case NVPTX::StoreParamI16: case NVPTX::StoreParamI32: - case NVPTX::StoreParamI64: case NVPTX::StoreParamI8: - case NVPTX::StoreParamS32I8: case NVPTX::StoreParamU32I8: - case NVPTX::StoreParamS32I16: case NVPTX::StoreParamU32I16: - case NVPTX::StoreParamScalar2F32: case NVPTX::StoreParamScalar2F64: - case NVPTX::StoreParamScalar2I16: case NVPTX::StoreParamScalar2I32: - case NVPTX::StoreParamScalar2I64: case NVPTX::StoreParamScalar2I8: - case NVPTX::StoreParamScalar4F32: case NVPTX::StoreParamScalar4I16: - case NVPTX::StoreParamScalar4I32: case NVPTX::StoreParamScalar4I8: - case NVPTX::StoreParamV2F32: case NVPTX::StoreParamV2F64: - case NVPTX::StoreParamV2I16: case NVPTX::StoreParamV2I32: - case NVPTX::StoreParamV2I64: case NVPTX::StoreParamV2I8: - case NVPTX::StoreParamV4F32: case NVPTX::StoreParamV4I16: - case NVPTX::StoreParamV4I32: case NVPTX::StoreParamV4I8: - case NVPTX::StoreRetvalF32: case NVPTX::StoreRetvalF64: - case NVPTX::StoreRetvalI16: case NVPTX::StoreRetvalI32: - case NVPTX::StoreRetvalI64: case NVPTX::StoreRetvalI8: - case NVPTX::StoreRetvalScalar2F32: case NVPTX::StoreRetvalScalar2F64: - case NVPTX::StoreRetvalScalar2I16: case NVPTX::StoreRetvalScalar2I32: - case NVPTX::StoreRetvalScalar2I64: case NVPTX::StoreRetvalScalar2I8: - case NVPTX::StoreRetvalScalar4F32: case NVPTX::StoreRetvalScalar4I16: - case NVPTX::StoreRetvalScalar4I32: case NVPTX::StoreRetvalScalar4I8: - case NVPTX::StoreRetvalV2F32: case NVPTX::StoreRetvalV2F64: - case NVPTX::StoreRetvalV2I16: case NVPTX::StoreRetvalV2I32: - case NVPTX::StoreRetvalV2I64: case NVPTX::StoreRetvalV2I8: - case NVPTX::StoreRetvalV4F32: case NVPTX::StoreRetvalV4I16: - case NVPTX::StoreRetvalV4I32: case NVPTX::StoreRetvalV4I8: - case NVPTX::LastCallArgF32: case NVPTX::LastCallArgF64: - case NVPTX::LastCallArgI16: case NVPTX::LastCallArgI32: - case NVPTX::LastCallArgI32imm: case NVPTX::LastCallArgI64: - case NVPTX::LastCallArgI8: case NVPTX::LastCallArgParam: - case NVPTX::LoadParamMemF32: case NVPTX::LoadParamMemF64: - case NVPTX::LoadParamMemI16: case NVPTX::LoadParamMemI32: - case NVPTX::LoadParamMemI64: case NVPTX::LoadParamMemI8: - case NVPTX::LoadParamRegF32: case NVPTX::LoadParamRegF64: - case NVPTX::LoadParamRegI16: case NVPTX::LoadParamRegI32: - case NVPTX::LoadParamRegI64: case NVPTX::LoadParamRegI8: - case NVPTX::LoadParamScalar2F32: case NVPTX::LoadParamScalar2F64: - case NVPTX::LoadParamScalar2I16: case NVPTX::LoadParamScalar2I32: - case NVPTX::LoadParamScalar2I64: case NVPTX::LoadParamScalar2I8: - case NVPTX::LoadParamScalar4F32: case NVPTX::LoadParamScalar4I16: - case NVPTX::LoadParamScalar4I32: case NVPTX::LoadParamScalar4I8: - case NVPTX::LoadParamV2F32: case NVPTX::LoadParamV2F64: - case NVPTX::LoadParamV2I16: case NVPTX::LoadParamV2I32: - case NVPTX::LoadParamV2I64: case NVPTX::LoadParamV2I8: - case NVPTX::LoadParamV4F32: case NVPTX::LoadParamV4I16: - case NVPTX::LoadParamV4I32: case NVPTX::LoadParamV4I8: - case NVPTX::PrototypeInst: case NVPTX::DBG_VALUE: - return true; + llvm_unreachable("Operand type not supported."); } - return false; } -// Force static initialization. -extern "C" void LLVMInitializeNVPTXBackendAsmPrinter() { - RegisterAsmPrinter X(TheNVPTXTarget32); - RegisterAsmPrinter Y(TheNVPTXTarget64); -} +void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum, + raw_ostream &O, const char *Modifier) { + printOperand(MI, opNum, O); + if (Modifier && !strcmp(Modifier, "add")) { + O << ", "; + printOperand(MI, opNum + 1, O); + } else { + if (MI->getOperand(opNum + 1).isImm() && + MI->getOperand(opNum + 1).getImm() == 0) + return; // don't print ',0' or '+0' + O << "+"; + printOperand(MI, opNum + 1, O); + } +} void NVPTXAsmPrinter::emitSrcInText(StringRef filename, unsigned line) { std::stringstream temp; - LineReader * reader = this->getReader(filename.str()); + LineReader *reader = this->getReader(filename); temp << "\n//"; temp << filename.str(); temp << ":"; @@ -2095,32 +2087,29 @@ void NVPTXAsmPrinter::emitSrcInText(StringRef filename, unsigned line) { temp << " "; temp << reader->readLine(line); temp << "\n"; - this->OutStreamer.EmitRawText(Twine(temp.str())); + this->OutStreamer->EmitRawText(temp.str()); } - LineReader *NVPTXAsmPrinter::getReader(std::string filename) { - if (reader == NULL) { - reader = new LineReader(filename); + if (!reader) { + reader = new LineReader(filename); } if (reader->fileName() != filename) { delete reader; - reader = new LineReader(filename); + reader = new LineReader(filename); } return reader; } - -std::string -LineReader::readLine(unsigned lineNum) { +std::string LineReader::readLine(unsigned lineNum) { if (lineNum < theCurLine) { theCurLine = 0; - fstr.seekg(0,std::ios::beg); + fstr.seekg(0, std::ios::beg); } while (theCurLine < lineNum) { - fstr.getline(buff,500); + fstr.getline(buff, 500); theCurLine++; } return buff;