NVPTXAsmPrinter.cpp Source File (original) (raw)

88#include

89#include

90#include

91#include

93using namespace llvm;

95#define DEPOTNAME "__local_depot"

99static void

103 Globals.insert(GV);

104 return;

105 }

106

108 for (const auto &O : U->operands())

110}

111

112

113

114

115static void

120

121 if (Visited.count(GV))

122 return;

123

124

125 if (!Visiting.insert(GV).second)

126 report_fatal_error("Circular dependency found in global variable set");

127

128

130 for (const auto &O : GV->operands())

132

135

136

139 Visiting.erase(GV);

140}

141

142void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {

143 NVPTX_MC::verifyInstructionPredicates(MI->getOpcode(),

145

146 MCInst Inst;

147 lowerToMCInst(MI, Inst);

149}

150

153

154 if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {

155 const MachineOperand &MO = MI->getOperand(0);

158 return;

159 }

160

161 for (const auto MO : MI->operands())

163}

164

167 default:

183

185 default:

187 break;

200 }

201 break;

202 }

203 }

204}

205

206unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {

208 const TargetRegisterClass *RC = MRI->getRegClass(Reg);

209

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

211 unsigned RegNum = RegMap[Reg];

212

213

214

215 unsigned Ret = 0;

216 if (RC == &NVPTX::B1RegClass) {

217 Ret = (1 << 28);

218 } else if (RC == &NVPTX::B16RegClass) {

219 Ret = (2 << 28);

220 } else if (RC == &NVPTX::B32RegClass) {

221 Ret = (3 << 28);

222 } else if (RC == &NVPTX::B64RegClass) {

223 Ret = (4 << 28);

224 } else if (RC == &NVPTX::B128RegClass) {

225 Ret = (7 << 28);

226 } else {

228 }

229

230

231 Ret |= (RegNum & 0x0FFFFFFF);

232 return Ret;

233 } else {

234

235

236 return Reg & 0x0FFFFFFF;

237 }

238}

239

241 const MCExpr *Expr;

244}

245

248 const NVPTXSubtarget &STI = TM.getSubtarget(*F);

250

251 Type *Ty = F->getReturnType();

253 return;

254 O << " (";

255

256 auto PrintScalarRetVal = [&](unsigned Size) {

258 };

260 const unsigned TotalSize = DL.getTypeAllocSize(Ty);

261 const Align RetAlignment = TLI->getFunctionArgumentAlignment(

262 F, Ty, AttributeList::ReturnIndex, DL);

263 O << ".param .align " << RetAlignment.value() << " .b8 func_retval0["

264 << TotalSize << "]";

268 PrintScalarRetVal(ITy->getBitWidth());

270 PrintScalarRetVal(TLI->getPointerTy(DL).getSizeInBits());

271 } else

273 O << ") ";

274}

275

276void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,

278 const Function &F = MF.getFunction();

279 printReturnValStr(&F, O);

280}

281

282

283

284bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(

287

289 return false;

290

291

292

293

294 for (const MachineBasicBlock *PMBB : MBB.predecessors()) {

296

297 continue;

298 }

299 if (const BasicBlock *PBB = PMBB->getBasicBlock()) {

300 if (MDNode *LoopID =

301 PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {

303 return true;

304 if (MDNode *UnrollCountMD =

307 ->isOne())

308 return true;

309 }

310 }

311 }

312 }

313 return false;

314}

315

318 if (isLoopHeaderOfNoUnroll(MBB))

319 OutStreamer->emitRawText(StringRef("\t.pragma \"nounroll\";\n"));

320}

321

323 SmallString<128> Str;

324 raw_svector_ostream O(Str);

325

326 if (!GlobalsEmitted) {

327 emitGlobals(*MF->getFunction().getParent());

328 GlobalsEmitted = true;

329 }

330

331

332 MRI = &MF->getRegInfo();

333 F = &MF->getFunction();

334 emitLinkageDirective(F, O);

336 O << ".entry ";

337 else {

338 O << ".func ";

339 printReturnValStr(*MF, O);

340 }

341

343

344 emitFunctionParamList(F, O);

345 O << "\n";

346

348 emitKernelFunctionDirectives(*F, O);

349

351 O << ".noreturn";

352

354

355 VRegMapping.clear();

356

357 OutStreamer->emitRawText(StringRef("{\n"));

358 setAndEmitFunctionVirtualRegisters(*MF);

359 encodeDebugInfoRegisterNumbers(*MF);

360

361 if (const DISubprogram *SP = MF->getFunction().getSubprogram()) {

363 if (->getUnit()->isDebugDirectivesOnly())

365 }

366}

367

370

371

372

373

374

376 return Result;

377}

378

384}

385

387 VRegMapping.clear();

388}

389

393 return OutContext.getOrCreateSymbol(Str);

394}

395

396void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const {

397 Register RegNo = MI->getOperand(0).getReg();

401 } else {

405 }

407}

408

409void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F,

411

412

413

415 if (!ReqNTID.empty())

416 O << formatv(".reqntid {0:$[, ]}\n",

417 make_range(ReqNTID.begin(), ReqNTID.end()));

418

420 if (!MaxNTID.empty())

421 O << formatv(".maxntid {0:$[, ]}\n",

422 make_range(MaxNTID.begin(), MaxNTID.end()));

423

425 O << ".minnctapersm " << *Mincta << "\n";

426

427 if (const auto Maxnreg = getMaxNReg(F))

428 O << ".maxnreg " << *Maxnreg << "\n";

429

430

431

432 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);

433 const NVPTXSubtarget *STI = &NTM.getSubtarget(F);

434

438

439 if (!ClusterDim.empty()) {

440

441 if (!BlocksAreClusters)

442 O << ".explicitcluster\n";

443

444 if (ClusterDim[0] != 0) {

446 "cluster_dim_x != 0 implies cluster_dim_y and cluster_dim_z "

447 "should be non-zero as well");

448

449 O << formatv(".reqnctapercluster {0:$[, ]}\n",

450 make_range(ClusterDim.begin(), ClusterDim.end()));

451 } else {

453 "cluster_dim_x == 0 implies cluster_dim_y and cluster_dim_z "

454 "should be 0 as well");

455 }

456 }

457

458 if (BlocksAreClusters) {

459 LLVMContext &Ctx = F.getContext();

460 if (ReqNTID.empty() || ClusterDim.empty())

461 Ctx.diagnose(DiagnosticInfoUnsupported(

462 F, "blocksareclusters requires reqntid and cluster_dim attributes",

463 F.getSubprogram()));

465 Ctx.diagnose(DiagnosticInfoUnsupported(

466 F, "blocksareclusters requires PTX version >= 9.0",

467 F.getSubprogram()));

468 else

469 O << ".blocksareclusters\n";

470 }

