1 //===-- NVPTXAsmPrinter.cpp - NVPTX LLVM assembly writer ------------------===//
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
7 //===----------------------------------------------------------------------===//
9 // This file contains a printer that converts from our internal representation
10 // of machine-dependent LLVM code to NVPTX assembly language.
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"
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"
94 #define DEPOTNAME "__local_depot"
96 /// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
99 DiscoverDependentGlobals(const Value *V,
100 DenseSet<const GlobalVariable *> &Globals) {
101 if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(V))
104 if (const User *U = dyn_cast<User>(V)) {
105 for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) {
106 DiscoverDependentGlobals(U->getOperand(i), Globals);
112 /// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable
113 /// instances to be emitted, but only after any dependents have been added
116 VisitGlobalVariableForEmission(const GlobalVariable *GV,
117 SmallVectorImpl<const GlobalVariable *> &Order,
118 DenseSet<const GlobalVariable *> &Visited,
119 DenseSet<const GlobalVariable *> &Visiting) {
120 // Have we already visited this one?
121 if (Visited.count(GV))
124 // Do we have a circular dependency?
125 if (!Visiting.insert(GV).second)
126 report_fatal_error("Circular dependency found in global variable set");
128 // Make sure we visit all dependents first
129 DenseSet<const GlobalVariable *> Others;
130 for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i)
131 DiscoverDependentGlobals(GV->getOperand(i), Others);
133 for (DenseSet<const GlobalVariable *>::iterator I = Others.begin(),
136 VisitGlobalVariableForEmission(*I, Order, Visited, Visiting);
138 // Now we can visit ourself
144 void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) {
146 lowerToMCInst(MI, Inst);
147 EmitToStreamer(*OutStreamer, Inst);
150 // Handle symbol backtracking for targets that do not support image handles
151 bool NVPTXAsmPrinter::lowerImageHandleOperand(const MachineInstr *MI,
152 unsigned OpNo, MCOperand &MCOp) {
153 const MachineOperand &MO = MI->getOperand(OpNo);
154 const MCInstrDesc &MCID = MI->getDesc();
156 if (MCID.TSFlags & NVPTXII::IsTexFlag) {
157 // This is a texture fetch, so operand 4 is a texref and operand 5 is
159 if (OpNo == 4 && MO.isImm()) {
160 lowerImageHandleSymbol(MO.getImm(), MCOp);
163 if (OpNo == 5 && MO.isImm() && !(MCID.TSFlags & NVPTXII::IsTexModeUnifiedFlag)) {
164 lowerImageHandleSymbol(MO.getImm(), MCOp);
169 } else if (MCID.TSFlags & NVPTXII::IsSuldMask) {
171 1 << (((MCID.TSFlags & NVPTXII::IsSuldMask) >> NVPTXII::IsSuldShift) - 1);
173 // For a surface load of vector size N, the Nth operand will be the surfref
174 if (OpNo == VecSize && MO.isImm()) {
175 lowerImageHandleSymbol(MO.getImm(), MCOp);
180 } else if (MCID.TSFlags & NVPTXII::IsSustFlag) {
181 // This is a surface store, so operand 0 is a surfref
182 if (OpNo == 0 && MO.isImm()) {
183 lowerImageHandleSymbol(MO.getImm(), MCOp);
188 } else if (MCID.TSFlags & NVPTXII::IsSurfTexQueryFlag) {
189 // This is a query, so operand 1 is a surfref/texref
190 if (OpNo == 1 && MO.isImm()) {
191 lowerImageHandleSymbol(MO.getImm(), MCOp);
201 void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) {
203 LLVMTargetMachine &TM = const_cast<LLVMTargetMachine&>(MF->getTarget());
204 NVPTXTargetMachine &nvTM = static_cast<NVPTXTargetMachine&>(TM);
205 const NVPTXMachineFunctionInfo *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
206 const char *Sym = MFI->getImageHandleSymbol(Index);
207 std::string *SymNamePtr =
208 nvTM.getManagedStrPool()->getManagedString(Sym);
209 MCOp = GetSymbolRef(OutContext.getOrCreateSymbol(StringRef(*SymNamePtr)));
212 void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
213 OutMI.setOpcode(MI->getOpcode());
214 // Special: Do not mangle symbol operand of CALL_PROTOTYPE
215 if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
216 const MachineOperand &MO = MI->getOperand(0);
217 OutMI.addOperand(GetSymbolRef(
218 OutContext.getOrCreateSymbol(Twine(MO.getSymbolName()))));
222 const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
223 for (unsigned i = 0, e = MI->getNumOperands(); i != e; ++i) {
224 const MachineOperand &MO = MI->getOperand(i);
227 if (!STI.hasImageHandles()) {
228 if (lowerImageHandleOperand(MI, i, MCOp)) {
229 OutMI.addOperand(MCOp);
234 if (lowerOperand(MO, MCOp))
235 OutMI.addOperand(MCOp);
239 bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO,
241 switch (MO.getType()) {
242 default: llvm_unreachable("unknown operand type");
243 case MachineOperand::MO_Register:
244 MCOp = MCOperand::createReg(encodeVirtualRegister(MO.getReg()));
246 case MachineOperand::MO_Immediate:
247 MCOp = MCOperand::createImm(MO.getImm());
249 case MachineOperand::MO_MachineBasicBlock:
250 MCOp = MCOperand::createExpr(MCSymbolRefExpr::create(
251 MO.getMBB()->getSymbol(), OutContext));
253 case MachineOperand::MO_ExternalSymbol:
254 MCOp = GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
256 case MachineOperand::MO_GlobalAddress:
257 MCOp = GetSymbolRef(getSymbol(MO.getGlobal()));
259 case MachineOperand::MO_FPImmediate: {
260 const ConstantFP *Cnt = MO.getFPImm();
261 const APFloat &Val = Cnt->getValueAPF();
263 switch (Cnt->getType()->getTypeID()) {
264 default: report_fatal_error("Unsupported FP type"); break;
266 MCOp = MCOperand::createExpr(
267 NVPTXFloatMCExpr::createConstantFPHalf(Val, OutContext));
269 case Type::FloatTyID:
270 MCOp = MCOperand::createExpr(
271 NVPTXFloatMCExpr::createConstantFPSingle(Val, OutContext));
273 case Type::DoubleTyID:
274 MCOp = MCOperand::createExpr(
275 NVPTXFloatMCExpr::createConstantFPDouble(Val, OutContext));
284 unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
285 if (Register::isVirtualRegister(Reg)) {
286 const TargetRegisterClass *RC = MRI->getRegClass(Reg);
288 DenseMap<unsigned, unsigned> &RegMap = VRegMapping[RC];
289 unsigned RegNum = RegMap[Reg];
291 // Encode the register class in the upper 4 bits
292 // Must be kept in sync with NVPTXInstPrinter::printRegName
294 if (RC == &NVPTX::Int1RegsRegClass) {
296 } else if (RC == &NVPTX::Int16RegsRegClass) {
298 } else if (RC == &NVPTX::Int32RegsRegClass) {
300 } else if (RC == &NVPTX::Int64RegsRegClass) {
302 } else if (RC == &NVPTX::Float32RegsRegClass) {
304 } else if (RC == &NVPTX::Float64RegsRegClass) {
306 } else if (RC == &NVPTX::Float16RegsRegClass) {
308 } else if (RC == &NVPTX::Float16x2RegsRegClass) {
311 report_fatal_error("Bad register class");
314 // Insert the vreg number
315 Ret |= (RegNum & 0x0FFFFFFF);
318 // Some special-use registers are actually physical registers.
319 // Encode this as the register class ID of 0 and the real register ID.
320 return Reg & 0x0FFFFFFF;
324 MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
326 Expr = MCSymbolRefExpr::create(Symbol, MCSymbolRefExpr::VK_None,
328 return MCOperand::createExpr(Expr);
331 void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
332 const DataLayout &DL = getDataLayout();
333 const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
334 const TargetLowering *TLI = STI.getTargetLowering();
336 Type *Ty = F->getReturnType();
338 bool isABI = (STI.getSmVersion() >= 20);
340 if (Ty->getTypeID() == Type::VoidTyID)
346 if (Ty->isFloatingPointTy() || (Ty->isIntegerTy() && !Ty->isIntegerTy(128))) {
348 if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
349 size = ITy->getBitWidth();
351 assert(Ty->isFloatingPointTy() && "Floating point type expected here");
352 size = Ty->getPrimitiveSizeInBits();
354 // PTX ABI requires all scalar return values to be at least 32
355 // bits in size. fp16 normally uses .b16 as its storage type in
356 // PTX, so its size must be adjusted here, too.
360 O << ".param .b" << size << " func_retval0";
361 } else if (isa<PointerType>(Ty)) {
362 O << ".param .b" << TLI->getPointerTy(DL).getSizeInBits()
364 } else if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
365 unsigned totalsz = DL.getTypeAllocSize(Ty);
366 unsigned retAlignment = 0;
367 if (!getAlign(*F, 0, retAlignment))
368 retAlignment = DL.getABITypeAlignment(Ty);
369 O << ".param .align " << retAlignment << " .b8 func_retval0[" << totalsz
372 llvm_unreachable("Unknown return type");
374 SmallVector<EVT, 16> vtparts;
375 ComputeValueVTs(*TLI, DL, Ty, vtparts);
377 for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
379 EVT elemtype = vtparts[i];
380 if (vtparts[i].isVector()) {
381 elems = vtparts[i].getVectorNumElements();
382 elemtype = vtparts[i].getVectorElementType();
385 for (unsigned j = 0, je = elems; j != je; ++j) {
386 unsigned sz = elemtype.getSizeInBits();
387 if (elemtype.isInteger() && (sz < 32))
389 O << ".reg .b" << sz << " func_retval" << idx;
401 void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
403 const Function &F = MF.getFunction();
404 printReturnValStr(&F, O);
407 // Return true if MBB is the header of a loop marked with
408 // llvm.loop.unroll.disable.
409 // TODO: consider "#pragma unroll 1" which is equivalent to "#pragma nounroll".
410 bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
411 const MachineBasicBlock &MBB) const {
412 MachineLoopInfo &LI = getAnalysis<MachineLoopInfo>();
413 // We insert .pragma "nounroll" only to the loop header.
414 if (!LI.isLoopHeader(&MBB))
417 // llvm.loop.unroll.disable is marked on the back edges of a loop. Therefore,
418 // we iterate through each back edge of the loop with header MBB, and check
419 // whether its metadata contains llvm.loop.unroll.disable.
420 for (auto I = MBB.pred_begin(); I != MBB.pred_end(); ++I) {
421 const MachineBasicBlock *PMBB = *I;
422 if (LI.getLoopFor(PMBB) != LI.getLoopFor(&MBB)) {
423 // Edges from other loops to MBB are not back edges.
426 if (const BasicBlock *PBB = PMBB->getBasicBlock()) {
428 PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
429 if (GetUnrollMetadata(LoopID, "llvm.loop.unroll.disable"))
437 void NVPTXAsmPrinter::EmitBasicBlockStart(const MachineBasicBlock &MBB) {
438 AsmPrinter::EmitBasicBlockStart(MBB);
439 if (isLoopHeaderOfNoUnroll(MBB))
440 OutStreamer->EmitRawText(StringRef("\t.pragma \"nounroll\";\n"));
443 void NVPTXAsmPrinter::EmitFunctionEntryLabel() {
444 SmallString<128> Str;
445 raw_svector_ostream O(Str);
447 if (!GlobalsEmitted) {
448 emitGlobals(*MF->getFunction().getParent());
449 GlobalsEmitted = true;
453 MRI = &MF->getRegInfo();
454 F = &MF->getFunction();
455 emitLinkageDirective(F, O);
456 if (isKernelFunction(*F))
460 printReturnValStr(*MF, O);
463 CurrentFnSym->print(O, MAI);
465 emitFunctionParamList(*MF, O);
467 if (isKernelFunction(*F))
468 emitKernelFunctionDirectives(*F, O);
470 OutStreamer->EmitRawText(O.str());
473 // Emit open brace for function body.
474 OutStreamer->EmitRawText(StringRef("{\n"));
475 setAndEmitFunctionVirtualRegisters(*MF);
476 // Emit initial .loc debug directive for correct relocation symbol data.
477 if (MMI && MMI->hasDebugInfo())
478 emitInitialRawDwarfLocDirective(*MF);
481 bool NVPTXAsmPrinter::runOnMachineFunction(MachineFunction &F) {
482 bool Result = AsmPrinter::runOnMachineFunction(F);
483 // Emit closing brace for the body of function F.
484 // The closing brace must be emitted here because we need to emit additional
485 // debug labels/data after the last basic block.
486 // We need to emit the closing brace here because we don't have function that
487 // finished emission of the function body.
488 OutStreamer->EmitRawText(StringRef("}\n"));
492 void NVPTXAsmPrinter::EmitFunctionBodyStart() {
493 SmallString<128> Str;
494 raw_svector_ostream O(Str);
495 emitDemotedVars(&MF->getFunction(), O);
496 OutStreamer->EmitRawText(O.str());
499 void NVPTXAsmPrinter::EmitFunctionBodyEnd() {
503 const MCSymbol *NVPTXAsmPrinter::getFunctionFrameSymbol() const {
504 SmallString<128> Str;
505 raw_svector_ostream(Str) << DEPOTNAME << getFunctionNumber();
506 return OutContext.getOrCreateSymbol(Str);
509 void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const {
510 Register RegNo = MI->getOperand(0).getReg();
511 if (Register::isVirtualRegister(RegNo)) {
512 OutStreamer->AddComment(Twine("implicit-def: ") +
513 getVirtualRegisterName(RegNo));
515 const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
516 OutStreamer->AddComment(Twine("implicit-def: ") +
517 STI.getRegisterInfo()->getName(RegNo));
519 OutStreamer->AddBlankLine();
522 void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F,
523 raw_ostream &O) const {
524 // If the NVVM IR has some of reqntid* specified, then output
525 // the reqntid directive, and set the unspecified ones to 1.
526 // If none of reqntid* is specified, don't output reqntid directive.
527 unsigned reqntidx, reqntidy, reqntidz;
528 bool specified = false;
529 if (!getReqNTIDx(F, reqntidx))
533 if (!getReqNTIDy(F, reqntidy))
537 if (!getReqNTIDz(F, reqntidz))
543 O << ".reqntid " << reqntidx << ", " << reqntidy << ", " << reqntidz
546 // If the NVVM IR has some of maxntid* specified, then output
547 // the maxntid directive, and set the unspecified ones to 1.
548 // If none of maxntid* is specified, don't output maxntid directive.
549 unsigned maxntidx, maxntidy, maxntidz;
551 if (!getMaxNTIDx(F, maxntidx))
555 if (!getMaxNTIDy(F, maxntidy))
559 if (!getMaxNTIDz(F, maxntidz))
565 O << ".maxntid " << maxntidx << ", " << maxntidy << ", " << maxntidz
569 if (getMinCTASm(F, mincta))
570 O << ".minnctapersm " << mincta << "\n";
573 if (getMaxNReg(F, maxnreg))
574 O << ".maxnreg " << maxnreg << "\n";
578 NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const {
579 const TargetRegisterClass *RC = MRI->getRegClass(Reg);
582 raw_string_ostream NameStr(Name);
584 VRegRCMap::const_iterator I = VRegMapping.find(RC);
585 assert(I != VRegMapping.end() && "Bad register class");
586 const DenseMap<unsigned, unsigned> &RegMap = I->second;
588 VRegMap::const_iterator VI = RegMap.find(Reg);
589 assert(VI != RegMap.end() && "Bad virtual register");
590 unsigned MappedVR = VI->second;
592 NameStr << getNVPTXRegClassStr(RC) << MappedVR;
598 void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr,
600 O << getVirtualRegisterName(vr);
603 void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) {
604 emitLinkageDirective(F, O);
605 if (isKernelFunction(*F))
609 printReturnValStr(F, O);
610 getSymbol(F)->print(O, MAI);
612 emitFunctionParamList(F, O);
616 static bool usedInGlobalVarDef(const Constant *C) {
620 if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(C)) {
621 return GV->getName() != "llvm.used";
624 for (const User *U : C->users())
625 if (const Constant *C = dyn_cast<Constant>(U))
626 if (usedInGlobalVarDef(C))
632 static bool usedInOneFunc(const User *U, Function const *&oneFunc) {
633 if (const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
634 if (othergv->getName() == "llvm.used")
638 if (const Instruction *instr = dyn_cast<Instruction>(U)) {
639 if (instr->getParent() && instr->getParent()->getParent()) {
640 const Function *curFunc = instr->getParent()->getParent();
641 if (oneFunc && (curFunc != oneFunc))
649 for (const User *UU : U->users())
650 if (!usedInOneFunc(UU, oneFunc))
656 /* Find out if a global variable can be demoted to local scope.
657 * Currently, this is valid for CUDA shared variables, which have local
658 * scope and global lifetime. So the conditions to check are :
659 * 1. Is the global variable in shared address space?
660 * 2. Does it have internal linkage?
661 * 3. Is the global variable referenced only in one function?
663 static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) {
664 if (!gv->hasInternalLinkage())
666 PointerType *Pty = gv->getType();
667 if (Pty->getAddressSpace() != ADDRESS_SPACE_SHARED)
670 const Function *oneFunc = nullptr;
672 bool flag = usedInOneFunc(gv, oneFunc);
681 static bool useFuncSeen(const Constant *C,
682 DenseMap<const Function *, bool> &seenMap) {
683 for (const User *U : C->users()) {
684 if (const Constant *cu = dyn_cast<Constant>(U)) {
685 if (useFuncSeen(cu, seenMap))
687 } else if (const Instruction *I = dyn_cast<Instruction>(U)) {
688 const BasicBlock *bb = I->getParent();
691 const Function *caller = bb->getParent();
694 if (seenMap.find(caller) != seenMap.end())
701 void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {
702 DenseMap<const Function *, bool> seenMap;
703 for (Module::const_iterator FI = M.begin(), FE = M.end(); FI != FE; ++FI) {
704 const Function *F = &*FI;
706 if (F->getAttributes().hasFnAttribute("nvptx-libcall-callee")) {
707 emitDeclaration(F, O);
711 if (F->isDeclaration()) {
714 if (F->getIntrinsicID())
716 emitDeclaration(F, O);
719 for (const User *U : F->users()) {
720 if (const Constant *C = dyn_cast<Constant>(U)) {
721 if (usedInGlobalVarDef(C)) {
722 // The use is in the initialization of a global variable
723 // that is a function pointer, so print a declaration
724 // for the original function
725 emitDeclaration(F, O);
728 // Emit a declaration of this function if the function that
729 // uses this constant expr has already been seen.
730 if (useFuncSeen(C, seenMap)) {
731 emitDeclaration(F, O);
736 if (!isa<Instruction>(U))
738 const Instruction *instr = cast<Instruction>(U);
739 const BasicBlock *bb = instr->getParent();
742 const Function *caller = bb->getParent();
746 // If a caller has already been seen, then the caller is
747 // appearing in the module before the callee. so print out
748 // a declaration for the callee.
749 if (seenMap.find(caller) != seenMap.end()) {
750 emitDeclaration(F, O);
758 static bool isEmptyXXStructor(GlobalVariable *GV) {
759 if (!GV) return true;
760 const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());
761 if (!InitList) return true; // Not an array; we don't know how to parse.
762 return InitList->getNumOperands() == 0;
765 bool NVPTXAsmPrinter::doInitialization(Module &M) {
766 // Construct a default subtarget off of the TargetMachine defaults. The
767 // rest of NVPTX isn't friendly to change subtargets per function and
768 // so the default TargetMachine will have all of the options.
769 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
770 const auto* STI = static_cast<const NVPTXSubtarget*>(NTM.getSubtargetImpl());
772 if (M.alias_size()) {
773 report_fatal_error("Module has aliases, which NVPTX does not support.");
774 return true; // error
776 if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_ctors"))) {
778 "Module has a nontrivial global ctor, which NVPTX does not support.");
779 return true; // error
781 if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_dtors"))) {
783 "Module has a nontrivial global dtor, which NVPTX does not support.");
784 return true; // error
787 SmallString<128> Str1;
788 raw_svector_ostream OS1(Str1);
790 // We need to call the parent's one explicitly.
791 bool Result = AsmPrinter::doInitialization(M);
793 // Emit header before any dwarf directives are emitted below.
794 emitHeader(M, OS1, *STI);
795 OutStreamer->EmitRawText(OS1.str());
797 // Emit module-level inline asm if it exists.
798 if (!M.getModuleInlineAsm().empty()) {
799 OutStreamer->AddComment("Start of file scope inline assembly");
800 OutStreamer->AddBlankLine();
801 OutStreamer->EmitRawText(StringRef(M.getModuleInlineAsm()));
802 OutStreamer->AddBlankLine();
803 OutStreamer->AddComment("End of file scope inline assembly");
804 OutStreamer->AddBlankLine();
807 GlobalsEmitted = false;
812 void NVPTXAsmPrinter::emitGlobals(const Module &M) {
813 SmallString<128> Str2;
814 raw_svector_ostream OS2(Str2);
816 emitDeclarations(M, OS2);
818 // As ptxas does not support forward references of globals, we need to first
819 // sort the list of module-level globals in def-use order. We visit each
820 // global variable in order, and ensure that we emit it *after* its dependent
821 // globals. We use a little extra memory maintaining both a set and a list to
822 // have fast searches while maintaining a strict ordering.
823 SmallVector<const GlobalVariable *, 8> Globals;
824 DenseSet<const GlobalVariable *> GVVisited;
825 DenseSet<const GlobalVariable *> GVVisiting;
827 // Visit each global variable, in order
828 for (const GlobalVariable &I : M.globals())
829 VisitGlobalVariableForEmission(&I, Globals, GVVisited, GVVisiting);
831 assert(GVVisited.size() == M.getGlobalList().size() &&
832 "Missed a global variable");
833 assert(GVVisiting.size() == 0 && "Did not fully process a global variable");
835 // Print out module-level global variables in proper order
836 for (unsigned i = 0, e = Globals.size(); i != e; ++i)
837 printModuleLevelGV(Globals[i], OS2);
841 OutStreamer->EmitRawText(OS2.str());
844 void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
845 const NVPTXSubtarget &STI) {
847 O << "// Generated by LLVM NVPTX Back-End\n";
851 unsigned PTXVersion = STI.getPTXVersion();
852 O << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n";
855 O << STI.getTargetName();
857 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
858 if (NTM.getDrvInterface() == NVPTX::NVCL)
859 O << ", texmode_independent";
861 bool HasFullDebugInfo = false;
862 for (DICompileUnit *CU : M.debug_compile_units()) {
863 switch(CU->getEmissionKind()) {
864 case DICompileUnit::NoDebug:
865 case DICompileUnit::DebugDirectivesOnly:
867 case DICompileUnit::LineTablesOnly:
868 case DICompileUnit::FullDebug:
869 HasFullDebugInfo = true;
872 if (HasFullDebugInfo)
875 if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo)
880 O << ".address_size ";
890 bool NVPTXAsmPrinter::doFinalization(Module &M) {
891 bool HasDebugInfo = MMI && MMI->hasDebugInfo();
893 // If we did not emit any functions, then the global declarations have not
895 if (!GlobalsEmitted) {
897 GlobalsEmitted = true;
900 // XXX Temproarily remove global variables so that doFinalization() will not
901 // emit them again (global variables are emitted at beginning).
903 Module::GlobalListType &global_list = M.getGlobalList();
904 int i, n = global_list.size();
905 GlobalVariable **gv_array = new GlobalVariable *[n];
907 // first, back-up GlobalVariable in gv_array
909 for (Module::global_iterator I = global_list.begin(), E = global_list.end();
913 // second, empty global_list
914 while (!global_list.empty())
915 global_list.remove(global_list.begin());
917 // call doFinalization
918 bool ret = AsmPrinter::doFinalization(M);
920 // now we restore global variables
921 for (i = 0; i < n; i++)
922 global_list.insert(global_list.end(), gv_array[i]);
924 clearAnnotationCache(&M);
927 // Close the last emitted section
929 static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
930 ->closeLastSection();
931 // Emit empty .debug_loc section for better support of the empty files.
932 OutStreamer->EmitRawText("\t.section\t.debug_loc\t{\t}");
935 // Output last DWARF .file directives, if any.
936 static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
937 ->outputDwarfFileDirectives();
941 //bool Result = AsmPrinter::doFinalization(M);
942 // Instead of calling the parents doFinalization, we may
943 // clone parents doFinalization and customize here.
944 // Currently, we if NVISA out the EmitGlobals() in
945 // parent's doFinalization, which is too intrusive.
947 // Same for the doInitialization.
951 // This function emits appropriate linkage directives for
952 // functions and global variables.
954 // extern function declaration -> .extern
955 // extern function definition -> .visible
956 // external global variable with init -> .visible
957 // external without init -> .extern
958 // appending -> not allowed, assert.
959 // for any linkage other than
960 // internal, private, linker_private,
961 // linker_private_weak, linker_private_weak_def_auto,
964 void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V,
966 if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA) {
967 if (V->hasExternalLinkage()) {
968 if (isa<GlobalVariable>(V)) {
969 const GlobalVariable *GVar = cast<GlobalVariable>(V);
971 if (GVar->hasInitializer())
976 } else if (V->isDeclaration())
980 } else if (V->hasAppendingLinkage()) {
982 msg.append("Error: ");
983 msg.append("Symbol ");
985 msg.append(V->getName());
986 msg.append("has unsupported appending linkage type");
987 llvm_unreachable(msg.c_str());
988 } else if (!V->hasInternalLinkage() &&
989 !V->hasPrivateLinkage()) {
995 void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
997 bool processDemoted) {
999 if (GVar->hasSection()) {
1000 if (GVar->getSection() == "llvm.metadata")
1004 // Skip LLVM intrinsic global variables
1005 if (GVar->getName().startswith("llvm.") ||
1006 GVar->getName().startswith("nvvm."))
1009 const DataLayout &DL = getDataLayout();
1011 // GlobalVariables are always constant pointers themselves.
1012 PointerType *PTy = GVar->getType();
1013 Type *ETy = GVar->getValueType();
1015 if (GVar->hasExternalLinkage()) {
1016 if (GVar->hasInitializer())
1020 } else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() ||
1021 GVar->hasAvailableExternallyLinkage() ||
1022 GVar->hasCommonLinkage()) {
1026 if (isTexture(*GVar)) {
1027 O << ".global .texref " << getTextureName(*GVar) << ";\n";
1031 if (isSurface(*GVar)) {
1032 O << ".global .surfref " << getSurfaceName(*GVar) << ";\n";
1036 if (GVar->isDeclaration()) {
1037 // (extern) declarations, no definition or initializer
1038 // Currently the only known declaration is for an automatic __local
1039 // (.shared) promoted to global.
1040 emitPTXGlobalVariable(GVar, O);
1045 if (isSampler(*GVar)) {
1046 O << ".global .samplerref " << getSamplerName(*GVar);
1048 const Constant *Initializer = nullptr;
1049 if (GVar->hasInitializer())
1050 Initializer = GVar->getInitializer();
1051 const ConstantInt *CI = nullptr;
1053 CI = dyn_cast<ConstantInt>(Initializer);
1055 unsigned sample = CI->getZExtValue();
1060 addr = ((sample & __CLK_ADDRESS_MASK) >> __CLK_ADDRESS_BASE);
1062 O << "addr_mode_" << i << " = ";
1068 O << "clamp_to_border";
1071 O << "clamp_to_edge";
1082 O << "filter_mode = ";
1083 switch ((sample & __CLK_FILTER_MASK) >> __CLK_FILTER_BASE) {
1091 llvm_unreachable("Anisotropic filtering is not supported");
1096 if (!((sample & __CLK_NORMALIZED_MASK) >> __CLK_NORMALIZED_BASE)) {
1097 O << ", force_unnormalized_coords = 1";
1106 if (GVar->hasPrivateLinkage()) {
1107 if (strncmp(GVar->getName().data(), "unrollpragma", 12) == 0)
1110 // FIXME - need better way (e.g. Metadata) to avoid generating this global
1111 if (strncmp(GVar->getName().data(), "filename", 8) == 0)
1113 if (GVar->use_empty())
1117 const Function *demotedFunc = nullptr;
1118 if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) {
1119 O << "// " << GVar->getName() << " has been demoted\n";
1120 if (localDecls.find(demotedFunc) != localDecls.end())
1121 localDecls[demotedFunc].push_back(GVar);
1123 std::vector<const GlobalVariable *> temp;
1124 temp.push_back(GVar);
1125 localDecls[demotedFunc] = temp;
1131 emitPTXAddressSpace(PTy->getAddressSpace(), O);
1133 if (isManaged(*GVar)) {
1134 O << " .attribute(.managed)";
1137 if (GVar->getAlignment() == 0)
1138 O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
1140 O << " .align " << GVar->getAlignment();
1142 if (ETy->isFloatingPointTy() || ETy->isPointerTy() ||
1143 (ETy->isIntegerTy() && ETy->getScalarSizeInBits() <= 64)) {
1145 // Special case: ABI requires that we use .u8 for predicates
1146 if (ETy->isIntegerTy(1))
1149 O << getPTXFundamentalTypeStr(ETy, false);
1151 getSymbol(GVar)->print(O, MAI);
1153 // Ptx allows variable initilization only for constant and global state
1155 if (GVar->hasInitializer()) {
1156 if ((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1157 (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) {
1158 const Constant *Initializer = GVar->getInitializer();
1159 // 'undef' is treated as there is no value specified.
1160 if (!Initializer->isNullValue() && !isa<UndefValue>(Initializer)) {
1162 printScalarConstant(Initializer, O);
1165 // The frontend adds zero-initializer to device and constant variables
1166 // that don't have an initial value, and UndefValue to shared
1167 // variables, so skip warning for this case.
1168 if (!GVar->getInitializer()->isNullValue() &&
1169 !isa<UndefValue>(GVar->getInitializer())) {
1170 report_fatal_error("initial value of '" + GVar->getName() +
1171 "' is not allowed in addrspace(" +
1172 Twine(PTy->getAddressSpace()) + ")");
1177 unsigned int ElementSize = 0;
1179 // Although PTX has direct support for struct type and array type and
1180 // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
1181 // targets that support these high level field accesses. Structs, arrays
1182 // and vectors are lowered into arrays of bytes.
1183 switch (ETy->getTypeID()) {
1184 case Type::IntegerTyID: // Integers larger than 64 bits
1185 case Type::StructTyID:
1186 case Type::ArrayTyID:
1187 case Type::VectorTyID:
1188 ElementSize = DL.getTypeStoreSize(ETy);
1189 // Ptx allows variable initilization only for constant and
1190 // global state spaces.
1191 if (((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1192 (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) &&
1193 GVar->hasInitializer()) {
1194 const Constant *Initializer = GVar->getInitializer();
1195 if (!isa<UndefValue>(Initializer) && !Initializer->isNullValue()) {
1196 AggBuffer aggBuffer(ElementSize, O, *this);
1197 bufferAggregateConstant(Initializer, &aggBuffer);
1198 if (aggBuffer.numSymbols) {
1199 if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit()) {
1201 getSymbol(GVar)->print(O, MAI);
1203 O << ElementSize / 8;
1206 getSymbol(GVar)->print(O, MAI);
1208 O << ElementSize / 4;
1213 getSymbol(GVar)->print(O, MAI);
1223 getSymbol(GVar)->print(O, MAI);
1232 getSymbol(GVar)->print(O, MAI);
1241 llvm_unreachable("type not supported yet");
1247 void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) {
1248 if (localDecls.find(f) == localDecls.end())
1251 std::vector<const GlobalVariable *> &gvars = localDecls[f];
1253 for (unsigned i = 0, e = gvars.size(); i != e; ++i) {
1254 O << "\t// demoted variable\n\t";
1255 printModuleLevelGV(gvars[i], O, true);
1259 void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
1260 raw_ostream &O) const {
1261 switch (AddressSpace) {
1262 case ADDRESS_SPACE_LOCAL:
1265 case ADDRESS_SPACE_GLOBAL:
1268 case ADDRESS_SPACE_CONST:
1271 case ADDRESS_SPACE_SHARED:
1275 report_fatal_error("Bad address space found while emitting PTX: " +
1276 llvm::Twine(AddressSpace));
1282 NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {
1283 switch (Ty->getTypeID()) {
1285 llvm_unreachable("unexpected type");
1287 case Type::IntegerTyID: {
1288 unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
1291 else if (NumBits <= 64) {
1292 std::string name = "u";
1293 return name + utostr(NumBits);
1295 llvm_unreachable("Integer too large");
1300 case Type::HalfTyID:
1301 // fp16 is stored as .b16 for compatibility with pre-sm_53 PTX assembly.
1303 case Type::FloatTyID:
1305 case Type::DoubleTyID:
1307 case Type::PointerTyID:
1308 if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit())
1318 llvm_unreachable("unexpected type");
1322 void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,
1324 const DataLayout &DL = getDataLayout();
1326 // GlobalVariables are always constant pointers themselves.
1327 Type *ETy = GVar->getValueType();
1330 emitPTXAddressSpace(GVar->getType()->getAddressSpace(), O);
1331 if (GVar->getAlignment() == 0)
1332 O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
1334 O << " .align " << GVar->getAlignment();
1336 // Special case for i128
1337 if (ETy->isIntegerTy(128)) {
1339 getSymbol(GVar)->print(O, MAI);
1344 if (ETy->isFloatingPointTy() || ETy->isIntOrPtrTy()) {
1346 O << getPTXFundamentalTypeStr(ETy);
1348 getSymbol(GVar)->print(O, MAI);
1352 int64_t ElementSize = 0;
1354 // Although PTX has direct support for struct type and array type and LLVM IR
1355 // is very similar to PTX, the LLVM CodeGen does not support for targets that
1356 // support these high level field accesses. Structs and arrays are lowered
1357 // into arrays of bytes.
1358 switch (ETy->getTypeID()) {
1359 case Type::StructTyID:
1360 case Type::ArrayTyID:
1361 case Type::VectorTyID:
1362 ElementSize = DL.getTypeStoreSize(ETy);
1364 getSymbol(GVar)->print(O, MAI);
1372 llvm_unreachable("type not supported yet");
1376 static unsigned int getOpenCLAlignment(const DataLayout &DL, Type *Ty) {
1377 if (Ty->isSingleValueType())
1378 return DL.getPrefTypeAlignment(Ty);
1380 auto *ATy = dyn_cast<ArrayType>(Ty);
1382 return getOpenCLAlignment(DL, ATy->getElementType());
1384 auto *STy = dyn_cast<StructType>(Ty);
1386 unsigned int alignStruct = 1;
1387 // Go through each element of the struct and find the
1388 // largest alignment.
1389 for (unsigned i = 0, e = STy->getNumElements(); i != e; i++) {
1390 Type *ETy = STy->getElementType(i);
1391 unsigned int align = getOpenCLAlignment(DL, ETy);
1392 if (align > alignStruct)
1393 alignStruct = align;
1398 auto *FTy = dyn_cast<FunctionType>(Ty);
1400 return DL.getPointerPrefAlignment().value();
1401 return DL.getPrefTypeAlignment(Ty);
1404 void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I,
1405 int paramIndex, raw_ostream &O) {
1406 getSymbol(I->getParent())->print(O, MAI);
1407 O << "_param_" << paramIndex;
1410 void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
1411 const DataLayout &DL = getDataLayout();
1412 const AttributeList &PAL = F->getAttributes();
1413 const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
1414 const TargetLowering *TLI = STI.getTargetLowering();
1415 Function::const_arg_iterator I, E;
1416 unsigned paramIndex = 0;
1418 bool isKernelFunc = isKernelFunction(*F);
1419 bool isABI = (STI.getSmVersion() >= 20);
1420 bool hasImageHandles = STI.hasImageHandles();
1421 MVT thePointerTy = TLI->getPointerTy(DL);
1423 if (F->arg_empty()) {
1430 for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) {
1431 Type *Ty = I->getType();
1438 // Handle image/sampler parameters
1439 if (isKernelFunction(*F)) {
1440 if (isSampler(*I) || isImage(*I)) {
1442 std::string sname = I->getName();
1443 if (isImageWriteOnly(*I) || isImageReadWrite(*I)) {
1444 if (hasImageHandles)
1445 O << "\t.param .u64 .ptr .surfref ";
1447 O << "\t.param .surfref ";
1448 CurrentFnSym->print(O, MAI);
1449 O << "_param_" << paramIndex;
1451 else { // Default image is read_only
1452 if (hasImageHandles)
1453 O << "\t.param .u64 .ptr .texref ";
1455 O << "\t.param .texref ";
1456 CurrentFnSym->print(O, MAI);
1457 O << "_param_" << paramIndex;
1460 if (hasImageHandles)
1461 O << "\t.param .u64 .ptr .samplerref ";
1463 O << "\t.param .samplerref ";
1464 CurrentFnSym->print(O, MAI);
1465 O << "_param_" << paramIndex;
1471 if (!PAL.hasParamAttribute(paramIndex, Attribute::ByVal)) {
1472 if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
1473 // Just print .param .align <a> .b8 .param[size];
1474 // <a> = PAL.getparamalignment
1475 // size = typeallocsize of element type
1476 const Align align = DL.getValueOrABITypeAlignment(
1477 PAL.getParamAlignment(paramIndex), Ty);
1479 unsigned sz = DL.getTypeAllocSize(Ty);
1480 O << "\t.param .align " << align.value() << " .b8 ";
1481 printParamName(I, paramIndex, O);
1482 O << "[" << sz << "]";
1487 auto *PTy = dyn_cast<PointerType>(Ty);
1490 // Special handling for pointer arguments to kernel
1491 O << "\t.param .u" << thePointerTy.getSizeInBits() << " ";
1493 if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() !=
1495 Type *ETy = PTy->getElementType();
1496 int addrSpace = PTy->getAddressSpace();
1497 switch (addrSpace) {
1501 case ADDRESS_SPACE_CONST:
1502 O << ".ptr .const ";
1504 case ADDRESS_SPACE_SHARED:
1505 O << ".ptr .shared ";
1507 case ADDRESS_SPACE_GLOBAL:
1508 O << ".ptr .global ";
1511 O << ".align " << (int)getOpenCLAlignment(DL, ETy) << " ";
1513 printParamName(I, paramIndex, O);
1517 // non-pointer scalar to kernel func
1519 // Special case: predicate operands become .u8 types
1520 if (Ty->isIntegerTy(1))
1523 O << getPTXFundamentalTypeStr(Ty);
1525 printParamName(I, paramIndex, O);
1528 // Non-kernel function, just print .param .b<size> for ABI
1529 // and .reg .b<size> for non-ABI
1531 if (isa<IntegerType>(Ty)) {
1532 sz = cast<IntegerType>(Ty)->getBitWidth();
1535 } else if (isa<PointerType>(Ty))
1536 sz = thePointerTy.getSizeInBits();
1537 else if (Ty->isHalfTy())
1538 // PTX ABI requires all scalar parameters to be at least 32
1539 // bits in size. fp16 normally uses .b16 as its storage type
1540 // in PTX, so its size must be adjusted here, too.
1543 sz = Ty->getPrimitiveSizeInBits();
1545 O << "\t.param .b" << sz << " ";
1547 O << "\t.reg .b" << sz << " ";
1548 printParamName(I, paramIndex, O);
1552 // param has byVal attribute. So should be a pointer
1553 auto *PTy = dyn_cast<PointerType>(Ty);
1554 assert(PTy && "Param with byval attribute should be a pointer type");
1555 Type *ETy = PTy->getElementType();
1557 if (isABI || isKernelFunc) {
1558 // Just print .param .align <a> .b8 .param[size];
1559 // <a> = PAL.getparamalignment
1560 // size = typeallocsize of element type
1562 DL.getValueOrABITypeAlignment(PAL.getParamAlignment(paramIndex), ETy);
1563 // Work around a bug in ptxas. When PTX code takes address of
1564 // byval parameter with alignment < 4, ptxas generates code to
1565 // spill argument into memory. Alas on sm_50+ ptxas generates
1566 // SASS code that fails with misaligned access. To work around
1567 // the problem, make sure that we align byval parameters by at
1568 // least 4. Matching change must be made in LowerCall() where we
1569 // prepare parameters for the call.
1571 // TODO: this will need to be undone when we get to support multi-TU
1572 // device-side compilation as it breaks ABI compatibility with nvcc.
1573 // Hopefully ptxas bug is fixed by then.
1574 if (!isKernelFunc && align < Align(4))
1576 unsigned sz = DL.getTypeAllocSize(ETy);
1577 O << "\t.param .align " << align.value() << " .b8 ";
1578 printParamName(I, paramIndex, O);
1579 O << "[" << sz << "]";
1582 // Split the ETy into constituent parts and
1583 // print .param .b<size> <name> for each part.
1584 // Further, if a part is vector, print the above for
1585 // each vector element.
1586 SmallVector<EVT, 16> vtparts;
1587 ComputeValueVTs(*TLI, DL, ETy, vtparts);
1588 for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
1590 EVT elemtype = vtparts[i];
1591 if (vtparts[i].isVector()) {
1592 elems = vtparts[i].getVectorNumElements();
1593 elemtype = vtparts[i].getVectorElementType();
1596 for (unsigned j = 0, je = elems; j != je; ++j) {
1597 unsigned sz = elemtype.getSizeInBits();
1598 if (elemtype.isInteger() && (sz < 32))
1600 O << "\t.reg .b" << sz << " ";
1601 printParamName(I, paramIndex, O);
1617 void NVPTXAsmPrinter::emitFunctionParamList(const MachineFunction &MF,
1619 const Function &F = MF.getFunction();
1620 emitFunctionParamList(&F, O);
1623 void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
1624 const MachineFunction &MF) {
1625 SmallString<128> Str;
1626 raw_svector_ostream O(Str);
1628 // Map the global virtual register number to a register class specific
1629 // virtual register number starting from 1 with that class.
1630 const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo();
1631 //unsigned numRegClasses = TRI->getNumRegClasses();
1633 // Emit the Fake Stack Object
1634 const MachineFrameInfo &MFI = MF.getFrameInfo();
1635 int NumBytes = (int) MFI.getStackSize();
1637 O << "\t.local .align " << MFI.getMaxAlignment() << " .b8 \t" << DEPOTNAME
1638 << getFunctionNumber() << "[" << NumBytes << "];\n";
1639 if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
1640 O << "\t.reg .b64 \t%SP;\n";
1641 O << "\t.reg .b64 \t%SPL;\n";
1643 O << "\t.reg .b32 \t%SP;\n";
1644 O << "\t.reg .b32 \t%SPL;\n";
1648 // Go through all virtual registers to establish the mapping between the
1650 // register number and the per class virtual register number.
1651 // We use the per class virtual register number in the ptx output.
1652 unsigned int numVRs = MRI->getNumVirtRegs();
1653 for (unsigned i = 0; i < numVRs; i++) {
1654 unsigned int vr = Register::index2VirtReg(i);
1655 const TargetRegisterClass *RC = MRI->getRegClass(vr);
1656 DenseMap<unsigned, unsigned> ®map = VRegMapping[RC];
1657 int n = regmap.size();
1658 regmap.insert(std::make_pair(vr, n + 1));
1661 // Emit register declarations
1662 // @TODO: Extract out the real register usage
1663 // O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n";
1664 // O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n";
1665 // O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n";
1666 // O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n";
1667 // O << "\t.reg .s64 %rd<" << NVPTXNumRegisters << ">;\n";
1668 // O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n";
1669 // O << "\t.reg .f64 %fd<" << NVPTXNumRegisters << ">;\n";
1671 // Emit declaration of the virtual registers or 'physical' registers for
1672 // each register class
1673 for (unsigned i=0; i< TRI->getNumRegClasses(); i++) {
1674 const TargetRegisterClass *RC = TRI->getRegClass(i);
1675 DenseMap<unsigned, unsigned> ®map = VRegMapping[RC];
1676 std::string rcname = getNVPTXRegClassName(RC);
1677 std::string rcStr = getNVPTXRegClassStr(RC);
1678 int n = regmap.size();
1680 // Only declare those registers that may be used.
1682 O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1)
1687 OutStreamer->EmitRawText(O.str());
1690 void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) {
1691 APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
1693 unsigned int numHex;
1696 if (Fp->getType()->getTypeID() == Type::FloatTyID) {
1699 APF.convert(APFloat::IEEEsingle(), APFloat::rmNearestTiesToEven, &ignored);
1700 } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
1703 APF.convert(APFloat::IEEEdouble(), APFloat::rmNearestTiesToEven, &ignored);
1705 llvm_unreachable("unsupported fp type");
1707 APInt API = APF.bitcastToAPInt();
1708 O << lead << format_hex_no_prefix(API.getZExtValue(), numHex, /*Upper=*/true);
1711 void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
1712 if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1713 O << CI->getValue();
1716 if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
1717 printFPConstant(CFP, O);
1720 if (isa<ConstantPointerNull>(CPV)) {
1724 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1725 bool IsNonGenericPointer = false;
1726 if (GVar->getType()->getAddressSpace() != 0) {
1727 IsNonGenericPointer = true;
1729 if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
1731 getSymbol(GVar)->print(O, MAI);
1734 getSymbol(GVar)->print(O, MAI);
1738 if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1739 const Value *v = Cexpr->stripPointerCasts();
1740 PointerType *PTy = dyn_cast<PointerType>(Cexpr->getType());
1741 bool IsNonGenericPointer = false;
1742 if (PTy && PTy->getAddressSpace() != 0) {
1743 IsNonGenericPointer = true;
1745 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
1746 if (EmitGeneric && !isa<Function>(v) && !IsNonGenericPointer) {
1748 getSymbol(GVar)->print(O, MAI);
1751 getSymbol(GVar)->print(O, MAI);
1755 lowerConstant(CPV)->print(O, MAI);
1759 llvm_unreachable("Not scalar type found in printScalarConstant()");
1762 // These utility functions assure we get the right sequence of bytes for a given
1763 // type even for big-endian machines
1764 template <typename T> static void ConvertIntToBytes(unsigned char *p, T val) {
1765 int64_t vp = (int64_t)val;
1766 for (unsigned i = 0; i < sizeof(T); ++i) {
1767 p[i] = (unsigned char)vp;
1771 static void ConvertFloatToBytes(unsigned char *p, float val) {
1772 int32_t *vp = (int32_t *)&val;
1773 for (unsigned i = 0; i < sizeof(int32_t); ++i) {
1774 p[i] = (unsigned char)*vp;
1778 static void ConvertDoubleToBytes(unsigned char *p, double val) {
1779 int64_t *vp = (int64_t *)&val;
1780 for (unsigned i = 0; i < sizeof(int64_t); ++i) {
1781 p[i] = (unsigned char)*vp;
1786 void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes,
1787 AggBuffer *aggBuffer) {
1788 const DataLayout &DL = getDataLayout();
1790 if (isa<UndefValue>(CPV) || CPV->isNullValue()) {
1791 int s = DL.getTypeAllocSize(CPV->getType());
1794 aggBuffer->addZeros(s);
1798 unsigned char ptr[8];
1799 switch (CPV->getType()->getTypeID()) {
1801 case Type::IntegerTyID: {
1802 Type *ETy = CPV->getType();
1803 if (ETy == Type::getInt8Ty(CPV->getContext())) {
1804 unsigned char c = (unsigned char)cast<ConstantInt>(CPV)->getZExtValue();
1805 ConvertIntToBytes<>(ptr, c);
1806 aggBuffer->addBytes(ptr, 1, Bytes);
1807 } else if (ETy == Type::getInt16Ty(CPV->getContext())) {
1808 short int16 = (short)cast<ConstantInt>(CPV)->getZExtValue();
1809 ConvertIntToBytes<>(ptr, int16);
1810 aggBuffer->addBytes(ptr, 2, Bytes);
1811 } else if (ETy == Type::getInt32Ty(CPV->getContext())) {
1812 if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
1813 int int32 = (int)(constInt->getZExtValue());
1814 ConvertIntToBytes<>(ptr, int32);
1815 aggBuffer->addBytes(ptr, 4, Bytes);
1817 } else if (const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1818 if (const auto *constInt = dyn_cast_or_null<ConstantInt>(
1819 ConstantFoldConstant(Cexpr, DL))) {
1820 int int32 = (int)(constInt->getZExtValue());
1821 ConvertIntToBytes<>(ptr, int32);
1822 aggBuffer->addBytes(ptr, 4, Bytes);
1825 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1826 Value *v = Cexpr->getOperand(0)->stripPointerCasts();
1827 aggBuffer->addSymbol(v, Cexpr->getOperand(0));
1828 aggBuffer->addZeros(4);
1832 llvm_unreachable("unsupported integer const type");
1833 } else if (ETy == Type::getInt64Ty(CPV->getContext())) {
1834 if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
1835 long long int64 = (long long)(constInt->getZExtValue());
1836 ConvertIntToBytes<>(ptr, int64);
1837 aggBuffer->addBytes(ptr, 8, Bytes);
1839 } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1840 if (const auto *constInt = dyn_cast_or_null<ConstantInt>(
1841 ConstantFoldConstant(Cexpr, DL))) {
1842 long long int64 = (long long)(constInt->getZExtValue());
1843 ConvertIntToBytes<>(ptr, int64);
1844 aggBuffer->addBytes(ptr, 8, Bytes);
1847 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1848 Value *v = Cexpr->getOperand(0)->stripPointerCasts();
1849 aggBuffer->addSymbol(v, Cexpr->getOperand(0));
1850 aggBuffer->addZeros(8);
1854 llvm_unreachable("unsupported integer const type");
1856 llvm_unreachable("unsupported integer const type");
1859 case Type::HalfTyID:
1860 case Type::FloatTyID:
1861 case Type::DoubleTyID: {
1862 const auto *CFP = cast<ConstantFP>(CPV);
1863 Type *Ty = CFP->getType();
1864 if (Ty == Type::getHalfTy(CPV->getContext())) {
1865 APInt API = CFP->getValueAPF().bitcastToAPInt();
1866 uint16_t float16 = API.getLoBits(16).getZExtValue();
1867 ConvertIntToBytes<>(ptr, float16);
1868 aggBuffer->addBytes(ptr, 2, Bytes);
1869 } else if (Ty == Type::getFloatTy(CPV->getContext())) {
1870 float float32 = (float) CFP->getValueAPF().convertToFloat();
1871 ConvertFloatToBytes(ptr, float32);
1872 aggBuffer->addBytes(ptr, 4, Bytes);
1873 } else if (Ty == Type::getDoubleTy(CPV->getContext())) {
1874 double float64 = CFP->getValueAPF().convertToDouble();
1875 ConvertDoubleToBytes(ptr, float64);
1876 aggBuffer->addBytes(ptr, 8, Bytes);
1878 llvm_unreachable("unsupported fp const type");
1882 case Type::PointerTyID: {
1883 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1884 aggBuffer->addSymbol(GVar, GVar);
1885 } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1886 const Value *v = Cexpr->stripPointerCasts();
1887 aggBuffer->addSymbol(v, Cexpr);
1889 unsigned int s = DL.getTypeAllocSize(CPV->getType());
1890 aggBuffer->addZeros(s);
1894 case Type::ArrayTyID:
1895 case Type::VectorTyID:
1896 case Type::StructTyID: {
1897 if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
1898 int ElementSize = DL.getTypeAllocSize(CPV->getType());
1899 bufferAggregateConstant(CPV, aggBuffer);
1900 if (Bytes > ElementSize)
1901 aggBuffer->addZeros(Bytes - ElementSize);
1902 } else if (isa<ConstantAggregateZero>(CPV))
1903 aggBuffer->addZeros(Bytes);
1905 llvm_unreachable("Unexpected Constant type");
1910 llvm_unreachable("unsupported type");
1914 void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV,
1915 AggBuffer *aggBuffer) {
1916 const DataLayout &DL = getDataLayout();
1919 // Integers of arbitrary width
1920 if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1921 APInt Val = CI->getValue();
1922 for (unsigned I = 0, E = DL.getTypeAllocSize(CPV->getType()); I < E; ++I) {
1923 uint8_t Byte = Val.getLoBits(8).getZExtValue();
1924 aggBuffer->addBytes(&Byte, 1, 1);
1931 if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
1932 if (CPV->getNumOperands())
1933 for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i)
1934 bufferLEByte(cast<Constant>(CPV->getOperand(i)), 0, aggBuffer);
1938 if (const ConstantDataSequential *CDS =
1939 dyn_cast<ConstantDataSequential>(CPV)) {
1940 if (CDS->getNumElements())
1941 for (unsigned i = 0; i < CDS->getNumElements(); ++i)
1942 bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
1947 if (isa<ConstantStruct>(CPV)) {
1948 if (CPV->getNumOperands()) {
1949 StructType *ST = cast<StructType>(CPV->getType());
1950 for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) {
1952 Bytes = DL.getStructLayout(ST)->getElementOffset(0) +
1953 DL.getTypeAllocSize(ST) -
1954 DL.getStructLayout(ST)->getElementOffset(i);
1956 Bytes = DL.getStructLayout(ST)->getElementOffset(i + 1) -
1957 DL.getStructLayout(ST)->getElementOffset(i);
1958 bufferLEByte(cast<Constant>(CPV->getOperand(i)), Bytes, aggBuffer);
1963 llvm_unreachable("unsupported constant type in printAggregateConstant()");
1966 /// lowerConstantForGV - Return an MCExpr for the given Constant. This is mostly
1967 /// a copy from AsmPrinter::lowerConstant, except customized to only handle
1968 /// expressions that are representable in PTX and create
1969 /// NVPTXGenericMCSymbolRefExpr nodes for addrspacecast instructions.
1971 NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV, bool ProcessingGeneric) {
1972 MCContext &Ctx = OutContext;
1974 if (CV->isNullValue() || isa<UndefValue>(CV))
1975 return MCConstantExpr::create(0, Ctx);
1977 if (const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
1978 return MCConstantExpr::create(CI->getZExtValue(), Ctx);
1980 if (const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
1981 const MCSymbolRefExpr *Expr =
1982 MCSymbolRefExpr::create(getSymbol(GV), Ctx);
1983 if (ProcessingGeneric) {
1984 return NVPTXGenericMCSymbolRefExpr::create(Expr, Ctx);
1990 const ConstantExpr *CE = dyn_cast<ConstantExpr>(CV);
1992 llvm_unreachable("Unknown constant value to lower!");
1995 switch (CE->getOpcode()) {
1997 // If the code isn't optimized, there may be outstanding folding
1998 // opportunities. Attempt to fold the expression using DataLayout as a
1999 // last resort before giving up.
2000 if (Constant *C = ConstantFoldConstant(CE, getDataLayout()))
2002 return lowerConstantForGV(C, ProcessingGeneric);
2004 // Otherwise report the problem to the user.
2007 raw_string_ostream OS(S);
2008 OS << "Unsupported expression in static initializer: ";
2009 CE->printAsOperand(OS, /*PrintType=*/false,
2010 !MF ? nullptr : MF->getFunction().getParent());
2011 report_fatal_error(OS.str());
2014 case Instruction::AddrSpaceCast: {
2015 // Strip the addrspacecast and pass along the operand
2016 PointerType *DstTy = cast<PointerType>(CE->getType());
2017 if (DstTy->getAddressSpace() == 0) {
2018 return lowerConstantForGV(cast<const Constant>(CE->getOperand(0)), true);
2021 raw_string_ostream OS(S);
2022 OS << "Unsupported expression in static initializer: ";
2023 CE->printAsOperand(OS, /*PrintType=*/ false,
2024 !MF ? nullptr : MF->getFunction().getParent());
2025 report_fatal_error(OS.str());
2028 case Instruction::GetElementPtr: {
2029 const DataLayout &DL = getDataLayout();
2031 // Generate a symbolic expression for the byte address
2032 APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0);
2033 cast<GEPOperator>(CE)->accumulateConstantOffset(DL, OffsetAI);
2035 const MCExpr *Base = lowerConstantForGV(CE->getOperand(0),
2040 int64_t Offset = OffsetAI.getSExtValue();
2041 return MCBinaryExpr::createAdd(Base, MCConstantExpr::create(Offset, Ctx),
2045 case Instruction::Trunc:
2046 // We emit the value and depend on the assembler to truncate the generated
2047 // expression properly. This is important for differences between
2048 // blockaddress labels. Since the two labels are in the same function, it
2049 // is reasonable to treat their delta as a 32-bit value.
2051 case Instruction::BitCast:
2052 return lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
2054 case Instruction::IntToPtr: {
2055 const DataLayout &DL = getDataLayout();
2057 // Handle casts to pointers by changing them into casts to the appropriate
2058 // integer type. This promotes constant folding and simplifies this code.
2059 Constant *Op = CE->getOperand(0);
2060 Op = ConstantExpr::getIntegerCast(Op, DL.getIntPtrType(CV->getType()),
2062 return lowerConstantForGV(Op, ProcessingGeneric);
2065 case Instruction::PtrToInt: {
2066 const DataLayout &DL = getDataLayout();
2068 // Support only foldable casts to/from pointers that can be eliminated by
2069 // changing the pointer to the appropriately sized integer type.
2070 Constant *Op = CE->getOperand(0);
2071 Type *Ty = CE->getType();
2073 const MCExpr *OpExpr = lowerConstantForGV(Op, ProcessingGeneric);
2075 // We can emit the pointer value into this slot if the slot is an
2076 // integer slot equal to the size of the pointer.
2077 if (DL.getTypeAllocSize(Ty) == DL.getTypeAllocSize(Op->getType()))
2080 // Otherwise the pointer is smaller than the resultant integer, mask off
2081 // the high bits so we are sure to get a proper truncation if the input is
2083 unsigned InBits = DL.getTypeAllocSizeInBits(Op->getType());
2084 const MCExpr *MaskExpr = MCConstantExpr::create(~0ULL >> (64-InBits), Ctx);
2085 return MCBinaryExpr::createAnd(OpExpr, MaskExpr, Ctx);
2088 // The MC library also has a right-shift operator, but it isn't consistently
2089 // signed or unsigned between different targets.
2090 case Instruction::Add: {
2091 const MCExpr *LHS = lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
2092 const MCExpr *RHS = lowerConstantForGV(CE->getOperand(1), ProcessingGeneric);
2093 switch (CE->getOpcode()) {
2094 default: llvm_unreachable("Unknown binary operator constant cast expr");
2095 case Instruction::Add: return MCBinaryExpr::createAdd(LHS, RHS, Ctx);
2101 // Copy of MCExpr::print customized for NVPTX
2102 void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) {
2103 switch (Expr.getKind()) {
2104 case MCExpr::Target:
2105 return cast<MCTargetExpr>(&Expr)->printImpl(OS, MAI);
2106 case MCExpr::Constant:
2107 OS << cast<MCConstantExpr>(Expr).getValue();
2110 case MCExpr::SymbolRef: {
2111 const MCSymbolRefExpr &SRE = cast<MCSymbolRefExpr>(Expr);
2112 const MCSymbol &Sym = SRE.getSymbol();
2117 case MCExpr::Unary: {
2118 const MCUnaryExpr &UE = cast<MCUnaryExpr>(Expr);
2119 switch (UE.getOpcode()) {
2120 case MCUnaryExpr::LNot: OS << '!'; break;
2121 case MCUnaryExpr::Minus: OS << '-'; break;
2122 case MCUnaryExpr::Not: OS << '~'; break;
2123 case MCUnaryExpr::Plus: OS << '+'; break;
2125 printMCExpr(*UE.getSubExpr(), OS);
2129 case MCExpr::Binary: {
2130 const MCBinaryExpr &BE = cast<MCBinaryExpr>(Expr);
2132 // Only print parens around the LHS if it is non-trivial.
2133 if (isa<MCConstantExpr>(BE.getLHS()) || isa<MCSymbolRefExpr>(BE.getLHS()) ||
2134 isa<NVPTXGenericMCSymbolRefExpr>(BE.getLHS())) {
2135 printMCExpr(*BE.getLHS(), OS);
2138 printMCExpr(*BE.getLHS(), OS);
2142 switch (BE.getOpcode()) {
2143 case MCBinaryExpr::Add:
2144 // Print "X-42" instead of "X+-42".
2145 if (const MCConstantExpr *RHSC = dyn_cast<MCConstantExpr>(BE.getRHS())) {
2146 if (RHSC->getValue() < 0) {
2147 OS << RHSC->getValue();
2154 default: llvm_unreachable("Unhandled binary operator");
2157 // Only print parens around the LHS if it is non-trivial.
2158 if (isa<MCConstantExpr>(BE.getRHS()) || isa<MCSymbolRefExpr>(BE.getRHS())) {
2159 printMCExpr(*BE.getRHS(), OS);
2162 printMCExpr(*BE.getRHS(), OS);
2169 llvm_unreachable("Invalid expression kind!");
2172 /// PrintAsmOperand - Print out an operand for an inline asm expression.
2174 bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
2175 const char *ExtraCode, raw_ostream &O) {
2176 if (ExtraCode && ExtraCode[0]) {
2177 if (ExtraCode[1] != 0)
2178 return true; // Unknown modifier.
2180 switch (ExtraCode[0]) {
2182 // See if this is a generic print operand
2183 return AsmPrinter::PrintAsmOperand(MI, OpNo, ExtraCode, O);
2189 printOperand(MI, OpNo, O);
2194 bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI,
2196 const char *ExtraCode,
2198 if (ExtraCode && ExtraCode[0])
2199 return true; // Unknown modifier
2202 printMemOperand(MI, OpNo, O);
2208 void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum,
2210 const MachineOperand &MO = MI->getOperand(opNum);
2211 switch (MO.getType()) {
2212 case MachineOperand::MO_Register:
2213 if (Register::isPhysicalRegister(MO.getReg())) {
2214 if (MO.getReg() == NVPTX::VRDepot)
2215 O << DEPOTNAME << getFunctionNumber();
2217 O << NVPTXInstPrinter::getRegisterName(MO.getReg());
2219 emitVirtualRegister(MO.getReg(), O);
2223 case MachineOperand::MO_Immediate:
2227 case MachineOperand::MO_FPImmediate:
2228 printFPConstant(MO.getFPImm(), O);
2231 case MachineOperand::MO_GlobalAddress:
2232 PrintSymbolOperand(MO, O);
2235 case MachineOperand::MO_MachineBasicBlock:
2236 MO.getMBB()->getSymbol()->print(O, MAI);
2240 llvm_unreachable("Operand type not supported.");
2244 void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum,
2245 raw_ostream &O, const char *Modifier) {
2246 printOperand(MI, opNum, O);
2248 if (Modifier && strcmp(Modifier, "add") == 0) {
2250 printOperand(MI, opNum + 1, O);
2252 if (MI->getOperand(opNum + 1).isImm() &&
2253 MI->getOperand(opNum + 1).getImm() == 0)
2254 return; // don't print ',0' or '+0'
2256 printOperand(MI, opNum + 1, O);
2260 // Force static initialization.
2261 extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXAsmPrinter() {
2262 RegisterAsmPrinter<NVPTXAsmPrinter> X(getTheNVPTXTarget32());
2263 RegisterAsmPrinter<NVPTXAsmPrinter> Y(getTheNVPTXTarget64());