1 //===-- NVPTXAsmPrinter.cpp - NVPTX LLVM assembly writer ------------------===//
3 // The LLVM Compiler Infrastructure
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
8 //===----------------------------------------------------------------------===//
10 // This file contains a printer that converts from our internal representation
11 // of machine-dependent LLVM code to NVPTX assembly language.
13 //===----------------------------------------------------------------------===//
15 #include "NVPTXAsmPrinter.h"
17 #include "NVPTXInstrInfo.h"
18 #include "NVPTXTargetMachine.h"
19 #include "NVPTXRegisterInfo.h"
20 #include "NVPTXUtilities.h"
21 #include "MCTargetDesc/NVPTXMCAsmInfo.h"
22 #include "NVPTXNumRegisters.h"
23 #include "../lib/CodeGen/AsmPrinter/DwarfDebug.h" // FIXME: layering violation!
24 #include "llvm/ADT/StringExtras.h"
25 #include "llvm/DebugInfo.h"
26 #include "llvm/Function.h"
27 #include "llvm/GlobalVariable.h"
28 #include "llvm/Module.h"
29 #include "llvm/CodeGen/Analysis.h"
30 #include "llvm/CodeGen/MachineRegisterInfo.h"
31 #include "llvm/CodeGen/MachineFrameInfo.h"
32 #include "llvm/CodeGen/MachineModuleInfo.h"
33 #include "llvm/MC/MCStreamer.h"
34 #include "llvm/MC/MCSymbol.h"
35 #include "llvm/Target/Mangler.h"
36 #include "llvm/Target/TargetLoweringObjectFile.h"
37 #include "llvm/Support/TargetRegistry.h"
38 #include "llvm/Support/ErrorHandling.h"
39 #include "llvm/Support/FormattedStream.h"
40 #include "llvm/DerivedTypes.h"
41 #include "llvm/Support/TimeValue.h"
42 #include "llvm/Support/CommandLine.h"
43 #include "llvm/Analysis/ConstantFolding.h"
44 #include "llvm/Support/Path.h"
45 #include "llvm/Assembly/Writer.h"
46 #include "cl_common_defines.h"
51 #include "NVPTXGenAsmWriter.inc"
53 bool RegAllocNilUsed = true;
55 #define DEPOTNAME "__local_depot"
58 EmitLineNumbers("nvptx-emit-line-numbers",
59 cl::desc("NVPTX Specific: Emit Line numbers even without -G"),
63 bool InterleaveSrcInPtx = false;
66 static cl::opt<bool, true>InterleaveSrc("nvptx-emit-src",
68 cl::desc("NVPTX Specific: Emit source line in ptx file"),
69 cl::location(llvm::InterleaveSrcInPtx));
74 // @TODO: This is a copy from AsmPrinter.cpp. The function is static, so we
75 // cannot just link to the existing version.
76 /// LowerConstant - Lower the specified LLVM Constant to an MCExpr.
78 using namespace nvptx;
79 const MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) {
80 MCContext &Ctx = AP.OutContext;
82 if (CV->isNullValue() || isa<UndefValue>(CV))
83 return MCConstantExpr::Create(0, Ctx);
85 if (const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
86 return MCConstantExpr::Create(CI->getZExtValue(), Ctx);
88 if (const GlobalValue *GV = dyn_cast<GlobalValue>(CV))
89 return MCSymbolRefExpr::Create(AP.Mang->getSymbol(GV), Ctx);
91 if (const BlockAddress *BA = dyn_cast<BlockAddress>(CV))
92 return MCSymbolRefExpr::Create(AP.GetBlockAddressSymbol(BA), Ctx);
94 const ConstantExpr *CE = dyn_cast<ConstantExpr>(CV);
96 llvm_unreachable("Unknown constant value to lower!");
99 switch (CE->getOpcode()) {
101 // If the code isn't optimized, there may be outstanding folding
102 // opportunities. Attempt to fold the expression using TargetData as a
103 // last resort before giving up.
105 ConstantFoldConstantExpression(CE, AP.TM.getTargetData()))
107 return LowerConstant(C, AP);
109 // Otherwise report the problem to the user.
112 raw_string_ostream OS(S);
113 OS << "Unsupported expression in static initializer: ";
114 WriteAsOperand(OS, CE, /*PrintType=*/false,
115 !AP.MF ? 0 : AP.MF->getFunction()->getParent());
116 report_fatal_error(OS.str());
118 case Instruction::GetElementPtr: {
119 const TargetData &TD = *AP.TM.getTargetData();
120 // Generate a symbolic expression for the byte address
121 const Constant *PtrVal = CE->getOperand(0);
122 SmallVector<Value*, 8> IdxVec(CE->op_begin()+1, CE->op_end());
123 int64_t Offset = TD.getIndexedOffset(PtrVal->getType(), IdxVec);
125 const MCExpr *Base = LowerConstant(CE->getOperand(0), AP);
129 // Truncate/sext the offset to the pointer size.
130 if (TD.getPointerSizeInBits() != 64) {
131 int SExtAmount = 64-TD.getPointerSizeInBits();
132 Offset = (Offset << SExtAmount) >> SExtAmount;
135 return MCBinaryExpr::CreateAdd(Base, MCConstantExpr::Create(Offset, Ctx),
139 case Instruction::Trunc:
140 // We emit the value and depend on the assembler to truncate the generated
141 // expression properly. This is important for differences between
142 // blockaddress labels. Since the two labels are in the same function, it
143 // is reasonable to treat their delta as a 32-bit value.
145 case Instruction::BitCast:
146 return LowerConstant(CE->getOperand(0), AP);
148 case Instruction::IntToPtr: {
149 const TargetData &TD = *AP.TM.getTargetData();
150 // Handle casts to pointers by changing them into casts to the appropriate
151 // integer type. This promotes constant folding and simplifies this code.
152 Constant *Op = CE->getOperand(0);
153 Op = ConstantExpr::getIntegerCast(Op, TD.getIntPtrType(CV->getContext()),
155 return LowerConstant(Op, AP);
158 case Instruction::PtrToInt: {
159 const TargetData &TD = *AP.TM.getTargetData();
160 // Support only foldable casts to/from pointers that can be eliminated by
161 // changing the pointer to the appropriately sized integer type.
162 Constant *Op = CE->getOperand(0);
163 Type *Ty = CE->getType();
165 const MCExpr *OpExpr = LowerConstant(Op, AP);
167 // We can emit the pointer value into this slot if the slot is an
168 // integer slot equal to the size of the pointer.
169 if (TD.getTypeAllocSize(Ty) == TD.getTypeAllocSize(Op->getType()))
172 // Otherwise the pointer is smaller than the resultant integer, mask off
173 // the high bits so we are sure to get a proper truncation if the input is
175 unsigned InBits = TD.getTypeAllocSizeInBits(Op->getType());
176 const MCExpr *MaskExpr = MCConstantExpr::Create(~0ULL >> (64-InBits), Ctx);
177 return MCBinaryExpr::CreateAnd(OpExpr, MaskExpr, Ctx);
180 // The MC library also has a right-shift operator, but it isn't consistently
181 // signed or unsigned between different targets.
182 case Instruction::Add:
183 case Instruction::Sub:
184 case Instruction::Mul:
185 case Instruction::SDiv:
186 case Instruction::SRem:
187 case Instruction::Shl:
188 case Instruction::And:
189 case Instruction::Or:
190 case Instruction::Xor: {
191 const MCExpr *LHS = LowerConstant(CE->getOperand(0), AP);
192 const MCExpr *RHS = LowerConstant(CE->getOperand(1), AP);
193 switch (CE->getOpcode()) {
194 default: llvm_unreachable("Unknown binary operator constant cast expr");
195 case Instruction::Add: return MCBinaryExpr::CreateAdd(LHS, RHS, Ctx);
196 case Instruction::Sub: return MCBinaryExpr::CreateSub(LHS, RHS, Ctx);
197 case Instruction::Mul: return MCBinaryExpr::CreateMul(LHS, RHS, Ctx);
198 case Instruction::SDiv: return MCBinaryExpr::CreateDiv(LHS, RHS, Ctx);
199 case Instruction::SRem: return MCBinaryExpr::CreateMod(LHS, RHS, Ctx);
200 case Instruction::Shl: return MCBinaryExpr::CreateShl(LHS, RHS, Ctx);
201 case Instruction::And: return MCBinaryExpr::CreateAnd(LHS, RHS, Ctx);
202 case Instruction::Or: return MCBinaryExpr::CreateOr (LHS, RHS, Ctx);
203 case Instruction::Xor: return MCBinaryExpr::CreateXor(LHS, RHS, Ctx);
210 void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI)
212 if (!EmitLineNumbers)
217 DebugLoc curLoc = MI.getDebugLoc();
219 if (prevDebugLoc.isUnknown() && curLoc.isUnknown())
222 if (prevDebugLoc == curLoc)
225 prevDebugLoc = curLoc;
227 if (curLoc.isUnknown())
231 const MachineFunction *MF = MI.getParent()->getParent();
232 //const TargetMachine &TM = MF->getTarget();
234 const LLVMContext &ctx = MF->getFunction()->getContext();
235 DIScope Scope(curLoc.getScope(ctx));
240 StringRef fileName(Scope.getFilename());
241 StringRef dirName(Scope.getDirectory());
242 SmallString<128> FullPathName = dirName;
243 if (!dirName.empty() && !sys::path::is_absolute(fileName)) {
244 sys::path::append(FullPathName, fileName);
245 fileName = FullPathName.str();
248 if (filenameMap.find(fileName.str()) == filenameMap.end())
252 // Emit the line from the source file.
253 if (llvm::InterleaveSrcInPtx)
254 this->emitSrcInText(fileName.str(), curLoc.getLine());
256 std::stringstream temp;
257 temp << "\t.loc " << filenameMap[fileName.str()]
258 << " " << curLoc.getLine() << " " << curLoc.getCol();
259 OutStreamer.EmitRawText(Twine(temp.str().c_str()));
262 void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) {
263 SmallString<128> Str;
264 raw_svector_ostream OS(Str);
265 if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)
266 emitLineNumberAsDotLoc(*MI);
267 printInstruction(MI, OS);
268 OutStreamer.EmitRawText(OS.str());
271 void NVPTXAsmPrinter::printReturnValStr(const Function *F,
274 const TargetData *TD = TM.getTargetData();
275 const TargetLowering *TLI = TM.getTargetLowering();
277 Type *Ty = F->getReturnType();
279 bool isABI = (nvptxSubtarget.getSmVersion() >= 20);
281 if (Ty->getTypeID() == Type::VoidTyID)
287 if (Ty->isPrimitiveType() || Ty->isIntegerTy()) {
289 if (const IntegerType *ITy = dyn_cast<IntegerType>(Ty)) {
290 size = ITy->getBitWidth();
291 if (size < 32) size = 32;
293 assert(Ty->isFloatingPointTy() &&
294 "Floating point type expected here");
295 size = Ty->getPrimitiveSizeInBits();
298 O << ".param .b" << size << " func_retval0";
300 else if (isa<PointerType>(Ty)) {
301 O << ".param .b" << TLI->getPointerTy().getSizeInBits()
304 if ((Ty->getTypeID() == Type::StructTyID) ||
305 isa<VectorType>(Ty)) {
306 SmallVector<EVT, 16> vtparts;
307 ComputeValueVTs(*TLI, Ty, vtparts);
308 unsigned totalsz = 0;
309 for (unsigned i=0,e=vtparts.size(); i!=e; ++i) {
311 EVT elemtype = vtparts[i];
312 if (vtparts[i].isVector()) {
313 elems = vtparts[i].getVectorNumElements();
314 elemtype = vtparts[i].getVectorElementType();
316 for (unsigned j=0, je=elems; j!=je; ++j) {
317 unsigned sz = elemtype.getSizeInBits();
318 if (elemtype.isInteger() && (sz < 8)) sz = 8;
322 unsigned retAlignment = 0;
323 if (!llvm::getAlign(*F, 0, retAlignment))
324 retAlignment = TD->getABITypeAlignment(Ty);
325 O << ".param .align "
327 << " .b8 func_retval0["
331 "Unknown return type");
334 SmallVector<EVT, 16> vtparts;
335 ComputeValueVTs(*TLI, Ty, vtparts);
337 for (unsigned i=0,e=vtparts.size(); i!=e; ++i) {
339 EVT elemtype = vtparts[i];
340 if (vtparts[i].isVector()) {
341 elems = vtparts[i].getVectorNumElements();
342 elemtype = vtparts[i].getVectorElementType();
345 for (unsigned j=0, je=elems; j!=je; ++j) {
346 unsigned sz = elemtype.getSizeInBits();
347 if (elemtype.isInteger() && (sz < 32)) sz = 32;
348 O << ".reg .b" << sz << " func_retval" << idx;
349 if (j<je-1) O << ", ";
360 void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
362 const Function *F = MF.getFunction();
363 printReturnValStr(F, O);
366 void NVPTXAsmPrinter::EmitFunctionEntryLabel() {
367 SmallString<128> Str;
368 raw_svector_ostream O(Str);
371 MRI = &MF->getRegInfo();
372 F = MF->getFunction();
373 emitLinkageDirective(F,O);
374 if (llvm::isKernelFunction(*F))
378 printReturnValStr(*MF, O);
383 emitFunctionParamList(*MF, O);
385 if (llvm::isKernelFunction(*F))
386 emitKernelFunctionDirectives(*F, O);
388 OutStreamer.EmitRawText(O.str());
390 prevDebugLoc = DebugLoc();
393 void NVPTXAsmPrinter::EmitFunctionBodyStart() {
394 const TargetRegisterInfo &TRI = *TM.getRegisterInfo();
395 unsigned numRegClasses = TRI.getNumRegClasses();
396 VRidGlobal2LocalMap = new std::map<unsigned, unsigned>[numRegClasses+1];
397 OutStreamer.EmitRawText(StringRef("{\n"));
398 setAndEmitFunctionVirtualRegisters(*MF);
400 SmallString<128> Str;
401 raw_svector_ostream O(Str);
402 emitDemotedVars(MF->getFunction(), O);
403 OutStreamer.EmitRawText(O.str());
406 void NVPTXAsmPrinter::EmitFunctionBodyEnd() {
407 OutStreamer.EmitRawText(StringRef("}\n"));
408 delete []VRidGlobal2LocalMap;
413 NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function& F,
414 raw_ostream &O) const {
415 // If the NVVM IR has some of reqntid* specified, then output
416 // the reqntid directive, and set the unspecified ones to 1.
417 // If none of reqntid* is specified, don't output reqntid directive.
418 unsigned reqntidx, reqntidy, reqntidz;
419 bool specified = false;
420 if (llvm::getReqNTIDx(F, reqntidx) == false) reqntidx = 1;
421 else specified = true;
422 if (llvm::getReqNTIDy(F, reqntidy) == false) reqntidy = 1;
423 else specified = true;
424 if (llvm::getReqNTIDz(F, reqntidz) == false) reqntidz = 1;
425 else specified = true;
428 O << ".reqntid " << reqntidx << ", "
429 << reqntidy << ", " << reqntidz << "\n";
431 // If the NVVM IR has some of maxntid* specified, then output
432 // the maxntid directive, and set the unspecified ones to 1.
433 // If none of maxntid* is specified, don't output maxntid directive.
434 unsigned maxntidx, maxntidy, maxntidz;
436 if (llvm::getMaxNTIDx(F, maxntidx) == false) maxntidx = 1;
437 else specified = true;
438 if (llvm::getMaxNTIDy(F, maxntidy) == false) maxntidy = 1;
439 else specified = true;
440 if (llvm::getMaxNTIDz(F, maxntidz) == false) maxntidz = 1;
441 else specified = true;
444 O << ".maxntid " << maxntidx << ", "
445 << maxntidy << ", " << maxntidz << "\n";
448 if (llvm::getMinCTASm(F, mincta))
449 O << ".minnctapersm " << mincta << "\n";
453 NVPTXAsmPrinter::getVirtualRegisterName(unsigned vr, bool isVec,
455 const TargetRegisterClass * RC = MRI->getRegClass(vr);
456 unsigned id = RC->getID();
458 std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[id];
459 unsigned mapped_vr = regmap[vr];
462 O << getNVPTXRegClassStr(RC) << mapped_vr;
465 // Vector virtual register
466 if (getNVPTXVectorSize(RC) == 4)
468 << getNVPTXRegClassStr(RC) << mapped_vr << "_0, "
469 << getNVPTXRegClassStr(RC) << mapped_vr << "_1, "
470 << getNVPTXRegClassStr(RC) << mapped_vr << "_2, "
471 << getNVPTXRegClassStr(RC) << mapped_vr << "_3"
473 else if (getNVPTXVectorSize(RC) == 2)
475 << getNVPTXRegClassStr(RC) << mapped_vr << "_0, "
476 << getNVPTXRegClassStr(RC) << mapped_vr << "_1"
479 llvm_unreachable("Unsupported vector size");
483 NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr, bool isVec,
485 getVirtualRegisterName(vr, isVec, O);
488 void NVPTXAsmPrinter::printVecModifiedImmediate(const MachineOperand &MO,
489 const char *Modifier,
491 static const char vecelem[] = {'0', '1', '2', '3', '0', '1', '2', '3'};
492 int Imm = (int)MO.getImm();
493 if(0 == strcmp(Modifier, "vecelem"))
494 O << "_" << vecelem[Imm];
495 else if(0 == strcmp(Modifier, "vecv4comm1")) {
496 if((Imm < 0) || (Imm > 3))
499 else if(0 == strcmp(Modifier, "vecv4comm2")) {
500 if((Imm < 4) || (Imm > 7))
503 else if(0 == strcmp(Modifier, "vecv4pos")) {
505 O << "_" << vecelem[Imm%4];
507 else if(0 == strcmp(Modifier, "vecv2comm1")) {
508 if((Imm < 0) || (Imm > 1))
511 else if(0 == strcmp(Modifier, "vecv2comm2")) {
512 if((Imm < 2) || (Imm > 3))
515 else if(0 == strcmp(Modifier, "vecv2pos")) {
517 O << "_" << vecelem[Imm%2];
520 llvm_unreachable("Unknown Modifier on immediate operand");
523 void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum,
524 raw_ostream &O, const char *Modifier) {
525 const MachineOperand &MO = MI->getOperand(opNum);
526 switch (MO.getType()) {
527 case MachineOperand::MO_Register:
528 if (TargetRegisterInfo::isPhysicalRegister(MO.getReg())) {
529 if (MO.getReg() == NVPTX::VRDepot)
530 O << DEPOTNAME << getFunctionNumber();
532 O << getRegisterName(MO.getReg());
535 emitVirtualRegister(MO.getReg(), false, O);
537 if (strcmp(Modifier, "vecfull") == 0)
538 emitVirtualRegister(MO.getReg(), true, O);
541 "Don't know how to handle the modifier on virtual register.");
546 case MachineOperand::MO_Immediate:
549 else if (strstr(Modifier, "vec") == Modifier)
550 printVecModifiedImmediate(MO, Modifier, O);
552 llvm_unreachable("Don't know how to handle modifier on immediate operand");
555 case MachineOperand::MO_FPImmediate:
556 printFPConstant(MO.getFPImm(), O);
559 case MachineOperand::MO_GlobalAddress:
560 O << *Mang->getSymbol(MO.getGlobal());
563 case MachineOperand::MO_ExternalSymbol: {
564 const char * symbname = MO.getSymbolName();
565 if (strstr(symbname, ".PARAM") == symbname) {
567 sscanf(symbname+6, "%u[];", &index);
568 printParamName(index, O);
570 else if (strstr(symbname, ".HLPPARAM") == symbname) {
572 sscanf(symbname+9, "%u[];", &index);
573 O << *CurrentFnSym << "_param_" << index << "_offset";
580 case MachineOperand::MO_MachineBasicBlock:
581 O << *MO.getMBB()->getSymbol();
585 llvm_unreachable("Operand type not supported.");
589 void NVPTXAsmPrinter::
590 printImplicitDef(const MachineInstr *MI, raw_ostream &O) const {
592 O << "\t// Implicit def :";
593 //printOperand(MI, 0);
598 void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum,
599 raw_ostream &O, const char *Modifier) {
600 printOperand(MI, opNum, O);
602 if (Modifier && !strcmp(Modifier, "add")) {
604 printOperand(MI, opNum+1, O);
606 if (MI->getOperand(opNum+1).isImm() &&
607 MI->getOperand(opNum+1).getImm() == 0)
608 return; // don't print ',0' or '+0'
610 printOperand(MI, opNum+1, O);
614 void NVPTXAsmPrinter::printLdStCode(const MachineInstr *MI, int opNum,
615 raw_ostream &O, const char *Modifier)
618 const MachineOperand &MO = MI->getOperand(opNum);
619 int Imm = (int)MO.getImm();
620 if (!strcmp(Modifier, "volatile")) {
623 } else if (!strcmp(Modifier, "addsp")) {
625 case NVPTX::PTXLdStInstCode::GLOBAL: O << ".global"; break;
626 case NVPTX::PTXLdStInstCode::SHARED: O << ".shared"; break;
627 case NVPTX::PTXLdStInstCode::LOCAL: O << ".local"; break;
628 case NVPTX::PTXLdStInstCode::PARAM: O << ".param"; break;
629 case NVPTX::PTXLdStInstCode::CONSTANT: O << ".const"; break;
630 case NVPTX::PTXLdStInstCode::GENERIC:
631 if (!nvptxSubtarget.hasGenericLdSt())
635 assert("wrong value");
638 else if (!strcmp(Modifier, "sign")) {
639 if (Imm==NVPTX::PTXLdStInstCode::Signed)
641 else if (Imm==NVPTX::PTXLdStInstCode::Unsigned)
646 else if (!strcmp(Modifier, "vec")) {
647 if (Imm==NVPTX::PTXLdStInstCode::V2)
649 else if (Imm==NVPTX::PTXLdStInstCode::V4)
653 assert("unknown modifier");
656 assert("unknown modifier");
659 void NVPTXAsmPrinter::emitDeclaration (const Function *F, raw_ostream &O) {
661 emitLinkageDirective(F,O);
662 if (llvm::isKernelFunction(*F))
666 printReturnValStr(F, O);
667 O << *CurrentFnSym << "\n";
668 emitFunctionParamList(F, O);
672 static bool usedInGlobalVarDef(const Constant *C)
677 if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(C)) {
678 if (GV->getName().str() == "llvm.used")
683 for (Value::const_use_iterator ui=C->use_begin(), ue=C->use_end();
685 const Constant *C = dyn_cast<Constant>(*ui);
686 if (usedInGlobalVarDef(C))
692 static bool usedInOneFunc(const User *U, Function const *&oneFunc)
694 if (const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
695 if (othergv->getName().str() == "llvm.used")
699 if (const Instruction *instr = dyn_cast<Instruction>(U)) {
700 if (instr->getParent() && instr->getParent()->getParent()) {
701 const Function *curFunc = instr->getParent()->getParent();
702 if (oneFunc && (curFunc != oneFunc))
711 if (const MDNode *md = dyn_cast<MDNode>(U))
712 if (md->hasName() && ((md->getName().str() == "llvm.dbg.gv") ||
713 (md->getName().str() == "llvm.dbg.sp")))
717 for (User::const_use_iterator ui=U->use_begin(), ue=U->use_end();
719 if (usedInOneFunc(*ui, oneFunc) == false)
725 /* Find out if a global variable can be demoted to local scope.
726 * Currently, this is valid for CUDA shared variables, which have local
727 * scope and global lifetime. So the conditions to check are :
728 * 1. Is the global variable in shared address space?
729 * 2. Does it have internal linkage?
730 * 3. Is the global variable referenced only in one function?
732 static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) {
733 if (gv->hasInternalLinkage() == false)
735 const PointerType *Pty = gv->getType();
736 if (Pty->getAddressSpace() != llvm::ADDRESS_SPACE_SHARED)
739 const Function *oneFunc = 0;
741 bool flag = usedInOneFunc(gv, oneFunc);
750 static bool useFuncSeen(const Constant *C,
751 llvm::DenseMap<const Function *, bool> &seenMap) {
752 for (Value::const_use_iterator ui=C->use_begin(), ue=C->use_end();
754 if (const Constant *cu = dyn_cast<Constant>(*ui)) {
755 if (useFuncSeen(cu, seenMap))
757 } else if (const Instruction *I = dyn_cast<Instruction>(*ui)) {
758 const BasicBlock *bb = I->getParent();
760 const Function *caller = bb->getParent();
761 if (!caller) continue;
762 if (seenMap.find(caller) != seenMap.end())
769 void NVPTXAsmPrinter::emitDeclarations (Module &M, raw_ostream &O) {
770 llvm::DenseMap<const Function *, bool> seenMap;
771 for (Module::const_iterator FI=M.begin(), FE=M.end();
773 const Function *F = FI;
775 if (F->isDeclaration()) {
778 if (F->getIntrinsicID())
780 CurrentFnSym = Mang->getSymbol(F);
781 emitDeclaration(F, O);
784 for (Value::const_use_iterator iter=F->use_begin(),
785 iterEnd=F->use_end(); iter!=iterEnd; ++iter) {
786 if (const Constant *C = dyn_cast<Constant>(*iter)) {
787 if (usedInGlobalVarDef(C)) {
788 // The use is in the initialization of a global variable
789 // that is a function pointer, so print a declaration
790 // for the original function
791 CurrentFnSym = Mang->getSymbol(F);
792 emitDeclaration(F, O);
795 // Emit a declaration of this function if the function that
796 // uses this constant expr has already been seen.
797 if (useFuncSeen(C, seenMap)) {
798 CurrentFnSym = Mang->getSymbol(F);
799 emitDeclaration(F, O);
804 if (!isa<Instruction>(*iter)) continue;
805 const Instruction *instr = cast<Instruction>(*iter);
806 const BasicBlock *bb = instr->getParent();
808 const Function *caller = bb->getParent();
809 if (!caller) continue;
811 // If a caller has already been seen, then the caller is
812 // appearing in the module before the callee. so print out
813 // a declaration for the callee.
814 if (seenMap.find(caller) != seenMap.end()) {
815 CurrentFnSym = Mang->getSymbol(F);
816 emitDeclaration(F, O);
824 void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) {
825 DebugInfoFinder DbgFinder;
826 DbgFinder.processModule(M);
829 for (DebugInfoFinder::iterator I = DbgFinder.compile_unit_begin(),
830 E = DbgFinder.compile_unit_end(); I != E; ++I) {
831 DICompileUnit DIUnit(*I);
832 StringRef Filename(DIUnit.getFilename());
833 StringRef Dirname(DIUnit.getDirectory());
834 SmallString<128> FullPathName = Dirname;
835 if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {
836 sys::path::append(FullPathName, Filename);
837 Filename = FullPathName.str();
839 if (filenameMap.find(Filename.str()) != filenameMap.end())
841 filenameMap[Filename.str()] = i;
842 OutStreamer.EmitDwarfFileDirective(i, "", Filename.str());
846 for (DebugInfoFinder::iterator I = DbgFinder.subprogram_begin(),
847 E = DbgFinder.subprogram_end(); I != E; ++I) {
849 StringRef Filename(SP.getFilename());
850 StringRef Dirname(SP.getDirectory());
851 SmallString<128> FullPathName = Dirname;
852 if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {
853 sys::path::append(FullPathName, Filename);
854 Filename = FullPathName.str();
856 if (filenameMap.find(Filename.str()) != filenameMap.end())
858 filenameMap[Filename.str()] = i;
863 bool NVPTXAsmPrinter::doInitialization (Module &M) {
865 SmallString<128> Str1;
866 raw_svector_ostream OS1(Str1);
868 MMI = getAnalysisIfAvailable<MachineModuleInfo>();
869 MMI->AnalyzeModule(M);
871 // We need to call the parent's one explicitly.
872 //bool Result = AsmPrinter::doInitialization(M);
874 // Initialize TargetLoweringObjectFile.
875 const_cast<TargetLoweringObjectFile&>(getObjFileLowering())
876 .Initialize(OutContext, TM);
878 Mang = new Mangler(OutContext, *TM.getTargetData());
880 // Emit header before any dwarf directives are emitted below.
882 OutStreamer.EmitRawText(OS1.str());
885 // Already commented out
886 //bool Result = AsmPrinter::doInitialization(M);
889 if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)
890 recordAndEmitFilenames(M);
892 SmallString<128> Str2;
893 raw_svector_ostream OS2(Str2);
895 emitDeclarations(M, OS2);
897 // Print out module-level global variables here.
898 for (Module::global_iterator I = M.global_begin(), E = M.global_end();
900 printModuleLevelGV(I, OS2);
904 OutStreamer.EmitRawText(OS2.str());
905 return false; // success
908 void NVPTXAsmPrinter::emitHeader (Module &M, raw_ostream &O) {
910 O << "// Generated by LLVM NVPTX Back-End\n";
914 O << ".version 3.0\n";
917 O << nvptxSubtarget.getTargetName();
919 if (nvptxSubtarget.getDrvInterface() == NVPTX::NVCL)
920 O << ", texmode_independent";
921 if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) {
922 if (!nvptxSubtarget.hasDouble())
923 O << ", map_f64_to_f32";
926 if (MAI->doesSupportDebugInformation())
931 O << ".address_size ";
932 if (nvptxSubtarget.is64Bit())
941 bool NVPTXAsmPrinter::doFinalization(Module &M) {
942 // XXX Temproarily remove global variables so that doFinalization() will not
943 // emit them again (global variables are emitted at beginning).
945 Module::GlobalListType &global_list = M.getGlobalList();
946 int i, n = global_list.size();
947 GlobalVariable **gv_array = new GlobalVariable* [n];
949 // first, back-up GlobalVariable in gv_array
951 for (Module::global_iterator I = global_list.begin(), E = global_list.end();
955 // second, empty global_list
956 while (!global_list.empty())
957 global_list.remove(global_list.begin());
959 // call doFinalization
960 bool ret = AsmPrinter::doFinalization(M);
962 // now we restore global variables
963 for (i = 0; i < n; i ++)
964 global_list.insert(global_list.end(), gv_array[i]);
970 //bool Result = AsmPrinter::doFinalization(M);
971 // Instead of calling the parents doFinalization, we may
972 // clone parents doFinalization and customize here.
973 // Currently, we if NVISA out the EmitGlobals() in
974 // parent's doFinalization, which is too intrusive.
976 // Same for the doInitialization.
980 // This function emits appropriate linkage directives for
981 // functions and global variables.
983 // extern function declaration -> .extern
984 // extern function definition -> .visible
985 // external global variable with init -> .visible
986 // external without init -> .extern
987 // appending -> not allowed, assert.
989 void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue* V, raw_ostream &O)
991 if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) {
992 if (V->hasExternalLinkage()) {
993 if (isa<GlobalVariable>(V)) {
994 const GlobalVariable *GVar = cast<GlobalVariable>(V);
996 if (GVar->hasInitializer())
1001 } else if (V->isDeclaration())
1005 } else if (V->hasAppendingLinkage()) {
1007 msg.append("Error: ");
1008 msg.append("Symbol ");
1010 msg.append(V->getName().str());
1011 msg.append("has unsupported appending linkage type");
1012 llvm_unreachable(msg.c_str());
1018 void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O,
1019 bool processDemoted) {
1022 if (GVar->hasSection()) {
1023 if (GVar->getSection() == "llvm.metadata")
1027 const TargetData *TD = TM.getTargetData();
1029 // GlobalVariables are always constant pointers themselves.
1030 const PointerType *PTy = GVar->getType();
1031 Type *ETy = PTy->getElementType();
1033 if (GVar->hasExternalLinkage()) {
1034 if (GVar->hasInitializer())
1040 if (llvm::isTexture(*GVar)) {
1041 O << ".global .texref " << llvm::getTextureName(*GVar) << ";\n";
1045 if (llvm::isSurface(*GVar)) {
1046 O << ".global .surfref " << llvm::getSurfaceName(*GVar) << ";\n";
1050 if (GVar->isDeclaration()) {
1051 // (extern) declarations, no definition or initializer
1052 // Currently the only known declaration is for an automatic __local
1053 // (.shared) promoted to global.
1054 emitPTXGlobalVariable(GVar, O);
1059 if (llvm::isSampler(*GVar)) {
1060 O << ".global .samplerref " << llvm::getSamplerName(*GVar);
1062 Constant *Initializer = NULL;
1063 if (GVar->hasInitializer())
1064 Initializer = GVar->getInitializer();
1065 ConstantInt *CI = NULL;
1067 CI = dyn_cast<ConstantInt>(Initializer);
1069 unsigned sample=CI->getZExtValue();
1073 for (int i =0, addr=((sample & __CLK_ADDRESS_MASK ) >>
1074 __CLK_ADDRESS_BASE) ; i < 3 ; i++) {
1075 O << "addr_mode_" << i << " = ";
1077 case 0: O << "wrap"; break;
1078 case 1: O << "clamp_to_border"; break;
1079 case 2: O << "clamp_to_edge"; break;
1080 case 3: O << "wrap"; break;
1081 case 4: O << "mirror"; break;
1085 O << "filter_mode = ";
1086 switch (( sample & __CLK_FILTER_MASK ) >> __CLK_FILTER_BASE ) {
1087 case 0: O << "nearest"; break;
1088 case 1: O << "linear"; break;
1089 case 2: assert ( 0 && "Anisotropic filtering is not supported");
1090 default: O << "nearest"; break;
1092 if (!(( sample &__CLK_NORMALIZED_MASK ) >> __CLK_NORMALIZED_BASE)) {
1093 O << ", force_unnormalized_coords = 1";
1102 if (GVar->hasPrivateLinkage()) {
1104 if (!strncmp(GVar->getName().data(), "unrollpragma", 12))
1107 // FIXME - need better way (e.g. Metadata) to avoid generating this global
1108 if (!strncmp(GVar->getName().data(), "filename", 8))
1110 if (GVar->use_empty())
1114 const Function *demotedFunc = 0;
1115 if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) {
1116 O << "// " << GVar->getName().str() << " has been demoted\n";
1117 if (localDecls.find(demotedFunc) != localDecls.end())
1118 localDecls[demotedFunc].push_back(GVar);
1120 std::vector<GlobalVariable *> temp;
1121 temp.push_back(GVar);
1122 localDecls[demotedFunc] = temp;
1128 emitPTXAddressSpace(PTy->getAddressSpace(), O);
1129 if (GVar->getAlignment() == 0)
1130 O << " .align " << (int) TD->getPrefTypeAlignment(ETy);
1132 O << " .align " << GVar->getAlignment();
1135 if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa<PointerType>(ETy)) {
1137 O << getPTXFundamentalTypeStr(ETy, false);
1139 O << *Mang->getSymbol(GVar);
1141 // Ptx allows variable initilization only for constant and global state
1143 if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) ||
1144 (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) ||
1145 (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST))
1146 && GVar->hasInitializer()) {
1147 Constant *Initializer = GVar->getInitializer();
1148 if (!Initializer->isNullValue()) {
1150 printScalarConstant(Initializer, O);
1154 unsigned int ElementSize =0;
1156 // Although PTX has direct support for struct type and array type and
1157 // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
1158 // targets that support these high level field accesses. Structs, arrays
1159 // and vectors are lowered into arrays of bytes.
1160 switch (ETy->getTypeID()) {
1161 case Type::StructTyID:
1162 case Type::ArrayTyID:
1163 case Type::VectorTyID:
1164 ElementSize = TD->getTypeStoreSize(ETy);
1165 // Ptx allows variable initilization only for constant and
1166 // global state spaces.
1167 if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) ||
1168 (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) ||
1169 (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST))
1170 && GVar->hasInitializer()) {
1171 Constant *Initializer = GVar->getInitializer();
1172 if (!isa<UndefValue>(Initializer) &&
1173 !Initializer->isNullValue()) {
1174 AggBuffer aggBuffer(ElementSize, O, *this);
1175 bufferAggregateConstant(Initializer, &aggBuffer);
1176 if (aggBuffer.numSymbols) {
1177 if (nvptxSubtarget.is64Bit()) {
1178 O << " .u64 " << *Mang->getSymbol(GVar) <<"[" ;
1182 O << " .u32 " << *Mang->getSymbol(GVar) <<"[" ;
1188 O << " .b8 " << *Mang->getSymbol(GVar) <<"[" ;
1197 O << " .b8 " << *Mang->getSymbol(GVar) ;
1206 O << " .b8 " << *Mang->getSymbol(GVar);
1215 assert( 0 && "type not supported yet");
1222 void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) {
1223 if (localDecls.find(f) == localDecls.end())
1226 std::vector<GlobalVariable *> &gvars = localDecls[f];
1228 for (unsigned i=0, e=gvars.size(); i!=e; ++i) {
1229 O << "\t// demoted variable\n\t";
1230 printModuleLevelGV(gvars[i], O, true);
1234 void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
1235 raw_ostream &O) const {
1236 switch (AddressSpace) {
1237 case llvm::ADDRESS_SPACE_LOCAL:
1240 case llvm::ADDRESS_SPACE_GLOBAL:
1243 case llvm::ADDRESS_SPACE_CONST:
1244 // This logic should be consistent with that in
1245 // getCodeAddrSpace() (NVPTXISelDATToDAT.cpp)
1246 if (nvptxSubtarget.hasGenericLdSt())
1251 case llvm::ADDRESS_SPACE_CONST_NOT_GEN:
1254 case llvm::ADDRESS_SPACE_SHARED:
1258 llvm_unreachable("unexpected address space");
1262 std::string NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty,
1263 bool useB4PTR) const {
1264 switch (Ty->getTypeID()) {
1266 llvm_unreachable("unexpected type");
1268 case Type::IntegerTyID: {
1269 unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
1272 else if (NumBits <= 64) {
1273 std::string name = "u";
1274 return name + utostr(NumBits);
1276 llvm_unreachable("Integer too large");
1281 case Type::FloatTyID:
1283 case Type::DoubleTyID:
1285 case Type::PointerTyID:
1286 if (nvptxSubtarget.is64Bit())
1287 if (useB4PTR) return "b64";
1290 if (useB4PTR) return "b32";
1293 llvm_unreachable("unexpected type");
1297 void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable* GVar,
1300 const TargetData *TD = TM.getTargetData();
1302 // GlobalVariables are always constant pointers themselves.
1303 const PointerType *PTy = GVar->getType();
1304 Type *ETy = PTy->getElementType();
1307 emitPTXAddressSpace(PTy->getAddressSpace(), O);
1308 if (GVar->getAlignment() == 0)
1309 O << " .align " << (int) TD->getPrefTypeAlignment(ETy);
1311 O << " .align " << GVar->getAlignment();
1313 if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa<PointerType>(ETy)) {
1315 O << getPTXFundamentalTypeStr(ETy);
1317 O << *Mang->getSymbol(GVar);
1321 int64_t ElementSize =0;
1323 // Although PTX has direct support for struct type and array type and LLVM IR
1324 // is very similar to PTX, the LLVM CodeGen does not support for targets that
1325 // support these high level field accesses. Structs and arrays are lowered
1326 // into arrays of bytes.
1327 switch (ETy->getTypeID()) {
1328 case Type::StructTyID:
1329 case Type::ArrayTyID:
1330 case Type::VectorTyID:
1331 ElementSize = TD->getTypeStoreSize(ETy);
1332 O << " .b8 " << *Mang->getSymbol(GVar) <<"[" ;
1334 O << itostr(ElementSize) ;
1339 assert( 0 && "type not supported yet");
1346 getOpenCLAlignment(const TargetData *TD,
1348 if (Ty->isPrimitiveType() || Ty->isIntegerTy() || isa<PointerType>(Ty))
1349 return TD->getPrefTypeAlignment(Ty);
1351 const ArrayType *ATy = dyn_cast<ArrayType>(Ty);
1353 return getOpenCLAlignment(TD, ATy->getElementType());
1355 const VectorType *VTy = dyn_cast<VectorType>(Ty);
1357 Type *ETy = VTy->getElementType();
1358 unsigned int numE = VTy->getNumElements();
1359 unsigned int alignE = TD->getPrefTypeAlignment(ETy);
1366 const StructType *STy = dyn_cast<StructType>(Ty);
1368 unsigned int alignStruct = 1;
1369 // Go through each element of the struct and find the
1370 // largest alignment.
1371 for (unsigned i=0, e=STy->getNumElements(); i != e; i++) {
1372 Type *ETy = STy->getElementType(i);
1373 unsigned int align = getOpenCLAlignment(TD, ETy);
1374 if (align > alignStruct)
1375 alignStruct = align;
1380 const FunctionType *FTy = dyn_cast<FunctionType>(Ty);
1382 return TD->getPointerPrefAlignment();
1383 return TD->getPrefTypeAlignment(Ty);
1386 void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I,
1387 int paramIndex, raw_ostream &O) {
1388 if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) ||
1389 (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA))
1390 O << *CurrentFnSym << "_param_" << paramIndex;
1392 std::string argName = I->getName();
1393 const char *p = argName.c_str();
1404 void NVPTXAsmPrinter::printParamName(int paramIndex, raw_ostream &O) {
1405 Function::const_arg_iterator I, E;
1408 if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) ||
1409 (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) {
1410 O << *CurrentFnSym << "_param_" << paramIndex;
1414 for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, i++) {
1415 if (i==paramIndex) {
1416 printParamName(I, paramIndex, O);
1420 llvm_unreachable("paramIndex out of bound");
1423 void NVPTXAsmPrinter::emitFunctionParamList(const Function *F,
1425 const TargetData *TD = TM.getTargetData();
1426 const AttrListPtr &PAL = F->getAttributes();
1427 const TargetLowering *TLI = TM.getTargetLowering();
1428 Function::const_arg_iterator I, E;
1429 unsigned paramIndex = 0;
1431 bool isKernelFunc = llvm::isKernelFunction(*F);
1432 bool isABI = (nvptxSubtarget.getSmVersion() >= 20);
1433 MVT thePointerTy = TLI->getPointerTy();
1437 for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) {
1438 const Type *Ty = I->getType();
1445 // Handle image/sampler parameters
1446 if (llvm::isSampler(*I) || llvm::isImage(*I)) {
1447 if (llvm::isImage(*I)) {
1448 std::string sname = I->getName();
1449 if (llvm::isImageWriteOnly(*I))
1450 O << "\t.param .surfref " << *CurrentFnSym << "_param_" << paramIndex;
1451 else // Default image is read_only
1452 O << "\t.param .texref " << *CurrentFnSym << "_param_" << paramIndex;
1454 else // Should be llvm::isSampler(*I)
1455 O << "\t.param .samplerref " << *CurrentFnSym << "_param_"
1460 if (PAL.paramHasAttr(paramIndex+1, Attribute::ByVal) == false) {
1462 const PointerType *PTy = dyn_cast<PointerType>(Ty);
1465 // Special handling for pointer arguments to kernel
1466 O << "\t.param .u" << thePointerTy.getSizeInBits() << " ";
1468 if (nvptxSubtarget.getDrvInterface() != NVPTX::CUDA) {
1469 Type *ETy = PTy->getElementType();
1470 int addrSpace = PTy->getAddressSpace();
1475 case llvm::ADDRESS_SPACE_CONST_NOT_GEN:
1476 O << ".ptr .const ";
1478 case llvm::ADDRESS_SPACE_SHARED:
1479 O << ".ptr .shared ";
1481 case llvm::ADDRESS_SPACE_GLOBAL:
1482 case llvm::ADDRESS_SPACE_CONST:
1483 O << ".ptr .global ";
1486 O << ".align " << (int)getOpenCLAlignment(TD, ETy) << " ";
1488 printParamName(I, paramIndex, O);
1492 // non-pointer scalar to kernel func
1494 << getPTXFundamentalTypeStr(Ty) << " ";
1495 printParamName(I, paramIndex, O);
1498 // Non-kernel function, just print .param .b<size> for ABI
1499 // and .reg .b<size> for non ABY
1501 if (isa<IntegerType>(Ty)) {
1502 sz = cast<IntegerType>(Ty)->getBitWidth();
1503 if (sz < 32) sz = 32;
1505 else if (isa<PointerType>(Ty))
1506 sz = thePointerTy.getSizeInBits();
1508 sz = Ty->getPrimitiveSizeInBits();
1510 O << "\t.param .b" << sz << " ";
1512 O << "\t.reg .b" << sz << " ";
1513 printParamName(I, paramIndex, O);
1517 // param has byVal attribute. So should be a pointer
1518 const PointerType *PTy = dyn_cast<PointerType>(Ty);
1520 "Param with byval attribute should be a pointer type");
1521 Type *ETy = PTy->getElementType();
1523 if (isABI || isKernelFunc) {
1524 // Just print .param .b8 .align <a> .param[size];
1525 // <a> = PAL.getparamalignment
1526 // size = typeallocsize of element type
1527 unsigned align = PAL.getParamAlignment(paramIndex+1);
1528 unsigned sz = TD->getTypeAllocSize(ETy);
1529 O << "\t.param .align " << align
1531 printParamName(I, paramIndex, O);
1532 O << "[" << sz << "]";
1535 // Split the ETy into constituent parts and
1536 // print .param .b<size> <name> for each part.
1537 // Further, if a part is vector, print the above for
1538 // each vector element.
1539 SmallVector<EVT, 16> vtparts;
1540 ComputeValueVTs(*TLI, ETy, vtparts);
1541 for (unsigned i=0,e=vtparts.size(); i!=e; ++i) {
1543 EVT elemtype = vtparts[i];
1544 if (vtparts[i].isVector()) {
1545 elems = vtparts[i].getVectorNumElements();
1546 elemtype = vtparts[i].getVectorElementType();
1549 for (unsigned j=0,je=elems; j!=je; ++j) {
1550 unsigned sz = elemtype.getSizeInBits();
1551 if (elemtype.isInteger() && (sz < 32)) sz = 32;
1552 O << "\t.reg .b" << sz << " ";
1553 printParamName(I, paramIndex, O);
1554 if (j<je-1) O << ",\n";
1568 void NVPTXAsmPrinter::emitFunctionParamList(const MachineFunction &MF,
1570 const Function *F = MF.getFunction();
1571 emitFunctionParamList(F, O);
1575 void NVPTXAsmPrinter::
1576 setAndEmitFunctionVirtualRegisters(const MachineFunction &MF) {
1577 SmallString<128> Str;
1578 raw_svector_ostream O(Str);
1580 // Map the global virtual register number to a register class specific
1581 // virtual register number starting from 1 with that class.
1582 const TargetRegisterInfo *TRI = MF.getTarget().getRegisterInfo();
1583 //unsigned numRegClasses = TRI->getNumRegClasses();
1585 // Emit the Fake Stack Object
1586 const MachineFrameInfo *MFI = MF.getFrameInfo();
1587 int NumBytes = (int) MFI->getStackSize();
1589 O << "\t.local .align " << MFI->getMaxAlignment() << " .b8 \t"
1591 << getFunctionNumber() << "[" << NumBytes << "];\n";
1592 if (nvptxSubtarget.is64Bit()) {
1593 O << "\t.reg .b64 \t%SP;\n";
1594 O << "\t.reg .b64 \t%SPL;\n";
1597 O << "\t.reg .b32 \t%SP;\n";
1598 O << "\t.reg .b32 \t%SPL;\n";
1602 // Go through all virtual registers to establish the mapping between the
1604 // register number and the per class virtual register number.
1605 // We use the per class virtual register number in the ptx output.
1606 unsigned int numVRs = MRI->getNumVirtRegs();
1607 for (unsigned i=0; i< numVRs; i++) {
1608 unsigned int vr = TRI->index2VirtReg(i);
1609 const TargetRegisterClass *RC = MRI->getRegClass(vr);
1610 std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[RC->getID()];
1611 int n = regmap.size();
1612 regmap.insert(std::make_pair(vr, n+1));
1615 // Emit register declarations
1616 // @TODO: Extract out the real register usage
1617 O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n";
1618 O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n";
1619 O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n";
1620 O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n";
1621 O << "\t.reg .s64 %rl<" << NVPTXNumRegisters << ">;\n";
1622 O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n";
1623 O << "\t.reg .f64 %fl<" << NVPTXNumRegisters << ">;\n";
1625 // Emit declaration of the virtual registers or 'physical' registers for
1626 // each register class
1627 //for (unsigned i=0; i< numRegClasses; i++) {
1628 // std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[i];
1629 // const TargetRegisterClass *RC = TRI->getRegClass(i);
1630 // std::string rcname = getNVPTXRegClassName(RC);
1631 // std::string rcStr = getNVPTXRegClassStr(RC);
1632 // //int n = regmap.size();
1633 // if (!isNVPTXVectorRegClass(RC)) {
1634 // O << "\t.reg " << rcname << " \t" << rcStr << "<"
1635 // << NVPTXNumRegisters << ">;\n";
1638 // Only declare those registers that may be used. And do not emit vector
1640 // they are all elementized to scalar registers.
1641 //if (n && !isNVPTXVectorRegClass(RC)) {
1642 // if (RegAllocNilUsed) {
1643 // O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1)
1647 // O << "\t.reg " << rcname << " \t" << StrToUpper(rcStr)
1648 // << "<" << 32 << ">;\n";
1653 OutStreamer.EmitRawText(O.str());
1657 void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) {
1658 APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
1660 unsigned int numHex;
1663 if (Fp->getType()->getTypeID()==Type::FloatTyID) {
1666 APF.convert(APFloat::IEEEsingle, APFloat::rmNearestTiesToEven,
1668 } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
1671 APF.convert(APFloat::IEEEdouble, APFloat::rmNearestTiesToEven,
1674 llvm_unreachable("unsupported fp type");
1676 APInt API = APF.bitcastToAPInt();
1677 std::string hexstr(utohexstr(API.getZExtValue()));
1679 if (hexstr.length() < numHex)
1680 O << std::string(numHex - hexstr.length(), '0');
1681 O << utohexstr(API.getZExtValue());
1684 void NVPTXAsmPrinter::printScalarConstant(Constant *CPV, raw_ostream &O) {
1685 if (ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
1686 O << CI->getValue();
1689 if (ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
1690 printFPConstant(CFP, O);
1693 if (isa<ConstantPointerNull>(CPV)) {
1697 if (GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1698 O << *Mang->getSymbol(GVar);
1701 if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1702 Value *v = Cexpr->stripPointerCasts();
1703 if (GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
1704 O << *Mang->getSymbol(GVar);
1707 O << *LowerConstant(CPV, *this);
1711 llvm_unreachable("Not scalar type found in printScalarConstant()");
1715 void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes,
1716 AggBuffer *aggBuffer) {
1718 const TargetData *TD = TM.getTargetData();
1720 if (isa<UndefValue>(CPV) || CPV->isNullValue()) {
1721 int s = TD->getTypeAllocSize(CPV->getType());
1724 aggBuffer->addZeros(s);
1729 switch (CPV->getType()->getTypeID()) {
1731 case Type::IntegerTyID: {
1732 const Type *ETy = CPV->getType();
1733 if ( ETy == Type::getInt8Ty(CPV->getContext()) ){
1735 (unsigned char)(dyn_cast<ConstantInt>(CPV))->getZExtValue();
1737 aggBuffer->addBytes(ptr, 1, Bytes);
1738 } else if ( ETy == Type::getInt16Ty(CPV->getContext()) ) {
1740 (short)(dyn_cast<ConstantInt>(CPV))->getZExtValue();
1741 ptr = (unsigned char*)&int16;
1742 aggBuffer->addBytes(ptr, 2, Bytes);
1743 } else if ( ETy == Type::getInt32Ty(CPV->getContext()) ) {
1744 if (ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
1745 int int32 =(int)(constInt->getZExtValue());
1746 ptr = (unsigned char*)&int32;
1747 aggBuffer->addBytes(ptr, 4, Bytes);
1749 } else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1750 if (ConstantInt *constInt =
1751 dyn_cast<ConstantInt>(ConstantFoldConstantExpression(
1753 int int32 =(int)(constInt->getZExtValue());
1754 ptr = (unsigned char*)&int32;
1755 aggBuffer->addBytes(ptr, 4, Bytes);
1758 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1759 Value *v = Cexpr->getOperand(0)->stripPointerCasts();
1760 aggBuffer->addSymbol(v);
1761 aggBuffer->addZeros(4);
1765 llvm_unreachable("unsupported integer const type");
1766 } else if (ETy == Type::getInt64Ty(CPV->getContext()) ) {
1767 if (ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
1768 long long int64 =(long long)(constInt->getZExtValue());
1769 ptr = (unsigned char*)&int64;
1770 aggBuffer->addBytes(ptr, 8, Bytes);
1772 } else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1773 if (ConstantInt *constInt = dyn_cast<ConstantInt>(
1774 ConstantFoldConstantExpression(Cexpr, TD))) {
1775 long long int64 =(long long)(constInt->getZExtValue());
1776 ptr = (unsigned char*)&int64;
1777 aggBuffer->addBytes(ptr, 8, Bytes);
1780 if (Cexpr->getOpcode() == Instruction::PtrToInt) {
1781 Value *v = Cexpr->getOperand(0)->stripPointerCasts();
1782 aggBuffer->addSymbol(v);
1783 aggBuffer->addZeros(8);
1787 llvm_unreachable("unsupported integer const type");
1789 llvm_unreachable("unsupported integer const type");
1792 case Type::FloatTyID:
1793 case Type::DoubleTyID: {
1794 ConstantFP *CFP = dyn_cast<ConstantFP>(CPV);
1795 const Type* Ty = CFP->getType();
1796 if (Ty == Type::getFloatTy(CPV->getContext())) {
1797 float float32 = (float)CFP->getValueAPF().convertToFloat();
1798 ptr = (unsigned char*)&float32;
1799 aggBuffer->addBytes(ptr, 4, Bytes);
1800 } else if (Ty == Type::getDoubleTy(CPV->getContext())) {
1801 double float64 = CFP->getValueAPF().convertToDouble();
1802 ptr = (unsigned char*)&float64;
1803 aggBuffer->addBytes(ptr, 8, Bytes);
1806 llvm_unreachable("unsupported fp const type");
1810 case Type::PointerTyID: {
1811 if (GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
1812 aggBuffer->addSymbol(GVar);
1814 else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
1815 Value *v = Cexpr->stripPointerCasts();
1816 aggBuffer->addSymbol(v);
1818 unsigned int s = TD->getTypeAllocSize(CPV->getType());
1819 aggBuffer->addZeros(s);
1823 case Type::ArrayTyID:
1824 case Type::VectorTyID:
1825 case Type::StructTyID: {
1826 if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV) ||
1827 isa<ConstantStruct>(CPV)) {
1828 int ElementSize = TD->getTypeAllocSize(CPV->getType());
1829 bufferAggregateConstant(CPV, aggBuffer);
1830 if ( Bytes > ElementSize )
1831 aggBuffer->addZeros(Bytes-ElementSize);
1833 else if (isa<ConstantAggregateZero>(CPV))
1834 aggBuffer->addZeros(Bytes);
1836 llvm_unreachable("Unexpected Constant type");
1841 llvm_unreachable("unsupported type");
1845 void NVPTXAsmPrinter::bufferAggregateConstant(Constant *CPV,
1846 AggBuffer *aggBuffer) {
1847 const TargetData *TD = TM.getTargetData();
1851 if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
1852 if (CPV->getNumOperands())
1853 for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i)
1854 bufferLEByte(cast<Constant>(CPV->getOperand(i)), 0, aggBuffer);
1858 if (const ConstantDataSequential *CDS =
1859 dyn_cast<ConstantDataSequential>(CPV)) {
1860 if (CDS->getNumElements())
1861 for (unsigned i = 0; i < CDS->getNumElements(); ++i)
1862 bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
1868 if (isa<ConstantStruct>(CPV)) {
1869 if (CPV->getNumOperands()) {
1870 StructType *ST = cast<StructType>(CPV->getType());
1871 for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) {
1873 Bytes = TD->getStructLayout(ST)->getElementOffset(0) +
1874 TD->getTypeAllocSize(ST)
1875 - TD->getStructLayout(ST)->getElementOffset(i);
1877 Bytes = TD->getStructLayout(ST)->getElementOffset(i+1) -
1878 TD->getStructLayout(ST)->getElementOffset(i);
1879 bufferLEByte(cast<Constant>(CPV->getOperand(i)), Bytes,
1885 llvm_unreachable("unsupported constant type in printAggregateConstant()");
1888 // buildTypeNameMap - Run through symbol table looking for type names.
1892 bool NVPTXAsmPrinter::isImageType(const Type *Ty) {
1894 std::map<const Type *, std::string>::iterator PI = TypeNameMap.find(Ty);
1896 if (PI != TypeNameMap.end() &&
1897 (!PI->second.compare("struct._image1d_t") ||
1898 !PI->second.compare("struct._image2d_t") ||
1899 !PI->second.compare("struct._image3d_t")))
1905 /// PrintAsmOperand - Print out an operand for an inline asm expression.
1907 bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
1908 unsigned AsmVariant,
1909 const char *ExtraCode,
1911 if (ExtraCode && ExtraCode[0]) {
1912 if (ExtraCode[1] != 0) return true; // Unknown modifier.
1914 switch (ExtraCode[0]) {
1916 // See if this is a generic print operand
1917 return AsmPrinter::PrintAsmOperand(MI, OpNo, AsmVariant, ExtraCode, O);
1923 printOperand(MI, OpNo, O);
1928 bool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI,
1930 unsigned AsmVariant,
1931 const char *ExtraCode,
1933 if (ExtraCode && ExtraCode[0])
1934 return true; // Unknown modifier
1937 printMemOperand(MI, OpNo, O);
1943 bool NVPTXAsmPrinter::ignoreLoc(const MachineInstr &MI)
1945 switch(MI.getOpcode()) {
1948 case NVPTX::CallArgBeginInst: case NVPTX::CallArgEndInst0:
1949 case NVPTX::CallArgEndInst1: case NVPTX::CallArgF32:
1950 case NVPTX::CallArgF64: case NVPTX::CallArgI16:
1951 case NVPTX::CallArgI32: case NVPTX::CallArgI32imm:
1952 case NVPTX::CallArgI64: case NVPTX::CallArgI8:
1953 case NVPTX::CallArgParam: case NVPTX::CallVoidInst:
1954 case NVPTX::CallVoidInstReg: case NVPTX::Callseq_End:
1955 case NVPTX::CallVoidInstReg64:
1956 case NVPTX::DeclareParamInst: case NVPTX::DeclareRetMemInst:
1957 case NVPTX::DeclareRetRegInst: case NVPTX::DeclareRetScalarInst:
1958 case NVPTX::DeclareScalarParamInst: case NVPTX::DeclareScalarRegInst:
1959 case NVPTX::StoreParamF32: case NVPTX::StoreParamF64:
1960 case NVPTX::StoreParamI16: case NVPTX::StoreParamI32:
1961 case NVPTX::StoreParamI64: case NVPTX::StoreParamI8:
1962 case NVPTX::StoreParamS32I8: case NVPTX::StoreParamU32I8:
1963 case NVPTX::StoreParamS32I16: case NVPTX::StoreParamU32I16:
1964 case NVPTX::StoreParamScalar2F32: case NVPTX::StoreParamScalar2F64:
1965 case NVPTX::StoreParamScalar2I16: case NVPTX::StoreParamScalar2I32:
1966 case NVPTX::StoreParamScalar2I64: case NVPTX::StoreParamScalar2I8:
1967 case NVPTX::StoreParamScalar4F32: case NVPTX::StoreParamScalar4I16:
1968 case NVPTX::StoreParamScalar4I32: case NVPTX::StoreParamScalar4I8:
1969 case NVPTX::StoreParamV2F32: case NVPTX::StoreParamV2F64:
1970 case NVPTX::StoreParamV2I16: case NVPTX::StoreParamV2I32:
1971 case NVPTX::StoreParamV2I64: case NVPTX::StoreParamV2I8:
1972 case NVPTX::StoreParamV4F32: case NVPTX::StoreParamV4I16:
1973 case NVPTX::StoreParamV4I32: case NVPTX::StoreParamV4I8:
1974 case NVPTX::StoreRetvalF32: case NVPTX::StoreRetvalF64:
1975 case NVPTX::StoreRetvalI16: case NVPTX::StoreRetvalI32:
1976 case NVPTX::StoreRetvalI64: case NVPTX::StoreRetvalI8:
1977 case NVPTX::StoreRetvalScalar2F32: case NVPTX::StoreRetvalScalar2F64:
1978 case NVPTX::StoreRetvalScalar2I16: case NVPTX::StoreRetvalScalar2I32:
1979 case NVPTX::StoreRetvalScalar2I64: case NVPTX::StoreRetvalScalar2I8:
1980 case NVPTX::StoreRetvalScalar4F32: case NVPTX::StoreRetvalScalar4I16:
1981 case NVPTX::StoreRetvalScalar4I32: case NVPTX::StoreRetvalScalar4I8:
1982 case NVPTX::StoreRetvalV2F32: case NVPTX::StoreRetvalV2F64:
1983 case NVPTX::StoreRetvalV2I16: case NVPTX::StoreRetvalV2I32:
1984 case NVPTX::StoreRetvalV2I64: case NVPTX::StoreRetvalV2I8:
1985 case NVPTX::StoreRetvalV4F32: case NVPTX::StoreRetvalV4I16:
1986 case NVPTX::StoreRetvalV4I32: case NVPTX::StoreRetvalV4I8:
1987 case NVPTX::LastCallArgF32: case NVPTX::LastCallArgF64:
1988 case NVPTX::LastCallArgI16: case NVPTX::LastCallArgI32:
1989 case NVPTX::LastCallArgI32imm: case NVPTX::LastCallArgI64:
1990 case NVPTX::LastCallArgI8: case NVPTX::LastCallArgParam:
1991 case NVPTX::LoadParamMemF32: case NVPTX::LoadParamMemF64:
1992 case NVPTX::LoadParamMemI16: case NVPTX::LoadParamMemI32:
1993 case NVPTX::LoadParamMemI64: case NVPTX::LoadParamMemI8:
1994 case NVPTX::LoadParamRegF32: case NVPTX::LoadParamRegF64:
1995 case NVPTX::LoadParamRegI16: case NVPTX::LoadParamRegI32:
1996 case NVPTX::LoadParamRegI64: case NVPTX::LoadParamRegI8:
1997 case NVPTX::LoadParamScalar2F32: case NVPTX::LoadParamScalar2F64:
1998 case NVPTX::LoadParamScalar2I16: case NVPTX::LoadParamScalar2I32:
1999 case NVPTX::LoadParamScalar2I64: case NVPTX::LoadParamScalar2I8:
2000 case NVPTX::LoadParamScalar4F32: case NVPTX::LoadParamScalar4I16:
2001 case NVPTX::LoadParamScalar4I32: case NVPTX::LoadParamScalar4I8:
2002 case NVPTX::LoadParamV2F32: case NVPTX::LoadParamV2F64:
2003 case NVPTX::LoadParamV2I16: case NVPTX::LoadParamV2I32:
2004 case NVPTX::LoadParamV2I64: case NVPTX::LoadParamV2I8:
2005 case NVPTX::LoadParamV4F32: case NVPTX::LoadParamV4I16:
2006 case NVPTX::LoadParamV4I32: case NVPTX::LoadParamV4I8:
2007 case NVPTX::PrototypeInst: case NVPTX::DBG_VALUE:
2013 // Force static initialization.
2014 extern "C" void LLVMInitializeNVPTXBackendAsmPrinter() {
2015 RegisterAsmPrinter<NVPTXAsmPrinter> X(TheNVPTXTarget32);
2016 RegisterAsmPrinter<NVPTXAsmPrinter> Y(TheNVPTXTarget64);
2020 void NVPTXAsmPrinter::emitSrcInText(StringRef filename, unsigned line) {
2021 std::stringstream temp;
2022 LineReader * reader = this->getReader(filename.str());
2024 temp << filename.str();
2028 temp << reader->readLine(line);
2030 this->OutStreamer.EmitRawText(Twine(temp.str()));
2034 LineReader *NVPTXAsmPrinter::getReader(std::string filename) {
2035 if (reader == NULL) {
2036 reader = new LineReader(filename);
2039 if (reader->fileName() != filename) {
2041 reader = new LineReader(filename);
2049 LineReader::readLine(unsigned lineNum) {
2050 if (lineNum < theCurLine) {
2052 fstr.seekg(0,std::ios::beg);
2054 while (theCurLine < lineNum) {
2055 fstr.getline(buff,500);
2061 // Force static initialization.
2062 extern "C" void LLVMInitializeNVPTXAsmPrinter() {
2063 RegisterAsmPrinter<NVPTXAsmPrinter> X(TheNVPTXTarget32);
2064 RegisterAsmPrinter<NVPTXAsmPrinter> Y(TheNVPTXTarget64);