/build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

Bug Summary

File:	llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
Warning:	line 1544, column 17 Called C++ object pointer is null
Annotated Source Code

Press '?' to see keyboard shortcuts
Show analyzer invocation
clang -cc1 -cc1 -triple x86_64-pc-linux-gnu -analyze -disable-free -disable-llvm-verifier -discard-value-names -main-file-name NVPTXAsmPrinter.cpp -analyzer-store=region -analyzer-opt-analyze-nested-blocks -analyzer-checker=core -analyzer-checker=apiModeling -analyzer-checker=unix -analyzer-checker=deadcode -analyzer-checker=cplusplus -analyzer-checker=security.insecureAPI.UncheckedReturn -analyzer-checker=security.insecureAPI.getpw -analyzer-checker=security.insecureAPI.gets -analyzer-checker=security.insecureAPI.mktemp -analyzer-checker=security.insecureAPI.mkstemp -analyzer-checker=security.insecureAPI.vfork -analyzer-checker=nullability.NullPassedToNonnull -analyzer-checker=nullability.NullReturnedFromNonnull -analyzer-output plist -w -setup-static-analyzer -analyzer-config-compatibility-mode=true -mrelocation-model pic -pic-level 2 -mframe-pointer=none -fmath-errno -fno-rounding-math -mconstructor-aliases -munwind-tables -target-cpu x86-64 -tune-cpu generic -debugger-tuning=gdb -ffunction-sections -fdata-sections -fcoverage-compilation-dir=/build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/lib/Target/NVPTX -resource-dir /usr/lib/llvm-14/lib/clang/14.0.0 -D _GNU_SOURCE -D __STDC_CONSTANT_MACROS -D __STDC_FORMAT_MACROS -D __STDC_LIMIT_MACROS -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/lib/Target/NVPTX -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/llvm/lib/Target/NVPTX -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/include -I /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/llvm/include -D NDEBUG -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/x86_64-linux-gnu/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/backward -internal-isystem /usr/lib/llvm-14/lib/clang/14.0.0/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../x86_64-linux-gnu/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -O2 -Wno-unused-parameter -Wwrite-strings -Wno-missing-field-initializers -Wno-long-long -Wno-maybe-uninitialized -Wno-class-memaccess -Wno-redundant-move -Wno-pessimizing-move -Wno-noexcept-type -Wno-comment -std=c++14 -fdeprecated-macro -fdebug-compilation-dir=/build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/build-llvm/lib/Target/NVPTX -fdebug-prefix-map=/build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e=. -ferror-limit 19 -fvisibility hidden -fvisibility-inlines-hidden -stack-protector 2 -fgnuc-version=4.2.1 -vectorize-loops -vectorize-slp -analyzer-output=html -analyzer-config stable-report-filename=true -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/scan-build-2021-09-04-040900-46481-1 -x c++ /build/llvm-toolchain-snapshot-14~++20210903100615+fd66b44ec19e/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
1//===-- NVPTXAsmPrinter.cpp - NVPTX LLVM assembly writer ------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This file contains a printer that converts from our internal representation
10// of machine-dependent LLVM code to NVPTX assembly language.
11//
12//===----------------------------------------------------------------------===//

14#include "NVPTXAsmPrinter.h"
15#include "MCTargetDesc/NVPTXBaseInfo.h"
16#include "MCTargetDesc/NVPTXInstPrinter.h"
17#include "MCTargetDesc/NVPTXMCAsmInfo.h"
18#include "MCTargetDesc/NVPTXTargetStreamer.h"
19#include "NVPTX.h"
20#include "NVPTXMCExpr.h"
21#include "NVPTXMachineFunctionInfo.h"
22#include "NVPTXRegisterInfo.h"
23#include "NVPTXSubtarget.h"
24#include "NVPTXTargetMachine.h"
25#include "NVPTXUtilities.h"
26#include "TargetInfo/NVPTXTargetInfo.h"
27#include "cl_common_defines.h"
28#include "llvm/ADT/APFloat.h"
29#include "llvm/ADT/APInt.h"
30#include "llvm/ADT/DenseMap.h"
31#include "llvm/ADT/DenseSet.h"
32#include "llvm/ADT/SmallString.h"
33#include "llvm/ADT/SmallVector.h"
34#include "llvm/ADT/StringExtras.h"
35#include "llvm/ADT/StringRef.h"
36#include "llvm/ADT/Triple.h"
37#include "llvm/ADT/Twine.h"
38#include "llvm/Analysis/ConstantFolding.h"
39#include "llvm/CodeGen/Analysis.h"
40#include "llvm/CodeGen/MachineBasicBlock.h"
41#include "llvm/CodeGen/MachineFrameInfo.h"
42#include "llvm/CodeGen/MachineFunction.h"
43#include "llvm/CodeGen/MachineInstr.h"
44#include "llvm/CodeGen/MachineLoopInfo.h"
45#include "llvm/CodeGen/MachineModuleInfo.h"
46#include "llvm/CodeGen/MachineOperand.h"
47#include "llvm/CodeGen/MachineRegisterInfo.h"
48#include "llvm/CodeGen/TargetLowering.h"
49#include "llvm/CodeGen/TargetRegisterInfo.h"
50#include "llvm/CodeGen/ValueTypes.h"
51#include "llvm/IR/Attributes.h"
52#include "llvm/IR/BasicBlock.h"
53#include "llvm/IR/Constant.h"
54#include "llvm/IR/Constants.h"
55#include "llvm/IR/DataLayout.h"
56#include "llvm/IR/DebugInfo.h"
57#include "llvm/IR/DebugInfoMetadata.h"
58#include "llvm/IR/DebugLoc.h"
59#include "llvm/IR/DerivedTypes.h"
60#include "llvm/IR/Function.h"
61#include "llvm/IR/GlobalValue.h"
62#include "llvm/IR/GlobalVariable.h"
63#include "llvm/IR/Instruction.h"
64#include "llvm/IR/LLVMContext.h"
65#include "llvm/IR/Module.h"
66#include "llvm/IR/Operator.h"
67#include "llvm/IR/Type.h"
68#include "llvm/IR/User.h"
69#include "llvm/MC/MCExpr.h"
70#include "llvm/MC/MCInst.h"
71#include "llvm/MC/MCInstrDesc.h"
72#include "llvm/MC/MCStreamer.h"
73#include "llvm/MC/MCSymbol.h"
74#include "llvm/Support/Casting.h"
75#include "llvm/Support/CommandLine.h"
76#include "llvm/Support/ErrorHandling.h"
77#include "llvm/Support/MachineValueType.h"
78#include "llvm/Support/Path.h"
79#include "llvm/Support/TargetRegistry.h"
80#include "llvm/Support/raw_ostream.h"
81#include "llvm/Target/TargetLoweringObjectFile.h"
82#include "llvm/Target/TargetMachine.h"
83#include "llvm/Transforms/Utils/UnrollLoop.h"
84#include <cassert>
85#include <cstdint>
86#include <cstring>
87#include <new>
88#include <string>
89#include <utility>
90#include <vector>

92using namespace llvm;

94#define DEPOTNAME"__local_depot" "__local_depot"

96/// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
97/// depends.
98static void
99DiscoverDependentGlobals(const Value *V,
                       DenseSet<const GlobalVariable *> &Globals) {
if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(V))
  Globals.insert(GV);
else {
  if (const User *U = dyn_cast<User>(V)) {
    for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) {
      DiscoverDependentGlobals(U->getOperand(i), Globals);
    }
  }
}
110}

112/// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable
113/// instances to be emitted, but only after any dependents have been added
114/// first.s
115static void
116VisitGlobalVariableForEmission(const GlobalVariable *GV,
                             SmallVectorImpl<const GlobalVariable *> &Order,
                             DenseSet<const GlobalVariable *> &Visited,
                             DenseSet<const GlobalVariable *> &Visiting) {
// Have we already visited this one?
if (Visited.count(GV))
  return;

// Do we have a circular dependency?
if (!Visiting.insert(GV).second)
  report_fatal_error("Circular dependency found in global variable set");

// Make sure we visit all dependents first
DenseSet<const GlobalVariable *> Others;
for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i)
  DiscoverDependentGlobals(GV->getOperand(i), Others);

for (DenseSet<const GlobalVariable *>::iterator I = Others.begin(),
                                                E = Others.end();
     I != E; ++I)
  VisitGlobalVariableForEmission(*I, Order, Visited, Visiting);

// Now we can visit ourself
Order.push_back(GV);
Visited.insert(GV);
Visiting.erase(GV);
142}

144void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {
MCInst Inst;
lowerToMCInst(MI, Inst);
EmitToStreamer(*OutStreamer, Inst);
148}

150// Handle symbol backtracking for targets that do not support image handles
151bool NVPTXAsmPrinter::lowerImageHandleOperand(const MachineInstr *MI,
                                         unsigned OpNo, MCOperand &MCOp) {
const MachineOperand &MO = MI->getOperand(OpNo);
const MCInstrDesc &MCID = MI->getDesc();

if (MCID.TSFlags & NVPTXII::IsTexFlag) {
  // This is a texture fetch, so operand 4 is a texref and operand 5 is
  // a samplerref
  if (OpNo == 4 && MO.isImm()) {
    lowerImageHandleSymbol(MO.getImm(), MCOp);
    return true;
  }
  if (OpNo == 5 && MO.isImm() && !(MCID.TSFlags & NVPTXII::IsTexModeUnifiedFlag)) {
    lowerImageHandleSymbol(MO.getImm(), MCOp);
    return true;
  }

  return false;
} else if (MCID.TSFlags & NVPTXII::IsSuldMask) {
  unsigned VecSize =
    1 << (((MCID.TSFlags & NVPTXII::IsSuldMask) >> NVPTXII::IsSuldShift) - 1);

  // For a surface load of vector size N, the Nth operand will be the surfref
  if (OpNo == VecSize && MO.isImm()) {
    lowerImageHandleSymbol(MO.getImm(), MCOp);
    return true;
  }

  return false;
} else if (MCID.TSFlags & NVPTXII::IsSustFlag) {
  // This is a surface store, so operand 0 is a surfref
  if (OpNo == 0 && MO.isImm()) {
    lowerImageHandleSymbol(MO.getImm(), MCOp);
    return true;
  }

  return false;
} else if (MCID.TSFlags & NVPTXII::IsSurfTexQueryFlag) {
  // This is a query, so operand 1 is a surfref/texref
  if (OpNo == 1 && MO.isImm()) {
    lowerImageHandleSymbol(MO.getImm(), MCOp);
    return true;
  }

  return false;
}

return false;
199}

