X-Git-Url: http://demsky.eecs.uci.edu/git/?a=blobdiff_plain;f=lib%2FTarget%2FNVPTX%2FNVPTXAsmPrinter.cpp;h=187b88c1d54adf075060667070b22eca0677f5d9;hb=d6663f565ca0b4400e4f07b78405e3ff6eb246f3;hp=b417d644b63209cc8dcb46d0f411661b4d9e37a1;hpb=82767327c59ede1f8663ec9b9a64a668993d501f;p=oota-llvm.git diff --git a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index b417d644b63..187b88c1d54 100644 --- a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -13,9 +13,11 @@ //===----------------------------------------------------------------------===// #include "NVPTXAsmPrinter.h" +#include "InstPrinter/NVPTXInstPrinter.h" #include "MCTargetDesc/NVPTXMCAsmInfo.h" #include "NVPTX.h" #include "NVPTXInstrInfo.h" +#include "NVPTXMachineFunctionInfo.h" #include "NVPTXMCExpr.h" #include "NVPTXRegisterInfo.h" #include "NVPTXTargetMachine.h" @@ -23,15 +25,15 @@ #include "cl_common_defines.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Analysis/ConstantFolding.h" -#include "llvm/Assembly/Writer.h" #include "llvm/CodeGen/Analysis.h" #include "llvm/CodeGen/MachineFrameInfo.h" #include "llvm/CodeGen/MachineModuleInfo.h" #include "llvm/CodeGen/MachineRegisterInfo.h" -#include "llvm/DebugInfo.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" @@ -42,26 +44,21 @@ #include "llvm/Support/Path.h" #include "llvm/Support/TargetRegistry.h" #include "llvm/Support/TimeValue.h" -#include "llvm/Target/Mangler.h" #include "llvm/Target/TargetLoweringObjectFile.h" #include using namespace llvm; -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::opt -InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, +static cl::opt +InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, cl::Hidden, cl::desc("NVPTX Specific: Emit source line in ptx file"), - cl::location(llvm::InterleaveSrcInPtx)); + cl::init(false)); namespace { /// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V @@ -129,13 +126,13 @@ const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { return MCConstantExpr::Create(CI->getZExtValue(), Ctx); if (const GlobalValue *GV = dyn_cast(CV)) - return MCSymbolRefExpr::Create(AP.Mang->getSymbol(GV), Ctx); + return MCSymbolRefExpr::Create(AP.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) + if (!CE) llvm_unreachable("Unknown constant value to lower!"); switch (CE->getOpcode()) { @@ -152,10 +149,25 @@ const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { 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()); + CE->printAsOperand(OS, /*PrintType=*/ false, + !AP.MF ? nullptr : AP.MF->getFunction()->getParent()); report_fatal_error(OS.str()); } + case Instruction::AddrSpaceCast: { + // Strip any addrspace(1)->addrspace(0) addrspace casts. These will be + // handled by the generic() logic in the MCExpr printer + PointerType *DstTy = cast(CE->getType()); + PointerType *SrcTy = cast(CE->getOperand(0)->getType()); + if (SrcTy->getAddressSpace() == 1 && DstTy->getAddressSpace() == 0) { + return LowerConstant(cast(CE->getOperand(0)), AP); + } + std::string S; + raw_string_ostream OS(S); + OS << "Unsupported expression in static initializer: "; + CE->printAsOperand(OS, /*PrintType=*/ false, + !AP.MF ? nullptr : 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 @@ -294,7 +306,7 @@ void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) { return; // Emit the line from the source file. - if (llvm::InterleaveSrcInPtx) + if (InterleaveSrc) this->emitSrcInText(fileName.str(), curLoc.getLine()); std::stringstream temp; @@ -311,16 +323,95 @@ void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) { MCInst Inst; lowerToMCInst(MI, Inst); - OutStreamer.EmitInstruction(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()); + const NVPTXSubtarget &ST = TM.getSubtarget(); + + // 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 (!ST.hasImageHandles()) { + if (lowerImageHandleOperand(MI, i, MCOp)) { + OutMI.addOperand(MCOp); + continue; + } + } + if (lowerOperand(MO, MCOp)) OutMI.addOperand(MCOp); } @@ -341,10 +432,10 @@ bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO, MO.getMBB()->getSymbol(), OutContext)); break; case MachineOperand::MO_ExternalSymbol: - MCOp = GetSymbolRef(MO, GetExternalSymbolSymbol(MO.getSymbolName())); + MCOp = GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName())); break; case MachineOperand::MO_GlobalAddress: - MCOp = GetSymbolRef(MO, Mang->getSymbol(MO.getGlobal())); + MCOp = GetSymbolRef(getSymbol(MO.getGlobal())); break; case MachineOperand::MO_FPImmediate: { const ConstantFP *Cnt = MO.getFPImm(); @@ -368,45 +459,45 @@ bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO, } unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) { - const TargetRegisterClass *RC = MRI->getRegClass(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"); + } - 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 = 0; - } else if (RC == &NVPTX::Int16RegsRegClass) { - Ret = (1 << 28); - } else if (RC == &NVPTX::Int32RegsRegClass) { - Ret = (2 << 28); - } else if (RC == &NVPTX::Int64RegsRegClass) { - Ret = (3 << 28); - } else if (RC == &NVPTX::Float32RegsRegClass) { - Ret = (4 << 28); - } else if (RC == &NVPTX::Float64RegsRegClass) { - Ret = (5 << 28); + // Insert the vreg number + Ret |= (RegNum & 0x0FFFFFFF); + return Ret; } else { - report_fatal_error("Bad register class"); + // 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; } - - // Insert the vreg number - Ret |= (RegNum & 0x0FFFFFFF); - return Ret; } -MCOperand NVPTXAsmPrinter::GetSymbolRef(const MachineOperand &MO, - const MCSymbol *Symbol) { +MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) { const MCExpr *Expr; - switch (MO.getTargetFlags()) { - default: { - Expr = MCSymbolRefExpr::Create(Symbol, MCSymbolRefExpr::VK_None, - OutContext); - break; - } - } + Expr = MCSymbolRefExpr::Create(Symbol, MCSymbolRefExpr::VK_None, + OutContext); return MCOperand::CreateExpr(Expr); } @@ -424,7 +515,7 @@ void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) { 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(); @@ -441,23 +532,7 @@ void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) { << " 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 totalsz = TD->getTypeAllocSize(Ty); unsigned retAlignment = 0; if (!llvm::getAlign(*F, 0, retAlignment)) retAlignment = TD->getABITypeAlignment(Ty); @@ -549,6 +624,19 @@ void NVPTXAsmPrinter::EmitFunctionBodyEnd() { VRegMapping.clear(); } +void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const { + unsigned RegNo = MI->getOperand(0).getReg(); + const TargetRegisterInfo *TRI = TM.getRegisterInfo(); + if (TRI->isVirtualRegister(RegNo)) { + OutStreamer.AddComment(Twine("implicit-def: ") + + getVirtualRegisterName(RegNo)); + } else { + OutStreamer.AddComment(Twine("implicit-def: ") + + TM.getRegisterInfo()->getName(RegNo)); + } + OutStreamer.AddBlankLine(); +} + void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, raw_ostream &O) const { // If the NVVM IR has some of reqntid* specified, then output @@ -600,23 +688,30 @@ void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, O << ".minnctapersm " << mincta << "\n"; } -void NVPTXAsmPrinter::getVirtualRegisterName(unsigned vr, bool isVec, - raw_ostream &O) { - const TargetRegisterClass *RC = MRI->getRegClass(vr); +std::string +NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const { + const TargetRegisterClass *RC = MRI->getRegClass(Reg); - DenseMap ®map = VRegMapping[RC]; - unsigned mapped_vr = regmap[vr]; + std::string Name; + raw_string_ostream NameStr(Name); - if (!isVec) { - O << getNVPTXRegClassStr(RC) << mapped_vr; - return; - } - report_fatal_error("Bad register!"); + 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, +void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr, raw_ostream &O) { - getVirtualRegisterName(vr, isVec, O); + O << getVirtualRegisterName(vr); } void NVPTXAsmPrinter::printVecModifiedImmediate( @@ -659,7 +754,7 @@ void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) { else O << ".func "; printReturnValStr(F, O); - O << *Mang->getSymbol(F) << "\n"; + O << *getSymbol(F) << "\n"; emitFunctionParamList(F, O); O << ";\n"; } @@ -674,12 +769,11 @@ static bool usedInGlobalVarDef(const Constant *C) { 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; } @@ -705,11 +799,10 @@ static bool usedInOneFunc(const User *U, Function const *&oneFunc) { (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) == false) return false; - } + return true; } @@ -727,7 +820,7 @@ static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) { 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) @@ -740,12 +833,11 @@ 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; @@ -772,10 +864,8 @@ void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) { 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 @@ -791,9 +881,9 @@ void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) { } } - if (!isa(*iter)) + if (!isa(U)) continue; - const Instruction *instr = cast(*iter); + const Instruction *instr = cast(U); const BasicBlock *bb = instr->getParent(); if (!bb) continue; @@ -818,10 +908,7 @@ void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) { 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); + for (DICompileUnit DIUnit : DbgFinder.compile_units()) { StringRef Filename(DIUnit.getFilename()); StringRef Dirname(DIUnit.getDirectory()); SmallString<128> FullPathName = Dirname; @@ -836,10 +923,7 @@ void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) { ++i; } - for (DebugInfoFinder::iterator I = DbgFinder.subprogram_begin(), - E = DbgFinder.subprogram_end(); - I != E; ++I) { - DISubprogram SP(*I); + for (DISubprogram SP : DbgFinder.subprograms()) { StringRef Filename(SP.getFilename()); StringRef Dirname(SP.getDirectory()); SmallString<128> FullPathName = Dirname; @@ -869,7 +953,7 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { const_cast(getObjFileLowering()) .Initialize(OutContext, TM); - Mang = new Mangler(OutContext, &TM); + Mang = new Mangler(TM.getDataLayout()); // Emit header before any dwarf directives are emitted below. emitHeader(M, OS1); @@ -996,6 +1080,8 @@ bool NVPTXAsmPrinter::doFinalization(Module &M) { for (i = 0; i < n; i++) global_list.insert(global_list.end(), gv_array[i]); + clearAnnotationCache(&M); + delete[] gv_array; return ret; @@ -1017,6 +1103,10 @@ bool NVPTXAsmPrinter::doFinalization(Module &M) { // external global variable with init -> .visible // external without init -> .extern // appending -> not allowed, assert. +// 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) { @@ -1042,6 +1132,9 @@ void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V, msg.append(V->getName().str()); msg.append("has unsupported appending linkage type"); llvm_unreachable(msg.c_str()); + } else if (!V->hasInternalLinkage() && + !V->hasPrivateLinkage()) { + O << ".weak "; } } } @@ -1052,10 +1145,15 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, // 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. @@ -1067,6 +1165,10 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, O << ".visible "; else O << ".extern "; + } else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() || + GVar->hasAvailableExternallyLinkage() || + GVar->hasCommonLinkage()) { + O << ".weak "; } if (llvm::isTexture(*GVar)) { @@ -1091,10 +1193,10 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, if (llvm::isSampler(*GVar)) { O << ".global .samplerref " << llvm::getSamplerName(*GVar); - const Constant *Initializer = NULL; + const Constant *Initializer = nullptr; if (GVar->hasInitializer()) Initializer = GVar->getInitializer(); - const ConstantInt *CI = NULL; + const ConstantInt *CI = nullptr; if (Initializer) CI = dyn_cast(Initializer); if (CI) { @@ -1134,7 +1236,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, O << "linear"; break; case 2: - assert(0 && "Anisotropic filtering is not supported"); + llvm_unreachable("Anisotropic filtering is not supported"); default: O << "nearest"; break; @@ -1161,7 +1263,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, return; } - const Function *demotedFunc = 0; + const Function *demotedFunc = nullptr; if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) { O << "// " << GVar->getName().str() << " has been demoted\n"; if (localDecls.find(demotedFunc) != localDecls.end()) @@ -1176,12 +1278,17 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, 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->isSingleValueType()) { O << " ."; // Special case: ABI requires that we use .u8 for predicates if (ETy->isIntegerTy(1)) @@ -1189,17 +1296,28 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, 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)) && - GVar->hasInitializer()) { - const 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 spefied. + 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().str() + + "' is not allowed in addrspace(" + + llvm::utostr_32(PTy->getAddressSpace()) + ")"; + report_fatal_error(warnMsg.c_str()); + } } } } else { @@ -1225,15 +1343,15 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, bufferAggregateConstant(Initializer, &aggBuffer); if (aggBuffer.numSymbols) { if (nvptxSubtarget.is64Bit()) { - O << " .u64 " << *Mang->getSymbol(GVar) << "["; + O << " .u64 " << *getSymbol(GVar) << "["; O << ElementSize / 8; } else { - O << " .u32 " << *Mang->getSymbol(GVar) << "["; + O << " .u32 " << *getSymbol(GVar) << "["; O << ElementSize / 4; } O << "]"; } else { - O << " .b8 " << *Mang->getSymbol(GVar) << "["; + O << " .b8 " << *getSymbol(GVar) << "["; O << ElementSize; O << "]"; } @@ -1241,7 +1359,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, aggBuffer.print(); O << "}"; } else { - O << " .b8 " << *Mang->getSymbol(GVar); + O << " .b8 " << *getSymbol(GVar); if (ElementSize) { O << "["; O << ElementSize; @@ -1249,7 +1367,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, } } } else { - O << " .b8 " << *Mang->getSymbol(GVar); + O << " .b8 " << *getSymbol(GVar); if (ElementSize) { O << "["; O << ElementSize; @@ -1258,7 +1376,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, } break; default: - assert(0 && "type not supported yet"); + llvm_unreachable("type not supported yet"); } } @@ -1333,7 +1451,7 @@ NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty, bool useB4PTR) const { return "u32"; } llvm_unreachable("unexpected type"); - return NULL; + return nullptr; } void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar, @@ -1352,11 +1470,11 @@ void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar, else O << " .align " << GVar->getAlignment(); - if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa(ETy)) { + if (ETy->isSingleValueType()) { O << " ."; O << getPTXFundamentalTypeStr(ETy); O << " "; - O << *Mang->getSymbol(GVar); + O << *getSymbol(GVar); return; } @@ -1371,20 +1489,20 @@ 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 << "]"; break; default: - assert(0 && "type not supported yet"); + llvm_unreachable("type not supported yet"); } return; } static unsigned int getOpenCLAlignment(const DataLayout *TD, Type *Ty) { - if (Ty->isPrimitiveType() || Ty->isIntegerTy() || isa(Ty)) + if (Ty->isSingleValueType()) return TD->getPrefTypeAlignment(Ty); const ArrayType *ATy = dyn_cast(Ty); @@ -1426,7 +1544,7 @@ void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I, int paramIndex, raw_ostream &O) { if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) || (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) - O << *Mang->getSymbol(I->getParent()) << "_param_" << paramIndex; + O << *getSymbol(I->getParent()) << "_param_" << paramIndex; else { std::string argName = I->getName(); const char *p = argName.c_str(); @@ -1481,24 +1599,38 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { 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 " << *Mang->getSymbol(F) << "_param_" - << paramIndex; - else // Default image is read_only - O << "\t.param .texref " << *Mang->getSymbol(F) << "_param_" - << paramIndex; - } else // Should be llvm::isSampler(*I) - O << "\t.param .samplerref " << *Mang->getSymbol(F) << "_param_" - << paramIndex; - continue; + 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; + } } if (PAL.hasAttribute(paramIndex + 1, Attribute::ByVal) == false) { - if (Ty->isVectorTy()) { - // Just print .param .b8 .align .param[size]; + 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); @@ -1554,7 +1686,7 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &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(); @@ -1578,7 +1710,7 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { 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); @@ -1676,9 +1808,9 @@ void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters( // 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 .s64 %rd<" << NVPTXNumRegisters << ">;\n"; // O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n"; - // O << "\t.reg .f64 %fl<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .f64 %fd<" << NVPTXNumRegisters << ">;\n"; // Emit declaration of the virtual registers or 'physical' registers for // each register class @@ -1738,13 +1870,35 @@ void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) { return; } if (const GlobalValue *GVar = dyn_cast(CPV)) { - O << *Mang->getSymbol(GVar); + 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 (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)) { - O << *Mang->getSymbol(GVar); + if (EmitGeneric && !isa(v) && !IsNonGenericPointer) { + O << "generic("; + O << *getSymbol(GVar); + O << ")"; + } else { + O << *getSymbol(GVar); + } return; } else { O << *LowerConstant(CPV, *this); @@ -1862,7 +2016,7 @@ void NVPTXAsmPrinter::bufferLEByte(const 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) @@ -1992,6 +2146,101 @@ bool NVPTXAsmPrinter::ignoreLoc(const MachineInstr &MI) { 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) { + if (ExtraCode && ExtraCode[0]) { + if (ExtraCode[1] != 0) + return true; // Unknown modifier. + + switch (ExtraCode[0]) { + default: + // See if this is a generic print operand + return AsmPrinter::PrintAsmOperand(MI, OpNo, AsmVariant, ExtraCode, O); + case 'r': + break; + } + } + + printOperand(MI, OpNo, O); + + return false; +} + +bool NVPTXAsmPrinter::PrintAsmMemoryOperand( + const MachineInstr *MI, unsigned OpNo, unsigned AsmVariant, + const char *ExtraCode, raw_ostream &O) { + if (ExtraCode && ExtraCode[0]) + return true; // Unknown modifier + + O << '['; + printMemOperand(MI, OpNo, O); + O << ']'; + + return false; +} + +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: + llvm_unreachable("Operand type not supported."); + } +} + +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); + } +} + + // Force static initialization. extern "C" void LLVMInitializeNVPTXBackendAsmPrinter() { RegisterAsmPrinter X(TheNVPTXTarget32); @@ -2012,7 +2261,7 @@ void NVPTXAsmPrinter::emitSrcInText(StringRef filename, unsigned line) { } LineReader *NVPTXAsmPrinter::getReader(std::string filename) { - if (reader == NULL) { + if (!reader) { reader = new LineReader(filename); }