471

473 O << ".maxclusterrank " << *Maxclusterrank << "\n";

474 }

475}

476

479

480 std::string Name;

482

484 assert(I != VRegMapping.end() && "Bad register class");

486

488 assert(VI != RegMap.end() && "Bad virtual register");

489 unsigned MappedVR = VI->second;

490

492

493 return Name;

494}

495

496void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr,

499}

500

501void NVPTXAsmPrinter::emitAliasDeclaration(const GlobalAlias *GA,

506 "NVPTX aliasee must be a non-kernel function definition");

507

511

512 emitDeclarationWithName(F, getSymbol(GA), O);

513}

514

516 emitDeclarationWithName(F, getSymbol(F), O);

517}

518

519void NVPTXAsmPrinter::emitDeclarationWithName(const Function *F, MCSymbol *S,

521 emitLinkageDirective(F, O);

523 O << ".entry ";

524 else

525 O << ".func ";

526 printReturnValStr(F, O);

528 O << "\n";

529 emitFunctionParamList(F, O);

530 O << "\n";

532 O << ".noreturn";

533 O << ";\n";

534}

535

537 if ()

538 return false;

539

541 return GV->getName() != "llvm.used";

542

543 for (const User *U : C->users())

546 return true;

547

548 return false;

549}

550

553 if (OtherGV->getName() == "llvm.used")

554 return true;

555

557 if (const Function *CurFunc = I->getFunction()) {

558 if (OneFunc && (CurFunc != OneFunc))

559 return false;

560 OneFunc = CurFunc;

561 return true;

562 }

563 return false;

564 }

565

566 for (const User *UU : U->users())

568 return false;

569

570 return true;

571}

572

573

574

575

576

577

578

579

582 return false;

584 return false;

585

586 const Function *oneFunc = nullptr;

587

589 if (!flag)

590 return false;

591 if (!oneFunc)

592 return false;

593 f = oneFunc;

594 return true;

595}

596

599 for (const User *U : C->users()) {

602 return true;

604 if (const Function *Caller = I->getFunction())

605 if (SeenSet.contains(Caller))

606 return true;

607 }

608 }

609 return false;

610}

611

612void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {

613 SmallPtrSet<const Function *, 32> SeenSet;

614 for (const Function &F : M) {

615 if (F.getAttributes().hasFnAttr("nvptx-libcall-callee")) {

616 emitDeclaration(&F, O);

617 continue;

618 }

619

620 if (F.isDeclaration()) {

621 if (F.use_empty())

622 continue;

623 if (F.getIntrinsicID())

624 continue;

625 emitDeclaration(&F, O);

626 continue;

627 }

628 for (const User *U : F.users()) {

631

632

633

634 emitDeclaration(&F, O);

635 break;

636 }

637

638

640 emitDeclaration(&F, O);

641 break;

642 }

643 }

644

646 continue;

648 if (!Caller)

649 continue;

650

651

652

653

654 if (SeenSet.contains(Caller)) {

655 emitDeclaration(&F, O);

656 break;

657 }

658 }

660 }

661 for (const GlobalAlias &GA : M.aliases())

662 emitAliasDeclaration(&GA, O);

663}

664

665void NVPTXAsmPrinter::emitStartOfAsmFile(Module &M) {

666

667

668

669 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);

671 SmallString<128> Str1;

672 raw_svector_ostream OS1(Str1);

673

674

675 emitHeader(M, OS1, *STI);

677}

678

684

685

687

688 GlobalsEmitted = false;

689

690 return Result;

691}

692

693void NVPTXAsmPrinter::emitGlobals(const Module &M) {

696

697 emitDeclarations(M, OS2);

698

699

700

701

702

703

707

708

711

712 assert(GVVisited.size() == M.global_size() && "Missed a global variable");

713 assert(GVVisiting.size() == 0 && "Did not fully process a global variable");

714

717

718

720 printModuleLevelGV(GV, OS2, false, STI);

721

722 OS2 << '\n';

723

725}

726

727void NVPTXAsmPrinter::emitGlobalAlias(const Module &M, const GlobalAlias &GA) {

730

732

734 << ";\n";

735

737}

738

742

743 O << "//\n"

744 "// Generated by LLVM NVPTX Back-End\n"

745 "//\n"

746 "\n"

747 << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n"

749

750 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);

752 O << ", texmode_independent";

753

754 bool HasFullDebugInfo = false;

755 for (DICompileUnit *CU : M.debug_compile_units()) {

756 switch(CU->getEmissionKind()) {

759 break;

762 HasFullDebugInfo = true;

763 break;

764 }

765 if (HasFullDebugInfo)

766 break;

767 }

768 if (HasFullDebugInfo)

769 O << ", debug";

770

771 O << "\n"

772 << ".address_size " << (NTM.is64Bit() ? "64" : "32") << "\n"

773 << "\n";

774}

775

777

778

779 if (!GlobalsEmitted) {

780 emitGlobals(M);

781 GlobalsEmitted = true;

782 }

783

784

786

788

789 auto *TS =

791

793 TS->closeLastSection();

794

795 OutStreamer->emitRawText("\t.section\t.debug_macinfo\t{\t}");

796 }

797

798

800

801 return ret;

802}

803

804

805

806

807

808

809

810

811

812

813

814

815

816

817void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V,

820 if (V->hasExternalLinkage()) {

822 O << (GVar->hasInitializer() ? ".visible " : ".extern ");

823 else if (V->isDeclaration())

824 O << ".extern ";

825 else

826 O << ".visible ";

827 } else if (V->hasAppendingLinkage()) {

829 "' has unsupported appending linkage type");

830 } else if (->hasInternalLinkage() && ->hasPrivateLinkage()) {

831 O << ".weak ";

832 }

833 }

834}

835

836void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,

839

841 if (GVar->getSection() == "llvm.metadata")

842 return;

843

844

847 return;

848

850

851

853

856 O << ".visible ";

857 else

858 O << ".extern ";

861 O << ".common ";

865 O << ".weak ";

866 }

867

869 O << ".global .texref " << getTextureName(*GVar) << ";\n";

870 return;

871 }

872

874 O << ".global .surfref " << getSurfaceName(*GVar) << ";\n";

875 return;

876 }

877

879

880

881

882 emitPTXGlobalVariable(GVar, O, STI);

883 O << ";\n";

884 return;

885 }

886

889

890 const Constant *Initializer = nullptr;

893 const ConstantInt *CI = nullptr;

894 if (Initializer)

