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"
15#include "MCTargetDesc/NVPTXBaseInfo.h"
16#include "MCTargetDesc/NVPTXInstPrinter.h"
17#include "MCTargetDesc/NVPTXMCAsmInfo.h"
18#include "MCTargetDesc/NVPTXTargetStreamer.h"
19#include "NVPTX.h"
20#include "NVPTXDwarfDebug.h"
21#include "NVPTXMCExpr.h"
22#include "NVPTXMachineFunctionInfo.h"
23#include "NVPTXRegisterInfo.h"
24#include "NVPTXSubtarget.h"
25#include "NVPTXTargetMachine.h"
26#include "NVPTXUtilities.h"
27#include "NVVMProperties.h"
28#include "TargetInfo/NVPTXTargetInfo.h"
29#include "cl_common_defines.h"
30#include "llvm/ADT/APFloat.h"
31#include "llvm/ADT/APInt.h"
32#include "llvm/ADT/ArrayRef.h"
33#include "llvm/ADT/DenseMap.h"
34#include "llvm/ADT/DenseSet.h"
35#include "llvm/ADT/STLExtras.h"
36#include "llvm/ADT/Sequence.h"
37#include "llvm/ADT/SmallString.h"
38#include "llvm/ADT/SmallVector.h"
39#include "llvm/ADT/StringExtras.h"
40#include "llvm/ADT/StringRef.h"
41#include "llvm/ADT/Twine.h"
42#include "llvm/ADT/iterator_range.h"
43#include "llvm/Analysis/ConstantFolding.h"
44#include "llvm/CodeGen/Analysis.h"
45#include "llvm/CodeGen/MachineBasicBlock.h"
46#include "llvm/CodeGen/MachineFrameInfo.h"
47#include "llvm/CodeGen/MachineFunction.h"
48#include "llvm/CodeGen/MachineInstr.h"
49#include "llvm/CodeGen/MachineLoopInfo.h"
50#include "llvm/CodeGen/MachineModuleInfo.h"
51#include "llvm/CodeGen/MachineOperand.h"
52#include "llvm/CodeGen/MachineRegisterInfo.h"
53#include "llvm/CodeGen/TargetRegisterInfo.h"
54#include "llvm/CodeGen/ValueTypes.h"
55#include "llvm/CodeGenTypes/MachineValueType.h"
56#include "llvm/IR/Argument.h"
57#include "llvm/IR/Attributes.h"
58#include "llvm/IR/BasicBlock.h"
59#include "llvm/IR/Constant.h"
60#include "llvm/IR/Constants.h"
61#include "llvm/IR/DataLayout.h"
62#include "llvm/IR/DebugInfo.h"
63#include "llvm/IR/DebugInfoMetadata.h"
64#include "llvm/IR/DebugLoc.h"
65#include "llvm/IR/DerivedTypes.h"
66#include "llvm/IR/Function.h"
67#include "llvm/IR/GlobalAlias.h"
68#include "llvm/IR/GlobalValue.h"
69#include "llvm/IR/GlobalVariable.h"
70#include "llvm/IR/Instruction.h"
71#include "llvm/IR/LLVMContext.h"
72#include "llvm/IR/Module.h"
73#include "llvm/IR/Operator.h"
74#include "llvm/IR/Type.h"
75#include "llvm/IR/User.h"
76#include "llvm/MC/MCExpr.h"
77#include "llvm/MC/MCInst.h"
78#include "llvm/MC/MCInstrDesc.h"
79#include "llvm/MC/MCStreamer.h"
80#include "llvm/MC/MCSymbol.h"
81#include "llvm/MC/TargetRegistry.h"
82#include "llvm/Support/Alignment.h"
83#include "llvm/Support/Casting.h"
84#include "llvm/Support/Compiler.h"
85#include "llvm/Support/Endian.h"
86#include "llvm/Support/ErrorHandling.h"
87#include "llvm/Support/NativeFormatting.h"
88#include "llvm/Support/raw_ostream.h"
89#include "llvm/Target/TargetLoweringObjectFile.h"
90#include "llvm/Target/TargetMachine.h"
91#include "llvm/Transforms/Utils/UnrollLoop.h"
92#include <cassert>
93#include <cstdint>
94#include <cstring>
95#include <string>
96
97using namespace llvm;
98
99#define DEPOTNAME "__local_depot"
100
101static StringRef getTextureName(const Value &V) {
102 assert(V.hasName() && "Found texture variable with no name");
103 return V.getName();
104}
105
106static StringRef getSurfaceName(const Value &V) {
107 assert(V.hasName() && "Found surface variable with no name");
108 return V.getName();
109}
110
111static StringRef getSamplerName(const Value &V) {
112 assert(V.hasName() && "Found sampler variable with no name");
113 return V.getName();
114}
115
116/// Emits initial debug location directive.
117static void emitInitialRawDwarfLocDirective(const MachineFunction &MF,
118 DwarfDebug *DD,
119 MCStreamer &OutStreamer) {
120 if (!DD)
121 return;
122
123 assert(OutStreamer.hasRawTextSupport() && "Expected assembly output mode.");
124 // This is NVPTX specific and it's unclear why.
125 // PR51079: If we have code without debug information we need to give up.
126 const DISubprogram *SP = MF.getFunction().getSubprogram();
127 if (!SP)
128 return;
129 assert(SP->getUnit());
130 // NoDebug and DebugDirectivesOnly do not require emitting the initial loc
131 // directive. NoDebug does not require any debug directives and the initial
132 // loc directive is not needed for DebugDirectivesOnly as it is redundant
133 // assuming this is a non-empty function.
134 if (SP->getUnit()->isDebugDirectivesOnly() || SP->getUnit()->isNoDebug())
135 return;
136
137 (void)DD->emitInitialLocDirective(MF, /*CUID=*/0);
138}
139
140/// discoverDependentGlobals - Return a set of GlobalVariables on which \p V
141/// depends.
142static void
143discoverDependentGlobals(const Value *V,
144 DenseSet<const GlobalVariable *> &Globals) {
145 if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(Val: V)) {
146 Globals.insert(V: GV);
147 return;
148 }
149
150 if (const User *U = dyn_cast<User>(Val: V))
151 for (const auto &O : U->operands())
152 discoverDependentGlobals(V: O, Globals);
153}
154
155/// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable
156/// instances to be emitted, but only after any dependents have been added
157/// first.s
158static void
159VisitGlobalVariableForEmission(const GlobalVariable *GV,
160 SmallVectorImpl<const GlobalVariable *> &Order,
161 DenseSet<const GlobalVariable *> &Visited,
162 DenseSet<const GlobalVariable *> &Visiting) {
163 // Have we already visited this one?
164 if (Visited.count(V: GV))
165 return;
166
167 // Do we have a circular dependency?
168 if (!Visiting.insert(V: GV).second)
169 report_fatal_error(reason: "Circular dependency found in global variable set");
170
171 // Make sure we visit all dependents first
172 DenseSet<const GlobalVariable *> Others;
173 for (const auto &O : GV->operands())
174 discoverDependentGlobals(V: O, Globals&: Others);
175
176 for (const GlobalVariable *GV : Others)
177 VisitGlobalVariableForEmission(GV, Order, Visited, Visiting);
178
179 // Now we can visit ourself
180 Order.push_back(Elt: GV);
181 Visited.insert(V: GV);
182 Visiting.erase(V: GV);
183}
184
185void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {
186 NVPTX_MC::verifyInstructionPredicates(Opcode: MI->getOpcode(),
187 Features: getSubtargetInfo().getFeatureBits());
188
189 MCInst Inst;
190 lowerToMCInst(MI, OutMI&: Inst);
191 EmitToStreamer(S&: *OutStreamer, Inst);
192}
193
194void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
195 OutMI.setOpcode(MI->getOpcode());
196 for (const auto MO : MI->operands())
197 OutMI.addOperand(Op: lowerOperand(MO));
198}
199
200MCOperand NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
201 switch (MO.getType()) {
202 default:
203 llvm_unreachable("unknown operand type");
204 case MachineOperand::MO_Register:
205 return MCOperand::createReg(Reg: encodeVirtualRegister(Reg: MO.getReg()));
206 case MachineOperand::MO_Immediate:
207 return MCOperand::createImm(Val: MO.getImm());
208 case MachineOperand::MO_MachineBasicBlock:
209 return MCOperand::createExpr(
210 Val: MCSymbolRefExpr::create(Symbol: MO.getMBB()->getSymbol(), Ctx&: OutContext));
211 case MachineOperand::MO_ExternalSymbol:
212 return GetSymbolRef(Symbol: GetExternalSymbolSymbol(Sym: MO.getSymbolName()));
213 case MachineOperand::MO_GlobalAddress:
214 return GetSymbolRef(Symbol: getSymbol(GV: MO.getGlobal()));
215 case MachineOperand::MO_FPImmediate: {
216 const ConstantFP *Cnt = MO.getFPImm();
217 const APFloat &Val = Cnt->getValueAPF();
218
219 switch (Cnt->getType()->getTypeID()) {
220 default:
221 report_fatal_error(reason: "Unsupported FP type");
222 break;
223 case Type::HalfTyID:
224 return MCOperand::createExpr(
225 Val: NVPTXFloatMCExpr::createConstantFPHalf(Flt: Val, Ctx&: OutContext));
226 case Type::BFloatTyID:
227 return MCOperand::createExpr(
228 Val: NVPTXFloatMCExpr::createConstantBFPHalf(Flt: Val, Ctx&: OutContext));
229 case Type::FloatTyID:
230 return MCOperand::createExpr(
231 Val: NVPTXFloatMCExpr::createConstantFPSingle(Flt: Val, Ctx&: OutContext));
232 case Type::DoubleTyID:
233 return MCOperand::createExpr(
234 Val: NVPTXFloatMCExpr::createConstantFPDouble(Flt: Val, Ctx&: OutContext));
235 }
236 break;
237 }
238 }
239}
240
241unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
242 if (Register::isVirtualRegister(Reg)) {
243 const TargetRegisterClass *RC = MRI->getRegClass(Reg);
244
245 DenseMap<unsigned, unsigned> &RegMap = VRegMapping[RC];
246 unsigned RegNum = RegMap[Reg];
247
248 // Encode the register class in the upper 4 bits
249 // Must be kept in sync with NVPTXInstPrinter::printRegName
250 unsigned Ret = 0;
251 if (RC == &NVPTX::B1RegClass) {
252 Ret = (1 << 28);
253 } else if (RC == &NVPTX::B16RegClass) {
254 Ret = (2 << 28);
255 } else if (RC == &NVPTX::B32RegClass) {
256 Ret = (3 << 28);
257 } else if (RC == &NVPTX::B64RegClass) {
258 Ret = (4 << 28);
259 } else if (RC == &NVPTX::B128RegClass) {
260 Ret = (7 << 28);
261 } else {
262 report_fatal_error(reason: "Bad register class");
263 }
264
265 // Insert the vreg number
266 Ret |= (RegNum & 0x0FFFFFFF);
267 return Ret;
268 } else {
269 // Some special-use registers are actually physical registers.
270 // Encode this as the register class ID of 0 and the real register ID.
271 return Reg & 0x0FFFFFFF;
272 }
273}
274
275MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
276 const MCExpr *Expr;
277 Expr = MCSymbolRefExpr::create(Symbol, Ctx&: OutContext);
278 return MCOperand::createExpr(Val: Expr);
279}
280
281void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
282 const DataLayout &DL = getDataLayout();
283 const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(F: *F);
284 const auto *TLI = cast<NVPTXTargetLowering>(Val: STI.getTargetLowering());
285
286 Type *Ty = F->getReturnType();
287 // A void or zero-sized return type (e.g. an empty struct) produces no return
288 // parameter.
289 if (Ty->isVoidTy() || Ty->isEmptyTy())
290 return;
291 O << " (";
292
293 auto PrintScalarRetVal = [&](unsigned Size) {
294 O << ".param .b" << promoteScalarArgumentSize(size: Size) << " func_retval0";
295 };
296 if (shouldPassAsArray(Ty)) {
297 const unsigned TotalSize = DL.getTypeAllocSize(Ty);
298 const Align RetAlignment =
299 getPTXParamAlign(F, Ty, AttrIdx: AttributeList::ReturnIndex, DL);
300 O << ".param .align " << RetAlignment.value() << " .b8 func_retval0["
301 << TotalSize << "]";
302 } else if (Ty->isFloatingPointTy()) {
303 PrintScalarRetVal(Ty->getPrimitiveSizeInBits());
304 } else if (auto *ITy = dyn_cast<IntegerType>(Val: Ty)) {
305 PrintScalarRetVal(ITy->getBitWidth());
306 } else if (isa<PointerType>(Val: Ty)) {
307 PrintScalarRetVal(TLI->getPointerTy(DL).getSizeInBits());
308 } else
309 llvm_unreachable("Unknown return type");
310 O << ") ";
311}
312
313void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
314 raw_ostream &O) {
315 const Function &F = MF.getFunction();
316 printReturnValStr(F: &F, O);
317}
318
319void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
320 unsigned UniqueCallSite,
321 raw_ostream &O) const {
322 const DataLayout &DL = getDataLayout();
323 const NVPTXSubtarget &STI = MF->getSubtarget<NVPTXSubtarget>();
324 const auto *TLI = cast<NVPTXTargetLowering>(Val: STI.getTargetLowering());
325 const auto PtrVT = TLI->getPointerTy(DL);
326 Type *RetTy = CB.getFunctionType()->getReturnType();
327
328 O << "prototype_" << UniqueCallSite << " : .callprototype ";
329
330 if (RetTy->isVoidTy() || RetTy->isEmptyTy()) {
331 O << "()";
332 } else {
333 O << "(";
334 if (shouldPassAsArray(Ty: RetTy)) {
335 const Align RetAlign =
336 getPTXParamAlign(CB: &CB, Ty: RetTy, AttrIdx: AttributeList::ReturnIndex, DL);
337 O << ".param .align " << RetAlign.value() << " .b8 _["
338 << DL.getTypeAllocSize(Ty: RetTy) << "]";
339 } else if (RetTy->isFloatingPointTy() || RetTy->isIntegerTy()) {
340 unsigned size = 0;
341 if (auto *ITy = dyn_cast<IntegerType>(Val: RetTy)) {
342 size = ITy->getBitWidth();
343 } else {
344 assert(RetTy->isFloatingPointTy() &&
345 "Floating point type expected here");
346 size = RetTy->getPrimitiveSizeInBits();
347 }
348 // PTX ABI requires all scalar return values to be at least 32
349 // bits in size. fp16 normally uses .b16 as its storage type in
350 // PTX, so its size must be adjusted here, too.
351 size = promoteScalarArgumentSize(size);
352
353 O << ".param .b" << size << " _";
354 } else if (isa<PointerType>(Val: RetTy)) {
355 O << ".param .b" << PtrVT.getSizeInBits() << " _";
356 } else {
357 llvm_unreachable("Unknown return type");
358 }
359 O << ") ";
360 }
361 O << "_ (";
362
363 auto MakeArg = [&](const unsigned I) {
364 Type *Ty = CB.getArgOperand(i: I)->getType();
365
366 if (CB.paramHasAttr(ArgNo: I, Kind: Attribute::ByVal)) {
367 Type *ETy = CB.getParamByValType(ArgNo: I);
368 Align ParamByValAlign = getDeviceByValParamAlign(
369 CB: &CB, ArgTy: ETy, AttrIdx: I + AttributeList::FirstArgIndex, DL);
370
371 O << ".param .align " << ParamByValAlign.value() << " .b8 _["
372 << DL.getTypeAllocSize(Ty: ETy) << "]";
373 return;
374 }
375
376 if (shouldPassAsArray(Ty)) {
377 Align ParamAlign =
378 getPTXParamAlign(CB: &CB, Ty, AttrIdx: I + AttributeList::FirstArgIndex, DL);
379 O << ".param .align " << ParamAlign.value() << " .b8 _["
380 << DL.getTypeAllocSize(Ty) << "]";
381 return;
382 }
383 // scalar type
384 unsigned sz = 0;
385 if (auto *ITy = dyn_cast<IntegerType>(Val: Ty)) {
386 sz = promoteScalarArgumentSize(size: ITy->getBitWidth());
387 } else if (isa<PointerType>(Val: Ty)) {
388 sz = PtrVT.getSizeInBits();
389 } else {
390 sz = Ty->getPrimitiveSizeInBits();
391 }
392 O << ".param .b" << sz << " _";
393 };
394
395 const FunctionType *FTy = CB.getFunctionType();
396 const unsigned NumArgs = FTy->getNumParams();
397
398 // Zero-sized arguments (e.g. empty structs) are not passed and so do not
399 // appear in the prototype.
400 const auto NonEmptyArgs = make_filter_range(Range: seq(Size: NumArgs), Pred: [&](unsigned I) {
401 return !CB.getArgOperand(i: I)->getType()->isEmptyTy();
402 });
403
404 interleave(c: NonEmptyArgs, os&: O, each_fn: MakeArg, separator: ", ");
405
406 if (FTy->isVarArg() && CB.arg_size() > NumArgs)
407 O << (NonEmptyArgs.empty() ? "" : ",") << " .param .align "
408 << STI.getMaxRequiredAlignment() << " .b8 _[]";
409
410 O << ")";
411 if (shouldEmitPTXNoReturn(V: &CB, TM))
412 O << " .noreturn";
413 O << ";\n";
414}
415
416// Return true if MBB is the header of a loop marked with
417// llvm.loop.unroll.disable or llvm.loop.unroll.count=1.
418bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
419 const MachineBasicBlock &MBB) const {
420 MachineLoopInfo &LI = getAnalysis<MachineLoopInfoWrapperPass>().getLI();
421 // We insert .pragma "nounroll" only to the loop header.
422 if (!LI.isLoopHeader(BB: &MBB))
423 return false;
424
425 // llvm.loop.unroll.disable is marked on the back edges of a loop. Therefore,
426 // we iterate through each back edge of the loop with header MBB, and check
427 // whether its metadata contains llvm.loop.unroll.disable.
428 for (const MachineBasicBlock *PMBB : MBB.predecessors()) {
429 if (LI.getLoopFor(BB: PMBB) != LI.getLoopFor(BB: &MBB)) {
430 // Edges from other loops to MBB are not back edges.
431 continue;
432 }
433 if (const BasicBlock *PBB = PMBB->getBasicBlock()) {
434 if (MDNode *LoopID =
435 PBB->getTerminator()->getMetadata(KindID: LLVMContext::MD_loop)) {
436 if (GetUnrollMetadata(LoopID, Name: "llvm.loop.unroll.disable"))
437 return true;
438 if (MDNode *UnrollCountMD =
439 GetUnrollMetadata(LoopID, Name: "llvm.loop.unroll.count")) {
440 if (mdconst::extract<ConstantInt>(MD: UnrollCountMD->getOperand(I: 1))
441 ->isOne())
442 return true;
443 }
444 }
445 }
446 }
447 return false;
448}
449
450void NVPTXAsmPrinter::emitBasicBlockStart(const MachineBasicBlock &MBB) {
451 AsmPrinter::emitBasicBlockStart(MBB);
452 if (isLoopHeaderOfNoUnroll(MBB))
453 OutStreamer->emitRawText(String: StringRef("\t.pragma \"nounroll\";\n"));
454}
455
456void NVPTXAsmPrinter::emitFunctionEntryLabel() {
457 SmallString<128> Str;
458 raw_svector_ostream O(Str);
459
460 if (!GlobalsEmitted) {
461 emitGlobals(M: *MF->getFunction().getParent());
462 GlobalsEmitted = true;
463 }
464
465 // Set up
466 MRI = &MF->getRegInfo();
467 F = &MF->getFunction();
468 emitLinkageDirective(V: F, O);
469 if (isKernelFunction(F: *F))
470 O << ".entry ";
471 else {
472 O << ".func ";
473 printReturnValStr(MF: *MF, O);
474 }
475
476 CurrentFnSym->print(OS&: O, MAI);
477
478 emitFunctionParamList(F, O);
479 O << "\n";
480
481 if (isKernelFunction(F: *F))
482 emitKernelFunctionDirectives(F: *F, O);
483
484 if (shouldEmitPTXNoReturn(V: F, TM))
485 O << ".noreturn";
486
487 OutStreamer->emitRawText(String: O.str());
488
489 VRegMapping.clear();
490 // Emit open brace for function body.
491 OutStreamer->emitRawText(String: StringRef("{\n"));
492 setAndEmitFunctionVirtualRegisters(*MF);
493 encodeDebugInfoRegisterNumbers(MF: *MF);
494 // Emit initial .loc debug directive for correct relocation symbol data.
495 emitInitialRawDwarfLocDirective(MF: *MF, DD: getDwarfDebug(), OutStreamer&: *OutStreamer);
496}
497
498bool NVPTXAsmPrinter::runOnMachineFunction(MachineFunction &F) {
499 bool Result = AsmPrinter::runOnMachineFunction(MF&: F);
500 // Emit closing brace for the body of function F.
501 // The closing brace must be emitted here because we need to emit additional
502 // debug labels/data after the last basic block.
503 // We need to emit the closing brace here because we don't have function that
504 // finished emission of the function body.
505 OutStreamer->emitRawText(String: StringRef("}\n"));
506 return Result;
507}
508
509void NVPTXAsmPrinter::emitFunctionBodyStart() {
510 SmallString<128> Str;
511 raw_svector_ostream O(Str);
512 emitDemotedVars(&MF->getFunction(), O);
513
514 const auto *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
515 for (const auto &[Id, CB] : MFI->getCallPrototypes())
516 emitCallPrototype(CB: *CB, UniqueCallSite: Id, O);
517
518 OutStreamer->emitRawText(String: O.str());
519}
520
521void NVPTXAsmPrinter::emitFunctionBodyEnd() {
522 VRegMapping.clear();
523}
524
525const MCSymbol *NVPTXAsmPrinter::getFunctionFrameSymbol() const {
526 SmallString<128> Str;
527 raw_svector_ostream(Str) << DEPOTNAME << getFunctionNumber();
528 return OutContext.getOrCreateSymbol(Name: Str);
529}
530
531void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const {
532 Register RegNo = MI->getOperand(i: 0).getReg();
533 if (RegNo.isVirtual()) {
534 OutStreamer->AddComment(T: Twine("implicit-def: ") +
535 getVirtualRegisterName(RegNo));
536 } else {
537 const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
538 OutStreamer->AddComment(T: Twine("implicit-def: ") +
539 STI.getRegisterInfo()->getName(RegNo));
540 }
541 OutStreamer->addBlankLine();
542}
543
544void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F,
545 raw_ostream &O) const {
546 // If the NVVM IR has some of reqntid* specified, then output
547 // the reqntid directive, and set the unspecified ones to 1.
548 // If none of Reqntid* is specified, don't output reqntid directive.
549 const auto ReqNTID = getReqNTID(F);
550 if (!ReqNTID.empty())
551 O << formatv(Fmt: ".reqntid {0:$[, ]}\n",
552 Vals: make_range(x: ReqNTID.begin(), y: ReqNTID.end()));
553
554 const auto MaxNTID = getMaxNTID(F);
555 if (!MaxNTID.empty())
556 O << formatv(Fmt: ".maxntid {0:$[, ]}\n",
557 Vals: make_range(x: MaxNTID.begin(), y: MaxNTID.end()));
558
559 if (const auto Mincta = getMinCTASm(F))
560 O << ".minnctapersm " << *Mincta << "\n";
561
562 if (const auto Maxnreg = getMaxNReg(F))
563 O << ".maxnreg " << *Maxnreg << "\n";
564
565 // .maxclusterrank directive requires SM_90 or higher, make sure that we
566 // filter it out for lower SM versions, as it causes a hard ptxas crash.
567 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
568 const NVPTXSubtarget *STI = &NTM.getSubtarget<NVPTXSubtarget>(F);
569
570 if (STI->getSmVersion() >= 90) {
571 const auto ClusterDim = getClusterDim(F);
572 const bool BlocksAreClusters = hasBlocksAreClusters(F);
573
574 if (!ClusterDim.empty()) {
575
576 if (!BlocksAreClusters)
577 O << ".explicitcluster\n";
578
579 if (ClusterDim[0] != 0) {
580 assert(llvm::all_of(ClusterDim, not_equal_to(0)) &&
581 "cluster_dim_x != 0 implies cluster_dim_y and cluster_dim_z "
582 "should be non-zero as well");
583
584 O << formatv(Fmt: ".reqnctapercluster {0:$[, ]}\n",
585 Vals: make_range(x: ClusterDim.begin(), y: ClusterDim.end()));
586 } else {
587 assert(llvm::all_of(ClusterDim, equal_to(0)) &&
588 "cluster_dim_x == 0 implies cluster_dim_y and cluster_dim_z "
589 "should be 0 as well");
590 }
591 }
592
593 if (BlocksAreClusters) {
594 LLVMContext &Ctx = F.getContext();
595 if (ReqNTID.empty() || ClusterDim.empty())
596 Ctx.diagnose(DI: DiagnosticInfoUnsupported(
597 F, "blocksareclusters requires reqntid and cluster_dim attributes",
598 F.getSubprogram()));
599 else if (STI->getPTXVersion() < 90)
600 Ctx.diagnose(DI: DiagnosticInfoUnsupported(
601 F, "blocksareclusters requires PTX version >= 9.0",
602 F.getSubprogram()));
603 else
604 O << ".blocksareclusters\n";
605 }
606
607 if (const auto Maxclusterrank = getMaxClusterRank(F))
608 O << ".maxclusterrank " << *Maxclusterrank << "\n";
609 }
610}
611
612std::string NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const {
613 const TargetRegisterClass *RC = MRI->getRegClass(Reg);
614
615 std::string Name;
616 raw_string_ostream NameStr(Name);
617
618 VRegRCMap::const_iterator I = VRegMapping.find(Val: RC);
619 assert(I != VRegMapping.end() && "Bad register class");
620 const DenseMap<unsigned, unsigned> &RegMap = I->second;
621
622 VRegMap::const_iterator VI = RegMap.find(Val: Reg);
623 assert(VI != RegMap.end() && "Bad virtual register");
624 unsigned MappedVR = VI->second;
625
626 NameStr << getNVPTXRegClassStr(RC) << MappedVR;
627
628 return Name;
629}
630
631void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr,
632 raw_ostream &O) {
633 O << getVirtualRegisterName(Reg: vr);
634}
635
636void NVPTXAsmPrinter::emitAliasDeclaration(const GlobalAlias *GA,
637 raw_ostream &O) {
638 const Function *F = dyn_cast_or_null<Function>(Val: GA->getAliaseeObject());
639 if (!F || isKernelFunction(F: *F) || F->isDeclaration())
640 report_fatal_error(
641 reason: "NVPTX aliasee must be a non-kernel function definition");
642
643 if (GA->hasLinkOnceLinkage() || GA->hasWeakLinkage() ||
644 GA->hasAvailableExternallyLinkage() || GA->hasCommonLinkage())
645 report_fatal_error(reason: "NVPTX aliasee must not be '.weak'");
646
647 emitDeclarationWithName(F, getSymbol(GV: GA), O);
648}
649
650void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) {
651 emitDeclarationWithName(F, getSymbol(GV: F), O);
652}
653
654void NVPTXAsmPrinter::emitDeclarationWithName(const Function *F, MCSymbol *S,
655 raw_ostream &O) {
656 emitLinkageDirective(V: F, O);
657 if (isKernelFunction(F: *F))
658 O << ".entry ";
659 else
660 O << ".func ";
661 printReturnValStr(F, O);
662 S->print(OS&: O, MAI);
663 O << "\n";
664 emitFunctionParamList(F, O);
665 O << "\n";
666 if (shouldEmitPTXNoReturn(V: F, TM))
667 O << ".noreturn";
668 O << ";\n";
669}
670
671static bool usedInGlobalVarDef(const Constant *C) {
672 if (!C)
673 return false;
674
675 if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(Val: C))
676 return GV->getName() != "llvm.used";
677
678 for (const User *U : C->users())
679 if (const Constant *C = dyn_cast<Constant>(Val: U))
680 if (usedInGlobalVarDef(C))
681 return true;
682
683 return false;
684}
685
686static bool usedInOneFunc(const User *U, Function const *&OneFunc) {
687 if (const GlobalVariable *OtherGV = dyn_cast<GlobalVariable>(Val: U))
688 if (OtherGV->getName() == "llvm.used")
689 return true;
690
691 if (const Instruction *I = dyn_cast<Instruction>(Val: U)) {
692 if (const Function *CurFunc = I->getFunction()) {
693 if (OneFunc && (CurFunc != OneFunc))
694 return false;
695 OneFunc = CurFunc;
696 return true;
697 }
698 return false;
699 }
700
701 for (const User *UU : U->users())
702 if (!usedInOneFunc(U: UU, OneFunc))
703 return false;
704
705 return true;
706}
707
708/* Find out if a global variable can be demoted to local scope.
709 * Currently, this is valid for CUDA shared variables, which have local
710 * scope and global lifetime. So the conditions to check are :
711 * 1. Is the global variable in shared address space?
712 * 2. Does it have local linkage?
713 * 3. Is the global variable referenced only in one function?
714 */
715static bool canDemoteGlobalVar(const GlobalVariable *GV, Function const *&f) {
716 if (!GV->hasLocalLinkage())
717 return false;
718 if (GV->getAddressSpace() != ADDRESS_SPACE_SHARED)
719 return false;
720
721 const Function *oneFunc = nullptr;
722
723 bool flag = usedInOneFunc(U: GV, OneFunc&: oneFunc);
724 if (!flag)
725 return false;
726 if (!oneFunc)
727 return false;
728 f = oneFunc;
729 return true;
730}
731
732static bool useFuncSeen(const Constant *C,
733 const SmallPtrSetImpl<const Function *> &SeenSet) {
734 for (const User *U : C->users()) {
735 if (const Constant *cu = dyn_cast<Constant>(Val: U)) {
736 if (useFuncSeen(C: cu, SeenSet))
737 return true;
738 } else if (const Instruction *I = dyn_cast<Instruction>(Val: U)) {
739 if (const Function *Caller = I->getFunction())
740 if (SeenSet.contains(Ptr: Caller))
741 return true;
742 }
743 }
744 return false;
745}
746
747void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {
748 SmallPtrSet<const Function *, 32> SeenSet;
749 for (const Function &F : M) {
750 if (F.getAttributes().hasFnAttr(Kind: "nvptx-libcall-callee")) {
751 emitDeclaration(F: &F, O);
752 continue;
753 }
754
755 if (F.isDeclaration()) {
756 if (F.use_empty())
757 continue;
758 if (F.getIntrinsicID())
759 continue;
760 // An unrecognized intrinsic would produce an invalid PTX declaration. Let
761 // the user know that, and skip it.
762 if (F.isIntrinsic()) {
763 LLVMContext &Ctx = F.getContext();
764 Ctx.diagnose(DI: DiagnosticInfoUnsupported(
765 F, "unknown intrinsic '" + F.getName() +
766 "' cannot be lowered by the NVPTX backend"));
767 continue;
768 }
769 emitDeclaration(F: &F, O);
770 continue;
771 }
772 for (const User *U : F.users()) {
773 if (const Constant *C = dyn_cast<Constant>(Val: U)) {
774 if (usedInGlobalVarDef(C)) {
775 // The use is in the initialization of a global variable
776 // that is a function pointer, so print a declaration
777 // for the original function
778 emitDeclaration(F: &F, O);
779 break;
780 }
781 // Emit a declaration of this function if the function that
782 // uses this constant expr has already been seen.
783 if (useFuncSeen(C, SeenSet)) {
784 emitDeclaration(F: &F, O);
785 break;
786 }
787 }
788
789 if (!isa<Instruction>(Val: U))
790 continue;
791 const Function *Caller = cast<Instruction>(Val: U)->getFunction();
792 if (!Caller)
793 continue;
794
795 // If a caller has already been seen, then the caller is
796 // appearing in the module before the callee. so print out
797 // a declaration for the callee.
798 if (SeenSet.contains(Ptr: Caller)) {
799 emitDeclaration(F: &F, O);
800 break;
801 }
802 }
803 SeenSet.insert(Ptr: &F);
804 }
805 for (const GlobalAlias &GA : M.aliases())
806 emitAliasDeclaration(GA: &GA, O);
807}
808
809void NVPTXAsmPrinter::emitStartOfAsmFile(Module &M) {
810 // Construct a default subtarget off of the TargetMachine defaults. The
811 // rest of NVPTX isn't friendly to change subtargets per function and
812 // so the default TargetMachine will have all of the options.
813 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
814 const NVPTXSubtarget *STI = NTM.getSubtargetImpl();
815
816 // Emit header before any dwarf directives are emitted below.
817 emitHeader(M, STI: *STI);
818}
819
820/// Create NVPTX-specific DwarfDebug handler.
821DwarfDebug *NVPTXAsmPrinter::createDwarfDebug() {
822 return new NVPTXDwarfDebug(this);
823}
824
825bool NVPTXAsmPrinter::doInitialization(Module &M) {
826 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
827 const NVPTXSubtarget &STI = *NTM.getSubtargetImpl();
828 if (M.alias_size() && (STI.getPTXVersion() < 63 || STI.getSmVersion() < 30))
829 report_fatal_error(reason: ".alias requires PTX version >= 6.3 and sm_30");
830
831 // We need to call the parent's one explicitly.
832 bool Result = AsmPrinter::doInitialization(M);
833
834 GlobalsEmitted = false;
835
836 return Result;
837}
838
839void NVPTXAsmPrinter::emitGlobals(const Module &M) {
840 SmallString<128> Str2;
841 raw_svector_ostream OS2(Str2);
842
843 emitDeclarations(M, O&: OS2);
844
845 // As ptxas does not support forward references of globals, we need to first
846 // sort the list of module-level globals in def-use order. We visit each
847 // global variable in order, and ensure that we emit it *after* its dependent
848 // globals. We use a little extra memory maintaining both a set and a list to
849 // have fast searches while maintaining a strict ordering.
850 SmallVector<const GlobalVariable *, 8> Globals;
851 DenseSet<const GlobalVariable *> GVVisited;
852 DenseSet<const GlobalVariable *> GVVisiting;
853
854 // Visit each global variable, in order
855 for (const GlobalVariable &I : M.globals())
856 VisitGlobalVariableForEmission(GV: &I, Order&: Globals, Visited&: GVVisited, Visiting&: GVVisiting);
857
858 assert(GVVisited.size() == M.global_size() && "Missed a global variable");
859 assert(GVVisiting.size() == 0 && "Did not fully process a global variable");
860
861 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
862 const NVPTXSubtarget &STI = *NTM.getSubtargetImpl();
863
864 // Print out module-level global variables in proper order
865 for (const GlobalVariable *GV : Globals)
866 printModuleLevelGV(GVar: GV, O&: OS2, /*ProcessDemoted=*/processDemoted: false, STI);
867
868 OS2 << '\n';
869
870 OutStreamer->emitRawText(String: OS2.str());
871}
872
873void NVPTXAsmPrinter::emitGlobalAlias(const Module &M, const GlobalAlias &GA) {
874 SmallString<128> Str;
875 raw_svector_ostream OS(Str);
876
877 MCSymbol *Name = getSymbol(GV: &GA);
878
879 OS << ".alias " << Name->getName() << ", " << GA.getAliaseeObject()->getName()
880 << ";\n";
881
882 OutStreamer->emitRawText(String: OS.str());
883}
884
885NVPTXTargetStreamer *NVPTXAsmPrinter::getTargetStreamer() const {
886 return static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer());
887}
888
889static bool hasFullDebugInfo(Module &M) {
890 for (DICompileUnit *CU : M.debug_compile_units()) {
891 switch(CU->getEmissionKind()) {
892 case DICompileUnit::NoDebug:
893 case DICompileUnit::DebugDirectivesOnly:
894 break;
895 case DICompileUnit::LineTablesOnly:
896 case DICompileUnit::FullDebug:
897 return true;
898 }
899 }
900
901 return false;
902}
903
904void NVPTXAsmPrinter::emitHeader(Module &M, const NVPTXSubtarget &STI) {
905 auto *TS = getTargetStreamer();
906
907 TS->emitBanner();
908
909 const unsigned PTXVersion = STI.getPTXVersion();
910 TS->emitVersionDirective(PTXVersion);
911
912 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
913 bool TexModeIndependent = NTM.getDrvInterface() == NVPTX::NVCL;
914
915 TS->emitTargetDirective(Target: STI.getTargetName(), TexModeIndependent,
916 HasDebug: hasFullDebugInfo(M));
917 TS->emitAddressSizeDirective(AddrSize: M.getDataLayout().getPointerSizeInBits());
918}
919
920bool NVPTXAsmPrinter::doFinalization(Module &M) {
921 // If we did not emit any functions, then the global declarations have not
922 // yet been emitted.
923 if (!GlobalsEmitted) {
924 emitGlobals(M);
925 GlobalsEmitted = true;
926 }
927
928 // call doFinalization
929 bool ret = AsmPrinter::doFinalization(M);
930
931 clearAnnotationCache(&M);
932
933 auto *TS =
934 static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer());
935 // Close the last emitted section
936 if (hasDebugInfo()) {
937 TS->closeLastSection();
938 // Emit empty .debug_macinfo section for better support of the empty files.
939 OutStreamer->emitRawText(String: "\t.section\t.debug_macinfo\t{\t}");
940 }
941
942 // Output last DWARF .file directives, if any.
943 TS->outputDwarfFileDirectives();
944
945 return ret;
946}
947
948// This function emits appropriate linkage directives for
949// functions and global variables.
950//
951// extern function declaration -> .extern
952// extern function definition -> .visible
953// external global variable with init -> .visible
954// external without init -> .extern
955// appending -> not allowed, assert.
956// for any linkage other than
957// internal, private, linker_private,
958// linker_private_weak, linker_private_weak_def_auto,
959// we emit -> .weak.
960
961void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V,
962 raw_ostream &O) {
963 if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA) {
964 if (V->hasExternalLinkage()) {
965 if (const auto *GVar = dyn_cast<GlobalVariable>(Val: V))
966 O << (GVar->hasInitializer() ? ".visible " : ".extern ");
967 else if (V->isDeclaration())
968 O << ".extern ";
969 else
970 O << ".visible ";
971 } else if (V->hasAppendingLinkage()) {
972 report_fatal_error(reason: "Symbol '" + (V->hasName() ? V->getName() : "") +
973 "' has unsupported appending linkage type");
974 } else if (!V->hasInternalLinkage() && !V->hasPrivateLinkage()) {
975 O << ".weak ";
976 }
977 }
978}
979
980void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
981 raw_ostream &O, bool ProcessDemoted,
982 const NVPTXSubtarget &STI) {
983 // Skip meta data
984 if (GVar->hasSection())
985 if (GVar->getSection() == "llvm.metadata")
986 return;
987
988 // Skip LLVM intrinsic global variables
989 if (GVar->getName().starts_with(Prefix: "llvm.") ||
990 GVar->getName().starts_with(Prefix: "nvvm."))
991 return;
992
993 const DataLayout &DL = getDataLayout();
994
995 // GlobalVariables are always constant pointers themselves.
996 Type *ETy = GVar->getValueType();
997
998 if (GVar->hasExternalLinkage()) {
999 if (GVar->hasInitializer())
1000 O << ".visible ";
1001 else
1002 O << ".extern ";
1003 } else if (STI.getPTXVersion() >= 50 && GVar->hasCommonLinkage() &&
1004 GVar->getAddressSpace() == ADDRESS_SPACE_GLOBAL) {
1005 O << ".common ";
1006 } else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() ||
1007 GVar->hasAvailableExternallyLinkage() ||
1008 GVar->hasCommonLinkage()) {
1009 O << ".weak ";
1010 }
1011
1012 const PTXOpaqueType OpaqueType = getPTXOpaqueType(*GVar);
1013
1014 if (OpaqueType == PTXOpaqueType::Texture) {
1015 O << ".global .texref " << getTextureName(V: *GVar) << ";\n";
1016 return;
1017 }
1018
1019 if (OpaqueType == PTXOpaqueType::Surface) {
1020 O << ".global .surfref " << getSurfaceName(V: *GVar) << ";\n";
1021 return;
1022 }
1023
1024 if (GVar->isDeclaration()) {
1025 // (extern) declarations, no definition or initializer
1026 // Currently the only known declaration is for an automatic __local
1027 // (.shared) promoted to global.
1028 emitPTXGlobalVariable(GVar, O, STI);
1029 O << ";\n";
1030 return;
1031 }
1032
1033 if (OpaqueType == PTXOpaqueType::Sampler) {
1034 O << ".global .samplerref " << getSamplerName(V: *GVar);
1035
1036 const Constant *Initializer = nullptr;
1037 if (GVar->hasInitializer())
1038 Initializer = GVar->getInitializer();
1039 const ConstantInt *CI = nullptr;
1040 if (Initializer)
1041 CI = dyn_cast<ConstantInt>(Val: Initializer);
1042 if (CI) {
1043 unsigned sample = CI->getZExtValue();
1044
1045 O << " = { ";
1046
1047 for (int i = 0,
1048 addr = ((sample & __CLK_ADDRESS_MASK) >> __CLK_ADDRESS_BASE);
1049 i < 3; i++) {
1050 O << "addr_mode_" << i << " = ";
1051 switch (addr) {
1052 case 0:
1053 O << "wrap";
1054 break;
1055 case 1:
1056 O << "clamp_to_border";
1057 break;
1058 case 2:
1059 O << "clamp_to_edge";
1060 break;
1061 case 3:
1062 O << "wrap";
1063 break;
1064 case 4:
1065 O << "mirror";
1066 break;
1067 }
1068 O << ", ";
1069 }
1070 O << "filter_mode = ";
1071 switch ((sample & __CLK_FILTER_MASK) >> __CLK_FILTER_BASE) {
1072 case 0:
1073 O << "nearest";
1074 break;
1075 case 1:
1076 O << "linear";
1077 break;
1078 case 2:
1079 llvm_unreachable("Anisotropic filtering is not supported");
1080 default:
1081 O << "nearest";
1082 break;
1083 }
1084 if (!((sample & __CLK_NORMALIZED_MASK) >> __CLK_NORMALIZED_BASE)) {
1085 O << ", force_unnormalized_coords = 1";
1086 }
1087 O << " }";
1088 }
1089
1090 O << ";\n";
1091 return;
1092 }
1093
1094 if (GVar->hasPrivateLinkage()) {
1095 if (GVar->getName().starts_with(Prefix: "unrollpragma"))
1096 return;
1097
1098 // FIXME - need better way (e.g. Metadata) to avoid generating this global
1099 if (GVar->getName().starts_with(Prefix: "filename"))
1100 return;
1101 if (GVar->use_empty())
1102 return;
1103 }
1104
1105 const Function *DemotedFunc = nullptr;
1106 if (!ProcessDemoted && canDemoteGlobalVar(GV: GVar, f&: DemotedFunc)) {
1107 O << "// " << GVar->getName() << " has been demoted\n";
1108 localDecls[DemotedFunc].push_back(x: GVar);
1109 return;
1110 }
1111
1112 O << ".";
1113 emitPTXAddressSpace(AddressSpace: GVar->getAddressSpace(), O);
1114
1115 if (isManaged(*GVar)) {
1116 if (STI.getPTXVersion() < 40 || STI.getSmVersion() < 30)
1117 report_fatal_error(
1118 reason: ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1119 O << " .attribute(.managed)";
1120 }
1121
1122 O << " .align "
1123 << GVar->getAlign().value_or(u: DL.getPrefTypeAlign(Ty: ETy)).value();
1124
1125 if (ETy->isPointerTy() || ((ETy->isIntegerTy() || ETy->isFloatingPointTy()) &&
1126 ETy->getScalarSizeInBits() <= 64)) {
1127 O << " .";
1128 // Special case: ABI requires that we use .u8 for predicates
1129 if (ETy->isIntegerTy(BitWidth: 1))
1130 O << "u8";
1131 else
1132 O << getPTXFundamentalTypeStr(Ty: ETy, false);
1133 O << " ";
1134 getSymbol(GV: GVar)->print(OS&: O, MAI);
1135
1136 // Ptx allows variable initilization only for constant and global state
1137 // spaces.
1138 if (GVar->hasInitializer()) {
1139 if ((GVar->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1140 (GVar->getAddressSpace() == ADDRESS_SPACE_CONST)) {
1141 const Constant *Initializer = GVar->getInitializer();
1142 // 'undef' is treated as there is no value specified.
1143 if (!Initializer->isNullValue() && !isa<UndefValue>(Val: Initializer)) {
1144 O << " = ";
1145 printScalarConstant(CPV: Initializer, O);
1146 }
1147 } else {
1148 // The frontend adds zero-initializer to device and constant variables
1149 // that don't have an initial value, and UndefValue to shared
1150 // variables, so skip warning for this case.
1151 if (!GVar->getInitializer()->isNullValue() &&
1152 !isa<UndefValue>(Val: GVar->getInitializer())) {
1153 report_fatal_error(reason: "initial value of '" + GVar->getName() +
1154 "' is not allowed in addrspace(" +
1155 Twine(GVar->getAddressSpace()) + ")");
1156 }
1157 }
1158 }
1159 } else {
1160 // Although PTX has direct support for struct type and array type and
1161 // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
1162 // targets that support these high level field accesses. Structs, arrays
1163 // and vectors are lowered into arrays of bytes.
1164 switch (ETy->getTypeID()) {
1165 case Type::IntegerTyID: // Integers larger than 64 bits
1166 case Type::FP128TyID:
1167 case Type::StructTyID:
1168 case Type::ArrayTyID:
1169 case Type::FixedVectorTyID: {
1170 const uint64_t ElementSize = DL.getTypeStoreSize(Ty: ETy);
1171 // Ptx allows variable initilization only for constant and
1172 // global state spaces.
1173 if (((GVar->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
1174 (GVar->getAddressSpace() == ADDRESS_SPACE_CONST)) &&
1175 GVar->hasInitializer()) {
1176 const Constant *Initializer = GVar->getInitializer();
1177 if (!isa<UndefValue>(Val: Initializer) && !Initializer->isNullValue()) {
1178 AggBuffer aggBuffer(ElementSize, *this);
1179 bufferAggregateConstant(CV: Initializer, aggBuffer: &aggBuffer);
1180 if (aggBuffer.numSymbols()) {
1181 const unsigned int ptrSize = MAI.getCodePointerSize();
1182 if (ElementSize % ptrSize ||
1183 !aggBuffer.allSymbolsAligned(ptrSize)) {
1184 // Print in bytes and use the mask() operator for pointers.
1185 if (!STI.hasMaskOperator())
1186 report_fatal_error(
1187 reason: "initialized packed aggregate with pointers '" +
1188 GVar->getName() +
1189 "' requires at least PTX ISA version 7.1");
1190 O << " .u8 ";
1191 getSymbol(GV: GVar)->print(OS&: O, MAI);
1192 O << "[" << ElementSize << "] = {";
1193 aggBuffer.printBytes(os&: O);
1194 O << "}";
1195 } else {
1196 O << " .u" << ptrSize * 8 << " ";
1197 getSymbol(GV: GVar)->print(OS&: O, MAI);
1198 O << "[" << ElementSize / ptrSize << "] = {";
1199 aggBuffer.printWords(os&: O);
1200 O << "}";
1201 }
1202 } else {
1203 O << " .b8 ";
1204 getSymbol(GV: GVar)->print(OS&: O, MAI);
1205 O << "[" << ElementSize << "] = {";
1206 aggBuffer.printBytes(os&: O);
1207 O << "}";
1208 }
1209 } else {
1210 O << " .b8 ";
1211 getSymbol(GV: GVar)->print(OS&: O, MAI);
1212 if (ElementSize)
1213 O << "[" << ElementSize << "]";
1214 }
1215 } else {
1216 O << " .b8 ";
1217 getSymbol(GV: GVar)->print(OS&: O, MAI);
1218 if (ElementSize)
1219 O << "[" << ElementSize << "]";
1220 }
1221 break;
1222 }
1223 default:
1224 llvm_unreachable("type not supported yet");
1225 }
1226 }
1227 O << ";\n";
1228}
1229
1230void NVPTXAsmPrinter::AggBuffer::printSymbol(unsigned nSym, raw_ostream &os) {
1231 const Value *v = Symbols[nSym];
1232 const Value *v0 = SymbolsBeforeStripping[nSym];
1233 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(Val: v)) {
1234 MCSymbol *Name = AP.getSymbol(GV: GVar);
1235 PointerType *PTy = dyn_cast<PointerType>(Val: v0->getType());
1236 // Is v0 a generic pointer?
1237 bool isGenericPointer = PTy && PTy->getAddressSpace() == 0;
1238 if (EmitGeneric && isGenericPointer && !isa<Function>(Val: v)) {
1239 os << "generic(";
1240 Name->print(OS&: os, MAI: AP.MAI);
1241 os << ")";
1242 } else {
1243 Name->print(OS&: os, MAI: AP.MAI);
1244 }
1245 } else if (const ConstantExpr *CExpr = dyn_cast<ConstantExpr>(Val: v0)) {
1246 const MCExpr *Expr = AP.lowerConstantForGV(CV: CExpr, ProcessingGeneric: false);
1247 AP.printMCExpr(Expr: *Expr, OS&: os);
1248 } else
1249 llvm_unreachable("symbol type unknown");
1250}
1251
1252void NVPTXAsmPrinter::AggBuffer::printBytes(raw_ostream &os) {
1253 unsigned int ptrSize = AP.MAI.getCodePointerSize();
1254 // Do not emit trailing zero initializers. They will be zero-initialized by
1255 // ptxas. This saves on both space requirements for the generated PTX and on
1256 // memory use by ptxas. (See:
1257 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#global-state-space)
1258 unsigned int InitializerCount = Size;
1259 // TODO: symbols make this harder, but it would still be good to trim trailing
1260 // 0s for aggs with symbols as well.
1261 if (numSymbols() == 0)
1262 while (InitializerCount >= 1 && !buffer[InitializerCount - 1])
1263 InitializerCount--;
1264
1265 symbolPosInBuffer.push_back(Elt: InitializerCount);
1266 unsigned int nSym = 0;
1267 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1268 for (unsigned int pos = 0; pos < InitializerCount;) {
1269 if (pos)
1270 os << ", ";
1271 if (pos != nextSymbolPos) {
1272 os << (unsigned int)buffer[pos];
1273 ++pos;
1274 continue;
1275 }
1276 // Generate a per-byte mask() operator for the symbol, which looks like:
1277 // .global .u8 addr[] = {0xFF(foo), 0xFF00(foo), 0xFF0000(foo), ...};
1278 // See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#initializers
1279 std::string symText;
1280 llvm::raw_string_ostream oss(symText);
1281 printSymbol(nSym, os&: oss);
1282 for (unsigned i = 0; i < ptrSize; ++i) {
1283 if (i)
1284 os << ", ";
1285 llvm::write_hex(S&: os, N: 0xFFULL << i * 8, Style: HexPrintStyle::PrefixUpper);
1286 os << "(" << symText << ")";
1287 }
1288 pos += ptrSize;
1289 nextSymbolPos = symbolPosInBuffer[++nSym];
1290 assert(nextSymbolPos >= pos);
1291 }
1292}
1293
1294void NVPTXAsmPrinter::AggBuffer::printWords(raw_ostream &os) {
1295 unsigned int ptrSize = AP.MAI.getCodePointerSize();
1296 symbolPosInBuffer.push_back(Elt: Size);
1297 unsigned int nSym = 0;
1298 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1299 assert(nextSymbolPos % ptrSize == 0);
1300 for (unsigned int pos = 0; pos < Size; pos += ptrSize) {
1301 if (pos)
1302 os << ", ";
1303 if (pos == nextSymbolPos) {
1304 printSymbol(nSym, os);
1305 nextSymbolPos = symbolPosInBuffer[++nSym];
1306 assert(nextSymbolPos % ptrSize == 0);
1307 assert(nextSymbolPos >= pos + ptrSize);
1308 } else if (ptrSize == 4)
1309 os << support::endian::read32le(P: &buffer[pos]);
1310 else
1311 os << support::endian::read64le(P: &buffer[pos]);
1312 }
1313}
1314
1315void NVPTXAsmPrinter::emitDemotedVars(const Function *F, raw_ostream &O) {
1316 auto It = localDecls.find(x: F);
1317 if (It == localDecls.end())
1318 return;
1319
1320 ArrayRef<const GlobalVariable *> GVars = It->second;
1321
1322 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
1323 const NVPTXSubtarget &STI = *NTM.getSubtargetImpl();
1324
1325 for (const GlobalVariable *GV : GVars) {
1326 O << "\t// demoted variable\n\t";
1327 printModuleLevelGV(GVar: GV, O, /*processDemoted=*/ProcessDemoted: true, STI);
1328 }
1329}
1330
1331void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
1332 raw_ostream &O) const {
1333 switch (AddressSpace) {
1334 case ADDRESS_SPACE_LOCAL:
1335 O << "local";
1336 break;
1337 case ADDRESS_SPACE_GLOBAL:
1338 O << "global";
1339 break;
1340 case ADDRESS_SPACE_CONST:
1341 O << "const";
1342 break;
1343 case ADDRESS_SPACE_SHARED:
1344 O << "shared";
1345 break;
1346 default:
1347 report_fatal_error(reason: "Bad address space found while emitting PTX: " +
1348 llvm::Twine(AddressSpace));
1349 break;
1350 }
1351}
1352
1353std::string
1354NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {
1355 switch (Ty->getTypeID()) {
1356 case Type::IntegerTyID: {
1357 unsigned NumBits = cast<IntegerType>(Val: Ty)->getBitWidth();
1358 if (NumBits == 1)
1359 return "pred";
1360 if (NumBits <= 64) {
1361 std::string name = "u";
1362 return name + utostr(X: NumBits);
1363 }
1364 llvm_unreachable("Integer too large");
1365 break;
1366 }
1367 case Type::BFloatTyID:
1368 case Type::HalfTyID:
1369 // fp16 and bf16 are stored as .b16 for compatibility with pre-sm_53
1370 // PTX assembly.
1371 return "b16";
1372 case Type::FloatTyID:
1373 return "f32";
1374 case Type::DoubleTyID:
1375 return "f64";
1376 case Type::PointerTyID: {
1377 unsigned PtrSize = TM.getPointerSizeInBits(AS: Ty->getPointerAddressSpace());
1378 assert((PtrSize == 64 || PtrSize == 32) && "Unexpected pointer size");
1379
1380 if (PtrSize == 64)
1381 if (useB4PTR)
1382 return "b64";
1383 else
1384 return "u64";
1385 else if (useB4PTR)
1386 return "b32";
1387 else
1388 return "u32";
1389 }
1390 default:
1391 break;
1392 }
1393 llvm_unreachable("unexpected type");
1394}
1395
1396void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,
1397 raw_ostream &O,
1398 const NVPTXSubtarget &STI) {
1399 const DataLayout &DL = getDataLayout();
1400
1401 // GlobalVariables are always constant pointers themselves.
1402 Type *ETy = GVar->getValueType();
1403
1404 O << ".";
1405 emitPTXAddressSpace(AddressSpace: GVar->getType()->getAddressSpace(), O);
1406 if (isManaged(*GVar)) {
1407 if (STI.getPTXVersion() < 40 || STI.getSmVersion() < 30)
1408 report_fatal_error(
1409 reason: ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1410
1411 O << " .attribute(.managed)";
1412 }
1413 O << " .align "
1414 << GVar->getAlign().value_or(u: DL.getPrefTypeAlign(Ty: ETy)).value();
1415
1416 // Special case for i128/fp128
1417 if (ETy->getScalarSizeInBits() == 128) {
1418 O << " .b8 ";
1419 getSymbol(GV: GVar)->print(OS&: O, MAI);
1420 O << "[16]";
1421 return;
1422 }
1423
1424 if (ETy->isFloatingPointTy() || ETy->isIntOrPtrTy()) {
1425 O << " ." << getPTXFundamentalTypeStr(Ty: ETy) << " ";
1426 getSymbol(GV: GVar)->print(OS&: O, MAI);
1427 return;
1428 }
1429
1430 int64_t ElementSize = 0;
1431
1432 // Although PTX has direct support for struct type and array type and LLVM IR
1433 // is very similar to PTX, the LLVM CodeGen does not support for targets that
1434 // support these high level field accesses. Structs and arrays are lowered
1435 // into arrays of bytes.
1436 switch (ETy->getTypeID()) {
1437 case Type::StructTyID:
1438 case Type::ArrayTyID:
1439 case Type::FixedVectorTyID:
1440 ElementSize = DL.getTypeStoreSize(Ty: ETy);
1441 O << " .b8 ";
1442 getSymbol(GV: GVar)->print(OS&: O, MAI);
1443 O << "[";
1444 if (ElementSize) {
1445 O << ElementSize;
1446 }
1447 O << "]";
1448 break;
1449 default:
1450 llvm_unreachable("type not supported yet");
1451 }
1452}
1453
1454void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
1455 const DataLayout &DL = getDataLayout();
1456 const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(F: *F);
1457 const auto *TLI = cast<NVPTXTargetLowering>(Val: STI.getTargetLowering());
1458 const NVPTXMachineFunctionInfo *MFI =
1459 MF ? MF->getInfo<NVPTXMachineFunctionInfo>() : nullptr;
1460
1461 bool IsFirst = true;
1462 const bool IsKernelFunc = isKernelFunction(F: *F);
1463
1464 // Zero-sized arguments (e.g. empty structs) do not produce a parameter.
1465 // Number the emitted parameters contiguously, skipping the zero-sized ones,
1466 // so that the names match those used in LowerFormalArguments and the
1467 // contiguous numbering used by callers (see LowerCall).
1468 const auto NonEmptyArgs =
1469 make_filter_range(Range: F->args(), Pred: [](const Argument &Arg) {
1470 return !Arg.getType()->isEmptyTy();
1471 });
1472
1473 if (NonEmptyArgs.empty() && !F->isVarArg()) {
1474 O << "()";
1475 return;
1476 }
1477
1478 O << "(\n";
1479
1480 for (const auto &[ParamIndex, Arg] : enumerate(First: NonEmptyArgs)) {
1481 Type *Ty = Arg.getType();
1482 const std::string ParamSym = TLI->getParamName(F, Idx: ParamIndex);
1483
1484 if (!IsFirst)
1485 O << ",\n";
1486
1487 IsFirst = false;
1488
1489 // Handle image/sampler parameters
1490 if (IsKernelFunc) {
1491 const PTXOpaqueType ArgOpaqueType = getPTXOpaqueType(Arg);
1492 if (ArgOpaqueType != PTXOpaqueType::None) {
1493 const bool EmitImgPtr = !MFI || !MFI->checkImageHandleSymbol(Symbol: ParamSym);
1494 O << "\t.param ";
1495 if (EmitImgPtr)
1496 O << ".u64 .ptr ";
1497
1498 switch (ArgOpaqueType) {
1499 case PTXOpaqueType::Sampler:
1500 O << ".samplerref ";
1501 break;
1502 case PTXOpaqueType::Texture:
1503 O << ".texref ";
1504 break;
1505 case PTXOpaqueType::Surface:
1506 O << ".surfref ";
1507 break;
1508 case PTXOpaqueType::None:
1509 llvm_unreachable("handled above");
1510 }
1511 O << ParamSym;
1512 continue;
1513 }
1514 }
1515
1516 if (Arg.hasByValAttr()) {
1517 // param has byVal attribute.
1518 Type *ETy = Arg.getParamByValType();
1519 assert(ETy && "Param should have byval type");
1520
1521 // Print .param .align <a> .b8 .param[size];
1522 // <a> = optimal alignment for the element type; always multiple of
1523 // PAL.getParamAlignment
1524 // size = typeallocsize of element type
1525 const unsigned ParamIdx = Arg.getArgNo() + AttributeList::FirstArgIndex;
1526 const Align OptimalAlign =
1527 IsKernelFunc ? getPTXParamAlign(F, Ty: ETy, AttrIdx: ParamIdx, DL)
1528 : getDeviceByValParamAlign(F, ArgTy: ETy, AttrIdx: ParamIdx, DL);
1529
1530 O << "\t.param .align " << OptimalAlign.value() << " .b8 " << ParamSym
1531 << "[" << DL.getTypeAllocSize(Ty: ETy) << "]";
1532 continue;
1533 }
1534
1535 if (shouldPassAsArray(Ty)) {
1536 // Just print .param .align <a> .b8 .param[size];
1537 // <a> = optimal alignment for the element type; always multiple of
1538 // PAL.getParamAlignment
1539 // size = typeallocsize of element type
1540 Align OptimalAlign = getPTXParamAlign(
1541 F, Ty, AttrIdx: Arg.getArgNo() + AttributeList::FirstArgIndex, DL);
1542
1543 O << "\t.param .align " << OptimalAlign.value() << " .b8 " << ParamSym
1544 << "[" << DL.getTypeAllocSize(Ty) << "]";
1545
1546 continue;
1547 }
1548 // Just a scalar
1549 auto *PTy = dyn_cast<PointerType>(Val: Ty);
1550 unsigned PTySizeInBits = 0;
1551 if (PTy) {
1552 PTySizeInBits =
1553 TLI->getPointerTy(DL, AS: PTy->getAddressSpace()).getSizeInBits();
1554 assert(PTySizeInBits && "Invalid pointer size");
1555 }
1556
1557 if (IsKernelFunc) {
1558 if (PTy) {
1559 O << "\t.param .u" << PTySizeInBits << " .ptr";
1560
1561 switch (PTy->getAddressSpace()) {
1562 default:
1563 break;
1564 case ADDRESS_SPACE_GLOBAL:
1565 O << " .global";
1566 break;
1567 case ADDRESS_SPACE_SHARED:
1568 O << " .shared";
1569 break;
1570 case ADDRESS_SPACE_CONST:
1571 O << " .const";
1572 break;
1573 case ADDRESS_SPACE_LOCAL:
1574 O << " .local";
1575 break;
1576 }
1577
1578 O << " .align " << Arg.getParamAlign().valueOrOne().value() << " "
1579 << ParamSym;
1580 continue;
1581 }
1582
1583 // non-pointer scalar to kernel func
1584 O << "\t.param .";
1585 // Special case: predicate operands become .u8 types
1586 if (Ty->isIntegerTy(BitWidth: 1))
1587 O << "u8";
1588 else
1589 O << getPTXFundamentalTypeStr(Ty);
1590 O << " " << ParamSym;
1591 continue;
1592 }
1593 // Non-kernel function, just print .param .b<size> for ABI
1594 // and .reg .b<size> for non-ABI
1595 unsigned Size;
1596 if (auto *ITy = dyn_cast<IntegerType>(Val: Ty)) {
1597 Size = promoteScalarArgumentSize(size: ITy->getBitWidth());
1598 } else if (PTy) {
1599 assert(PTySizeInBits && "Invalid pointer size");
1600 Size = PTySizeInBits;
1601 } else
1602 Size = Ty->getPrimitiveSizeInBits();
1603 O << "\t.param .b" << Size << " " << ParamSym;
1604 }
1605
1606 if (F->isVarArg()) {
1607 if (!IsFirst)
1608 O << ",\n";
1609 O << "\t.param .align " << STI.getMaxRequiredAlignment() << " .b8 "
1610 << TLI->getParamName(F, /* vararg */ Idx: -1) << "[]";
1611 }
1612
1613 O << "\n)";
1614}
1615
1616void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
1617 const MachineFunction &MF) {
1618 SmallString<128> Str;
1619 raw_svector_ostream O(Str);
1620
1621 // Map the global virtual register number to a register class specific
1622 // virtual register number starting from 1 with that class.
1623 const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo();
1624
1625 // Emit the Fake Stack Object
1626 const MachineFrameInfo &MFI = MF.getFrameInfo();
1627 int64_t NumBytes = MFI.getStackSize();
1628 if (NumBytes) {
1629 O << "\t.local .align " << MFI.getMaxAlign().value() << " .b8 \t"
1630 << DEPOTNAME << getFunctionNumber() << "[" << NumBytes << "];\n";
1631 if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
1632 O << "\t.reg .b64 \t%SP;\n"
1633 << "\t.reg .b64 \t%SPL;\n";
1634 } else {
1635 O << "\t.reg .b32 \t%SP;\n"
1636 << "\t.reg .b32 \t%SPL;\n";
1637 }
1638 }
1639
1640 // Go through all virtual registers to establish the mapping between the
1641 // global virtual
1642 // register number and the per class virtual register number.
1643 // We use the per class virtual register number in the ptx output.
1644 for (unsigned I : llvm::seq(Size: MRI->getNumVirtRegs())) {
1645 Register VR = Register::index2VirtReg(Index: I);
1646 if (MRI->use_empty(RegNo: VR) && MRI->def_empty(RegNo: VR))
1647 continue;
1648 auto &RCRegMap = VRegMapping[MRI->getRegClass(Reg: VR)];
1649 RCRegMap[VR] = RCRegMap.size() + 1;
1650 }
1651
1652 // Emit declaration of the virtual registers or 'physical' registers for
1653 // each register class
1654 for (const TargetRegisterClass &RC : TRI->regclasses()) {
1655 const unsigned N = VRegMapping[&RC].size();
1656
1657 // Only declare those registers that may be used.
1658 if (N) {
1659 const StringRef RCName = getNVPTXRegClassName(RC: &RC);
1660 const StringRef RCStr = getNVPTXRegClassStr(RC: &RC);
1661 O << "\t.reg " << RCName << " \t" << RCStr << "<" << (N + 1) << ">;\n";
1662 }
1663 }
1664
1665 OutStreamer->emitRawText(String: O.str());
1666}
1667
1668/// Translate virtual register numbers in DebugInfo locations to their printed
1669/// encodings, as used by CUDA-GDB.
1670void NVPTXAsmPrinter::encodeDebugInfoRegisterNumbers(
1671 const MachineFunction &MF) {
1672 const NVPTXSubtarget &STI = MF.getSubtarget<NVPTXSubtarget>();
1673 const NVPTXRegisterInfo *registerInfo = STI.getRegisterInfo();
1674
1675 // Clear the old mapping, and add the new one. This mapping is used after the
1676 // printing of the current function is complete, but before the next function
1677 // is printed.
1678 registerInfo->clearDebugRegisterMap();
1679
1680 for (auto &classMap : VRegMapping) {
1681 for (auto &registerMapping : classMap.getSecond()) {
1682 auto reg = registerMapping.getFirst();
1683 registerInfo->addToDebugRegisterMap(preEncodedVirtualRegister: reg, RegisterName: getVirtualRegisterName(Reg: reg));
1684 }
1685 }
1686}
1687
1688void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp,
1689 raw_ostream &O) const {
1690 APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
1691 bool ignored;
1692 unsigned int numHex;
1693 const char *lead;
1694
1695 if (Fp->getType()->getTypeID() == Type::FloatTyID) {
1696 numHex = 8;
1697 lead = "0f";
1698 APF.convert(ToSemantics: APFloat::IEEEsingle(), RM: APFloat::rmNearestTiesToEven, losesInfo: &ignored);
1699 } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
1700 numHex = 16;
1701 lead = "0d";
1702 APF.convert(ToSemantics: APFloat::IEEEdouble(), RM: APFloat::rmNearestTiesToEven, losesInfo: &ignored);
1703 } else
1704 llvm_unreachable("unsupported fp type");
1705
1706 APInt API = APF.bitcastToAPInt();
1707 O << lead << format_hex_no_prefix(N: API.getZExtValue(), Width: numHex, /*Upper=*/true);
1708}
1709
1710void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
1711 if (const ConstantInt *CI = dyn_cast<ConstantInt>(Val: CPV)) {
1712 O << CI->getValue();
1713 return;
1714 }
1715 if (const ConstantFP *CFP = dyn_cast<ConstantFP>(Val: CPV)) {
1716 printFPConstant(Fp: CFP, O);
1717 return;
1718 }
1719 if (isa<ConstantPointerNull>(Val: CPV)) {
1720 O << "0";
1721 return;
1722 }
1723 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(Val: CPV)) {
1724 const bool IsNonGenericPointer = GVar->getAddressSpace() != 0;
1725 if (EmitGeneric && !isa<Function>(Val: CPV) && !IsNonGenericPointer) {
1726 O << "generic(";
1727 getSymbol(GV: GVar)->print(OS&: O, MAI);
1728 O << ")";
1729 } else {
1730 getSymbol(GV: GVar)->print(OS&: O, MAI);
1731 }
1732 return;
1733 }
1734 if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(Val: CPV)) {
1735 const MCExpr *E = lowerConstantForGV(CV: cast<Constant>(Val: Cexpr), ProcessingGeneric: false);
1736 printMCExpr(Expr: *E, OS&: O);
1737 return;
1738 }
1739 llvm_unreachable("Not scalar type found in printScalarConstant()");
1740}
1741
1742void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes,
1743 AggBuffer *AggBuffer) {
1744 const DataLayout &DL = getDataLayout();
1745 int AllocSize = DL.getTypeAllocSize(Ty: CPV->getType());
1746 if (isa<UndefValue>(Val: CPV) || CPV->isNullValue()) {
1747 // Non-zero Bytes indicates that we need to zero-fill everything. Otherwise,
1748 // only the space allocated by CPV.
1749 AggBuffer->addZeros(Num: Bytes ? Bytes : AllocSize);
1750 return;
1751 }
1752
1753 // Helper for filling AggBuffer with APInts.
1754 auto AddIntToBuffer = [AggBuffer, Bytes](const APInt &Val) {
1755 size_t NumBytes = (Val.getBitWidth() + 7) / 8;
1756 SmallVector<unsigned char, 16> Buf(NumBytes);
1757 // `extractBitsAsZExtValue` does not allow the extraction of bits beyond the
1758 // input's bit width, and i1 arrays may not have a length that is a multuple
1759 // of 8. We handle the last byte separately, so we never request out of
1760 // bounds bits.
1761 for (unsigned I = 0; I < NumBytes - 1; ++I) {
1762 Buf[I] = Val.extractBitsAsZExtValue(numBits: 8, bitPosition: I * 8);
1763 }
1764 size_t LastBytePosition = (NumBytes - 1) * 8;
1765 size_t LastByteBits = Val.getBitWidth() - LastBytePosition;
1766 Buf[NumBytes - 1] =
1767 Val.extractBitsAsZExtValue(numBits: LastByteBits, bitPosition: LastBytePosition);
1768 AggBuffer->addBytes(Ptr: Buf.data(), Num: NumBytes, Bytes);
1769 };
1770
1771 switch (CPV->getType()->getTypeID()) {
1772 case Type::IntegerTyID:
1773 if (const auto *CI = dyn_cast<ConstantInt>(Val: CPV)) {
1774 AddIntToBuffer(CI->getValue());
1775 break;
1776 }
1777 if (const auto *Cexpr = dyn_cast<ConstantExpr>(Val: CPV)) {
1778 if (const auto *CI =
1779 dyn_cast<ConstantInt>(Val: ConstantFoldConstant(C: Cexpr, DL))) {
1780 AddIntToBuffer(CI->getValue());
1781 break;
1782 }
1783 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1784 Value *V = Cexpr->getOperand(i_nocapture: 0)->stripPointerCasts();
1785 AggBuffer->addSymbol(GVar: V, GVarBeforeStripping: Cexpr->getOperand(i_nocapture: 0));
1786 AggBuffer->addZeros(Num: AllocSize);
1787 break;
1788 }
1789 // A symbol-relative integer whose offset is applied outside the
1790 // ptrtoint, e.g. add(ptrtoint(@g), C). It can't fold to a ConstantInt
1791 // because it references a symbol; emit it through lowerConstantForGV, the
1792 // same path scalar symbol-relative integer globals use.
1793 AggBuffer->addSymbol(GVar: Cexpr, GVarBeforeStripping: Cexpr);
1794 AggBuffer->addZeros(Num: AllocSize);
1795 break;
1796 }
1797 llvm_unreachable("unsupported integer const type");
1798 break;
1799
1800 case Type::HalfTyID:
1801 case Type::BFloatTyID:
1802 case Type::FloatTyID:
1803 case Type::DoubleTyID:
1804 AddIntToBuffer(cast<ConstantFP>(Val: CPV)->getValueAPF().bitcastToAPInt());
1805 break;
1806
1807 case Type::PointerTyID: {
1808 if (const GlobalValue *GVar = dyn_cast<GlobalValue>(Val: CPV)) {
1809 AggBuffer->addSymbol(GVar, GVarBeforeStripping: GVar);
1810 } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(Val: CPV)) {
1811 const Value *v = Cexpr->stripPointerCasts();
1812 AggBuffer->addSymbol(GVar: v, GVarBeforeStripping: Cexpr);
1813 }
1814 AggBuffer->addZeros(Num: AllocSize);
1815 break;
1816 }
1817
1818 case Type::ArrayTyID:
1819 case Type::FixedVectorTyID:
1820 case Type::StructTyID: {
1821 if (isa<ConstantAggregate>(Val: CPV) || isa<ConstantDataSequential>(Val: CPV)) {
1822 // bufferAggregateConstant doesn't emit tail-padding, i.e. it writes
1823 // `store_size` bytes, not `alloc_size` bytes. Do it ourselves here.
1824 unsigned StartPos = AggBuffer->getCurpos();
1825 bufferAggregateConstant(CV: CPV, aggBuffer: AggBuffer);
1826 unsigned Written = AggBuffer->getCurpos() - StartPos;
1827 unsigned SlotSize = std::max<int>(a: Bytes, b: AllocSize);
1828 if (SlotSize > Written)
1829 AggBuffer->addZeros(Num: SlotSize - Written);
1830 } else if (isa<ConstantAggregateZero>(Val: CPV))
1831 AggBuffer->addZeros(Num: Bytes);
1832 else
1833 llvm_unreachable("Unexpected Constant type");
1834 break;
1835 }
1836
1837 default:
1838 llvm_unreachable("unsupported type");
1839 }
1840}
1841
1842void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV,
1843 AggBuffer *aggBuffer) {
1844 const DataLayout &DL = getDataLayout();
1845
1846 auto ExtendBuffer = [](APInt Val, AggBuffer *Buffer) {
1847 unsigned NumBytes = divideCeil(Numerator: Val.getBitWidth(), Denominator: 8);
1848 for (unsigned I : llvm::seq(Size: NumBytes)) {
1849 unsigned NumBits = std::min(a: 8u, b: Val.getBitWidth() - I * 8);
1850 Buffer->addByte(Byte: Val.extractBitsAsZExtValue(numBits: NumBits, bitPosition: I * 8));
1851 }
1852 };
1853
1854 // Integer or floating point vector splats.
1855 if (isa<ConstantInt, ConstantFP>(Val: CPV)) {
1856 if (auto *VTy = dyn_cast<FixedVectorType>(Val: CPV->getType())) {
1857 for (unsigned I : llvm::seq(Size: VTy->getNumElements()))
1858 bufferLEByte(CPV: CPV->getAggregateElement(Elt: I), Bytes: 0, AggBuffer: aggBuffer);
1859 return;
1860 }
1861 }
1862
1863 // Integers of arbitrary width
1864 if (const ConstantInt *CI = dyn_cast<ConstantInt>(Val: CPV)) {
1865 assert(CI->getType()->isIntegerTy() && "Expected integer constant!");
1866 ExtendBuffer(CI->getValue(), aggBuffer);
1867 return;
1868 }
1869
1870 // f128
1871 if (const ConstantFP *CFP = dyn_cast<ConstantFP>(Val: CPV)) {
1872 assert(CFP->getType()->isFloatingPointTy() && "Expected fp constant!");
1873 if (CFP->getType()->isFP128Ty()) {
1874 ExtendBuffer(CFP->getValueAPF().bitcastToAPInt(), aggBuffer);
1875 return;
1876 }
1877 }
1878
1879 // Buffer arrays one element at a time.
1880 if (isa<ConstantArray>(Val: CPV)) {
1881 for (const auto &Op : CPV->operands())
1882 bufferLEByte(CPV: cast<Constant>(Val: Op), Bytes: 0, AggBuffer: aggBuffer);
1883 return;
1884 }
1885
1886 // Constant vectors
1887 if (const auto *CVec = dyn_cast<ConstantVector>(Val: CPV)) {
1888 bufferAggregateConstVec(CV: CVec, aggBuffer);
1889 return;
1890 }
1891
1892 if (const auto *CDS = dyn_cast<ConstantDataSequential>(Val: CPV)) {
1893 for (unsigned I : llvm::seq(Size: CDS->getNumElements()))
1894 bufferLEByte(CPV: cast<Constant>(Val: CDS->getElementAsConstant(i: I)), Bytes: 0, AggBuffer: aggBuffer);
1895 return;
1896 }
1897
1898 if (isa<ConstantStruct>(Val: CPV)) {
1899 if (CPV->getNumOperands()) {
1900 StructType *ST = cast<StructType>(Val: CPV->getType());
1901 for (unsigned I : llvm::seq(Size: CPV->getNumOperands())) {
1902 int EndOffset = (I + 1 == CPV->getNumOperands())
1903 ? DL.getStructLayout(Ty: ST)->getElementOffset(Idx: 0) +
1904 DL.getTypeAllocSize(Ty: ST)
1905 : DL.getStructLayout(Ty: ST)->getElementOffset(Idx: I + 1);
1906 int Bytes = EndOffset - DL.getStructLayout(Ty: ST)->getElementOffset(Idx: I);
1907 bufferLEByte(CPV: cast<Constant>(Val: CPV->getOperand(i: I)), Bytes, AggBuffer: aggBuffer);
1908 }
1909 }
1910 return;
1911 }
1912 llvm_unreachable("unsupported constant type in printAggregateConstant()");
1913}
1914
1915void NVPTXAsmPrinter::bufferAggregateConstVec(const ConstantVector *CV,
1916 AggBuffer *aggBuffer) {
1917 unsigned NumElems = CV->getType()->getNumElements();
1918 const unsigned BuffSize = aggBuffer->getBufferSize();
1919
1920 // Buffer one element at a time if we have allocated enough buffer space.
1921 if (BuffSize >= NumElems) {
1922 for (const auto &Op : CV->operands())
1923 bufferLEByte(CPV: cast<Constant>(Val: Op), Bytes: 0, AggBuffer: aggBuffer);
1924 return;
1925 }
1926
1927 // Sub-byte datatypes will have more elements than bytes allocated for the
1928 // buffer. Merge consecutive elements to form a full byte. We expect that 8 %
1929 // sub-byte-elem-size should be 0 and current expected usage is for i4 (for
1930 // e2m1-fp4 types).
1931 Type *ElemTy = CV->getType()->getElementType();
1932 assert(ElemTy->isIntegerTy() && "Expected integer data type.");
1933 unsigned ElemTySize = ElemTy->getPrimitiveSizeInBits();
1934 assert(ElemTySize < 8 && "Expected sub-byte data type.");
1935 assert(8 % ElemTySize == 0 && "Element type size must evenly divide a byte.");
1936 // Number of elements to merge to form a full byte.
1937 unsigned NumElemsPerByte = 8 / ElemTySize;
1938 unsigned NumCompleteBytes = NumElems / NumElemsPerByte;
1939 unsigned NumTailElems = NumElems % NumElemsPerByte;
1940
1941 // Helper lambda to constant-fold sub-vector of sub-byte type elements into
1942 // i8. Start and end indices of the sub-vector is provided, along with number
1943 // of padding zeros if required.
1944 auto ConvertSubCVtoInt8 = [this, &ElemTy](const ConstantVector *CV,
1945 unsigned Start, unsigned End,
1946 unsigned NumPaddingZeros = 0) {
1947 // Collect elements to create sub-vector.
1948 SmallVector<Constant *, 8> SubCVElems;
1949 for (unsigned I : llvm::seq(Begin: Start, End))
1950 SubCVElems.push_back(Elt: CV->getAggregateElement(Elt: I));
1951
1952 // Optionally pad with zeros.
1953 if (NumPaddingZeros)
1954 SubCVElems.append(NumInputs: NumPaddingZeros, Elt: ConstantInt::getNullValue(Ty: ElemTy));
1955
1956 auto SubCV = ConstantVector::get(V: SubCVElems);
1957 Type *Int8Ty = IntegerType::get(C&: SubCV->getContext(), NumBits: 8);
1958
1959 // Merge elements of the sub-vector using ConstantFolding.
1960 ConstantInt *MergedElem =
1961 dyn_cast_or_null<ConstantInt>(Val: ConstantFoldConstant(
1962 C: ConstantExpr::getBitCast(C: const_cast<Constant *>(SubCV), Ty: Int8Ty),
1963 DL: getDataLayout()));
1964
1965 if (!MergedElem)
1966 report_fatal_error(
1967 reason: "Cannot lower vector global with unusual element type");
1968
1969 return MergedElem;
1970 };
1971
1972 // Iterate through elements of vector one chunk at a time and buffer that
1973 // chunk.
1974 for (unsigned ByteIdx : llvm::seq(Size: NumCompleteBytes))
1975 bufferLEByte(CPV: ConvertSubCVtoInt8(CV, ByteIdx * NumElemsPerByte,
1976 (ByteIdx + 1) * NumElemsPerByte),
1977 Bytes: 0, AggBuffer: aggBuffer);
1978
1979 // For unevenly sized vectors add tail padding zeros.
1980 if (NumTailElems > 0)
1981 bufferLEByte(CPV: ConvertSubCVtoInt8(CV, NumElems - NumTailElems, NumElems,
1982 NumElemsPerByte - NumTailElems),
1983 Bytes: 0, AggBuffer: aggBuffer);
1984}
1985
1986/// lowerConstantForGV - Return an MCExpr for the given Constant. This is mostly
1987/// a copy from AsmPrinter::lowerConstant, except customized to only handle
1988/// expressions that are representable in PTX and create
1989/// NVPTXGenericMCSymbolRefExpr nodes for addrspacecast instructions.
1990const MCExpr *
1991NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV,
1992 bool ProcessingGeneric) const {
1993 MCContext &Ctx = OutContext;
1994
1995 if (CV->isNullValue() || isa<UndefValue>(Val: CV))
1996 return MCConstantExpr::create(Value: 0, Ctx);
1997
1998 if (const ConstantInt *CI = dyn_cast<ConstantInt>(Val: CV))
1999 return MCConstantExpr::create(Value: CI->getZExtValue(), Ctx);
2000
2001 if (const GlobalValue *GV = dyn_cast<GlobalValue>(Val: CV)) {
2002 const MCSymbolRefExpr *Expr = MCSymbolRefExpr::create(Symbol: getSymbol(GV), Ctx);
2003 if (ProcessingGeneric)
2004 return NVPTXGenericMCSymbolRefExpr::create(SymExpr: Expr, Ctx);
2005 return Expr;
2006 }
2007
2008 const ConstantExpr *CE = dyn_cast<ConstantExpr>(Val: CV);
2009 if (!CE) {
2010 llvm_unreachable("Unknown constant value to lower!");
2011 }
2012
2013 switch (CE->getOpcode()) {
2014 default:
2015 break; // Error
2016
2017 case Instruction::AddrSpaceCast: {
2018 // Strip the addrspacecast and pass along the operand
2019 PointerType *DstTy = cast<PointerType>(Val: CE->getType());
2020 if (DstTy->getAddressSpace() == 0)
2021 return lowerConstantForGV(CV: cast<const Constant>(Val: CE->getOperand(i_nocapture: 0)), ProcessingGeneric: true);
2022
2023 break; // Error
2024 }
2025
2026 case Instruction::GetElementPtr: {
2027 const DataLayout &DL = getDataLayout();
2028
2029 // Generate a symbolic expression for the byte address
2030 APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0);
2031 cast<GEPOperator>(Val: CE)->accumulateConstantOffset(DL, Offset&: OffsetAI);
2032
2033 const MCExpr *Base = lowerConstantForGV(CV: CE->getOperand(i_nocapture: 0),
2034 ProcessingGeneric);
2035 if (!OffsetAI)
2036 return Base;
2037
2038 int64_t Offset = OffsetAI.getSExtValue();
2039 return MCBinaryExpr::createAdd(LHS: Base, RHS: MCConstantExpr::create(Value: Offset, Ctx),
2040 Ctx);
2041 }
2042
2043 case Instruction::Trunc:
2044 // We emit the value and depend on the assembler to truncate the generated
2045 // expression properly. This is important for differences between
2046 // blockaddress labels. Since the two labels are in the same function, it
2047 // is reasonable to treat their delta as a 32-bit value.
2048 [[fallthrough]];
2049 case Instruction::BitCast:
2050 return lowerConstantForGV(CV: CE->getOperand(i_nocapture: 0), ProcessingGeneric);
2051
2052 case Instruction::IntToPtr: {
2053 const DataLayout &DL = getDataLayout();
2054
2055 // Handle casts to pointers by changing them into casts to the appropriate
2056 // integer type. This promotes constant folding and simplifies this code.
2057 Constant *Op = CE->getOperand(i_nocapture: 0);
2058 Op = ConstantFoldIntegerCast(C: Op, DestTy: DL.getIntPtrType(CV->getType()),
2059 /*IsSigned*/ false, DL);
2060 if (Op)
2061 return lowerConstantForGV(CV: Op, ProcessingGeneric);
2062
2063 break; // Error
2064 }
2065
2066 case Instruction::PtrToInt: {
2067 const DataLayout &DL = getDataLayout();
2068
2069 // Support only foldable casts to/from pointers that can be eliminated by
2070 // changing the pointer to the appropriately sized integer type.
2071 Constant *Op = CE->getOperand(i_nocapture: 0);
2072 Type *Ty = CE->getType();
2073
2074 const MCExpr *OpExpr = lowerConstantForGV(CV: Op, ProcessingGeneric);
2075
2076 // We can emit the pointer value into this slot if the slot is an
2077 // integer slot equal to the size of the pointer.
2078 if (DL.getTypeAllocSize(Ty) == DL.getTypeAllocSize(Ty: Op->getType()))
2079 return OpExpr;
2080
2081 // Otherwise the pointer is smaller than the resultant integer, mask off
2082 // the high bits so we are sure to get a proper truncation if the input is
2083 // a constant expr.
2084 unsigned InBits = DL.getTypeAllocSizeInBits(Ty: Op->getType());
2085 const MCExpr *MaskExpr = MCConstantExpr::create(Value: ~0ULL >> (64-InBits), Ctx);
2086 return MCBinaryExpr::createAnd(LHS: OpExpr, RHS: MaskExpr, Ctx);
2087 }
2088
2089 // The MC library also has a right-shift operator, but it isn't consistently
2090 // signed or unsigned between different targets.
2091 case Instruction::Add: {
2092 const MCExpr *LHS = lowerConstantForGV(CV: CE->getOperand(i_nocapture: 0), ProcessingGeneric);
2093 const MCExpr *RHS = lowerConstantForGV(CV: CE->getOperand(i_nocapture: 1), ProcessingGeneric);
2094 switch (CE->getOpcode()) {
2095 default: llvm_unreachable("Unknown binary operator constant cast expr");
2096 case Instruction::Add: return MCBinaryExpr::createAdd(LHS, RHS, Ctx);
2097 }
2098 }
2099 }
2100
2101 // If the code isn't optimized, there may be outstanding folding
2102 // opportunities. Attempt to fold the expression using DataLayout as a
2103 // last resort before giving up.
2104 Constant *C = ConstantFoldConstant(C: CE, DL: getDataLayout());
2105 if (C != CE)
2106 return lowerConstantForGV(CV: C, ProcessingGeneric);
2107
2108 // Otherwise report the problem to the user.
2109 std::string S;
2110 raw_string_ostream OS(S);
2111 OS << "Unsupported expression in static initializer: ";
2112 CE->printAsOperand(O&: OS, /*PrintType=*/false,
2113 M: !MF ? nullptr : MF->getFunction().getParent());
2114 report_fatal_error(reason: Twine(OS.str()));
2115}
2116
2117void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) const {
2118 OutContext.getAsmInfo().printExpr(OS, Expr);
2119}
2120
2121/// PrintAsmOperand - Print out an operand for an inline asm expression.
2122///
2123bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
2124 const char *ExtraCode, raw_ostream &O) {
2125 if (ExtraCode && ExtraCode[0]) {
2126 if (ExtraCode[1] != 0)
2127 return true; // Unknown modifier.
2128
2129 switch (ExtraCode[0]) {
2130 default:
2131 // See if this is a generic print operand
2132 return AsmPrinter::PrintAsmOperand(MI, OpNo, ExtraCode, OS&: O);
2133 case 'r':
2134 break;
2135 }
2136 }
2137
2138 printOperand(MI, OpNum: OpNo, O);
2139
2140 return false;
2141}
2142
2143bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI,
2144 unsigned OpNo,
2145 const char *ExtraCode,
2146 raw_ostream &O) {
2147 if (ExtraCode && ExtraCode[0])
2148 return true; // Unknown modifier
2149
2150 O << '[';
2151 printMemOperand(MI, OpNum: OpNo, O);
2152 O << ']';
2153
2154 return false;
2155}
2156
2157void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, unsigned OpNum,
2158 raw_ostream &O) {
2159 const MachineOperand &MO = MI->getOperand(i: OpNum);
2160 switch (MO.getType()) {
2161 case MachineOperand::MO_Register:
2162 if (MO.getReg().isPhysical()) {
2163 if (MO.getReg() == NVPTX::VRDepot)
2164 O << DEPOTNAME << getFunctionNumber();
2165 else
2166 O << NVPTXInstPrinter::getRegisterName(Reg: MO.getReg());
2167 } else {
2168 emitVirtualRegister(vr: MO.getReg(), O);
2169 }
2170 break;
2171
2172 case MachineOperand::MO_Immediate:
2173 O << MO.getImm();
2174 break;
2175
2176 case MachineOperand::MO_FPImmediate:
2177 printFPConstant(Fp: MO.getFPImm(), O);
2178 break;
2179
2180 case MachineOperand::MO_GlobalAddress:
2181 PrintSymbolOperand(MO, OS&: O);
2182 break;
2183
2184 case MachineOperand::MO_MachineBasicBlock:
2185 MO.getMBB()->getSymbol()->print(OS&: O, MAI);
2186 break;
2187
2188 default:
2189 llvm_unreachable("Operand type not supported.");
2190 }
2191}
2192
2193void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, unsigned OpNum,
2194 raw_ostream &O, const char *Modifier) {
2195 printOperand(MI, OpNum, O);
2196
2197 if (Modifier && strcmp(s1: Modifier, s2: "add") == 0) {
2198 O << ", ";
2199 printOperand(MI, OpNum: OpNum + 1, O);
2200 } else {
2201 if (MI->getOperand(i: OpNum + 1).isImm() &&
2202 MI->getOperand(i: OpNum + 1).getImm() == 0)
2203 return; // don't print ',0' or '+0'
2204 O << "+";
2205 printOperand(MI, OpNum: OpNum + 1, O);
2206 }
2207}
2208
2209/// Returns true if \p Line begins with an alphabetic character or underscore,
2210/// indicating it is a PTX instruction that should receive a .loc directive.
2211static bool isPTXInstruction(StringRef Line) {
2212 StringRef Trimmed = Line.ltrim();
2213 return !Trimmed.empty() &&
2214 (std::isalpha(static_cast<unsigned char>(Trimmed[0])) ||
2215 Trimmed[0] == '_');
2216}
2217
2218/// Returns the DILocation for an inline asm MachineInstr if debug line info
2219/// should be emitted, or nullptr otherwise.
2220static const DILocation *getInlineAsmDebugLoc(const MachineInstr *MI) {
2221 if (!MI || !MI->getDebugLoc())
2222 return nullptr;
2223 const DISubprogram *SP = MI->getMF()->getFunction().getSubprogram();
2224 if (!SP || SP->getUnit()->getEmissionKind() == DICompileUnit::NoDebug)
2225 return nullptr;
2226 const DILocation *DL = MI->getDebugLoc();
2227 if (!DL->getFile() || !DL->getLine())
2228 return nullptr;
2229 return DL;
2230}
2231
2232namespace {
2233struct InlineAsmInliningContext {
2234 MCSymbol *FuncNameSym = nullptr;
2235 unsigned FileIA = 0;
2236 unsigned LineIA = 0;
2237 unsigned ColIA = 0;
2238
2239 bool hasInlinedAt() const { return FuncNameSym != nullptr; }
2240};
2241} // namespace
2242
2243/// Resolves the enhanced-lineinfo inlining context for an inline asm debug
2244/// location. Returns a default (empty) context if inlining info is unavailable.
2245static InlineAsmInliningContext
2246getInlineAsmInliningContext(const DILocation *DL, const MachineFunction &MF,
2247 NVPTXDwarfDebug *NVDD, MCStreamer &Streamer,
2248 unsigned CUID) {
2249 InlineAsmInliningContext Ctx;
2250 const DILocation *InlinedAt = DL->getInlinedAt();
2251 if (!InlinedAt || !InlinedAt->getFile() || !NVDD ||
2252 !NVDD->isEnhancedLineinfo(MF))
2253 return Ctx;
2254 const auto *SubProg = getDISubprogram(Scope: DL->getScope());
2255 if (!SubProg)
2256 return Ctx;
2257 Ctx.FuncNameSym = NVDD->getOrCreateFuncNameSymbol(LinkageName: SubProg->getLinkageName());
2258 Ctx.FileIA = Streamer.emitDwarfFileDirective(
2259 FileNo: 0, Directory: InlinedAt->getFile()->getDirectory(),
2260 Filename: InlinedAt->getFile()->getFilename(), Checksum: std::nullopt, Source: std::nullopt, CUID);
2261 Ctx.LineIA = InlinedAt->getLine();
2262 Ctx.ColIA = InlinedAt->getColumn();
2263 return Ctx;
2264}
2265
2266void NVPTXAsmPrinter::emitInlineAsm(StringRef Str, const MCSubtargetInfo &STI,
2267 const MCTargetOptions &MCOptions,
2268 const MDNode *LocMDNode,
2269 InlineAsm::AsmDialect Dialect,
2270 const MachineInstr *MI) {
2271 assert(!Str.empty() && "Can't emit empty inline asm block");
2272 if (Str.back() == 0)
2273 Str = Str.substr(Start: 0, N: Str.size() - 1);
2274
2275 auto emitAsmStr = [&](StringRef AsmStr) {
2276 emitInlineAsmStart();
2277 OutStreamer->emitRawText(String: AsmStr);
2278 emitInlineAsmEnd(StartInfo: STI, EndInfo: nullptr, MI);
2279 };
2280
2281 const DILocation *DL = getInlineAsmDebugLoc(MI);
2282 if (!DL) {
2283 emitAsmStr(Str);
2284 return;
2285 }
2286
2287 const DIFile *File = DL->getFile();
2288 unsigned Line = DL->getLine();
2289 const unsigned Column = DL->getColumn();
2290 const unsigned CUID = OutStreamer->getContext().getDwarfCompileUnitID();
2291 const unsigned FileNumber = OutStreamer->emitDwarfFileDirective(
2292 FileNo: 0, Directory: File->getDirectory(), Filename: File->getFilename(), Checksum: std::nullopt, Source: std::nullopt,
2293 CUID);
2294
2295 auto *NVDD = static_cast<NVPTXDwarfDebug *>(getDwarfDebug());
2296 InlineAsmInliningContext InlineCtx =
2297 getInlineAsmInliningContext(DL, MF: *MI->getMF(), NVDD, Streamer&: *OutStreamer, CUID);
2298
2299 SmallVector<StringRef, 16> Lines;
2300 Str.split(A&: Lines, Separator: '\n');
2301 emitInlineAsmStart();
2302 for (const StringRef &L : Lines) {
2303 StringRef RTrimmed = L.rtrim(Char: '\r');
2304 if (isPTXInstruction(Line: L)) {
2305 if (InlineCtx.hasInlinedAt()) {
2306 OutStreamer->emitDwarfLocDirectiveWithInlinedAt(
2307 FileNo: FileNumber, Line, Column, FileIA: InlineCtx.FileIA, LineIA: InlineCtx.LineIA,
2308 ColumnIA: InlineCtx.ColIA, Sym: InlineCtx.FuncNameSym, DWARF2_FLAG_IS_STMT, Isa: 0, Discriminator: 0,
2309 FileName: File->getFilename());
2310 } else {
2311 OutStreamer->emitDwarfLocDirective(FileNo: FileNumber, Line, Column,
2312 DWARF2_FLAG_IS_STMT, Isa: 0, Discriminator: 0,
2313 FileName: File->getFilename());
2314 }
2315 }
2316 OutStreamer->emitRawText(String: RTrimmed);
2317 ++Line;
2318 }
2319 emitInlineAsmEnd(StartInfo: STI, EndInfo: nullptr, MI);
2320}
2321
2322char NVPTXAsmPrinter::ID = 0;
2323
2324INITIALIZE_PASS(NVPTXAsmPrinter, "nvptx-asm-printer", "NVPTX Assembly Printer",
2325 false, false)
2326
2327// Force static initialization.
2328extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void
2329LLVMInitializeNVPTXAsmPrinter() {
2330 RegisterAsmPrinter<NVPTXAsmPrinter> X(getTheNVPTXTarget32());
2331 RegisterAsmPrinter<NVPTXAsmPrinter> Y(getTheNVPTXTarget64());
2332}
2333