201void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) {
// Ewwww
LLVMTargetMachine &TM = const_cast<LLVMTargetMachine&>(MF->getTarget());
NVPTXTargetMachine &nvTM = static_cast<NVPTXTargetMachine&>(TM);
const NVPTXMachineFunctionInfo *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
const char *Sym = MFI->getImageHandleSymbol(Index);
std::string *SymNamePtr =
  nvTM.getManagedStrPool()->getManagedString(Sym);
MCOp = GetSymbolRef(OutContext.getOrCreateSymbol(StringRef(*SymNamePtr)));
210}

212void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
OutMI.setOpcode(MI->getOpcode());
// Special: Do not mangle symbol operand of CALL_PROTOTYPE
if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
  const MachineOperand &MO = MI->getOperand(0);
  OutMI.addOperand(GetSymbolRef(
    OutContext.getOrCreateSymbol(Twine(MO.getSymbolName()))));
  return;
}

const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
for (unsigned i = 0, e = MI->getNumOperands(); i != e; ++i) {
  const MachineOperand &MO = MI->getOperand(i);

  MCOperand MCOp;
  if (!STI.hasImageHandles()) {
    if (lowerImageHandleOperand(MI, i, MCOp)) {
      OutMI.addOperand(MCOp);
      continue;
    }
  }

  if (lowerOperand(MO, MCOp))
    OutMI.addOperand(MCOp);
}
237}

239bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO,
                                 MCOperand &MCOp) {
switch (MO.getType()) {
default: llvm_unreachable("unknown operand type")__builtin_unreachable();
case MachineOperand::MO_Register:
  MCOp = MCOperand::createReg(encodeVirtualRegister(MO.getReg()));
  break;
case MachineOperand::MO_Immediate:
  MCOp = MCOperand::createImm(MO.getImm());
  break;
case MachineOperand::MO_MachineBasicBlock:
  MCOp = MCOperand::createExpr(MCSymbolRefExpr::create(
      MO.getMBB()->getSymbol(), OutContext));
  break;
case MachineOperand::MO_ExternalSymbol:
  MCOp = GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
  break;
case MachineOperand::MO_GlobalAddress:
  MCOp = GetSymbolRef(getSymbol(MO.getGlobal()));
  break;
case MachineOperand::MO_FPImmediate: {
  const ConstantFP *Cnt = MO.getFPImm();
  const APFloat &Val = Cnt->getValueAPF();

  switch (Cnt->getType()->getTypeID()) {
  default: report_fatal_error("Unsupported FP type"); break;
  case Type::HalfTyID:
    MCOp = MCOperand::createExpr(
      NVPTXFloatMCExpr::createConstantFPHalf(Val, OutContext));
    break;
  case Type::FloatTyID:
    MCOp = MCOperand::createExpr(
      NVPTXFloatMCExpr::createConstantFPSingle(Val, OutContext));
    break;
  case Type::DoubleTyID:
    MCOp = MCOperand::createExpr(
      NVPTXFloatMCExpr::createConstantFPDouble(Val, OutContext));
    break;
  }
  break;
}
}
return true;
282}

284unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
if (Register::isVirtualRegister(Reg)) {
  const TargetRegisterClass *RC = MRI->getRegClass(Reg);

  DenseMap<unsigned, unsigned> &RegMap = VRegMapping[RC];
  unsigned RegNum = RegMap[Reg];

  // Encode the register class in the upper 4 bits
  // Must be kept in sync with NVPTXInstPrinter::printRegName
  unsigned Ret = 0;
  if (RC == &NVPTX::Int1RegsRegClass) {
    Ret = (1 << 28);
  } else if (RC == &NVPTX::Int16RegsRegClass) {
    Ret = (2 << 28);
  } else if (RC == &NVPTX::Int32RegsRegClass) {
    Ret = (3 << 28);
  } else if (RC == &NVPTX::Int64RegsRegClass) {
    Ret = (4 << 28);
  } else if (RC == &NVPTX::Float32RegsRegClass) {
    Ret = (5 << 28);
  } else if (RC == &NVPTX::Float64RegsRegClass) {
    Ret = (6 << 28);
  } else if (RC == &NVPTX::Float16RegsRegClass) {
    Ret = (7 << 28);
  } else if (RC == &NVPTX::Float16x2RegsRegClass) {
    Ret = (8 << 28);
  } else {
    report_fatal_error("Bad register class");
  }

  // Insert the vreg number
  Ret |= (RegNum & 0x0FFFFFFF);
  return Ret;
} else {
  // Some special-use registers are actually physical registers.
  // Encode this as the register class ID of 0 and the real register ID.
  return Reg & 0x0FFFFFFF;
}
322}

324MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
const MCExpr *Expr;
Expr = MCSymbolRefExpr::create(Symbol, MCSymbolRefExpr::VK_None,
                               OutContext);
return MCOperand::createExpr(Expr);
329}

331void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
const DataLayout &DL = getDataLayout();
const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
const TargetLowering *TLI = STI.getTargetLowering();

Type *Ty = F->getReturnType();

bool isABI = (STI.getSmVersion() >= 20);

if (Ty->getTypeID() == Type::VoidTyID)
  return;

O << " (";

if (isABI) {
  if (Ty->isFloatingPointTy() || (Ty->isIntegerTy() && !Ty->isIntegerTy(128))) {
    unsigned size = 0;
    if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
      size = ITy->getBitWidth();
    } else {
      assert(Ty->isFloatingPointTy() && "Floating point type expected here")(static_cast<void> (0));
      size = Ty->getPrimitiveSizeInBits();
    }
    // PTX ABI requires all scalar return values to be at least 32
    // bits in size.  fp16 normally uses .b16 as its storage type in
    // PTX, so its size must be adjusted here, too.
    if (size < 32)
      size = 32;

    O << ".param .b" << size << " func_retval0";
  } else if (isa<PointerType>(Ty)) {
    O << ".param .b" << TLI->getPointerTy(DL).getSizeInBits()
      << " func_retval0";
  } else if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
    unsigned totalsz = DL.getTypeAllocSize(Ty);
    unsigned retAlignment = 0;
    if (!getAlign(*F, 0, retAlignment))
      retAlignment = DL.getABITypeAlignment(Ty);
    O << ".param .align " << retAlignment << " .b8 func_retval0[" << totalsz
      << "]";
  } else
    llvm_unreachable("Unknown return type")__builtin_unreachable();
} else {
  SmallVector<EVT, 16> vtparts;
  ComputeValueVTs(*TLI, DL, Ty, vtparts);
  unsigned idx = 0;
  for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
    unsigned elems = 1;
    EVT elemtype = vtparts[i];
    if (vtparts[i].isVector()) {
      elems = vtparts[i].getVectorNumElements();
      elemtype = vtparts[i].getVectorElementType();
    }

    for (unsigned j = 0, je = elems; j != je; ++j) {
      unsigned sz = elemtype.getSizeInBits();
      if (elemtype.isInteger() && (sz < 32))
        sz = 32;
      O << ".reg .b" << sz << " func_retval" << idx;
      if (j < je - 1)
        O << ", ";
      ++idx;
    }
    if (i < e - 1)
      O << ", ";
  }
}
O << ") ";
399}

401void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
                                      raw_ostream &O) {
const Function &F = MF.getFunction();
printReturnValStr(&F, O);
405}

407// Return true if MBB is the header of a loop marked with
408// llvm.loop.unroll.disable.
409// TODO: consider "#pragma unroll 1" which is equivalent to "#pragma nounroll".
410bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
  const MachineBasicBlock &MBB) const {
MachineLoopInfo &LI = getAnalysis<MachineLoopInfo>();
// We insert .pragma "nounroll" only to the loop header.
if (!LI.isLoopHeader(&MBB))
  return false;

// llvm.loop.unroll.disable is marked on the back edges of a loop. Therefore,
// we iterate through each back edge of the loop with header MBB, and check
// whether its metadata contains llvm.loop.unroll.disable.
for (auto I = MBB.pred_begin(); I != MBB.pred_end(); ++I) {
  const MachineBasicBlock *PMBB = *I;
  if (LI.getLoopFor(PMBB) != LI.getLoopFor(&MBB)) {
    // Edges from other loops to MBB are not back edges.
    continue;
  }
  if (const BasicBlock *PBB = PMBB->getBasicBlock()) {
    if (MDNode *LoopID =
            PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
      if (GetUnrollMetadata(LoopID, "llvm.loop.unroll.disable"))
        return true;
    }
  }
}
return false;
435}

437void NVPTXAsmPrinter::emitBasicBlockStart(const MachineBasicBlock &MBB) {
AsmPrinter::emitBasicBlockStart(MBB);
if (isLoopHeaderOfNoUnroll(MBB))
  OutStreamer->emitRawText(StringRef("\t.pragma \"nounroll\";\n"));
441}