896 if (CI) {

898

899 O << " = { ";

900

901 for (int i = 0,

903 i < 3; i++) {

904 O << "addr_mode_" << i << " = ";

905 switch (addr) {

906 case 0:

907 O << "wrap";

908 break;

909 case 1:

910 O << "clamp_to_border";

911 break;

912 case 2:

913 O << "clamp_to_edge";

914 break;

915 case 3:

916 O << "wrap";

917 break;

918 case 4:

919 O << "mirror";

920 break;

921 }

922 O << ", ";

923 }

924 O << "filter_mode = ";

926 case 0:

927 O << "nearest";

928 break;

929 case 1:

930 O << "linear";

931 break;

932 case 2:

934 default:

935 O << "nearest";

936 break;

937 }

939 O << ", force_unnormalized_coords = 1";

940 }

941 O << " }";

942 }

943

944 O << ";\n";

945 return;

946 }

947

950 return;

951

952

954 return;

956 return;

957 }

958

959 const Function *DemotedFunc = nullptr;

961 O << "// " << GVar->getName() << " has been demoted\n";

962 localDecls[DemotedFunc].push_back(GVar);

963 return;

964 }

965

966 O << ".";

968

972 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");

973 O << " .attribute(.managed)";

974 }

975

976 O << " .align "

977 << GVar->getAlign().value_or(DL.getPrefTypeAlign(ETy)).value();

978

981 O << " .";

982

984 O << "u8";

985 else

986 O << getPTXFundamentalTypeStr(ETy, false);

987 O << " ";

989

990

991

996

998 O << " = ";

999 printScalarConstant(Initializer, O);

1000 }

1001 } else {

1002

1003

1004

1008 "' is not allowed in addrspace(" +

1010 }

1011 }

1012 }

1013 } else {

1014

1015

1016

1017

1024 const uint64_t ElementSize = DL.getTypeStoreSize(ETy);

1025

1026

1032 AggBuffer aggBuffer(ElementSize, *this);

1033 bufferAggregateConstant(Initializer, &aggBuffer);

1034 if (aggBuffer.numSymbols()) {

1035 const unsigned int ptrSize = MAI->getCodePointerSize();

1036 if (ElementSize % ptrSize ||

1037 !aggBuffer.allSymbolsAligned(ptrSize)) {

1038

1041 "initialized packed aggregate with pointers '" +

1043 "' requires at least PTX ISA version 7.1");

1044 O << " .u8 ";

1046 O << "[" << ElementSize << "] = {";

1047 aggBuffer.printBytes(O);

1048 O << "}";

1049 } else {

1050 O << " .u" << ptrSize * 8 << " ";

1052 O << "[" << ElementSize / ptrSize << "] = {";

1053 aggBuffer.printWords(O);

1054 O << "}";

1055 }

1056 } else {

1057 O << " .b8 ";

1059 O << "[" << ElementSize << "] = {";

1060 aggBuffer.printBytes(O);

1061 O << "}";

1062 }

1063 } else {

1064 O << " .b8 ";

1066 if (ElementSize)

1067 O << "[" << ElementSize << "]";

1068 }

1069 } else {

1070 O << " .b8 ";

1072 if (ElementSize)

1073 O << "[" << ElementSize << "]";

1074 }

1075 break;

1076 }

1077 default:

1079 }

1080 }

1081 O << ";\n";

1082}

1083

1084void NVPTXAsmPrinter::AggBuffer::printSymbol(unsigned nSym, raw_ostream &os) {

1085 const Value *v = Symbols[nSym];

1086 const Value *v0 = SymbolsBeforeStripping[nSym];

1090

1091 bool isGenericPointer = PTy && PTy->getAddressSpace() == 0;

1092 if (EmitGeneric && isGenericPointer && isa<Function> (v)) {

1093 os << "generic(";

1094 Name->print(os, AP.MAI);

1095 os << ")";

1096 } else {

1097 Name->print(os, AP.MAI);

1098 }

1100 const MCExpr *Expr = AP.lowerConstantForGV(CExpr, false);

1101 AP.printMCExpr(*Expr, os);

1102 } else

1104}

1105

1106void NVPTXAsmPrinter::AggBuffer::printBytes(raw_ostream &os) {

1107 unsigned int ptrSize = AP.MAI->getCodePointerSize();

1108

1109

1110

1111

1112 unsigned int InitializerCount = size;

1113

1114

1115 if (numSymbols() == 0)

1116 while (InitializerCount >= 1 && !buffer[InitializerCount - 1])

1117 InitializerCount--;

1118

1119 symbolPosInBuffer.push_back(InitializerCount);

1120 unsigned int nSym = 0;

1121 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];

1122 for (unsigned int pos = 0; pos < InitializerCount;) {

1123 if (pos)

1124 os << ", ";

1125 if (pos != nextSymbolPos) {

1126 os << (unsigned int)buffer[pos];

1127 ++pos;

1128 continue;

1129 }

1130

1131

1132

1133 std::string symText;

1134 llvm::raw_string_ostream oss(symText);

1135 printSymbol(nSym, oss);

1136 for (unsigned i = 0; i < ptrSize; ++i) {

1137 if (i)

1138 os << ", ";

1140 os << "(" << symText << ")";

1141 }

1142 pos += ptrSize;

1143 nextSymbolPos = symbolPosInBuffer[++nSym];

1144 assert(nextSymbolPos >= pos);

1145 }

1146}

1147

1148void NVPTXAsmPrinter::AggBuffer::printWords(raw_ostream &os) {

1149 unsigned int ptrSize = AP.MAI->getCodePointerSize();

1150 symbolPosInBuffer.push_back(size);

1151 unsigned int nSym = 0;

1152 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];

1153 assert(nextSymbolPos % ptrSize == 0);

1154 for (unsigned int pos = 0; pos < size; pos += ptrSize) {

1155 if (pos)

1156 os << ", ";

1157 if (pos == nextSymbolPos) {

1158 printSymbol(nSym, os);

1159 nextSymbolPos = symbolPosInBuffer[++nSym];

1160 assert(nextSymbolPos % ptrSize == 0);

1161 assert(nextSymbolPos >= pos + ptrSize);

1162 } else if (ptrSize == 4)

1164 else

1166 }

1167}

1168

1169void NVPTXAsmPrinter::emitDemotedVars(const Function *F, raw_ostream &O) {

1170 auto It = localDecls.find(F);

1171 if (It == localDecls.end())

1172 return;

1173

1175

1176 const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);

1178

1179 for (const GlobalVariable *GV : GVars) {

1180 O << "\t// demoted variable\n\t";

1181 printModuleLevelGV(GV, O, true, STI);

1182 }

1183}

1184

1185void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,

