X-Git-Url: http://demsky.eecs.uci.edu/git/?a=blobdiff_plain;f=lib%2FTarget%2FNVPTX%2FNVPTXAsmPrinter.cpp;h=3bbea400e53e441d0a787508b4064be092a4e23f;hb=344593ce6c914b3576328cf6bdf790cd2ea083af;hp=af6fc094331f8f75507e12802f33f7c0dd076d86;hpb=77f268945e4d4b777102d4a2089cb86137042883;p=oota-llvm.git diff --git a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index af6fc094331..3bbea400e53 100644 --- a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -17,8 +17,8 @@ #include "MCTargetDesc/NVPTXMCAsmInfo.h" #include "NVPTX.h" #include "NVPTXInstrInfo.h" -#include "NVPTXMachineFunctionInfo.h" #include "NVPTXMCExpr.h" +#include "NVPTXMachineFunctionInfo.h" #include "NVPTXRegisterInfo.h" #include "NVPTXTargetMachine.h" #include "NVPTXUtilities.h" @@ -27,6 +27,7 @@ #include "llvm/Analysis/ConstantFolding.h" #include "llvm/CodeGen/Analysis.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" @@ -36,6 +37,7 @@ #include "llvm/IR/Mangler.h" #include "llvm/IR/Module.h" #include "llvm/IR/Operator.h" +#include "llvm/MC/MCInst.h" #include "llvm/MC/MCStreamer.h" #include "llvm/MC/MCSymbol.h" #include "llvm/Support/CommandLine.h" @@ -45,6 +47,7 @@ #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; @@ -88,12 +91,9 @@ void VisitGlobalVariableForEmission( 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; for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i) @@ -111,159 +111,6 @@ void VisitGlobalVariableForEmission( } } -// @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.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: "; - CE->printAsOperand(OS, /*PrintType=*/ false, - !AP.MF ? 0 : 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 ? 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 - APInt OffsetAI(TD.getPointerSizeInBits(), 0); - cast(CE)->accumulateConstantOffset(TD, OffsetAI); - - const MCExpr *Base = LowerConstant(CE->getOperand(0), AP); - if (!OffsetAI) - return Base; - - int64_t Offset = OffsetAI.getSExtValue(); - 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) { if (!EmitLineNumbers) return; @@ -272,7 +119,7 @@ void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) { DebugLoc curLoc = MI.getDebugLoc(); - if (prevDebugLoc.isUnknown() && curLoc.isUnknown()) + if (!prevDebugLoc && !curLoc) return; if (prevDebugLoc == curLoc) @@ -280,303 +127,94 @@ void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) { prevDebugLoc = curLoc; - if (curLoc.isUnknown()) + if (!curLoc) return; - const MachineFunction *MF = MI.getParent()->getParent(); - //const TargetMachine &TM = MF->getTarget(); - - const LLVMContext &ctx = MF->getFunction()->getContext(); - DIScope Scope(curLoc.getScope(ctx)); - - assert((!Scope || Scope.isScope()) && - "Scope of a DebugLoc should be null or a DIScope."); + auto *Scope = cast_or_null(curLoc.getScope()); if (!Scope) 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 (InterleaveSrc) - this->emitSrcInText(fileName.str(), curLoc.getLine()); + this->emitSrcInText(fileName, curLoc.getLine()); std::stringstream temp; - temp << "\t.loc " << filenameMap[fileName.str()] << " " << curLoc.getLine() + temp << "\t.loc " << filenameMap[fileName] << " " << curLoc.getLine() << " " << curLoc.getCol(); - OutStreamer.EmitRawText(Twine(temp.str().c_str())); + 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); MCInst Inst; lowerToMCInst(MI, Inst); - EmitToStreamer(OutStreamer, 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(); - switch (MI->getOpcode()) { - default: return false; - case NVPTX::TEX_1D_F32_I32: - case NVPTX::TEX_1D_F32_F32: - case NVPTX::TEX_1D_F32_F32_LEVEL: - case NVPTX::TEX_1D_F32_F32_GRAD: - case NVPTX::TEX_1D_I32_I32: - case NVPTX::TEX_1D_I32_F32: - case NVPTX::TEX_1D_I32_F32_LEVEL: - case NVPTX::TEX_1D_I32_F32_GRAD: - case NVPTX::TEX_1D_ARRAY_F32_I32: - case NVPTX::TEX_1D_ARRAY_F32_F32: - case NVPTX::TEX_1D_ARRAY_F32_F32_LEVEL: - case NVPTX::TEX_1D_ARRAY_F32_F32_GRAD: - case NVPTX::TEX_1D_ARRAY_I32_I32: - case NVPTX::TEX_1D_ARRAY_I32_F32: - case NVPTX::TEX_1D_ARRAY_I32_F32_LEVEL: - case NVPTX::TEX_1D_ARRAY_I32_F32_GRAD: - case NVPTX::TEX_2D_F32_I32: - case NVPTX::TEX_2D_F32_F32: - case NVPTX::TEX_2D_F32_F32_LEVEL: - case NVPTX::TEX_2D_F32_F32_GRAD: - case NVPTX::TEX_2D_I32_I32: - case NVPTX::TEX_2D_I32_F32: - case NVPTX::TEX_2D_I32_F32_LEVEL: - case NVPTX::TEX_2D_I32_F32_GRAD: - case NVPTX::TEX_2D_ARRAY_F32_I32: - case NVPTX::TEX_2D_ARRAY_F32_F32: - case NVPTX::TEX_2D_ARRAY_F32_F32_LEVEL: - case NVPTX::TEX_2D_ARRAY_F32_F32_GRAD: - case NVPTX::TEX_2D_ARRAY_I32_I32: - case NVPTX::TEX_2D_ARRAY_I32_F32: - case NVPTX::TEX_2D_ARRAY_I32_F32_LEVEL: - case NVPTX::TEX_2D_ARRAY_I32_F32_GRAD: - case NVPTX::TEX_3D_F32_I32: - case NVPTX::TEX_3D_F32_F32: - case NVPTX::TEX_3D_F32_F32_LEVEL: - case NVPTX::TEX_3D_F32_F32_GRAD: - case NVPTX::TEX_3D_I32_I32: - case NVPTX::TEX_3D_I32_F32: - case NVPTX::TEX_3D_I32_F32_LEVEL: - case NVPTX::TEX_3D_I32_F32_GRAD: - { + 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) { + if (OpNo == 4 && MO.isImm()) { lowerImageHandleSymbol(MO.getImm(), MCOp); return true; } - if (OpNo == 5) { + if (OpNo == 5 && MO.isImm() && !(MCID.TSFlags & NVPTXII::IsTexModeUnifiedFlag)) { lowerImageHandleSymbol(MO.getImm(), MCOp); return true; } return false; - } - case NVPTX::SULD_1D_I8_TRAP: - case NVPTX::SULD_1D_I16_TRAP: - case NVPTX::SULD_1D_I32_TRAP: - case NVPTX::SULD_1D_ARRAY_I8_TRAP: - case NVPTX::SULD_1D_ARRAY_I16_TRAP: - case NVPTX::SULD_1D_ARRAY_I32_TRAP: - case NVPTX::SULD_2D_I8_TRAP: - case NVPTX::SULD_2D_I16_TRAP: - case NVPTX::SULD_2D_I32_TRAP: - case NVPTX::SULD_2D_ARRAY_I8_TRAP: - case NVPTX::SULD_2D_ARRAY_I16_TRAP: - case NVPTX::SULD_2D_ARRAY_I32_TRAP: - case NVPTX::SULD_3D_I8_TRAP: - case NVPTX::SULD_3D_I16_TRAP: - case NVPTX::SULD_3D_I32_TRAP: { - // This is a V1 surface load, so operand 1 is a surfref - if (OpNo == 1) { - lowerImageHandleSymbol(MO.getImm(), MCOp); - return true; - } + } else if (MCID.TSFlags & NVPTXII::IsSuldMask) { + unsigned VecSize = + 1 << (((MCID.TSFlags & NVPTXII::IsSuldMask) >> NVPTXII::IsSuldShift) - 1); - return false; - } - case NVPTX::SULD_1D_V2I8_TRAP: - case NVPTX::SULD_1D_V2I16_TRAP: - case NVPTX::SULD_1D_V2I32_TRAP: - case NVPTX::SULD_1D_ARRAY_V2I8_TRAP: - case NVPTX::SULD_1D_ARRAY_V2I16_TRAP: - case NVPTX::SULD_1D_ARRAY_V2I32_TRAP: - case NVPTX::SULD_2D_V2I8_TRAP: - case NVPTX::SULD_2D_V2I16_TRAP: - case NVPTX::SULD_2D_V2I32_TRAP: - case NVPTX::SULD_2D_ARRAY_V2I8_TRAP: - case NVPTX::SULD_2D_ARRAY_V2I16_TRAP: - case NVPTX::SULD_2D_ARRAY_V2I32_TRAP: - case NVPTX::SULD_3D_V2I8_TRAP: - case NVPTX::SULD_3D_V2I16_TRAP: - case NVPTX::SULD_3D_V2I32_TRAP: { - // This is a V2 surface load, so operand 2 is a surfref - if (OpNo == 2) { + // 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; - } - case NVPTX::SULD_1D_V4I8_TRAP: - case NVPTX::SULD_1D_V4I16_TRAP: - case NVPTX::SULD_1D_V4I32_TRAP: - case NVPTX::SULD_1D_ARRAY_V4I8_TRAP: - case NVPTX::SULD_1D_ARRAY_V4I16_TRAP: - case NVPTX::SULD_1D_ARRAY_V4I32_TRAP: - case NVPTX::SULD_2D_V4I8_TRAP: - case NVPTX::SULD_2D_V4I16_TRAP: - case NVPTX::SULD_2D_V4I32_TRAP: - case NVPTX::SULD_2D_ARRAY_V4I8_TRAP: - case NVPTX::SULD_2D_ARRAY_V4I16_TRAP: - case NVPTX::SULD_2D_ARRAY_V4I32_TRAP: - case NVPTX::SULD_3D_V4I8_TRAP: - case NVPTX::SULD_3D_V4I16_TRAP: - case NVPTX::SULD_3D_V4I32_TRAP: { - // This is a V4 surface load, so operand 4 is a surfref - if (OpNo == 4) { - lowerImageHandleSymbol(MO.getImm(), MCOp); - return true; - } - - return false; - } - case NVPTX::SUST_B_1D_B8_TRAP: - case NVPTX::SUST_B_1D_B16_TRAP: - case NVPTX::SUST_B_1D_B32_TRAP: - case NVPTX::SUST_B_1D_V2B8_TRAP: - case NVPTX::SUST_B_1D_V2B16_TRAP: - case NVPTX::SUST_B_1D_V2B32_TRAP: - case NVPTX::SUST_B_1D_V4B8_TRAP: - case NVPTX::SUST_B_1D_V4B16_TRAP: - case NVPTX::SUST_B_1D_V4B32_TRAP: - case NVPTX::SUST_B_1D_ARRAY_B8_TRAP: - case NVPTX::SUST_B_1D_ARRAY_B16_TRAP: - case NVPTX::SUST_B_1D_ARRAY_B32_TRAP: - case NVPTX::SUST_B_1D_ARRAY_V2B8_TRAP: - case NVPTX::SUST_B_1D_ARRAY_V2B16_TRAP: - case NVPTX::SUST_B_1D_ARRAY_V2B32_TRAP: - case NVPTX::SUST_B_1D_ARRAY_V4B8_TRAP: - case NVPTX::SUST_B_1D_ARRAY_V4B16_TRAP: - case NVPTX::SUST_B_1D_ARRAY_V4B32_TRAP: - case NVPTX::SUST_B_2D_B8_TRAP: - case NVPTX::SUST_B_2D_B16_TRAP: - case NVPTX::SUST_B_2D_B32_TRAP: - case NVPTX::SUST_B_2D_V2B8_TRAP: - case NVPTX::SUST_B_2D_V2B16_TRAP: - case NVPTX::SUST_B_2D_V2B32_TRAP: - case NVPTX::SUST_B_2D_V4B8_TRAP: - case NVPTX::SUST_B_2D_V4B16_TRAP: - case NVPTX::SUST_B_2D_V4B32_TRAP: - case NVPTX::SUST_B_2D_ARRAY_B8_TRAP: - case NVPTX::SUST_B_2D_ARRAY_B16_TRAP: - case NVPTX::SUST_B_2D_ARRAY_B32_TRAP: - case NVPTX::SUST_B_2D_ARRAY_V2B8_TRAP: - case NVPTX::SUST_B_2D_ARRAY_V2B16_TRAP: - case NVPTX::SUST_B_2D_ARRAY_V2B32_TRAP: - case NVPTX::SUST_B_2D_ARRAY_V4B8_TRAP: - case NVPTX::SUST_B_2D_ARRAY_V4B16_TRAP: - case NVPTX::SUST_B_2D_ARRAY_V4B32_TRAP: - case NVPTX::SUST_B_3D_B8_TRAP: - case NVPTX::SUST_B_3D_B16_TRAP: - case NVPTX::SUST_B_3D_B32_TRAP: - case NVPTX::SUST_B_3D_V2B8_TRAP: - case NVPTX::SUST_B_3D_V2B16_TRAP: - case NVPTX::SUST_B_3D_V2B32_TRAP: - case NVPTX::SUST_B_3D_V4B8_TRAP: - case NVPTX::SUST_B_3D_V4B16_TRAP: - case NVPTX::SUST_B_3D_V4B32_TRAP: - case NVPTX::SUST_P_1D_B8_TRAP: - case NVPTX::SUST_P_1D_B16_TRAP: - case NVPTX::SUST_P_1D_B32_TRAP: - case NVPTX::SUST_P_1D_V2B8_TRAP: - case NVPTX::SUST_P_1D_V2B16_TRAP: - case NVPTX::SUST_P_1D_V2B32_TRAP: - case NVPTX::SUST_P_1D_V4B8_TRAP: - case NVPTX::SUST_P_1D_V4B16_TRAP: - case NVPTX::SUST_P_1D_V4B32_TRAP: - case NVPTX::SUST_P_1D_ARRAY_B8_TRAP: - case NVPTX::SUST_P_1D_ARRAY_B16_TRAP: - case NVPTX::SUST_P_1D_ARRAY_B32_TRAP: - case NVPTX::SUST_P_1D_ARRAY_V2B8_TRAP: - case NVPTX::SUST_P_1D_ARRAY_V2B16_TRAP: - case NVPTX::SUST_P_1D_ARRAY_V2B32_TRAP: - case NVPTX::SUST_P_1D_ARRAY_V4B8_TRAP: - case NVPTX::SUST_P_1D_ARRAY_V4B16_TRAP: - case NVPTX::SUST_P_1D_ARRAY_V4B32_TRAP: - case NVPTX::SUST_P_2D_B8_TRAP: - case NVPTX::SUST_P_2D_B16_TRAP: - case NVPTX::SUST_P_2D_B32_TRAP: - case NVPTX::SUST_P_2D_V2B8_TRAP: - case NVPTX::SUST_P_2D_V2B16_TRAP: - case NVPTX::SUST_P_2D_V2B32_TRAP: - case NVPTX::SUST_P_2D_V4B8_TRAP: - case NVPTX::SUST_P_2D_V4B16_TRAP: - case NVPTX::SUST_P_2D_V4B32_TRAP: - case NVPTX::SUST_P_2D_ARRAY_B8_TRAP: - case NVPTX::SUST_P_2D_ARRAY_B16_TRAP: - case NVPTX::SUST_P_2D_ARRAY_B32_TRAP: - case NVPTX::SUST_P_2D_ARRAY_V2B8_TRAP: - case NVPTX::SUST_P_2D_ARRAY_V2B16_TRAP: - case NVPTX::SUST_P_2D_ARRAY_V2B32_TRAP: - case NVPTX::SUST_P_2D_ARRAY_V4B8_TRAP: - case NVPTX::SUST_P_2D_ARRAY_V4B16_TRAP: - case NVPTX::SUST_P_2D_ARRAY_V4B32_TRAP: - case NVPTX::SUST_P_3D_B8_TRAP: - case NVPTX::SUST_P_3D_B16_TRAP: - case NVPTX::SUST_P_3D_B32_TRAP: - case NVPTX::SUST_P_3D_V2B8_TRAP: - case NVPTX::SUST_P_3D_V2B16_TRAP: - case NVPTX::SUST_P_3D_V2B32_TRAP: - case NVPTX::SUST_P_3D_V4B8_TRAP: - case NVPTX::SUST_P_3D_V4B16_TRAP: - case NVPTX::SUST_P_3D_V4B32_TRAP: { + } else if (MCID.TSFlags & NVPTXII::IsSustFlag) { // This is a surface store, so operand 0 is a surfref - if (OpNo == 0) { + if (OpNo == 0 && MO.isImm()) { lowerImageHandleSymbol(MO.getImm(), MCOp); return true; } return false; - } - case NVPTX::TXQ_CHANNEL_ORDER: - case NVPTX::TXQ_CHANNEL_DATA_TYPE: - case NVPTX::TXQ_WIDTH: - case NVPTX::TXQ_HEIGHT: - case NVPTX::TXQ_DEPTH: - case NVPTX::TXQ_ARRAY_SIZE: - case NVPTX::TXQ_NUM_SAMPLES: - case NVPTX::TXQ_NUM_MIPMAP_LEVELS: - case NVPTX::SUQ_CHANNEL_ORDER: - case NVPTX::SUQ_CHANNEL_DATA_TYPE: - case NVPTX::SUQ_WIDTH: - case NVPTX::SUQ_HEIGHT: - case NVPTX::SUQ_DEPTH: - case NVPTX::SUQ_ARRAY_SIZE: { + } else if (MCID.TSFlags & NVPTXII::IsSurfTexQueryFlag) { // This is a query, so operand 1 is a surfref/texref - if (OpNo == 1) { + if (OpNo == 1 && MO.isImm()) { lowerImageHandleSymbol(MO.getImm(), MCOp); return true; } return false; } - } + + return false; } void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) { @@ -587,19 +225,17 @@ void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) { const char *Sym = MFI->getImageHandleSymbol(Index); std::string *SymNamePtr = nvTM.getManagedStrPool()->getManagedString(Sym); - MCOp = GetSymbolRef(OutContext.GetOrCreateSymbol( + 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())))); + OutContext.getOrCreateSymbol(Twine(MO.getSymbolName())))); return; } @@ -607,7 +243,7 @@ void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) { const MachineOperand &MO = MI->getOperand(i); MCOperand MCOp; - if (!ST.hasImageHandles()) { + if (!nvptxSubtarget->hasImageHandles()) { if (lowerImageHandleOperand(MI, i, MCOp)) { OutMI.addOperand(MCOp); continue; @@ -624,13 +260,13 @@ bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO, switch (MO.getType()) { default: llvm_unreachable("unknown operand type"); case MachineOperand::MO_Register: - MCOp = MCOperand::CreateReg(encodeVirtualRegister(MO.getReg())); + MCOp = MCOperand::createReg(encodeVirtualRegister(MO.getReg())); break; case MachineOperand::MO_Immediate: - MCOp = MCOperand::CreateImm(MO.getImm()); + MCOp = MCOperand::createImm(MO.getImm()); break; case MachineOperand::MO_MachineBasicBlock: - MCOp = MCOperand::CreateExpr(MCSymbolRefExpr::Create( + MCOp = MCOperand::createExpr(MCSymbolRefExpr::Create( MO.getMBB()->getSymbol(), OutContext)); break; case MachineOperand::MO_ExternalSymbol: @@ -646,11 +282,11 @@ bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO, switch (Cnt->getType()->getTypeID()) { default: report_fatal_error("Unsupported FP type"); break; case Type::FloatTyID: - MCOp = MCOperand::CreateExpr( + MCOp = MCOperand::createExpr( NVPTXFloatMCExpr::CreateConstantFPSingle(Val, OutContext)); break; case Type::DoubleTyID: - MCOp = MCOperand::CreateExpr( + MCOp = MCOperand::createExpr( NVPTXFloatMCExpr::CreateConstantFPDouble(Val, OutContext)); break; } @@ -700,16 +336,16 @@ MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) { const MCExpr *Expr; Expr = MCSymbolRefExpr::Create(Symbol, MCSymbolRefExpr::VK_None, OutContext); - return MCOperand::CreateExpr(Expr); + 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; @@ -732,33 +368,15 @@ void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) { } 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"); - } + } 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); @@ -794,6 +412,42 @@ void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF, printReturnValStr(F, O); } +// Return true if MBB is the header of a loop marked with +// llvm.loop.unroll.disable. +// TODO: consider "#pragma unroll 1" which is equivalent to "#pragma nounroll". +bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll( + const MachineBasicBlock &MBB) const { + MachineLoopInfo &LI = getAnalysis(); + // 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); @@ -821,38 +475,37 @@ void NVPTXAsmPrinter::EmitFunctionEntryLabel() { if (llvm::isKernelFunction(*F)) emitKernelFunctionDirectives(*F, O); - OutStreamer.EmitRawText(O.str()); + OutStreamer->EmitRawText(O.str()); prevDebugLoc = DebugLoc(); } void NVPTXAsmPrinter::EmitFunctionBodyStart() { VRegMapping.clear(); - OutStreamer.EmitRawText(StringRef("{\n")); + 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")); + OutStreamer->EmitRawText(StringRef("}\n")); 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)); + if (TargetRegisterInfo::isVirtualRegister(RegNo)) { + OutStreamer->AddComment(Twine("implicit-def: ") + + getVirtualRegisterName(RegNo)); } else { - OutStreamer.AddComment(Twine("implicit-def: ") + - TM.getRegisterInfo()->getName(RegNo)); + OutStreamer->AddComment(Twine("implicit-def: ") + + nvptxSubtarget->getRegisterInfo()->getName(RegNo)); } - OutStreamer.AddBlankLine(); + OutStreamer->AddBlankLine(); } void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, @@ -862,15 +515,15 @@ void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, // If none of reqntid* is specified, don't output reqntid directive. unsigned reqntidx, reqntidy, reqntidz; bool specified = false; - if (llvm::getReqNTIDx(F, reqntidx) == false) + if (!llvm::getReqNTIDx(F, reqntidx)) reqntidx = 1; else specified = true; - if (llvm::getReqNTIDy(F, reqntidy) == false) + if (!llvm::getReqNTIDy(F, reqntidy)) reqntidy = 1; else specified = true; - if (llvm::getReqNTIDz(F, reqntidz) == false) + if (!llvm::getReqNTIDz(F, reqntidz)) reqntidz = 1; else specified = true; @@ -884,15 +537,15 @@ void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, // If none of maxntid* is specified, don't output maxntid directive. unsigned maxntidx, maxntidy, maxntidz; specified = false; - if (llvm::getMaxNTIDx(F, maxntidx) == false) + if (!llvm::getMaxNTIDx(F, maxntidx)) maxntidx = 1; else specified = true; - if (llvm::getMaxNTIDy(F, maxntidy) == false) + if (!llvm::getMaxNTIDy(F, maxntidy)) maxntidy = 1; else specified = true; - if (llvm::getMaxNTIDz(F, maxntidz) == false) + if (!llvm::getMaxNTIDz(F, maxntidz)) maxntidz = 1; else specified = true; @@ -982,7 +635,7 @@ static bool usedInGlobalVarDef(const Constant *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; } @@ -997,7 +650,7 @@ static bool usedInGlobalVarDef(const Constant *C) { 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; } @@ -1012,13 +665,8 @@ static bool usedInOneFunc(const User *U, Function const *&oneFunc) { 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 (const User *UU : U->users()) - if (usedInOneFunc(UU, oneFunc) == false) + if (!usedInOneFunc(UU, oneFunc)) return false; return true; @@ -1032,16 +680,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; @@ -1126,37 +774,45 @@ void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) { DbgFinder.processModule(M); unsigned i = 1; - for (DICompileUnit DIUnit : DbgFinder.compile_units()) { - StringRef Filename(DIUnit.getFilename()); - StringRef Dirname(DIUnit.getDirectory()); + for (const DICompileUnit *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 (DISubprogram SP : DbgFinder.subprograms()) { - StringRef Filename(SP.getFilename()); - StringRef Dirname(SP.getDirectory()); + for (DISubprogram *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) { + // 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); @@ -1174,23 +830,24 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { 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(); + 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; @@ -1228,25 +885,27 @@ void NVPTXAsmPrinter::emitGlobals(const Module &M) { OS2 << '\n'; - OutStreamer.EmitRawText(OS2.str()); + 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"; } @@ -1256,7 +915,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"; @@ -1266,7 +925,6 @@ 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) { @@ -1321,10 +979,14 @@ 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) { - if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) { + if (static_cast(TM).getDrvInterface() == NVPTX::CUDA) { if (V->hasExternalLinkage()) { if (isa(V)) { const GlobalVariable *GVar = cast(V); @@ -1343,9 +1005,12 @@ void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V, 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 "; } } } @@ -1356,10 +1021,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. @@ -1371,6 +1041,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)) { @@ -1395,10 +1069,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) { @@ -1438,7 +1112,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; @@ -1465,9 +1139,9 @@ 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"; + O << "// " << GVar->getName() << " has been demoted\n"; if (localDecls.find(demotedFunc) != localDecls.end()) localDecls[demotedFunc].push_back(GVar); else { @@ -1480,12 +1154,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->isSingleValueType()) { + if (ETy->isFloatingPointTy() || ETy->isIntegerTy() || ETy->isPointerTy()) { O << " ."; // Special case: ABI requires that we use .u8 for predicates if (ETy->isIntegerTy(1)) @@ -1497,13 +1176,25 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *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 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 { @@ -1528,7 +1219,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, AggBuffer aggBuffer(ElementSize, O, *this); bufferAggregateConstant(Initializer, &aggBuffer); if (aggBuffer.numSymbols) { - if (nvptxSubtarget.is64Bit()) { + if (static_cast(TM).is64Bit()) { O << " .u64 " << *getSymbol(GVar) << "["; O << ElementSize / 8; } else { @@ -1562,7 +1253,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, } break; default: - assert(0 && "type not supported yet"); + llvm_unreachable("type not supported yet"); } } @@ -1626,7 +1317,7 @@ NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty, bool useB4PTR) const { case Type::DoubleTyID: return "f64"; case Type::PointerTyID: - if (nvptxSubtarget.is64Bit()) + if (static_cast(TM).is64Bit()) if (useB4PTR) return "b64"; else @@ -1637,7 +1328,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, @@ -1656,7 +1347,7 @@ void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar, else O << " .align " << GVar->getAlignment(); - if (ETy->isSingleValueType()) { + if (ETy->isFloatingPointTy() || ETy->isIntegerTy() || ETy->isPointerTy()) { O << " ."; O << getPTXFundamentalTypeStr(ETy); O << " "; @@ -1682,7 +1373,7 @@ void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar, O << "]"; break; default: - assert(0 && "type not supported yet"); + llvm_unreachable("type not supported yet"); } return; } @@ -1695,17 +1386,6 @@ static unsigned int getOpenCLAlignment(const DataLayout *TD, Type *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; @@ -1728,50 +1408,22 @@ static unsigned int getOpenCLAlignment(const DataLayout *TD, Type *Ty) { void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I, int paramIndex, raw_ostream &O) { - if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) || - (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) - O << *getSymbol(I->getParent()) << "_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) { const DataLayout *TD = TM.getDataLayout(); const AttributeSet &PAL = F->getAttributes(); - const TargetLowering *TLI = TM.getTargetLowering(); + 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"; @@ -1790,21 +1442,21 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { if (isImage(*I)) { std::string sname = I->getName(); if (isImageWriteOnly(*I) || isImageReadWrite(*I)) { - if (nvptxSubtarget.hasImageHandles()) + 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()) + if (nvptxSubtarget->hasImageHandles()) O << "\t.param .u64 .ptr .texref "; else O << "\t.param .texref "; O << *CurrentFnSym << "_param_" << paramIndex; } } else { - if (nvptxSubtarget.hasImageHandles()) + if (nvptxSubtarget->hasImageHandles()) O << "\t.param .u64 .ptr .samplerref "; else O << "\t.param .samplerref "; @@ -1814,7 +1466,7 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { } } - if (PAL.hasAttribute(paramIndex + 1, Attribute::ByVal) == false) { + if (!PAL.hasAttribute(paramIndex + 1, Attribute::ByVal)) { if (Ty->isAggregateType() || Ty->isVectorTy()) { // Just print .param .align .b8 .param[size]; // = PAL.getparamalignment @@ -1837,7 +1489,8 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { // 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) { @@ -1957,7 +1610,7 @@ void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters( // 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 @@ -1966,7 +1619,7 @@ void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters( if (NumBytes) { O << "\t.local .align " << MFI->getMaxAlignment() << " .b8 \t" << DEPOTNAME << getFunctionNumber() << "[" << NumBytes << "];\n"; - if (nvptxSubtarget.is64Bit()) { + if (static_cast(MF.getTarget()).is64Bit()) { O << "\t.reg .b64 \t%SP;\n"; O << "\t.reg .b64 \t%SPL;\n"; } else { @@ -1994,9 +1647,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 @@ -2014,7 +1667,7 @@ void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters( } } - OutStreamer.EmitRawText(O.str()); + OutStreamer->EmitRawText(O.str()); } void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) { @@ -2087,7 +1740,7 @@ void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) { } return; } else { - O << *LowerConstant(CPV, *this); + O << *lowerConstant(CPV); return; } } @@ -2113,12 +1766,11 @@ void NVPTXAsmPrinter::bufferLEByte(const 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(); + 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(); + short int16 = (short)cast(CPV)->getZExtValue(); ptr = (unsigned char *)&int16; aggBuffer->addBytes(ptr, 2, Bytes); } else if (ETy == Type::getInt32Ty(CPV->getContext())) { @@ -2129,7 +1781,7 @@ void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, break; } else if (const ConstantExpr *Cexpr = dyn_cast(CPV)) { if (const ConstantInt *constInt = dyn_cast( - ConstantFoldConstantExpression(Cexpr, TD))) { + ConstantFoldConstantExpression(Cexpr, *TD))) { int int32 = (int)(constInt->getZExtValue()); ptr = (unsigned char *)&int32; aggBuffer->addBytes(ptr, 4, Bytes); @@ -2137,7 +1789,7 @@ void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, } if (Cexpr->getOpcode() == Instruction::PtrToInt) { Value *v = Cexpr->getOperand(0)->stripPointerCasts(); - aggBuffer->addSymbol(v); + aggBuffer->addSymbol(v, Cexpr->getOperand(0)); aggBuffer->addZeros(4); break; } @@ -2151,7 +1803,7 @@ void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, break; } else if (const ConstantExpr *Cexpr = dyn_cast(CPV)) { if (const ConstantInt *constInt = dyn_cast( - ConstantFoldConstantExpression(Cexpr, TD))) { + ConstantFoldConstantExpression(Cexpr, *TD))) { long long int64 = (long long)(constInt->getZExtValue()); ptr = (unsigned char *)&int64; aggBuffer->addBytes(ptr, 8, Bytes); @@ -2159,7 +1811,7 @@ void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, } if (Cexpr->getOpcode() == Instruction::PtrToInt) { Value *v = Cexpr->getOperand(0)->stripPointerCasts(); - aggBuffer->addSymbol(v); + aggBuffer->addSymbol(v, Cexpr->getOperand(0)); aggBuffer->addZeros(8); break; } @@ -2188,10 +1840,10 @@ void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, } case Type::PointerTyID: { if (const GlobalValue *GVar = dyn_cast(CPV)) { - aggBuffer->addSymbol(GVar); + aggBuffer->addSymbol(GVar, GVar); } else if (const ConstantExpr *Cexpr = dyn_cast(CPV)) { const Value *v = Cexpr->stripPointerCasts(); - aggBuffer->addSymbol(v); + aggBuffer->addSymbol(v, Cexpr); } unsigned int s = TD->getTypeAllocSize(CPV->getType()); aggBuffer->addZeros(s); @@ -2332,6 +1984,212 @@ bool NVPTXAsmPrinter::ignoreLoc(const MachineInstr &MI) { return false; } +/// lowerConstantForGV - Return an MCExpr for the given Constant. This is mostly +/// a copy from AsmPrinter::lowerConstant, except customized to only handle +/// expressions that are representable in PTX and create +/// NVPTXGenericMCSymbolRefExpr nodes for addrspacecast instructions. +const MCExpr * +NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV, bool ProcessingGeneric) { + MCContext &Ctx = 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)) { + const MCSymbolRefExpr *Expr = + MCSymbolRefExpr::Create(getSymbol(GV), Ctx); + if (ProcessingGeneric) { + return NVPTXGenericMCSymbolRefExpr::Create(Expr, Ctx); + } else { + return Expr; + } + } + + const ConstantExpr *CE = dyn_cast(CV); + if (!CE) { + 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, *TM.getDataLayout())) + if (C != CE) + return lowerConstantForGV(C, ProcessingGeneric); + + // Otherwise report the problem to the user. + { + std::string S; + raw_string_ostream OS(S); + OS << "Unsupported expression in static initializer: "; + CE->printAsOperand(OS, /*PrintType=*/false, + !MF ? nullptr : MF->getFunction()->getParent()); + report_fatal_error(OS.str()); + } + + case Instruction::AddrSpaceCast: { + // Strip the addrspacecast and pass along the operand + PointerType *DstTy = cast(CE->getType()); + if (DstTy->getAddressSpace() == 0) { + return lowerConstantForGV(cast(CE->getOperand(0)), true); + } + std::string S; + raw_string_ostream OS(S); + OS << "Unsupported expression in static initializer: "; + CE->printAsOperand(OS, /*PrintType=*/ false, + !MF ? 0 : MF->getFunction()->getParent()); + report_fatal_error(OS.str()); + } + + case Instruction::GetElementPtr: { + const DataLayout &DL = *TM.getDataLayout(); + + // Generate a symbolic expression for the byte address + APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0); + cast(CE)->accumulateConstantOffset(DL, OffsetAI); + + const MCExpr *Base = lowerConstantForGV(CE->getOperand(0), + ProcessingGeneric); + if (!OffsetAI) + return Base; + + int64_t Offset = OffsetAI.getSExtValue(); + 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 lowerConstantForGV(CE->getOperand(0), ProcessingGeneric); + + case Instruction::IntToPtr: { + const DataLayout &DL = *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, DL.getIntPtrType(CV->getType()), + false/*ZExt*/); + return lowerConstantForGV(Op, ProcessingGeneric); + } + + case Instruction::PtrToInt: { + const DataLayout &DL = *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 = lowerConstantForGV(Op, ProcessingGeneric); + + // We can emit the pointer value into this slot if the slot is an + // integer slot equal to the size of the pointer. + if (DL.getTypeAllocSize(Ty) == DL.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 = DL.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: { + const MCExpr *LHS = lowerConstantForGV(CE->getOperand(0), ProcessingGeneric); + const MCExpr *RHS = lowerConstantForGV(CE->getOperand(1), ProcessingGeneric); + switch (CE->getOpcode()) { + default: llvm_unreachable("Unknown binary operator constant cast expr"); + case Instruction::Add: return MCBinaryExpr::CreateAdd(LHS, RHS, Ctx); + } + } + } +} + +// Copy of MCExpr::print customized for NVPTX +void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) { + switch (Expr.getKind()) { + case MCExpr::Target: + return cast(&Expr)->PrintImpl(OS); + case MCExpr::Constant: + OS << cast(Expr).getValue(); + return; + + case MCExpr::SymbolRef: { + const MCSymbolRefExpr &SRE = cast(Expr); + const MCSymbol &Sym = SRE.getSymbol(); + OS << Sym; + return; + } + + case MCExpr::Unary: { + const MCUnaryExpr &UE = cast(Expr); + switch (UE.getOpcode()) { + case MCUnaryExpr::LNot: OS << '!'; break; + case MCUnaryExpr::Minus: OS << '-'; break; + case MCUnaryExpr::Not: OS << '~'; break; + case MCUnaryExpr::Plus: OS << '+'; break; + } + printMCExpr(*UE.getSubExpr(), OS); + return; + } + + case MCExpr::Binary: { + const MCBinaryExpr &BE = cast(Expr); + + // Only print parens around the LHS if it is non-trivial. + if (isa(BE.getLHS()) || isa(BE.getLHS()) || + isa(BE.getLHS())) { + printMCExpr(*BE.getLHS(), OS); + } else { + OS << '('; + printMCExpr(*BE.getLHS(), OS); + OS<< ')'; + } + + switch (BE.getOpcode()) { + case MCBinaryExpr::Add: + // Print "X-42" instead of "X+-42". + if (const MCConstantExpr *RHSC = dyn_cast(BE.getRHS())) { + if (RHSC->getValue() < 0) { + OS << RHSC->getValue(); + return; + } + } + + OS << '+'; + break; + default: llvm_unreachable("Unhandled binary operator"); + } + + // Only print parens around the LHS if it is non-trivial. + if (isa(BE.getRHS()) || isa(BE.getRHS())) { + printMCExpr(*BE.getRHS(), OS); + } else { + OS << '('; + printMCExpr(*BE.getRHS(), OS); + OS << ')'; + } + return; + } + } + + llvm_unreachable("Invalid expression kind!"); +} + /// PrintAsmOperand - Print out an operand for an inline asm expression. /// bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, @@ -2426,16 +2284,9 @@ void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum, } } - -// Force static initialization. -extern "C" void LLVMInitializeNVPTXBackendAsmPrinter() { - RegisterAsmPrinter X(TheNVPTXTarget32); - RegisterAsmPrinter Y(TheNVPTXTarget64); -} - 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 << ":"; @@ -2443,11 +2294,11 @@ 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) { + if (!reader) { reader = new LineReader(filename); }