443void NVPTXAsmPrinter::emitFunctionEntryLabel() {
SmallString<128> Str;
raw_svector_ostream O(Str);

if (!GlobalsEmitted) {
  emitGlobals(*MF->getFunction().getParent());
  GlobalsEmitted = true;
}

// Set up
MRI = &MF->getRegInfo();
F = &MF->getFunction();
emitLinkageDirective(F, O);
if (isKernelFunction(*F))
  O << ".entry ";
else {
  O << ".func ";
  printReturnValStr(*MF, O);
}

CurrentFnSym->print(O, MAI);

emitFunctionParamList(*MF, O);

if (isKernelFunction(*F))
  emitKernelFunctionDirectives(*F, O);

OutStreamer->emitRawText(O.str());

VRegMapping.clear();
// Emit open brace for function body.
OutStreamer->emitRawText(StringRef("{\n"));
setAndEmitFunctionVirtualRegisters(*MF);
// Emit initial .loc debug directive for correct relocation symbol data.
if (MMI && MMI->hasDebugInfo())
  emitInitialRawDwarfLocDirective(*MF);
479}

481bool NVPTXAsmPrinter::runOnMachineFunction(MachineFunction &F) {
bool Result = AsmPrinter::runOnMachineFunction(F);
// Emit closing brace for the body of function F.
// The closing brace must be emitted here because we need to emit additional
// debug labels/data after the last basic block.
// We need to emit the closing brace here because we don't have function that
// finished emission of the function body.
OutStreamer->emitRawText(StringRef("}\n"));
return Result;
490}

492void NVPTXAsmPrinter::emitFunctionBodyStart() {
SmallString<128> Str;
raw_svector_ostream O(Str);
emitDemotedVars(&MF->getFunction(), O);
OutStreamer->emitRawText(O.str());
497}

499void NVPTXAsmPrinter::emitFunctionBodyEnd() {
VRegMapping.clear();
501}

503const MCSymbol *NVPTXAsmPrinter::getFunctionFrameSymbol() const {
  SmallString<128> Str;
  raw_svector_ostream(Str) << DEPOTNAME"__local_depot" << getFunctionNumber();
  return OutContext.getOrCreateSymbol(Str);
507}

509void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const {
Register RegNo = MI->getOperand(0).getReg();
if (Register::isVirtualRegister(RegNo)) {
  OutStreamer->AddComment(Twine("implicit-def: ") +
                          getVirtualRegisterName(RegNo));
} else {
  const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
  OutStreamer->AddComment(Twine("implicit-def: ") +
                          STI.getRegisterInfo()->getName(RegNo));
}
OutStreamer->AddBlankLine();
520}

522void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F,
                                                 raw_ostream &O) const {
// If the NVVM IR has some of reqntid* specified, then output
// the reqntid directive, and set the unspecified ones to 1.
// If none of reqntid* is specified, don't output reqntid directive.
unsigned reqntidx, reqntidy, reqntidz;
bool specified = false;
if (!getReqNTIDx(F, reqntidx))
  reqntidx = 1;
else
  specified = true;
if (!getReqNTIDy(F, reqntidy))
  reqntidy = 1;
else
  specified = true;
if (!getReqNTIDz(F, reqntidz))
  reqntidz = 1;
else
  specified = true;

if (specified)
  O << ".reqntid " << reqntidx << ", " << reqntidy << ", " << reqntidz
    << "\n";

// If the NVVM IR has some of maxntid* specified, then output
// the maxntid directive, and set the unspecified ones to 1.
// If none of maxntid* is specified, don't output maxntid directive.
unsigned maxntidx, maxntidy, maxntidz;
specified = false;
if (!getMaxNTIDx(F, maxntidx))
  maxntidx = 1;
else
  specified = true;
if (!getMaxNTIDy(F, maxntidy))
  maxntidy = 1;
else
  specified = true;
if (!getMaxNTIDz(F, maxntidz))
  maxntidz = 1;
else
  specified = true;

if (specified)
  O << ".maxntid " << maxntidx << ", " << maxntidy << ", " << maxntidz
    << "\n";

unsigned mincta;
if (getMinCTASm(F, mincta))
  O << ".minnctapersm " << mincta << "\n";

unsigned maxnreg;
if (getMaxNReg(F, maxnreg))
  O << ".maxnreg " << maxnreg << "\n";
575}

577std::string
578NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const {
const TargetRegisterClass *RC = MRI->getRegClass(Reg);

std::string Name;
raw_string_ostream NameStr(Name);

VRegRCMap::const_iterator I = VRegMapping.find(RC);
assert(I != VRegMapping.end() && "Bad register class")(static_cast<void> (0));
const DenseMap<unsigned, unsigned> &RegMap = I->second;

VRegMap::const_iterator VI = RegMap.find(Reg);
assert(VI != RegMap.end() && "Bad virtual register")(static_cast<void> (0));
unsigned MappedVR = VI->second;

NameStr << getNVPTXRegClassStr(RC) << MappedVR;

NameStr.flush();
return Name;
596}

598void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr,
                                        raw_ostream &O) {
O << getVirtualRegisterName(vr);
601}

603void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) {
emitLinkageDirective(F, O);
if (isKernelFunction(*F))
10
←
Assuming the condition is false→
11
←
Taking false branch→
  O << ".entry ";
else
  O << ".func ";
printReturnValStr(F, O);
getSymbol(F)->print(O, MAI);
O << "\n";
emitFunctionParamList(F, O);
12
←
Calling 'NVPTXAsmPrinter::emitFunctionParamList'→
O << ";\n";
614}

616static bool usedInGlobalVarDef(const Constant *C) {
if (!C)
  return false;

if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(C)) {
  return GV->getName() != "llvm.used";
}

for (const User *U : C->users())
  if (const Constant *C = dyn_cast<Constant>(U))
    if (usedInGlobalVarDef(C))
      return true;

return false;
630}

632static bool usedInOneFunc(const User *U, Function const *&oneFunc) {
if (const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
  if (othergv->getName() == "llvm.used")
    return true;
}

if (const Instruction *instr = dyn_cast<Instruction>(U)) {
  if (instr->getParent() && instr->getParent()->getParent()) {
    const Function *curFunc = instr->getParent()->getParent();
    if (oneFunc && (curFunc != oneFunc))
      return false;
    oneFunc = curFunc;
    return true;
  } else
    return false;
}

for (const User *UU : U->users())
  if (!usedInOneFunc(UU, oneFunc))
    return false;

return true;
654}

656/* Find out if a global variable can be demoted to local scope.
* Currently, this is valid for CUDA shared variables, which have local
* scope and global lifetime. So the conditions to check are :
* 1. Is the global variable in shared address space?
* 2. Does it have internal linkage?
* 3. Is the global variable referenced only in one function?
*/
663static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) {
if (!gv->hasInternalLinkage())
  return false;
PointerType *Pty = gv->getType();
if (Pty->getAddressSpace() != ADDRESS_SPACE_SHARED)
  return false;

const Function *oneFunc = nullptr;

bool flag = usedInOneFunc(gv, oneFunc);
if (!flag)
  return false;
if (!oneFunc)
  return false;
f = oneFunc;
return true;
679}

681static bool useFuncSeen(const Constant *C,
                      DenseMap<const Function *, bool> &seenMap) {
for (const User *U : C->users()) {
  if (const Constant *cu = dyn_cast<Constant>(U)) {
    if (useFuncSeen(cu, seenMap))
      return true;
  } else if (const Instruction *I = dyn_cast<Instruction>(U)) {
    const BasicBlock *bb = I->getParent();
    if (!bb)
      continue;
    const Function *caller = bb->getParent();
    if (!caller)
      continue;
    if (seenMap.find(caller) != seenMap.end())
      return true;
  }
}
return false;
699}

701void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {
DenseMap<const Function *, bool> seenMap;
for (Module::const_iterator FI = M.begin(), FE = M.end(); FI != FE; ++FI) {
6
←
Loop condition is true.  Entering loop body→
  const Function *F = &*FI;

  if (F->getAttributes().hasFnAttr("nvptx-libcall-callee")) {
7
←
Assuming the condition is true→
8
←
Taking true branch→
    emitDeclaration(F, O);
9
←
Calling 'NVPTXAsmPrinter::emitDeclaration'→
    continue;
  }

  if (F->isDeclaration()) {
    if (F->use_empty())
      continue;
    if (F->getIntrinsicID())
      continue;
    emitDeclaration(F, O);
    continue;
  }
  for (const User *U : F->users()) {
    if (const Constant *C = dyn_cast<Constant>(U)) {
      if (usedInGlobalVarDef(C)) {
        // The use is in the initialization of a global variable
        // that is a function pointer, so print a declaration
        // for the original function
        emitDeclaration(F, O);
        break;
      }
      // Emit a declaration of this function if the function that
      // uses this constant expr has already been seen.
      if (useFuncSeen(C, seenMap)) {
        emitDeclaration(F, O);
        break;
      }
    }

    if (!isa<Instruction>(U))
      continue;
    const Instruction *instr = cast<Instruction>(U);
    const BasicBlock *bb = instr->getParent();
    if (!bb)
      continue;
    const Function *caller = bb->getParent();
    if (!caller)
      continue;

    // If a caller has already been seen, then the caller is
    // appearing in the module before the callee. so print out
    // a declaration for the callee.
    if (seenMap.find(caller) != seenMap.end()) {
      emitDeclaration(F, O);
      break;
    }
  }
  seenMap[F] = true;
}
756}

758static bool isEmptyXXStructor(GlobalVariable *GV) {
if (!GV) return true;
const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());
if (!InitList) return true;  // Not an array; we don't know how to parse.
return InitList->getNumOperands() == 0;
763}