1186 raw_ostream &O) const {

1189 O << "local";

1190 break;

1192 O << "global";

1193 break;

1195 O << "const";

1196 break;

1198 O << "shared";

1199 break;

1200 default:

1203 break;

1204 }

1205}

1206

1207std::string

1208NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {

1212 if (NumBits == 1)

1213 return "pred";

1214 if (NumBits <= 64) {

1215 std::string name = "u";

1217 }

1219 break;

1220 }

1223

1224

1225 return "b16";

1227 return "f32";

1229 return "f64";

1232 assert((PtrSize == 64 || PtrSize == 32) && "Unexpected pointer size");

1233

1234 if (PtrSize == 64)

1235 if (useB4PTR)

1236 return "b64";

1237 else

1238 return "u64";

1239 else if (useB4PTR)

1240 return "b32";

1241 else

1242 return "u32";

1243 }

1244 default:

1245 break;

1246 }

1248}

1249

1250void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,

1251 raw_ostream &O,

1252 const NVPTXSubtarget &STI) {

1254

1255

1257

1258 O << ".";

1263 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");

1264

1265 O << " .attribute(.managed)";

1266 }

1267 O << " .align "

1268 << GVar->getAlign().value_or(DL.getPrefTypeAlign(ETy)).value();

1269

1270

1272 O << " .b8 ";

1274 O << "[16]";

1275 return;

1276 }

1277

1279 O << " ." << getPTXFundamentalTypeStr(ETy) << " ";

1281 return;

1282 }

1283

1284 int64_t ElementSize = 0;

1285

1286

1287

1288

1289

1294 ElementSize = DL.getTypeStoreSize(ETy);

1295 O << " .b8 ";

1297 O << "[";

1298 if (ElementSize) {

1299 O << ElementSize;

1300 }

1301 O << "]";

1302 break;

1303 default:

1305 }

1306}

1307

1308void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {

1310 const NVPTXSubtarget &STI = TM.getSubtarget(*F);

1312 const NVPTXMachineFunctionInfo *MFI =

1313 MF ? MF->getInfo() : nullptr;

1314

1315 bool IsFirst = true;

1317

1318 if (F->arg_empty() && !F->isVarArg()) {

1319 O << "()";

1320 return;

1321 }

1322

1323 O << "(\n";

1324

1325 for (const Argument &Arg : F->args()) {

1326 Type *Ty = Arg.getType();

1327 const std::string ParamSym = TLI->getParamName(F, Arg.getArgNo());

1328

1329 if (!IsFirst)

1330 O << ",\n";

1331

1332 IsFirst = false;

1333

1334

1335 if (IsKernelFunc) {

1336 const bool IsSampler = isSampler(Arg);

1337 const bool IsTexture = !IsSampler && isImageReadOnly(Arg);

1338 const bool IsSurface = !IsSampler && !IsTexture &&

1340 if (IsSampler || IsTexture || IsSurface) {

1342 O << "\t.param ";

1343 if (EmitImgPtr)

1344 O << ".u64 .ptr ";

1345

1346 if (IsSampler)

1347 O << ".samplerref ";

1348 else if (IsTexture)

1349 O << ".texref ";

1350 else

1351 O << ".surfref ";

1352 O << ParamSym;

1353 continue;

1354 }

1355 }

1356

1357 auto GetOptimalAlignForParam = [TLI, &DL, F, &Arg](Type *Ty) -> Align {

1358 if (MaybeAlign StackAlign =

1359 getAlign(*F, Arg.getArgNo() + AttributeList::FirstArgIndex))

1360 return StackAlign.value();

1361

1362 Align TypeAlign = TLI->getFunctionParamOptimizedAlign(F, Ty, DL);

1363 MaybeAlign ParamAlign =

1364 Arg.hasByValAttr() ? Arg.getParamAlign() : MaybeAlign();

1365 return std::max(TypeAlign, ParamAlign.valueOrOne());

1366 };

1367

1368 if (Arg.hasByValAttr()) {

1369

1370 Type *ETy = Arg.getParamByValType();

1371 assert(ETy && "Param should have byval type");

1372

1373

1374

1375

1376

1377 const Align OptimalAlign =

1378 IsKernelFunc ? GetOptimalAlignForParam(ETy)

1379 : TLI->getFunctionByValParamAlign(

1380 F, ETy, Arg.getParamAlign().valueOrOne(), DL);

1381

1382 O << "\t.param .align " << OptimalAlign.value() << " .b8 " << ParamSym

1383 << "[" << DL.getTypeAllocSize(ETy) << "]";

1384 continue;

1385 }

1386

1388

1389

1390

1391

1392 Align OptimalAlign = GetOptimalAlignForParam(Ty);

1393

1394 O << "\t.param .align " << OptimalAlign.value() << " .b8 " << ParamSym

1395 << "[" << DL.getTypeAllocSize(Ty) << "]";

1396

1397 continue;

1398 }

1399

1401 unsigned PTySizeInBits = 0;

1402 if (PTy) {

1403 PTySizeInBits =

1404 TLI->getPointerTy(DL, PTy->getAddressSpace()).getSizeInBits();

1405 assert(PTySizeInBits && "Invalid pointer size");

1406 }

1407

1408 if (IsKernelFunc) {

1409 if (PTy) {

1410 O << "\t.param .u" << PTySizeInBits << " .ptr";

1411

1412 switch (PTy->getAddressSpace()) {

1413 default:

1414 break;

1416 O << " .global";

1417 break;

1419 O << " .shared";

1420 break;

1422 O << " .const";

1423 break;

1425 O << " .local";

1426 break;

1427 }

1428

1429 O << " .align " << Arg.getParamAlign().valueOrOne().value() << " "

1430 << ParamSym;

1431 continue;

1432 }

1433

1434

1435 O << "\t.param .";

1436

1438 O << "u8";

1439 else

1440 O << getPTXFundamentalTypeStr(Ty);

1441 O << " " << ParamSym;

1442 continue;

1443 }

1444

1445

1446 unsigned Size;

1449 } else if (PTy) {

1450 assert(PTySizeInBits && "Invalid pointer size");

1451 Size = PTySizeInBits;

1452 } else

1454 O << "\t.param .b" << Size << " " << ParamSym;

1455 }

1456

1457 if (F->isVarArg()) {

1458 if (!IsFirst)

1459 O << ",\n";

1461 << TLI->getParamName(F, -1) << "[]";

1462 }

1463

1464 O << "\n)";

1465}

1466

