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