765void NVPTXAsmPrinter::emitStartOfAsmFile(Module &M) {
// Construct a default subtarget off of the TargetMachine defaults. The
// rest of NVPTX isn't friendly to change subtargets per function and
// so the default TargetMachine will have all of the options.
const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
const auto* STI = static_cast<const NVPTXSubtarget*>(NTM.getSubtargetImpl());
SmallString<128> Str1;
raw_svector_ostream OS1(Str1);

// Emit header before any dwarf directives are emitted below.
emitHeader(M, OS1, *STI);
OutStreamer->emitRawText(OS1.str());
777}

779bool NVPTXAsmPrinter::doInitialization(Module &M) {
if (M.alias_size()) {
  report_fatal_error("Module has aliases, which NVPTX does not support.");
  return true; // error
}
if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_ctors"))) {
  report_fatal_error(
      "Module has a nontrivial global ctor, which NVPTX does not support.");
  return true;  // error
}
if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_dtors"))) {
  report_fatal_error(
      "Module has a nontrivial global dtor, which NVPTX does not support.");
  return true;  // error
}

// We need to call the parent's one explicitly.
bool Result = AsmPrinter::doInitialization(M);

GlobalsEmitted = false;

return Result;
801}

803void NVPTXAsmPrinter::emitGlobals(const Module &M) {
SmallString<128> Str2;
raw_svector_ostream OS2(Str2);

emitDeclarations(M, OS2);
5
←
Calling 'NVPTXAsmPrinter::emitDeclarations'→

// As ptxas does not support forward references of globals, we need to first
// sort the list of module-level globals in def-use order. We visit each
// global variable in order, and ensure that we emit it *after* its dependent
// globals. We use a little extra memory maintaining both a set and a list to
// have fast searches while maintaining a strict ordering.
SmallVector<const GlobalVariable *, 8> Globals;
DenseSet<const GlobalVariable *> GVVisited;
DenseSet<const GlobalVariable *> GVVisiting;

// Visit each global variable, in order
for (const GlobalVariable &I : M.globals())
  VisitGlobalVariableForEmission(&I, Globals, GVVisited, GVVisiting);

assert(GVVisited.size() == M.getGlobalList().size() &&(static_cast<void> (0))
       "Missed a global variable")(static_cast<void> (0));
assert(GVVisiting.size() == 0 && "Did not fully process a global variable")(static_cast<void> (0));

// Print out module-level global variables in proper order
for (unsigned i = 0, e = Globals.size(); i != e; ++i)
  printModuleLevelGV(Globals[i], OS2);

OS2 << '\n';

OutStreamer->emitRawText(OS2.str());
833}

835void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
                               const NVPTXSubtarget &STI) {
O << "//\n";
O << "// Generated by LLVM NVPTX Back-End\n";
O << "//\n";
O << "\n";

unsigned PTXVersion = STI.getPTXVersion();
O << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n";

O << ".target ";
O << STI.getTargetName();

const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
if (NTM.getDrvInterface() == NVPTX::NVCL)
  O << ", texmode_independent";

bool HasFullDebugInfo = false;
for (DICompileUnit *CU : M.debug_compile_units()) {
  switch(CU->getEmissionKind()) {
  case DICompileUnit::NoDebug:
  case DICompileUnit::DebugDirectivesOnly:
    break;
  case DICompileUnit::LineTablesOnly:
  case DICompileUnit::FullDebug:
    HasFullDebugInfo = true;
    break;
  }
  if (HasFullDebugInfo)
    break;
}
if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo)
  O << ", debug";

O << "\n";

O << ".address_size ";
if (NTM.is64Bit())
  O << "64";
else
  O << "32";
O << "\n";

O << "\n";
879}

881bool NVPTXAsmPrinter::doFinalization(Module &M) {
bool HasDebugInfo = MMI && MMI->hasDebugInfo();
1
Assuming field 'MMI' is null→

// If we did not emit any functions, then the global declarations have not
// yet been emitted.
if (!GlobalsEmitted) {
2
←
Assuming field 'GlobalsEmitted' is false→
3
←
Taking true branch→
  emitGlobals(M);
4
←
Calling 'NVPTXAsmPrinter::emitGlobals'→
  GlobalsEmitted = true;
}

// XXX Temproarily remove global variables so that doFinalization() will not
// emit them again (global variables are emitted at beginning).

Module::GlobalListType &global_list = M.getGlobalList();
int i, n = global_list.size();
GlobalVariable **gv_array = new GlobalVariable *[n];

// first, back-up GlobalVariable in gv_array
i = 0;
for (Module::global_iterator I = global_list.begin(), E = global_list.end();
     I != E; ++I)
  gv_array[i++] = &*I;

// second, empty global_list
while (!global_list.empty())
  global_list.remove(global_list.begin());

// call doFinalization
bool ret = AsmPrinter::doFinalization(M);

// now we restore global variables
for (i = 0; i < n; i++)
  global_list.insert(global_list.end(), gv_array[i]);

clearAnnotationCache(&M);

delete[] gv_array;
// Close the last emitted section
if (HasDebugInfo) {
  static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
      ->closeLastSection();
  // Emit empty .debug_loc section for better support of the empty files.
  OutStreamer->emitRawText("\t.section\t.debug_loc\t{\t}");
}

// Output last DWARF .file directives, if any.
static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
    ->outputDwarfFileDirectives();

return ret;

//bool Result = AsmPrinter::doFinalization(M);
// Instead of calling the parents doFinalization, we may
// clone parents doFinalization and customize here.
// Currently, we if NVISA out the EmitGlobals() in
// parent's doFinalization, which is too intrusive.
//
// Same for the doInitialization.
//return Result;
940}

942// This function emits appropriate linkage directives for
943// functions and global variables.
944//
945// extern function declaration            -> .extern
946// extern function definition             -> .visible
947// external global variable with init     -> .visible
948// external without init                  -> .extern
949// appending                              -> not allowed, assert.
950// for any linkage other than
951// internal, private, linker_private,
952// linker_private_weak, linker_private_weak_def_auto,
953// we emit                                -> .weak.

955void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V,
                                         raw_ostream &O) {
if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA) {
  if (V->hasExternalLinkage()) {
    if (isa<GlobalVariable>(V)) {
      const GlobalVariable *GVar = cast<GlobalVariable>(V);
      if (GVar) {
        if (GVar->hasInitializer())
          O << ".visible ";
        else
          O << ".extern ";
      }
    } else if (V->isDeclaration())
      O << ".extern ";
    else
      O << ".visible ";
  } else if (V->hasAppendingLinkage()) {
    std::string msg;
    msg.append("Error: ");
    msg.append("Symbol ");
    if (V->hasName())
      msg.append(std::string(V->getName()));
    msg.append("has unsupported appending linkage type");
    llvm_unreachable(msg.c_str())__builtin_unreachable();
  } else if (!V->hasInternalLinkage() &&
             !V->hasPrivateLinkage()) {
    O << ".weak ";
  }
}
984}

986void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
                                       raw_ostream &O,
                                       bool processDemoted) {
// Skip meta data
if (GVar->hasSection()) {
  if (GVar->getSection() == "llvm.metadata")
    return;
}

// Skip LLVM intrinsic global variables
if (GVar->getName().startswith("llvm.") ||
    GVar->getName().startswith("nvvm."))
  return;

const DataLayout &DL = getDataLayout();

// GlobalVariables are always constant pointers themselves.
PointerType *PTy = GVar->getType();
Type *ETy = GVar->getValueType();

if (GVar->hasExternalLinkage()) {
  if (GVar->hasInitializer())
    O << ".visible ";
  else
    O << ".extern ";
} else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() ||
           GVar->hasAvailableExternallyLinkage() ||
           GVar->hasCommonLinkage()) {
  O << ".weak ";
}

if (isTexture(*GVar)) {
  O << ".global .texref " << getTextureName(*GVar) << ";\n";
  return;
}

if (isSurface(*GVar)) {
  O << ".global .surfref " << getSurfaceName(*GVar) << ";\n";
  return;
}

if (GVar->isDeclaration()) {
  // (extern) declarations, no definition or initializer
  // Currently the only known declaration is for an automatic __local
  // (.shared) promoted to global.
  emitPTXGlobalVariable(GVar, O);
  O << ";\n";
  return;
}

if (isSampler(*GVar)) {
  O << ".global .samplerref " << getSamplerName(*GVar);

  const Constant *Initializer = nullptr;
  if (GVar->hasInitializer())
    Initializer = GVar->getInitializer();
  const ConstantInt *CI = nullptr;
  if (Initializer)
    CI = dyn_cast<ConstantInt>(Initializer);
  if (CI) {
    unsigned sample = CI->getZExtValue();

    O << " = { ";

    for (int i = 0,
             addr = ((sample & __CLK_ADDRESS_MASK) >> __CLK_ADDRESS_BASE);
         i < 3; i++) {
      O << "addr_mode_" << i << " = ";
      switch (addr) {
      case 0:
        O << "wrap";
        break;
      case 1:
        O << "clamp_to_border";
        break;
      case 2:
        O << "clamp_to_edge";
        break;
      case 3:
        O << "wrap";
        break;
      case 4:
        O << "mirror";
        break;
      }
      O << ", ";
    }
    O << "filter_mode = ";
    switch ((sample & __CLK_FILTER_MASK) >> __CLK_FILTER_BASE) {
    case 0:
      O << "nearest";
      break;
    case 1:
      O << "linear";
      break;
    case 2:
      llvm_unreachable("Anisotropic filtering is not supported")__builtin_unreachable();
    default:
      O << "nearest";
      break;
    }
    if (!((sample & __CLK_NORMALIZED_MASK) >> __CLK_NORMALIZED_BASE)) {
      O << ", force_unnormalized_coords = 1";
    }
    O << " }";
  }

  O << ";\n";
  return;
}