1467void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(

1468 const MachineFunction &MF) {

1469 SmallString<128> Str;

1470 raw_svector_ostream O(Str);

1471

1472

1473

1474 const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo();

1475

1476

1477 const MachineFrameInfo &MFI = MF.getFrameInfo();

1479 if (NumBytes) {

1480 O << "\t.local .align " << MFI.getMaxAlign().value() << " .b8 \t"

1482 if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {

1483 O << "\t.reg .b64 \t%SP;\n"

1484 << "\t.reg .b64 \t%SPL;\n";

1485 } else {

1486 O << "\t.reg .b32 \t%SP;\n"

1487 << "\t.reg .b32 \t%SPL;\n";

1488 }

1489 }

1490

1491

1492

1493

1494

1495 for (unsigned I : llvm::seq(MRI->getNumVirtRegs())) {

1497 if (MRI->use_empty(VR) && MRI->def_empty(VR))

1498 continue;

1499 auto &RCRegMap = VRegMapping[MRI->getRegClass(VR)];

1500 RCRegMap[VR] = RCRegMap.size() + 1;

1501 }

1502

1503

1504

1505 for (const TargetRegisterClass *RC : TRI->regclasses()) {

1506 const unsigned N = VRegMapping[RC].size();

1507

1508

1509 if (N) {

1512 O << "\t.reg " << RCName << " \t" << RCStr << "<" << (N + 1) << ">;\n";

1513 }

1514 }

1515

1517}

1518

1519

1520

1521void NVPTXAsmPrinter::encodeDebugInfoRegisterNumbers(

1522 const MachineFunction &MF) {

1523 const NVPTXSubtarget &STI = MF.getSubtarget();

1524 const NVPTXRegisterInfo *registerInfo = STI.getRegisterInfo();

1525

1526

1527

1528

1530

1531 for (auto &classMap : VRegMapping) {

1532 for (auto &registerMapping : classMap.getSecond()) {

1533 auto reg = registerMapping.getFirst();

1535 }

1536 }

1537}

1538

1539void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp,

1540 raw_ostream &O) const {

1542 bool ignored;

1543 unsigned int numHex;

1544 const char *lead;

1545

1547 numHex = 8;

1548 lead = "0f";

1551 numHex = 16;

1552 lead = "0d";

1554 } else

1556

1559}

1560

1561void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {

1564 return;

1565 }

1567 printFPConstant(CFP, O);

1568 return;

1569 }

1571 O << "0";

1572 return;

1573 }

1575 const bool IsNonGenericPointer = GVar->getAddressSpace() != 0;

1576 if (EmitGeneric && isa<Function> (CPV) && !IsNonGenericPointer) {

1577 O << "generic(";

1579 O << ")";

1580 } else {

1582 }

1583 return;

1584 }

1586 const MCExpr *E = lowerConstantForGV(cast(Cexpr), false);

1587 printMCExpr(*E, O);

1588 return;

1589 }

1590 llvm_unreachable("Not scalar type found in printScalarConstant()");

1591}

1592

1593void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes,

1594 AggBuffer *AggBuffer) {

1596 int AllocSize = DL.getTypeAllocSize(CPV->getType());

1598

1599

1600 AggBuffer->addZeros(Bytes ? Bytes : AllocSize);

1601 return;

1602 }

1603

1604

1605 auto AddIntToBuffer = [AggBuffer, Bytes](const APInt &Val) {

1606 size_t NumBytes = (Val.getBitWidth() + 7) / 8;

1608

1609

1610

1611

1612 for (unsigned I = 0; I < NumBytes - 1; ++I) {

1613 Buf[I] = Val.extractBitsAsZExtValue(8, I * 8);

1614 }

1615 size_t LastBytePosition = (NumBytes - 1) * 8;

1616 size_t LastByteBits = Val.getBitWidth() - LastBytePosition;

1617 Buf[NumBytes - 1] =

1618 Val.extractBitsAsZExtValue(LastByteBits, LastBytePosition);

1619 AggBuffer->addBytes(Buf.data(), NumBytes, Bytes);

1620 };

1621

1625 AddIntToBuffer(CI->getValue());

1626 break;

1627 }

1629 if (const auto *CI =

1631 AddIntToBuffer(CI->getValue());

1632 break;

1633 }

1634 if (Cexpr->getOpcode() == Instruction::PtrToInt) {

1635 Value *V = Cexpr->getOperand(0)->stripPointerCasts();

1636 AggBuffer->addSymbol(V, Cexpr->getOperand(0));

1637 AggBuffer->addZeros(AllocSize);

1638 break;

1639 }

1640 }

1642 break;

1643

1648 AddIntToBuffer(cast(CPV)->getValueAPF().bitcastToAPInt());

1649 break;

1650

1653 AggBuffer->addSymbol(GVar, GVar);

1655 const Value *v = Cexpr->stripPointerCasts();

1656 AggBuffer->addSymbol(v, Cexpr);

1657 }

1658 AggBuffer->addZeros(AllocSize);

1659 break;

1660 }

1661

1666 bufferAggregateConstant(CPV, AggBuffer);

1667 if (Bytes > AllocSize)

1668 AggBuffer->addZeros(Bytes - AllocSize);

1670 AggBuffer->addZeros(Bytes);

1671 else

1673 break;

1674 }

1675

1676 default:

1678 }

1679}

1680

1681void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV,

1682 AggBuffer *aggBuffer) {

1684

1685 auto ExtendBuffer = [](APInt Val, AggBuffer *Buffer) {

1688 };

1689

1690

1692 ExtendBuffer(CI->getValue(), aggBuffer);

1693 return;

1694 }

1695

1696

1698 if (CFP->getType()->isFP128Ty()) {

1699 ExtendBuffer(CFP->getValueAPF().bitcastToAPInt(), aggBuffer);

1700 return;

1701 }

1702 }

1703

1704

1706 for (const auto &Op : CPV->operands())

1708 return;

1709 }

1710

1712 for (unsigned I : llvm::seq(CDS->getNumElements()))

1713 bufferLEByte(cast(CDS->getElementAsConstant(I)), 0, aggBuffer);

1714 return;

1715 }

1716

1722 ? DL.getStructLayout(ST)->getElementOffset(0) +

1723 DL.getTypeAllocSize(ST)

1724 : DL.getStructLayout(ST)->getElementOffset(I + 1);

1725 int Bytes = EndOffset - DL.getStructLayout(ST)->getElementOffset(I);

1727 }

1728 }

1729 return;

1730 }

1731 llvm_unreachable("unsupported constant type in printAggregateConstant()");

1732}

1733

1734

1735

1736

1737

1738const MCExpr *

1739NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV,

1740 bool ProcessingGeneric) const {

1742

1745

1748

1751 if (ProcessingGeneric)

1753 return Expr;

1754 }

1755

1757 if (!CE) {

1759 }

1760

