95#define DEPOTNAME "__local_depot"
105 if (
const User *U = dyn_cast<User>(V)) {
106 for (
unsigned i = 0, e = U->getNumOperands(); i != e; ++i) {
122 if (Visited.
count(GV))
126 if (!Visiting.
insert(GV).second)
144 NVPTX_MC::verifyInstructionPredicates(
MI->getOpcode(),
148 lowerToMCInst(
MI, Inst);
153bool NVPTXAsmPrinter::lowerImageHandleOperand(
const MachineInstr *
MI,
161 if (OpNo == 4 && MO.
isImm()) {
162 lowerImageHandleSymbol(MO.
getImm(), MCOp);
166 lowerImageHandleSymbol(MO.
getImm(), MCOp);
176 if (OpNo == VecSize && MO.
isImm()) {
177 lowerImageHandleSymbol(MO.
getImm(), MCOp);
184 if (OpNo == 0 && MO.
isImm()) {
185 lowerImageHandleSymbol(MO.
getImm(), MCOp);
192 if (OpNo == 1 && MO.
isImm()) {
193 lowerImageHandleSymbol(MO.
getImm(), MCOp);
203void NVPTXAsmPrinter::lowerImageHandleSymbol(
unsigned Index,
MCOperand &MCOp) {
216 if (
MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
224 for (
unsigned i = 0, e =
MI->getNumOperands(); i !=
e; ++i) {
229 if (lowerImageHandleOperand(
MI, i, MCOp)) {
235 if (lowerOperand(MO, MCOp))
285unsigned NVPTXAsmPrinter::encodeVirtualRegister(
unsigned Reg) {
290 unsigned RegNum = RegMap[
Reg];
295 if (RC == &NVPTX::Int1RegsRegClass) {
297 }
else if (RC == &NVPTX::Int16RegsRegClass) {
299 }
else if (RC == &NVPTX::Int32RegsRegClass) {
301 }
else if (RC == &NVPTX::Int64RegsRegClass) {
303 }
else if (RC == &NVPTX::Float32RegsRegClass) {
305 }
else if (RC == &NVPTX::Float64RegsRegClass) {
307 }
else if (RC == &NVPTX::Float16RegsRegClass) {
309 }
else if (RC == &NVPTX::Float16x2RegsRegClass) {
316 Ret |= (RegNum & 0x0FFFFFFF);
321 return Reg & 0x0FFFFFFF;
337 Type *Ty =
F->getReturnType();
349 if (
auto *ITy = dyn_cast<IntegerType>(Ty)) {
350 size = ITy->getBitWidth();
360 O <<
".param .b" <<
size <<
" func_retval0";
361 }
else if (isa<PointerType>(Ty)) {
362 O <<
".param .b" << TLI->getPointerTy(
DL).getSizeInBits()
365 unsigned totalsz =
DL.getTypeAllocSize(Ty);
366 unsigned retAlignment = 0;
368 retAlignment = TLI->getFunctionParamOptimizedAlign(F, Ty,
DL).value();
369 O <<
".param .align " << retAlignment <<
" .b8 func_retval0[" << totalsz
377 for (
unsigned i = 0, e = vtparts.
size(); i != e; ++i) {
379 EVT elemtype = vtparts[i];
381 elems = vtparts[i].getVectorNumElements();
382 elemtype = vtparts[i].getVectorElementType();
385 for (
unsigned j = 0, je = elems; j != je; ++j) {
389 O <<
".reg .b" << sz <<
" func_retval" << idx;
404 printReturnValStr(&F, O);
409bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
424 if (
const BasicBlock *PBB = PMBB->getBasicBlock()) {
426 PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
429 if (
MDNode *UnrollCountMD =
431 if (mdconst::extract<ConstantInt>(UnrollCountMD->getOperand(1))
432 ->getZExtValue() == 1)
443 if (isLoopHeaderOfNoUnroll(
MBB))
447void NVPTXAsmPrinter::emitFunctionEntryLabel() {
451 if (!GlobalsEmitted) {
453 GlobalsEmitted =
true;
459 emitLinkageDirective(F, O);
464 printReturnValStr(*
MF, O);
469 emitFunctionParamList(*
MF, O);
472 emitKernelFunctionDirectives(*F, O);
482 setAndEmitFunctionVirtualRegisters(*
MF);
499void NVPTXAsmPrinter::emitFunctionBodyStart() {
506void NVPTXAsmPrinter::emitFunctionBodyEnd() {
516void NVPTXAsmPrinter::emitImplicitDef(
const MachineInstr *
MI)
const {
529void NVPTXAsmPrinter::emitKernelFunctionDirectives(
const Function &
F,
534 unsigned reqntidx, reqntidy, reqntidz;
535 bool specified =
false;
550 O <<
".reqntid " << reqntidx <<
", " << reqntidy <<
", " << reqntidz
556 unsigned maxntidx, maxntidy, maxntidz;
572 O <<
".maxntid " << maxntidx <<
", " << maxntidy <<
", " << maxntidz
577 O <<
".minnctapersm " << mincta <<
"\n";
581 O <<
".maxnreg " << maxnreg <<
"\n";
592 assert(
I != VRegMapping.
end() &&
"Bad register class");
596 assert(
VI != RegMap.
end() &&
"Bad virtual register");
597 unsigned MappedVR =
VI->second;
605void NVPTXAsmPrinter::emitVirtualRegister(
unsigned int vr,
611 emitLinkageDirective(F, O);
616 printReturnValStr(F, O);
619 emitFunctionParamList(F, O);
630 return GV->getName() !=
"llvm.used";
633 for (
const User *U :
C->users())
634 if (
const Constant *
C = dyn_cast<Constant>(U))
642 if (
const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
643 if (othergv->getName() ==
"llvm.used")
648 if (
instr->getParent() &&
instr->getParent()->getParent()) {
650 if (oneFunc && (curFunc != oneFunc))
692 for (
const User *U :
C->users()) {
693 if (
const Constant *cu = dyn_cast<Constant>(U)) {
696 }
else if (
const Instruction *
I = dyn_cast<Instruction>(U)) {
703 if (seenMap.
find(caller) != seenMap.
end())
713 if (
F.getAttributes().hasFnAttr(
"nvptx-libcall-callee")) {
714 emitDeclaration(&F, O);
718 if (
F.isDeclaration()) {
721 if (
F.getIntrinsicID())
723 emitDeclaration(&F, O);
726 for (
const User *U :
F.users()) {
727 if (
const Constant *
C = dyn_cast<Constant>(U)) {
732 emitDeclaration(&F, O);
738 emitDeclaration(&F, O);
743 if (!isa<Instruction>(U))
756 if (seenMap.
find(caller) != seenMap.
end()) {
757 emitDeclaration(&F, O);
766 if (!GV)
return true;
768 if (!InitList)
return true;
772void NVPTXAsmPrinter::emitStartOfAsmFile(
Module &M) {
782 emitHeader(M, OS1, *STI);
787 if (M.alias_size()) {
793 "Module has a nontrivial global ctor, which NVPTX does not support.");
798 "Module has a nontrivial global dtor, which NVPTX does not support.");
805 GlobalsEmitted =
false;
810void NVPTXAsmPrinter::emitGlobals(
const Module &M) {
814 emitDeclarations(M, OS2);
829 assert(GVVisited.
size() == M.getGlobalList().size() &&
830 "Missed a global variable");
831 assert(GVVisiting.
size() == 0 &&
"Did not fully process a global variable");
838 for (
unsigned i = 0, e = Globals.
size(); i != e; ++i)
839 printModuleLevelGV(Globals[i], OS2,
false, STI);
849 O <<
"// Generated by LLVM NVPTX Back-End\n";
854 O <<
".version " << (PTXVersion / 10) <<
"." << (PTXVersion % 10) <<
"\n";
861 O <<
", texmode_independent";
863 bool HasFullDebugInfo =
false;
865 switch(
CU->getEmissionKind()) {
871 HasFullDebugInfo =
true;
874 if (HasFullDebugInfo)
882 O <<
".address_size ";
897 if (!GlobalsEmitted) {
899 GlobalsEmitted =
true;
913 OutStreamer->emitRawText(
"\t.section\t.debug_loc\t{\t}");
917 TS->outputDwarfFileDirectives();
935void NVPTXAsmPrinter::emitLinkageDirective(
const GlobalValue *V,
939 if (isa<GlobalVariable>(V)) {
953 msg.append(
"Error: ");
954 msg.append(
"Symbol ");
956 msg.append(std::string(V->
getName()));
957 msg.append(
"has unsupported appending linkage type");
966void NVPTXAsmPrinter::printModuleLevelGV(
const GlobalVariable *GVar,
1011 emitPTXGlobalVariable(GVar, O, STI);
1019 const Constant *Initializer =
nullptr;
1024 CI = dyn_cast<ConstantInt>(Initializer);
1033 O <<
"addr_mode_" << i <<
" = ";
1039 O <<
"clamp_to_border";
1042 O <<
"clamp_to_edge";
1053 O <<
"filter_mode = ";
1068 O <<
", force_unnormalized_coords = 1";
1078 if (strncmp(GVar->
getName().
data(),
"unrollpragma", 12) == 0)
1082 if (strncmp(GVar->
getName().
data(),
"filename", 8) == 0)
1088 const Function *demotedFunc =
nullptr;
1090 O <<
"// " << GVar->
getName() <<
" has been demoted\n";
1091 if (localDecls.find(demotedFunc) != localDecls.end())
1092 localDecls[demotedFunc].push_back(GVar);
1094 std::vector<const GlobalVariable *> temp;
1095 temp.push_back(GVar);
1096 localDecls[demotedFunc] = temp;
1102 emitPTXAddressSpace(PTy->getAddressSpace(), O);
1107 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1109 O <<
" .attribute(.managed)";
1113 O <<
" .align " <<
A->value();
1115 O <<
" .align " << (int)
DL.getPrefTypeAlign(ETy).value();
1124 O << getPTXFundamentalTypeStr(ETy,
false);
1135 if (!Initializer->
isNullValue() && !isa<UndefValue>(Initializer)) {
1137 printScalarConstant(Initializer, O);
1146 "' is not allowed in addrspace(" +
1147 Twine(PTy->getAddressSpace()) +
")");
1152 unsigned int ElementSize = 0;
1163 ElementSize =
DL.getTypeStoreSize(ETy);
1170 if (!isa<UndefValue>(Initializer) && !Initializer->
isNullValue()) {
1171 AggBuffer aggBuffer(ElementSize, *
this);
1172 bufferAggregateConstant(Initializer, &aggBuffer);
1173 if (aggBuffer.numSymbols()) {
1175 if (ElementSize % ptrSize ||
1176 !aggBuffer.allSymbolsAligned(ptrSize)) {
1180 "initialized packed aggregate with pointers '" +
1182 "' requires at least PTX ISA version 7.1");
1185 O <<
"[" << ElementSize <<
"] = {";
1186 aggBuffer.printBytes(O);
1189 O <<
" .u" << ptrSize * 8 <<
" ";
1191 O <<
"[" << ElementSize / ptrSize <<
"] = {";
1192 aggBuffer.printWords(O);
1198 O <<
"[" << ElementSize <<
"] = {";
1199 aggBuffer.printBytes(O);
1228void NVPTXAsmPrinter::AggBuffer::printSymbol(
unsigned nSym,
raw_ostream &os) {
1229 const Value *v = Symbols[nSym];
1230 const Value *v0 = SymbolsBeforeStripping[nSym];
1231 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
1235 bool isGenericPointer = PTy && PTy->getAddressSpace() == 0;
1236 if (EmitGeneric && isGenericPointer && !isa<Function>(v)) {
1238 Name->print(os, AP.MAI);
1241 Name->print(os, AP.MAI);
1243 }
else if (
const ConstantExpr *CExpr = dyn_cast<ConstantExpr>(v0)) {
1244 const MCExpr *Expr = AP.lowerConstantForGV(cast<Constant>(CExpr),
false);
1245 AP.printMCExpr(*Expr, os);
1250void NVPTXAsmPrinter::AggBuffer::printBytes(
raw_ostream &os) {
1251 unsigned int ptrSize = AP.MAI->getCodePointerSize();
1252 symbolPosInBuffer.push_back(
size);
1253 unsigned int nSym = 0;
1254 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1255 for (
unsigned int pos = 0; pos <
size;) {
1258 if (pos != nextSymbolPos) {
1259 os << (
unsigned int)buffer[pos];
1266 std::string symText;
1268 printSymbol(nSym, oss);
1269 for (
unsigned i = 0; i < ptrSize; ++i) {
1273 os <<
"(" << symText <<
")";
1276 nextSymbolPos = symbolPosInBuffer[++nSym];
1277 assert(nextSymbolPos >= pos);
1281void NVPTXAsmPrinter::AggBuffer::printWords(
raw_ostream &os) {
1282 unsigned int ptrSize = AP.MAI->getCodePointerSize();
1283 symbolPosInBuffer.push_back(
size);
1284 unsigned int nSym = 0;
1285 unsigned int nextSymbolPos = symbolPosInBuffer[nSym];
1286 assert(nextSymbolPos % ptrSize == 0);
1287 for (
unsigned int pos = 0; pos <
size; pos += ptrSize) {
1290 if (pos == nextSymbolPos) {
1291 printSymbol(nSym, os);
1292 nextSymbolPos = symbolPosInBuffer[++nSym];
1293 assert(nextSymbolPos % ptrSize == 0);
1294 assert(nextSymbolPos >= pos + ptrSize);
1295 }
else if (ptrSize == 4)
1303 if (localDecls.find(f) == localDecls.end())
1306 std::vector<const GlobalVariable *> &gvars = localDecls[f];
1313 O <<
"\t// demoted variable\n\t";
1314 printModuleLevelGV(GV, O,
true, STI);
1318void NVPTXAsmPrinter::emitPTXAddressSpace(
unsigned int AddressSpace,
1341NVPTXAsmPrinter::getPTXFundamentalTypeStr(
Type *Ty,
bool useB4PTR)
const {
1344 unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
1347 else if (NumBits <= 64) {
1348 std::string
name =
"u";
1349 return name + utostr(NumBits);
1365 assert((PtrSize == 64 || PtrSize == 32) &&
"Unexpected pointer size");
1383void NVPTXAsmPrinter::emitPTXGlobalVariable(
const GlobalVariable *GVar,
1396 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1398 O <<
" .attribute(.managed)";
1401 O <<
" .align " <<
A->value();
1403 O <<
" .align " << (int)
DL.getPrefTypeAlign(ETy).value();
1415 O << getPTXFundamentalTypeStr(ETy);
1421 int64_t ElementSize = 0;
1431 ElementSize =
DL.getTypeStoreSize(ETy);
1448 O <<
"_param_" << paramIndex;
1458 unsigned paramIndex = 0;
1464 if (
F->arg_empty() && !
F->isVarArg()) {
1471 for (
I =
F->arg_begin(),
E =
F->arg_end();
I !=
E; ++
I, paramIndex++) {
1472 Type *Ty =
I->getType();
1483 std::string sname = std::string(
I->getName());
1485 if (hasImageHandles)
1486 O <<
"\t.param .u64 .ptr .surfref ";
1488 O <<
"\t.param .surfref ";
1490 O <<
"_param_" << paramIndex;
1493 if (hasImageHandles)
1494 O <<
"\t.param .u64 .ptr .texref ";
1496 O <<
"\t.param .texref ";
1498 O <<
"_param_" << paramIndex;
1501 if (hasImageHandles)
1502 O <<
"\t.param .u64 .ptr .samplerref ";
1504 O <<
"\t.param .samplerref ";
1506 O <<
"_param_" << paramIndex;
1512 auto getOptimalAlignForParam = [TLI, &
DL, &PAL,
F,
1514 Align TypeAlign = TLI->getFunctionParamOptimizedAlign(F, Ty,
DL);
1515 MaybeAlign ParamAlign = PAL.getParamAlignment(paramIndex);
1516 return std::max(TypeAlign, ParamAlign.
valueOrOne());
1519 if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
1525 Align OptimalAlign = getOptimalAlignForParam(Ty);
1527 O <<
"\t.param .align " << OptimalAlign.
value() <<
" .b8 ";
1528 printParamName(
I, paramIndex, O);
1529 O <<
"[" <<
DL.getTypeAllocSize(Ty) <<
"]";
1534 auto *PTy = dyn_cast<PointerType>(Ty);
1535 unsigned PTySizeInBits = 0;
1538 TLI->getPointerTy(
DL, PTy->getAddressSpace()).getSizeInBits();
1539 assert(PTySizeInBits &&
"Invalid pointer size");
1545 O <<
"\t.param .u" << PTySizeInBits <<
" ";
1549 int addrSpace = PTy->getAddressSpace();
1550 switch (addrSpace) {
1555 O <<
".ptr .const ";
1558 O <<
".ptr .shared ";
1561 O <<
".ptr .global ";
1564 Align ParamAlign =
I->getParamAlign().valueOrOne();
1565 O <<
".align " << ParamAlign.
value() <<
" ";
1567 printParamName(
I, paramIndex, O);
1577 O << getPTXFundamentalTypeStr(Ty);
1579 printParamName(
I, paramIndex, O);
1585 if (isa<IntegerType>(Ty)) {
1586 sz = cast<IntegerType>(Ty)->getBitWidth();
1589 assert(PTySizeInBits &&
"Invalid pointer size");
1599 O <<
"\t.param .b" << sz <<
" ";
1601 O <<
"\t.reg .b" << sz <<
" ";
1602 printParamName(
I, paramIndex, O);
1607 Type *ETy = PAL.getParamByValType(paramIndex);
1608 assert(ETy &&
"Param should have byval type");
1610 if (isABI || isKernelFunc) {
1615 Align OptimalAlign =
1617 ? getOptimalAlignForParam(ETy)
1618 : TLI->getFunctionByValParamAlign(
1619 F, ETy, PAL.getParamAlignment(paramIndex).valueOrOne(),
DL);
1621 unsigned sz =
DL.getTypeAllocSize(ETy);
1622 O <<
"\t.param .align " << OptimalAlign.
value() <<
" .b8 ";
1623 printParamName(
I, paramIndex, O);
1624 O <<
"[" << sz <<
"]";
1633 for (
unsigned i = 0, e = vtparts.
size(); i != e; ++i) {
1635 EVT elemtype = vtparts[i];
1637 elems = vtparts[i].getVectorNumElements();
1638 elemtype = vtparts[i].getVectorElementType();
1641 for (
unsigned j = 0, je = elems; j != je; ++j) {
1645 O <<
"\t.reg .b" << sz <<
" ";
1646 printParamName(
I, paramIndex, O);
1659 if (
F->isVarArg()) {
1674 emitFunctionParamList(&F, O);
1677void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
1694 O <<
"\t.reg .b64 \t%SP;\n";
1695 O <<
"\t.reg .b64 \t%SPL;\n";
1697 O <<
"\t.reg .b32 \t%SP;\n";
1698 O <<
"\t.reg .b32 \t%SPL;\n";
1707 for (
unsigned i = 0; i < numVRs; i++) {
1711 int n = regmap.
size();
1712 regmap.
insert(std::make_pair(vr, n + 1));
1727 for (
unsigned i=0; i<
TRI->getNumRegClasses(); i++) {
1732 int n = regmap.
size();
1736 O <<
"\t.reg " << rcname <<
" \t" << rcStr <<
"<" << (n+1)
1747 unsigned int numHex;
1766 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1770 if (
const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
1771 printFPConstant(CFP, O);
1774 if (isa<ConstantPointerNull>(CPV)) {
1778 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1779 bool IsNonGenericPointer =
false;
1781 IsNonGenericPointer =
true;
1783 if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
1792 if (
const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1793 const MCExpr *
E = lowerConstantForGV(cast<Constant>(Cexpr),
false);
1800void NVPTXAsmPrinter::bufferLEByte(
const Constant *CPV,
int Bytes,
1801 AggBuffer *AggBuffer) {
1803 int AllocSize =
DL.getTypeAllocSize(CPV->
getType());
1807 AggBuffer->addZeros(Bytes ? Bytes : AllocSize);
1812 auto AddIntToBuffer = [AggBuffer, Bytes](
const APInt &Val) {
1813 size_t NumBytes = (Val.getBitWidth() + 7) / 8;
1815 for (
unsigned I = 0;
I < NumBytes; ++
I) {
1816 Buf[
I] = Val.extractBitsAsZExtValue(8,
I * 8);
1818 AggBuffer->addBytes(Buf.data(), NumBytes, Bytes);
1823 if (
const auto CI = dyn_cast<ConstantInt>(CPV)) {
1827 if (
const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1828 if (
const auto *CI =
1833 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1835 AggBuffer->addSymbol(V, Cexpr->getOperand(0));
1836 AggBuffer->addZeros(AllocSize);
1847 AddIntToBuffer(cast<ConstantFP>(CPV)->getValueAPF().bitcastToAPInt());
1851 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1852 AggBuffer->addSymbol(GVar, GVar);
1853 }
else if (
const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1855 AggBuffer->addSymbol(v, Cexpr);
1857 AggBuffer->addZeros(AllocSize);
1864 if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
1865 bufferAggregateConstant(CPV, AggBuffer);
1866 if (Bytes > AllocSize)
1867 AggBuffer->addZeros(Bytes - AllocSize);
1868 }
else if (isa<ConstantAggregateZero>(CPV))
1869 AggBuffer->addZeros(Bytes);
1880void NVPTXAsmPrinter::bufferAggregateConstant(
const Constant *CPV,
1881 AggBuffer *aggBuffer) {
1886 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1888 for (
unsigned I = 0,
E =
DL.getTypeAllocSize(CPV->
getType());
I <
E; ++
I) {
1890 aggBuffer->addBytes(&Byte, 1, 1);
1897 if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
1900 bufferLEByte(cast<Constant>(CPV->
getOperand(i)), 0, aggBuffer);
1905 dyn_cast<ConstantDataSequential>(CPV)) {
1906 if (CDS->getNumElements())
1907 for (
unsigned i = 0; i < CDS->getNumElements(); ++i)
1908 bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
1913 if (isa<ConstantStruct>(CPV)) {
1918 Bytes =
DL.getStructLayout(ST)->getElementOffset(0) +
1919 DL.getTypeAllocSize(ST) -
1920 DL.getStructLayout(ST)->getElementOffset(i);
1922 Bytes =
DL.getStructLayout(ST)->getElementOffset(i + 1) -
1923 DL.getStructLayout(ST)->getElementOffset(i);
1924 bufferLEByte(cast<Constant>(CPV->
getOperand(i)), Bytes, aggBuffer);
1937NVPTXAsmPrinter::lowerConstantForGV(
const Constant *CV,
bool ProcessingGeneric) {
1943 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
1946 if (
const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
1949 if (ProcessingGeneric) {
1961 switch (
CE->getOpcode()) {
1968 return lowerConstantForGV(
C, ProcessingGeneric);
1973 OS <<
"Unsupported expression in static initializer: ";
1974 CE->printAsOperand(OS,
false,
1979 case Instruction::AddrSpaceCast: {
1982 if (DstTy->getAddressSpace() == 0) {
1983 return lowerConstantForGV(cast<const Constant>(
CE->getOperand(0)),
true);
1987 OS <<
"Unsupported expression in static initializer: ";
1988 CE->printAsOperand(OS,
false,
1993 case Instruction::GetElementPtr: {
1997 APInt OffsetAI(
DL.getPointerTypeSizeInBits(
CE->getType()), 0);
1998 cast<GEPOperator>(CE)->accumulateConstantOffset(
DL, OffsetAI);
2000 const MCExpr *
Base = lowerConstantForGV(
CE->getOperand(0),
2005 int64_t
Offset = OffsetAI.getSExtValue();
2010 case Instruction::Trunc:
2016 case Instruction::BitCast:
2017 return lowerConstantForGV(
CE->getOperand(0), ProcessingGeneric);
2019 case Instruction::IntToPtr: {
2027 return lowerConstantForGV(Op, ProcessingGeneric);
2030 case Instruction::PtrToInt: {
2036 Type *Ty =
CE->getType();
2038 const MCExpr *OpExpr = lowerConstantForGV(Op, ProcessingGeneric);
2042 if (
DL.getTypeAllocSize(Ty) ==
DL.getTypeAllocSize(
Op->getType()))
2048 unsigned InBits =
DL.getTypeAllocSizeInBits(
Op->getType());
2055 case Instruction::Add: {
2056 const MCExpr *
LHS = lowerConstantForGV(
CE->getOperand(0), ProcessingGeneric);
2057 const MCExpr *
RHS = lowerConstantForGV(
CE->getOperand(1), ProcessingGeneric);
2058 switch (
CE->getOpcode()) {
2070 return cast<MCTargetExpr>(&Expr)->printImpl(OS,
MAI);
2072 OS << cast<MCConstantExpr>(Expr).getValue();
2098 if (isa<MCConstantExpr>(BE.
getLHS()) || isa<MCSymbolRefExpr>(BE.
getLHS()) ||
2099 isa<NVPTXGenericMCSymbolRefExpr>(BE.
getLHS())) {
2100 printMCExpr(*BE.
getLHS(), OS);
2103 printMCExpr(*BE.
getLHS(), OS);
2111 if (RHSC->getValue() < 0) {
2112 OS << RHSC->getValue();
2123 if (isa<MCConstantExpr>(BE.
getRHS()) || isa<MCSymbolRefExpr>(BE.
getRHS())) {
2124 printMCExpr(*BE.
getRHS(), OS);
2127 printMCExpr(*BE.
getRHS(), OS);
2139bool NVPTXAsmPrinter::PrintAsmOperand(
const MachineInstr *
MI,
unsigned OpNo,
2141 if (ExtraCode && ExtraCode[0]) {
2142 if (ExtraCode[1] != 0)
2145 switch (ExtraCode[0]) {
2154 printOperand(
MI, OpNo, O);
2159bool NVPTXAsmPrinter::PrintAsmMemoryOperand(
const MachineInstr *
MI,
2161 const char *ExtraCode,
2163 if (ExtraCode && ExtraCode[0])
2167 printMemOperand(
MI, OpNo, O);
2173void NVPTXAsmPrinter::printOperand(
const MachineInstr *
MI,
int opNum,
2179 if (MO.
getReg() == NVPTX::VRDepot)
2184 emitVirtualRegister(MO.
getReg(), O);
2209void NVPTXAsmPrinter::printMemOperand(
const MachineInstr *
MI,
int opNum,
2211 printOperand(
MI, opNum, O);
2213 if (Modifier && strcmp(Modifier,
"add") == 0) {
2215 printOperand(
MI, opNum + 1, O);
2217 if (
MI->getOperand(opNum + 1).isImm() &&
2218 MI->getOperand(opNum + 1).getImm() == 0)
2221 printOperand(
MI, opNum + 1, O);
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This header is deprecated in favour of llvm/TargetParser/Triple.h.
This file declares a class to represent arbitrary precision floating point values and provide a varie...
This file implements a class to represent arbitrary precision integral constant values and operations...
This file contains the simple types necessary to represent the attributes associated with functions a...
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
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...
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 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)
return ToRemove size() > 0
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 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.
MachineModuleInfo * MMI
This is a pointer to the current MachineModuleInfo.
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(StringRef 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.
static Constant * getIntegerCast(Constant *C, Type *Ty, bool IsSigned)
Create a ZExt, Bitcast or Trunc for integer -> integer casts.
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.
A parsed version of the target data layout string in and methods for querying it.
iterator find(const_arg_type_t< KeyT > Val)
std::pair< iterator, bool > insert(const std::pair< KeyT, ValueT > &KV)
Implements a dense probed hash-table based set.
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 hasPrivateLinkage() const
Module * getParent()
Get the module that this global value is contained inside of...
bool hasInternalLinkage() const
PointerType * getType() const
Global values are always pointers.
bool hasWeakLinkage() const
bool hasCommonLinkage() const
bool hasAppendingLinkage() 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...
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.
bool isLoopHeader(const MachineBasicBlock *BB) const
True if the block is a loop header node.
MachineLoop * getLoopFor(const MachineBasicBlock *BB) const
Return the innermost loop that BB lives in.
bool hasDebugInfo() const
Returns true if valid debug info is present.
MachineOperand class - Representation of each machine instruction operand.
const GlobalValue * getGlobal() const
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 * 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.
bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
static Register index2VirtReg(unsigned Index)
Convert a 0-based index to a virtual register number.
static bool isVirtualRegister(unsigned Reg)
Return true if the specified register number is in the virtual register namespace.
bool isVirtual() const
Return true if the specified register number is in the virtual 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 startswith(StringRef Prefix) const
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.
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.
iterator_range< user_iterator > users()
const Value * stripPointerCasts() const
Strip off pointer casts, all-zero GEPs and address space casts.
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.
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)
std::string getSamplerName(const Value &val)
bool getAlign(const Function &F, unsigned index, unsigned &align)
bool getMinCTASm(const Function &F, unsigned &x)
std::string getNVPTXRegClassName(TargetRegisterClass const *RC)
bool isImage(const Value &val)
bool getMaxNTIDz(const Function &F, unsigned &z)
Constant * ConstantFoldConstant(const Constant *C, const DataLayout &DL, const TargetLibraryInfo *TLI=nullptr)
ConstantFoldConstant - Fold the constant using the specified DataLayout.
bool isManaged(const Value &val)
unsigned promoteScalarArgumentSize(unsigned size)
bool isSurface(const Value &val)
void clearAnnotationCache(const Module *Mod)
void report_fatal_error(Error Err, bool gen_crash_diag=true)
Report a serious error, calling any installed error handler.
std::string getSurfaceName(const Value &val)
bool getReqNTIDx(const Function &F, unsigned &x)
bool getReqNTIDy(const Function &F, unsigned &y)
FormattedNumber format_hex_no_prefix(uint64_t N, unsigned Width, bool Upper=false)
format_hex_no_prefix - Output N as a fixed width hexadecimal.
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)
std::string getNVPTXRegClassStr(TargetRegisterClass const *RC)
Target & getTheNVPTXTarget64()
bool isKernelFunction(const Function &F)
bool getReqNTIDz(const Function &F, unsigned &z)
bool getMaxNTIDx(const Function &F, unsigned &x)
void ComputeValueVTs(const TargetLowering &TLI, const DataLayout &DL, Type *Ty, SmallVectorImpl< EVT > &ValueVTs, SmallVectorImpl< uint64_t > *Offsets=nullptr, uint64_t StartingOffset=0)
ComputeValueVTs - Given an LLVM IR type, compute a sequence of EVTs that represent all the individual...
bool getMaxNTIDy(const Function &F, unsigned &y)
bool isSampler(const Value &val)
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,...