if (GVar->hasPrivateLinkage()) {
  if (strncmp(GVar->getName().data(), "unrollpragma", 12) == 0)
    return;

  // FIXME - need better way (e.g. Metadata) to avoid generating this global
  if (strncmp(GVar->getName().data(), "filename", 8) == 0)
    return;
  if (GVar->use_empty())
    return;
}

const Function *demotedFunc = nullptr;
if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) {
  O << "// " << GVar->getName() << " has been demoted\n";
  if (localDecls.find(demotedFunc) != localDecls.end())
    localDecls[demotedFunc].push_back(GVar);
  else {
    std::vector<const GlobalVariable *> temp;
    temp.push_back(GVar);
    localDecls[demotedFunc] = temp;
  }
  return;
}

O << ".";
emitPTXAddressSpace(PTy->getAddressSpace(), O);

if (isManaged(*GVar)) {
  O << " .attribute(.managed)";
}

if (GVar->getAlignment() == 0)
  O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
else
  O << " .align " << GVar->getAlignment();

if (ETy->isFloatingPointTy() || ETy->isPointerTy() ||
    (ETy->isIntegerTy() && ETy->getScalarSizeInBits() <= 64)) {
  O << " .";
  // Special case: ABI requires that we use .u8 for predicates
  if (ETy->isIntegerTy(1))
    O << "u8";
  else
    O << getPTXFundamentalTypeStr(ETy, false);
  O << " ";
  getSymbol(GVar)->print(O, MAI);

  // Ptx allows variable initilization only for constant and global state
  // spaces.
  if (GVar->hasInitializer()) {
    if ((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
        (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) {
      const Constant *Initializer = GVar->getInitializer();
      // 'undef' is treated as there is no value specified.
      if (!Initializer->isNullValue() && !isa<UndefValue>(Initializer)) {
        O << " = ";
        printScalarConstant(Initializer, O);
      }
    } else {
      // The frontend adds zero-initializer to device and constant variables
      // that don't have an initial value, and UndefValue to shared
      // variables, so skip warning for this case.
      if (!GVar->getInitializer()->isNullValue() &&
          !isa<UndefValue>(GVar->getInitializer())) {
        report_fatal_error("initial value of '" + GVar->getName() +
                           "' is not allowed in addrspace(" +
                           Twine(PTy->getAddressSpace()) + ")");
      }
    }
  }
} else {
  unsigned int ElementSize = 0;

  // Although PTX has direct support for struct type and array type and
  // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
  // targets that support these high level field accesses. Structs, arrays
  // and vectors are lowered into arrays of bytes.
  switch (ETy->getTypeID()) {
  case Type::IntegerTyID: // Integers larger than 64 bits
  case Type::StructTyID:
  case Type::ArrayTyID:
  case Type::FixedVectorTyID:
    ElementSize = DL.getTypeStoreSize(ETy);
    // Ptx allows variable initilization only for constant and
    // global state spaces.
    if (((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
         (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) &&
        GVar->hasInitializer()) {
      const Constant *Initializer = GVar->getInitializer();
      if (!isa<UndefValue>(Initializer) && !Initializer->isNullValue()) {
        AggBuffer aggBuffer(ElementSize, O, *this);
        bufferAggregateConstant(Initializer, &aggBuffer);
        if (aggBuffer.numSymbols) {
          if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit()) {
            O << " .u64 ";
            getSymbol(GVar)->print(O, MAI);
            O << "[";
            O << ElementSize / 8;
          } else {
            O << " .u32 ";
            getSymbol(GVar)->print(O, MAI);
            O << "[";
            O << ElementSize / 4;
          }
          O << "]";
        } else {
          O << " .b8 ";
          getSymbol(GVar)->print(O, MAI);
          O << "[";
          O << ElementSize;
          O << "]";
        }
        O << " = {";
        aggBuffer.print();
        O << "}";
      } else {
        O << " .b8 ";
        getSymbol(GVar)->print(O, MAI);
        if (ElementSize) {
          O << "[";
          O << ElementSize;
          O << "]";
        }
      }
    } else {
      O << " .b8 ";
      getSymbol(GVar)->print(O, MAI);
      if (ElementSize) {
        O << "[";
        O << ElementSize;
        O << "]";
      }
    }
    break;
  default:
    llvm_unreachable("type not supported yet")__builtin_unreachable();
  }
}
O << ";\n";
1236}

1238void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) {
if (localDecls.find(f) == localDecls.end())
  return;

std::vector<const GlobalVariable *> &gvars = localDecls[f];

for (unsigned i = 0, e = gvars.size(); i != e; ++i) {
  O << "\t// demoted variable\n\t";
  printModuleLevelGV(gvars[i], O, true);
}
1248}

1250void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
                                        raw_ostream &O) const {
switch (AddressSpace) {
case ADDRESS_SPACE_LOCAL:
  O << "local";
  break;
case ADDRESS_SPACE_GLOBAL:
  O << "global";
  break;
case ADDRESS_SPACE_CONST:
  O << "const";
  break;
case ADDRESS_SPACE_SHARED:
  O << "shared";
  break;
default:
  report_fatal_error("Bad address space found while emitting PTX: " +
                     llvm::Twine(AddressSpace));
  break;
}
1270}

1272std::string
1273NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {
switch (Ty->getTypeID()) {
case Type::IntegerTyID: {
  unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
  if (NumBits == 1)
    return "pred";
  else if (NumBits <= 64) {
    std::string name = "u";
    return name + utostr(NumBits);
  } else {
    llvm_unreachable("Integer too large")__builtin_unreachable();
    break;
  }
  break;
}
case Type::HalfTyID:
  // fp16 is stored as .b16 for compatibility with pre-sm_53 PTX assembly.
  return "b16";
case Type::FloatTyID:
  return "f32";
case Type::DoubleTyID:
  return "f64";
case Type::PointerTyID:
  if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit())
    if (useB4PTR)
      return "b64";
    else
      return "u64";
  else if (useB4PTR)
    return "b32";
  else
    return "u32";
default:
  break;
}
llvm_unreachable("unexpected type")__builtin_unreachable();
1309}

1311void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,
                                          raw_ostream &O) {
const DataLayout &DL = getDataLayout();

// GlobalVariables are always constant pointers themselves.
Type *ETy = GVar->getValueType();

O << ".";
emitPTXAddressSpace(GVar->getType()->getAddressSpace(), O);
if (GVar->getAlignment() == 0)
  O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
else
  O << " .align " << GVar->getAlignment();

// Special case for i128
if (ETy->isIntegerTy(128)) {
  O << " .b8 ";
  getSymbol(GVar)->print(O, MAI);
  O << "[16]";
  return;
}

if (ETy->isFloatingPointTy() || ETy->isIntOrPtrTy()) {
  O << " .";
  O << getPTXFundamentalTypeStr(ETy);
  O << " ";
  getSymbol(GVar)->print(O, MAI);
  return;
}

int64_t ElementSize = 0;

// Although PTX has direct support for struct type and array type and LLVM IR
// is very similar to PTX, the LLVM CodeGen does not support for targets that
// support these high level field accesses. Structs and arrays are lowered
// into arrays of bytes.
switch (ETy->getTypeID()) {
case Type::StructTyID:
case Type::ArrayTyID:
case Type::FixedVectorTyID:
  ElementSize = DL.getTypeStoreSize(ETy);
  O << " .b8 ";
  getSymbol(GVar)->print(O, MAI);
  O << "[";
  if (ElementSize) {
    O << ElementSize;
  }
  O << "]";
  break;
default:
  llvm_unreachable("type not supported yet")__builtin_unreachable();
}
1363}

1365static unsigned int getOpenCLAlignment(const DataLayout &DL, Type *Ty) {
if (Ty->isSingleValueType())
  return DL.getPrefTypeAlignment(Ty);

auto *ATy = dyn_cast<ArrayType>(Ty);
if (ATy)
  return getOpenCLAlignment(DL, ATy->getElementType());

auto *STy = dyn_cast<StructType>(Ty);
if (STy) {
  unsigned int alignStruct = 1;
  // Go through each element of the struct and find the
  // largest alignment.
  for (unsigned i = 0, e = STy->getNumElements(); i != e; i++) {
    Type *ETy = STy->getElementType(i);
    unsigned int align = getOpenCLAlignment(DL, ETy);
    if (align > alignStruct)
      alignStruct = align;
  }
  return alignStruct;
}

auto *FTy = dyn_cast<FunctionType>(Ty);
if (FTy)
  return DL.getPointerPrefAlignment().value();
return DL.getPrefTypeAlignment(Ty);
1391}

1393void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I,
                                   int paramIndex, raw_ostream &O) {
getSymbol(I->getParent())->print(O, MAI);
O << "_param_" << paramIndex;
1397}

1399void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
const DataLayout &DL = getDataLayout();
const AttributeList &PAL = F->getAttributes();
const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
const TargetLowering *TLI = STI.getTargetLowering();
Function::const_arg_iterator I, E;
unsigned paramIndex = 0;
bool first = true;
bool isKernelFunc = isKernelFunction(*F);
bool isABI = (STI.getSmVersion() >= 20);
13
←
Assuming the condition is false→
bool hasImageHandles = STI.hasImageHandles();
MVT thePointerTy = TLI->getPointerTy(DL);

if (F->arg_empty()) {
14
←
Assuming the condition is false→
15
←
Taking false branch→
  O << "()\n";
  return;
}

O << "(\n";