1761 switch (CE->getOpcode()) {

1762 default:

1763 break;

1764

1765 case Instruction::AddrSpaceCast: {

1766

1768 if (DstTy->getAddressSpace() == 0)

1770

1771 break;

1772 }

1773

1774 case Instruction::GetElementPtr: {

1776

1777

1778 APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0);

1780

1781 const MCExpr *Base = lowerConstantForGV(CE->getOperand(0),

1782 ProcessingGeneric);

1783 if (!OffsetAI)

1784 return Base;

1785

1786 int64_t Offset = OffsetAI.getSExtValue();

1788 Ctx);

1789 }

1790

1791 case Instruction::Trunc:

1792

1793

1794

1795

1796 [[fallthrough]];

1797 case Instruction::BitCast:

1798 return lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);

1799

1800 case Instruction::IntToPtr: {

1802

1803

1804

1807 false, DL);

1808 if (Op)

1809 return lowerConstantForGV(Op, ProcessingGeneric);

1810

1811 break;

1812 }

1813

1814 case Instruction::PtrToInt: {

1816

1817

1818

1820 Type *Ty = CE->getType();

1821

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

1823

1824

1825

1826 if (DL.getTypeAllocSize(Ty) == DL.getTypeAllocSize(Op->getType()))

1827 return OpExpr;

1828

1829

1830

1831

1832 unsigned InBits = DL.getTypeAllocSizeInBits(Op->getType());

1835 }

1836

1837

1838

1839 case Instruction::Add: {

1840 const MCExpr *LHS = lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);

1841 const MCExpr *RHS = lowerConstantForGV(CE->getOperand(1), ProcessingGeneric);

1842 switch (CE->getOpcode()) {

1843 default: llvm_unreachable("Unknown binary operator constant cast expr");

1845 }

1846 }

1847 }

1848

1849

1850

1851

1853 if (C != CE)

1854 return lowerConstantForGV(C, ProcessingGeneric);

1855

1856

1857 std::string S;

1858 raw_string_ostream OS(S);

1859 OS << "Unsupported expression in static initializer: ";

1860 CE->printAsOperand(OS, false,

1861 ? nullptr : MF->getFunction().getParent());

1863}

1864

1865void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) const {

1866 OutContext.getAsmInfo()->printExpr(OS, Expr);

1867}

1868

1869

1870

1871bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,

1872 const char *ExtraCode, raw_ostream &O) {

1873 if (ExtraCode && ExtraCode[0]) {

1874 if (ExtraCode[1] != 0)

1875 return true;

1876

1877 switch (ExtraCode[0]) {

1878 default:

1879

1881 case 'r':

1882 break;

1883 }

1884 }

1885

1886 printOperand(MI, OpNo, O);

1887

1888 return false;

1889}

1890

1891bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI,

1892 unsigned OpNo,

1893 const char *ExtraCode,

1894 raw_ostream &O) {

1895 if (ExtraCode && ExtraCode[0])

1896 return true;

1897

1898 O << '[';

1899 printMemOperand(MI, OpNo, O);

1900 O << ']';

1901

1902 return false;

1903}

1904

1905void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, unsigned OpNum,

1906 raw_ostream &O) {

1907 const MachineOperand &MO = MI->getOperand(OpNum);

1911 if (MO.getReg() == NVPTX::VRDepot)

1913 else

1915 } else {

1916 emitVirtualRegister(MO.getReg(), O);

1917 }

1918 break;

1919

1922 break;

1923

1925 printFPConstant(MO.getFPImm(), O);

1926 break;

1927

1930 break;

1931

1934 break;

1935

1936 default:

1938 }

1939}

1940

1941void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, unsigned OpNum,

1942 raw_ostream &O, const char *Modifier) {

1943 printOperand(MI, OpNum, O);

1944

1945 if (Modifier && strcmp(Modifier, "add") == 0) {

1946 O << ", ";

1947 printOperand(MI, OpNum + 1, O);

1948 } else {

1949 if (MI->getOperand(OpNum + 1).isImm() &&

1950 MI->getOperand(OpNum + 1).getImm() == 0)

1951 return;

1952 O << "+";

1953 printOperand(MI, OpNum + 1, O);

1954 }

1955}

1956

1958

1960 false, false)

1961

1962

1964LLVMInitializeNVPTXAsmPrinter() {

1967}

assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")

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...

MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL

This file contains the simple types necessary to represent the attributes associated with functions a...

static GCRegistry::Add< StatepointGC > D("statepoint-example", "an example strategy for statepoint")

static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")

#define LLVM_EXTERNAL_VISIBILITY

This file contains the declarations for the subclasses of Constant, which represent the different fla...

This file defines the DenseMap class.

This file defines the DenseSet and SmallDenseSet classes.

Module.h This file contains the declarations for the Module class.

Promote Memory to Register

#define DEPOTNAME

Definition NVPTXAsmPrinter.cpp:95

static void discoverDependentGlobals(const Value *V, DenseSet< const GlobalVariable * > &Globals)

discoverDependentGlobals - Return a set of GlobalVariables on which V depends.

Definition NVPTXAsmPrinter.cpp:100

static bool canDemoteGlobalVar(const GlobalVariable *GV, Function const *&f)

Definition NVPTXAsmPrinter.cpp:580

static bool useFuncSeen(const Constant *C, const SmallPtrSetImpl< const Function * > &SeenSet)

Definition NVPTXAsmPrinter.cpp:597

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,...

Definition NVPTXAsmPrinter.cpp:116

static bool usedInGlobalVarDef(const Constant *C)

Definition NVPTXAsmPrinter.cpp:536

static bool usedInOneFunc(const User *U, Function const *&OneFunc)

Definition NVPTXAsmPrinter.cpp:551

#define INITIALIZE_PASS(passName, arg, name, cfg, analysis)

This file defines the SmallString class.

This file defines the SmallVector class.

static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")

static TableGen::Emitter::OptClass< SkeletonEmitter > X("gen-skeleton-class", "Generate example skeleton class")

static const fltSemantics & IEEEsingle()

static const fltSemantics & IEEEdouble()

static constexpr roundingMode rmNearestTiesToEven

LLVM_ABI opStatus convert(const fltSemantics &ToSemantics, roundingMode RM, bool *losesInfo)

APInt bitcastToAPInt() const

uint64_t getZExtValue() const

Get zero extended value.

LLVM_ABI uint64_t extractBitsAsZExtValue(unsigned numBits, unsigned bitPosition) const

unsigned getBitWidth() const

Return the number of bits in the APInt.

MCSymbol * getSymbol(const GlobalValue *GV) const

void EmitToStreamer(MCStreamer &S, const MCInst &Inst)

TargetMachine & TM

Target machine description.

