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