for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) {
16
←
Assuming 'I' is not equal to 'E'→
17
←
Loop condition is true.  Entering loop body→
  Type *Ty = I->getType();

  if (!first17.1
'first' is true
)
18
←
Taking false branch→
    O << ",\n";

  first = false;

  // Handle image/sampler parameters
  if (isKernelFunction(*F)) {
19
←
Assuming the condition is false→
20
←
Taking false branch→
    if (isSampler(*I) || isImage(*I)) {
      if (isImage(*I)) {
        std::string sname = std::string(I->getName());
        if (isImageWriteOnly(*I) || isImageReadWrite(*I)) {
          if (hasImageHandles)
            O << "\t.param .u64 .ptr .surfref ";
          else
            O << "\t.param .surfref ";
          CurrentFnSym->print(O, MAI);
          O << "_param_" << paramIndex;
        }
        else { // Default image is read_only
          if (hasImageHandles)
            O << "\t.param .u64 .ptr .texref ";
          else
            O << "\t.param .texref ";
          CurrentFnSym->print(O, MAI);
          O << "_param_" << paramIndex;
        }
      } else {
        if (hasImageHandles)
          O << "\t.param .u64 .ptr .samplerref ";
        else
          O << "\t.param .samplerref ";
        CurrentFnSym->print(O, MAI);
        O << "_param_" << paramIndex;
      }
      continue;
    }
  }

  if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
21
←
Assuming the condition is false→
22
←
Taking false branch→
    if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
      // Just print .param .align <a> .b8 .param[size];
      // <a> = PAL.getparamalignment
      // size = typeallocsize of element type
      const Align align = DL.getValueOrABITypeAlignment(
          PAL.getParamAlignment(paramIndex), Ty);

      unsigned sz = DL.getTypeAllocSize(Ty);
      O << "\t.param .align " << align.value() << " .b8 ";
      printParamName(I, paramIndex, O);
      O << "[" << sz << "]";

      continue;
    }
    // Just a scalar
    auto *PTy = dyn_cast<PointerType>(Ty);
    if (isKernelFunc) {
      if (PTy) {
        // Special handling for pointer arguments to kernel
        O << "\t.param .u" << thePointerTy.getSizeInBits() << " ";

        if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() !=
            NVPTX::CUDA) {
          Type *ETy = PTy->getElementType();
          int addrSpace = PTy->getAddressSpace();
          switch (addrSpace) {
          default:
            O << ".ptr ";
            break;
          case ADDRESS_SPACE_CONST:
            O << ".ptr .const ";
            break;
          case ADDRESS_SPACE_SHARED:
            O << ".ptr .shared ";
            break;
          case ADDRESS_SPACE_GLOBAL:
            O << ".ptr .global ";
            break;
          }
          O << ".align " << (int)getOpenCLAlignment(DL, ETy) << " ";
        }
        printParamName(I, paramIndex, O);
        continue;
      }

      // non-pointer scalar to kernel func
      O << "\t.param .";
      // Special case: predicate operands become .u8 types
      if (Ty->isIntegerTy(1))
        O << "u8";
      else
        O << getPTXFundamentalTypeStr(Ty);
      O << " ";
      printParamName(I, paramIndex, O);
      continue;
    }
    // Non-kernel function, just print .param .b<size> for ABI
    // and .reg .b<size> for non-ABI
    unsigned sz = 0;
    if (isa<IntegerType>(Ty)) {
      sz = cast<IntegerType>(Ty)->getBitWidth();
      if (sz < 32)
        sz = 32;
    } else if (isa<PointerType>(Ty))
      sz = thePointerTy.getSizeInBits();
    else if (Ty->isHalfTy())
      // PTX ABI requires all scalar parameters to be at least 32
      // bits in size.  fp16 normally uses .b16 as its storage type
      // in PTX, so its size must be adjusted here, too.
      sz = 32;
    else
      sz = Ty->getPrimitiveSizeInBits();
    if (isABI)
      O << "\t.param .b" << sz << " ";
    else
      O << "\t.reg .b" << sz << " ";
    printParamName(I, paramIndex, O);
    continue;
  }

  // param has byVal attribute. So should be a pointer
  auto *PTy = dyn_cast<PointerType>(Ty);
23
←
Assuming 'Ty' is not a 'PointerType'→
24
←
'PTy' initialized to a null pointer value→
  assert(PTy && "Param with byval attribute should be a pointer type")(static_cast<void> (0));
  Type *ETy = PTy->getElementType();
25
←
Called C++ object pointer is null

  if (isABI || isKernelFunc) {
    // Just print .param .align <a> .b8 .param[size];
    // <a> = PAL.getparamalignment
    // size = typeallocsize of element type
    Align align =
        DL.getValueOrABITypeAlignment(PAL.getParamAlignment(paramIndex), ETy);
    // Work around a bug in ptxas. When PTX code takes address of
    // byval parameter with alignment < 4, ptxas generates code to
    // spill argument into memory. Alas on sm_50+ ptxas generates
    // SASS code that fails with misaligned access. To work around
    // the problem, make sure that we align byval parameters by at
    // least 4. Matching change must be made in LowerCall() where we
    // prepare parameters for the call.
    //
    // TODO: this will need to be undone when we get to support multi-TU
    // device-side compilation as it breaks ABI compatibility with nvcc.
    // Hopefully ptxas bug is fixed by then.
    if (!isKernelFunc && align < Align(4))
      align = Align(4);
    unsigned sz = DL.getTypeAllocSize(ETy);
    O << "\t.param .align " << align.value() << " .b8 ";
    printParamName(I, paramIndex, O);
    O << "[" << sz << "]";
    continue;
  } else {
    // Split the ETy into constituent parts and
    // print .param .b<size> <name> for each part.
    // Further, if a part is vector, print the above for
    // each vector element.
    SmallVector<EVT, 16> vtparts;
    ComputeValueVTs(*TLI, DL, ETy, vtparts);
    for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
      unsigned elems = 1;
      EVT elemtype = vtparts[i];
      if (vtparts[i].isVector()) {
        elems = vtparts[i].getVectorNumElements();
        elemtype = vtparts[i].getVectorElementType();
      }

      for (unsigned j = 0, je = elems; j != je; ++j) {
        unsigned sz = elemtype.getSizeInBits();
        if (elemtype.isInteger() && (sz < 32))
          sz = 32;
        O << "\t.reg .b" << sz << " ";
        printParamName(I, paramIndex, O);
        if (j < je - 1)
          O << ",\n";
        ++paramIndex;
      }
      if (i < e - 1)
        O << ",\n";
    }
    --paramIndex;
    continue;
  }
}

O << "\n)\n";
1604}

1606void NVPTXAsmPrinter::emitFunctionParamList(const MachineFunction &MF,
                                          raw_ostream &O) {
const Function &F = MF.getFunction();
emitFunctionParamList(&F, O);
1610}

1612void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
  const MachineFunction &MF) {
SmallString<128> Str;
raw_svector_ostream O(Str);

// Map the global virtual register number to a register class specific
// virtual register number starting from 1 with that class.
const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo();
//unsigned numRegClasses = TRI->getNumRegClasses();

// Emit the Fake Stack Object
const MachineFrameInfo &MFI = MF.getFrameInfo();
int NumBytes = (int) MFI.getStackSize();
if (NumBytes) {
  O << "\t.local .align " << MFI.getMaxAlign().value() << " .b8 \t"
    << DEPOTNAME"__local_depot" << getFunctionNumber() << "[" << NumBytes << "];\n";
  if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
    O << "\t.reg .b64 \t%SP;\n";
    O << "\t.reg .b64 \t%SPL;\n";
  } else {
    O << "\t.reg .b32 \t%SP;\n";
    O << "\t.reg .b32 \t%SPL;\n";
  }
}

// Go through all virtual registers to establish the mapping between the
// global virtual
// register number and the per class virtual register number.
// We use the per class virtual register number in the ptx output.
unsigned int numVRs = MRI->getNumVirtRegs();
for (unsigned i = 0; i < numVRs; i++) {
  unsigned int vr = Register::index2VirtReg(i);
  const TargetRegisterClass *RC = MRI->getRegClass(vr);
  DenseMap<unsigned, unsigned> &regmap = VRegMapping[RC];
  int n = regmap.size();
  regmap.insert(std::make_pair(vr, n + 1));
}

// Emit register declarations
// @TODO: Extract out the real register usage
// O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n";
// O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n";
// O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n";
// O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n";
// O << "\t.reg .s64 %rd<" << NVPTXNumRegisters << ">;\n";
// O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n";
// O << "\t.reg .f64 %fd<" << NVPTXNumRegisters << ">;\n";

// Emit declaration of the virtual registers or 'physical' registers for
// each register class
for (unsigned i=0; i< TRI->getNumRegClasses(); i++) {
  const TargetRegisterClass *RC = TRI->getRegClass(i);
  DenseMap<unsigned, unsigned> &regmap = VRegMapping[RC];
  std::string rcname = getNVPTXRegClassName(RC);
  std::string rcStr = getNVPTXRegClassStr(RC);
  int n = regmap.size();

  // Only declare those registers that may be used.
  if (n) {
     O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1)
       << ">;\n";
  }
}

OutStreamer->emitRawText(O.str());
1677}

1679void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) {
APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
bool ignored;
unsigned int numHex;
const char *lead;

if (Fp->getType()->getTypeID() == Type::FloatTyID) {
  numHex = 8;
  lead = "0f";
  APF.convert(APFloat::IEEEsingle(), APFloat::rmNearestTiesToEven, &ignored);
} else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
  numHex = 16;
  lead = "0d";
  APF.convert(APFloat::IEEEdouble(), APFloat::rmNearestTiesToEven, &ignored);
} else
  llvm_unreachable("unsupported fp type")__builtin_unreachable();

APInt API = APF.bitcastToAPInt();
O << lead << format_hex_no_prefix(API.getZExtValue(), numHex, /*Upper=*/true);
1698}