virtual void PrintSymbolOperand(const MachineOperand &MO, raw_ostream &OS)

Print the MachineOperand as a symbol.

const MCAsmInfo * MAI

Target Asm Printer information.

MachineFunction * MF

The current machine function.

bool hasDebugInfo() const

Returns true if valid debug info is present.

virtual void emitFunctionBodyStart()

Targets can override this to emit stuff before the first basic block in the function.

bool doInitialization(Module &M) override

Set up the AsmPrinter when we are working on a new module.

unsigned getFunctionNumber() const

Return a unique ID for the current function.

MCSymbol * CurrentFnSym

The symbol for the current function.

MCContext & OutContext

This is the context for the output file that we are streaming.

bool doFinalization(Module &M) override

Shut down the asmprinter.

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.

std::unique_ptr< MCStreamer > OutStreamer

This is the MCStreamer object for the file we are generating.

virtual void emitFunctionBodyEnd()

Targets can override this to emit stuff after the last basic block in the function.

const DataLayout & getDataLayout() const

Return information about data layout.

virtual void emitFunctionEntryLabel()

EmitFunctionEntryLabel - Emit the label that is the entrypoint for the function.

void emitInitialRawDwarfLocDirective(const MachineFunction &MF)

Emits inital debug location directive.

MCSymbol * GetExternalSymbolSymbol(const Twine &Sym) const

Return the MCSymbol for the specified ExternalSymbol.

const MCSubtargetInfo & getSubtargetInfo() const

Return information about subtarget.

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.

const APFloat & getValueAPF() const

uint64_t getZExtValue() const

Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...

const APInt & getValue() const

Return the constant as an APInt value reference.

This is an important base class in LLVM.

LLVM_ABI bool isNullValue() const

Return true if this is the value that would be returned by getNullValue.

iterator find(const_arg_type_t< KeyT > Val)

DenseMapIterator< KeyT, ValueT, KeyInfoT, BucketT, true > const_iterator

Implements a dense probed hash-table based set.

LLVM_ABI const GlobalObject * getAliaseeObject() const

StringRef getSection() const

Get the custom section of this global if it has one.

bool hasSection() const

Check if this global has a custom object file section.

bool hasLinkOnceLinkage() const

bool hasExternalLinkage() const

LLVM_ABI bool isDeclaration() const

Return true if the primary definition of this global value is outside of the current translation unit...

bool hasLocalLinkage() const

bool hasPrivateLinkage() const

unsigned getAddressSpace() const

PointerType * getType() const

Global values are always pointers.

bool hasWeakLinkage() const

bool hasCommonLinkage() const

bool hasAvailableExternallyLinkage() const

Type * getValueType() const

const Constant * getInitializer() const

getInitializer - Return the initializer for this global variable.

bool hasInitializer() const

Definitions have initializers, declarations don't.

MaybeAlign getAlign() const

Returns the alignment of the given variable.

LLVM_ABI void diagnose(const DiagnosticInfo &DI)

Report a message to the currently installed diagnostic handler.

bool isLoopHeader(const BlockT *BB) const

LoopT * getLoopFor(const BlockT *BB) const

Return the inner most loop that BB lives in.

static const MCBinaryExpr * createAdd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx, SMLoc Loc=SMLoc())

static const MCBinaryExpr * createAnd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx)

static LLVM_ABI const MCConstantExpr * create(int64_t Value, MCContext &Ctx, bool PrintInHex=false, unsigned SizeInBytes=0)

Instances of this class represent a single low-level machine instruction.

void addOperand(const MCOperand Op)

void setOpcode(unsigned Op)

Instances of this class represent operands of the MCInst class.

static MCOperand createExpr(const MCExpr *Val)

static MCOperand createReg(MCRegister Reg)

static MCOperand createImm(int64_t Val)

static const MCSymbolRefExpr * create(const MCSymbol *Symbol, MCContext &Ctx, SMLoc Loc=SMLoc())

MCSymbol - Instances of this class represent a symbol name in the MC file, and MCSymbols are created ...

LLVM_ABI void print(raw_ostream &OS, const MCAsmInfo *MAI) const

print - Print the value to the stream OS.

LLVM_ABI MCSymbol * getSymbol() const

Return the MCSymbol for this basic block.

iterator_range< pred_iterator > predecessors()

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...

Function & getFunction()

Return the LLVM function that this machine code represents.

Representation of each machine instruction.

MachineOperand class - Representation of each machine instruction operand.

const GlobalValue * getGlobal() const

MachineBasicBlock * getMBB() const

MachineOperandType getType() const

getType - Returns the MachineOperandType for this operand.

const char * getSymbolName() 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

@ MO_ExternalSymbol

Name of external global symbol.

@ MO_FPImmediate

Floating-point immediate operand.

A Module instance is used to store all the information related to an LLVM module.

bool doInitialization(Module &M) override

Set up the AsmPrinter when we are working on a new module.

Definition NVPTXAsmPrinter.cpp:679

bool runOnMachineFunction(MachineFunction &F) override

Emit the specified function out to the OutStreamer.

Definition NVPTXAsmPrinter.cpp:368

std::string getVirtualRegisterName(unsigned) const

Definition NVPTXAsmPrinter.cpp:477

bool doFinalization(Module &M) override

Shut down the asmprinter.

Definition NVPTXAsmPrinter.cpp:776

const MCSymbol * getFunctionFrameSymbol() const override

Return symbol for the function pseudo stack if the stack frame is not a register based.

Definition NVPTXAsmPrinter.cpp:390

static const NVPTXFloatMCExpr * createConstantBFPHalf(const APFloat &Flt, MCContext &Ctx)

static const NVPTXFloatMCExpr * createConstantFPHalf(const APFloat &Flt, MCContext &Ctx)

static const NVPTXFloatMCExpr * createConstantFPSingle(const APFloat &Flt, MCContext &Ctx)

static const NVPTXFloatMCExpr * createConstantFPDouble(const APFloat &Flt, MCContext &Ctx)

static const NVPTXGenericMCSymbolRefExpr * create(const MCSymbolRefExpr *SymExpr, MCContext &Ctx)

static const char * getRegisterName(MCRegister Reg)

bool checkImageHandleSymbol(StringRef Symbol) const

Check if the symbol has a mapping.

void clearDebugRegisterMap() const

const char * getName(unsigned RegNo) const

std::string getTargetName() 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...

Implments NVPTX-specific streamer.

void outputDwarfFileDirectives()

Outputs the list of the DWARF '.file' directives to the streamer.

AnalysisType & getAnalysis() const

getAnalysis() - This function is used by subclasses to get to the analysis information ...

unsigned getAddressSpace() const

Return the address space of the Pointer type.

Wrapper class representing virtual and physical registers.

