diff options
Diffstat (limited to 'lib/Target/NVPTX/NVPTXAsmPrinter.cpp')
-rw-r--r-- | lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 277 |
1 files changed, 37 insertions, 240 deletions
diff --git a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index decf02a..35ba4f1 100644 --- a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -88,12 +88,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<const GlobalVariable *> Others; for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i) @@ -140,7 +137,8 @@ const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { // 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 (Constant *C = ConstantFoldConstantExpression( + CE, AP.TM.getSubtargetImpl()->getDataLayout())) if (C != CE) return LowerConstant(C, AP); @@ -169,7 +167,7 @@ const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { report_fatal_error(OS.str()); } case Instruction::GetElementPtr: { - const DataLayout &TD = *AP.TM.getDataLayout(); + const DataLayout &TD = *AP.TM.getSubtargetImpl()->getDataLayout(); // Generate a symbolic expression for the byte address APInt OffsetAI(TD.getPointerSizeInBits(), 0); cast<GEPOperator>(CE)->accumulateConstantOffset(TD, OffsetAI); @@ -193,7 +191,7 @@ const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { return LowerConstant(CE->getOperand(0), AP); case Instruction::IntToPtr: { - const DataLayout &TD = *AP.TM.getDataLayout(); + const DataLayout &TD = *AP.TM.getSubtargetImpl()->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); @@ -203,7 +201,7 @@ const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { } case Instruction::PtrToInt: { - const DataLayout &TD = *AP.TM.getDataLayout(); + const DataLayout &TD = *AP.TM.getSubtargetImpl()->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); @@ -330,253 +328,51 @@ void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) { 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) { @@ -704,8 +500,8 @@ MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) { } void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) { - const DataLayout *TD = TM.getDataLayout(); - const TargetLowering *TLI = TM.getTargetLowering(); + const DataLayout *TD = TM.getSubtargetImpl()->getDataLayout(); + const TargetLowering *TLI = TM.getSubtargetImpl()->getTargetLowering(); Type *Ty = F->getReturnType(); @@ -828,13 +624,14 @@ void NVPTXAsmPrinter::EmitFunctionBodyEnd() { void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const { unsigned RegNo = MI->getOperand(0).getReg(); - const TargetRegisterInfo *TRI = TM.getRegisterInfo(); + const TargetRegisterInfo *TRI = TM.getSubtargetImpl()->getRegisterInfo(); if (TRI->isVirtualRegister(RegNo)) { OutStreamer.AddComment(Twine("implicit-def: ") + getVirtualRegisterName(RegNo)); } else { - OutStreamer.AddComment(Twine("implicit-def: ") + - TM.getRegisterInfo()->getName(RegNo)); + OutStreamer.AddComment( + Twine("implicit-def: ") + + TM.getSubtargetImpl()->getRegisterInfo()->getName(RegNo)); } OutStreamer.AddBlankLine(); } @@ -1155,7 +952,7 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { const_cast<TargetLoweringObjectFile &>(getObjFileLowering()) .Initialize(OutContext, TM); - Mang = new Mangler(TM.getDataLayout()); + Mang = new Mangler(TM.getSubtargetImpl()->getDataLayout()); // Emit header before any dwarf directives are emitted below. emitHeader(M, OS1); @@ -1356,7 +1153,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, GVar->getName().startswith("nvvm.")) return; - const DataLayout *TD = TM.getDataLayout(); + const DataLayout *TD = TM.getSubtargetImpl()->getDataLayout(); // GlobalVariables are always constant pointers themselves. const PointerType *PTy = GVar->getType(); @@ -1659,7 +1456,7 @@ NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty, bool useB4PTR) const { void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar, raw_ostream &O) { - const DataLayout *TD = TM.getDataLayout(); + const DataLayout *TD = TM.getSubtargetImpl()->getDataLayout(); // GlobalVariables are always constant pointers themselves. const PointerType *PTy = GVar->getType(); @@ -1780,9 +1577,9 @@ void NVPTXAsmPrinter::printParamName(int paramIndex, raw_ostream &O) { } void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { - const DataLayout *TD = TM.getDataLayout(); + const DataLayout *TD = TM.getSubtargetImpl()->getDataLayout(); const AttributeSet &PAL = F->getAttributes(); - const TargetLowering *TLI = TM.getTargetLowering(); + const TargetLowering *TLI = TM.getSubtargetImpl()->getTargetLowering(); Function::const_arg_iterator I, E; unsigned paramIndex = 0; bool first = true; @@ -1973,7 +1770,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 @@ -2010,9 +1807,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 @@ -2113,7 +1910,7 @@ void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) { void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, AggBuffer *aggBuffer) { - const DataLayout *TD = TM.getDataLayout(); + const DataLayout *TD = TM.getSubtargetImpl()->getDataLayout(); if (isa<UndefValue>(CPV) || CPV->isNullValue()) { int s = TD->getTypeAllocSize(CPV->getType()); @@ -2237,7 +2034,7 @@ void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV, AggBuffer *aggBuffer) { - const DataLayout *TD = TM.getDataLayout(); + const DataLayout *TD = TM.getSubtargetImpl()->getDataLayout(); int Bytes; // Old constants |