1700void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
  O << CI->getValue();
  return;
}
if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
  printFPConstant(CFP, O);
  return;
}
if (isa<ConstantPointerNull>(CPV)) {
  O << "0";
  return;
}
if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
  bool IsNonGenericPointer = false;
  if (GVar->getType()->getAddressSpace() != 0) {
    IsNonGenericPointer = true;
  }
  if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
    O << "generic(";
    getSymbol(GVar)->print(O, MAI);
    O << ")";
  } else {
    getSymbol(GVar)->print(O, MAI);
  }
  return;
}
if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
  const Value *v = Cexpr->stripPointerCasts();
  PointerType *PTy = dyn_cast<PointerType>(Cexpr->getType());
  bool IsNonGenericPointer = false;
  if (PTy && PTy->getAddressSpace() != 0) {
    IsNonGenericPointer = true;
  }
  if (const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
    if (EmitGeneric && !isa<Function>(v) && !IsNonGenericPointer) {
      O << "generic(";
      getSymbol(GVar)->print(O, MAI);
      O << ")";
    } else {
      getSymbol(GVar)->print(O, MAI);
    }
    return;
  } else {
    lowerConstant(CPV)->print(O, MAI);
    return;
  }
}
llvm_unreachable("Not scalar type found in printScalarConstant()")__builtin_unreachable();
1749}

1751// These utility functions assure we get the right sequence of bytes for a given
1752// type even for big-endian machines
1753template <typename T> static void ConvertIntToBytes(unsigned char *p, T val) {
int64_t vp = (int64_t)val;
for (unsigned i = 0; i < sizeof(T); ++i) {
  p[i] = (unsigned char)vp;
  vp >>= 8;
}
1759}
1760static void ConvertFloatToBytes(unsigned char *p, float val) {
int32_t *vp = (int32_t *)&val;
for (unsigned i = 0; i < sizeof(int32_t); ++i) {
  p[i] = (unsigned char)*vp;
  *vp >>= 8;
}
1766}
1767static void ConvertDoubleToBytes(unsigned char *p, double val) {
int64_t *vp = (int64_t *)&val;
for (unsigned i = 0; i < sizeof(int64_t); ++i) {
  p[i] = (unsigned char)*vp;
  *vp >>= 8;
}
1773}

1775void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes,
                                 AggBuffer *aggBuffer) {
const DataLayout &DL = getDataLayout();

if (isa<UndefValue>(CPV) || CPV->isNullValue()) {
  int s = DL.getTypeAllocSize(CPV->getType());
  if (s < Bytes)
    s = Bytes;
  aggBuffer->addZeros(s);
  return;
}

unsigned char ptr[8];
switch (CPV->getType()->getTypeID()) {

case Type::IntegerTyID: {
  Type *ETy = CPV->getType();
  if (ETy == Type::getInt8Ty(CPV->getContext())) {
    unsigned char c = (unsigned char)cast<ConstantInt>(CPV)->getZExtValue();
    ConvertIntToBytes<>(ptr, c);
    aggBuffer->addBytes(ptr, 1, Bytes);
  } else if (ETy == Type::getInt16Ty(CPV->getContext())) {
    short int16 = (short)cast<ConstantInt>(CPV)->getZExtValue();
    ConvertIntToBytes<>(ptr, int16);
    aggBuffer->addBytes(ptr, 2, Bytes);
  } else if (ETy == Type::getInt32Ty(CPV->getContext())) {
    if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
      int int32 = (int)(constInt->getZExtValue());
      ConvertIntToBytes<>(ptr, int32);
      aggBuffer->addBytes(ptr, 4, Bytes);
      break;
    } else if (const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
      if (const auto *constInt = dyn_cast<ConstantInt>(
              ConstantFoldConstant(Cexpr, DL))) {
        int int32 = (int)(constInt->getZExtValue());
        ConvertIntToBytes<>(ptr, int32);
        aggBuffer->addBytes(ptr, 4, Bytes);
        break;
      }
      if (Cexpr->getOpcode() == Instruction::PtrToInt) {
        Value *v = Cexpr->getOperand(0)->stripPointerCasts();
        aggBuffer->addSymbol(v, Cexpr->getOperand(0));
        aggBuffer->addZeros(4);
        break;
      }
    }
    llvm_unreachable("unsupported integer const type")__builtin_unreachable();
  } else if (ETy == Type::getInt64Ty(CPV->getContext())) {
    if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
      long long int64 = (long long)(constInt->getZExtValue());
      ConvertIntToBytes<>(ptr, int64);
      aggBuffer->addBytes(ptr, 8, Bytes);
      break;
    } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
      if (const auto *constInt = dyn_cast<ConstantInt>(
              ConstantFoldConstant(Cexpr, DL))) {
        long long int64 = (long long)(constInt->getZExtValue());
        ConvertIntToBytes<>(ptr, int64);
        aggBuffer->addBytes(ptr, 8, Bytes);
        break;
      }
      if (Cexpr->getOpcode() == Instruction::PtrToInt) {
        Value *v = Cexpr->getOperand(0)->stripPointerCasts();
        aggBuffer->addSymbol(v, Cexpr->getOperand(0));
        aggBuffer->addZeros(8);
        break;
      }
    }
    llvm_unreachable("unsupported integer const type")__builtin_unreachable();
  } else
    llvm_unreachable("unsupported integer const type")__builtin_unreachable();
  break;
}
case Type::HalfTyID:
case Type::FloatTyID:
case Type::DoubleTyID: {
  const auto *CFP = cast<ConstantFP>(CPV);
  Type *Ty = CFP->getType();
  if (Ty == Type::getHalfTy(CPV->getContext())) {
    APInt API = CFP->getValueAPF().bitcastToAPInt();
    uint16_t float16 = API.getLoBits(16).getZExtValue();
    ConvertIntToBytes<>(ptr, float16);
    aggBuffer->addBytes(ptr, 2, Bytes);
  } else if (Ty == Type::getFloatTy(CPV->getContext())) {
    float float32 = (float) CFP->getValueAPF().convertToFloat();
    ConvertFloatToBytes(ptr, float32);
    aggBuffer->addBytes(ptr, 4, Bytes);
  } else if (Ty == Type::getDoubleTy(CPV->getContext())) {
    double float64 = CFP->getValueAPF().convertToDouble();
    ConvertDoubleToBytes(ptr, float64);
    aggBuffer->addBytes(ptr, 8, Bytes);
  } else {
    llvm_unreachable("unsupported fp const type")__builtin_unreachable();
  }
  break;
}
case Type::PointerTyID: {
  if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
    aggBuffer->addSymbol(GVar, GVar);
  } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
    const Value *v = Cexpr->stripPointerCasts();
    aggBuffer->addSymbol(v, Cexpr);
  }
  unsigned int s = DL.getTypeAllocSize(CPV->getType());
  aggBuffer->addZeros(s);
  break;
}

case Type::ArrayTyID:
case Type::FixedVectorTyID:
case Type::StructTyID: {
  if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
    int ElementSize = DL.getTypeAllocSize(CPV->getType());
    bufferAggregateConstant(CPV, aggBuffer);
    if (Bytes > ElementSize)
      aggBuffer->addZeros(Bytes - ElementSize);
  } else if (isa<ConstantAggregateZero>(CPV))
    aggBuffer->addZeros(Bytes);
  else
    llvm_unreachable("Unexpected Constant type")__builtin_unreachable();
  break;
}

default:
  llvm_unreachable("unsupported type")__builtin_unreachable();
}
1901}

1903void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV,
                                            AggBuffer *aggBuffer) {
const DataLayout &DL = getDataLayout();
int Bytes;

// Integers of arbitrary width
if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
  APInt Val = CI->getValue();
  for (unsigned I = 0, E = DL.getTypeAllocSize(CPV->getType()); I < E; ++I) {
    uint8_t Byte = Val.getLoBits(8).getZExtValue();
    aggBuffer->addBytes(&Byte, 1, 1);
    Val.lshrInPlace(8);
  }
  return;
}

// Old constants
if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
  if (CPV->getNumOperands())
    for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i)
      bufferLEByte(cast<Constant>(CPV->getOperand(i)), 0, aggBuffer);
  return;
}

if (const ConstantDataSequential *CDS =
        dyn_cast<ConstantDataSequential>(CPV)) {
  if (CDS->getNumElements())
    for (unsigned i = 0; i < CDS->getNumElements(); ++i)
      bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
                   aggBuffer);
  return;
}

if (isa<ConstantStruct>(CPV)) {
  if (CPV->getNumOperands()) {
    StructType *ST = cast<StructType>(CPV->getType());
    for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) {
      if (i == (e - 1))
        Bytes = DL.getStructLayout(ST)->getElementOffset(0) +
                DL.getTypeAllocSize(ST) -
                DL.getStructLayout(ST)->getElementOffset(i);
      else
        Bytes = DL.getStructLayout(ST)->getElementOffset(i + 1) -
                DL.getStructLayout(ST)->getElementOffset(i);
      bufferLEByte(cast<Constant>(CPV->getOperand(i)), Bytes, aggBuffer);
    }
  }
  return;
}
llvm_unreachable("unsupported constant type in printAggregateConstant()")__builtin_unreachable();
1953}