static Register index2VirtReg(unsigned Index)

Convert a 0-based index to a virtual register number.

constexpr bool isVirtual() const

Return true if the specified register number is in the virtual register namespace.

static constexpr bool isVirtualRegister(unsigned Reg)

Return true if the specified register number is in the virtual register namespace.

constexpr bool isPhysical() const

Return true if the specified register number is in the physical register namespace.

A templated base class for SmallPtrSet which provides the typesafe interface that is common across al...

std::pair< iterator, bool > insert(PtrType Ptr)

Inserts Ptr if and only if there is no element in the container equal to Ptr.

bool contains(ConstPtrType Ptr) const

SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...

This class consists of common code factored out of the SmallVector class to reduce code duplication b...

void push_back(const T &Elt)

This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.

StringRef - Represent a constant reference to a string, i.e.

bool starts_with(StringRef Prefix) const

Check if this string starts with the given Prefix.

const STC & getSubtarget(const Function &F) const

This method returns a pointer to the specified type of TargetSubtargetInfo.

Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...

bool isPointerTy() const

True if this is an instance of PointerType.

LLVM_ABI unsigned getPointerAddressSpace() const

Get the address space of this pointer or pointer vector type.

@ HalfTyID

16-bit floating point type

@ VoidTyID

type with no size

@ FloatTyID

32-bit floating point type

@ IntegerTyID

Arbitrary bit width integers.

@ FixedVectorTyID

Fixed width SIMD vector type.

@ BFloatTyID

16-bit floating point type (7-bit significand)

@ DoubleTyID

64-bit floating point type

@ FP128TyID

128-bit floating point type (112-bit significand)

LLVM_ABI TypeSize getPrimitiveSizeInBits() const LLVM_READONLY

Return the basic size of this type if it is a primitive type.

LLVM_ABI unsigned getScalarSizeInBits() const LLVM_READONLY

If this is a vector type, return the getPrimitiveSizeInBits value for the element type.

bool isFloatingPointTy() const

Return true if this is one of the floating-point types.

bool isIntOrPtrTy() const

Return true if this is an integer type or a pointer type.

bool isIntegerTy() const

True if this is an instance of IntegerType.

TypeID getTypeID() const

Return the type id for the type.

Value * getOperand(unsigned i) const

unsigned getNumOperands() const

LLVM Value Representation.

Type * getType() const

All values are typed, get the type of this value.

iterator_range< user_iterator > users()

LLVM_ABI StringRef getName() const

Return a constant reference to the value's name.

std::pair< iterator, bool > insert(const ValueT &V)

bool erase(const ValueT &V)

size_type count(const_arg_type_t< ValueT > V) const

Return 1 if the specified key is in the set, 0 otherwise.

This class implements an extremely fast bulk output stream that can only output to a stream.

A raw_ostream that writes to an std::string.

A raw_ostream that writes to an SmallVector or SmallString.

This provides a very simple, boring adaptor for a begin and end iterator into a range type.

#define llvm_unreachable(msg)

Marks that the current location is not supposed to be reachable.

constexpr char Align[]

Key for Kernel::Arg::Metadata::mAlign.

@ C

The default llvm calling convention, compatible with C.

@ CE

Windows NT (Windows on ARM)

std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)

Extract a Value from Metadata.

uint64_t read64le(const void *P)

uint32_t read32le(const void *P)

This is an optimization pass for GlobalISel generic memory operations.

FunctionAddr VTableAddr Value

bool isManaged(const Value &V)

StringRef getNVPTXRegClassStr(TargetRegisterClass const *RC)

bool shouldEmitPTXNoReturn(const Value *V, const TargetMachine &TM)

bool all_of(R &&range, UnaryPredicate P)

Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.

MaybeAlign getAlign(const CallInst &I, unsigned Index)

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.

std::optional< unsigned > getMaxNReg(const Function &F)

decltype(auto) dyn_cast(const From &Val)

dyn_cast - Return the argument parameter cast to the specified type.

StringRef getSamplerName(const Value &V)

bool isImageReadWrite(const Value &V)

bool isImageReadOnly(const Value &V)

iterator_range< T > make_range(T x, T y)

Convenience function for iterating over sub-ranges.

std::string utostr(uint64_t X, bool isNeg=false)

std::optional< unsigned > getMinCTASm(const Function &F)

SmallVector< unsigned, 3 > getReqNTID(const Function &F)

LLVM_ABI Constant * ConstantFoldConstant(const Constant *C, const DataLayout &DL, const TargetLibraryInfo *TLI=nullptr)

ConstantFoldConstant - Fold the constant using the specified DataLayout.

auto dyn_cast_or_null(const Y &Val)

auto formatv(bool Validate, const char *Fmt, Ts &&...Vals)

bool isSampler(const Value &V)

unsigned promoteScalarArgumentSize(unsigned size)

void clearAnnotationCache(const Module *Mod)

LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)

bool shouldPassAsArray(Type *Ty)

StringRef getNVPTXRegClassName(TargetRegisterClass const *RC)

bool isSurface(const Value &V)

class LLVM_GSL_OWNER SmallVector

Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...

FormattedNumber format_hex_no_prefix(uint64_t N, unsigned Width, bool Upper=false)

format_hex_no_prefix - Output N as a fixed width hexadecimal.

bool isa(const From &Val)

isa - Return true if the parameter to the template is an instance of one of the template type argu...

std::optional< unsigned > getMaxClusterRank(const Function &F)

StringRef getTextureName(const Value &V)

SmallVector< unsigned, 3 > getMaxNTID(const Function &F)

LLVM_ABI void write_hex(raw_ostream &S, uint64_t N, HexPrintStyle Style, std::optional< size_t > Width=std::nullopt)

DWARFExpression::Operation Op

ArrayRef(const T &OneElt) -> ArrayRef< T >

StringRef getSurfaceName(const Value &V)

Target & getTheNVPTXTarget64()

bool isKernelFunction(const Function &F)

decltype(auto) cast(const From &Val)

cast - Return the argument parameter cast to the specified type.

bool isTexture(const Value &V)

bool isImageWriteOnly(const Value &V)

auto seq(T Begin, T End)

Iterate over an integral type from Begin up to - but not including - End.

bool hasBlocksAreClusters(const Function &F)

SmallVector< unsigned, 3 > getClusterDim(const Function &F)

LLVM_ABI 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...

LLVM_ABI 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...

Target & getTheNVPTXTarget32()

constexpr uint64_t value() const

This is a hole in the type system and should not be abused.

Align valueOrOne() const

For convenience, returns a valid alignment or 1 if undefined.

RegisterAsmPrinter - Helper template for registering a target specific assembly printer,...