1 //===-- NVPTXAsmPrinter.cpp - NVPTX LLVM assembly writer ------------------===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file contains a printer that converts from our internal representation
10 // of machine-dependent LLVM code to NVPTX assembly language.
11 //
12 //===----------------------------------------------------------------------===//
13 
14 #include "NVPTXAsmPrinter.h"
15 #include "MCTargetDesc/NVPTXBaseInfo.h"
16 #include "MCTargetDesc/NVPTXInstPrinter.h"
17 #include "MCTargetDesc/NVPTXMCAsmInfo.h"
18 #include "MCTargetDesc/NVPTXTargetStreamer.h"
19 #include "NVPTX.h"
20 #include "NVPTXMCExpr.h"
21 #include "NVPTXMachineFunctionInfo.h"
22 #include "NVPTXRegisterInfo.h"
23 #include "NVPTXSubtarget.h"
24 #include "NVPTXTargetMachine.h"
25 #include "NVPTXUtilities.h"
26 #include "TargetInfo/NVPTXTargetInfo.h"
27 #include "cl_common_defines.h"
28 #include "llvm/ADT/APFloat.h"
29 #include "llvm/ADT/APInt.h"
30 #include "llvm/ADT/DenseMap.h"
31 #include "llvm/ADT/DenseSet.h"
32 #include "llvm/ADT/SmallString.h"
33 #include "llvm/ADT/SmallVector.h"
34 #include "llvm/ADT/StringExtras.h"
35 #include "llvm/ADT/StringRef.h"
36 #include "llvm/ADT/Triple.h"
37 #include "llvm/ADT/Twine.h"
38 #include "llvm/Analysis/ConstantFolding.h"
39 #include "llvm/CodeGen/Analysis.h"
40 #include "llvm/CodeGen/MachineBasicBlock.h"
41 #include "llvm/CodeGen/MachineFrameInfo.h"
42 #include "llvm/CodeGen/MachineFunction.h"
43 #include "llvm/CodeGen/MachineInstr.h"
44 #include "llvm/CodeGen/MachineLoopInfo.h"
45 #include "llvm/CodeGen/MachineModuleInfo.h"
46 #include "llvm/CodeGen/MachineOperand.h"
47 #include "llvm/CodeGen/MachineRegisterInfo.h"
48 #include "llvm/CodeGen/TargetRegisterInfo.h"
49 #include "llvm/CodeGen/ValueTypes.h"
50 #include "llvm/IR/Attributes.h"
51 #include "llvm/IR/BasicBlock.h"
52 #include "llvm/IR/Constant.h"
53 #include "llvm/IR/Constants.h"
54 #include "llvm/IR/DataLayout.h"
55 #include "llvm/IR/DebugInfo.h"
56 #include "llvm/IR/DebugInfoMetadata.h"
57 #include "llvm/IR/DebugLoc.h"
58 #include "llvm/IR/DerivedTypes.h"
59 #include "llvm/IR/Function.h"
60 #include "llvm/IR/GlobalValue.h"
61 #include "llvm/IR/GlobalVariable.h"
62 #include "llvm/IR/Instruction.h"
63 #include "llvm/IR/LLVMContext.h"
64 #include "llvm/IR/Module.h"
65 #include "llvm/IR/Operator.h"
66 #include "llvm/IR/Type.h"
67 #include "llvm/IR/User.h"
68 #include "llvm/MC/MCExpr.h"
69 #include "llvm/MC/MCInst.h"
70 #include "llvm/MC/MCInstrDesc.h"
71 #include "llvm/MC/MCStreamer.h"
72 #include "llvm/MC/MCSymbol.h"
73 #include "llvm/MC/TargetRegistry.h"
74 #include "llvm/Support/Casting.h"
75 #include "llvm/Support/CommandLine.h"
76 #include "llvm/Support/ErrorHandling.h"
77 #include "llvm/Support/MachineValueType.h"
78 #include "llvm/Support/Path.h"
79 #include "llvm/Support/raw_ostream.h"
80 #include "llvm/Target/TargetLoweringObjectFile.h"
81 #include "llvm/Target/TargetMachine.h"
82 #include "llvm/Transforms/Utils/UnrollLoop.h"
83 #include <cassert>
84 #include <cstdint>
85 #include <cstring>
86 #include <new>
87 #include <string>
88 #include <utility>
89 #include <vector>
90 
91 using namespace llvm;
92 
93 #define DEPOTNAME "__local_depot"
94 
95 /// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
96 /// depends.
97 static void
98 DiscoverDependentGlobals(const Value *V,
99                          DenseSet<const GlobalVariable *> &Globals) {
100   if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(V))
101     Globals.insert(GV);
102   else {
103     if (const User *U = dyn_cast<User>(V)) {
104       for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) {
105         DiscoverDependentGlobals(U->getOperand(i), Globals);
106       }
107     }
108   }
109 }
110 
111 /// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable
112 /// instances to be emitted, but only after any dependents have been added
113 /// first.s
114 static void
115 VisitGlobalVariableForEmission(const GlobalVariable *GV,
116                                SmallVectorImpl<const GlobalVariable *> &Order,
117                                DenseSet<const GlobalVariable *> &Visited,
118                                DenseSet<const GlobalVariable *> &Visiting) {
119   // Have we already visited this one?
120   if (Visited.count(GV))
121     return;
122 
123   // Do we have a circular dependency?
124   if (!Visiting.insert(GV).second)
125     report_fatal_error("Circular dependency found in global variable set");
126 
127   // Make sure we visit all dependents first
128   DenseSet<const GlobalVariable *> Others;
129   for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i)
130     DiscoverDependentGlobals(GV->getOperand(i), Others);
131 
132   for (const GlobalVariable *GV : Others)
133     VisitGlobalVariableForEmission(GV, Order, Visited, Visiting);
134 
135   // Now we can visit ourself
136   Order.push_back(GV);
137   Visited.insert(GV);
138   Visiting.erase(GV);
139 }
140 
141 void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {
142   NVPTX_MC::verifyInstructionPredicates(MI->getOpcode(),
143                                         getSubtargetInfo().getFeatureBits());
144 
145   MCInst Inst;
146   lowerToMCInst(MI, Inst);
147   EmitToStreamer(*OutStreamer, Inst);
148 }
149 
150 // Handle symbol backtracking for targets that do not support image handles
151 bool NVPTXAsmPrinter::lowerImageHandleOperand(const MachineInstr *MI,
152                                            unsigned OpNo, MCOperand &MCOp) {
153   const MachineOperand &MO = MI->getOperand(OpNo);
154   const MCInstrDesc &MCID = MI->getDesc();
155 
156   if (MCID.TSFlags & NVPTXII::IsTexFlag) {
157     // This is a texture fetch, so operand 4 is a texref and operand 5 is
158     // a samplerref
159     if (OpNo == 4 && MO.isImm()) {
160       lowerImageHandleSymbol(MO.getImm(), MCOp);
161       return true;
162     }
163     if (OpNo == 5 && MO.isImm() && !(MCID.TSFlags & NVPTXII::IsTexModeUnifiedFlag)) {
164       lowerImageHandleSymbol(MO.getImm(), MCOp);
165       return true;
166     }
167 
168     return false;
169   } else if (MCID.TSFlags & NVPTXII::IsSuldMask) {
170     unsigned VecSize =
171       1 << (((MCID.TSFlags & NVPTXII::IsSuldMask) >> NVPTXII::IsSuldShift) - 1);
172 
173     // For a surface load of vector size N, the Nth operand will be the surfref
174     if (OpNo == VecSize && MO.isImm()) {
175       lowerImageHandleSymbol(MO.getImm(), MCOp);
176       return true;
177     }
178 
179     return false;
180   } else if (MCID.TSFlags & NVPTXII::IsSustFlag) {
181     // This is a surface store, so operand 0 is a surfref
182     if (OpNo == 0 && MO.isImm()) {
183       lowerImageHandleSymbol(MO.getImm(), MCOp);
184       return true;
185     }
186 
187     return false;
188   } else if (MCID.TSFlags & NVPTXII::IsSurfTexQueryFlag) {
189     // This is a query, so operand 1 is a surfref/texref
190     if (OpNo == 1 && MO.isImm()) {
191       lowerImageHandleSymbol(MO.getImm(), MCOp);
192       return true;
193     }
194 
195     return false;
196   }
197 
198   return false;
199 }
200 
201 void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) {
202   // Ewwww
203   LLVMTargetMachine &TM = const_cast<LLVMTargetMachine&>(MF->getTarget());
204   NVPTXTargetMachine &nvTM = static_cast<NVPTXTargetMachine&>(TM);
205   const NVPTXMachineFunctionInfo *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
206   const char *Sym = MFI->getImageHandleSymbol(Index);
207   std::string *SymNamePtr =
208     nvTM.getManagedStrPool()->getManagedString(Sym);
209   MCOp = GetSymbolRef(OutContext.getOrCreateSymbol(StringRef(*SymNamePtr)));
210 }
211 
212 void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
213   OutMI.setOpcode(MI->getOpcode());
214   // Special: Do not mangle symbol operand of CALL_PROTOTYPE
215   if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
216     const MachineOperand &MO = MI->getOperand(0);
217     OutMI.addOperand(GetSymbolRef(
218       OutContext.getOrCreateSymbol(Twine(MO.getSymbolName()))));
219     return;
220   }
221 
222   const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
223   for (unsigned i = 0, e = MI->getNumOperands(); i != e; ++i) {
224     const MachineOperand &MO = MI->getOperand(i);
225 
226     MCOperand MCOp;
227     if (!STI.hasImageHandles()) {
228       if (lowerImageHandleOperand(MI, i, MCOp)) {
229         OutMI.addOperand(MCOp);
230         continue;
231       }
232     }
233 
234     if (lowerOperand(MO, MCOp))
235       OutMI.addOperand(MCOp);
236   }
237 }
238 
239 bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO,
240                                    MCOperand &MCOp) {
241   switch (MO.getType()) {
242   default: llvm_unreachable("unknown operand type");
243   case MachineOperand::MO_Register:
244     MCOp = MCOperand::createReg(encodeVirtualRegister(MO.getReg()));
245     break;
246   case MachineOperand::MO_Immediate:
247     MCOp = MCOperand::createImm(MO.getImm());
248     break;
249   case MachineOperand::MO_MachineBasicBlock:
250     MCOp = MCOperand::createExpr(MCSymbolRefExpr::create(
251         MO.getMBB()->getSymbol(), OutContext));
252     break;
253   case MachineOperand::MO_ExternalSymbol:
254     MCOp = GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
255     break;
256   case MachineOperand::MO_GlobalAddress:
257     MCOp = GetSymbolRef(getSymbol(MO.getGlobal()));
258     break;
259   case MachineOperand::MO_FPImmediate: {
260     const ConstantFP *Cnt = MO.getFPImm();
261     const APFloat &Val = Cnt->getValueAPF();
262 
263     switch (Cnt->getType()->getTypeID()) {
264     default: report_fatal_error("Unsupported FP type"); break;
265     case Type::HalfTyID:
266       MCOp = MCOperand::createExpr(
267         NVPTXFloatMCExpr::createConstantFPHalf(Val, OutContext));
268       break;
269     case Type::FloatTyID:
270       MCOp = MCOperand::createExpr(
271         NVPTXFloatMCExpr::createConstantFPSingle(Val, OutContext));
272       break;
273     case Type::DoubleTyID:
274       MCOp = MCOperand::createExpr(
275         NVPTXFloatMCExpr::createConstantFPDouble(Val, OutContext));
276       break;
277     }
278     break;
279   }
280   }
281   return true;
282 }
283 
284 unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
285   if (Register::isVirtualRegister(Reg)) {
286     const TargetRegisterClass *RC = MRI->getRegClass(Reg);
287 
288     DenseMap<unsigned, unsigned> &RegMap = VRegMapping[RC];
289     unsigned RegNum = RegMap[Reg];
290 
291     // Encode the register class in the upper 4 bits
292     // Must be kept in sync with NVPTXInstPrinter::printRegName
293     unsigned Ret = 0;
294     if (RC == &NVPTX::Int1RegsRegClass) {
295       Ret = (1 << 28);
296     } else if (RC == &NVPTX::Int16RegsRegClass) {
297       Ret = (2 << 28);
298     } else if (RC == &NVPTX::Int32RegsRegClass) {
299       Ret = (3 << 28);
300     } else if (RC == &NVPTX::Int64RegsRegClass) {
301       Ret = (4 << 28);
302     } else if (RC == &NVPTX::Float32RegsRegClass) {
303       Ret = (5 << 28);
304     } else if (RC == &NVPTX::Float64RegsRegClass) {
305       Ret = (6 << 28);
306     } else if (RC == &NVPTX::Float16RegsRegClass) {
307       Ret = (7 << 28);
308     } else if (RC == &NVPTX::Float16x2RegsRegClass) {
309       Ret = (8 << 28);
310     } else {
311       report_fatal_error("Bad register class");
312     }
313 
314     // Insert the vreg number
315     Ret |= (RegNum & 0x0FFFFFFF);
316     return Ret;
317   } else {
318     // Some special-use registers are actually physical registers.
319     // Encode this as the register class ID of 0 and the real register ID.
320     return Reg & 0x0FFFFFFF;
321   }
322 }
323 
324 MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
325   const MCExpr *Expr;
326   Expr = MCSymbolRefExpr::create(Symbol, MCSymbolRefExpr::VK_None,
327                                  OutContext);
328   return MCOperand::createExpr(Expr);
329 }
330 
331 void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
332   const DataLayout &DL = getDataLayout();
333   const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
334   const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
335 
336   Type *Ty = F->getReturnType();
337 
338   bool isABI = (STI.getSmVersion() >= 20);
339 
340   if (Ty->getTypeID() == Type::VoidTyID)
341     return;
342 
343   O << " (";
344 
345   if (isABI) {
346     if (Ty->isFloatingPointTy() || (Ty->isIntegerTy() && !Ty->isIntegerTy(128))) {
347       unsigned size = 0;
348       if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
349         size = ITy->getBitWidth();
350       } else {
351         assert(Ty->isFloatingPointTy() && "Floating point type expected here");
352         size = Ty->getPrimitiveSizeInBits();
353       }
354       // PTX ABI requires all scalar return values to be at least 32
355       // bits in size.  fp16 normally uses .b16 as its storage type in
356       // PTX, so its size must be adjusted here, too.
357       if (size < 32)
358         size = 32;
359 
360       O << ".param .b" << size << " func_retval0";
361     } else if (isa<PointerType>(Ty)) {
362       O << ".param .b" << TLI->getPointerTy(DL).getSizeInBits()
363         << " func_retval0";
364     } else if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
365       unsigned totalsz = DL.getTypeAllocSize(Ty);
366       unsigned retAlignment = 0;
367       if (!getAlign(*F, 0, retAlignment))
368         retAlignment = TLI->getFunctionParamOptimizedAlign(F, Ty, DL).value();
369       O << ".param .align " << retAlignment << " .b8 func_retval0[" << totalsz
370         << "]";
371     } else
372       llvm_unreachable("Unknown return type");
373   } else {
374     SmallVector<EVT, 16> vtparts;
375     ComputeValueVTs(*TLI, DL, Ty, vtparts);
376     unsigned idx = 0;
377     for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
378       unsigned elems = 1;
379       EVT elemtype = vtparts[i];
380       if (vtparts[i].isVector()) {
381         elems = vtparts[i].getVectorNumElements();
382         elemtype = vtparts[i].getVectorElementType();
383       }
384 
385       for (unsigned j = 0, je = elems; j != je; ++j) {
386         unsigned sz = elemtype.getSizeInBits();
387         if (elemtype.isInteger() && (sz < 32))
388           sz = 32;
389         O << ".reg .b" << sz << " func_retval" << idx;
390         if (j < je - 1)
391           O << ", ";
392         ++idx;
393       }
394       if (i < e - 1)
395         O << ", ";
396     }
397   }
398   O << ") ";
399 }
400 
401 void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
402                                         raw_ostream &O) {
403   const Function &F = MF.getFunction();
404   printReturnValStr(&F, O);
405 }
406 
407 // Return true if MBB is the header of a loop marked with
408 // llvm.loop.unroll.disable.
409 // TODO: consider "#pragma unroll 1" which is equivalent to "#pragma nounroll".
410 bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
411     const MachineBasicBlock &MBB) const {
412   MachineLoopInfo &LI = getAnalysis<MachineLoopInfo>();
413   // We insert .pragma "nounroll" only to the loop header.
414   if (!LI.isLoopHeader(&MBB))
415     return false;
416 
417   // llvm.loop.unroll.disable is marked on the back edges of a loop. Therefore,
418   // we iterate through each back edge of the loop with header MBB, and check
419   // whether its metadata contains llvm.loop.unroll.disable.
420   for (const MachineBasicBlock *PMBB : MBB.predecessors()) {
421     if (LI.getLoopFor(PMBB) != LI.getLoopFor(&MBB)) {
422       // Edges from other loops to MBB are not back edges.
423       continue;
424     }
425     if (const BasicBlock *PBB = PMBB->getBasicBlock()) {
426       if (MDNode *LoopID =
427               PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
428         if (GetUnrollMetadata(LoopID, "llvm.loop.unroll.disable"))
429           return true;
430       }
431     }
432   }
433   return false;
434 }
435 
436 void NVPTXAsmPrinter::emitBasicBlockStart(const MachineBasicBlock &MBB) {
437   AsmPrinter::emitBasicBlockStart(MBB);
438   if (isLoopHeaderOfNoUnroll(MBB))
439     OutStreamer->emitRawText(StringRef("\t.pragma \"nounroll\";\n"));
440 }
441 
442 void NVPTXAsmPrinter::emitFunctionEntryLabel() {
443   SmallString<128> Str;
444   raw_svector_ostream O(Str);
445 
446   if (!GlobalsEmitted) {
447     emitGlobals(*MF->getFunction().getParent());
448     GlobalsEmitted = true;
449   }
450 
451   // Set up
452   MRI = &MF->getRegInfo();
453   F = &MF->getFunction();
454   emitLinkageDirective(F, O);
455   if (isKernelFunction(*F))
456     O << ".entry ";
457   else {
458     O << ".func ";
459     printReturnValStr(*MF, O);
460   }
461 
462   CurrentFnSym->print(O, MAI);
463 
464   emitFunctionParamList(*MF, O);
465 
466   if (isKernelFunction(*F))
467     emitKernelFunctionDirectives(*F, O);
468 
469   OutStreamer->emitRawText(O.str());
470 
471   VRegMapping.clear();
472   // Emit open brace for function body.
473   OutStreamer->emitRawText(StringRef("{\n"));
474   setAndEmitFunctionVirtualRegisters(*MF);
475   // Emit initial .loc debug directive for correct relocation symbol data.
476   if (MMI && MMI->hasDebugInfo())
477     emitInitialRawDwarfLocDirective(*MF);
478 }
479 
480 bool NVPTXAsmPrinter::runOnMachineFunction(MachineFunction &F) {
481   bool Result = AsmPrinter::runOnMachineFunction(F);
482   // Emit closing brace for the body of function F.
483   // The closing brace must be emitted here because we need to emit additional
484   // debug labels/data after the last basic block.
485   // We need to emit the closing brace here because we don't have function that
486   // finished emission of the function body.
487   OutStreamer->emitRawText(StringRef("}\n"));
488   return Result;
489 }
490 
491 void NVPTXAsmPrinter::emitFunctionBodyStart() {
492   SmallString<128> Str;
493   raw_svector_ostream O(Str);
494   emitDemotedVars(&MF->getFunction(), O);
495   OutStreamer->emitRawText(O.str());
496 }
497 
498 void NVPTXAsmPrinter::emitFunctionBodyEnd() {
499   VRegMapping.clear();
500 }
501 
502 const MCSymbol *NVPTXAsmPrinter::getFunctionFrameSymbol() const {
503     SmallString<128> Str;
504     raw_svector_ostream(Str) << DEPOTNAME << getFunctionNumber();
505     return OutContext.getOrCreateSymbol(Str);
506 }
507 
508 void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const {
509   Register RegNo = MI->getOperand(0).getReg();
510   if (Register::isVirtualRegister(RegNo)) {
511     OutStreamer->AddComment(Twine("implicit-def: ") +
512                             getVirtualRegisterName(RegNo));
513   } else {
514     const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
515     OutStreamer->AddComment(Twine("implicit-def: ") +
516                             STI.getRegisterInfo()->getName(RegNo));
517   }
518   OutStreamer->addBlankLine();
519 }
520 
521 void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F,
522                                                    raw_ostream &O) const {
523   // If the NVVM IR has some of reqntid* specified, then output
524   // the reqntid directive, and set the unspecified ones to 1.
525   // If none of reqntid* is specified, don't output reqntid directive.
526   unsigned reqntidx, reqntidy, reqntidz;
527   bool specified = false;
528   if (!getReqNTIDx(F, reqntidx))
529     reqntidx = 1;
530   else
531     specified = true;
532   if (!getReqNTIDy(F, reqntidy))
533     reqntidy = 1;
534   else
535     specified = true;
536   if (!getReqNTIDz(F, reqntidz))
537     reqntidz = 1;
538   else
539     specified = true;
540 
541   if (specified)
542     O << ".reqntid " << reqntidx << ", " << reqntidy << ", " << reqntidz
543       << "\n";
544 
545   // If the NVVM IR has some of maxntid* specified, then output
546   // the maxntid directive, and set the unspecified ones to 1.
547   // If none of maxntid* is specified, don't output maxntid directive.
548   unsigned maxntidx, maxntidy, maxntidz;
549   specified = false;
550   if (!getMaxNTIDx(F, maxntidx))
551     maxntidx = 1;
552   else
553     specified = true;
554   if (!getMaxNTIDy(F, maxntidy))
555     maxntidy = 1;
556   else
557     specified = true;
558   if (!getMaxNTIDz(F, maxntidz))
559     maxntidz = 1;
560   else
561     specified = true;
562 
563   if (specified)
564     O << ".maxntid " << maxntidx << ", " << maxntidy << ", " << maxntidz
565       << "\n";
566 
567   unsigned mincta;
568   if (getMinCTASm(F, mincta))
569     O << ".minnctapersm " << mincta << "\n";
570 
571   unsigned maxnreg;
572   if (getMaxNReg(F, maxnreg))
573     O << ".maxnreg " << maxnreg << "\n";
574 }
575 
576 std::string
577 NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const {
578   const TargetRegisterClass *RC = MRI->getRegClass(Reg);
579 
580   std::string Name;
581   raw_string_ostream NameStr(Name);
582 
583   VRegRCMap::const_iterator I = VRegMapping.find(RC);
584   assert(I != VRegMapping.end() && "Bad register class");
585   const DenseMap<unsigned, unsigned> &RegMap = I->second;
586 
587   VRegMap::const_iterator VI = RegMap.find(Reg);
588   assert(VI != RegMap.end() && "Bad virtual register");
589   unsigned MappedVR = VI->second;
590 
591   NameStr << getNVPTXRegClassStr(RC) << MappedVR;
592 
593   NameStr.flush();
594   return Name;
595 }
596 
597 void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr,
598                                           raw_ostream &O) {
599   O << getVirtualRegisterName(vr);
600 }
601 
602 void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) {
603   emitLinkageDirective(F, O);
604   if (isKernelFunction(*F))
605     O << ".entry ";
606   else
607     O << ".func ";
608   printReturnValStr(F, O);
609   getSymbol(F)->print(O, MAI);
610   O << "\n";
611   emitFunctionParamList(F, O);
612   O << ";\n";
613 }
614 
615 static bool usedInGlobalVarDef(const Constant *C) {
616   if (!C)
617     return false;
618 
619   if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(C)) {
620     return GV->getName() != "llvm.used";
621   }
622 
623   for (const User *U : C->users())
624     if (const Constant *C = dyn_cast<Constant>(U))
625       if (usedInGlobalVarDef(C))
626         return true;
627 
628   return false;
629 }
630 
631 static bool usedInOneFunc(const User *U, Function const *&oneFunc) {
632   if (const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
633     if (othergv->getName() == "llvm.used")
634       return true;
635   }
636 
637   if (const Instruction *instr = dyn_cast<Instruction>(U)) {
638     if (instr->getParent() && instr->getParent()->getParent()) {
639       const Function *curFunc = instr->getParent()->getParent();
640       if (oneFunc && (curFunc != oneFunc))
641         return false;
642       oneFunc = curFunc;
643       return true;
644     } else
645       return false;
646   }
647 
648   for (const User *UU : U->users())
649     if (!usedInOneFunc(UU, oneFunc))
650       return false;
651 
652   return true;
653 }
654 
655 /* Find out if a global variable can be demoted to local scope.
656  * Currently, this is valid for CUDA shared variables, which have local
657  * scope and global lifetime. So the conditions to check are :
658  * 1. Is the global variable in shared address space?
659  * 2. Does it have internal linkage?
660  * 3. Is the global variable referenced only in one function?
661  */
662 static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) {
663   if (!gv->hasInternalLinkage())
664     return false;
665   PointerType *Pty = gv->getType();
666   if (Pty->getAddressSpace() != ADDRESS_SPACE_SHARED)
667     return false;
668 
669   const Function *oneFunc = nullptr;
670 
671   bool flag = usedInOneFunc(gv, oneFunc);
672   if (!flag)
673     return false;
674   if (!oneFunc)
675     return false;
676   f = oneFunc;
677   return true;
678 }
679 
680 static bool useFuncSeen(const Constant *C,
681                         DenseMap<const Function *, bool> &seenMap) {
682   for (const User *U : C->users()) {
683     if (const Constant *cu = dyn_cast<Constant>(U)) {
684       if (useFuncSeen(cu, seenMap))
685         return true;
686     } else if (const Instruction *I = dyn_cast<Instruction>(U)) {
687       const BasicBlock *bb = I->getParent();
688       if (!bb)
689         continue;
690       const Function *caller = bb->getParent();
691       if (!caller)
692         continue;
693       if (seenMap.find(caller) != seenMap.end())
694         return true;
695     }
696   }
697   return false;
698 }
699 
700 void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {
701   DenseMap<const Function *, bool> seenMap;
702   for (const Function &F : M) {
703     if (F.getAttributes().hasFnAttr("nvptx-libcall-callee")) {
704       emitDeclaration(&F, O);
705       continue;
706     }
707 
708     if (F.isDeclaration()) {
709       if (F.use_empty())
710         continue;
711       if (F.getIntrinsicID())
712         continue;
713       emitDeclaration(&F, O);
714       continue;
715     }
716     for (const User *U : F.users()) {
717       if (const Constant *C = dyn_cast<Constant>(U)) {
718         if (usedInGlobalVarDef(C)) {
719           // The use is in the initialization of a global variable
720           // that is a function pointer, so print a declaration
721           // for the original function
722           emitDeclaration(&F, O);
723           break;
724         }
725         // Emit a declaration of this function if the function that
726         // uses this constant expr has already been seen.
727         if (useFuncSeen(C, seenMap)) {
728           emitDeclaration(&F, O);
729           break;
730         }
731       }
732 
733       if (!isa<Instruction>(U))
734         continue;
735       const Instruction *instr = cast<Instruction>(U);
736       const BasicBlock *bb = instr->getParent();
737       if (!bb)
738         continue;
739       const Function *caller = bb->getParent();
740       if (!caller)
741         continue;
742 
743       // If a caller has already been seen, then the caller is
744       // appearing in the module before the callee. so print out
745       // a declaration for the callee.
746       if (seenMap.find(caller) != seenMap.end()) {
747         emitDeclaration(&F, O);
748         break;
749       }
750     }
751     seenMap[&F] = true;
752   }
753 }
754 
755 static bool isEmptyXXStructor(GlobalVariable *GV) {
756   if (!GV) return true;
757   const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());
758   if (!InitList) return true;  // Not an array; we don't know how to parse.
759   return InitList->getNumOperands() == 0;
760 }
761 
762 void NVPTXAsmPrinter::emitStartOfAsmFile(Module &M) {
763   // Construct a default subtarget off of the TargetMachine defaults. The
764   // rest of NVPTX isn't friendly to change subtargets per function and
765   // so the default TargetMachine will have all of the options.
766   const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
767   const auto* STI = static_cast<const NVPTXSubtarget*>(NTM.getSubtargetImpl());
768   SmallString<128> Str1;
769   raw_svector_ostream OS1(Str1);
770 
771   // Emit header before any dwarf directives are emitted below.
772   emitHeader(M, OS1, *STI);
773   OutStreamer->emitRawText(OS1.str());
774 }
775 
776 bool NVPTXAsmPrinter::doInitialization(Module &M) {
777   if (M.alias_size()) {
778     report_fatal_error("Module has aliases, which NVPTX does not support.");
779     return true; // error
780   }
781   if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_ctors"))) {
782     report_fatal_error(
783         "Module has a nontrivial global ctor, which NVPTX does not support.");
784     return true;  // error
785   }
786   if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_dtors"))) {
787     report_fatal_error(
788         "Module has a nontrivial global dtor, which NVPTX does not support.");
789     return true;  // error
790   }
791 
792   // We need to call the parent's one explicitly.
793   bool Result = AsmPrinter::doInitialization(M);
794 
795   GlobalsEmitted = false;
796 
797   return Result;
798 }
799 
800 void NVPTXAsmPrinter::emitGlobals(const Module &M) {
801   SmallString<128> Str2;
802   raw_svector_ostream OS2(Str2);
803 
804   emitDeclarations(M, OS2);
805 
806   // As ptxas does not support forward references of globals, we need to first
807   // sort the list of module-level globals in def-use order. We visit each
808   // global variable in order, and ensure that we emit it *after* its dependent
809   // globals. We use a little extra memory maintaining both a set and a list to
810   // have fast searches while maintaining a strict ordering.
811   SmallVector<const GlobalVariable *, 8> Globals;
812   DenseSet<const GlobalVariable *> GVVisited;
813   DenseSet<const GlobalVariable *> GVVisiting;
814 
815   // Visit each global variable, in order
816   for (const GlobalVariable &I : M.globals())
817     VisitGlobalVariableForEmission(&I, Globals, GVVisited, GVVisiting);
818 
819   assert(GVVisited.size() == M.getGlobalList().size() &&
820          "Missed a global variable");
821   assert(GVVisiting.size() == 0 && "Did not fully process a global variable");
822 
823   const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
824   const NVPTXSubtarget &STI =
825       *static_cast<const NVPTXSubtarget *>(NTM.getSubtargetImpl());
826 
827   // Print out module-level global variables in proper order
828   for (unsigned i = 0, e = Globals.size(); i != e; ++i)
829     printModuleLevelGV(Globals[i], OS2, /*processDemoted=*/false, STI);
830 
831   OS2 << '\n';
832 
833   OutStreamer->emitRawText(OS2.str());
834 }
835 
836 void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
837                                  const NVPTXSubtarget &STI) {
838   O << "//\n";
839   O << "// Generated by LLVM NVPTX Back-End\n";
840   O << "//\n";
841   O << "\n";
842 
843   unsigned PTXVersion = STI.getPTXVersion();
844   O << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n";
845 
846   O << ".target ";
847   O << STI.getTargetName();
848 
849   const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
850   if (NTM.getDrvInterface() == NVPTX::NVCL)
851     O << ", texmode_independent";
852 
853   bool HasFullDebugInfo = false;
854   for (DICompileUnit *CU : M.debug_compile_units()) {
855     switch(CU->getEmissionKind()) {
856     case DICompileUnit::NoDebug:
857     case DICompileUnit::DebugDirectivesOnly:
858       break;
859     case DICompileUnit::LineTablesOnly:
860     case DICompileUnit::FullDebug:
861       HasFullDebugInfo = true;
862       break;
863     }
864     if (HasFullDebugInfo)
865       break;
866   }
867   if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo)
868     O << ", debug";
869 
870   O << "\n";
871 
872   O << ".address_size ";
873   if (NTM.is64Bit())
874     O << "64";
875   else
876     O << "32";
877   O << "\n";
878 
879   O << "\n";
880 }
881 
882 bool NVPTXAsmPrinter::doFinalization(Module &M) {
883   bool HasDebugInfo = MMI && MMI->hasDebugInfo();
884 
885   // If we did not emit any functions, then the global declarations have not
886   // yet been emitted.
887   if (!GlobalsEmitted) {
888     emitGlobals(M);
889     GlobalsEmitted = true;
890   }
891 
892   // call doFinalization
893   bool ret = AsmPrinter::doFinalization(M);
894 
895   clearAnnotationCache(&M);
896 
897   if (auto *TS = static_cast<NVPTXTargetStreamer *>(
898           OutStreamer->getTargetStreamer())) {
899     // Close the last emitted section
900     if (HasDebugInfo) {
901       TS->closeLastSection();
902       // Emit empty .debug_loc section for better support of the empty files.
903       OutStreamer->emitRawText("\t.section\t.debug_loc\t{\t}");
904     }
905 
906     // Output last DWARF .file directives, if any.
907     TS->outputDwarfFileDirectives();
908   }
909 
910   return ret;
911 
912   //bool Result = AsmPrinter::doFinalization(M);
913   // Instead of calling the parents doFinalization, we may
914   // clone parents doFinalization and customize here.
915   // Currently, we if NVISA out the EmitGlobals() in
916   // parent's doFinalization, which is too intrusive.
917   //
918   // Same for the doInitialization.
919   //return Result;
920 }
921 
922 // This function emits appropriate linkage directives for
923 // functions and global variables.
924 //
925 // extern function declaration            -> .extern
926 // extern function definition             -> .visible
927 // external global variable with init     -> .visible
928 // external without init                  -> .extern
929 // appending                              -> not allowed, assert.
930 // for any linkage other than
931 // internal, private, linker_private,
932 // linker_private_weak, linker_private_weak_def_auto,
933 // we emit                                -> .weak.
934 
935 void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V,
936                                            raw_ostream &O) {
937   if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA) {
938     if (V->hasExternalLinkage()) {
939       if (isa<GlobalVariable>(V)) {
940         const GlobalVariable *GVar = cast<GlobalVariable>(V);
941         if (GVar) {
942           if (GVar->hasInitializer())
943             O << ".visible ";
944           else
945             O << ".extern ";
946         }
947       } else if (V->isDeclaration())
948         O << ".extern ";
949       else
950         O << ".visible ";
951     } else if (V->hasAppendingLinkage()) {
952       std::string msg;
953       msg.append("Error: ");
954       msg.append("Symbol ");
955       if (V->hasName())
956         msg.append(std::string(V->getName()));
957       msg.append("has unsupported appending linkage type");
958       llvm_unreachable(msg.c_str());
959     } else if (!V->hasInternalLinkage() &&
960                !V->hasPrivateLinkage()) {
961       O << ".weak ";
962     }
963   }
964 }
965 
966 void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
967                                          raw_ostream &O, bool processDemoted,
968                                          const NVPTXSubtarget &STI) {
969   // Skip meta data
970   if (GVar->hasSection()) {
971     if (GVar->getSection() == "llvm.metadata")
972       return;
973   }
974 
975   // Skip LLVM intrinsic global variables
976   if (GVar->getName().startswith("llvm.") ||
977       GVar->getName().startswith("nvvm."))
978     return;
979 
980   const DataLayout &DL = getDataLayout();
981 
982   // GlobalVariables are always constant pointers themselves.
983   PointerType *PTy = GVar->getType();
984   Type *ETy = GVar->getValueType();
985 
986   if (GVar->hasExternalLinkage()) {
987     if (GVar->hasInitializer())
988       O << ".visible ";
989     else
990       O << ".extern ";
991   } else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() ||
992              GVar->hasAvailableExternallyLinkage() ||
993              GVar->hasCommonLinkage()) {
994     O << ".weak ";
995   }
996 
997   if (isTexture(*GVar)) {
998     O << ".global .texref " << getTextureName(*GVar) << ";\n";
999     return;
1000   }
1001 
1002   if (isSurface(*GVar)) {
1003     O << ".global .surfref " << getSurfaceName(*GVar) << ";\n";
1004     return;
1005   }
1006 
1007   if (GVar->isDeclaration()) {
1008     // (extern) declarations, no definition or initializer
1009     // Currently the only known declaration is for an automatic __local
1010     // (.shared) promoted to global.
1011     emitPTXGlobalVariable(GVar, O, STI);
1012     O << ";\n";
1013     return;
1014   }
1015 
1016   if (isSampler(*GVar)) {
1017     O << ".global .samplerref " << getSamplerName(*GVar);
1018 
1019     const Constant *Initializer = nullptr;
1020     if (GVar->hasInitializer())
1021       Initializer = GVar->getInitializer();
1022     const ConstantInt *CI = nullptr;
1023     if (Initializer)
1024       CI = dyn_cast<ConstantInt>(Initializer);
1025     if (CI) {
1026       unsigned sample = CI->getZExtValue();
1027 
1028       O << " = { ";
1029 
1030       for (int i = 0,
1031                addr = ((sample & __CLK_ADDRESS_MASK) >> __CLK_ADDRESS_BASE);
1032            i < 3; i++) {
1033         O << "addr_mode_" << i << " = ";
1034         switch (addr) {
1035         case 0:
1036           O << "wrap";
1037           break;
1038         case 1:
1039           O << "clamp_to_border";
1040           break;
1041         case 2:
1042           O << "clamp_to_edge";
1043           break;
1044         case 3:
1045           O << "wrap";
1046           break;
1047         case 4:
1048           O << "mirror";
1049           break;
1050         }
1051         O << ", ";
1052       }
1053       O << "filter_mode = ";
1054       switch ((sample & __CLK_FILTER_MASK) >> __CLK_FILTER_BASE) {
1055       case 0:
1056         O << "nearest";
1057         break;
1058       case 1:
1059         O << "linear";
1060         break;
1061       case 2:
1062         llvm_unreachable("Anisotropic filtering is not supported");
1063       default:
1064         O << "nearest";
1065         break;
1066       }
1067       if (!((sample & __CLK_NORMALIZED_MASK) >> __CLK_NORMALIZED_BASE)) {
1068         O << ", force_unnormalized_coords = 1";
1069       }
1070       O << " }";
1071     }
1072 
1073     O << ";\n";
1074     return;
1075   }
1076 
1077   if (GVar->hasPrivateLinkage()) {
1078     if (strncmp(GVar->getName().data(), "unrollpragma", 12) == 0)
1079       return;
1080 
1081     // FIXME - need better way (e.g. Metadata) to avoid generating this global
1082     if (strncmp(GVar->getName().data(), "filename", 8) == 0)
1083       return;
1084     if (GVar->use_empty())
1085       return;
1086   }
1087 
1088   const Function *demotedFunc = nullptr;
1089   if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) {
1090     O << "// " << GVar->getName() << " has been demoted\n";
1091     if (localDecls.find(demotedFunc) != localDecls.end())
1092       localDecls[demotedFunc].push_back(GVar);
1093     else {
1094       std::vector<const GlobalVariable *> temp;
1095       temp.push_back(GVar);
1096       localDecls[demotedFunc] = temp;
1097     }
1098     return;
1099   }
1100 
1101   O << ".";
1102   emitPTXAddressSpace(PTy->getAddressSpace(), O);
1103 
1104   if (isManaged(*GVar)) {
1105     if (STI.getPTXVersion() < 40 || STI.getSmVersion() < 30) {
1106       report_fatal_error(
1107           ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1108     }
1109     O << " .attribute(.managed)";
1110   }
1111 
1112   if (MaybeAlign A = GVar->getAlign())
1113     O << " .align " << A->value();
1114   else
1115     O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
1116 
1117   if (ETy->isFloatingPointTy() || ETy->isPointerTy() ||
1118       (ETy->isIntegerTy() && ETy->getScalarSizeInBits() <= 64)) {
1119     O << " .";
1120     // Special case: ABI requires that we use .u8 for predicates
1121     if (ETy->isIntegerTy(1))
1122       O << "u8";
1123     else
1124       O << getPTXFundamentalTypeStr(ETy, false);
1125     O << " ";
1126     getSymbol(GVar)->print(O, MAI);
1127 
1128     // Ptx allows variable initilization only for constant and global state
1129     // spaces.
1130     if (GVar->hasInitializer()) {
1131       if ((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1132           (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) {
1133         const Constant *Initializer = GVar->getInitializer();
1134         // 'undef' is treated as there is no value specified.
1135         if (!Initializer->isNullValue() && !isa<UndefValue>(Initializer)) {
1136           O << " = ";
1137           printScalarConstant(Initializer, O);
1138         }
1139       } else {
1140         // The frontend adds zero-initializer to device and constant variables
1141         // that don't have an initial value, and UndefValue to shared
1142         // variables, so skip warning for this case.
1143         if (!GVar->getInitializer()->isNullValue() &&
1144             !isa<UndefValue>(GVar->getInitializer())) {
1145           report_fatal_error("initial value of '" + GVar->getName() +
1146                              "' is not allowed in addrspace(" +
1147                              Twine(PTy->getAddressSpace()) + ")");
1148         }
1149       }
1150     }
1151   } else {
1152     unsigned int ElementSize = 0;
1153 
1154     // Although PTX has direct support for struct type and array type and
1155     // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
1156     // targets that support these high level field accesses. Structs, arrays
1157     // and vectors are lowered into arrays of bytes.
1158     switch (ETy->getTypeID()) {
1159     case Type::IntegerTyID: // Integers larger than 64 bits
1160     case Type::StructTyID:
1161     case Type::ArrayTyID:
1162     case Type::FixedVectorTyID:
1163       ElementSize = DL.getTypeStoreSize(ETy);
1164       // Ptx allows variable initilization only for constant and
1165       // global state spaces.
1166       if (((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1167            (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) &&
1168           GVar->hasInitializer()) {
1169         const Constant *Initializer = GVar->getInitializer();
1170         if (!isa<UndefValue>(Initializer) && !Initializer->isNullValue()) {
1171           AggBuffer aggBuffer(ElementSize, O, *this);
1172           bufferAggregateConstant(Initializer, &aggBuffer);
1173           if (aggBuffer.numSymbols) {
1174             if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit()) {
1175               O << " .u64 ";
1176               getSymbol(GVar)->print(O, MAI);
1177               O << "[";
1178               O << ElementSize / 8;
1179             } else {
1180               O << " .u32 ";
1181               getSymbol(GVar)->print(O, MAI);
1182               O << "[";
1183               O << ElementSize / 4;
1184             }
1185             O << "]";
1186           } else {
1187             O << " .b8 ";
1188             getSymbol(GVar)->print(O, MAI);
1189             O << "[";
1190             O << ElementSize;
1191             O << "]";
1192           }
1193           O << " = {";
1194           aggBuffer.print();
1195           O << "}";
1196         } else {
1197           O << " .b8 ";
1198           getSymbol(GVar)->print(O, MAI);
1199           if (ElementSize) {
1200             O << "[";
1201             O << ElementSize;
1202             O << "]";
1203           }
1204         }
1205       } else {
1206         O << " .b8 ";
1207         getSymbol(GVar)->print(O, MAI);
1208         if (ElementSize) {
1209           O << "[";
1210           O << ElementSize;
1211           O << "]";
1212         }
1213       }
1214       break;
1215     default:
1216       llvm_unreachable("type not supported yet");
1217     }
1218   }
1219   O << ";\n";
1220 }
1221 
1222 void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) {
1223   if (localDecls.find(f) == localDecls.end())
1224     return;
1225 
1226   std::vector<const GlobalVariable *> &gvars = localDecls[f];
1227 
1228   const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
1229   const NVPTXSubtarget &STI =
1230       *static_cast<const NVPTXSubtarget *>(NTM.getSubtargetImpl());
1231 
1232   for (const GlobalVariable *GV : gvars) {
1233     O << "\t// demoted variable\n\t";
1234     printModuleLevelGV(GV, O, /*processDemoted=*/true, STI);
1235   }
1236 }
1237 
1238 void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
1239                                           raw_ostream &O) const {
1240   switch (AddressSpace) {
1241   case ADDRESS_SPACE_LOCAL:
1242     O << "local";
1243     break;
1244   case ADDRESS_SPACE_GLOBAL:
1245     O << "global";
1246     break;
1247   case ADDRESS_SPACE_CONST:
1248     O << "const";
1249     break;
1250   case ADDRESS_SPACE_SHARED:
1251     O << "shared";
1252     break;
1253   default:
1254     report_fatal_error("Bad address space found while emitting PTX: " +
1255                        llvm::Twine(AddressSpace));
1256     break;
1257   }
1258 }
1259 
1260 std::string
1261 NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {
1262   switch (Ty->getTypeID()) {
1263   case Type::IntegerTyID: {
1264     unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
1265     if (NumBits == 1)
1266       return "pred";
1267     else if (NumBits <= 64) {
1268       std::string name = "u";
1269       return name + utostr(NumBits);
1270     } else {
1271       llvm_unreachable("Integer too large");
1272       break;
1273     }
1274     break;
1275   }
1276   case Type::HalfTyID:
1277     // fp16 is stored as .b16 for compatibility with pre-sm_53 PTX assembly.
1278     return "b16";
1279   case Type::FloatTyID:
1280     return "f32";
1281   case Type::DoubleTyID:
1282     return "f64";
1283   case Type::PointerTyID:
1284     if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit())
1285       if (useB4PTR)
1286         return "b64";
1287       else
1288         return "u64";
1289     else if (useB4PTR)
1290       return "b32";
1291     else
1292       return "u32";
1293   default:
1294     break;
1295   }
1296   llvm_unreachable("unexpected type");
1297 }
1298 
1299 void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,
1300                                             raw_ostream &O,
1301                                             const NVPTXSubtarget &STI) {
1302   const DataLayout &DL = getDataLayout();
1303 
1304   // GlobalVariables are always constant pointers themselves.
1305   Type *ETy = GVar->getValueType();
1306 
1307   O << ".";
1308   emitPTXAddressSpace(GVar->getType()->getAddressSpace(), O);
1309   if (isManaged(*GVar)) {
1310     if (STI.getPTXVersion() < 40 || STI.getSmVersion() < 30) {
1311       report_fatal_error(
1312           ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1313     }
1314     O << " .attribute(.managed)";
1315   }
1316   if (MaybeAlign A = GVar->getAlign())
1317     O << " .align " << A->value();
1318   else
1319     O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
1320 
1321   // Special case for i128
1322   if (ETy->isIntegerTy(128)) {
1323     O << " .b8 ";
1324     getSymbol(GVar)->print(O, MAI);
1325     O << "[16]";
1326     return;
1327   }
1328 
1329   if (ETy->isFloatingPointTy() || ETy->isIntOrPtrTy()) {
1330     O << " .";
1331     O << getPTXFundamentalTypeStr(ETy);
1332     O << " ";
1333     getSymbol(GVar)->print(O, MAI);
1334     return;
1335   }
1336 
1337   int64_t ElementSize = 0;
1338 
1339   // Although PTX has direct support for struct type and array type and LLVM IR
1340   // is very similar to PTX, the LLVM CodeGen does not support for targets that
1341   // support these high level field accesses. Structs and arrays are lowered
1342   // into arrays of bytes.
1343   switch (ETy->getTypeID()) {
1344   case Type::StructTyID:
1345   case Type::ArrayTyID:
1346   case Type::FixedVectorTyID:
1347     ElementSize = DL.getTypeStoreSize(ETy);
1348     O << " .b8 ";
1349     getSymbol(GVar)->print(O, MAI);
1350     O << "[";
1351     if (ElementSize) {
1352       O << ElementSize;
1353     }
1354     O << "]";
1355     break;
1356   default:
1357     llvm_unreachable("type not supported yet");
1358   }
1359 }
1360 
1361 void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I,
1362                                      int paramIndex, raw_ostream &O) {
1363   getSymbol(I->getParent())->print(O, MAI);
1364   O << "_param_" << paramIndex;
1365 }
1366 
1367 void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
1368   const DataLayout &DL = getDataLayout();
1369   const AttributeList &PAL = F->getAttributes();
1370   const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
1371   const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
1372 
1373   Function::const_arg_iterator I, E;
1374   unsigned paramIndex = 0;
1375   bool first = true;
1376   bool isKernelFunc = isKernelFunction(*F);
1377   bool isABI = (STI.getSmVersion() >= 20);
1378   bool hasImageHandles = STI.hasImageHandles();
1379   MVT thePointerTy = TLI->getPointerTy(DL);
1380 
1381   if (F->arg_empty()) {
1382     O << "()\n";
1383     return;
1384   }
1385 
1386   O << "(\n";
1387 
1388   for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) {
1389     Type *Ty = I->getType();
1390 
1391     if (!first)
1392       O << ",\n";
1393 
1394     first = false;
1395 
1396     // Handle image/sampler parameters
1397     if (isKernelFunction(*F)) {
1398       if (isSampler(*I) || isImage(*I)) {
1399         if (isImage(*I)) {
1400           std::string sname = std::string(I->getName());
1401           if (isImageWriteOnly(*I) || isImageReadWrite(*I)) {
1402             if (hasImageHandles)
1403               O << "\t.param .u64 .ptr .surfref ";
1404             else
1405               O << "\t.param .surfref ";
1406             CurrentFnSym->print(O, MAI);
1407             O << "_param_" << paramIndex;
1408           }
1409           else { // Default image is read_only
1410             if (hasImageHandles)
1411               O << "\t.param .u64 .ptr .texref ";
1412             else
1413               O << "\t.param .texref ";
1414             CurrentFnSym->print(O, MAI);
1415             O << "_param_" << paramIndex;
1416           }
1417         } else {
1418           if (hasImageHandles)
1419             O << "\t.param .u64 .ptr .samplerref ";
1420           else
1421             O << "\t.param .samplerref ";
1422           CurrentFnSym->print(O, MAI);
1423           O << "_param_" << paramIndex;
1424         }
1425         continue;
1426       }
1427     }
1428 
1429     auto getOptimalAlignForParam = [TLI, &DL, &PAL, F,
1430                                     paramIndex](Type *Ty) -> Align {
1431       Align TypeAlign = TLI->getFunctionParamOptimizedAlign(F, Ty, DL);
1432       MaybeAlign ParamAlign = PAL.getParamAlignment(paramIndex);
1433       return std::max(TypeAlign, ParamAlign.valueOrOne());
1434     };
1435 
1436     if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
1437       if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
1438         // Just print .param .align <a> .b8 .param[size];
1439         // <a>  = optimal alignment for the element type; always multiple of
1440         //        PAL.getParamAlignment
1441         // size = typeallocsize of element type
1442         Align OptimalAlign = getOptimalAlignForParam(Ty);
1443 
1444         O << "\t.param .align " << OptimalAlign.value() << " .b8 ";
1445         printParamName(I, paramIndex, O);
1446         O << "[" << DL.getTypeAllocSize(Ty) << "]";
1447 
1448         continue;
1449       }
1450       // Just a scalar
1451       auto *PTy = dyn_cast<PointerType>(Ty);
1452       if (isKernelFunc) {
1453         if (PTy) {
1454           // Special handling for pointer arguments to kernel
1455           O << "\t.param .u" << thePointerTy.getSizeInBits() << " ";
1456 
1457           if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() !=
1458               NVPTX::CUDA) {
1459             int addrSpace = PTy->getAddressSpace();
1460             switch (addrSpace) {
1461             default:
1462               O << ".ptr ";
1463               break;
1464             case ADDRESS_SPACE_CONST:
1465               O << ".ptr .const ";
1466               break;
1467             case ADDRESS_SPACE_SHARED:
1468               O << ".ptr .shared ";
1469               break;
1470             case ADDRESS_SPACE_GLOBAL:
1471               O << ".ptr .global ";
1472               break;
1473             }
1474             Align ParamAlign = I->getParamAlign().valueOrOne();
1475             O << ".align " << ParamAlign.value() << " ";
1476           }
1477           printParamName(I, paramIndex, O);
1478           continue;
1479         }
1480 
1481         // non-pointer scalar to kernel func
1482         O << "\t.param .";
1483         // Special case: predicate operands become .u8 types
1484         if (Ty->isIntegerTy(1))
1485           O << "u8";
1486         else
1487           O << getPTXFundamentalTypeStr(Ty);
1488         O << " ";
1489         printParamName(I, paramIndex, O);
1490         continue;
1491       }
1492       // Non-kernel function, just print .param .b<size> for ABI
1493       // and .reg .b<size> for non-ABI
1494       unsigned sz = 0;
1495       if (isa<IntegerType>(Ty)) {
1496         sz = cast<IntegerType>(Ty)->getBitWidth();
1497         if (sz < 32)
1498           sz = 32;
1499       } else if (isa<PointerType>(Ty))
1500         sz = thePointerTy.getSizeInBits();
1501       else if (Ty->isHalfTy())
1502         // PTX ABI requires all scalar parameters to be at least 32
1503         // bits in size.  fp16 normally uses .b16 as its storage type
1504         // in PTX, so its size must be adjusted here, too.
1505         sz = 32;
1506       else
1507         sz = Ty->getPrimitiveSizeInBits();
1508       if (isABI)
1509         O << "\t.param .b" << sz << " ";
1510       else
1511         O << "\t.reg .b" << sz << " ";
1512       printParamName(I, paramIndex, O);
1513       continue;
1514     }
1515 
1516     // param has byVal attribute.
1517     Type *ETy = PAL.getParamByValType(paramIndex);
1518     assert(ETy && "Param should have byval type");
1519 
1520     if (isABI || isKernelFunc) {
1521       // Just print .param .align <a> .b8 .param[size];
1522       // <a>  = optimal alignment for the element type; always multiple of
1523       //        PAL.getParamAlignment
1524       // size = typeallocsize of element type
1525       Align OptimalAlign = getOptimalAlignForParam(ETy);
1526 
1527       // Work around a bug in ptxas. When PTX code takes address of
1528       // byval parameter with alignment < 4, ptxas generates code to
1529       // spill argument into memory. Alas on sm_50+ ptxas generates
1530       // SASS code that fails with misaligned access. To work around
1531       // the problem, make sure that we align byval parameters by at
1532       // least 4. Matching change must be made in LowerCall() where we
1533       // prepare parameters for the call.
1534       //
1535       // TODO: this will need to be undone when we get to support multi-TU
1536       // device-side compilation as it breaks ABI compatibility with nvcc.
1537       // Hopefully ptxas bug is fixed by then.
1538       if (!isKernelFunc && OptimalAlign < Align(4))
1539         OptimalAlign = Align(4);
1540       unsigned sz = DL.getTypeAllocSize(ETy);
1541       O << "\t.param .align " << OptimalAlign.value() << " .b8 ";
1542       printParamName(I, paramIndex, O);
1543       O << "[" << sz << "]";
1544       continue;
1545     } else {
1546       // Split the ETy into constituent parts and
1547       // print .param .b<size> <name> for each part.
1548       // Further, if a part is vector, print the above for
1549       // each vector element.
1550       SmallVector<EVT, 16> vtparts;
1551       ComputeValueVTs(*TLI, DL, ETy, vtparts);
1552       for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
1553         unsigned elems = 1;
1554         EVT elemtype = vtparts[i];
1555         if (vtparts[i].isVector()) {
1556           elems = vtparts[i].getVectorNumElements();
1557           elemtype = vtparts[i].getVectorElementType();
1558         }
1559 
1560         for (unsigned j = 0, je = elems; j != je; ++j) {
1561           unsigned sz = elemtype.getSizeInBits();
1562           if (elemtype.isInteger() && (sz < 32))
1563             sz = 32;
1564           O << "\t.reg .b" << sz << " ";
1565           printParamName(I, paramIndex, O);
1566           if (j < je - 1)
1567             O << ",\n";
1568           ++paramIndex;
1569         }
1570         if (i < e - 1)
1571           O << ",\n";
1572       }
1573       --paramIndex;
1574       continue;
1575     }
1576   }
1577 
1578   O << "\n)\n";
1579 }
1580 
1581 void NVPTXAsmPrinter::emitFunctionParamList(const MachineFunction &MF,
1582                                             raw_ostream &O) {
1583   const Function &F = MF.getFunction();
1584   emitFunctionParamList(&F, O);
1585 }
1586 
1587 void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
1588     const MachineFunction &MF) {
1589   SmallString<128> Str;
1590   raw_svector_ostream O(Str);
1591 
1592   // Map the global virtual register number to a register class specific
1593   // virtual register number starting from 1 with that class.
1594   const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo();
1595   //unsigned numRegClasses = TRI->getNumRegClasses();
1596 
1597   // Emit the Fake Stack Object
1598   const MachineFrameInfo &MFI = MF.getFrameInfo();
1599   int NumBytes = (int) MFI.getStackSize();
1600   if (NumBytes) {
1601     O << "\t.local .align " << MFI.getMaxAlign().value() << " .b8 \t"
1602       << DEPOTNAME << getFunctionNumber() << "[" << NumBytes << "];\n";
1603     if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
1604       O << "\t.reg .b64 \t%SP;\n";
1605       O << "\t.reg .b64 \t%SPL;\n";
1606     } else {
1607       O << "\t.reg .b32 \t%SP;\n";
1608       O << "\t.reg .b32 \t%SPL;\n";
1609     }
1610   }
1611 
1612   // Go through all virtual registers to establish the mapping between the
1613   // global virtual
1614   // register number and the per class virtual register number.
1615   // We use the per class virtual register number in the ptx output.
1616   unsigned int numVRs = MRI->getNumVirtRegs();
1617   for (unsigned i = 0; i < numVRs; i++) {
1618     Register vr = Register::index2VirtReg(i);
1619     const TargetRegisterClass *RC = MRI->getRegClass(vr);
1620     DenseMap<unsigned, unsigned> &regmap = VRegMapping[RC];
1621     int n = regmap.size();
1622     regmap.insert(std::make_pair(vr, n + 1));
1623   }
1624 
1625   // Emit register declarations
1626   // @TODO: Extract out the real register usage
1627   // O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n";
1628   // O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n";
1629   // O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n";
1630   // O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n";
1631   // O << "\t.reg .s64 %rd<" << NVPTXNumRegisters << ">;\n";
1632   // O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n";
1633   // O << "\t.reg .f64 %fd<" << NVPTXNumRegisters << ">;\n";
1634 
1635   // Emit declaration of the virtual registers or 'physical' registers for
1636   // each register class
1637   for (unsigned i=0; i< TRI->getNumRegClasses(); i++) {
1638     const TargetRegisterClass *RC = TRI->getRegClass(i);
1639     DenseMap<unsigned, unsigned> &regmap = VRegMapping[RC];
1640     std::string rcname = getNVPTXRegClassName(RC);
1641     std::string rcStr = getNVPTXRegClassStr(RC);
1642     int n = regmap.size();
1643 
1644     // Only declare those registers that may be used.
1645     if (n) {
1646        O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1)
1647          << ">;\n";
1648     }
1649   }
1650 
1651   OutStreamer->emitRawText(O.str());
1652 }
1653 
1654 void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) {
1655   APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
1656   bool ignored;
1657   unsigned int numHex;
1658   const char *lead;
1659 
1660   if (Fp->getType()->getTypeID() == Type::FloatTyID) {
1661     numHex = 8;
1662     lead = "0f";
1663     APF.convert(APFloat::IEEEsingle(), APFloat::rmNearestTiesToEven, &ignored);
1664   } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
1665     numHex = 16;
1666     lead = "0d";
1667     APF.convert(APFloat::IEEEdouble(), APFloat::rmNearestTiesToEven, &ignored);
1668   } else
1669     llvm_unreachable("unsupported fp type");
1670 
1671   APInt API = APF.bitcastToAPInt();
1672   O << lead << format_hex_no_prefix(API.getZExtValue(), numHex, /*Upper=*/true);
1673 }
1674 
1675 void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
1676   if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1677     O << CI->getValue();
1678     return;
1679   }
1680   if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
1681     printFPConstant(CFP, O);
1682     return;
1683   }
1684   if (isa<ConstantPointerNull>(CPV)) {
1685     O << "0";
1686     return;
1687   }
1688   if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1689     bool IsNonGenericPointer = false;
1690     if (GVar->getType()->getAddressSpace() != 0) {
1691       IsNonGenericPointer = true;
1692     }
1693     if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
1694       O << "generic(";
1695       getSymbol(GVar)->print(O, MAI);
1696       O << ")";
1697     } else {
1698       getSymbol(GVar)->print(O, MAI);
1699     }
1700     return;
1701   }
1702   if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1703     const Value *v = Cexpr->stripPointerCasts();
1704     PointerType *PTy = dyn_cast<PointerType>(Cexpr->getType());
1705     bool IsNonGenericPointer = false;
1706     if (PTy && PTy->getAddressSpace() != 0) {
1707       IsNonGenericPointer = true;
1708     }
1709     if (const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
1710       if (EmitGeneric && !isa<Function>(v) && !IsNonGenericPointer) {
1711         O << "generic(";
1712         getSymbol(GVar)->print(O, MAI);
1713         O << ")";
1714       } else {
1715         getSymbol(GVar)->print(O, MAI);
1716       }
1717       return;
1718     } else {
1719       lowerConstant(CPV)->print(O, MAI);
1720       return;
1721     }
1722   }
1723   llvm_unreachable("Not scalar type found in printScalarConstant()");
1724 }
1725 
1726 void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes,
1727                                    AggBuffer *AggBuffer) {
1728   const DataLayout &DL = getDataLayout();
1729   int AllocSize = DL.getTypeAllocSize(CPV->getType());
1730   if (isa<UndefValue>(CPV) || CPV->isNullValue()) {
1731     // Non-zero Bytes indicates that we need to zero-fill everything. Otherwise,
1732     // only the space allocated by CPV.
1733     AggBuffer->addZeros(Bytes ? Bytes : AllocSize);
1734     return;
1735   }
1736 
1737   // Helper for filling AggBuffer with APInts.
1738   auto AddIntToBuffer = [AggBuffer, Bytes](const APInt &Val) {
1739     size_t NumBytes = (Val.getBitWidth() + 7) / 8;
1740     SmallVector<unsigned char, 16> Buf(NumBytes);
1741     for (unsigned I = 0; I < NumBytes; ++I) {
1742       Buf[I] = Val.extractBitsAsZExtValue(8, I * 8);
1743     }
1744     AggBuffer->addBytes(Buf.data(), NumBytes, Bytes);
1745   };
1746 
1747   switch (CPV->getType()->getTypeID()) {
1748   case Type::IntegerTyID:
1749     if (const auto CI = dyn_cast<ConstantInt>(CPV)) {
1750       AddIntToBuffer(CI->getValue());
1751       break;
1752     }
1753     if (const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1754       if (const auto *CI =
1755               dyn_cast<ConstantInt>(ConstantFoldConstant(Cexpr, DL))) {
1756         AddIntToBuffer(CI->getValue());
1757         break;
1758       }
1759       if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1760         Value *V = Cexpr->getOperand(0)->stripPointerCasts();
1761         AggBuffer->addSymbol(V, Cexpr->getOperand(0));
1762         AggBuffer->addZeros(AllocSize);
1763         break;
1764       }
1765     }
1766     llvm_unreachable("unsupported integer const type");
1767     break;
1768 
1769   case Type::HalfTyID:
1770   case Type::FloatTyID:
1771   case Type::DoubleTyID:
1772     AddIntToBuffer(cast<ConstantFP>(CPV)->getValueAPF().bitcastToAPInt());
1773     break;
1774 
1775   case Type::PointerTyID: {
1776     if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1777       AggBuffer->addSymbol(GVar, GVar);
1778     } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1779       const Value *v = Cexpr->stripPointerCasts();
1780       AggBuffer->addSymbol(v, Cexpr);
1781     }
1782     AggBuffer->addZeros(AllocSize);
1783     break;
1784   }
1785 
1786   case Type::ArrayTyID:
1787   case Type::FixedVectorTyID:
1788   case Type::StructTyID: {
1789     if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
1790       bufferAggregateConstant(CPV, AggBuffer);
1791       if (Bytes > AllocSize)
1792         AggBuffer->addZeros(Bytes - AllocSize);
1793     } else if (isa<ConstantAggregateZero>(CPV))
1794       AggBuffer->addZeros(Bytes);
1795     else
1796       llvm_unreachable("Unexpected Constant type");
1797     break;
1798   }
1799 
1800   default:
1801     llvm_unreachable("unsupported type");
1802   }
1803 }
1804 
1805 void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV,
1806                                               AggBuffer *aggBuffer) {
1807   const DataLayout &DL = getDataLayout();
1808   int Bytes;
1809 
1810   // Integers of arbitrary width
1811   if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1812     APInt Val = CI->getValue();
1813     for (unsigned I = 0, E = DL.getTypeAllocSize(CPV->getType()); I < E; ++I) {
1814       uint8_t Byte = Val.getLoBits(8).getZExtValue();
1815       aggBuffer->addBytes(&Byte, 1, 1);
1816       Val.lshrInPlace(8);
1817     }
1818     return;
1819   }
1820 
1821   // Old constants
1822   if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
1823     if (CPV->getNumOperands())
1824       for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i)
1825         bufferLEByte(cast<Constant>(CPV->getOperand(i)), 0, aggBuffer);
1826     return;
1827   }
1828 
1829   if (const ConstantDataSequential *CDS =
1830           dyn_cast<ConstantDataSequential>(CPV)) {
1831     if (CDS->getNumElements())
1832       for (unsigned i = 0; i < CDS->getNumElements(); ++i)
1833         bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
1834                      aggBuffer);
1835     return;
1836   }
1837 
1838   if (isa<ConstantStruct>(CPV)) {
1839     if (CPV->getNumOperands()) {
1840       StructType *ST = cast<StructType>(CPV->getType());
1841       for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) {
1842         if (i == (e - 1))
1843           Bytes = DL.getStructLayout(ST)->getElementOffset(0) +
1844                   DL.getTypeAllocSize(ST) -
1845                   DL.getStructLayout(ST)->getElementOffset(i);
1846         else
1847           Bytes = DL.getStructLayout(ST)->getElementOffset(i + 1) -
1848                   DL.getStructLayout(ST)->getElementOffset(i);
1849         bufferLEByte(cast<Constant>(CPV->getOperand(i)), Bytes, aggBuffer);
1850       }
1851     }
1852     return;
1853   }
1854   llvm_unreachable("unsupported constant type in printAggregateConstant()");
1855 }
1856 
1857 /// lowerConstantForGV - Return an MCExpr for the given Constant.  This is mostly
1858 /// a copy from AsmPrinter::lowerConstant, except customized to only handle
1859 /// expressions that are representable in PTX and create
1860 /// NVPTXGenericMCSymbolRefExpr nodes for addrspacecast instructions.
1861 const MCExpr *
1862 NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV, bool ProcessingGeneric) {
1863   MCContext &Ctx = OutContext;
1864 
1865   if (CV->isNullValue() || isa<UndefValue>(CV))
1866     return MCConstantExpr::create(0, Ctx);
1867 
1868   if (const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
1869     return MCConstantExpr::create(CI->getZExtValue(), Ctx);
1870 
1871   if (const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
1872     const MCSymbolRefExpr *Expr =
1873       MCSymbolRefExpr::create(getSymbol(GV), Ctx);
1874     if (ProcessingGeneric) {
1875       return NVPTXGenericMCSymbolRefExpr::create(Expr, Ctx);
1876     } else {
1877       return Expr;
1878     }
1879   }
1880 
1881   const ConstantExpr *CE = dyn_cast<ConstantExpr>(CV);
1882   if (!CE) {
1883     llvm_unreachable("Unknown constant value to lower!");
1884   }
1885 
1886   switch (CE->getOpcode()) {
1887   default: {
1888     // If the code isn't optimized, there may be outstanding folding
1889     // opportunities. Attempt to fold the expression using DataLayout as a
1890     // last resort before giving up.
1891     Constant *C = ConstantFoldConstant(CE, getDataLayout());
1892     if (C != CE)
1893       return lowerConstantForGV(C, ProcessingGeneric);
1894 
1895     // Otherwise report the problem to the user.
1896     std::string S;
1897     raw_string_ostream OS(S);
1898     OS << "Unsupported expression in static initializer: ";
1899     CE->printAsOperand(OS, /*PrintType=*/false,
1900                    !MF ? nullptr : MF->getFunction().getParent());
1901     report_fatal_error(Twine(OS.str()));
1902   }
1903 
1904   case Instruction::AddrSpaceCast: {
1905     // Strip the addrspacecast and pass along the operand
1906     PointerType *DstTy = cast<PointerType>(CE->getType());
1907     if (DstTy->getAddressSpace() == 0) {
1908       return lowerConstantForGV(cast<const Constant>(CE->getOperand(0)), true);
1909     }
1910     std::string S;
1911     raw_string_ostream OS(S);
1912     OS << "Unsupported expression in static initializer: ";
1913     CE->printAsOperand(OS, /*PrintType=*/ false,
1914                        !MF ? nullptr : MF->getFunction().getParent());
1915     report_fatal_error(Twine(OS.str()));
1916   }
1917 
1918   case Instruction::GetElementPtr: {
1919     const DataLayout &DL = getDataLayout();
1920 
1921     // Generate a symbolic expression for the byte address
1922     APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0);
1923     cast<GEPOperator>(CE)->accumulateConstantOffset(DL, OffsetAI);
1924 
1925     const MCExpr *Base = lowerConstantForGV(CE->getOperand(0),
1926                                             ProcessingGeneric);
1927     if (!OffsetAI)
1928       return Base;
1929 
1930     int64_t Offset = OffsetAI.getSExtValue();
1931     return MCBinaryExpr::createAdd(Base, MCConstantExpr::create(Offset, Ctx),
1932                                    Ctx);
1933   }
1934 
1935   case Instruction::Trunc:
1936     // We emit the value and depend on the assembler to truncate the generated
1937     // expression properly.  This is important for differences between
1938     // blockaddress labels.  Since the two labels are in the same function, it
1939     // is reasonable to treat their delta as a 32-bit value.
1940     LLVM_FALLTHROUGH;
1941   case Instruction::BitCast:
1942     return lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
1943 
1944   case Instruction::IntToPtr: {
1945     const DataLayout &DL = getDataLayout();
1946 
1947     // Handle casts to pointers by changing them into casts to the appropriate
1948     // integer type.  This promotes constant folding and simplifies this code.
1949     Constant *Op = CE->getOperand(0);
1950     Op = ConstantExpr::getIntegerCast(Op, DL.getIntPtrType(CV->getType()),
1951                                       false/*ZExt*/);
1952     return lowerConstantForGV(Op, ProcessingGeneric);
1953   }
1954 
1955   case Instruction::PtrToInt: {
1956     const DataLayout &DL = getDataLayout();
1957 
1958     // Support only foldable casts to/from pointers that can be eliminated by
1959     // changing the pointer to the appropriately sized integer type.
1960     Constant *Op = CE->getOperand(0);
1961     Type *Ty = CE->getType();
1962 
1963     const MCExpr *OpExpr = lowerConstantForGV(Op, ProcessingGeneric);
1964 
1965     // We can emit the pointer value into this slot if the slot is an
1966     // integer slot equal to the size of the pointer.
1967     if (DL.getTypeAllocSize(Ty) == DL.getTypeAllocSize(Op->getType()))
1968       return OpExpr;
1969 
1970     // Otherwise the pointer is smaller than the resultant integer, mask off
1971     // the high bits so we are sure to get a proper truncation if the input is
1972     // a constant expr.
1973     unsigned InBits = DL.getTypeAllocSizeInBits(Op->getType());
1974     const MCExpr *MaskExpr = MCConstantExpr::create(~0ULL >> (64-InBits), Ctx);
1975     return MCBinaryExpr::createAnd(OpExpr, MaskExpr, Ctx);
1976   }
1977 
1978   // The MC library also has a right-shift operator, but it isn't consistently
1979   // signed or unsigned between different targets.
1980   case Instruction::Add: {
1981     const MCExpr *LHS = lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
1982     const MCExpr *RHS = lowerConstantForGV(CE->getOperand(1), ProcessingGeneric);
1983     switch (CE->getOpcode()) {
1984     default: llvm_unreachable("Unknown binary operator constant cast expr");
1985     case Instruction::Add: return MCBinaryExpr::createAdd(LHS, RHS, Ctx);
1986     }
1987   }
1988   }
1989 }
1990 
1991 // Copy of MCExpr::print customized for NVPTX
1992 void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) {
1993   switch (Expr.getKind()) {
1994   case MCExpr::Target:
1995     return cast<MCTargetExpr>(&Expr)->printImpl(OS, MAI);
1996   case MCExpr::Constant:
1997     OS << cast<MCConstantExpr>(Expr).getValue();
1998     return;
1999 
2000   case MCExpr::SymbolRef: {
2001     const MCSymbolRefExpr &SRE = cast<MCSymbolRefExpr>(Expr);
2002     const MCSymbol &Sym = SRE.getSymbol();
2003     Sym.print(OS, MAI);
2004     return;
2005   }
2006 
2007   case MCExpr::Unary: {
2008     const MCUnaryExpr &UE = cast<MCUnaryExpr>(Expr);
2009     switch (UE.getOpcode()) {
2010     case MCUnaryExpr::LNot:  OS << '!'; break;
2011     case MCUnaryExpr::Minus: OS << '-'; break;
2012     case MCUnaryExpr::Not:   OS << '~'; break;
2013     case MCUnaryExpr::Plus:  OS << '+'; break;
2014     }
2015     printMCExpr(*UE.getSubExpr(), OS);
2016     return;
2017   }
2018 
2019   case MCExpr::Binary: {
2020     const MCBinaryExpr &BE = cast<MCBinaryExpr>(Expr);
2021 
2022     // Only print parens around the LHS if it is non-trivial.
2023     if (isa<MCConstantExpr>(BE.getLHS()) || isa<MCSymbolRefExpr>(BE.getLHS()) ||
2024         isa<NVPTXGenericMCSymbolRefExpr>(BE.getLHS())) {
2025       printMCExpr(*BE.getLHS(), OS);
2026     } else {
2027       OS << '(';
2028       printMCExpr(*BE.getLHS(), OS);
2029       OS<< ')';
2030     }
2031 
2032     switch (BE.getOpcode()) {
2033     case MCBinaryExpr::Add:
2034       // Print "X-42" instead of "X+-42".
2035       if (const MCConstantExpr *RHSC = dyn_cast<MCConstantExpr>(BE.getRHS())) {
2036         if (RHSC->getValue() < 0) {
2037           OS << RHSC->getValue();
2038           return;
2039         }
2040       }
2041 
2042       OS <<  '+';
2043       break;
2044     default: llvm_unreachable("Unhandled binary operator");
2045     }
2046 
2047     // Only print parens around the LHS if it is non-trivial.
2048     if (isa<MCConstantExpr>(BE.getRHS()) || isa<MCSymbolRefExpr>(BE.getRHS())) {
2049       printMCExpr(*BE.getRHS(), OS);
2050     } else {
2051       OS << '(';
2052       printMCExpr(*BE.getRHS(), OS);
2053       OS << ')';
2054     }
2055     return;
2056   }
2057   }
2058 
2059   llvm_unreachable("Invalid expression kind!");
2060 }
2061 
2062 /// PrintAsmOperand - Print out an operand for an inline asm expression.
2063 ///
2064 bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
2065                                       const char *ExtraCode, raw_ostream &O) {
2066   if (ExtraCode && ExtraCode[0]) {
2067     if (ExtraCode[1] != 0)
2068       return true; // Unknown modifier.
2069 
2070     switch (ExtraCode[0]) {
2071     default:
2072       // See if this is a generic print operand
2073       return AsmPrinter::PrintAsmOperand(MI, OpNo, ExtraCode, O);
2074     case 'r':
2075       break;
2076     }
2077   }
2078 
2079   printOperand(MI, OpNo, O);
2080 
2081   return false;
2082 }
2083 
2084 bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI,
2085                                             unsigned OpNo,
2086                                             const char *ExtraCode,
2087                                             raw_ostream &O) {
2088   if (ExtraCode && ExtraCode[0])
2089     return true; // Unknown modifier
2090 
2091   O << '[';
2092   printMemOperand(MI, OpNo, O);
2093   O << ']';
2094 
2095   return false;
2096 }
2097 
2098 void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum,
2099                                    raw_ostream &O) {
2100   const MachineOperand &MO = MI->getOperand(opNum);
2101   switch (MO.getType()) {
2102   case MachineOperand::MO_Register:
2103     if (Register::isPhysicalRegister(MO.getReg())) {
2104       if (MO.getReg() == NVPTX::VRDepot)
2105         O << DEPOTNAME << getFunctionNumber();
2106       else
2107         O << NVPTXInstPrinter::getRegisterName(MO.getReg());
2108     } else {
2109       emitVirtualRegister(MO.getReg(), O);
2110     }
2111     break;
2112 
2113   case MachineOperand::MO_Immediate:
2114     O << MO.getImm();
2115     break;
2116 
2117   case MachineOperand::MO_FPImmediate:
2118     printFPConstant(MO.getFPImm(), O);
2119     break;
2120 
2121   case MachineOperand::MO_GlobalAddress:
2122     PrintSymbolOperand(MO, O);
2123     break;
2124 
2125   case MachineOperand::MO_MachineBasicBlock:
2126     MO.getMBB()->getSymbol()->print(O, MAI);
2127     break;
2128 
2129   default:
2130     llvm_unreachable("Operand type not supported.");
2131   }
2132 }
2133 
2134 void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum,
2135                                       raw_ostream &O, const char *Modifier) {
2136   printOperand(MI, opNum, O);
2137 
2138   if (Modifier && strcmp(Modifier, "add") == 0) {
2139     O << ", ";
2140     printOperand(MI, opNum + 1, O);
2141   } else {
2142     if (MI->getOperand(opNum + 1).isImm() &&
2143         MI->getOperand(opNum + 1).getImm() == 0)
2144       return; // don't print ',0' or '+0'
2145     O << "+";
2146     printOperand(MI, opNum + 1, O);
2147   }
2148 }
2149 
2150 // Force static initialization.
2151 extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXAsmPrinter() {
2152   RegisterAsmPrinter<NVPTXAsmPrinter> X(getTheNVPTXTarget32());
2153   RegisterAsmPrinter<NVPTXAsmPrinter> Y(getTheNVPTXTarget64());
2154 }
2155