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