1955/// lowerConstantForGV - Return an MCExpr for the given Constant.  This is mostly
1956/// a copy from AsmPrinter::lowerConstant, except customized to only handle
1957/// expressions that are representable in PTX and create
1958/// NVPTXGenericMCSymbolRefExpr nodes for addrspacecast instructions.
1959const MCExpr *
1960NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV, bool ProcessingGeneric) {
MCContext &Ctx = OutContext;

if (CV->isNullValue() || isa<UndefValue>(CV))
  return MCConstantExpr::create(0, Ctx);

if (const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
  return MCConstantExpr::create(CI->getZExtValue(), Ctx);

if (const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
  const MCSymbolRefExpr *Expr =
    MCSymbolRefExpr::create(getSymbol(GV), Ctx);
  if (ProcessingGeneric) {
    return NVPTXGenericMCSymbolRefExpr::create(Expr, Ctx);
  } else {
    return Expr;
  }
}

const ConstantExpr *CE = dyn_cast<ConstantExpr>(CV);
if (!CE) {
  llvm_unreachable("Unknown constant value to lower!")__builtin_unreachable();
}

switch (CE->getOpcode()) {
default: {
  // If the code isn't optimized, there may be outstanding folding
  // opportunities. Attempt to fold the expression using DataLayout as a
  // last resort before giving up.
  Constant *C = ConstantFoldConstant(CE, getDataLayout());
  if (C != CE)
    return lowerConstantForGV(C, ProcessingGeneric);

  // Otherwise report the problem to the user.
  std::string S;
  raw_string_ostream OS(S);
  OS << "Unsupported expression in static initializer: ";
  CE->printAsOperand(OS, /*PrintType=*/false,
                 !MF ? nullptr : MF->getFunction().getParent());
  report_fatal_error(OS.str());
}

case Instruction::AddrSpaceCast: {
  // Strip the addrspacecast and pass along the operand
  PointerType *DstTy = cast<PointerType>(CE->getType());
  if (DstTy->getAddressSpace() == 0) {
    return lowerConstantForGV(cast<const Constant>(CE->getOperand(0)), true);
  }
  std::string S;
  raw_string_ostream OS(S);
  OS << "Unsupported expression in static initializer: ";
  CE->printAsOperand(OS, /*PrintType=*/ false,
                     !MF ? nullptr : MF->getFunction().getParent());
  report_fatal_error(OS.str());
}

case Instruction::GetElementPtr: {
  const DataLayout &DL = getDataLayout();

  // Generate a symbolic expression for the byte address
  APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0);
  cast<GEPOperator>(CE)->accumulateConstantOffset(DL, OffsetAI);

  const MCExpr *Base = lowerConstantForGV(CE->getOperand(0),
                                          ProcessingGeneric);
  if (!OffsetAI)
    return Base;

  int64_t Offset = OffsetAI.getSExtValue();
  return MCBinaryExpr::createAdd(Base, MCConstantExpr::create(Offset, Ctx),
                                 Ctx);
}

case Instruction::Trunc:
  // We emit the value and depend on the assembler to truncate the generated
  // expression properly.  This is important for differences between
  // blockaddress labels.  Since the two labels are in the same function, it
  // is reasonable to treat their delta as a 32-bit value.
  LLVM_FALLTHROUGH[[gnu::fallthrough]];
case Instruction::BitCast:
  return lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);

case Instruction::IntToPtr: {
  const DataLayout &DL = getDataLayout();

  // Handle casts to pointers by changing them into casts to the appropriate
  // integer type.  This promotes constant folding and simplifies this code.
  Constant *Op = CE->getOperand(0);
  Op = ConstantExpr::getIntegerCast(Op, DL.getIntPtrType(CV->getType()),
                                    false/*ZExt*/);
  return lowerConstantForGV(Op, ProcessingGeneric);
}

case Instruction::PtrToInt: {
  const DataLayout &DL = getDataLayout();

  // Support only foldable casts to/from pointers that can be eliminated by
  // changing the pointer to the appropriately sized integer type.
  Constant *Op = CE->getOperand(0);
  Type *Ty = CE->getType();

  const MCExpr *OpExpr = lowerConstantForGV(Op, ProcessingGeneric);

  // We can emit the pointer value into this slot if the slot is an
  // integer slot equal to the size of the pointer.
  if (DL.getTypeAllocSize(Ty) == DL.getTypeAllocSize(Op->getType()))
    return OpExpr;

  // Otherwise the pointer is smaller than the resultant integer, mask off
  // the high bits so we are sure to get a proper truncation if the input is
  // a constant expr.
  unsigned InBits = DL.getTypeAllocSizeInBits(Op->getType());
  const MCExpr *MaskExpr = MCConstantExpr::create(~0ULL >> (64-InBits), Ctx);
  return MCBinaryExpr::createAnd(OpExpr, MaskExpr, Ctx);
}

// The MC library also has a right-shift operator, but it isn't consistently
// signed or unsigned between different targets.
case Instruction::Add: {
  const MCExpr *LHS = lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
  const MCExpr *RHS = lowerConstantForGV(CE->getOperand(1), ProcessingGeneric);
  switch (CE->getOpcode()) {
  default: llvm_unreachable("Unknown binary operator constant cast expr")__builtin_unreachable();
  case Instruction::Add: return MCBinaryExpr::createAdd(LHS, RHS, Ctx);
  }
}
}
2087}

2089// Copy of MCExpr::print customized for NVPTX
2090void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) {
switch (Expr.getKind()) {
case MCExpr::Target:
  return cast<MCTargetExpr>(&Expr)->printImpl(OS, MAI);
case MCExpr::Constant:
  OS << cast<MCConstantExpr>(Expr).getValue();
  return;

case MCExpr::SymbolRef: {
  const MCSymbolRefExpr &SRE = cast<MCSymbolRefExpr>(Expr);
  const MCSymbol &Sym = SRE.getSymbol();
  Sym.print(OS, MAI);
  return;
}

case MCExpr::Unary: {
  const MCUnaryExpr &UE = cast<MCUnaryExpr>(Expr);
  switch (UE.getOpcode()) {
  case MCUnaryExpr::LNot:  OS << '!'; break;
  case MCUnaryExpr::Minus: OS << '-'; break;
  case MCUnaryExpr::Not:   OS << '~'; break;
  case MCUnaryExpr::Plus:  OS << '+'; break;
  }
  printMCExpr(*UE.getSubExpr(), OS);
  return;
}

case MCExpr::Binary: {
  const MCBinaryExpr &BE = cast<MCBinaryExpr>(Expr);

  // Only print parens around the LHS if it is non-trivial.
  if (isa<MCConstantExpr>(BE.getLHS()) || isa<MCSymbolRefExpr>(BE.getLHS()) ||
      isa<NVPTXGenericMCSymbolRefExpr>(BE.getLHS())) {
    printMCExpr(*BE.getLHS(), OS);
  } else {
    OS << '(';
    printMCExpr(*BE.getLHS(), OS);
    OS<< ')';
  }

  switch (BE.getOpcode()) {
  case MCBinaryExpr::Add:
    // Print "X-42" instead of "X+-42".
    if (const MCConstantExpr *RHSC = dyn_cast<MCConstantExpr>(BE.getRHS())) {
      if (RHSC->getValue() < 0) {
        OS << RHSC->getValue();
        return;
      }
    }

    OS <<  '+';
    break;
  default: llvm_unreachable("Unhandled binary operator")__builtin_unreachable();
  }

  // Only print parens around the LHS if it is non-trivial.
  if (isa<MCConstantExpr>(BE.getRHS()) || isa<MCSymbolRefExpr>(BE.getRHS())) {
    printMCExpr(*BE.getRHS(), OS);
  } else {
    OS << '(';
    printMCExpr(*BE.getRHS(), OS);
    OS << ')';
  }
  return;
}
}

llvm_unreachable("Invalid expression kind!")__builtin_unreachable();
2158}

2160/// PrintAsmOperand - Print out an operand for an inline asm expression.
2161///
2162bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
                                    const char *ExtraCode, raw_ostream &O) {
if (ExtraCode && ExtraCode[0]) {
  if (ExtraCode[1] != 0)
    return true; // Unknown modifier.

  switch (ExtraCode[0]) {
  default:
    // See if this is a generic print operand
    return AsmPrinter::PrintAsmOperand(MI, OpNo, ExtraCode, O);
  case 'r':
    break;
  }
}

printOperand(MI, OpNo, O);

return false;
2180}

2182bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI,
                                          unsigned OpNo,
                                          const char *ExtraCode,
                                          raw_ostream &O) {
if (ExtraCode && ExtraCode[0])
  return true; // Unknown modifier

O << '[';
printMemOperand(MI, OpNo, O);
O << ']';

return false;
2194}

2196void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum,
                                 raw_ostream &O) {
const MachineOperand &MO = MI->getOperand(opNum);
switch (MO.getType()) {
case MachineOperand::MO_Register:
  if (Register::isPhysicalRegister(MO.getReg())) {
    if (MO.getReg() == NVPTX::VRDepot)
      O << DEPOTNAME"__local_depot" << getFunctionNumber();
    else
      O << NVPTXInstPrinter::getRegisterName(MO.getReg());
  } else {
    emitVirtualRegister(MO.getReg(), O);
  }
  break;

case MachineOperand::MO_Immediate:
  O << MO.getImm();
  break;

case MachineOperand::MO_FPImmediate:
  printFPConstant(MO.getFPImm(), O);
  break;

case MachineOperand::MO_GlobalAddress:
  PrintSymbolOperand(MO, O);
  break;

case MachineOperand::MO_MachineBasicBlock:
  MO.getMBB()->getSymbol()->print(O, MAI);
  break;

default:
  llvm_unreachable("Operand type not supported.")__builtin_unreachable();
}
2230}

2232void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum,
                                    raw_ostream &O, const char *Modifier) {
printOperand(MI, opNum, O);

if (Modifier && strcmp(Modifier, "add") == 0) {
  O << ", ";
  printOperand(MI, opNum + 1, O);
} else {
  if (MI->getOperand(opNum + 1).isImm() &&
      MI->getOperand(opNum + 1).getImm() == 0)
    return; // don't print ',0' or '+0'
  O << "+";
  printOperand(MI, opNum + 1, O);
}
2246}

2248// Force static initialization.
2249extern "C" LLVM_EXTERNAL_VISIBILITY__attribute__ ((visibility("default"))) void LLVMInitializeNVPTXAsmPrinter() {
RegisterAsmPrinter<NVPTXAsmPrinter> X(getTheNVPTXTarget32());
RegisterAsmPrinter<NVPTXAsmPrinter> Y(getTheNVPTXTarget64());
2252}