LLVM 19.0.0git
NVPTXAsmPrinter.cpp
Go to the documentation of this file.
1//===-- NVPTXAsmPrinter.cpp - NVPTX LLVM assembly writer ------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This file contains a printer that converts from our internal representation
10// of machine-dependent LLVM code to NVPTX assembly language.
11//
12//===----------------------------------------------------------------------===//
13
14#include "NVPTXAsmPrinter.h"
19#include "NVPTX.h"
20#include "NVPTXMCExpr.h"
22#include "NVPTXRegisterInfo.h"
23#include "NVPTXSubtarget.h"
24#include "NVPTXTargetMachine.h"
25#include "NVPTXUtilities.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"
35#include "llvm/ADT/StringRef.h"
36#include "llvm/ADT/Twine.h"
50#include "llvm/IR/Attributes.h"
51#include "llvm/IR/BasicBlock.h"
52#include "llvm/IR/Constant.h"
53#include "llvm/IR/Constants.h"
54#include "llvm/IR/DataLayout.h"
55#include "llvm/IR/DebugInfo.h"
57#include "llvm/IR/DebugLoc.h"
59#include "llvm/IR/Function.h"
60#include "llvm/IR/GlobalAlias.h"
61#include "llvm/IR/GlobalValue.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"
77#include "llvm/Support/Endian.h"
80#include "llvm/Support/Path.h"
86#include <cassert>
87#include <cstdint>
88#include <cstring>
89#include <new>
90#include <string>
91#include <utility>
92#include <vector>
93
94using namespace llvm;
95
96static cl::opt<bool>
97 LowerCtorDtor("nvptx-lower-global-ctor-dtor",
98 cl::desc("Lower GPU ctor / dtors to globals on the device."),
99 cl::init(false), cl::Hidden);
100
101#define DEPOTNAME "__local_depot"
102
103/// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
104/// depends.
105static void
108 if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(V))
109 Globals.insert(GV);
110 else {
111 if (const User *U = dyn_cast<User>(V)) {
112 for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) {
113 DiscoverDependentGlobals(U->getOperand(i), Globals);
114 }
115 }
116 }
117}
118
119/// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable
120/// instances to be emitted, but only after any dependents have been added
121/// first.s
122static void
127 // Have we already visited this one?
128 if (Visited.count(GV))
129 return;
130
131 // Do we have a circular dependency?
132 if (!Visiting.insert(GV).second)
133 report_fatal_error("Circular dependency found in global variable set");
134
135 // Make sure we visit all dependents first
137 for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i)
138 DiscoverDependentGlobals(GV->getOperand(i), Others);
139
140 for (const GlobalVariable *GV : Others)
141 VisitGlobalVariableForEmission(GV, Order, Visited, Visiting);
142
143 // Now we can visit ourself
144 Order.push_back(GV);
145 Visited.insert(GV);
146 Visiting.erase(GV);
147}
148
149void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {
150 NVPTX_MC::verifyInstructionPredicates(MI->getOpcode(),
151 getSubtargetInfo().getFeatureBits());
152
153 MCInst Inst;
154 lowerToMCInst(MI, Inst);
156}
157
158// Handle symbol backtracking for targets that do not support image handles
159bool NVPTXAsmPrinter::lowerImageHandleOperand(const MachineInstr *MI,
160 unsigned OpNo, MCOperand &MCOp) {
161 const MachineOperand &MO = MI->getOperand(OpNo);
162 const MCInstrDesc &MCID = MI->getDesc();
163
164 if (MCID.TSFlags & NVPTXII::IsTexFlag) {
165 // This is a texture fetch, so operand 4 is a texref and operand 5 is
166 // a samplerref
167 if (OpNo == 4 && MO.isImm()) {
168 lowerImageHandleSymbol(MO.getImm(), MCOp);
169 return true;
170 }
171 if (OpNo == 5 && MO.isImm() && !(MCID.TSFlags & NVPTXII::IsTexModeUnifiedFlag)) {
172 lowerImageHandleSymbol(MO.getImm(), MCOp);
173 return true;
174 }
175
176 return false;
177 } else if (MCID.TSFlags & NVPTXII::IsSuldMask) {
178 unsigned VecSize =
179 1 << (((MCID.TSFlags & NVPTXII::IsSuldMask) >> NVPTXII::IsSuldShift) - 1);
180
181 // For a surface load of vector size N, the Nth operand will be the surfref
182 if (OpNo == VecSize && MO.isImm()) {
183 lowerImageHandleSymbol(MO.getImm(), MCOp);
184 return true;
185 }
186
187 return false;
188 } else if (MCID.TSFlags & NVPTXII::IsSustFlag) {
189 // This is a surface store, so operand 0 is a surfref
190 if (OpNo == 0 && MO.isImm()) {
191 lowerImageHandleSymbol(MO.getImm(), MCOp);
192 return true;
193 }
194
195 return false;
196 } else if (MCID.TSFlags & NVPTXII::IsSurfTexQueryFlag) {
197 // This is a query, so operand 1 is a surfref/texref
198 if (OpNo == 1 && MO.isImm()) {
199 lowerImageHandleSymbol(MO.getImm(), MCOp);
200 return true;
201 }
202
203 return false;
204 }
205
206 return false;
207}
208
209void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) {
210 // Ewwww
212 NVPTXTargetMachine &nvTM = static_cast<NVPTXTargetMachine&>(TM);
214 const char *Sym = MFI->getImageHandleSymbol(Index);
215 StringRef SymName = nvTM.getStrPool().save(Sym);
216 MCOp = GetSymbolRef(OutContext.getOrCreateSymbol(SymName));
217}
218
219void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
220 OutMI.setOpcode(MI->getOpcode());
221 // Special: Do not mangle symbol operand of CALL_PROTOTYPE
222 if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
223 const MachineOperand &MO = MI->getOperand(0);
224 OutMI.addOperand(GetSymbolRef(
226 return;
227 }
228
229 const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
230 for (unsigned i = 0, e = MI->getNumOperands(); i != e; ++i) {
231 const MachineOperand &MO = MI->getOperand(i);
232
233 MCOperand MCOp;
234 if (!STI.hasImageHandles()) {
235 if (lowerImageHandleOperand(MI, i, MCOp)) {
236 OutMI.addOperand(MCOp);
237 continue;
238 }
239 }
240
241 if (lowerOperand(MO, MCOp))
242 OutMI.addOperand(MCOp);
243 }
244}
245
246bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO,
247 MCOperand &MCOp) {
248 switch (MO.getType()) {
249 default: llvm_unreachable("unknown operand type");
251 MCOp = MCOperand::createReg(encodeVirtualRegister(MO.getReg()));
252 break;
254 MCOp = MCOperand::createImm(MO.getImm());
255 break;
258 MO.getMBB()->getSymbol(), OutContext));
259 break;
261 MCOp = GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
262 break;
264 MCOp = GetSymbolRef(getSymbol(MO.getGlobal()));
265 break;
267 const ConstantFP *Cnt = MO.getFPImm();
268 const APFloat &Val = Cnt->getValueAPF();
269
270 switch (Cnt->getType()->getTypeID()) {
271 default: report_fatal_error("Unsupported FP type"); break;
272 case Type::HalfTyID:
275 break;
276 case Type::BFloatTyID:
279 break;
280 case Type::FloatTyID:
283 break;
284 case Type::DoubleTyID:
287 break;
288 }
289 break;
290 }
291 }
292 return true;
293}
294
295unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
297 const TargetRegisterClass *RC = MRI->getRegClass(Reg);
298
299 DenseMap<unsigned, unsigned> &RegMap = VRegMapping[RC];
300 unsigned RegNum = RegMap[Reg];
301
302 // Encode the register class in the upper 4 bits
303 // Must be kept in sync with NVPTXInstPrinter::printRegName
304 unsigned Ret = 0;
305 if (RC == &NVPTX::Int1RegsRegClass) {
306 Ret = (1 << 28);
307 } else if (RC == &NVPTX::Int16RegsRegClass) {
308 Ret = (2 << 28);
309 } else if (RC == &NVPTX::Int32RegsRegClass) {
310 Ret = (3 << 28);
311 } else if (RC == &NVPTX::Int64RegsRegClass) {
312 Ret = (4 << 28);
313 } else if (RC == &NVPTX::Float32RegsRegClass) {
314 Ret = (5 << 28);
315 } else if (RC == &NVPTX::Float64RegsRegClass) {
316 Ret = (6 << 28);
317 } else {
318 report_fatal_error("Bad register class");
319 }
320
321 // Insert the vreg number
322 Ret |= (RegNum & 0x0FFFFFFF);
323 return Ret;
324 } else {
325 // Some special-use registers are actually physical registers.
326 // Encode this as the register class ID of 0 and the real register ID.
327 return Reg & 0x0FFFFFFF;
328 }
329}
330
331MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
332 const MCExpr *Expr;
334 OutContext);
335 return MCOperand::createExpr(Expr);
336}
337
338static bool ShouldPassAsArray(Type *Ty) {
339 return Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128) ||
340 Ty->isHalfTy() || Ty->isBFloatTy();
341}
342
343void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
344 const DataLayout &DL = getDataLayout();
346 const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
347
348 Type *Ty = F->getReturnType();
349
350 bool isABI = (STI.getSmVersion() >= 20);
351
352 if (Ty->getTypeID() == Type::VoidTyID)
353 return;
354 O << " (";
355
356 if (isABI) {
357 if ((Ty->isFloatingPointTy() || Ty->isIntegerTy()) &&
358 !ShouldPassAsArray(Ty)) {
359 unsigned size = 0;
360 if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
361 size = ITy->getBitWidth();
362 } else {
363 assert(Ty->isFloatingPointTy() && "Floating point type expected here");
365 }
367 O << ".param .b" << size << " func_retval0";
368 } else if (isa<PointerType>(Ty)) {
369 O << ".param .b" << TLI->getPointerTy(DL).getSizeInBits()
370 << " func_retval0";
371 } else if (ShouldPassAsArray(Ty)) {
372 unsigned totalsz = DL.getTypeAllocSize(Ty);
373 unsigned retAlignment = 0;
374 if (!getAlign(*F, 0, retAlignment))
375 retAlignment = TLI->getFunctionParamOptimizedAlign(F, Ty, DL).value();
376 O << ".param .align " << retAlignment << " .b8 func_retval0[" << totalsz
377 << "]";
378 } else
379 llvm_unreachable("Unknown return type");
380 } else {
381 SmallVector<EVT, 16> vtparts;
382 ComputeValueVTs(*TLI, DL, Ty, vtparts);
383 unsigned idx = 0;
384 for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
385 unsigned elems = 1;
386 EVT elemtype = vtparts[i];
387 if (vtparts[i].isVector()) {
388 elems = vtparts[i].getVectorNumElements();
389 elemtype = vtparts[i].getVectorElementType();
390 }
391
392 for (unsigned j = 0, je = elems; j != je; ++j) {
393 unsigned sz = elemtype.getSizeInBits();
394 if (elemtype.isInteger())
396 O << ".reg .b" << sz << " func_retval" << idx;
397 if (j < je - 1)
398 O << ", ";
399 ++idx;
400 }
401 if (i < e - 1)
402 O << ", ";
403 }
404 }
405 O << ") ";
406}
407
408void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
409 raw_ostream &O) {
410 const Function &F = MF.getFunction();
411 printReturnValStr(&F, O);
412}
413
414// Return true if MBB is the header of a loop marked with
415// llvm.loop.unroll.disable or llvm.loop.unroll.count=1.
416bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
417 const MachineBasicBlock &MBB) const {
418 MachineLoopInfo &LI = getAnalysis<MachineLoopInfo>();
419 // We insert .pragma "nounroll" only to the loop header.
420 if (!LI.isLoopHeader(&MBB))
421 return false;
422
423 // llvm.loop.unroll.disable is marked on the back edges of a loop. Therefore,
424 // we iterate through each back edge of the loop with header MBB, and check
425 // whether its metadata contains llvm.loop.unroll.disable.
426 for (const MachineBasicBlock *PMBB : MBB.predecessors()) {
427 if (LI.getLoopFor(PMBB) != LI.getLoopFor(&MBB)) {
428 // Edges from other loops to MBB are not back edges.
429 continue;
430 }
431 if (const BasicBlock *PBB = PMBB->getBasicBlock()) {
432 if (MDNode *LoopID =
433 PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
434 if (GetUnrollMetadata(LoopID, "llvm.loop.unroll.disable"))
435 return true;
436 if (MDNode *UnrollCountMD =
437 GetUnrollMetadata(LoopID, "llvm.loop.unroll.count")) {
438 if (mdconst::extract<ConstantInt>(UnrollCountMD->getOperand(1))
439 ->isOne())
440 return true;
441 }
442 }
443 }
444 }
445 return false;
446}
447
448void NVPTXAsmPrinter::emitBasicBlockStart(const MachineBasicBlock &MBB) {
450 if (isLoopHeaderOfNoUnroll(MBB))
451 OutStreamer->emitRawText(StringRef("\t.pragma \"nounroll\";\n"));
452}
453
454void NVPTXAsmPrinter::emitFunctionEntryLabel() {
457
458 if (!GlobalsEmitted) {
459 emitGlobals(*MF->getFunction().getParent());
460 GlobalsEmitted = true;
461 }
462
463 // Set up
464 MRI = &MF->getRegInfo();
465 F = &MF->getFunction();
466 emitLinkageDirective(F, O);
467 if (isKernelFunction(*F))
468 O << ".entry ";
469 else {
470 O << ".func ";
471 printReturnValStr(*MF, O);
472 }
473
475
476 emitFunctionParamList(F, O);
477 O << "\n";
478
479 if (isKernelFunction(*F))
480 emitKernelFunctionDirectives(*F, O);
481
483 O << ".noreturn";
484
485 OutStreamer->emitRawText(O.str());
486
487 VRegMapping.clear();
488 // Emit open brace for function body.
489 OutStreamer->emitRawText(StringRef("{\n"));
490 setAndEmitFunctionVirtualRegisters(*MF);
491 // Emit initial .loc debug directive for correct relocation symbol data.
492 if (const DISubprogram *SP = MF->getFunction().getSubprogram()) {
493 assert(SP->getUnit());
494 if (!SP->getUnit()->isDebugDirectivesOnly() && MMI && MMI->hasDebugInfo())
496 }
497}
498
500 bool Result = AsmPrinter::runOnMachineFunction(F);
501 // Emit closing brace for the body of function F.
502 // The closing brace must be emitted here because we need to emit additional
503 // debug labels/data after the last basic block.
504 // We need to emit the closing brace here because we don't have function that
505 // finished emission of the function body.
506 OutStreamer->emitRawText(StringRef("}\n"));
507 return Result;
508}
509
510void NVPTXAsmPrinter::emitFunctionBodyStart() {
512 raw_svector_ostream O(Str);
513 emitDemotedVars(&MF->getFunction(), O);
514 OutStreamer->emitRawText(O.str());
515}
516
517void NVPTXAsmPrinter::emitFunctionBodyEnd() {
518 VRegMapping.clear();
519}
520
524 return OutContext.getOrCreateSymbol(Str);
525}
526
527void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const {
528 Register RegNo = MI->getOperand(0).getReg();
529 if (RegNo.isVirtual()) {
530 OutStreamer->AddComment(Twine("implicit-def: ") +
532 } else {
533 const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
534 OutStreamer->AddComment(Twine("implicit-def: ") +
535 STI.getRegisterInfo()->getName(RegNo));
536 }
537 OutStreamer->addBlankLine();
538}
539
540void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F,
541 raw_ostream &O) const {
542 // If the NVVM IR has some of reqntid* specified, then output
543 // the reqntid directive, and set the unspecified ones to 1.
544 // If none of Reqntid* is specified, don't output reqntid directive.
545 unsigned Reqntidx, Reqntidy, Reqntidz;
546 Reqntidx = Reqntidy = Reqntidz = 1;
547 bool ReqSpecified = false;
548 ReqSpecified |= getReqNTIDx(F, Reqntidx);
549 ReqSpecified |= getReqNTIDy(F, Reqntidy);
550 ReqSpecified |= getReqNTIDz(F, Reqntidz);
551
552 if (ReqSpecified)
553 O << ".reqntid " << Reqntidx << ", " << Reqntidy << ", " << Reqntidz
554 << "\n";
555
556 // If the NVVM IR has some of maxntid* specified, then output
557 // the maxntid directive, and set the unspecified ones to 1.
558 // If none of maxntid* is specified, don't output maxntid directive.
559 unsigned Maxntidx, Maxntidy, Maxntidz;
560 Maxntidx = Maxntidy = Maxntidz = 1;
561 bool MaxSpecified = false;
562 MaxSpecified |= getMaxNTIDx(F, Maxntidx);
563 MaxSpecified |= getMaxNTIDy(F, Maxntidy);
564 MaxSpecified |= getMaxNTIDz(F, Maxntidz);
565
566 if (MaxSpecified)
567 O << ".maxntid " << Maxntidx << ", " << Maxntidy << ", " << Maxntidz
568 << "\n";
569
570 unsigned Mincta = 0;
571 if (getMinCTASm(F, Mincta))
572 O << ".minnctapersm " << Mincta << "\n";
573
574 unsigned Maxnreg = 0;
575 if (getMaxNReg(F, Maxnreg))
576 O << ".maxnreg " << Maxnreg << "\n";
577
578 // .maxclusterrank directive requires SM_90 or higher, make sure that we
579 // filter it out for lower SM versions, as it causes a hard ptxas crash.
580 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
581 const auto *STI = static_cast<const NVPTXSubtarget *>(NTM.getSubtargetImpl());
582 unsigned Maxclusterrank = 0;
583 if (getMaxClusterRank(F, Maxclusterrank) && STI->getSmVersion() >= 90)
584 O << ".maxclusterrank " << Maxclusterrank << "\n";
585}
586
587std::string NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const {
588 const TargetRegisterClass *RC = MRI->getRegClass(Reg);
589
590 std::string Name;
591 raw_string_ostream NameStr(Name);
592
593 VRegRCMap::const_iterator I = VRegMapping.find(RC);
594 assert(I != VRegMapping.end() && "Bad register class");
595 const DenseMap<unsigned, unsigned> &RegMap = I->second;
596
597 VRegMap::const_iterator VI = RegMap.find(Reg);
598 assert(VI != RegMap.end() && "Bad virtual register");
599 unsigned MappedVR = VI->second;
600
601 NameStr << getNVPTXRegClassStr(RC) << MappedVR;
602
603 NameStr.flush();
604 return Name;
605}
606
607void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr,
608 raw_ostream &O) {
609 O << getVirtualRegisterName(vr);
610}
611
612void NVPTXAsmPrinter::emitAliasDeclaration(const GlobalAlias *GA,
613 raw_ostream &O) {
614 const Function *F = dyn_cast_or_null<Function>(GA->getAliaseeObject());
615 if (!F || isKernelFunction(*F) || F->isDeclaration())
617 "NVPTX aliasee must be a non-kernel function definition");
618
619 if (GA->hasLinkOnceLinkage() || GA->hasWeakLinkage() ||
621 report_fatal_error("NVPTX aliasee must not be '.weak'");
622
623 emitDeclarationWithName(F, getSymbol(GA), O);
624}
625
626void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) {
627 emitDeclarationWithName(F, getSymbol(F), O);
628}
629
630void NVPTXAsmPrinter::emitDeclarationWithName(const Function *F, MCSymbol *S,
631 raw_ostream &O) {
632 emitLinkageDirective(F, O);
633 if (isKernelFunction(*F))
634 O << ".entry ";
635 else
636 O << ".func ";
637 printReturnValStr(F, O);
638 S->print(O, MAI);
639 O << "\n";
640 emitFunctionParamList(F, O);
641 O << "\n";
643 O << ".noreturn";
644 O << ";\n";
645}
646
647static bool usedInGlobalVarDef(const Constant *C) {
648 if (!C)
649 return false;
650
651 if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(C)) {
652 return GV->getName() != "llvm.used";
653 }
654
655 for (const User *U : C->users())
656 if (const Constant *C = dyn_cast<Constant>(U))
658 return true;
659
660 return false;
661}
662
663static bool usedInOneFunc(const User *U, Function const *&oneFunc) {
664 if (const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
665 if (othergv->getName() == "llvm.used")
666 return true;
667 }
668
669 if (const Instruction *instr = dyn_cast<Instruction>(U)) {
670 if (instr->getParent() && instr->getParent()->getParent()) {
671 const Function *curFunc = instr->getParent()->getParent();
672 if (oneFunc && (curFunc != oneFunc))
673 return false;
674 oneFunc = curFunc;
675 return true;
676 } else
677 return false;
678 }
679
680 for (const User *UU : U->users())
681 if (!usedInOneFunc(UU, oneFunc))
682 return false;
683
684 return true;
685}
686
687/* Find out if a global variable can be demoted to local scope.
688 * Currently, this is valid for CUDA shared variables, which have local
689 * scope and global lifetime. So the conditions to check are :
690 * 1. Is the global variable in shared address space?
691 * 2. Does it have local linkage?
692 * 3. Is the global variable referenced only in one function?
693 */
694static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) {
695 if (!gv->hasLocalLinkage())
696 return false;
697 PointerType *Pty = gv->getType();
698 if (Pty->getAddressSpace() != ADDRESS_SPACE_SHARED)
699 return false;
700
701 const Function *oneFunc = nullptr;
702
703 bool flag = usedInOneFunc(gv, oneFunc);
704 if (!flag)
705 return false;
706 if (!oneFunc)
707 return false;
708 f = oneFunc;
709 return true;
710}
711
712static bool useFuncSeen(const Constant *C,
714 for (const User *U : C->users()) {
715 if (const Constant *cu = dyn_cast<Constant>(U)) {
716 if (useFuncSeen(cu, seenMap))
717 return true;
718 } else if (const Instruction *I = dyn_cast<Instruction>(U)) {
719 const BasicBlock *bb = I->getParent();
720 if (!bb)
721 continue;
722 const Function *caller = bb->getParent();
723 if (!caller)
724 continue;
725 if (seenMap.contains(caller))
726 return true;
727 }
728 }
729 return false;
730}
731
732void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {
734 for (const Function &F : M) {
735 if (F.getAttributes().hasFnAttr("nvptx-libcall-callee")) {
736 emitDeclaration(&F, O);
737 continue;
738 }
739
740 if (F.isDeclaration()) {
741 if (F.use_empty())
742 continue;
743 if (F.getIntrinsicID())
744 continue;
745 emitDeclaration(&F, O);
746 continue;
747 }
748 for (const User *U : F.users()) {
749 if (const Constant *C = dyn_cast<Constant>(U)) {
750 if (usedInGlobalVarDef(C)) {
751 // The use is in the initialization of a global variable
752 // that is a function pointer, so print a declaration
753 // for the original function
754 emitDeclaration(&F, O);
755 break;
756 }
757 // Emit a declaration of this function if the function that
758 // uses this constant expr has already been seen.
759 if (useFuncSeen(C, seenMap)) {
760 emitDeclaration(&F, O);
761 break;
762 }
763 }
764
765 if (!isa<Instruction>(U))
766 continue;
767 const Instruction *instr = cast<Instruction>(U);
768 const BasicBlock *bb = instr->getParent();
769 if (!bb)
770 continue;
771 const Function *caller = bb->getParent();
772 if (!caller)
773 continue;
774
775 // If a caller has already been seen, then the caller is
776 // appearing in the module before the callee. so print out
777 // a declaration for the callee.
778 if (seenMap.contains(caller)) {
779 emitDeclaration(&F, O);
780 break;
781 }
782 }
783 seenMap[&F] = true;
784 }
785 for (const GlobalAlias &GA : M.aliases())
786 emitAliasDeclaration(&GA, O);
787}
788
790 if (!GV) return true;
791 const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());
792 if (!InitList) return true; // Not an array; we don't know how to parse.
793 return InitList->getNumOperands() == 0;
794}
795
796void NVPTXAsmPrinter::emitStartOfAsmFile(Module &M) {
797 // Construct a default subtarget off of the TargetMachine defaults. The
798 // rest of NVPTX isn't friendly to change subtargets per function and
799 // so the default TargetMachine will have all of the options.
800 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
801 const auto* STI = static_cast<const NVPTXSubtarget*>(NTM.getSubtargetImpl());
802 SmallString<128> Str1;
803 raw_svector_ostream OS1(Str1);
804
805 // Emit header before any dwarf directives are emitted below.
806 emitHeader(M, OS1, *STI);
807 OutStreamer->emitRawText(OS1.str());
808}
809
811 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
812 const NVPTXSubtarget &STI =
813 *static_cast<const NVPTXSubtarget *>(NTM.getSubtargetImpl());
814 if (M.alias_size() && (STI.getPTXVersion() < 63 || STI.getSmVersion() < 30))
815 report_fatal_error(".alias requires PTX version >= 6.3 and sm_30");
816
817 // OpenMP supports NVPTX global constructors and destructors.
818 bool IsOpenMP = M.getModuleFlag("openmp") != nullptr;
819
820 if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_ctors")) &&
821 !LowerCtorDtor && !IsOpenMP) {
823 "Module has a nontrivial global ctor, which NVPTX does not support.");
824 return true; // error
825 }
826 if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_dtors")) &&
827 !LowerCtorDtor && !IsOpenMP) {
829 "Module has a nontrivial global dtor, which NVPTX does not support.");
830 return true; // error
831 }
832
833 // We need to call the parent's one explicitly.
834 bool Result = AsmPrinter::doInitialization(M);
835
836 GlobalsEmitted = false;
837
838 return Result;
839}
840
841void NVPTXAsmPrinter::emitGlobals(const Module &M) {
842 SmallString<128> Str2;
843 raw_svector_ostream OS2(Str2);
844
845 emitDeclarations(M, OS2);
846
847 // As ptxas does not support forward references of globals, we need to first
848 // sort the list of module-level globals in def-use order. We visit each
849 // global variable in order, and ensure that we emit it *after* its dependent
850 // globals. We use a little extra memory maintaining both a set and a list to
851 // have fast searches while maintaining a strict ordering.
855
856 // Visit each global variable, in order
857 for (const GlobalVariable &I : M.globals())
858 VisitGlobalVariableForEmission(&I, Globals, GVVisited, GVVisiting);
859
860 assert(GVVisited.size() == M.global_size() && "Missed a global variable");
861 assert(GVVisiting.size() == 0 && "Did not fully process a global variable");
862
863 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
864 const NVPTXSubtarget &STI =
865 *static_cast<const NVPTXSubtarget *>(NTM.getSubtargetImpl());
866
867 // Print out module-level global variables in proper order
868 for (unsigned i = 0, e = Globals.size(); i != e; ++i)
869 printModuleLevelGV(Globals[i], OS2, /*processDemoted=*/false, STI);
870
871 OS2 << '\n';
872
873 OutStreamer->emitRawText(OS2.str());
874}
875
876void NVPTXAsmPrinter::emitGlobalAlias(const Module &M, const GlobalAlias &GA) {
879
880 MCSymbol *Name = getSymbol(&GA);
881
882 OS << ".alias " << Name->getName() << ", " << GA.getAliaseeObject()->getName()
883 << ";\n";
884
885 OutStreamer->emitRawText(OS.str());
886}
887
888void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
889 const NVPTXSubtarget &STI) {
890 O << "//\n";
891 O << "// Generated by LLVM NVPTX Back-End\n";
892 O << "//\n";
893 O << "\n";
894
895 unsigned PTXVersion = STI.getPTXVersion();
896 O << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n";
897
898 O << ".target ";
899 O << STI.getTargetName();
900
901 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
902 if (NTM.getDrvInterface() == NVPTX::NVCL)
903 O << ", texmode_independent";
904
905 bool HasFullDebugInfo = false;
906 for (DICompileUnit *CU : M.debug_compile_units()) {
907 switch(CU->getEmissionKind()) {
910 break;
913 HasFullDebugInfo = true;
914 break;
915 }
916 if (HasFullDebugInfo)
917 break;
918 }
919 if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo)
920 O << ", debug";
921
922 O << "\n";
923
924 O << ".address_size ";
925 if (NTM.is64Bit())
926 O << "64";
927 else
928 O << "32";
929 O << "\n";
930
931 O << "\n";
932}
933
935 bool HasDebugInfo = MMI && MMI->hasDebugInfo();
936
937 // If we did not emit any functions, then the global declarations have not
938 // yet been emitted.
939 if (!GlobalsEmitted) {
940 emitGlobals(M);
941 GlobalsEmitted = true;
942 }
943
944 // call doFinalization
945 bool ret = AsmPrinter::doFinalization(M);
946
948
949 auto *TS =
950 static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer());
951 // Close the last emitted section
952 if (HasDebugInfo) {
953 TS->closeLastSection();
954 // Emit empty .debug_loc section for better support of the empty files.
955 OutStreamer->emitRawText("\t.section\t.debug_loc\t{\t}");
956 }
957
958 // Output last DWARF .file directives, if any.
959 TS->outputDwarfFileDirectives();
960
961 return ret;
962}
963
964// This function emits appropriate linkage directives for
965// functions and global variables.
966//
967// extern function declaration -> .extern
968// extern function definition -> .visible
969// external global variable with init -> .visible
970// external without init -> .extern
971// appending -> not allowed, assert.
972// for any linkage other than
973// internal, private, linker_private,
974// linker_private_weak, linker_private_weak_def_auto,
975// we emit -> .weak.
976
977void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V,
978 raw_ostream &O) {
979 if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA) {
980 if (V->hasExternalLinkage()) {
981 if (isa<GlobalVariable>(V)) {
982 const GlobalVariable *GVar = cast<GlobalVariable>(V);
983 if (GVar) {
984 if (GVar->hasInitializer())
985 O << ".visible ";
986 else
987 O << ".extern ";
988 }
989 } else if (V->isDeclaration())
990 O << ".extern ";
991 else
992 O << ".visible ";
993 } else if (V->hasAppendingLinkage()) {
994 std::string msg;
995 msg.append("Error: ");
996 msg.append("Symbol ");
997 if (V->hasName())
998 msg.append(std::string(V->getName()));
999 msg.append("has unsupported appending linkage type");
1000 llvm_unreachable(msg.c_str());
1001 } else if (!V->hasInternalLinkage() &&
1002 !V->hasPrivateLinkage()) {
1003 O << ".weak ";
1004 }
1005 }
1006}
1007
1008void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
1009 raw_ostream &O, bool processDemoted,
1010 const NVPTXSubtarget &STI) {
1011 // Skip meta data
1012 if (GVar->hasSection()) {
1013 if (GVar->getSection() == "llvm.metadata")
1014 return;
1015 }
1016
1017 // Skip LLVM intrinsic global variables
1018 if (GVar->getName().starts_with("llvm.") ||
1019 GVar->getName().starts_with("nvvm."))
1020 return;
1021
1022 const DataLayout &DL = getDataLayout();
1023
1024 // GlobalVariables are always constant pointers themselves.
1025 PointerType *PTy = GVar->getType();
1026 Type *ETy = GVar->getValueType();
1027
1028 if (GVar->hasExternalLinkage()) {
1029 if (GVar->hasInitializer())
1030 O << ".visible ";
1031 else
1032 O << ".extern ";
1033 } else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() ||
1035 GVar->hasCommonLinkage()) {
1036 O << ".weak ";
1037 }
1038
1039 if (isTexture(*GVar)) {
1040 O << ".global .texref " << getTextureName(*GVar) << ";\n";
1041 return;
1042 }
1043
1044 if (isSurface(*GVar)) {
1045 O << ".global .surfref " << getSurfaceName(*GVar) << ";\n";
1046 return;
1047 }
1048
1049 if (GVar->isDeclaration()) {
1050 // (extern) declarations, no definition or initializer
1051 // Currently the only known declaration is for an automatic __local
1052 // (.shared) promoted to global.
1053 emitPTXGlobalVariable(GVar, O, STI);
1054 O << ";\n";
1055 return;
1056 }
1057
1058 if (isSampler(*GVar)) {
1059 O << ".global .samplerref " << getSamplerName(*GVar);
1060
1061 const Constant *Initializer = nullptr;
1062 if (GVar->hasInitializer())
1063 Initializer = GVar->getInitializer();
1064 const ConstantInt *CI = nullptr;
1065 if (Initializer)
1066 CI = dyn_cast<ConstantInt>(Initializer);
1067 if (CI) {
1068 unsigned sample = CI->getZExtValue();
1069
1070 O << " = { ";
1071
1072 for (int i = 0,
1073 addr = ((sample & __CLK_ADDRESS_MASK) >> __CLK_ADDRESS_BASE);
1074 i < 3; i++) {
1075 O << "addr_mode_" << i << " = ";
1076 switch (addr) {
1077 case 0:
1078 O << "wrap";
1079 break;
1080 case 1:
1081 O << "clamp_to_border";
1082 break;
1083 case 2:
1084 O << "clamp_to_edge";
1085 break;
1086 case 3:
1087 O << "wrap";
1088 break;
1089 case 4:
1090 O << "mirror";
1091 break;
1092 }
1093 O << ", ";
1094 }
1095 O << "filter_mode = ";
1096 switch ((sample & __CLK_FILTER_MASK) >> __CLK_FILTER_BASE) {
1097 case 0:
1098 O << "nearest";
1099 break;
1100 case 1:
1101 O << "linear";
1102 break;
1103 case 2:
1104 llvm_unreachable("Anisotropic filtering is not supported");
1105 default:
1106 O << "nearest";
1107 break;
1108 }
1109 if (!((sample & __CLK_NORMALIZED_MASK) >> __CLK_NORMALIZED_BASE)) {
1110 O << ", force_unnormalized_coords = 1";
1111 }
1112 O << " }";
1113 }
1114
1115 O << ";\n";
1116 return;
1117 }
1118
1119 if (GVar->hasPrivateLinkage()) {
1120 if (strncmp(GVar->getName().data(), "unrollpragma", 12) == 0)
1121 return;
1122
1123 // FIXME - need better way (e.g. Metadata) to avoid generating this global
1124 if (strncmp(GVar->getName().data(), "filename", 8) == 0)
1125 return;
1126 if (GVar->use_empty())
1127 return;
1128 }
1129
1130 const Function *demotedFunc = nullptr;
1131 if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) {
1132 O << "// " << GVar->getName() << " has been demoted\n";
1133 if (localDecls.find(demotedFunc) != localDecls.end())
1134 localDecls[demotedFunc].push_back(GVar);
1135 else {
1136 std::vector<const GlobalVariable *> temp;
1137 temp.push_back(GVar);
1138 localDecls[demotedFunc] = temp;
1139 }
1140 return;
1141 }
1142
1143 O << ".";
1144 emitPTXAddressSpace(PTy->getAddressSpace(), O);
1145
1146 if (isManaged(*GVar)) {
1147 if (STI.getPTXVersion() < 40 || STI.getSmVersion() < 30) {
1149 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1150 }
1151 O << " .attribute(.managed)";
1152 }
1153
1154 if (MaybeAlign A = GVar->getAlign())
1155 O << " .align " << A->value();
1156 else
1157 O << " .align " << (int)DL.getPrefTypeAlign(ETy).value();
1158
1159 if (ETy->isFloatingPointTy() || ETy->isPointerTy() ||
1160 (ETy->isIntegerTy() && ETy->getScalarSizeInBits() <= 64)) {
1161 O << " .";
1162 // Special case: ABI requires that we use .u8 for predicates
1163 if (ETy->isIntegerTy(1))
1164 O << "u8";
1165 else
1166 O << getPTXFundamentalTypeStr(ETy, false);
1167 O << " ";
1168 getSymbol(GVar)->print(O, MAI);
1169
1170 // Ptx allows variable initilization only for constant and global state
1171 // spaces.
1172 if (GVar->hasInitializer()) {
1173 if ((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1174 (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) {
1175 const Constant *Initializer = GVar->getInitializer();
1176 // 'undef' is treated as there is no value specified.
1177 if (!Initializer->isNullValue() && !isa<UndefValue>(Initializer)) {
1178 O << " = ";
1179 printScalarConstant(Initializer, O);
1180 }
1181 } else {
1182 // The frontend adds zero-initializer to device and constant variables
1183 // that don't have an initial value, and UndefValue to shared
1184 // variables, so skip warning for this case.
1185 if (!GVar->getInitializer()->isNullValue() &&
1186 !isa<UndefValue>(GVar->getInitializer())) {
1187 report_fatal_error("initial value of '" + GVar->getName() +
1188 "' is not allowed in addrspace(" +
1189 Twine(PTy->getAddressSpace()) + ")");
1190 }
1191 }
1192 }
1193 } else {
1194 uint64_t ElementSize = 0;
1195
1196 // Although PTX has direct support for struct type and array type and
1197 // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
1198 // targets that support these high level field accesses. Structs, arrays
1199 // and vectors are lowered into arrays of bytes.
1200 switch (ETy->getTypeID()) {
1201 case Type::IntegerTyID: // Integers larger than 64 bits
1202 case Type::StructTyID:
1203 case Type::ArrayTyID:
1205 ElementSize = DL.getTypeStoreSize(ETy);
1206 // Ptx allows variable initilization only for constant and
1207 // global state spaces.
1208 if (((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1209 (PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) &&
1210 GVar->hasInitializer()) {
1211 const Constant *Initializer = GVar->getInitializer();
1212 if (!isa<UndefValue>(Initializer) && !Initializer->isNullValue()) {
1213 AggBuffer aggBuffer(ElementSize, *this);
1214 bufferAggregateConstant(Initializer, &aggBuffer);
1215 if (aggBuffer.numSymbols()) {
1216 unsigned int ptrSize = MAI->getCodePointerSize();
1217 if (ElementSize % ptrSize ||
1218 !aggBuffer.allSymbolsAligned(ptrSize)) {
1219 // Print in bytes and use the mask() operator for pointers.
1220 if (!STI.hasMaskOperator())
1222 "initialized packed aggregate with pointers '" +
1223 GVar->getName() +
1224 "' requires at least PTX ISA version 7.1");
1225 O << " .u8 ";
1226 getSymbol(GVar)->print(O, MAI);
1227 O << "[" << ElementSize << "] = {";
1228 aggBuffer.printBytes(O);
1229 O << "}";
1230 } else {
1231 O << " .u" << ptrSize * 8 << " ";
1232 getSymbol(GVar)->print(O, MAI);
1233 O << "[" << ElementSize / ptrSize << "] = {";
1234 aggBuffer.printWords(O);
1235 O << "}";
1236 }
1237 } else {
1238 O << " .b8 ";
1239 getSymbol(GVar)->print(O, MAI);
1240 O << "[" << ElementSize << "] = {";
1241 aggBuffer.printBytes(O);
1242 O << "}";
1243 }
1244 } else {
1245 O << " .b8 ";
1246 getSymbol(GVar)->print(O, MAI);
1247 if (ElementSize) {
1248 O << "[";
1249 O << ElementSize;
1250 O << "]";
1251 }
1252 }
1253 } else {
1254 O << " .b8 ";
1255 getSymbol(GVar)->print(O, MAI);
1256 if (ElementSize) {
1257 O << "[";
1258 O << ElementSize;
1259 O << "]";
1260 }
1261 }
1262 break;
1263 default:
1264 llvm_unreachable("type not supported yet");
1265 }
1266 }
1267 O << ";\n";
1268}
1269
1270void NVPTXAsmPrinter::AggBuffer::printSymbol(unsigned nSym, raw_ostream &os) {
1271 const Value *v = Symbols[nSym];
1272 const Value *v0 = SymbolsBeforeStripping[nSym];
1273 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
1274 MCSymbol *Name = AP.getSymbol(GVar);
1275 PointerType *PTy = dyn_cast<PointerType>(v0->getType());
1276 // Is v0 a generic pointer?
1277 bool isGenericPointer = PTy && PTy->getAddressSpace() == 0;
1278 if (EmitGeneric && isGenericPointer && !isa<Function>(v)) {
1279 os << "generic(";
1280 Name->print(os, AP.MAI);
1281 os << ")";
1282 } else {
1283 Name->print(os, AP.MAI);
1284 }
1285 } else if (const ConstantExpr *CExpr = dyn_cast<ConstantExpr>(v0)) {
1286 const MCExpr *Expr = AP.lowerConstantForGV(cast<Constant>(CExpr), false);
1287 AP.printMCExpr(*Expr, os);
1288 } else
1289 llvm_unreachable("symbol type unknown");
1290}
1291
1292void NVPTXAsmPrinter::AggBuffer::printBytes(raw_ostream &os) {
1293 unsigned int ptrSize = AP.MAI->getCodePointerSize();
1294 // Do not emit trailing zero initializers. They will be zero-initialized by
1295 // ptxas. This saves on both space requirements for the generated PTX and on
1296 // memory use by ptxas. (See:
1297 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#global-state-space)
1298 unsigned int InitializerCount = size;
1299 // TODO: symbols make this harder, but it would still be good to trim trailing
1300 // 0s for aggs with symbols as well.
1301 if (numSymbols() == 0)
1302 while (InitializerCount >= 1 && !buffer[InitializerCount - 1])
1303 InitializerCount--;
1304
1305 symbolPosInBuffer.push_back(InitializerCount);
1306 unsigned int nSym = 0;
1307 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1308 for (unsigned int pos = 0; pos < InitializerCount;) {
1309 if (pos)
1310 os << ", ";
1311 if (pos != nextSymbolPos) {
1312 os << (unsigned int)buffer[pos];
1313 ++pos;
1314 continue;
1315 }
1316 // Generate a per-byte mask() operator for the symbol, which looks like:
1317 // .global .u8 addr[] = {0xFF(foo), 0xFF00(foo), 0xFF0000(foo), ...};
1318 // See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#initializers
1319 std::string symText;
1320 llvm::raw_string_ostream oss(symText);
1321 printSymbol(nSym, oss);
1322 for (unsigned i = 0; i < ptrSize; ++i) {
1323 if (i)
1324 os << ", ";
1325 llvm::write_hex(os, 0xFFULL << i * 8, HexPrintStyle::PrefixUpper);
1326 os << "(" << symText << ")";
1327 }
1328 pos += ptrSize;
1329 nextSymbolPos = symbolPosInBuffer[++nSym];
1330 assert(nextSymbolPos >= pos);
1331 }
1332}
1333
1334void NVPTXAsmPrinter::AggBuffer::printWords(raw_ostream &os) {
1335 unsigned int ptrSize = AP.MAI->getCodePointerSize();
1336 symbolPosInBuffer.push_back(size);
1337 unsigned int nSym = 0;
1338 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1339 assert(nextSymbolPos % ptrSize == 0);
1340 for (unsigned int pos = 0; pos < size; pos += ptrSize) {
1341 if (pos)
1342 os << ", ";
1343 if (pos == nextSymbolPos) {
1344 printSymbol(nSym, os);
1345 nextSymbolPos = symbolPosInBuffer[++nSym];
1346 assert(nextSymbolPos % ptrSize == 0);
1347 assert(nextSymbolPos >= pos + ptrSize);
1348 } else if (ptrSize == 4)
1349 os << support::endian::read32le(&buffer[pos]);
1350 else
1351 os << support::endian::read64le(&buffer[pos]);
1352 }
1353}
1354
1355void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) {
1356 if (localDecls.find(f) == localDecls.end())
1357 return;
1358
1359 std::vector<const GlobalVariable *> &gvars = localDecls[f];
1360
1361 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
1362 const NVPTXSubtarget &STI =
1363 *static_cast<const NVPTXSubtarget *>(NTM.getSubtargetImpl());
1364
1365 for (const GlobalVariable *GV : gvars) {
1366 O << "\t// demoted variable\n\t";
1367 printModuleLevelGV(GV, O, /*processDemoted=*/true, STI);
1368 }
1369}
1370
1371void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
1372 raw_ostream &O) const {
1373 switch (AddressSpace) {
1375 O << "local";
1376 break;
1378 O << "global";
1379 break;
1381 O << "const";
1382 break;
1384 O << "shared";
1385 break;
1386 default:
1387 report_fatal_error("Bad address space found while emitting PTX: " +
1389 break;
1390 }
1391}
1392
1393std::string
1394NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {
1395 switch (Ty->getTypeID()) {
1396 case Type::IntegerTyID: {
1397 unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
1398 if (NumBits == 1)
1399 return "pred";
1400 else if (NumBits <= 64) {
1401 std::string name = "u";
1402 return name + utostr(NumBits);
1403 } else {
1404 llvm_unreachable("Integer too large");
1405 break;
1406 }
1407 break;
1408 }
1409 case Type::BFloatTyID:
1410 case Type::HalfTyID:
1411 // fp16 and bf16 are stored as .b16 for compatibility with pre-sm_53
1412 // PTX assembly.
1413 return "b16";
1414 case Type::FloatTyID:
1415 return "f32";
1416 case Type::DoubleTyID:
1417 return "f64";
1418 case Type::PointerTyID: {
1419 unsigned PtrSize = TM.getPointerSizeInBits(Ty->getPointerAddressSpace());
1420 assert((PtrSize == 64 || PtrSize == 32) && "Unexpected pointer size");
1421
1422 if (PtrSize == 64)
1423 if (useB4PTR)
1424 return "b64";
1425 else
1426 return "u64";
1427 else if (useB4PTR)
1428 return "b32";
1429 else
1430 return "u32";
1431 }
1432 default:
1433 break;
1434 }
1435 llvm_unreachable("unexpected type");
1436}
1437
1438void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,
1439 raw_ostream &O,
1440 const NVPTXSubtarget &STI) {
1441 const DataLayout &DL = getDataLayout();
1442
1443 // GlobalVariables are always constant pointers themselves.
1444 Type *ETy = GVar->getValueType();
1445
1446 O << ".";
1447 emitPTXAddressSpace(GVar->getType()->getAddressSpace(), O);
1448 if (isManaged(*GVar)) {
1449 if (STI.getPTXVersion() < 40 || STI.getSmVersion() < 30) {
1451 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1452 }
1453 O << " .attribute(.managed)";
1454 }
1455 if (MaybeAlign A = GVar->getAlign())
1456 O << " .align " << A->value();
1457 else
1458 O << " .align " << (int)DL.getPrefTypeAlign(ETy).value();
1459
1460 // Special case for i128
1461 if (ETy->isIntegerTy(128)) {
1462 O << " .b8 ";
1463 getSymbol(GVar)->print(O, MAI);
1464 O << "[16]";
1465 return;
1466 }
1467
1468 if (ETy->isFloatingPointTy() || ETy->isIntOrPtrTy()) {
1469 O << " .";
1470 O << getPTXFundamentalTypeStr(ETy);
1471 O << " ";
1472 getSymbol(GVar)->print(O, MAI);
1473 return;
1474 }
1475
1476 int64_t ElementSize = 0;
1477
1478 // Although PTX has direct support for struct type and array type and LLVM IR
1479 // is very similar to PTX, the LLVM CodeGen does not support for targets that
1480 // support these high level field accesses. Structs and arrays are lowered
1481 // into arrays of bytes.
1482 switch (ETy->getTypeID()) {
1483 case Type::StructTyID:
1484 case Type::ArrayTyID:
1486 ElementSize = DL.getTypeStoreSize(ETy);
1487 O << " .b8 ";
1488 getSymbol(GVar)->print(O, MAI);
1489 O << "[";
1490 if (ElementSize) {
1491 O << ElementSize;
1492 }
1493 O << "]";
1494 break;
1495 default:
1496 llvm_unreachable("type not supported yet");
1497 }
1498}
1499
1500void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
1501 const DataLayout &DL = getDataLayout();
1502 const AttributeList &PAL = F->getAttributes();
1503 const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
1504 const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
1505
1507 unsigned paramIndex = 0;
1508 bool first = true;
1509 bool isKernelFunc = isKernelFunction(*F);
1510 bool isABI = (STI.getSmVersion() >= 20);
1511 bool hasImageHandles = STI.hasImageHandles();
1512
1513 if (F->arg_empty() && !F->isVarArg()) {
1514 O << "()";
1515 return;
1516 }
1517
1518 O << "(\n";
1519
1520 for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) {
1521 Type *Ty = I->getType();
1522
1523 if (!first)
1524 O << ",\n";
1525
1526 first = false;
1527
1528 // Handle image/sampler parameters
1529 if (isKernelFunction(*F)) {
1530 if (isSampler(*I) || isImage(*I)) {
1531 if (isImage(*I)) {
1532 if (isImageWriteOnly(*I) || isImageReadWrite(*I)) {
1533 if (hasImageHandles)
1534 O << "\t.param .u64 .ptr .surfref ";
1535 else
1536 O << "\t.param .surfref ";
1537 O << TLI->getParamName(F, paramIndex);
1538 }
1539 else { // Default image is read_only
1540 if (hasImageHandles)
1541 O << "\t.param .u64 .ptr .texref ";
1542 else
1543 O << "\t.param .texref ";
1544 O << TLI->getParamName(F, paramIndex);
1545 }
1546 } else {
1547 if (hasImageHandles)
1548 O << "\t.param .u64 .ptr .samplerref ";
1549 else
1550 O << "\t.param .samplerref ";
1551 O << TLI->getParamName(F, paramIndex);
1552 }
1553 continue;
1554 }
1555 }
1556
1557 auto getOptimalAlignForParam = [TLI, &DL, &PAL, F,
1558 paramIndex](Type *Ty) -> Align {
1559 Align TypeAlign = TLI->getFunctionParamOptimizedAlign(F, Ty, DL);
1560 MaybeAlign ParamAlign = PAL.getParamAlignment(paramIndex);
1561 return std::max(TypeAlign, ParamAlign.valueOrOne());
1562 };
1563
1564 if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
1565 if (ShouldPassAsArray(Ty)) {
1566 // Just print .param .align <a> .b8 .param[size];
1567 // <a> = optimal alignment for the element type; always multiple of
1568 // PAL.getParamAlignment
1569 // size = typeallocsize of element type
1570 Align OptimalAlign = getOptimalAlignForParam(Ty);
1571
1572 O << "\t.param .align " << OptimalAlign.value() << " .b8 ";
1573 O << TLI->getParamName(F, paramIndex);
1574 O << "[" << DL.getTypeAllocSize(Ty) << "]";
1575
1576 continue;
1577 }
1578 // Just a scalar
1579 auto *PTy = dyn_cast<PointerType>(Ty);
1580 unsigned PTySizeInBits = 0;
1581 if (PTy) {
1582 PTySizeInBits =
1583 TLI->getPointerTy(DL, PTy->getAddressSpace()).getSizeInBits();
1584 assert(PTySizeInBits && "Invalid pointer size");
1585 }
1586
1587 if (isKernelFunc) {
1588 if (PTy) {
1589 // Special handling for pointer arguments to kernel
1590 O << "\t.param .u" << PTySizeInBits << " ";
1591
1592 if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() !=
1593 NVPTX::CUDA) {
1594 int addrSpace = PTy->getAddressSpace();
1595 switch (addrSpace) {
1596 default:
1597 O << ".ptr ";
1598 break;
1600 O << ".ptr .const ";
1601 break;
1603 O << ".ptr .shared ";
1604 break;
1606 O << ".ptr .global ";
1607 break;
1608 }
1609 Align ParamAlign = I->getParamAlign().valueOrOne();
1610 O << ".align " << ParamAlign.value() << " ";
1611 }
1612 O << TLI->getParamName(F, paramIndex);
1613 continue;
1614 }
1615
1616 // non-pointer scalar to kernel func
1617 O << "\t.param .";
1618 // Special case: predicate operands become .u8 types
1619 if (Ty->isIntegerTy(1))
1620 O << "u8";
1621 else
1622 O << getPTXFundamentalTypeStr(Ty);
1623 O << " ";
1624 O << TLI->getParamName(F, paramIndex);
1625 continue;
1626 }
1627 // Non-kernel function, just print .param .b<size> for ABI
1628 // and .reg .b<size> for non-ABI
1629 unsigned sz = 0;
1630 if (isa<IntegerType>(Ty)) {
1631 sz = cast<IntegerType>(Ty)->getBitWidth();
1633 } else if (PTy) {
1634 assert(PTySizeInBits && "Invalid pointer size");
1635 sz = PTySizeInBits;
1636 } else
1637 sz = Ty->getPrimitiveSizeInBits();
1638 if (isABI)
1639 O << "\t.param .b" << sz << " ";
1640 else
1641 O << "\t.reg .b" << sz << " ";
1642 O << TLI->getParamName(F, paramIndex);
1643 continue;
1644 }
1645
1646 // param has byVal attribute.
1647 Type *ETy = PAL.getParamByValType(paramIndex);
1648 assert(ETy && "Param should have byval type");
1649
1650 if (isABI || isKernelFunc) {
1651 // Just print .param .align <a> .b8 .param[size];
1652 // <a> = optimal alignment for the element type; always multiple of
1653 // PAL.getParamAlignment
1654 // size = typeallocsize of element type
1655 Align OptimalAlign =
1656 isKernelFunc
1657 ? getOptimalAlignForParam(ETy)
1658 : TLI->getFunctionByValParamAlign(
1659 F, ETy, PAL.getParamAlignment(paramIndex).valueOrOne(), DL);
1660
1661 unsigned sz = DL.getTypeAllocSize(ETy);
1662 O << "\t.param .align " << OptimalAlign.value() << " .b8 ";
1663 O << TLI->getParamName(F, paramIndex);
1664 O << "[" << sz << "]";
1665 continue;
1666 } else {
1667 // Split the ETy into constituent parts and
1668 // print .param .b<size> <name> for each part.
1669 // Further, if a part is vector, print the above for
1670 // each vector element.
1671 SmallVector<EVT, 16> vtparts;
1672 ComputeValueVTs(*TLI, DL, ETy, vtparts);
1673 for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
1674 unsigned elems = 1;
1675 EVT elemtype = vtparts[i];
1676 if (vtparts[i].isVector()) {
1677 elems = vtparts[i].getVectorNumElements();
1678 elemtype = vtparts[i].getVectorElementType();
1679 }
1680
1681 for (unsigned j = 0, je = elems; j != je; ++j) {
1682 unsigned sz = elemtype.getSizeInBits();
1683 if (elemtype.isInteger())
1685 O << "\t.reg .b" << sz << " ";
1686 O << TLI->getParamName(F, paramIndex);
1687 if (j < je - 1)
1688 O << ",\n";
1689 ++paramIndex;
1690 }
1691 if (i < e - 1)
1692 O << ",\n";
1693 }
1694 --paramIndex;
1695 continue;
1696 }
1697 }
1698
1699 if (F->isVarArg()) {
1700 if (!first)
1701 O << ",\n";
1702 O << "\t.param .align " << STI.getMaxRequiredAlignment();
1703 O << " .b8 ";
1704 O << TLI->getParamName(F, /* vararg */ -1) << "[]";
1705 }
1706
1707 O << "\n)";
1708}
1709
1710void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
1711 const MachineFunction &MF) {
1712 SmallString<128> Str;
1714
1715 // Map the global virtual register number to a register class specific
1716 // virtual register number starting from 1 with that class.
1718 //unsigned numRegClasses = TRI->getNumRegClasses();
1719
1720 // Emit the Fake Stack Object
1721 const MachineFrameInfo &MFI = MF.getFrameInfo();
1722 int64_t NumBytes = MFI.getStackSize();
1723 if (NumBytes) {
1724 O << "\t.local .align " << MFI.getMaxAlign().value() << " .b8 \t"
1725 << DEPOTNAME << getFunctionNumber() << "[" << NumBytes << "];\n";
1726 if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
1727 O << "\t.reg .b64 \t%SP;\n";
1728 O << "\t.reg .b64 \t%SPL;\n";
1729 } else {
1730 O << "\t.reg .b32 \t%SP;\n";
1731 O << "\t.reg .b32 \t%SPL;\n";
1732 }
1733 }
1734
1735 // Go through all virtual registers to establish the mapping between the
1736 // global virtual
1737 // register number and the per class virtual register number.
1738 // We use the per class virtual register number in the ptx output.
1739 unsigned int numVRs = MRI->getNumVirtRegs();
1740 for (unsigned i = 0; i < numVRs; i++) {
1742 const TargetRegisterClass *RC = MRI->getRegClass(vr);
1743 DenseMap<unsigned, unsigned> &regmap = VRegMapping[RC];
1744 int n = regmap.size();
1745 regmap.insert(std::make_pair(vr, n + 1));
1746 }
1747
1748 // Emit register declarations
1749 // @TODO: Extract out the real register usage
1750 // O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n";
1751 // O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n";
1752 // O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n";
1753 // O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n";
1754 // O << "\t.reg .s64 %rd<" << NVPTXNumRegisters << ">;\n";
1755 // O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n";
1756 // O << "\t.reg .f64 %fd<" << NVPTXNumRegisters << ">;\n";
1757
1758 // Emit declaration of the virtual registers or 'physical' registers for
1759 // each register class
1760 for (unsigned i=0; i< TRI->getNumRegClasses(); i++) {
1761 const TargetRegisterClass *RC = TRI->getRegClass(i);
1762 DenseMap<unsigned, unsigned> &regmap = VRegMapping[RC];
1763 std::string rcname = getNVPTXRegClassName(RC);
1764 std::string rcStr = getNVPTXRegClassStr(RC);
1765 int n = regmap.size();
1766
1767 // Only declare those registers that may be used.
1768 if (n) {
1769 O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1)
1770 << ">;\n";
1771 }
1772 }
1773
1774 OutStreamer->emitRawText(O.str());
1775}
1776
1777void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) {
1778 APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
1779 bool ignored;
1780 unsigned int numHex;
1781 const char *lead;
1782
1783 if (Fp->getType()->getTypeID() == Type::FloatTyID) {
1784 numHex = 8;
1785 lead = "0f";
1787 } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
1788 numHex = 16;
1789 lead = "0d";
1791 } else
1792 llvm_unreachable("unsupported fp type");
1793
1794 APInt API = APF.bitcastToAPInt();
1795 O << lead << format_hex_no_prefix(API.getZExtValue(), numHex, /*Upper=*/true);
1796}
1797
1798void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
1799 if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1800 O << CI->getValue();
1801 return;
1802 }
1803 if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
1804 printFPConstant(CFP, O);
1805 return;
1806 }
1807 if (isa<ConstantPointerNull>(CPV)) {
1808 O << "0";
1809 return;
1810 }
1811 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1812 bool IsNonGenericPointer = false;
1813 if (GVar->getType()->getAddressSpace() != 0) {
1814 IsNonGenericPointer = true;
1815 }
1816 if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
1817 O << "generic(";
1818 getSymbol(GVar)->print(O, MAI);
1819 O << ")";
1820 } else {
1821 getSymbol(GVar)->print(O, MAI);
1822 }
1823 return;
1824 }
1825 if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1826 const MCExpr *E = lowerConstantForGV(cast<Constant>(Cexpr), false);
1827 printMCExpr(*E, O);
1828 return;
1829 }
1830 llvm_unreachable("Not scalar type found in printScalarConstant()");
1831}
1832
1833void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes,
1834 AggBuffer *AggBuffer) {
1835 const DataLayout &DL = getDataLayout();
1836 int AllocSize = DL.getTypeAllocSize(CPV->getType());
1837 if (isa<UndefValue>(CPV) || CPV->isNullValue()) {
1838 // Non-zero Bytes indicates that we need to zero-fill everything. Otherwise,
1839 // only the space allocated by CPV.
1840 AggBuffer->addZeros(Bytes ? Bytes : AllocSize);
1841 return;
1842 }
1843
1844 // Helper for filling AggBuffer with APInts.
1845 auto AddIntToBuffer = [AggBuffer, Bytes](const APInt &Val) {
1846 size_t NumBytes = (Val.getBitWidth() + 7) / 8;
1847 SmallVector<unsigned char, 16> Buf(NumBytes);
1848 for (unsigned I = 0; I < NumBytes; ++I) {
1849 Buf[I] = Val.extractBitsAsZExtValue(8, I * 8);
1850 }
1851 AggBuffer->addBytes(Buf.data(), NumBytes, Bytes);
1852 };
1853
1854 switch (CPV->getType()->getTypeID()) {
1855 case Type::IntegerTyID:
1856 if (const auto CI = dyn_cast<ConstantInt>(CPV)) {
1857 AddIntToBuffer(CI->getValue());
1858 break;
1859 }
1860 if (const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1861 if (const auto *CI =
1862 dyn_cast<ConstantInt>(ConstantFoldConstant(Cexpr, DL))) {
1863 AddIntToBuffer(CI->getValue());
1864 break;
1865 }
1866 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1867 Value *V = Cexpr->getOperand(0)->stripPointerCasts();
1868 AggBuffer->addSymbol(V, Cexpr->getOperand(0));
1869 AggBuffer->addZeros(AllocSize);
1870 break;
1871 }
1872 }
1873 llvm_unreachable("unsupported integer const type");
1874 break;
1875
1876 case Type::HalfTyID:
1877 case Type::BFloatTyID:
1878 case Type::FloatTyID:
1879 case Type::DoubleTyID:
1880 AddIntToBuffer(cast<ConstantFP>(CPV)->getValueAPF().bitcastToAPInt());
1881 break;
1882
1883 case Type::PointerTyID: {
1884 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1885 AggBuffer->addSymbol(GVar, GVar);
1886 } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1887 const Value *v = Cexpr->stripPointerCasts();
1888 AggBuffer->addSymbol(v, Cexpr);
1889 }
1890 AggBuffer->addZeros(AllocSize);
1891 break;
1892 }
1893
1894 case Type::ArrayTyID:
1896 case Type::StructTyID: {
1897 if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
1898 bufferAggregateConstant(CPV, AggBuffer);
1899 if (Bytes > AllocSize)
1900 AggBuffer->addZeros(Bytes - AllocSize);
1901 } else if (isa<ConstantAggregateZero>(CPV))
1902 AggBuffer->addZeros(Bytes);
1903 else
1904 llvm_unreachable("Unexpected Constant type");
1905 break;
1906 }
1907
1908 default:
1909 llvm_unreachable("unsupported type");
1910 }
1911}
1912
1913void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV,
1914 AggBuffer *aggBuffer) {
1915 const DataLayout &DL = getDataLayout();
1916 int Bytes;
1917
1918 // Integers of arbitrary width
1919 if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1920 APInt Val = CI->getValue();
1921 for (unsigned I = 0, E = DL.getTypeAllocSize(CPV->getType()); I < E; ++I) {
1922 uint8_t Byte = Val.getLoBits(8).getZExtValue();
1923 aggBuffer->addBytes(&Byte, 1, 1);
1924 Val.lshrInPlace(8);
1925 }
1926 return;
1927 }
1928
1929 // Old constants
1930 if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
1931 if (CPV->getNumOperands())
1932 for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i)
1933 bufferLEByte(cast<Constant>(CPV->getOperand(i)), 0, aggBuffer);
1934 return;
1935 }
1936
1937 if (const ConstantDataSequential *CDS =
1938 dyn_cast<ConstantDataSequential>(CPV)) {
1939 if (CDS->getNumElements())
1940 for (unsigned i = 0; i < CDS->getNumElements(); ++i)
1941 bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
1942 aggBuffer);
1943 return;
1944 }
1945
1946 if (isa<ConstantStruct>(CPV)) {
1947 if (CPV->getNumOperands()) {
1948 StructType *ST = cast<StructType>(CPV->getType());
1949 for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) {
1950 if (i == (e - 1))
1951 Bytes = DL.getStructLayout(ST)->getElementOffset(0) +
1952 DL.getTypeAllocSize(ST) -
1953 DL.getStructLayout(ST)->getElementOffset(i);
1954 else
1955 Bytes = DL.getStructLayout(ST)->getElementOffset(i + 1) -
1956 DL.getStructLayout(ST)->getElementOffset(i);
1957 bufferLEByte(cast<Constant>(CPV->getOperand(i)), Bytes, aggBuffer);
1958 }
1959 }
1960 return;
1961 }
1962 llvm_unreachable("unsupported constant type in printAggregateConstant()");
1963}
1964
1965/// lowerConstantForGV - Return an MCExpr for the given Constant. This is mostly
1966/// a copy from AsmPrinter::lowerConstant, except customized to only handle
1967/// expressions that are representable in PTX and create
1968/// NVPTXGenericMCSymbolRefExpr nodes for addrspacecast instructions.
1969const MCExpr *
1970NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV, bool ProcessingGeneric) {
1971 MCContext &Ctx = OutContext;
1972
1973 if (CV->isNullValue() || isa<UndefValue>(CV))
1974 return MCConstantExpr::create(0, Ctx);
1975
1976 if (const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
1977 return MCConstantExpr::create(CI->getZExtValue(), Ctx);
1978
1979 if (const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
1980 const MCSymbolRefExpr *Expr =
1982 if (ProcessingGeneric) {
1983 return NVPTXGenericMCSymbolRefExpr::create(Expr, Ctx);
1984 } else {
1985 return Expr;
1986 }
1987 }
1988
1989 const ConstantExpr *CE = dyn_cast<ConstantExpr>(CV);
1990 if (!CE) {
1991 llvm_unreachable("Unknown constant value to lower!");
1992 }
1993
1994 switch (CE->getOpcode()) {
1995 default:
1996 break; // Error
1997
1998 case Instruction::AddrSpaceCast: {
1999 // Strip the addrspacecast and pass along the operand
2000 PointerType *DstTy = cast<PointerType>(CE->getType());
2001 if (DstTy->getAddressSpace() == 0)
2002 return lowerConstantForGV(cast<const Constant>(CE->getOperand(0)), true);
2003
2004 break; // Error
2005 }
2006
2007 case Instruction::GetElementPtr: {
2008 const DataLayout &DL = getDataLayout();
2009
2010 // Generate a symbolic expression for the byte address
2011 APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0);
2012 cast<GEPOperator>(CE)->accumulateConstantOffset(DL, OffsetAI);
2013
2014 const MCExpr *Base = lowerConstantForGV(CE->getOperand(0),
2015 ProcessingGeneric);
2016 if (!OffsetAI)
2017 return Base;
2018
2019 int64_t Offset = OffsetAI.getSExtValue();
2021 Ctx);
2022 }
2023
2024 case Instruction::Trunc:
2025 // We emit the value and depend on the assembler to truncate the generated
2026 // expression properly. This is important for differences between
2027 // blockaddress labels. Since the two labels are in the same function, it
2028 // is reasonable to treat their delta as a 32-bit value.
2029 [[fallthrough]];
2030 case Instruction::BitCast:
2031 return lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
2032
2033 case Instruction::IntToPtr: {
2034 const DataLayout &DL = getDataLayout();
2035
2036 // Handle casts to pointers by changing them into casts to the appropriate
2037 // integer type. This promotes constant folding and simplifies this code.
2038 Constant *Op = CE->getOperand(0);
2039 Op = ConstantFoldIntegerCast(Op, DL.getIntPtrType(CV->getType()),
2040 /*IsSigned*/ false, DL);
2041 if (Op)
2042 return lowerConstantForGV(Op, ProcessingGeneric);
2043
2044 break; // Error
2045 }
2046
2047 case Instruction::PtrToInt: {
2048 const DataLayout &DL = getDataLayout();
2049
2050 // Support only foldable casts to/from pointers that can be eliminated by
2051 // changing the pointer to the appropriately sized integer type.
2052 Constant *Op = CE->getOperand(0);
2053 Type *Ty = CE->getType();
2054
2055 const MCExpr *OpExpr = lowerConstantForGV(Op, ProcessingGeneric);
2056
2057 // We can emit the pointer value into this slot if the slot is an
2058 // integer slot equal to the size of the pointer.
2059 if (DL.getTypeAllocSize(Ty) == DL.getTypeAllocSize(Op->getType()))
2060 return OpExpr;
2061
2062 // Otherwise the pointer is smaller than the resultant integer, mask off
2063 // the high bits so we are sure to get a proper truncation if the input is
2064 // a constant expr.
2065 unsigned InBits = DL.getTypeAllocSizeInBits(Op->getType());
2066 const MCExpr *MaskExpr = MCConstantExpr::create(~0ULL >> (64-InBits), Ctx);
2067 return MCBinaryExpr::createAnd(OpExpr, MaskExpr, Ctx);
2068 }
2069
2070 // The MC library also has a right-shift operator, but it isn't consistently
2071 // signed or unsigned between different targets.
2072 case Instruction::Add: {
2073 const MCExpr *LHS = lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
2074 const MCExpr *RHS = lowerConstantForGV(CE->getOperand(1), ProcessingGeneric);
2075 switch (CE->getOpcode()) {
2076 default: llvm_unreachable("Unknown binary operator constant cast expr");
2077 case Instruction::Add: return MCBinaryExpr::createAdd(LHS, RHS, Ctx);
2078 }
2079 }
2080 }
2081
2082 // If the code isn't optimized, there may be outstanding folding
2083 // opportunities. Attempt to fold the expression using DataLayout as a
2084 // last resort before giving up.
2086 if (C != CE)
2087 return lowerConstantForGV(C, ProcessingGeneric);
2088
2089 // Otherwise report the problem to the user.
2090 std::string S;
2092 OS << "Unsupported expression in static initializer: ";
2093 CE->printAsOperand(OS, /*PrintType=*/false,
2094 !MF ? nullptr : MF->getFunction().getParent());
2095 report_fatal_error(Twine(OS.str()));
2096}
2097
2098// Copy of MCExpr::print customized for NVPTX
2099void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) {
2100 switch (Expr.getKind()) {
2101 case MCExpr::Target:
2102 return cast<MCTargetExpr>(&Expr)->printImpl(OS, MAI);
2103 case MCExpr::Constant:
2104 OS << cast<MCConstantExpr>(Expr).getValue();
2105 return;
2106
2107 case MCExpr::SymbolRef: {
2108 const MCSymbolRefExpr &SRE = cast<MCSymbolRefExpr>(Expr);
2109 const MCSymbol &Sym = SRE.getSymbol();
2110 Sym.print(OS, MAI);
2111 return;
2112 }
2113
2114 case MCExpr::Unary: {
2115 const MCUnaryExpr &UE = cast<MCUnaryExpr>(Expr);
2116 switch (UE.getOpcode()) {
2117 case MCUnaryExpr::LNot: OS << '!'; break;
2118 case MCUnaryExpr::Minus: OS << '-'; break;
2119 case MCUnaryExpr::Not: OS << '~'; break;
2120 case MCUnaryExpr::Plus: OS << '+'; break;
2121 }
2122 printMCExpr(*UE.getSubExpr(), OS);
2123 return;
2124 }
2125
2126 case MCExpr::Binary: {
2127 const MCBinaryExpr &BE = cast<MCBinaryExpr>(Expr);
2128
2129 // Only print parens around the LHS if it is non-trivial.
2130 if (isa<MCConstantExpr>(BE.getLHS()) || isa<MCSymbolRefExpr>(BE.getLHS()) ||
2131 isa<NVPTXGenericMCSymbolRefExpr>(BE.getLHS())) {
2132 printMCExpr(*BE.getLHS(), OS);
2133 } else {
2134 OS << '(';
2135 printMCExpr(*BE.getLHS(), OS);
2136 OS<< ')';
2137 }
2138
2139 switch (BE.getOpcode()) {
2140 case MCBinaryExpr::Add:
2141 // Print "X-42" instead of "X+-42".
2142 if (const MCConstantExpr *RHSC = dyn_cast<MCConstantExpr>(BE.getRHS())) {
2143 if (RHSC->getValue() < 0) {
2144 OS << RHSC->getValue();
2145 return;
2146 }
2147 }
2148
2149 OS << '+';
2150 break;
2151 default: llvm_unreachable("Unhandled binary operator");
2152 }
2153
2154 // Only print parens around the LHS if it is non-trivial.
2155 if (isa<MCConstantExpr>(BE.getRHS()) || isa<MCSymbolRefExpr>(BE.getRHS())) {
2156 printMCExpr(*BE.getRHS(), OS);
2157 } else {
2158 OS << '(';
2159 printMCExpr(*BE.getRHS(), OS);
2160 OS << ')';
2161 }
2162 return;
2163 }
2164 }
2165
2166 llvm_unreachable("Invalid expression kind!");
2167}
2168
2169/// PrintAsmOperand - Print out an operand for an inline asm expression.
2170///
2171bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
2172 const char *ExtraCode, raw_ostream &O) {
2173 if (ExtraCode && ExtraCode[0]) {
2174 if (ExtraCode[1] != 0)
2175 return true; // Unknown modifier.
2176
2177 switch (ExtraCode[0]) {
2178 default:
2179 // See if this is a generic print operand
2180 return AsmPrinter::PrintAsmOperand(MI, OpNo, ExtraCode, O);
2181 case 'r':
2182 break;
2183 }
2184 }
2185
2186 printOperand(MI, OpNo, O);
2187
2188 return false;
2189}
2190
2191bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI,
2192 unsigned OpNo,
2193 const char *ExtraCode,
2194 raw_ostream &O) {
2195 if (ExtraCode && ExtraCode[0])
2196 return true; // Unknown modifier
2197
2198 O << '[';
2199 printMemOperand(MI, OpNo, O);
2200 O << ']';
2201
2202 return false;
2203}
2204
2205void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, unsigned OpNum,
2206 raw_ostream &O) {
2207 const MachineOperand &MO = MI->getOperand(OpNum);
2208 switch (MO.getType()) {
2210 if (MO.getReg().isPhysical()) {
2211 if (MO.getReg() == NVPTX::VRDepot)
2213 else
2215 } else {
2216 emitVirtualRegister(MO.getReg(), O);
2217 }
2218 break;
2219
2221 O << MO.getImm();
2222 break;
2223
2225 printFPConstant(MO.getFPImm(), O);
2226 break;
2227
2229 PrintSymbolOperand(MO, O);
2230 break;
2231
2233 MO.getMBB()->getSymbol()->print(O, MAI);
2234 break;
2235
2236 default:
2237 llvm_unreachable("Operand type not supported.");
2238 }
2239}
2240
2241void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, unsigned OpNum,
2242 raw_ostream &O, const char *Modifier) {
2243 printOperand(MI, OpNum, O);
2244
2245 if (Modifier && strcmp(Modifier, "add") == 0) {
2246 O << ", ";
2247 printOperand(MI, OpNum + 1, O);
2248 } else {
2249 if (MI->getOperand(OpNum + 1).isImm() &&
2250 MI->getOperand(OpNum + 1).getImm() == 0)
2251 return; // don't print ',0' or '+0'
2252 O << "+";
2253 printOperand(MI, OpNum + 1, O);
2254 }
2255}
2256
2257// Force static initialization.
2261}
MachineBasicBlock & MBB
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static cl::opt< bool > LowerCtorDtor("amdgpu-lower-global-ctor-dtor", cl::desc("Lower GPU ctor / dtors to globals on the device."), cl::init(true), cl::Hidden)
This file declares a class to represent arbitrary precision floating point values and provide a varie...
This file implements a class to represent arbitrary precision integral constant values and operations...
This file contains the simple types necessary to represent the attributes associated with functions a...
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
#define LLVM_EXTERNAL_VISIBILITY
Definition: Compiler.h:135
This file contains the declarations for the subclasses of Constant, which represent the different fla...
Looks at all the uses of the given value Returns the Liveness deduced from the uses of this value Adds all uses that cause the result to be MaybeLive to MaybeLiveRetUses If the result is MaybeLiveUses might be modified but its content should be ignored(since it might not be complete). DeadArgumentEliminationPass
This file defines the DenseMap class.
This file defines the DenseSet and SmallDenseSet classes.
std::string Name
Symbol * Sym
Definition: ELF_riscv.cpp:479
static GCMetadataPrinterRegistry::Add< ErlangGCPrinter > X("erlang", "erlang-compatible garbage collector")
IRTranslator LLVM IR MI
#define F(x, y, z)
Definition: MD5.cpp:55
#define I(x, y, z)
Definition: MD5.cpp:58
unsigned const TargetRegisterInfo * TRI
Module.h This file contains the declarations for the Module class.
static bool isEmptyXXStructor(GlobalVariable *GV)
#define DEPOTNAME
static bool usedInOneFunc(const User *U, Function const *&oneFunc)
static void VisitGlobalVariableForEmission(const GlobalVariable *GV, SmallVectorImpl< const GlobalVariable * > &Order, DenseSet< const GlobalVariable * > &Visited, DenseSet< const GlobalVariable * > &Visiting)
VisitGlobalVariableForEmission - Add GV to the list of GlobalVariable instances to be emitted,...
LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXAsmPrinter()
static bool usedInGlobalVarDef(const Constant *C)
static bool useFuncSeen(const Constant *C, DenseMap< const Function *, bool > &seenMap)
static cl::opt< bool > LowerCtorDtor("nvptx-lower-global-ctor-dtor", cl::desc("Lower GPU ctor / dtors to globals on the device."), cl::init(false), cl::Hidden)
static bool ShouldPassAsArray(Type *Ty)
static void DiscoverDependentGlobals(const Value *V, DenseSet< const GlobalVariable * > &Globals)
DiscoverDependentGlobals - Return a set of GlobalVariables on which V depends.
static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f)
static GCMetadataPrinterRegistry::Add< OcamlGCMetadataPrinter > Y("ocaml", "ocaml 3.10-compatible collector")
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
static const char * name
Definition: SMEABIPass.cpp:49
raw_pwrite_stream & OS
This file defines the SmallString class.
This file defines the SmallVector class.
This file contains some functions that are useful when dealing with strings.
Value * RHS
Value * LHS
@ __CLK_ADDRESS_BASE
@ __CLK_FILTER_BASE
@ __CLK_NORMALIZED_BASE
@ __CLK_NORMALIZED_MASK
@ __CLK_ADDRESS_MASK
@ __CLK_FILTER_MASK
opStatus convert(const fltSemantics &ToSemantics, roundingMode RM, bool *losesInfo)
Definition: APFloat.cpp:5196
APInt bitcastToAPInt() const
Definition: APFloat.h:1210
Class for arbitrary precision integers.
Definition: APInt.h:76
APInt getLoBits(unsigned numBits) const
Compute an APInt containing numBits lowbits from this APInt.
Definition: APInt.cpp:613
uint64_t getZExtValue() const
Get zero extended value.
Definition: APInt.h:1491
void lshrInPlace(unsigned ShiftAmt)
Logical right-shift this APInt by ShiftAmt in place.
Definition: APInt.h:836
This class represents an incoming formal argument to a Function.
Definition: Argument.h:31
MCSymbol * getSymbol(const GlobalValue *GV) const
Definition: AsmPrinter.cpp:700
void EmitToStreamer(MCStreamer &S, const MCInst &Inst)
Definition: AsmPrinter.cpp:418
TargetMachine & TM
Target machine description.
Definition: AsmPrinter.h:87
virtual void PrintSymbolOperand(const MachineOperand &MO, raw_ostream &OS)
Print the MachineOperand as a symbol.
const MCAsmInfo * MAI
Target Asm Printer information.
Definition: AsmPrinter.h:90
MachineFunction * MF
The current machine function.
Definition: AsmPrinter.h:102
bool doInitialization(Module &M) override
Set up the AsmPrinter when we are working on a new module.
Definition: AsmPrinter.cpp:449
unsigned getFunctionNumber() const
Return a unique ID for the current function.
Definition: AsmPrinter.cpp:394
MCSymbol * CurrentFnSym
The symbol for the current function.
Definition: AsmPrinter.h:121
MachineModuleInfo * MMI
This is a pointer to the current MachineModuleInfo.
Definition: AsmPrinter.h:105
MCContext & OutContext
This is the context for the output file that we are streaming.
Definition: AsmPrinter.h:94
bool doFinalization(Module &M) override
Shut down the asmprinter.
MCSymbol * GetExternalSymbolSymbol(Twine Sym) const
Return the MCSymbol for the specified ExternalSymbol.
virtual void emitBasicBlockStart(const MachineBasicBlock &MBB)
Targets can override this to emit stuff at the start of a basic block.
bool runOnMachineFunction(MachineFunction &MF) override
Emit the specified function out to the OutStreamer.
Definition: AsmPrinter.h:395
std::unique_ptr< MCStreamer > OutStreamer
This is the MCStreamer object for the file we are generating.
Definition: AsmPrinter.h:99
const DataLayout & getDataLayout() const
Return information about data layout.
Definition: AsmPrinter.cpp:402
void emitInitialRawDwarfLocDirective(const MachineFunction &MF)
Emits inital debug location directive.
Definition: AsmPrinter.cpp:422
const MCSubtargetInfo & getSubtargetInfo() const
Return information about subtarget.
Definition: AsmPrinter.cpp:413
virtual bool PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, const char *ExtraCode, raw_ostream &OS)
Print the specified operand of MI, an INLINEASM instruction, using the specified assembler variant.
LLVM Basic Block Representation.
Definition: BasicBlock.h:60
const Function * getParent() const
Return the enclosing method, or null if none.
Definition: BasicBlock.h:205
ConstantArray - Constant Array Declarations.
Definition: Constants.h:423
ConstantDataSequential - A vector or array constant whose element type is a simple 1/2/4/8-byte integ...
Definition: Constants.h:583
A constant value that is initialized with an expression using other constant values.
Definition: Constants.h:1017
ConstantFP - Floating Point Values [float, double].
Definition: Constants.h:268
const APFloat & getValueAPF() const
Definition: Constants.h:311
This is the shared class of boolean and integer constants.
Definition: Constants.h:80
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
Definition: Constants.h:154
const APInt & getValue() const
Return the constant as an APInt value reference.
Definition: Constants.h:145
This is an important base class in LLVM.
Definition: Constant.h:41
bool isNullValue() const
Return true if this is the value that would be returned by getNullValue.
Definition: Constants.cpp:90
Subprogram description.
This class represents an Operation in the Expression.
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:110
iterator find(const_arg_type_t< KeyT > Val)
Definition: DenseMap.h:155
unsigned size() const
Definition: DenseMap.h:99
iterator end()
Definition: DenseMap.h:84
bool contains(const_arg_type_t< KeyT > Val) const
Return true if the specified key is in the map, false otherwise.
Definition: DenseMap.h:145
std::pair< iterator, bool > insert(const std::pair< KeyT, ValueT > &KV)
Definition: DenseMap.h:220
Implements a dense probed hash-table based set.
Definition: DenseSet.h:271
DISubprogram * getSubprogram() const
Get the attached subprogram.
Definition: Metadata.cpp:1831
const GlobalObject * getAliaseeObject() const
Definition: Globals.cpp:556
StringRef getSection() const
Get the custom section of this global if it has one.
Definition: GlobalObject.h:118
MaybeAlign getAlign() const
Returns the alignment of the given variable or function.
Definition: GlobalObject.h:80
bool hasSection() const
Check if this global has a custom object file section.
Definition: GlobalObject.h:110
bool hasLinkOnceLinkage() const
Definition: GlobalValue.h:514
bool hasExternalLinkage() const
Definition: GlobalValue.h:510
bool isDeclaration() const
Return true if the primary definition of this global value is outside of the current translation unit...
Definition: Globals.cpp:274
bool hasLocalLinkage() const
Definition: GlobalValue.h:527
bool hasPrivateLinkage() const
Definition: GlobalValue.h:526
Module * getParent()
Get the module that this global value is contained inside of...
Definition: GlobalValue.h:655
PointerType * getType() const
Global values are always pointers.
Definition: GlobalValue.h:294
bool hasWeakLinkage() const
Definition: GlobalValue.h:521
bool hasCommonLinkage() const
Definition: GlobalValue.h:531
bool hasAvailableExternallyLinkage() const
Definition: GlobalValue.h:511
Type * getValueType() const
Definition: GlobalValue.h:296
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasInitializer() const
Definitions have initializers, declarations don't.
This class describes a target machine that is implemented with the LLVM target-independent code gener...
unsigned getCodePointerSize() const
Get the code pointer size in bytes.
Definition: MCAsmInfo.h:546
Binary assembler expressions.
Definition: MCExpr.h:492
const MCExpr * getLHS() const
Get the left-hand side expression of the binary operator.
Definition: MCExpr.h:639
const MCExpr * getRHS() const
Get the right-hand side expression of the binary operator.
Definition: MCExpr.h:642
static const MCBinaryExpr * createAnd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx)
Definition: MCExpr.h:541
static const MCBinaryExpr * createAdd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx)
Definition: MCExpr.h:536
Opcode getOpcode() const
Get the kind of this binary expression.
Definition: MCExpr.h:636
@ Add
Addition.
Definition: MCExpr.h:495
static const MCConstantExpr * create(int64_t Value, MCContext &Ctx, bool PrintInHex=false, unsigned SizeInBytes=0)
Definition: MCExpr.cpp:194
Context object for machine code objects.
Definition: MCContext.h:76
MCSymbol * getOrCreateSymbol(const Twine &Name)
Lookup the symbol inside with the specified Name.
Definition: MCContext.cpp:200
Base class for the full range of assembler expressions which are needed for parsing.
Definition: MCExpr.h:35
@ Unary
Unary expressions.
Definition: MCExpr.h:41
@ Constant
Constant expressions.
Definition: MCExpr.h:39
@ SymbolRef
References to labels and assigned expressions.
Definition: MCExpr.h:40
@ Target
Target specific expression.
Definition: MCExpr.h:42
@ Binary
Binary expressions.
Definition: MCExpr.h:38
ExprKind getKind() const
Definition: MCExpr.h:81
Instances of this class represent a single low-level machine instruction.
Definition: MCInst.h:184
void addOperand(const MCOperand Op)
Definition: MCInst.h:210
void setOpcode(unsigned Op)
Definition: MCInst.h:197
Describe properties that are true of each instruction in the target description file.
Definition: MCInstrDesc.h:198
Instances of this class represent operands of the MCInst class.
Definition: MCInst.h:36
static MCOperand createReg(unsigned Reg)
Definition: MCInst.h:134
static MCOperand createExpr(const MCExpr *Val)
Definition: MCInst.h:162
static MCOperand createImm(int64_t Val)
Definition: MCInst.h:141
Represent a reference to a symbol from inside an expression.
Definition: MCExpr.h:192
const MCSymbol & getSymbol() const
Definition: MCExpr.h:410
static const MCSymbolRefExpr * create(const MCSymbol *Symbol, MCContext &Ctx)
Definition: MCExpr.h:397
MCSymbol - Instances of this class represent a symbol name in the MC file, and MCSymbols are created ...
Definition: MCSymbol.h:40
void print(raw_ostream &OS, const MCAsmInfo *MAI) const
print - Print the value to the stream OS.
Definition: MCSymbol.cpp:58
Unary assembler expressions.
Definition: MCExpr.h:436
Opcode getOpcode() const
Get the kind of this unary expression.
Definition: MCExpr.h:479
@ Minus
Unary minus.
Definition: MCExpr.h:440
@ Plus
Unary plus.
Definition: MCExpr.h:442
@ Not
Bitwise negation.
Definition: MCExpr.h:441
@ LNot
Logical negation.
Definition: MCExpr.h:439
const MCExpr * getSubExpr() const
Get the child of this unary expression.
Definition: MCExpr.h:482
Metadata node.
Definition: Metadata.h:1067
MCSymbol * getSymbol() const
Return the MCSymbol for this basic block.
iterator_range< pred_iterator > predecessors()
The MachineFrameInfo class represents an abstract stack frame until prolog/epilog code is inserted.
uint64_t getStackSize() const
Return the number of bytes that must be allocated to hold all of the fixed size frame objects.
Align getMaxAlign() const
Return the alignment in bytes that this function must be aligned to, which is greater than the defaul...
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
MachineFrameInfo & getFrameInfo()
getFrameInfo - Return the frame info object for the current function.
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
const LLVMTargetMachine & getTarget() const
getTarget - Return the target machine this machine code is compiled with
Ty * getInfo()
getInfo - Keep track of various per-function pieces of information for backends that would like to do...
Representation of each machine instruction.
Definition: MachineInstr.h:69
bool isLoopHeader(const MachineBasicBlock *BB) const
True if the block is a loop header node.
MachineLoop * getLoopFor(const MachineBasicBlock *BB) const
Return the innermost loop that BB lives in.
bool hasDebugInfo() const
Returns true if valid debug info is present.
MachineOperand class - Representation of each machine instruction operand.
const GlobalValue * getGlobal() const
int64_t getImm() const
MachineBasicBlock * getMBB() const
bool isImm() const
isImm - Tests if this is a MO_Immediate operand.
MachineOperandType getType() const
getType - Returns the MachineOperandType for this operand.
const char * getSymbolName() const
Register getReg() const
getReg - Returns the register number.
const ConstantFP * getFPImm() const
@ MO_Immediate
Immediate operand.
@ MO_GlobalAddress
Address of a global value.
@ MO_MachineBasicBlock
MachineBasicBlock reference.
@ MO_Register
Register operand.
@ MO_ExternalSymbol
Name of external global symbol.
@ MO_FPImmediate
Floating-point immediate operand.
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
A Module instance is used to store all the information related to an LLVM module.
Definition: Module.h:65
bool doInitialization(Module &M) override
Set up the AsmPrinter when we are working on a new module.
bool runOnMachineFunction(MachineFunction &F) override
Emit the specified function out to the OutStreamer.
std::string getVirtualRegisterName(unsigned) const
bool doFinalization(Module &M) override
Shut down the asmprinter.
const MCSymbol * getFunctionFrameSymbol() const override
Return symbol for the function pseudo stack if the stack frame is not a register based.
static const NVPTXFloatMCExpr * createConstantBFPHalf(const APFloat &Flt, MCContext &Ctx)
Definition: NVPTXMCExpr.h:44
static const NVPTXFloatMCExpr * createConstantFPHalf(const APFloat &Flt, MCContext &Ctx)
Definition: NVPTXMCExpr.h:49
static const NVPTXFloatMCExpr * createConstantFPSingle(const APFloat &Flt, MCContext &Ctx)
Definition: NVPTXMCExpr.h:54
static const NVPTXFloatMCExpr * createConstantFPDouble(const APFloat &Flt, MCContext &Ctx)
Definition: NVPTXMCExpr.h:59
static const NVPTXGenericMCSymbolRefExpr * create(const MCSymbolRefExpr *SymExpr, MCContext &Ctx)
Definition: NVPTXMCExpr.cpp:59
static const char * getRegisterName(MCRegister Reg)
const char * getImageHandleSymbol(unsigned Idx) const
Returns the symbol name at the given index.
const char * getName(unsigned RegNo) const
std::string getTargetName() const
bool hasImageHandles() const
unsigned getMaxRequiredAlignment() const
bool hasMaskOperator() const
const NVPTXTargetLowering * getTargetLowering() const override
unsigned getPTXVersion() const
const NVPTXRegisterInfo * getRegisterInfo() const override
unsigned int getSmVersion() const
NVPTX::DrvInterface getDrvInterface() const
const NVPTXSubtarget * getSubtargetImpl(const Function &) const override
Virtual method implemented by subclasses that returns a reference to that target's TargetSubtargetInf...
UniqueStringSaver & getStrPool() const
Implments NVPTX-specific streamer.
void closeLastSection()
Close last section.
unsigned getAddressSpace() const
Return the address space of the Pointer type.
Definition: DerivedTypes.h:679
Wrapper class representing virtual and physical registers.
Definition: Register.h:19
static Register index2VirtReg(unsigned Index)
Convert a 0-based index to a virtual register number.
Definition: Register.h:84
constexpr bool isVirtual() const
Return true if the specified register number is in the virtual register namespace.
Definition: Register.h:91
static constexpr bool isVirtualRegister(unsigned Reg)
Return true if the specified register number is in the virtual register namespace.
Definition: Register.h:71
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
Definition: Register.h:95
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
Definition: SmallString.h:26
size_t size() const
Definition: SmallVector.h:91
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
Definition: SmallVector.h:586
void push_back(const T &Elt)
Definition: SmallVector.h:426
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1209
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:50
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
Definition: StringRef.h:257
constexpr const char * data() const
data - Get a pointer to the start of the string (which may not be null terminated).
Definition: StringRef.h:131
Class to represent struct types.
Definition: DerivedTypes.h:216
const STC & getSubtarget(const Function &F) const
This method returns a pointer to the specified type of TargetSubtargetInfo.
unsigned getPointerSizeInBits(unsigned AS) const
TargetRegisterInfo base class - We assume that the target defines a static array of TargetRegisterDes...
virtual const TargetRegisterInfo * getRegisterInfo() const
getRegisterInfo - If register information is available, return it.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
Definition: Twine.h:81
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
bool isVectorTy() const
True if this is an instance of VectorType.
Definition: Type.h:265
bool isPointerTy() const
True if this is an instance of PointerType.
Definition: Type.h:255
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
Definition: Type.h:146
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
@ ArrayTyID
Arrays.
Definition: Type.h:75
@ HalfTyID
16-bit floating point type
Definition: Type.h:56
@ VoidTyID
type with no size
Definition: Type.h:63
@ FloatTyID
32-bit floating point type
Definition: Type.h:58
@ StructTyID
Structures.
Definition: Type.h:74
@ IntegerTyID
Arbitrary bit width integers.
Definition: Type.h:71
@ FixedVectorTyID
Fixed width SIMD vector type.
Definition: Type.h:76
@ BFloatTyID
16-bit floating point type (7-bit significand)
Definition: Type.h:57
@ DoubleTyID
64-bit floating point type
Definition: Type.h:59
@ PointerTyID
Pointers.
Definition: Type.h:73
unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
bool isAggregateType() const
Return true if the type is an aggregate type.
Definition: Type.h:295
bool isHalfTy() const
Return true if this is 'half', a 16-bit IEEE fp type.
Definition: Type.h:143
bool isFloatingPointTy() const
Return true if this is one of the floating-point types.
Definition: Type.h:185
bool isIntOrPtrTy() const
Return true if this is an integer type or a pointer type.
Definition: Type.h:243
bool isIntegerTy() const
True if this is an instance of IntegerType.
Definition: Type.h:228
TypeID getTypeID() const
Return the type id for the type.
Definition: Type.h:137
TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
StringRef save(const char *S)
Definition: StringSaver.h:52
Value * getOperand(unsigned i) const
Definition: User.h:169
unsigned getNumOperands() const
Definition: User.h:191
LLVM Value Representation.
Definition: Value.h:74
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:255
bool use_empty() const
Definition: Value.h:344
StringRef getName() const
Return a constant reference to the value's name.
Definition: Value.cpp:309
std::pair< iterator, bool > insert(const ValueT &V)
Definition: DenseSet.h:206
size_type size() const
Definition: DenseSet.h:81
bool erase(const ValueT &V)
Definition: DenseSet.h:101
size_type count(const_arg_type_t< ValueT > V) const
Return 1 if the specified key is in the set, 0 otherwise.
Definition: DenseSet.h:97
This class implements an extremely fast bulk output stream that can only output to a stream.
Definition: raw_ostream.h:52
A raw_ostream that writes to an std::string.
Definition: raw_ostream.h:660
A raw_ostream that writes to an SmallVector or SmallString.
Definition: raw_ostream.h:690
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ C
The default llvm calling convention, compatible with C.
Definition: CallingConv.h:34
LegalityPredicate isVector(unsigned TypeIdx)
True iff the specified type index is a vector.
@ NVCL
Definition: NVPTX.h:79
@ CUDA
Definition: NVPTX.h:80
@ CE
Windows NT (Windows on ARM)
Reg
All possible values of the reg field in the ModR/M byte.
initializer< Ty > init(const Ty &Val)
Definition: CommandLine.h:450
constexpr double e
Definition: MathExtras.h:31
uint64_t read64le(const void *P)
Definition: Endian.h:412
uint32_t read32le(const void *P)
Definition: Endian.h:409
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:18
@ Offset
Definition: DWP.cpp:456
bool shouldEmitPTXNoReturn(const Value *V, const TargetMachine &TM)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
Definition: STLExtras.h:1689
std::string getSamplerName(const Value &val)
AddressSpace
Definition: NVPTXBaseInfo.h:21
@ ADDRESS_SPACE_LOCAL
Definition: NVPTXBaseInfo.h:26
@ ADDRESS_SPACE_CONST
Definition: NVPTXBaseInfo.h:25
@ ADDRESS_SPACE_GLOBAL
Definition: NVPTXBaseInfo.h:23
@ ADDRESS_SPACE_SHARED
Definition: NVPTXBaseInfo.h:24
bool getAlign(const Function &F, unsigned index, unsigned &align)
bool getMinCTASm(const Function &F, unsigned &x)
std::string getNVPTXRegClassName(TargetRegisterClass const *RC)
bool isImage(const Value &val)
bool getMaxNTIDz(const Function &F, unsigned &z)
Constant * ConstantFoldConstant(const Constant *C, const DataLayout &DL, const TargetLibraryInfo *TLI=nullptr)
ConstantFoldConstant - Fold the constant using the specified DataLayout.
bool isManaged(const Value &val)
unsigned promoteScalarArgumentSize(unsigned size)
bool isSurface(const Value &val)
void clearAnnotationCache(const Module *Mod)
void report_fatal_error(Error Err, bool gen_crash_diag=true)
Report a serious error, calling any installed error handler.
Definition: Error.cpp:156
std::string getSurfaceName(const Value &val)
bool getReqNTIDx(const Function &F, unsigned &x)
bool getReqNTIDy(const Function &F, unsigned &y)
FormattedNumber format_hex_no_prefix(uint64_t N, unsigned Width, bool Upper=false)
format_hex_no_prefix - Output N as a fixed width hexadecimal.
Definition: Format.h:200
bool getMaxNReg(const Function &F, unsigned &x)
bool isTexture(const Value &val)
bool isImageWriteOnly(const Value &val)
bool isImageReadWrite(const Value &val)
void write_hex(raw_ostream &S, uint64_t N, HexPrintStyle Style, std::optional< size_t > Width=std::nullopt)
std::string getTextureName(const Value &val)
void ComputeValueVTs(const TargetLowering &TLI, const DataLayout &DL, Type *Ty, SmallVectorImpl< EVT > &ValueVTs, SmallVectorImpl< EVT > *MemVTs, SmallVectorImpl< TypeSize > *Offsets=nullptr, TypeSize StartingOffset=TypeSize::getZero())
ComputeValueVTs - Given an LLVM IR type, compute a sequence of EVTs that represent all the individual...
Definition: Analysis.cpp:79
std::string getNVPTXRegClassStr(TargetRegisterClass const *RC)
Target & getTheNVPTXTarget64()
bool isKernelFunction(const Function &F)
bool getReqNTIDz(const Function &F, unsigned &z)
bool getMaxNTIDx(const Function &F, unsigned &x)
bool getMaxNTIDy(const Function &F, unsigned &y)
bool isSampler(const Value &val)
Constant * ConstantFoldIntegerCast(Constant *C, Type *DestTy, bool IsSigned, const DataLayout &DL)
Constant fold a zext, sext or trunc, depending on IsSigned and whether the DestTy is wider or narrowe...
bool getMaxClusterRank(const Function &F, unsigned &x)
MDNode * GetUnrollMetadata(MDNode *LoopID, StringRef Name)
Given an llvm.loop loop id metadata node, returns the loop hint metadata node with the given name (fo...
Definition: LoopUnroll.cpp:927
Target & getTheNVPTXTarget32()
static const fltSemantics & IEEEsingle() LLVM_READNONE
Definition: APFloat.cpp:249
static constexpr roundingMode rmNearestTiesToEven
Definition: APFloat.h:230
static const fltSemantics & IEEEdouble() LLVM_READNONE
Definition: APFloat.cpp:250
This struct is a compact representation of a valid (non-zero power of two) alignment.
Definition: Alignment.h:39
uint64_t value() const
This is a hole in the type system and should not be abused.
Definition: Alignment.h:85
Extended Value Type.
Definition: ValueTypes.h:34
TypeSize getSizeInBits() const
Return the size of the specified value type in bits.
Definition: ValueTypes.h:358
bool isInteger() const
Return true if this is an integer or a vector integer type.
Definition: ValueTypes.h:151
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
Definition: Alignment.h:117
Align valueOrOne() const
For convenience, returns a valid alignment or 1 if undefined.
Definition: Alignment.h:141
RegisterAsmPrinter - Helper template for registering a target specific assembly printer,...