99 cl::desc(
"Lower GPU ctor / dtors to globals on the device."),
102#define DEPOTNAME "__local_depot"
112 if (
const User *U = dyn_cast<User>(V)) {
113 for (
unsigned i = 0, e = U->getNumOperands(); i != e; ++i) {
129 if (Visited.
count(GV))
133 if (!Visiting.
insert(GV).second)
151 NVPTX_MC::verifyInstructionPredicates(
MI->getOpcode(),
155 lowerToMCInst(
MI, Inst);
160bool NVPTXAsmPrinter::lowerImageHandleOperand(
const MachineInstr *
MI,
168 if (OpNo == 4 && MO.
isImm()) {
169 lowerImageHandleSymbol(MO.
getImm(), MCOp);
173 lowerImageHandleSymbol(MO.
getImm(), MCOp);
183 if (OpNo == VecSize && MO.
isImm()) {
184 lowerImageHandleSymbol(MO.
getImm(), MCOp);
191 if (OpNo == 0 && MO.
isImm()) {
192 lowerImageHandleSymbol(MO.
getImm(), MCOp);
199 if (OpNo == 1 && MO.
isImm()) {
200 lowerImageHandleSymbol(MO.
getImm(), MCOp);
210void NVPTXAsmPrinter::lowerImageHandleSymbol(
unsigned Index,
MCOperand &MCOp) {
223 if (
MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
231 for (
unsigned i = 0, e =
MI->getNumOperands(); i !=
e; ++i) {
236 if (lowerImageHandleOperand(
MI, i, MCOp)) {
242 if (lowerOperand(MO, MCOp))
296unsigned NVPTXAsmPrinter::encodeVirtualRegister(
unsigned Reg) {
301 unsigned RegNum = RegMap[
Reg];
306 if (RC == &NVPTX::Int1RegsRegClass) {
308 }
else if (RC == &NVPTX::Int16RegsRegClass) {
310 }
else if (RC == &NVPTX::Int32RegsRegClass) {
312 }
else if (RC == &NVPTX::Int64RegsRegClass) {
314 }
else if (RC == &NVPTX::Float32RegsRegClass) {
316 }
else if (RC == &NVPTX::Float64RegsRegClass) {
318 }
else if (RC == &NVPTX::Int128RegsRegClass) {
325 Ret |= (RegNum & 0x0FFFFFFF);
330 return Reg & 0x0FFFFFFF;
351 Type *Ty =
F->getReturnType();
363 if (
auto *ITy = dyn_cast<IntegerType>(Ty)) {
364 size = ITy->getBitWidth();
370 O <<
".param .b" <<
size <<
" func_retval0";
371 }
else if (isa<PointerType>(Ty)) {
372 O <<
".param .b" << TLI->getPointerTy(
DL).getSizeInBits()
375 unsigned totalsz =
DL.getTypeAllocSize(Ty);
376 Align RetAlignment = TLI->getFunctionArgumentAlignment(
378 O <<
".param .align " << RetAlignment.
value() <<
" .b8 func_retval0["
386 for (
unsigned i = 0, e = vtparts.
size(); i != e; ++i) {
388 EVT elemtype = vtparts[i];
390 elems = vtparts[i].getVectorNumElements();
391 elemtype = vtparts[i].getVectorElementType();
394 for (
unsigned j = 0, je = elems;
j != je; ++
j) {
398 O <<
".reg .b" << sz <<
" func_retval" << idx;
413 printReturnValStr(&F, O);
418bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
420 MachineLoopInfo &LI = getAnalysis<MachineLoopInfoWrapperPass>().getLI();
433 if (
const BasicBlock *PBB = PMBB->getBasicBlock()) {
435 PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
438 if (
MDNode *UnrollCountMD =
440 if (mdconst::extract<ConstantInt>(UnrollCountMD->getOperand(1))
452 if (isLoopHeaderOfNoUnroll(
MBB))
456void NVPTXAsmPrinter::emitFunctionEntryLabel() {
460 if (!GlobalsEmitted) {
462 GlobalsEmitted =
true;
468 emitLinkageDirective(F, O);
473 printReturnValStr(*
MF, O);
478 emitFunctionParamList(F, O);
482 emitKernelFunctionDirectives(*F, O);
492 setAndEmitFunctionVirtualRegisters(*
MF);
496 if (!SP->getUnit()->isDebugDirectivesOnly())
512void NVPTXAsmPrinter::emitFunctionBodyStart() {
519void NVPTXAsmPrinter::emitFunctionBodyEnd() {
529void NVPTXAsmPrinter::emitImplicitDef(
const MachineInstr *
MI)
const {
542void NVPTXAsmPrinter::emitKernelFunctionDirectives(
const Function &
F,
551 if (Reqntidx || Reqntidy || Reqntidz)
552 O <<
".reqntid " << Reqntidx.value_or(1) <<
", " << Reqntidy.value_or(1)
553 <<
", " << Reqntidz.value_or(1) <<
"\n";
562 if (Maxntidx || Maxntidy || Maxntidz)
563 O <<
".maxntid " << Maxntidx.value_or(1) <<
", " << Maxntidy.value_or(1)
564 <<
", " << Maxntidz.value_or(1) <<
"\n";
568 O <<
".minnctapersm " << Mincta <<
"\n";
570 unsigned Maxnreg = 0;
572 O <<
".maxnreg " << Maxnreg <<
"\n";
578 unsigned Maxclusterrank = 0;
580 O <<
".maxclusterrank " << Maxclusterrank <<
"\n";
590 assert(
I != VRegMapping.
end() &&
"Bad register class");
594 assert(VI != RegMap.
end() &&
"Bad virtual register");
595 unsigned MappedVR = VI->second;
603void NVPTXAsmPrinter::emitVirtualRegister(
unsigned int vr,
608void NVPTXAsmPrinter::emitAliasDeclaration(
const GlobalAlias *GA,
613 "NVPTX aliasee must be a non-kernel function definition");
619 emitDeclarationWithName(F,
getSymbol(GA), O);
623 emitDeclarationWithName(F,
getSymbol(F), O);
628 emitLinkageDirective(F, O);
633 printReturnValStr(F, O);
636 emitFunctionParamList(F, O);
648 return GV->getName() !=
"llvm.used";
651 for (
const User *U :
C->users())
652 if (
const Constant *
C = dyn_cast<Constant>(U))
660 if (
const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
661 if (othergv->getName() ==
"llvm.used")
666 if (
instr->getParent() &&
instr->getParent()->getParent()) {
668 if (oneFunc && (curFunc != oneFunc))
676 for (
const User *UU : U->users())
710 for (
const User *U :
C->users()) {
711 if (
const Constant *cu = dyn_cast<Constant>(U)) {
714 }
else if (
const Instruction *
I = dyn_cast<Instruction>(U)) {
731 if (
F.getAttributes().hasFnAttr(
"nvptx-libcall-callee")) {
732 emitDeclaration(&F, O);
736 if (
F.isDeclaration()) {
739 if (
F.getIntrinsicID())
741 emitDeclaration(&F, O);
744 for (
const User *U :
F.users()) {
745 if (
const Constant *
C = dyn_cast<Constant>(U)) {
750 emitDeclaration(&F, O);
756 emitDeclaration(&F, O);
761 if (!isa<Instruction>(U))
775 emitDeclaration(&F, O);
782 emitAliasDeclaration(&GA, O);
786 if (!GV)
return true;
788 if (!InitList)
return true;
792void NVPTXAsmPrinter::emitStartOfAsmFile(
Module &M) {
802 emitHeader(M, OS1, *STI);
814 bool IsOpenMP = M.getModuleFlag(
"openmp") !=
nullptr;
819 "Module has a nontrivial global ctor, which NVPTX does not support.");
825 "Module has a nontrivial global dtor, which NVPTX does not support.");
832 GlobalsEmitted =
false;
837void NVPTXAsmPrinter::emitGlobals(
const Module &M) {
841 emitDeclarations(M, OS2);
856 assert(GVVisited.
size() == M.global_size() &&
"Missed a global variable");
857 assert(GVVisiting.
size() == 0 &&
"Did not fully process a global variable");
865 printModuleLevelGV(GV, OS2,
false, STI);
887 O <<
"// Generated by LLVM NVPTX Back-End\n";
892 O <<
".version " << (PTXVersion / 10) <<
"." << (PTXVersion % 10) <<
"\n";
899 O <<
", texmode_independent";
901 bool HasFullDebugInfo =
false;
903 switch(
CU->getEmissionKind()) {
909 HasFullDebugInfo =
true;
912 if (HasFullDebugInfo)
915 if (HasFullDebugInfo)
920 O <<
".address_size ";
933 if (!GlobalsEmitted) {
935 GlobalsEmitted =
true;
949 OutStreamer->emitRawText(
"\t.section\t.debug_loc\t{\t}");
953 TS->outputDwarfFileDirectives();
971void NVPTXAsmPrinter::emitLinkageDirective(
const GlobalValue *V,
974 if (V->hasExternalLinkage()) {
975 if (isa<GlobalVariable>(V)) {
983 }
else if (
V->isDeclaration())
987 }
else if (
V->hasAppendingLinkage()) {
989 msg.append(
"Error: ");
990 msg.append(
"Symbol ");
992 msg.append(std::string(
V->getName()));
993 msg.append(
"has unsupported appending linkage type");
995 }
else if (!
V->hasInternalLinkage() &&
996 !
V->hasPrivateLinkage()) {
1002void NVPTXAsmPrinter::printModuleLevelGV(
const GlobalVariable *GVar,
1049 emitPTXGlobalVariable(GVar, O, STI);
1057 const Constant *Initializer =
nullptr;
1062 CI = dyn_cast<ConstantInt>(Initializer);
1071 O <<
"addr_mode_" << i <<
" = ";
1077 O <<
"clamp_to_border";
1080 O <<
"clamp_to_edge";
1091 O <<
"filter_mode = ";
1106 O <<
", force_unnormalized_coords = 1";
1116 if (strncmp(GVar->
getName().
data(),
"unrollpragma", 12) == 0)
1120 if (strncmp(GVar->
getName().
data(),
"filename", 8) == 0)
1126 const Function *demotedFunc =
nullptr;
1128 O <<
"// " << GVar->
getName() <<
" has been demoted\n";
1129 if (localDecls.find(demotedFunc) != localDecls.end())
1130 localDecls[demotedFunc].push_back(GVar);
1132 std::vector<const GlobalVariable *> temp;
1133 temp.push_back(GVar);
1134 localDecls[demotedFunc] = temp;
1145 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1147 O <<
" .attribute(.managed)";
1151 O <<
" .align " <<
A->value();
1153 O <<
" .align " << (int)
DL.getPrefTypeAlign(ETy).value();
1162 O << getPTXFundamentalTypeStr(ETy,
false);
1173 if (!Initializer->
isNullValue() && !isa<UndefValue>(Initializer)) {
1175 printScalarConstant(Initializer, O);
1184 "' is not allowed in addrspace(" +
1201 ElementSize =
DL.getTypeStoreSize(ETy);
1208 if (!isa<UndefValue>(Initializer) && !Initializer->
isNullValue()) {
1209 AggBuffer aggBuffer(ElementSize, *
this);
1210 bufferAggregateConstant(Initializer, &aggBuffer);
1211 if (aggBuffer.numSymbols()) {
1213 if (ElementSize % ptrSize ||
1214 !aggBuffer.allSymbolsAligned(ptrSize)) {
1218 "initialized packed aggregate with pointers '" +
1220 "' requires at least PTX ISA version 7.1");
1223 O <<
"[" << ElementSize <<
"] = {";
1224 aggBuffer.printBytes(O);
1227 O <<
" .u" << ptrSize * 8 <<
" ";
1229 O <<
"[" << ElementSize / ptrSize <<
"] = {";
1230 aggBuffer.printWords(O);
1236 O <<
"[" << ElementSize <<
"] = {";
1237 aggBuffer.printBytes(O);
1266void NVPTXAsmPrinter::AggBuffer::printSymbol(
unsigned nSym,
raw_ostream &os) {
1267 const Value *
v = Symbols[nSym];
1268 const Value *v0 = SymbolsBeforeStripping[nSym];
1269 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
1273 bool isGenericPointer = PTy && PTy->getAddressSpace() == 0;
1274 if (EmitGeneric && isGenericPointer && !isa<Function>(v)) {
1276 Name->print(os, AP.MAI);
1279 Name->print(os, AP.MAI);
1281 }
else if (
const ConstantExpr *CExpr = dyn_cast<ConstantExpr>(v0)) {
1282 const MCExpr *Expr = AP.lowerConstantForGV(cast<Constant>(CExpr),
false);
1283 AP.printMCExpr(*Expr, os);
1288void NVPTXAsmPrinter::AggBuffer::printBytes(
raw_ostream &os) {
1289 unsigned int ptrSize = AP.MAI->getCodePointerSize();
1294 unsigned int InitializerCount =
size;
1297 if (numSymbols() == 0)
1298 while (InitializerCount >= 1 && !buffer[InitializerCount - 1])
1301 symbolPosInBuffer.push_back(InitializerCount);
1302 unsigned int nSym = 0;
1303 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1304 for (
unsigned int pos = 0; pos < InitializerCount;) {
1307 if (pos != nextSymbolPos) {
1308 os << (
unsigned int)buffer[pos];
1315 std::string symText;
1317 printSymbol(nSym, oss);
1318 for (
unsigned i = 0; i < ptrSize; ++i) {
1322 os <<
"(" << symText <<
")";
1325 nextSymbolPos = symbolPosInBuffer[++nSym];
1326 assert(nextSymbolPos >= pos);
1330void NVPTXAsmPrinter::AggBuffer::printWords(
raw_ostream &os) {
1331 unsigned int ptrSize = AP.MAI->getCodePointerSize();
1332 symbolPosInBuffer.push_back(size);
1333 unsigned int nSym = 0;
1334 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1335 assert(nextSymbolPos % ptrSize == 0);
1336 for (
unsigned int pos = 0; pos <
size; pos += ptrSize) {
1339 if (pos == nextSymbolPos) {
1340 printSymbol(nSym, os);
1341 nextSymbolPos = symbolPosInBuffer[++nSym];
1342 assert(nextSymbolPos % ptrSize == 0);
1343 assert(nextSymbolPos >= pos + ptrSize);
1344 }
else if (ptrSize == 4)
1352 if (localDecls.find(f) == localDecls.end())
1355 std::vector<const GlobalVariable *> &gvars = localDecls[
f];
1362 O <<
"\t// demoted variable\n\t";
1363 printModuleLevelGV(GV, O,
true, STI);
1367void NVPTXAsmPrinter::emitPTXAddressSpace(
unsigned int AddressSpace,
1390NVPTXAsmPrinter::getPTXFundamentalTypeStr(
Type *Ty,
bool useB4PTR)
const {
1393 unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
1396 else if (NumBits <= 64) {
1397 std::string
name =
"u";
1398 return name + utostr(NumBits);
1416 assert((PtrSize == 64 || PtrSize == 32) &&
"Unexpected pointer size");
1434void NVPTXAsmPrinter::emitPTXGlobalVariable(
const GlobalVariable *GVar,
1447 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1449 O <<
" .attribute(.managed)";
1452 O <<
" .align " <<
A->value();
1454 O <<
" .align " << (int)
DL.getPrefTypeAlign(ETy).value();
1466 O << getPTXFundamentalTypeStr(ETy);
1472 int64_t ElementSize = 0;
1482 ElementSize =
DL.getTypeStoreSize(ETy);
1503 unsigned paramIndex = 0;
1509 if (
F->arg_empty() && !
F->isVarArg()) {
1516 for (
I =
F->arg_begin(), E =
F->arg_end();
I != E; ++
I, paramIndex++) {
1517 Type *Ty =
I->getType();
1529 if (hasImageHandles)
1530 O <<
"\t.param .u64 .ptr .surfref ";
1532 O <<
"\t.param .surfref ";
1533 O << TLI->getParamName(F, paramIndex);
1536 if (hasImageHandles)
1537 O <<
"\t.param .u64 .ptr .texref ";
1539 O <<
"\t.param .texref ";
1540 O << TLI->getParamName(F, paramIndex);
1543 if (hasImageHandles)
1544 O <<
"\t.param .u64 .ptr .samplerref ";
1546 O <<
"\t.param .samplerref ";
1547 O << TLI->getParamName(F, paramIndex);
1553 auto getOptimalAlignForParam = [TLI, &
DL, &PAL,
F,
1557 return StackAlign.
value();
1559 Align TypeAlign = TLI->getFunctionParamOptimizedAlign(F, Ty,
DL);
1560 MaybeAlign ParamAlign = PAL.getParamAlignment(paramIndex);
1561 return std::max(TypeAlign, ParamAlign.
valueOrOne());
1564 if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
1570 Align OptimalAlign = getOptimalAlignForParam(Ty);
1572 O <<
"\t.param .align " << OptimalAlign.
value() <<
" .b8 ";
1573 O << TLI->getParamName(F, paramIndex);
1574 O <<
"[" <<
DL.getTypeAllocSize(Ty) <<
"]";
1579 auto *PTy = dyn_cast<PointerType>(Ty);
1580 unsigned PTySizeInBits = 0;
1583 TLI->getPointerTy(
DL, PTy->getAddressSpace()).getSizeInBits();
1584 assert(PTySizeInBits &&
"Invalid pointer size");
1590 O <<
"\t.param .u" << PTySizeInBits <<
" ";
1594 int addrSpace = PTy->getAddressSpace();
1595 switch (addrSpace) {
1600 O <<
".ptr .const ";
1603 O <<
".ptr .shared ";
1606 O <<
".ptr .global ";
1609 Align ParamAlign =
I->getParamAlign().valueOrOne();
1610 O <<
".align " << ParamAlign.
value() <<
" ";
1612 O << TLI->getParamName(F, paramIndex);
1622 O << getPTXFundamentalTypeStr(Ty);
1624 O << TLI->getParamName(F, paramIndex);
1630 if (isa<IntegerType>(Ty)) {
1631 sz = cast<IntegerType>(Ty)->getBitWidth();
1634 assert(PTySizeInBits &&
"Invalid pointer size");
1639 O <<
"\t.param .b" << sz <<
" ";
1641 O <<
"\t.reg .b" << sz <<
" ";
1642 O << TLI->getParamName(F, paramIndex);
1647 Type *ETy = PAL.getParamByValType(paramIndex);
1648 assert(ETy &&
"Param should have byval type");
1650 if (isABI || isKernelFunc) {
1655 Align OptimalAlign =
1657 ? getOptimalAlignForParam(ETy)
1658 : TLI->getFunctionByValParamAlign(
1659 F, ETy, PAL.getParamAlignment(paramIndex).valueOrOne(),
DL);
1661 unsigned sz =
DL.getTypeAllocSize(ETy);
1662 O <<
"\t.param .align " << OptimalAlign.
value() <<
" .b8 ";
1663 O << TLI->getParamName(F, paramIndex);
1664 O <<
"[" << sz <<
"]";
1673 for (
unsigned i = 0, e = vtparts.
size(); i != e; ++i) {
1675 EVT elemtype = vtparts[i];
1677 elems = vtparts[i].getVectorNumElements();
1678 elemtype = vtparts[i].getVectorElementType();
1681 for (
unsigned j = 0, je = elems;
j != je; ++
j) {
1685 O <<
"\t.reg .b" << sz <<
" ";
1686 O << TLI->getParamName(F, paramIndex);
1699 if (
F->isVarArg()) {
1704 O << TLI->getParamName(F, -1) <<
"[]";
1710void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
1727 O <<
"\t.reg .b64 \t%SP;\n";
1728 O <<
"\t.reg .b64 \t%SPL;\n";
1730 O <<
"\t.reg .b32 \t%SP;\n";
1731 O <<
"\t.reg .b32 \t%SPL;\n";
1740 for (
unsigned i = 0; i < numVRs; i++) {
1744 int n = regmap.
size();
1745 regmap.
insert(std::make_pair(vr, n + 1));
1760 for (
unsigned i=0; i<
TRI->getNumRegClasses(); i++) {
1765 int n = regmap.
size();
1769 O <<
"\t.reg " << rcname <<
" \t" << rcStr <<
"<" << (n+1)
1780 unsigned int numHex;
1799 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1803 if (
const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
1804 printFPConstant(CFP, O);
1807 if (isa<ConstantPointerNull>(CPV)) {
1811 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1812 bool IsNonGenericPointer =
false;
1814 IsNonGenericPointer =
true;
1816 if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
1825 if (
const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1826 const MCExpr *E = lowerConstantForGV(cast<Constant>(Cexpr),
false);
1833void NVPTXAsmPrinter::bufferLEByte(
const Constant *CPV,
int Bytes,
1834 AggBuffer *AggBuffer) {
1836 int AllocSize =
DL.getTypeAllocSize(CPV->
getType());
1840 AggBuffer->addZeros(Bytes ? Bytes : AllocSize);
1845 auto AddIntToBuffer = [AggBuffer, Bytes](
const APInt &Val) {
1846 size_t NumBytes = (Val.getBitWidth() + 7) / 8;
1852 for (
unsigned I = 0;
I < NumBytes - 1; ++
I) {
1853 Buf[
I] = Val.extractBitsAsZExtValue(8,
I * 8);
1855 size_t LastBytePosition = (NumBytes - 1) * 8;
1856 size_t LastByteBits = Val.getBitWidth() - LastBytePosition;
1858 Val.extractBitsAsZExtValue(LastByteBits, LastBytePosition);
1859 AggBuffer->addBytes(Buf.data(), NumBytes, Bytes);
1864 if (
const auto CI = dyn_cast<ConstantInt>(CPV)) {
1868 if (
const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1869 if (
const auto *CI =
1874 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1875 Value *
V = Cexpr->getOperand(0)->stripPointerCasts();
1876 AggBuffer->addSymbol(V, Cexpr->getOperand(0));
1877 AggBuffer->addZeros(AllocSize);
1888 AddIntToBuffer(cast<ConstantFP>(CPV)->getValueAPF().bitcastToAPInt());
1892 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1893 AggBuffer->addSymbol(GVar, GVar);
1894 }
else if (
const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1895 const Value *
v = Cexpr->stripPointerCasts();
1896 AggBuffer->addSymbol(v, Cexpr);
1898 AggBuffer->addZeros(AllocSize);
1905 if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
1906 bufferAggregateConstant(CPV, AggBuffer);
1907 if (Bytes > AllocSize)
1908 AggBuffer->addZeros(Bytes - AllocSize);
1909 }
else if (isa<ConstantAggregateZero>(CPV))
1910 AggBuffer->addZeros(Bytes);
1921void NVPTXAsmPrinter::bufferAggregateConstant(
const Constant *CPV,
1922 AggBuffer *aggBuffer) {
1927 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1929 for (
unsigned I = 0, E =
DL.getTypeAllocSize(CPV->
getType());
I < E; ++
I) {
1931 aggBuffer->addBytes(&Byte, 1, 1);
1938 if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
1941 bufferLEByte(cast<Constant>(CPV->
getOperand(i)), 0, aggBuffer);
1946 dyn_cast<ConstantDataSequential>(CPV)) {
1947 if (CDS->getNumElements())
1948 for (
unsigned i = 0; i < CDS->getNumElements(); ++i)
1949 bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
1954 if (isa<ConstantStruct>(CPV)) {
1959 Bytes =
DL.getStructLayout(ST)->getElementOffset(0) +
1960 DL.getTypeAllocSize(ST) -
1961 DL.getStructLayout(ST)->getElementOffset(i);
1963 Bytes =
DL.getStructLayout(ST)->getElementOffset(i + 1) -
1964 DL.getStructLayout(ST)->getElementOffset(i);
1965 bufferLEByte(cast<Constant>(CPV->
getOperand(i)), Bytes, aggBuffer);
1978NVPTXAsmPrinter::lowerConstantForGV(
const Constant *CV,
bool ProcessingGeneric) {
1984 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
1987 if (
const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
1990 if (ProcessingGeneric) {
2002 switch (
CE->getOpcode()) {
2006 case Instruction::AddrSpaceCast: {
2009 if (DstTy->getAddressSpace() == 0)
2010 return lowerConstantForGV(cast<const Constant>(
CE->getOperand(0)),
true);
2015 case Instruction::GetElementPtr: {
2019 APInt OffsetAI(
DL.getPointerTypeSizeInBits(
CE->getType()), 0);
2020 cast<GEPOperator>(CE)->accumulateConstantOffset(
DL, OffsetAI);
2022 const MCExpr *
Base = lowerConstantForGV(
CE->getOperand(0),
2027 int64_t
Offset = OffsetAI.getSExtValue();
2032 case Instruction::Trunc:
2038 case Instruction::BitCast:
2039 return lowerConstantForGV(
CE->getOperand(0), ProcessingGeneric);
2041 case Instruction::IntToPtr: {
2050 return lowerConstantForGV(
Op, ProcessingGeneric);
2055 case Instruction::PtrToInt: {
2061 Type *Ty =
CE->getType();
2063 const MCExpr *OpExpr = lowerConstantForGV(
Op, ProcessingGeneric);
2067 if (
DL.getTypeAllocSize(Ty) ==
DL.getTypeAllocSize(
Op->getType()))
2073 unsigned InBits =
DL.getTypeAllocSizeInBits(
Op->getType());
2080 case Instruction::Add: {
2081 const MCExpr *
LHS = lowerConstantForGV(
CE->getOperand(0), ProcessingGeneric);
2082 const MCExpr *
RHS = lowerConstantForGV(
CE->getOperand(1), ProcessingGeneric);
2083 switch (
CE->getOpcode()) {
2095 return lowerConstantForGV(
C, ProcessingGeneric);
2100 OS <<
"Unsupported expression in static initializer: ";
2101 CE->printAsOperand(
OS,
false,
2110 return cast<MCTargetExpr>(&Expr)->printImpl(
OS,
MAI);
2112 OS << cast<MCConstantExpr>(Expr).getValue();
2138 if (isa<MCConstantExpr>(BE.
getLHS()) || isa<MCSymbolRefExpr>(BE.
getLHS()) ||
2139 isa<NVPTXGenericMCSymbolRefExpr>(BE.
getLHS())) {
2151 if (RHSC->getValue() < 0) {
2152 OS << RHSC->getValue();
2163 if (isa<MCConstantExpr>(BE.
getRHS()) || isa<MCSymbolRefExpr>(BE.
getRHS())) {
2179bool NVPTXAsmPrinter::PrintAsmOperand(
const MachineInstr *
MI,
unsigned OpNo,
2181 if (ExtraCode && ExtraCode[0]) {
2182 if (ExtraCode[1] != 0)
2185 switch (ExtraCode[0]) {
2194 printOperand(
MI, OpNo, O);
2199bool NVPTXAsmPrinter::PrintAsmMemoryOperand(
const MachineInstr *
MI,
2201 const char *ExtraCode,
2203 if (ExtraCode && ExtraCode[0])
2207 printMemOperand(
MI, OpNo, O);
2213void NVPTXAsmPrinter::printOperand(
const MachineInstr *
MI,
unsigned OpNum,
2219 if (MO.
getReg() == NVPTX::VRDepot)
2224 emitVirtualRegister(MO.
getReg(), O);
2249void NVPTXAsmPrinter::printMemOperand(
const MachineInstr *
MI,
unsigned OpNum,
2251 printOperand(
MI, OpNum, O);
2253 if (Modifier && strcmp(Modifier,
"add") == 0) {
2255 printOperand(
MI, OpNum + 1, O);
2257 if (
MI->getOperand(OpNum + 1).isImm() &&
2258 MI->getOperand(OpNum + 1).getImm() == 0)
2261 printOperand(
MI, OpNum + 1, O);
static cl::opt< bool > LowerCtorDtor("amdgpu-lower-global-ctor-dtor", cl::desc("Lower GPU ctor / dtors to globals on the device."), cl::init(true), cl::Hidden)
This file declares a class to represent arbitrary precision floating point values and provide a varie...
This file implements a class to represent arbitrary precision integral constant values and operations...
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
#define LLVM_EXTERNAL_VISIBILITY
This file contains the declarations for the subclasses of Constant, which represent the different fla...
Looks at all the uses of the given value Returns the Liveness deduced from the uses of this value Adds all uses that cause the result to be MaybeLive to MaybeLiveRetUses If the result is MaybeLiveUses might be modified but its content should be ignored(since it might not be complete). DeadArgumentEliminationPass
This file defines the DenseMap class.
This file defines the DenseSet and SmallDenseSet classes.
static GCMetadataPrinterRegistry::Add< ErlangGCPrinter > X("erlang", "erlang-compatible garbage collector")
unsigned const TargetRegisterInfo * TRI
Module.h This file contains the declarations for the Module class.
static bool isEmptyXXStructor(GlobalVariable *GV)
static bool usedInOneFunc(const User *U, Function const *&oneFunc)
static void VisitGlobalVariableForEmission(const GlobalVariable *GV, SmallVectorImpl< const GlobalVariable * > &Order, DenseSet< const GlobalVariable * > &Visited, DenseSet< const GlobalVariable * > &Visiting)
VisitGlobalVariableForEmission - Add GV to the list of GlobalVariable instances to be emitted,...
LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXAsmPrinter()
static bool usedInGlobalVarDef(const Constant *C)
static bool useFuncSeen(const Constant *C, DenseMap< const Function *, bool > &seenMap)
static cl::opt< bool > LowerCtorDtor("nvptx-lower-global-ctor-dtor", cl::desc("Lower GPU ctor / dtors to globals on the device."), cl::init(false), cl::Hidden)
static bool ShouldPassAsArray(Type *Ty)
static void DiscoverDependentGlobals(const Value *V, DenseSet< const GlobalVariable * > &Globals)
DiscoverDependentGlobals - Return a set of GlobalVariables on which V depends.
static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f)
static GCMetadataPrinterRegistry::Add< OcamlGCMetadataPrinter > Y("ocaml", "ocaml 3.10-compatible collector")
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
This file defines the SmallString class.
This file defines the SmallVector class.
opStatus convert(const fltSemantics &ToSemantics, roundingMode RM, bool *losesInfo)
APInt bitcastToAPInt() const
Class for arbitrary precision integers.
APInt getLoBits(unsigned numBits) const
Compute an APInt containing numBits lowbits from this APInt.
uint64_t getZExtValue() const
Get zero extended value.
void lshrInPlace(unsigned ShiftAmt)
Logical right-shift this APInt by ShiftAmt in place.
This class represents an incoming formal argument to a Function.
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.
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.
MCSymbol * GetExternalSymbolSymbol(Twine Sym) const
Return the MCSymbol for the specified ExternalSymbol.
virtual void emitBasicBlockStart(const MachineBasicBlock &MBB)
Targets can override this to emit stuff at the start of a basic block.
bool runOnMachineFunction(MachineFunction &MF) override
Emit the specified function out to the OutStreamer.
std::unique_ptr< MCStreamer > OutStreamer
This is the MCStreamer object for the file we are generating.
const DataLayout & getDataLayout() const
Return information about data layout.
void emitInitialRawDwarfLocDirective(const MachineFunction &MF)
Emits inital debug location directive.
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.
LLVM Basic Block Representation.
const Function * getParent() const
Return the enclosing method, or null if none.
ConstantArray - Constant Array Declarations.
ConstantDataSequential - A vector or array constant whose element type is a simple 1/2/4/8-byte integ...
A constant value that is initialized with an expression using other constant values.
ConstantFP - Floating Point Values [float, double].
const APFloat & getValueAPF() const
This is the shared class of boolean and integer constants.
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.
bool isNullValue() const
Return true if this is the value that would be returned by getNullValue.
This class represents an Operation in the Expression.
A parsed version of the target data layout string in and methods for querying it.
iterator find(const_arg_type_t< KeyT > Val)
bool contains(const_arg_type_t< KeyT > Val) const
Return true if the specified key is in the map, false otherwise.
std::pair< iterator, bool > insert(const std::pair< KeyT, ValueT > &KV)
Implements a dense probed hash-table based set.
DISubprogram * getSubprogram() const
Get the attached subprogram.
const GlobalObject * getAliaseeObject() const
StringRef getSection() const
Get the custom section of this global if it has one.
MaybeAlign getAlign() const
Returns the alignment of the given variable or function.
bool hasSection() const
Check if this global has a custom object file section.
bool hasLinkOnceLinkage() const
bool hasExternalLinkage() const
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
Module * getParent()
Get the module that this global value is contained inside of...
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.
This class describes a target machine that is implemented with the LLVM target-independent code gener...
bool isLoopHeader(const BlockT *BB) const
LoopT * getLoopFor(const BlockT *BB) const
Return the inner most loop that BB lives in.
unsigned getCodePointerSize() const
Get the code pointer size in bytes.
Binary assembler expressions.
const MCExpr * getLHS() const
Get the left-hand side expression of the binary operator.
const MCExpr * getRHS() const
Get the right-hand side expression of the binary operator.
static const MCBinaryExpr * createAnd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx)
static const MCBinaryExpr * createAdd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx)
Opcode getOpcode() const
Get the kind of this binary expression.
static const MCConstantExpr * create(int64_t Value, MCContext &Ctx, bool PrintInHex=false, unsigned SizeInBytes=0)
Context object for machine code objects.
MCSymbol * getOrCreateSymbol(const Twine &Name)
Lookup the symbol inside with the specified Name.
Base class for the full range of assembler expressions which are needed for parsing.
@ Unary
Unary expressions.
@ Constant
Constant expressions.
@ SymbolRef
References to labels and assigned expressions.
@ Target
Target specific expression.
@ Binary
Binary expressions.
Instances of this class represent a single low-level machine instruction.
void addOperand(const MCOperand Op)
void setOpcode(unsigned Op)
Describe properties that are true of each instruction in the target description file.
Instances of this class represent operands of the MCInst class.
static MCOperand createReg(unsigned Reg)
static MCOperand createExpr(const MCExpr *Val)
static MCOperand createImm(int64_t Val)
Represent a reference to a symbol from inside an expression.
const MCSymbol & getSymbol() const
static const MCSymbolRefExpr * create(const MCSymbol *Symbol, MCContext &Ctx)
MCSymbol - Instances of this class represent a symbol name in the MC file, and MCSymbols are created ...
void print(raw_ostream &OS, const MCAsmInfo *MAI) const
print - Print the value to the stream OS.
Unary assembler expressions.
Opcode getOpcode() const
Get the kind of this unary expression.
const MCExpr * getSubExpr() const
Get the child of this unary expression.
MCSymbol * getSymbol() const
Return the MCSymbol for this basic block.
iterator_range< pred_iterator > predecessors()
The MachineFrameInfo class represents an abstract stack frame until prolog/epilog code is inserted.
uint64_t getStackSize() const
Return the number of bytes that must be allocated to hold all of the fixed size frame objects.
Align getMaxAlign() const
Return the alignment in bytes that this function must be aligned to, which is greater than the defaul...
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
MachineFrameInfo & getFrameInfo()
getFrameInfo - Return the frame info object for the current function.
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
const LLVMTargetMachine & getTarget() const
getTarget - Return the target machine this machine code is compiled with
Ty * getInfo()
getInfo - Keep track of various per-function pieces of information for backends that would like to do...
Representation of each machine instruction.
MachineOperand class - Representation of each machine instruction operand.
const GlobalValue * getGlobal() const
MachineBasicBlock * getMBB() const
bool isImm() const
isImm - Tests if this is a MO_Immediate operand.
MachineOperandType getType() const
getType - Returns the MachineOperandType for this operand.
const char * getSymbolName() const
Register getReg() const
getReg - Returns the register number.
const ConstantFP * getFPImm() const
@ MO_Immediate
Immediate operand.
@ MO_GlobalAddress
Address of a global value.
@ MO_MachineBasicBlock
MachineBasicBlock reference.
@ MO_Register
Register operand.
@ MO_ExternalSymbol
Name of external global symbol.
@ MO_FPImmediate
Floating-point immediate operand.
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
A Module instance is used to store all the information related to an LLVM module.
bool doInitialization(Module &M) override
Set up the AsmPrinter when we are working on a new module.
bool runOnMachineFunction(MachineFunction &F) override
Emit the specified function out to the OutStreamer.
std::string getVirtualRegisterName(unsigned) const
bool doFinalization(Module &M) override
Shut down the asmprinter.
const MCSymbol * getFunctionFrameSymbol() const override
Return symbol for the function pseudo stack if the stack frame is not a register based.
static const NVPTXFloatMCExpr * createConstantBFPHalf(const APFloat &Flt, MCContext &Ctx)
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)
const char * getImageHandleSymbol(unsigned Idx) const
Returns the symbol name at the given index.
const char * getName(unsigned RegNo) const
std::string getTargetName() const
bool hasImageHandles() const
unsigned getMaxRequiredAlignment() const
bool hasMaskOperator() const
const NVPTXTargetLowering * getTargetLowering() const override
unsigned getPTXVersion() const
const NVPTXRegisterInfo * getRegisterInfo() const override
unsigned int getSmVersion() const
NVPTX::DrvInterface getDrvInterface() const
const NVPTXSubtarget * getSubtargetImpl(const Function &) const override
Virtual method implemented by subclasses that returns a reference to that target's TargetSubtargetInf...
UniqueStringSaver & getStrPool() const
Implments NVPTX-specific streamer.
void closeLastSection()
Close last section.
unsigned getAddressSpace() const
Return the address space of the Pointer type.
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.
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.
constexpr const char * data() const
data - Get a pointer to the start of the string (which may not be null terminated).
Class to represent struct types.
const STC & getSubtarget(const Function &F) const
This method returns a pointer to the specified type of TargetSubtargetInfo.
unsigned getPointerSizeInBits(unsigned AS) const
TargetRegisterInfo base class - We assume that the target defines a static array of TargetRegisterDes...
virtual const TargetRegisterInfo * getRegisterInfo() const
getRegisterInfo - If register information is available, return it.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
bool isVectorTy() const
True if this is an instance of VectorType.
bool isPointerTy() const
True if this is an instance of PointerType.
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
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
unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
bool isAggregateType() const
Return true if the type is an aggregate type.
bool isHalfTy() const
Return true if this is 'half', a 16-bit IEEE fp 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.
TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
StringRef save(const char *S)
Value * getOperand(unsigned i) const
unsigned getNumOperands() const
LLVM Value Representation.
Type * getType() const
All values are typed, get the type of this value.
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.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ C
The default llvm calling convention, compatible with C.
LegalityPredicate isVector(unsigned TypeIdx)
True iff the specified type index is a vector.
@ CE
Windows NT (Windows on ARM)
Reg
All possible values of the reg field in the ModR/M byte.
initializer< Ty > init(const Ty &Val)
uint64_t read64le(const void *P)
uint32_t read32le(const void *P)
This is an optimization pass for GlobalISel generic memory operations.
bool shouldEmitPTXNoReturn(const Value *V, const TargetMachine &TM)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
std::optional< unsigned > getMaxNTIDy(const Function &F)
std::string getSamplerName(const Value &val)
bool getMinCTASm(const Function &F, unsigned &x)
std::string getNVPTXRegClassName(TargetRegisterClass const *RC)
bool isImage(const Value &val)
std::optional< unsigned > getMaxNTIDz(const Function &F)
MaybeAlign getAlign(const Function &F, unsigned Index)
std::optional< unsigned > getMaxNTIDx(const Function &F)
Constant * ConstantFoldConstant(const Constant *C, const DataLayout &DL, const TargetLibraryInfo *TLI=nullptr)
ConstantFoldConstant - Fold the constant using the specified DataLayout.
bool isManaged(const Value &val)
unsigned promoteScalarArgumentSize(unsigned size)
bool isSurface(const Value &val)
void clearAnnotationCache(const Module *Mod)
void report_fatal_error(Error Err, bool gen_crash_diag=true)
Report a serious error, calling any installed error handler.
std::string getSurfaceName(const Value &val)
std::optional< unsigned > getReqNTIDy(const Function &F)
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 getMaxNReg(const Function &F, unsigned &x)
bool isTexture(const Value &val)
bool isImageWriteOnly(const Value &val)
bool isImageReadWrite(const Value &val)
void write_hex(raw_ostream &S, uint64_t N, HexPrintStyle Style, std::optional< size_t > Width=std::nullopt)
std::string getTextureName(const Value &val)
void ComputeValueVTs(const TargetLowering &TLI, const DataLayout &DL, Type *Ty, SmallVectorImpl< EVT > &ValueVTs, SmallVectorImpl< EVT > *MemVTs, SmallVectorImpl< TypeSize > *Offsets=nullptr, TypeSize StartingOffset=TypeSize::getZero())
ComputeValueVTs - Given an LLVM IR type, compute a sequence of EVTs that represent all the individual...
std::string getNVPTXRegClassStr(TargetRegisterClass const *RC)
Target & getTheNVPTXTarget64()
bool isKernelFunction(const Function &F)
std::optional< unsigned > getReqNTIDz(const Function &F)
std::optional< unsigned > getReqNTIDx(const Function &F)
bool isSampler(const Value &val)
Constant * ConstantFoldIntegerCast(Constant *C, Type *DestTy, bool IsSigned, const DataLayout &DL)
Constant fold a zext, sext or trunc, depending on IsSigned and whether the DestTy is wider or narrowe...
bool getMaxClusterRank(const Function &F, unsigned &x)
MDNode * GetUnrollMetadata(MDNode *LoopID, StringRef Name)
Given an llvm.loop loop id metadata node, returns the loop hint metadata node with the given name (fo...
Target & getTheNVPTXTarget32()
static const fltSemantics & IEEEsingle() LLVM_READNONE
static constexpr roundingMode rmNearestTiesToEven
static const fltSemantics & IEEEdouble() LLVM_READNONE
This struct is a compact representation of a valid (non-zero power of two) alignment.
uint64_t value() const
This is a hole in the type system and should not be abused.
TypeSize getSizeInBits() const
Return the size of the specified value type in bits.
bool isInteger() const
Return true if this is an integer or a vector integer type.
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
Align valueOrOne() const
For convenience, returns a valid alignment or 1 if undefined.
RegisterAsmPrinter - Helper template for registering a target specific assembly printer,...