Go to the documentation of this file.
93 #define DEPOTNAME "__local_depot"
103 if (
const User *U = dyn_cast<User>(V)) {
104 for (
unsigned i = 0,
e = U->getNumOperands();
i !=
e; ++
i) {
120 if (Visited.
count(GV))
124 if (!Visiting.
insert(GV).second)
143 lowerToMCInst(
MI, Inst);
148 bool NVPTXAsmPrinter::lowerImageHandleOperand(
const MachineInstr *
MI,
156 if (OpNo == 4 && MO.
isImm()) {
157 lowerImageHandleSymbol(MO.
getImm(), MCOp);
161 lowerImageHandleSymbol(MO.
getImm(), MCOp);
172 lowerImageHandleSymbol(MO.
getImm(), MCOp);
179 if (OpNo == 0 && MO.
isImm()) {
180 lowerImageHandleSymbol(MO.
getImm(), MCOp);
187 if (OpNo == 1 && MO.
isImm()) {
188 lowerImageHandleSymbol(MO.
getImm(), MCOp);
198 void NVPTXAsmPrinter::lowerImageHandleSymbol(
unsigned Index,
MCOperand &MCOp) {
204 std::string *SymNamePtr =
212 if (
MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
220 for (
unsigned i = 0,
e =
MI->getNumOperands();
i !=
e; ++
i) {
225 if (lowerImageHandleOperand(
MI,
i, MCOp)) {
231 if (lowerOperand(MO, MCOp))
281 unsigned NVPTXAsmPrinter::encodeVirtualRegister(
unsigned Reg) {
286 unsigned RegNum = RegMap[
Reg];
291 if (RC == &NVPTX::Int1RegsRegClass) {
293 }
else if (RC == &NVPTX::Int16RegsRegClass) {
295 }
else if (RC == &NVPTX::Int32RegsRegClass) {
297 }
else if (RC == &NVPTX::Int64RegsRegClass) {
299 }
else if (RC == &NVPTX::Float32RegsRegClass) {
301 }
else if (RC == &NVPTX::Float64RegsRegClass) {
303 }
else if (RC == &NVPTX::Float16RegsRegClass) {
305 }
else if (RC == &NVPTX::Float16x2RegsRegClass) {
312 Ret |= (RegNum & 0x0FFFFFFF);
317 return Reg & 0x0FFFFFFF;
333 Type *Ty =
F->getReturnType();
345 if (
auto *ITy = dyn_cast<IntegerType>(Ty)) {
346 size = ITy->getBitWidth();
357 O <<
".param .b" <<
size <<
" func_retval0";
358 }
else if (isa<PointerType>(Ty)) {
359 O <<
".param .b" << TLI->getPointerTy(
DL).getSizeInBits()
362 unsigned totalsz =
DL.getTypeAllocSize(Ty);
363 unsigned retAlignment = 0;
365 retAlignment = TLI->getFunctionParamOptimizedAlign(
F, Ty,
DL).value();
366 O <<
".param .align " << retAlignment <<
" .b8 func_retval0[" << totalsz
374 for (
unsigned i = 0,
e = vtparts.size();
i !=
e; ++
i) {
376 EVT elemtype = vtparts[
i];
378 elems = vtparts[
i].getVectorNumElements();
379 elemtype = vtparts[
i].getVectorElementType();
382 for (
unsigned j = 0, je = elems;
j != je; ++
j) {
386 O <<
".reg .b" << sz <<
" func_retval" << idx;
401 printReturnValStr(&
F,
O);
407 bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
422 if (
const BasicBlock *PBB = PMBB->getBasicBlock()) {
424 PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
435 if (isLoopHeaderOfNoUnroll(
MBB))
439 void NVPTXAsmPrinter::emitFunctionEntryLabel() {
443 if (!GlobalsEmitted) {
445 GlobalsEmitted =
true;
451 emitLinkageDirective(
F,
O);
456 printReturnValStr(*
MF,
O);
461 emitFunctionParamList(*
MF,
O);
464 emitKernelFunctionDirectives(*
F,
O);
471 setAndEmitFunctionVirtualRegisters(*
MF);
488 void NVPTXAsmPrinter::emitFunctionBodyStart() {
495 void NVPTXAsmPrinter::emitFunctionBodyEnd() {
505 void NVPTXAsmPrinter::emitImplicitDef(
const MachineInstr *
MI)
const {
518 void NVPTXAsmPrinter::emitKernelFunctionDirectives(
const Function &
F,
523 unsigned reqntidx, reqntidy, reqntidz;
524 bool specified =
false;
539 O <<
".reqntid " << reqntidx <<
", " << reqntidy <<
", " << reqntidz
545 unsigned maxntidx, maxntidy, maxntidz;
561 O <<
".maxntid " << maxntidx <<
", " << maxntidy <<
", " << maxntidz
566 O <<
".minnctapersm " << mincta <<
"\n";
570 O <<
".maxnreg " << maxnreg <<
"\n";
581 assert(
I != VRegMapping.
end() &&
"Bad register class");
585 assert(
VI != RegMap.
end() &&
"Bad virtual register");
586 unsigned MappedVR =
VI->second;
594 void NVPTXAsmPrinter::emitVirtualRegister(
unsigned int vr,
600 emitLinkageDirective(
F,
O);
605 printReturnValStr(
F,
O);
608 emitFunctionParamList(
F,
O);
617 return GV->getName() !=
"llvm.used";
621 if (
const Constant *
C = dyn_cast<Constant>(U))
629 if (
const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
630 if (othergv->getName() ==
"llvm.used")
634 if (
const Instruction *instr = dyn_cast<Instruction>(U)) {
635 if (instr->getParent() && instr->getParent()->getParent()) {
637 if (oneFunc && (curFunc != oneFunc))
680 if (
const Constant *cu = dyn_cast<Constant>(U)) {
683 }
else if (
const Instruction *
I = dyn_cast<Instruction>(U)) {
700 if (
F.getAttributes().hasFnAttr(
"nvptx-libcall-callee")) {
701 emitDeclaration(&
F,
O);
705 if (
F.isDeclaration()) {
708 if (
F.getIntrinsicID())
710 emitDeclaration(&
F,
O);
713 for (
const User *U :
F.users()) {
714 if (
const Constant *
C = dyn_cast<Constant>(U)) {
719 emitDeclaration(&
F,
O);
725 emitDeclaration(&
F,
O);
730 if (!isa<Instruction>(U))
744 emitDeclaration(&
F,
O);
753 if (!GV)
return true;
755 if (!InitList)
return true;
759 void NVPTXAsmPrinter::emitStartOfAsmFile(
Module &
M) {
769 emitHeader(
M, OS1, *STI);
774 if (
M.alias_size()) {
780 "Module has a nontrivial global ctor, which NVPTX does not support.");
785 "Module has a nontrivial global dtor, which NVPTX does not support.");
792 GlobalsEmitted =
false;
797 void NVPTXAsmPrinter::emitGlobals(
const Module &
M) {
801 emitDeclarations(
M, OS2);
816 assert(GVVisited.
size() ==
M.getGlobalList().size() &&
817 "Missed a global variable");
818 assert(GVVisiting.
size() == 0 &&
"Did not fully process a global variable");
825 for (
unsigned i = 0,
e = Globals.size();
i !=
e; ++
i)
826 printModuleLevelGV(Globals[
i], OS2,
false, STI);
836 O <<
"// Generated by LLVM NVPTX Back-End\n";
841 O <<
".version " << (PTXVersion / 10) <<
"." << (PTXVersion % 10) <<
"\n";
848 O <<
", texmode_independent";
850 bool HasFullDebugInfo =
false;
852 switch(
CU->getEmissionKind()) {
858 HasFullDebugInfo =
true;
861 if (HasFullDebugInfo)
869 O <<
".address_size ";
884 if (!GlobalsEmitted) {
886 GlobalsEmitted =
true;
898 TS->closeLastSection();
900 OutStreamer->emitRawText(
"\t.section\t.debug_loc\t{\t}");
904 TS->outputDwarfFileDirectives();
932 void NVPTXAsmPrinter::emitLinkageDirective(
const GlobalValue *V,
936 if (isa<GlobalVariable>(V)) {
950 msg.append(
"Error: ");
951 msg.append(
"Symbol ");
953 msg.append(std::string(V->
getName()));
954 msg.append(
"has unsupported appending linkage type");
963 void NVPTXAsmPrinter::printModuleLevelGV(
const GlobalVariable *GVar,
1008 emitPTXGlobalVariable(GVar,
O, STI);
1021 CI = dyn_cast<ConstantInt>(Initializer);
1030 O <<
"addr_mode_" <<
i <<
" = ";
1036 O <<
"clamp_to_border";
1039 O <<
"clamp_to_edge";
1050 O <<
"filter_mode = ";
1065 O <<
", force_unnormalized_coords = 1";
1075 if (strncmp(GVar->
getName().
data(),
"unrollpragma", 12) == 0)
1079 if (strncmp(GVar->
getName().
data(),
"filename", 8) == 0)
1085 const Function *demotedFunc =
nullptr;
1087 O <<
"// " << GVar->
getName() <<
" has been demoted\n";
1088 if (localDecls.find(demotedFunc) != localDecls.end())
1089 localDecls[demotedFunc].push_back(GVar);
1091 std::vector<const GlobalVariable *> temp;
1092 temp.push_back(GVar);
1093 localDecls[demotedFunc] = temp;
1104 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1106 O <<
" .attribute(.managed)";
1110 O <<
" .align " <<
A->value();
1112 O <<
" .align " << (
int)
DL.getPrefTypeAlignment(ETy);
1121 O << getPTXFundamentalTypeStr(ETy,
false);
1132 if (!
Initializer->isNullValue() && !isa<UndefValue>(Initializer)) {
1134 printScalarConstant(Initializer,
O);
1143 "' is not allowed in addrspace(" +
1149 unsigned int ElementSize = 0;
1160 ElementSize =
DL.getTypeStoreSize(ETy);
1167 if (!isa<UndefValue>(Initializer) && !
Initializer->isNullValue()) {
1169 bufferAggregateConstant(Initializer, &aggBuffer);
1170 if (aggBuffer.numSymbols) {
1175 O << ElementSize / 8;
1180 O << ElementSize / 4;
1220 if (localDecls.find(
f) == localDecls.end())
1223 std::vector<const GlobalVariable *> &gvars = localDecls[
f];
1230 O <<
"\t// demoted variable\n\t";
1231 printModuleLevelGV(GV,
O,
true, STI);
1235 void NVPTXAsmPrinter::emitPTXAddressSpace(
unsigned int AddressSpace,
1258 NVPTXAsmPrinter::getPTXFundamentalTypeStr(
Type *Ty,
bool useB4PTR)
const {
1261 unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
1264 else if (NumBits <= 64) {
1265 std::string
name =
"u";
1266 return name + utostr(NumBits);
1296 void NVPTXAsmPrinter::emitPTXGlobalVariable(
const GlobalVariable *GVar,
1309 ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
1311 O <<
" .attribute(.managed)";
1314 O <<
" .align " <<
A->value();
1316 O <<
" .align " << (
int)
DL.getPrefTypeAlignment(ETy);
1328 O << getPTXFundamentalTypeStr(ETy);
1334 int64_t ElementSize = 0;
1344 ElementSize =
DL.getTypeStoreSize(ETy);
1361 O <<
"_param_" << paramIndex;
1371 unsigned paramIndex = 0;
1376 MVT thePointerTy = TLI->getPointerTy(
DL);
1378 if (
F->arg_empty()) {
1385 for (
I =
F->arg_begin(),
E =
F->arg_end();
I !=
E; ++
I, paramIndex++) {
1386 Type *Ty =
I->getType();
1397 std::string sname = std::string(
I->getName());
1399 if (hasImageHandles)
1400 O <<
"\t.param .u64 .ptr .surfref ";
1402 O <<
"\t.param .surfref ";
1404 O <<
"_param_" << paramIndex;
1407 if (hasImageHandles)
1408 O <<
"\t.param .u64 .ptr .texref ";
1410 O <<
"\t.param .texref ";
1412 O <<
"_param_" << paramIndex;
1415 if (hasImageHandles)
1416 O <<
"\t.param .u64 .ptr .samplerref ";
1418 O <<
"\t.param .samplerref ";
1420 O <<
"_param_" << paramIndex;
1426 auto getOptimalAlignForParam = [TLI, &
DL, &PAL,
F,
1428 Align TypeAlign = TLI->getFunctionParamOptimizedAlign(
F, Ty,
DL);
1429 MaybeAlign ParamAlign = PAL.getParamAlignment(paramIndex);
1430 return max(TypeAlign, ParamAlign);
1433 if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
1439 Align OptimalAlign = getOptimalAlignForParam(Ty);
1441 O <<
"\t.param .align " << OptimalAlign.
value() <<
" .b8 ";
1442 printParamName(
I, paramIndex,
O);
1443 O <<
"[" <<
DL.getTypeAllocSize(Ty) <<
"]";
1448 auto *PTy = dyn_cast<PointerType>(Ty);
1457 switch (addrSpace) {
1462 O <<
".ptr .const ";
1465 O <<
".ptr .shared ";
1468 O <<
".ptr .global ";
1471 Align ParamAlign =
I->getParamAlign().valueOrOne();
1472 O <<
".align " << ParamAlign.
value() <<
" ";
1474 printParamName(
I, paramIndex,
O);
1484 O << getPTXFundamentalTypeStr(Ty);
1486 printParamName(
I, paramIndex,
O);
1492 if (isa<IntegerType>(Ty)) {
1493 sz = cast<IntegerType>(Ty)->getBitWidth();
1496 }
else if (isa<PointerType>(Ty))
1506 O <<
"\t.param .b" << sz <<
" ";
1508 O <<
"\t.reg .b" << sz <<
" ";
1509 printParamName(
I, paramIndex,
O);
1514 Type *ETy = PAL.getParamByValType(paramIndex);
1515 assert(ETy &&
"Param should have byval type");
1517 if (isABI || isKernelFunc) {
1522 Align OptimalAlign = getOptimalAlignForParam(ETy);
1535 if (!isKernelFunc && OptimalAlign <
Align(4))
1536 OptimalAlign =
Align(4);
1537 unsigned sz =
DL.getTypeAllocSize(ETy);
1538 O <<
"\t.param .align " << OptimalAlign.
value() <<
" .b8 ";
1539 printParamName(
I, paramIndex,
O);
1540 O <<
"[" << sz <<
"]";
1549 for (
unsigned i = 0,
e = vtparts.size();
i !=
e; ++
i) {
1551 EVT elemtype = vtparts[
i];
1553 elems = vtparts[
i].getVectorNumElements();
1554 elemtype = vtparts[
i].getVectorElementType();
1557 for (
unsigned j = 0, je = elems;
j != je; ++
j) {
1561 O <<
"\t.reg .b" << sz <<
" ";
1562 printParamName(
I, paramIndex,
O);
1578 void NVPTXAsmPrinter::emitFunctionParamList(
const MachineFunction &MF,
1581 emitFunctionParamList(&
F,
O);
1584 void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
1601 O <<
"\t.reg .b64 \t%SP;\n";
1602 O <<
"\t.reg .b64 \t%SPL;\n";
1604 O <<
"\t.reg .b32 \t%SP;\n";
1605 O <<
"\t.reg .b32 \t%SPL;\n";
1614 for (
unsigned i = 0;
i < numVRs;
i++) {
1618 int n = regmap.
size();
1619 regmap.
insert(std::make_pair(vr,
n + 1));
1639 int n = regmap.
size();
1643 O <<
"\t.reg " << rcname <<
" \t" << rcStr <<
"<" << (
n+1)
1654 unsigned int numHex;
1673 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1677 if (
const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
1678 printFPConstant(CFP,
O);
1681 if (isa<ConstantPointerNull>(CPV)) {
1685 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1686 bool IsNonGenericPointer =
false;
1688 IsNonGenericPointer =
true;
1690 if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
1699 if (
const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1701 PointerType *PTy = dyn_cast<PointerType>(Cexpr->getType());
1702 bool IsNonGenericPointer =
false;
1704 IsNonGenericPointer =
true;
1706 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
1707 if (EmitGeneric && !isa<Function>(v) && !IsNonGenericPointer) {
1723 void NVPTXAsmPrinter::bufferLEByte(
const Constant *CPV,
int Bytes,
1724 AggBuffer *AggBuffer) {
1726 int AllocSize =
DL.getTypeAllocSize(CPV->
getType());
1730 AggBuffer->addZeros(Bytes ? Bytes : AllocSize);
1736 size_t NumBytes = (Val.getBitWidth() + 7) / 8;
1738 for (
unsigned I = 0;
I < NumBytes; ++
I) {
1739 Buf[
I] = Val.extractBitsAsZExtValue(8,
I * 8);
1741 AggBuffer->addBytes(Buf.data(), NumBytes, Bytes);
1746 if (
const auto CI = dyn_cast<ConstantInt>(CPV)) {
1750 if (
const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1751 if (
const auto *CI =
1756 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1758 AggBuffer->addSymbol(V, Cexpr->getOperand(0));
1769 AddIntToBuffer(cast<ConstantFP>(CPV)->getValueAPF().bitcastToAPInt());
1773 if (
const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1775 }
else if (
const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1786 if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
1787 bufferAggregateConstant(CPV,
AggBuffer);
1788 if (Bytes > AllocSize)
1790 }
else if (isa<ConstantAggregateZero>(CPV))
1802 void NVPTXAsmPrinter::bufferAggregateConstant(
const Constant *CPV,
1803 AggBuffer *aggBuffer) {
1808 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1810 for (
unsigned I = 0,
E =
DL.getTypeAllocSize(CPV->
getType());
I <
E; ++
I) {
1812 aggBuffer->addBytes(&Byte, 1, 1);
1819 if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
1822 bufferLEByte(cast<Constant>(CPV->
getOperand(
i)), 0, aggBuffer);
1827 dyn_cast<ConstantDataSequential>(CPV)) {
1828 if (CDS->getNumElements())
1829 for (
unsigned i = 0;
i < CDS->getNumElements(); ++
i)
1830 bufferLEByte(cast<Constant>(CDS->getElementAsConstant(
i)), 0,
1835 if (isa<ConstantStruct>(CPV)) {
1840 Bytes =
DL.getStructLayout(
ST)->getElementOffset(0) +
1841 DL.getTypeAllocSize(
ST) -
1842 DL.getStructLayout(
ST)->getElementOffset(
i);
1844 Bytes =
DL.getStructLayout(
ST)->getElementOffset(
i + 1) -
1845 DL.getStructLayout(
ST)->getElementOffset(
i);
1846 bufferLEByte(cast<Constant>(CPV->
getOperand(
i)), Bytes, aggBuffer);
1859 NVPTXAsmPrinter::lowerConstantForGV(
const Constant *CV,
bool ProcessingGeneric) {
1865 if (
const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
1868 if (
const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
1871 if (ProcessingGeneric) {
1883 switch (
CE->getOpcode()) {
1890 return lowerConstantForGV(
C, ProcessingGeneric);
1895 OS <<
"Unsupported expression in static initializer: ";
1896 CE->printAsOperand(OS,
false,
1901 case Instruction::AddrSpaceCast: {
1905 return lowerConstantForGV(cast<const Constant>(
CE->getOperand(0)),
true);
1909 OS <<
"Unsupported expression in static initializer: ";
1910 CE->printAsOperand(OS,
false,
1915 case Instruction::GetElementPtr: {
1919 APInt OffsetAI(
DL.getPointerTypeSizeInBits(
CE->getType()), 0);
1920 cast<GEPOperator>(CE)->accumulateConstantOffset(
DL, OffsetAI);
1922 const MCExpr *
Base = lowerConstantForGV(
CE->getOperand(0),
1927 int64_t
Offset = OffsetAI.getSExtValue();
1932 case Instruction::Trunc:
1938 case Instruction::BitCast:
1939 return lowerConstantForGV(
CE->getOperand(0), ProcessingGeneric);
1941 case Instruction::IntToPtr: {
1949 return lowerConstantForGV(
Op, ProcessingGeneric);
1952 case Instruction::PtrToInt: {
1958 Type *Ty =
CE->getType();
1960 const MCExpr *OpExpr = lowerConstantForGV(
Op, ProcessingGeneric);
1964 if (
DL.getTypeAllocSize(Ty) ==
DL.getTypeAllocSize(
Op->getType()))
1970 unsigned InBits =
DL.getTypeAllocSizeInBits(
Op->getType());
1978 const MCExpr *
LHS = lowerConstantForGV(
CE->getOperand(0), ProcessingGeneric);
1979 const MCExpr *
RHS = lowerConstantForGV(
CE->getOperand(1), ProcessingGeneric);
1980 switch (
CE->getOpcode()) {
1992 return cast<MCTargetExpr>(&Expr)->printImpl(OS,
MAI);
1994 OS << cast<MCConstantExpr>(Expr).getValue();
2020 if (isa<MCConstantExpr>(BE.
getLHS()) || isa<MCSymbolRefExpr>(BE.
getLHS()) ||
2021 isa<NVPTXGenericMCSymbolRefExpr>(BE.
getLHS())) {
2033 if (RHSC->getValue() < 0) {
2034 OS << RHSC->getValue();
2045 if (isa<MCConstantExpr>(BE.
getRHS()) || isa<MCSymbolRefExpr>(BE.
getRHS())) {
2061 bool NVPTXAsmPrinter::PrintAsmOperand(
const MachineInstr *
MI,
unsigned OpNo,
2063 if (ExtraCode && ExtraCode[0]) {
2064 if (ExtraCode[1] != 0)
2067 switch (ExtraCode[0]) {
2081 bool NVPTXAsmPrinter::PrintAsmMemoryOperand(
const MachineInstr *
MI,
2083 const char *ExtraCode,
2085 if (ExtraCode && ExtraCode[0])
2101 if (MO.
getReg() == NVPTX::VRDepot)
2106 emitVirtualRegister(MO.
getReg(),
O);
2135 if (Modifier && strcmp(Modifier,
"add") == 0) {
2139 if (
MI->getOperand(opNum + 1).isImm() &&
2140 MI->getOperand(opNum + 1).getImm() == 0)
bool doFinalization(Module &M) override
Shut down the asmprinter.
This class represents an incoming formal argument to a Function.
LLVM_NODISCARD bool startswith(StringRef Prefix) const
Check if this string starts with the given Prefix.
@ FloatTyID
32-bit floating point type
bool getReqNTIDx(const Function &F, unsigned &x)
@ MO_Immediate
Immediate operand.
@ DoubleTyID
64-bit floating point type
This is an optimization pass for GlobalISel generic memory operations.
MCSymbol - Instances of this class represent a symbol name in the MC file, and MCSymbols are created ...
We currently emits eax Perhaps this is what we really should generate is Is imull three or four cycles eax eax The current instruction priority is based on pattern complexity The former is more complex because it folds a load so the latter will not be emitted Perhaps we should use AddedComplexity to give LEA32r a higher priority We should always try to match LEA first since the LEA matching code does some estimate to determine whether the match is profitable if we care more about code then imull is better It s two bytes shorter than movl leal On a Pentium M
bool hasCommonLinkage() const
A parsed version of the target data layout string in and methods for querying it.
MachineLoop * getLoopFor(const MachineBasicBlock *BB) const
Return the innermost loop that BB lives in.
static const char * getRegisterName(unsigned RegNo)
static MCOperand createExpr(const MCExpr *Val)
Opcode getOpcode() const
Get the kind of this unary expression.
virtual const MCExpr * lowerConstant(const Constant *CV)
Lower the specified LLVM Constant to an MCExpr.
bool hasExternalLinkage() const
bool isPointerTy() const
True if this is an instance of PointerType.
static const fltSemantics & IEEEsingle() LLVM_READNONE
const GlobalValue * getGlobal() const
static MCOperand createImm(int64_t Val)
Context object for machine code objects.
@ VoidTyID
type with no size
static bool is64Bit(const char *name)
const NVPTXTargetLowering * getTargetLowering() const override
A raw_ostream that writes to an std::string.
bool isImageReadWrite(const Value &val)
StringRef getSection() const
Get the custom section of this global if it has one.
const MCAsmInfo * MAI
Target Asm Printer information.
static const MCConstantExpr * create(int64_t Value, MCContext &Ctx, bool PrintInHex=false, unsigned SizeInBytes=0)
bool hasWeakLinkage() const
const APInt & getValue() const
Return the constant as an APInt value reference.
static bool usedInGlobalVarDef(const Constant *C)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
MaybeAlign getAlign() const
Returns the alignment of the given variable or function.
static const NVPTXGenericMCSymbolRefExpr * create(const MCSymbolRefExpr *SymExpr, MCContext &Ctx)
static void DiscoverDependentGlobals(const Value *V, DenseSet< const GlobalVariable * > &Globals)
DiscoverDependentGlobals - Return a set of GlobalVariables on which V depends.
Reg
All possible values of the reg field in the ModR/M byte.
unsigned getAddressSpace() const
Return the address space of the Pointer type.
static bool usedInOneFunc(const User *U, Function const *&oneFunc)
static const MCBinaryExpr * createAnd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx)
TypeID getTypeID() const
Return the type id for the type.
unsigned int getSmVersion() const
bool doFinalization(Module &M) override
Shut down the asmprinter.
virtual const TargetRegisterInfo * getRegisterInfo() const
getRegisterInfo - If register information is available, return it.
bool hasDebugInfo() const
Returns true if valid debug info is present.
TargetRegisterInfo base class - We assume that the target defines a static array of TargetRegisterDes...
bool getMinCTASm(const Function &F, unsigned &x)
The instances of the Type class are immutable: once they are created, they are never changed.
virtual void PrintSymbolOperand(const MachineOperand &MO, raw_ostream &OS)
Print the MachineOperand as a symbol.
static const NVPTXFloatMCExpr * createConstantFPDouble(const APFloat &Flt, MCContext &Ctx)
static bool printOperand(raw_ostream &OS, const SelectionDAG *G, const SDValue Value)
bool isImage(const Value &val)
const APFloat & getValueAPF() const
to esp esp setne al movzbw ax esp setg cl movzbw cx cmove cx cl jne LBB1_2 esp ret(also really horrible code on ppc). This is due to the expand code for 64-bit compares. GCC produces multiple branches
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
static Register index2VirtReg(unsigned Index)
Convert a 0-based index to a virtual register number.
Instances of this class represent a single low-level machine instruction.
Binary assembler expressions.
int caller(int32 arg1, int32 arg2)
bool isFloatingPointTy() const
Return true if this is one of the six floating-point types.
FormattedNumber format_hex_no_prefix(uint64_t N, unsigned Width, bool Upper=false)
format_hex_no_prefix - Output N as a fixed width hexadecimal.
std::pair< iterator, bool > insert(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.
unsigned const TargetRegisterInfo * TRI
Unary assembler expressions.
bool isTexture(const Value &val)
@ MO_Register
Register operand.
bool hasInitializer() const
Definitions have initializers, declarations don't.
void EmitToStreamer(MCStreamer &S, const MCInst &Inst)
bool isSampler(const Value &val)
LLVM Basic Block Representation.
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...
void setOpcode(unsigned Op)
MCSymbol * getSymbol() const
Return the MCSymbol for this basic block.
static const NVPTXFloatMCExpr * createConstantFPSingle(const APFloat &Flt, MCContext &Ctx)
bool runOnMachineFunction(MachineFunction &F) override
Emit the specified function out to the OutStreamer.
This is the shared class of boolean and integer constants.
< i1 > br i1 label label bb bb
bool getReqNTIDy(const Function &F, unsigned &y)
static void printMCExpr(const MCExpr *E, raw_ostream &OS)
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
MCSymbol * getOrCreateSymbol(const Twine &Name)
Lookup the symbol inside with the specified Name.
ConstantDataSequential - A vector or array constant whose element type is a simple 1/2/4/8-byte integ...
bool isDeclaration() const
Return true if the primary definition of this global value is outside of the current translation unit...
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...
void lshrInPlace(unsigned ShiftAmt)
Logical right-shift this APInt by ShiftAmt in place.
bool hasAppendingLinkage() const
bool isNullValue() const
Return true if this is the value that would be returned by getNullValue.
Itanium Name Demangler i e convert the string _Z1fv into f()". You can also use the CRTP base ManglingParser to perform some simple analysis on the mangled name
@ MO_GlobalAddress
Address of a global value.
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
Ty * getInfo()
getInfo - Keep track of various per-function pieces of information for backends that would like to do...
@ Target
Target specific expression.
std::unique_ptr< MCStreamer > OutStreamer
This is the MCStreamer object for the file we are generating.
(vector float) vec_cmpeq(*A, *B) C
Clang compiles this i1 i64 store i64 i64 store i64 i64 store i64 i64 store i64 align Which gets codegen d xmm0 movaps rbp movaps rbp movaps rbp movaps rbp rbp rbp rbp rbp It would be better to have movq s of instead of the movaps s LLVM produces ret int
static GCMetadataPrinterRegistry::Add< OcamlGCMetadataPrinter > Y("ocaml", "ocaml 3.10-compatible collector")
static bool isPhysicalRegister(unsigned Reg)
Return true if the specified register number is in the physical register namespace.
bool hasPrivateLinkage() const
unsigned getPTXVersion() const
bool isVectorTy() const
True if this is an instance of VectorType.
const MCExpr * getRHS() const
Get the right-hand side expression of the binary operator.
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.
Describe properties that are true of each instruction in the target description file.
std::string getSamplerName(const Value &val)
MachineOperand class - Representation of each machine instruction operand.
ConstantArray - Constant Array Declarations.
ManagedStringPool * getManagedStrPool() const
bool getReqNTIDz(const Function &F, unsigned &z)
bool isInteger() const
Return true if this is an integer or a vector integer type.
unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
const NVPTXRegisterInfo * getRegisterInfo() const override
MCSymbol * CurrentFnSym
The symbol for the current function.
void report_fatal_error(Error Err, bool gen_crash_diag=true)
Report a serious error, calling any installed error handler.
uint64_t getZExtValue() const
Get zero extended value.
This class implements an extremely fast bulk output stream that can only output to a stream.
ConstantFP - Floating Point Values [float, double].
const MCSymbol & getSymbol() const
uint64_t getStackSize() const
Return the number of bytes that must be allocated to hold all of the fixed size frame objects.
bool hasInternalLinkage() const
APInt bitcastToAPInt() const
This struct is a compact representation of a valid (non-zero power of two) alignment.
void print(raw_ostream &OS, const MCAsmInfo *MAI) const
print - Print the value to the stream OS.
virtual void emitBasicBlockStart(const MachineBasicBlock &MBB)
Targets can override this to emit stuff at the start of a basic block.
bool isLoopHeader(const MachineBasicBlock *BB) const
True if the block is a loop header node.
@ Binary
Binary expressions.
static GCMetadataPrinterRegistry::Add< ErlangGCPrinter > X("erlang", "erlang-compatible garbage collector")
std::string getNVPTXRegClassStr(TargetRegisterClass const *RC)
void addOperand(const MCOperand Op)
std::string * getManagedString(const char *S)
Target & getTheNVPTXTarget64()
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
static const NVPTXFloatMCExpr * createConstantFPHalf(const APFloat &Flt, MCContext &Ctx)
bool isIntegerTy() const
True if this is an instance of IntegerType.
bool isKernelFunction(const Function &F)
Implements a dense probed hash-table based set.
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
const char * getName(unsigned RegNo) const
bool hasSection() const
Check if this global has a custom object file section.
bool hasAvailableExternallyLinkage() const
LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXAsmPrinter()
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasImageHandles() const
This is an important base class in LLVM.
TypeSize getSizeInBits() const
Return the size of the specified value type in bits.
const TargetRegisterClass * getRegClass(unsigned i) const
Returns the register class associated with the enumeration value.
Representation of each machine instruction.
Module * getParent()
Get the module that this global value is contained inside of...
#define LLVM_EXTERNAL_VISIBILITY
bool isIntOrPtrTy() const
Return true if this is an integer type or a pointer type.
bool doInitialization(Module &M) override
Set up the AsmPrinter when we are working on a new module.
static bool isEmptyXXStructor(GlobalVariable *GV)
Class to represent pointers.
static bool useFuncSeen(const Constant *C, DenseMap< const Function *, bool > &seenMap)
MachineOperandType getType() const
getType - Returns the MachineOperandType for this operand.
static void printMemOperand(raw_ostream &OS, const MachineMemOperand &MMO, const MachineFunction *MF, const Module *M, const MachineFrameInfo *MFI, const TargetInstrInfo *TII, LLVMContext &Ctx)
bool isHalfTy() const
Return true if this is 'half', a 16-bit IEEE fp type.
const ConstantFP * getFPImm() const
static bool isVirtualRegister(unsigned Reg)
Return true if the specified register number is in the virtual register namespace.
iterator find(const_arg_type_t< KeyT > Val)
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
MCSymbol * getSymbol(const GlobalValue *GV) const
TypeSize getSizeInBits() const
Returns the size of the specified MVT in bits.
const char * getImageHandleSymbol(unsigned Idx) const
Returns the symbol name at the given index.
MachineFrameInfo & getFrameInfo()
getFrameInfo - Return the frame info object for the current function.
MachineFunction * MF
The current machine function.
MCContext & OutContext
This is the context for the output file that we are streaming.
@ CE
Windows NT (Windows on ARM)
@ MO_FPImmediate
Floating-point immediate operand.
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,...
Register getReg() const
getReg - Returns the register number.
A Module instance is used to store all the information related to an LLVM module.
iterator_range< pred_iterator > predecessors()
Represent a reference to a symbol from inside an expression.
Class for arbitrary precision integers.
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.
bool runOnMachineFunction(MachineFunction &MF) override
Emit the specified function out to the OutStreamer.
static const MCBinaryExpr * createAdd(const MCExpr *LHS, const MCExpr *RHS, MCContext &Ctx)
std::string getVirtualRegisterName(unsigned) const
MCSymbol * GetExternalSymbolSymbol(StringRef Sym) const
Return the MCSymbol for the specified ExternalSymbol.
const MCExpr * getSubExpr() const
Get the child of this unary expression.
MachineBasicBlock * getMBB() const
Class to represent struct types.
static const fltSemantics & IEEEdouble() LLVM_READNONE
StringRef - Represent a constant reference to a string, i.e.
static MCOperand GetSymbolRef(const MachineOperand &MO, const MCSymbol *Symbol, HexagonAsmPrinter &Printer, bool MustExtend)
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
Type * getType() const
All values are typed, get the type of this value.
static MCOperand createReg(unsigned Reg)
@ MO_MachineBasicBlock
MachineBasicBlock reference.
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
add sub stmia L5 ldr r0 bl L_printf $stub Instead of a and a wouldn t it be better to do three moves *Return an aggregate type is even return S
static Constant * getIntegerCast(Constant *C, Type *Ty, bool IsSigned)
Create a ZExt, Bitcast or Trunc for integer -> integer casts.
#define LLVM_FALLTHROUGH
LLVM_FALLTHROUGH - Mark fallthrough cases in switch statements.
StringRef getName() const
Return a constant reference to the value's name.
std::pair< iterator, bool > insert(const std::pair< KeyT, ValueT > &KV)
unsigned const MachineRegisterInfo * MRI
Align getMaxAlign() const
Return the alignment in bytes that this function must be aligned to, which is greater than the defaul...
const Value * stripPointerCasts() const
Strip off pointer casts, all-zero GEPs and address space casts.
Wrapper class representing virtual and physical registers.
MachineModuleInfo * MMI
This is a pointer to the current MachineModuleInfo.
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
@ IntegerTyID
Arbitrary bit width integers.
static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f)
Implments NVPTX-specific streamer.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
Opcode getOpcode() const
Get the kind of this binary expression.
A constant value that is initialized with an expression using other constant values.
Function & getFunction()
Return the LLVM function that this machine code represents.
bool getMaxNTIDz(const Function &F, unsigned &z)
const LLVMTargetMachine & getTarget() const
getTarget - Return the target machine this machine code is compiled with
uint64_t value() const
This is a hole in the type system and should not be abused.
@ FixedVectorTyID
Fixed width SIMD vector type.
@ MO_ExternalSymbol
Name of external global symbol.
bool isSurface(const Value &val)
const NVPTXSubtarget * getSubtargetImpl(const Function &) const override
Virtual method implemented by subclasses that returns a reference to that target's TargetSubtargetInf...
APInt getLoBits(unsigned numBits) const
Compute an APInt containing numBits lowbits from this APInt.
This class describes a target machine that is implemented with the LLVM target-independent code gener...
static const MCSymbolRefExpr * create(const MCSymbol *Symbol, MCContext &Ctx)
unsigned getNumRegClasses() const
bool isImm() const
isImm - Tests if this is a MO_Immediate operand.
bool getAlign(const Function &F, unsigned index, unsigned &align)
std::string getSurfaceName(const Value &val)
bool getMaxNTIDy(const Function &F, unsigned &y)
bool isImageWriteOnly(const Value &val)
LegalityPredicate isVector(unsigned TypeIdx)
True iff the specified type index is a vector.
The MachineFrameInfo class represents an abstract stack frame until prolog/epilog code is inserted.
static constexpr roundingMode rmNearestTiesToEven
void clearAnnotationCache(const Module *Mod)
void print(raw_ostream &OS, const MCAsmInfo *MAI, bool InParens=false) const
opStatus convert(const fltSemantics &ToSemantics, roundingMode RM, bool *losesInfo)
unsigned getNumOperands() const
TargetMachine & TM
Target machine description.
std::string getTargetName() const
@ SymbolRef
References to labels and assigned expressions.
const char * getSymbolName() const
const STC & getSubtarget(const Function &F) const
This method returns a pointer to the specified type of TargetSubtargetInfo.
@ Unary
Unary expressions.
const BasicBlock * getParent() const
Align max(MaybeAlign Lhs, Align Rhs)
@ Constant
Constant expressions.
A raw_ostream that writes to an SmallVector or SmallString.
const DataLayout & getDataLayout() const
Return information about data layout.
bool erase(const ValueT &V)
const LLVM_NODISCARD char * data() const
data - Get a pointer to the start of the string (which may not be null terminated).
bool hasLinkOnceLinkage() const
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
Constant * ConstantFoldConstant(const Constant *C, const DataLayout &DL, const TargetLibraryInfo *TLI=nullptr)
ConstantFoldConstant - Fold the constant using the specified DataLayout.
PointerType * getType() const
Global values are always pointers.
Type * getValueType() const
Instances of this class represent operands of the MCInst class.
const MCSymbol * getFunctionFrameSymbol() const override
Return symbol for the function pseudo stack if the stack frame is not a register based.
void emitInitialRawDwarfLocDirective(const MachineFunction &MF)
Emits inital debug location directive.
bool isAggregateType() const
Return true if the type is an aggregate type.
NVPTX::DrvInterface getDrvInterface() const
RegisterAsmPrinter - Helper template for registering a target specific assembly printer,...
Value * getOperand(unsigned i) const
bool getMaxNTIDx(const Function &F, unsigned &x)
std::string getTextureName(const Value &val)
bool getMaxNReg(const Function &F, unsigned &x)
The same transformation can work with an even modulo with the addition of a and shrink the compare RHS by the same amount Unless the target supports that transformation probably isn t worthwhile The transformation can also easily be made to work with non zero equality for n
std::string getNVPTXRegClassName(TargetRegisterClass const *RC)
bool isManaged(const Value &val)
unsigned getFunctionNumber() const
Return a unique ID for the current function.
LLVM Value Representation.
Base class for the full range of assembler expressions which are needed for parsing.
@ HalfTyID
16-bit floating point type
iterator_range< user_iterator > users()
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.
TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
const MCExpr * getLHS() const
Get the left-hand side expression of the binary operator.
bool doInitialization(Module &M) override
Set up the AsmPrinter when we are working on a new module.
Target & getTheNVPTXTarget32()