119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman//===-- PTXAsmPrinter.cpp - PTX LLVM assembly writer ----------------------===// 219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// 319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// The LLVM Compiler Infrastructure 419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// 519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// This file is distributed under the University of Illinois Open Source 619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// License. See LICENSE.TXT for details. 719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// 819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman//===----------------------------------------------------------------------===// 919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// 1019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// This file contains a printer that converts from our internal representation 1119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// of machine-dependent LLVM code to PTX assembly language. 1219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// 1319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman//===----------------------------------------------------------------------===// 1419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 1519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#define DEBUG_TYPE "ptx-asm-printer" 1619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 1719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "PTX.h" 1819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "PTXAsmPrinter.h" 1919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "PTXMachineFunctionInfo.h" 2019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "PTXParamManager.h" 2119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "PTXRegisterInfo.h" 2219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "PTXTargetMachine.h" 2319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Argument.h" 2419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/DerivedTypes.h" 2519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Function.h" 2619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Module.h" 2719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/ADT/SmallString.h" 2819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/ADT/StringExtras.h" 2919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/ADT/Twine.h" 3019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Analysis/DebugInfo.h" 3119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/CodeGen/AsmPrinter.h" 3219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/CodeGen/MachineFrameInfo.h" 3319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/CodeGen/MachineInstr.h" 3419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/CodeGen/MachineRegisterInfo.h" 3519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/MC/MCContext.h" 3619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/MC/MCExpr.h" 3719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/MC/MCInst.h" 3819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/MC/MCStreamer.h" 3919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/MC/MCSymbol.h" 4019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Target/Mangler.h" 4119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Target/TargetLoweringObjectFile.h" 4219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Support/CommandLine.h" 4319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Support/Debug.h" 4419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Support/ErrorHandling.h" 4519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Support/MathExtras.h" 4619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Support/Path.h" 4719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Support/TargetRegistry.h" 4819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#include "llvm/Support/raw_ostream.h" 4919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 5019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanusing namespace llvm; 5119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 5219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanstatic const char PARAM_PREFIX[] = "__param_"; 5319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanstatic const char RETURN_PREFIX[] = "__ret_"; 5419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 5519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanstatic const char *getRegisterTypeName(unsigned RegNo, 5619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const MachineRegisterInfo& MRI) { 5719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const TargetRegisterClass *TRC = MRI.getRegClass(RegNo); 5819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 5919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#define TEST_REGCLS(cls, clsstr) \ 6019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (PTX::cls ## RegisterClass == TRC) return # clsstr; 6119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 6219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman TEST_REGCLS(RegPred, pred); 6319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman TEST_REGCLS(RegI16, b16); 6419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman TEST_REGCLS(RegI32, b32); 6519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman TEST_REGCLS(RegI64, b64); 6619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman TEST_REGCLS(RegF32, b32); 6719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman TEST_REGCLS(RegF64, b64); 6819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman#undef TEST_REGCLS 6919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 7019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman llvm_unreachable("Not in any register class!"); 7119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return NULL; 7219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 7319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 7419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanstatic const char *getStateSpaceName(unsigned addressSpace) { 7519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman switch (addressSpace) { 7619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman default: llvm_unreachable("Unknown state space"); 7719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case PTXStateSpace::Global: return "global"; 7819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case PTXStateSpace::Constant: return "const"; 7919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case PTXStateSpace::Local: return "local"; 8019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case PTXStateSpace::Parameter: return "param"; 8119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case PTXStateSpace::Shared: return "shared"; 8219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 8319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return NULL; 8419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 8519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 8619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanstatic const char *getTypeName(Type* type) { 8719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman while (true) { 8819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman switch (type->getTypeID()) { 8919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman default: llvm_unreachable("Unknown type"); 9019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case Type::FloatTyID: return ".f32"; 9119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case Type::DoubleTyID: return ".f64"; 9219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case Type::IntegerTyID: 9319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman switch (type->getPrimitiveSizeInBits()) { 9419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman default: llvm_unreachable("Unknown integer bit-width"); 9519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case 16: return ".u16"; 9619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case 32: return ".u32"; 9719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case 64: return ".u64"; 9819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 9919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case Type::ArrayTyID: 10019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case Type::PointerTyID: 10119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman type = dyn_cast<SequentialType>(type)->getElementType(); 10219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 10319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 10419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 10519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return NULL; 10619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 10719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 10819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanbool PTXAsmPrinter::doFinalization(Module &M) { 10919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // XXX Temproarily remove global variables so that doFinalization() will not 11019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // emit them again (global variables are emitted at beginning). 11119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 11219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman Module::GlobalListType &global_list = M.getGlobalList(); 11319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman int i, n = global_list.size(); 11419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman GlobalVariable **gv_array = new GlobalVariable* [n]; 11519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 11619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // first, back-up GlobalVariable in gv_array 11719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman i = 0; 11819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (Module::global_iterator I = global_list.begin(), E = global_list.end(); 11919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman I != E; ++I) 12019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman gv_array[i++] = &*I; 12119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 12219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // second, empty global_list 12319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman while (!global_list.empty()) 12419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman global_list.remove(global_list.begin()); 12519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 12619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // call doFinalization 12719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman bool ret = AsmPrinter::doFinalization(M); 12819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 12919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // now we restore global variables 13019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (i = 0; i < n; i ++) 13119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman global_list.insert(global_list.end(), gv_array[i]); 13219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 13319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman delete[] gv_array; 13419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return ret; 13519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 13619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 13719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanvoid PTXAsmPrinter::EmitStartOfAsmFile(Module &M) 13819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman{ 13919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>(); 14019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 14119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Emit the PTX .version and .target attributes 14219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine("\t.version " + ST.getPTXVersionString())); 14319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine("\t.target " + ST.getTargetString() + 14419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman (ST.supportsDouble() ? "" 14519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman : ", map_f64_to_f32"))); 14619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // .address_size directive is optional, but it must immediately follow 14719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // the .target directive if present within a module 14819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (ST.supportsPTX23()) { 14919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman std::string addrSize = ST.is64Bit() ? "64" : "32"; 15019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine("\t.address_size " + addrSize)); 15119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 15219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 15319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.AddBlankLine(); 15419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 15519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Define any .file directives 15619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman DebugInfoFinder DbgFinder; 15719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman DbgFinder.processModule(M); 15819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 15919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (DebugInfoFinder::iterator I = DbgFinder.compile_unit_begin(), 16019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman E = DbgFinder.compile_unit_end(); I != E; ++I) { 16119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman DICompileUnit DIUnit(*I); 16219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman StringRef FN = DIUnit.getFilename(); 16319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman StringRef Dir = DIUnit.getDirectory(); 16419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman GetOrCreateSourceID(FN, Dir); 16519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 16619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 16719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.AddBlankLine(); 16819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 16919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // declare global variables 17019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (Module::const_global_iterator i = M.global_begin(), e = M.global_end(); 17119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman i != e; ++i) 17219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman EmitVariableDeclaration(i); 17319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 17419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 17519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanvoid PTXAsmPrinter::EmitFunctionBodyStart() { 17619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine("{")); 17719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 17819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>(); 17919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PTXParamManager &PM = MFI->getParamManager(); 18019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 18119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Print register definitions 18219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman std::string regDefs; 18319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman unsigned numRegs; 18419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 18519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // pred 18619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman numRegs = MFI->getNumRegistersForClass(PTX::RegPredRegisterClass); 18719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if(numRegs > 0) { 18819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += "\t.reg .pred %p<"; 18919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += utostr(numRegs); 19019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += ">;\n"; 19119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 19219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 19319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // i16 19419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman numRegs = MFI->getNumRegistersForClass(PTX::RegI16RegisterClass); 19519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if(numRegs > 0) { 19619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += "\t.reg .b16 %rh<"; 19719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += utostr(numRegs); 19819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += ">;\n"; 19919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 20019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 20119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // i32 20219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman numRegs = MFI->getNumRegistersForClass(PTX::RegI32RegisterClass); 20319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if(numRegs > 0) { 20419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += "\t.reg .b32 %r<"; 20519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += utostr(numRegs); 20619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += ">;\n"; 20719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 20819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 20919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // i64 21019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman numRegs = MFI->getNumRegistersForClass(PTX::RegI64RegisterClass); 21119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if(numRegs > 0) { 21219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += "\t.reg .b64 %rd<"; 21319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += utostr(numRegs); 21419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += ">;\n"; 21519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 21619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 21719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // f32 21819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman numRegs = MFI->getNumRegistersForClass(PTX::RegF32RegisterClass); 21919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if(numRegs > 0) { 22019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += "\t.reg .f32 %f<"; 22119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += utostr(numRegs); 22219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += ">;\n"; 22319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 22419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 22519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // f64 22619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman numRegs = MFI->getNumRegistersForClass(PTX::RegF64RegisterClass); 22719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if(numRegs > 0) { 22819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += "\t.reg .f64 %fd<"; 22919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += utostr(numRegs); 23019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += ">;\n"; 23119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 23219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 23319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Local params 23419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (PTXParamManager::param_iterator i = PM.local_begin(), e = PM.local_end(); 23519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman i != e; ++i) { 23619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += "\t.param .b"; 23719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += utostr(PM.getParamSize(*i)); 23819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += " "; 23919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += PM.getParamName(*i); 24019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman regDefs += ";\n"; 24119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 24219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 24319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine(regDefs)); 24419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 24519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 24619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const MachineFrameInfo* FrameInfo = MF->getFrameInfo(); 24719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman DEBUG(dbgs() << "Have " << FrameInfo->getNumObjects() 24819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman << " frame object(s)\n"); 24919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (unsigned i = 0, e = FrameInfo->getNumObjects(); i != e; ++i) { 25019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman DEBUG(dbgs() << "Size of object: " << FrameInfo->getObjectSize(i) << "\n"); 25119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (FrameInfo->getObjectSize(i) > 0) { 25219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman std::string def = "\t.local .align "; 25319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += utostr(FrameInfo->getObjectAlignment(i)); 25419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += " .b8"; 25519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += " __local"; 25619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += utostr(i); 25719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += "["; 25819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += utostr(FrameInfo->getObjectSize(i)); // Convert to bits 25919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += "]"; 26019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman def += ";"; 26119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine(def)); 26219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 26319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 26419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 26519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman //unsigned Index = 1; 26619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Print parameter passing params 26719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman //for (PTXMachineFunctionInfo::param_iterator 26819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // i = MFI->paramBegin(), e = MFI->paramEnd(); i != e; ++i) { 26919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // std::string def = "\t.param .b"; 27019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // def += utostr(*i); 27119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // def += " __ret_"; 27219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // def += utostr(Index); 27319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Index++; 27419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // def += ";"; 27519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // OutStreamer.EmitRawText(Twine(def)); 27619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman //} 27719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 27819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 27919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanvoid PTXAsmPrinter::EmitFunctionBodyEnd() { 28019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine("}")); 28119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 28219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 28319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanvoid PTXAsmPrinter::EmitInstruction(const MachineInstr *MI) { 28419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCInst TmpInst; 28519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman LowerPTXMachineInstrToMCInst(MI, TmpInst, *this); 28619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitInstruction(TmpInst); 28719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 28819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 28919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanvoid PTXAsmPrinter::EmitVariableDeclaration(const GlobalVariable *gv) { 29019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Check to see if this is a special global used by LLVM, if so, emit it. 29119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (EmitSpecialLLVMGlobal(gv)) 29219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return; 29319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 29419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCSymbol *gvsym = Mang->getSymbol(gv); 29519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 29619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman assert(gvsym->isUndefined() && "Cannot define a symbol twice!"); 29719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 29819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman std::string decl; 29919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 30019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // check if it is defined in some other translation unit 30119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (gv->isDeclaration()) 30219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".extern "; 30319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 30419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // state space: e.g., .global 30519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "."; 30619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += getStateSpaceName(gv->getType()->getAddressSpace()); 30719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 30819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 30919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // alignment (optional) 31019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman unsigned alignment = gv->getAlignment(); 31119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (alignment != 0) { 31219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".align "; 31319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += utostr(gv->getAlignment()); 31419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 31519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 31619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 31719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 31819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (PointerType::classof(gv->getType())) { 31919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman PointerType* pointerTy = dyn_cast<PointerType>(gv->getType()); 32019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman Type* elementTy = pointerTy->getElementType(); 32119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 32219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".b8 "; 32319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += gvsym->getName(); 32419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "["; 32519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 32619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (elementTy->isArrayTy()) 32719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman { 32819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman assert(elementTy->isArrayTy() && "Only pointers to arrays are supported"); 32919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 33019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman ArrayType* arrayTy = dyn_cast<ArrayType>(elementTy); 33119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman elementTy = arrayTy->getElementType(); 33219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 33319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman unsigned numElements = arrayTy->getNumElements(); 33419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 33519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman while (elementTy->isArrayTy()) { 33619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 33719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman arrayTy = dyn_cast<ArrayType>(elementTy); 33819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman elementTy = arrayTy->getElementType(); 33919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 34019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman numElements *= arrayTy->getNumElements(); 34119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 34219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 34319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // FIXME: isPrimitiveType() == false for i16? 34419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman assert(elementTy->isSingleValueType() && 34519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman "Non-primitive types are not handled"); 34619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 34719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Compute the size of the array, in bytes. 34819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman uint64_t arraySize = (elementTy->getPrimitiveSizeInBits() >> 3) 34919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman * numElements; 35019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 35119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += utostr(arraySize); 35219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 35319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 35419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "]"; 35519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 35619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // handle string constants (assume ConstantArray means string) 35719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 35819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (gv->hasInitializer()) 35919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman { 36019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const Constant *C = gv->getInitializer(); 36119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (const ConstantArray *CA = dyn_cast<ConstantArray>(C)) 36219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman { 36319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " = {"; 36419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 36519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (unsigned i = 0, e = C->getNumOperands(); i != e; ++i) 36619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman { 36719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (i > 0) decl += ","; 36819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 36919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "0x" + 37019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman utohexstr(cast<ConstantInt>(CA->getOperand(i))->getZExtValue()); 37119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 37219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 37319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "}"; 37419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 37519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 37619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 37719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman else { 37819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Note: this is currently the fall-through case and most likely generates 37919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // incorrect code. 38019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += getTypeName(gv->getType()); 38119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 38219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 38319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += gvsym->getName(); 38419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 38519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (ArrayType::classof(gv->getType()) || 38619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman PointerType::classof(gv->getType())) 38719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "[]"; 38819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 38919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 39019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ";"; 39119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 39219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine(decl)); 39319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 39419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.AddBlankLine(); 39519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 39619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 39719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanvoid PTXAsmPrinter::EmitFunctionEntryLabel() { 39819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // The function label could have already been emitted if two symbols end up 39919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // conflicting due to asm renaming. Detect this and emit an error. 40019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (!CurrentFnSym->isUndefined()) { 40119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman report_fatal_error("'" + Twine(CurrentFnSym->getName()) + 40219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman "' label emitted multiple times to assembly file"); 40319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return; 40419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 40519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 40619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>(); 40719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PTXParamManager &PM = MFI->getParamManager(); 40819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const bool isKernel = MFI->isKernel(); 40919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>(); 41019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const MachineRegisterInfo& MRI = MF->getRegInfo(); 41119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 41219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman std::string decl = isKernel ? ".entry" : ".func"; 41319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 41419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (!isKernel) { 41519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " ("; 41619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (ST.useParamSpaceForDeviceArgs()) { 41719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (PTXParamManager::param_iterator i = PM.ret_begin(), e = PM.ret_end(), 41819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman b = i; i != e; ++i) { 41919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (i != b) { 42019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ", "; 42119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 42219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 42319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".param .b"; 42419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += utostr(PM.getParamSize(*i)); 42519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 42619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += PM.getParamName(*i); 42719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 42819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } else { 42919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (PTXMachineFunctionInfo::reg_iterator 43019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman i = MFI->retreg_begin(), e = MFI->retreg_end(), b = i; 43119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman i != e; ++i) { 43219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (i != b) { 43319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ", "; 43419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 43519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".reg ."; 43619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += getRegisterTypeName(*i, MRI); 43719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 43819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += MFI->getRegisterName(*i); 43919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 44019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 44119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ")"; 44219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 44319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 44419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Print function name 44519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 44619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += CurrentFnSym->getName().str(); 44719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 44819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " ("; 44919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 45019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const Function *F = MF->getFunction(); 45119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 45219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Print parameters 45319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (isKernel || ST.useParamSpaceForDeviceArgs()) { 45419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman /*for (PTXParamManager::param_iterator i = PM.arg_begin(), e = PM.arg_end(), 45519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman b = i; i != e; ++i) { 45619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (i != b) { 45719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ", "; 45819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 45919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 46019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".param .b"; 46119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += utostr(PM.getParamSize(*i)); 46219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 46319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += PM.getParamName(*i); 46419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman }*/ 46519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman int Counter = 1; 46619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (Function::const_arg_iterator i = F->arg_begin(), e = F->arg_end(), 46719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman b = i; i != e; ++i) { 46819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (i != b) 46919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ", "; 47019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const Type *ArgType = (*i).getType(); 47119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".param .b"; 47219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (ArgType->isPointerTy()) { 47319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (ST.is64Bit()) 47419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "64"; 47519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman else 47619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += "32"; 47719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } else { 47819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += utostr(ArgType->getPrimitiveSizeInBits()); 47919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 48019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (ArgType->isPointerTy() && ST.emitPtrAttribute()) { 48119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PointerType *PtrType = dyn_cast<const PointerType>(ArgType); 48219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " .ptr"; 48319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman switch (PtrType->getAddressSpace()) { 48419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman default: 48519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman llvm_unreachable("Unknown address space in argument"); 48619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case PTXStateSpace::Global: 48719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " .global"; 48819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 48919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case PTXStateSpace::Shared: 49019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " .shared"; 49119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 49219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 49319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 49419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " __param_"; 49519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += utostr(Counter++); 49619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 49719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } else { 49819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman for (PTXMachineFunctionInfo::reg_iterator 49919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman i = MFI->argreg_begin(), e = MFI->argreg_end(), b = i; 50019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman i != e; ++i) { 50119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (i != b) { 50219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ", "; 50319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 50419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 50519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ".reg ."; 50619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += getRegisterTypeName(*i, MRI); 50719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += " "; 50819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += MFI->getRegisterName(*i); 50919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 51019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 51119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman decl += ")"; 51219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 51319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitRawText(Twine(decl)); 51419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 51519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 51619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanunsigned PTXAsmPrinter::GetOrCreateSourceID(StringRef FileName, 51719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman StringRef DirName) { 51819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // If FE did not provide a file name, then assume stdin. 51919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (FileName.empty()) 52019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return GetOrCreateSourceID("<stdin>", StringRef()); 52119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 52219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // MCStream expects full path name as filename. 52319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (!DirName.empty() && !sys::path::is_absolute(FileName)) { 52419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman SmallString<128> FullPathName = DirName; 52519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman sys::path::append(FullPathName, FileName); 52619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Here FullPathName will be copied into StringMap by GetOrCreateSourceID. 52719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return GetOrCreateSourceID(StringRef(FullPathName), StringRef()); 52819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 52919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 53019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman StringMapEntry<unsigned> &Entry = SourceIdMap.GetOrCreateValue(FileName); 53119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman if (Entry.getValue()) 53219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return Entry.getValue(); 53319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 53419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman unsigned SrcId = SourceIdMap.size(); 53519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman Entry.setValue(SrcId); 53619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 53719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // Print out a .file directive to specify files for .loc directives. 53819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutStreamer.EmitDwarfFileDirective(SrcId, Entry.getKey()); 53919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 54019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return SrcId; 54119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 54219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 54319bac1e08be200c31efd26f0f5fd144c9b3eefd3John BaumanMCOperand PTXAsmPrinter::GetSymbolRef(const MachineOperand &MO, 54419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const MCSymbol *Symbol) { 54519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const MCExpr *Expr; 54619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman Expr = MCSymbolRefExpr::Create(Symbol, MCSymbolRefExpr::VK_None, OutContext); 54719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return MCOperand::CreateExpr(Expr); 54819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 54919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 55019bac1e08be200c31efd26f0f5fd144c9b3eefd3John BaumanMCOperand PTXAsmPrinter::lowerOperand(const MachineOperand &MO) { 55119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCOperand MCOp; 55219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>(); 55319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const MCExpr *Expr; 55419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman const char *RegSymbolName; 55519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman switch (MO.getType()) { 55619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman default: 55719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman llvm_unreachable("Unknown operand type"); 55819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case MachineOperand::MO_Register: 55919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // We create register operands as symbols, since the PTXInstPrinter class 56019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // has no way to map virtual registers back to a name without some ugly 56119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // hacks. 56219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman // FIXME: Figure out a better way to handle virtual register naming. 56319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman RegSymbolName = MFI->getRegisterName(MO.getReg()); 56419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman Expr = MCSymbolRefExpr::Create(RegSymbolName, MCSymbolRefExpr::VK_None, 56519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman OutContext); 56619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCOp = MCOperand::CreateExpr(Expr); 56719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 56819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case MachineOperand::MO_Immediate: 56919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCOp = MCOperand::CreateImm(MO.getImm()); 57019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 57119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case MachineOperand::MO_MachineBasicBlock: 57219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCOp = MCOperand::CreateExpr(MCSymbolRefExpr::Create( 57319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MO.getMBB()->getSymbol(), OutContext)); 57419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 57519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case MachineOperand::MO_GlobalAddress: 57619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCOp = GetSymbolRef(MO, Mang->getSymbol(MO.getGlobal())); 57719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 57819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case MachineOperand::MO_ExternalSymbol: 57919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCOp = GetSymbolRef(MO, GetExternalSymbolSymbol(MO.getSymbolName())); 58019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 58119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman case MachineOperand::MO_FPImmediate: 58219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman APFloat Val = MO.getFPImm()->getValueAPF(); 58319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman bool ignored; 58419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman Val.convert(APFloat::IEEEdouble, APFloat::rmTowardZero, &ignored); 58519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman MCOp = MCOperand::CreateFPImm(Val.convertToDouble()); 58619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman break; 58719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman } 58819bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 58919bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman return MCOp; 59019bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 59119bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 59219bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman// Force static initialization. 59319bac1e08be200c31efd26f0f5fd144c9b3eefd3John Baumanextern "C" void LLVMInitializePTXAsmPrinter() { 59419bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman RegisterAsmPrinter<PTXAsmPrinter> X(ThePTX32Target); 59519bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman RegisterAsmPrinter<PTXAsmPrinter> Y(ThePTX64Target); 59619bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman} 59719bac1e08be200c31efd26f0f5fd144c9b3eefd3John Bauman 598