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