1 files changed, 154 insertions, 190 deletions

diff --git a/contrib/llvm/lib/Target/PTX/PTXAsmPrinter.cpp b/contrib/llvm/lib/Target/PTX/PTXAsmPrinter.cpp
index 733744b..0b6ac7b 100644
--- a/contrib/llvm/lib/Target/PTX/PTXAsmPrinter.cpp
+++ b/contrib/llvm/lib/Target/PTX/PTXAsmPrinter.cpp
@@ -14,8 +14,8 @@
 
 #define DEBUG_TYPE "ptx-asm-printer"
 
-#include "PTX.h"
 #include "PTXAsmPrinter.h"
+#include "PTX.h"
 #include "PTXMachineFunctionInfo.h"
 #include "PTXParamManager.h"
 #include "PTXRegisterInfo.h"
@@ -25,7 +25,6 @@
 #include "llvm/Function.h"
 #include "llvm/Module.h"
 #include "llvm/ADT/SmallString.h"
-#include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/Twine.h"
 #include "llvm/Analysis/DebugInfo.h"
 #include "llvm/CodeGen/AsmPrinter.h"
@@ -52,23 +51,23 @@ using namespace llvm;
 static const char PARAM_PREFIX[] = "__param_";
 static const char RETURN_PREFIX[] = "__ret_";
 
-static const char *getRegisterTypeName(unsigned RegNo,
-                                       const MachineRegisterInfo& MRI) {
-  const TargetRegisterClass *TRC = MRI.getRegClass(RegNo);
-
-#define TEST_REGCLS(cls, clsstr) \
-  if (PTX::cls ## RegisterClass == TRC) return # clsstr;
-
-  TEST_REGCLS(RegPred, pred);
-  TEST_REGCLS(RegI16, b16);
-  TEST_REGCLS(RegI32, b32);
-  TEST_REGCLS(RegI64, b64);
-  TEST_REGCLS(RegF32, b32);
-  TEST_REGCLS(RegF64, b64);
-#undef TEST_REGCLS
-
-  llvm_unreachable("Not in any register class!");
-  return NULL;
+static const char *getRegisterTypeName(unsigned RegType) {
+  switch (RegType) {
+  default:
+    llvm_unreachable("Unknown register type");
+  case PTXRegisterType::Pred:
+    return ".pred";
+  case PTXRegisterType::B16:
+    return ".b16";
+  case PTXRegisterType::B32:
+    return ".b32";
+  case PTXRegisterType::B64:
+    return ".b64";
+  case PTXRegisterType::F32:
+    return ".f32";
+  case PTXRegisterType::F64:
+    return ".f64";
+  }
 }
 
 static const char *getStateSpaceName(unsigned addressSpace) {
@@ -80,7 +79,6 @@ static const char *getStateSpaceName(unsigned addressSpace) {
   case PTXStateSpace::Parameter: return "param";
   case PTXStateSpace::Shared:    return "shared";
   }
-  return NULL;
 }
 
 static const char *getTypeName(Type* type) {
@@ -139,15 +137,15 @@ void PTXAsmPrinter::EmitStartOfAsmFile(Module &M)
   const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>();
 
   // Emit the PTX .version and .target attributes
-  OutStreamer.EmitRawText(Twine("\t.version " + ST.getPTXVersionString()));
-  OutStreamer.EmitRawText(Twine("\t.target " + ST.getTargetString() +
+  OutStreamer.EmitRawText(Twine("\t.version ") + ST.getPTXVersionString());
+  OutStreamer.EmitRawText(Twine("\t.target ") + ST.getTargetString() +
                                 (ST.supportsDouble() ? ""
-                                                     : ", map_f64_to_f32")));
+                                                     : ", map_f64_to_f32"));
   // .address_size directive is optional, but it must immediately follow
   // the .target directive if present within a module
   if (ST.supportsPTX23()) {
-    std::string addrSize = ST.is64Bit() ? "64" : "32";
-    OutStreamer.EmitRawText(Twine("\t.address_size " + addrSize));
+    const char *addrSize = ST.is64Bit() ? "64" : "32";
+    OutStreamer.EmitRawText(Twine("\t.address_size ") + addrSize);
   }
 
   OutStreamer.AddBlankLine();
@@ -166,6 +164,11 @@ void PTXAsmPrinter::EmitStartOfAsmFile(Module &M)
 
   OutStreamer.AddBlankLine();
 
+  // declare external functions
+  for (Module::const_iterator i = M.begin(), e = M.end();
+       i != e; ++i)
+    EmitFunctionDeclaration(i);
+  
   // declare global variables
   for (Module::const_global_iterator i = M.global_begin(), e = M.global_end();
        i != e; ++i)
@@ -179,68 +182,47 @@ void PTXAsmPrinter::EmitFunctionBodyStart() {
   const PTXParamManager &PM = MFI->getParamManager();
 
   // Print register definitions
-  std::string regDefs;
+  SmallString<128> regDefs;
+  raw_svector_ostream os(regDefs);
   unsigned numRegs;
 
   // pred
-  numRegs = MFI->getNumRegistersForClass(PTX::RegPredRegisterClass);
-  if(numRegs > 0) {
-    regDefs += "\t.reg .pred %p<";
-    regDefs += utostr(numRegs);
-    regDefs += ">;\n";
-  }
+  numRegs = MFI->countRegisters(PTXRegisterType::Pred, PTXRegisterSpace::Reg);
+  if(numRegs > 0)
+    os << "\t.reg .pred %p<" << numRegs << ">;\n";
 
   // i16
-  numRegs = MFI->getNumRegistersForClass(PTX::RegI16RegisterClass);
-  if(numRegs > 0) {
-    regDefs += "\t.reg .b16 %rh<";
-    regDefs += utostr(numRegs);
-    regDefs += ">;\n";
-  }
+  numRegs = MFI->countRegisters(PTXRegisterType::B16, PTXRegisterSpace::Reg);
+  if(numRegs > 0)
+    os << "\t.reg .b16 %rh<" << numRegs << ">;\n";
 
   // i32
-  numRegs = MFI->getNumRegistersForClass(PTX::RegI32RegisterClass);
-  if(numRegs > 0) {
-    regDefs += "\t.reg .b32 %r<";
-    regDefs += utostr(numRegs);
-    regDefs += ">;\n";
-  }
+  numRegs = MFI->countRegisters(PTXRegisterType::B32, PTXRegisterSpace::Reg);
+  if(numRegs > 0)
+    os << "\t.reg .b32 %r<" << numRegs << ">;\n";
 
   // i64
-  numRegs = MFI->getNumRegistersForClass(PTX::RegI64RegisterClass);
-  if(numRegs > 0) {
-    regDefs += "\t.reg .b64 %rd<";
-    regDefs += utostr(numRegs);
-    regDefs += ">;\n";
-  }
+  numRegs = MFI->countRegisters(PTXRegisterType::B64, PTXRegisterSpace::Reg);
+  if(numRegs > 0)
+    os << "\t.reg .b64 %rd<" << numRegs << ">;\n";
 
   // f32
-  numRegs = MFI->getNumRegistersForClass(PTX::RegF32RegisterClass);
-  if(numRegs > 0) {
-    regDefs += "\t.reg .f32 %f<";
-    regDefs += utostr(numRegs);
-    regDefs += ">;\n";
-  }
+  numRegs = MFI->countRegisters(PTXRegisterType::F32, PTXRegisterSpace::Reg);
+  if(numRegs > 0)
+    os << "\t.reg .f32 %f<" << numRegs << ">;\n";
 
   // f64
-  numRegs = MFI->getNumRegistersForClass(PTX::RegF64RegisterClass);
-  if(numRegs > 0) {
-    regDefs += "\t.reg .f64 %fd<";
-    regDefs += utostr(numRegs);
-    regDefs += ">;\n";
-  }
+  numRegs = MFI->countRegisters(PTXRegisterType::F64, PTXRegisterSpace::Reg);
+  if(numRegs > 0)
+    os << "\t.reg .f64 %fd<" << numRegs << ">;\n";
 
   // Local params
   for (PTXParamManager::param_iterator i = PM.local_begin(), e = PM.local_end();
-       i != e; ++i) {
-    regDefs += "\t.param .b";
-    regDefs += utostr(PM.getParamSize(*i));
-    regDefs += " ";
-    regDefs += PM.getParamName(*i);
-    regDefs += ";\n";
-  }
+       i != e; ++i)
+    os << "\t.param .b" << PM.getParamSize(*i) << ' ' << PM.getParamName(*i)
+       << ";\n";
 
-  OutStreamer.EmitRawText(Twine(regDefs));
+  OutStreamer.EmitRawText(os.str());
 
 
   const MachineFrameInfo* FrameInfo = MF->getFrameInfo();
@@ -249,16 +231,13 @@ void PTXAsmPrinter::EmitFunctionBodyStart() {
   for (unsigned i = 0, e = FrameInfo->getNumObjects(); i != e; ++i) {
     DEBUG(dbgs() << "Size of object: " << FrameInfo->getObjectSize(i) << "\n");
     if (FrameInfo->getObjectSize(i) > 0) {
-      std::string def = "\t.local .align ";
-      def += utostr(FrameInfo->getObjectAlignment(i));
-      def += " .b8";
-      def += " __local";
-      def += utostr(i);
-      def += "[";
-      def += utostr(FrameInfo->getObjectSize(i)); // Convert to bits
-      def += "]";
-      def += ";";
-      OutStreamer.EmitRawText(Twine(def));
+      OutStreamer.EmitRawText("\t.local .align " +
+                              Twine(FrameInfo->getObjectAlignment(i)) +
+                              " .b8 __local" +
+                              Twine(i) +
+                              "[" +
+                              Twine(FrameInfo->getObjectSize(i)) +
+                              "];");
     }
   }
 
@@ -295,36 +274,27 @@ void PTXAsmPrinter::EmitVariableDeclaration(const GlobalVariable *gv) {
 
   assert(gvsym->isUndefined() && "Cannot define a symbol twice!");
 
-  std::string decl;
+  SmallString<128> decl;
+  raw_svector_ostream os(decl);
 
   // check if it is defined in some other translation unit
   if (gv->isDeclaration())
-    decl += ".extern ";
+    os << ".extern ";
 
   // state space: e.g., .global
-  decl += ".";
-  decl += getStateSpaceName(gv->getType()->getAddressSpace());
-  decl += " ";
+  os << '.' << getStateSpaceName(gv->getType()->getAddressSpace()) << ' ';
 
   // alignment (optional)
   unsigned alignment = gv->getAlignment();
-  if (alignment != 0) {
-    decl += ".align ";
-    decl += utostr(gv->getAlignment());
-    decl += " ";
-  }
+  if (alignment != 0)
+    os << ".align " << gv->getAlignment() << ' ';
 
 
   if (PointerType::classof(gv->getType())) {
     PointerType* pointerTy = dyn_cast<PointerType>(gv->getType());
     Type* elementTy = pointerTy->getElementType();
 
-    decl += ".b8 ";
-    decl += gvsym->getName();
-    decl += "[";
-
-    if (elementTy->isArrayTy())
-    {
+    if (elementTy->isArrayTy()) {
       assert(elementTy->isArrayTy() && "Only pointers to arrays are supported");
 
       ArrayType* arrayTy = dyn_cast<ArrayType>(elementTy);
@@ -333,7 +303,6 @@ void PTXAsmPrinter::EmitVariableDeclaration(const GlobalVariable *gv) {
       unsigned numElements = arrayTy->getNumElements();
 
       while (elementTy->isArrayTy()) {
-
         arrayTy = dyn_cast<ArrayType>(elementTy);
         elementTy = arrayTy->getElementType();
 
@@ -342,110 +311,91 @@ void PTXAsmPrinter::EmitVariableDeclaration(const GlobalVariable *gv) {
 
       // FIXME: isPrimitiveType() == false for i16?
       assert(elementTy->isSingleValueType() &&
-              "Non-primitive types are not handled");
+             "Non-primitive types are not handled");
 
-      // Compute the size of the array, in bytes.
-      uint64_t arraySize = (elementTy->getPrimitiveSizeInBits() >> 3)
-                        * numElements;
+      // Find the size of the element in bits
+      unsigned elementSize = elementTy->getPrimitiveSizeInBits();
 
-      decl += utostr(arraySize);
+      os << ".b" << elementSize << ' ' << gvsym->getName()
+         << '[' << numElements << ']';
+    } else {
+      os << ".b8" << gvsym->getName() << "[]";
     }
 
-    decl += "]";
-
     // handle string constants (assume ConstantArray means string)
-
-    if (gv->hasInitializer())
-    {
+    if (gv->hasInitializer()) {
       const Constant *C = gv->getInitializer();
-      if (const ConstantArray *CA = dyn_cast<ConstantArray>(C))
-      {
-        decl += " = {";
+      if (const ConstantArray *CA = dyn_cast<ConstantArray>(C)) {
+        os << " = {";
 
-        for (unsigned i = 0, e = C->getNumOperands(); i != e; ++i)
-        {
-          if (i > 0)   decl += ",";
+        for (unsigned i = 0, e = C->getNumOperands(); i != e; ++i) {
+          if (i > 0)
+            os << ',';
 
-          decl += "0x" +
-                utohexstr(cast<ConstantInt>(CA->getOperand(i))->getZExtValue());
+          os << "0x";
+          os.write_hex(cast<ConstantInt>(CA->getOperand(i))->getZExtValue());
         }
 
-        decl += "}";
+        os << '}';
       }
     }
-  }
-  else {
+  } else {
     // Note: this is currently the fall-through case and most likely generates
     //       incorrect code.
-    decl += getTypeName(gv->getType());
-    decl += " ";
-
-    decl += gvsym->getName();
+    os << getTypeName(gv->getType()) << ' ' << gvsym->getName();
 
-    if (ArrayType::classof(gv->getType()) ||
-        PointerType::classof(gv->getType()))
-      decl += "[]";
+    if (isa<ArrayType>(gv->getType()) || isa<PointerType>(gv->getType()))
+      os << "[]";
   }
 
-  decl += ";";
-
-  OutStreamer.EmitRawText(Twine(decl));
+  os << ';';
 
+  OutStreamer.EmitRawText(os.str());
   OutStreamer.AddBlankLine();
 }
 
 void PTXAsmPrinter::EmitFunctionEntryLabel() {
   // The function label could have already been emitted if two symbols end up
   // conflicting due to asm renaming.  Detect this and emit an error.
-  if (!CurrentFnSym->isUndefined()) {
+  if (!CurrentFnSym->isUndefined())
     report_fatal_error("'" + Twine(CurrentFnSym->getName()) +
                        "' label emitted multiple times to assembly file");
-    return;
-  }
 
   const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>();
   const PTXParamManager &PM = MFI->getParamManager();
   const bool isKernel = MFI->isKernel();
   const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>();
-  const MachineRegisterInfo& MRI = MF->getRegInfo();
 
-  std::string decl = isKernel ? ".entry" : ".func";
+  SmallString<128> decl;
+  raw_svector_ostream os(decl);
+  os << (isKernel ? ".entry" : ".func");
 
   if (!isKernel) {
-    decl += " (";
+    os << " (";
     if (ST.useParamSpaceForDeviceArgs()) {
       for (PTXParamManager::param_iterator i = PM.ret_begin(), e = PM.ret_end(),
            b = i; i != e; ++i) {
-        if (i != b) {
-          decl += ", ";
-        }
+        if (i != b)
+          os << ", ";
 
-        decl += ".param .b";
-        decl += utostr(PM.getParamSize(*i));
-        decl += " ";
-        decl += PM.getParamName(*i);
+        os << ".param .b" << PM.getParamSize(*i) << ' ' << PM.getParamName(*i);
       }
     } else {
       for (PTXMachineFunctionInfo::reg_iterator
            i = MFI->retreg_begin(), e = MFI->retreg_end(), b = i;
            i != e; ++i) {
-        if (i != b) {
-          decl += ", ";
-        }
-        decl += ".reg .";
-        decl += getRegisterTypeName(*i, MRI);
-        decl += " ";
-        decl += MFI->getRegisterName(*i);
+        if (i != b)
+          os << ", ";
+
+        os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' '
+           << MFI->getRegisterName(*i);
       }
     }
-    decl += ")";
+    os << ')';
   }
 
   // Print function name
-  decl += " ";
-  decl += CurrentFnSym->getName().str();
-
-  decl += " (";
+  os << ' ' << CurrentFnSym->getName() << " (";
 
   const Function *F = MF->getFunction();
 
@@ -453,63 +403,80 @@ void PTXAsmPrinter::EmitFunctionEntryLabel() {
   if (isKernel || ST.useParamSpaceForDeviceArgs()) {
     /*for (PTXParamManager::param_iterator i = PM.arg_begin(), e = PM.arg_end(),
          b = i; i != e; ++i) {
-      if (i != b) {
-        decl += ", ";
-      }
+      if (i != b)
+        os << ", ";
 
-      decl += ".param .b";
-      decl += utostr(PM.getParamSize(*i));
-      decl += " ";
-      decl += PM.getParamName(*i);
+      os << ".param .b" << PM.getParamSize(*i) << ' ' << PM.getParamName(*i);
     }*/
     int Counter = 1;
     for (Function::const_arg_iterator i = F->arg_begin(), e = F->arg_end(),
          b = i; i != e; ++i) {
       if (i != b)
-        decl += ", ";
+        os << ", ";
       const Type *ArgType = (*i).getType();
-      decl += ".param .b";
+      os << ".param .b";
       if (ArgType->isPointerTy()) {
         if (ST.is64Bit())
-          decl += "64";
+          os << "64";
         else
-          decl += "32";
+          os << "32";
       } else {
-        decl += utostr(ArgType->getPrimitiveSizeInBits());
+        os << ArgType->getPrimitiveSizeInBits();
       }
       if (ArgType->isPointerTy() && ST.emitPtrAttribute()) {
         const PointerType *PtrType = dyn_cast<const PointerType>(ArgType);
-        decl += " .ptr";
+        os << " .ptr";
         switch (PtrType->getAddressSpace()) {
         default:
           llvm_unreachable("Unknown address space in argument");
         case PTXStateSpace::Global:
-          decl += " .global";
+          os << " .global";
           break;
         case PTXStateSpace::Shared:
-          decl += " .shared";
+          os << " .shared";
           break;
         }
       }
-      decl += " __param_";
-      decl += utostr(Counter++);
+      os << " __param_" << Counter++;
     }
   } else {
     for (PTXMachineFunctionInfo::reg_iterator
          i = MFI->argreg_begin(), e = MFI->argreg_end(), b = i;
          i != e; ++i) {
-      if (i != b) {
-        decl += ", ";
-      }
+      if (i != b)
+        os << ", ";
 
-      decl += ".reg .";
-      decl += getRegisterTypeName(*i, MRI);
-      decl += " ";
-      decl += MFI->getRegisterName(*i);
+      os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' '
+         << MFI->getRegisterName(*i);
     }
   }
-  decl += ")";
+  os << ')';
 
+  OutStreamer.EmitRawText(os.str());
+}
+
+void PTXAsmPrinter::EmitFunctionDeclaration(const Function* func)
+{
+  const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>();
+	
+  std::string decl = "";
+
+  // hard-coded emission of extern vprintf function 
+  
+  if (func->getName() == "printf" || func->getName() == "puts") {		
+    decl += ".extern .func (.param .b32 __param_1) vprintf (.param .b";
+    if (ST.is64Bit())	
+      decl += "64";
+    else				
+      decl += "32";
+    decl += " __param_2, .param .b";
+    if (ST.is64Bit())	
+      decl += "64";
+    else				
+      decl += "32";
+    decl += " __param_3)\n";
+  }
+  
   OutStreamer.EmitRawText(Twine(decl));
 }
 
@@ -535,7 +502,7 @@ unsigned PTXAsmPrinter::GetOrCreateSourceID(StringRef FileName,
   Entry.setValue(SrcId);
 
   // Print out a .file directive to specify files for .loc directives.
-  OutStreamer.EmitDwarfFileDirective(SrcId, Entry.getKey());
+  OutStreamer.EmitDwarfFileDirective(SrcId, "", Entry.getKey());
 
   return SrcId;
 }
@@ -550,20 +517,18 @@ MCOperand PTXAsmPrinter::GetSymbolRef(const MachineOperand &MO,
 MCOperand PTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
   MCOperand MCOp;
   const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>();
-  const MCExpr *Expr;
-  const char *RegSymbolName;
+  unsigned EncodedReg;
   switch (MO.getType()) {
   default:
     llvm_unreachable("Unknown operand type");
   case MachineOperand::MO_Register:
-    // We create register operands as symbols, since the PTXInstPrinter class
-    // has no way to map virtual registers back to a name without some ugly
-    // hacks.
-    // FIXME: Figure out a better way to handle virtual register naming.
-    RegSymbolName = MFI->getRegisterName(MO.getReg());
-    Expr = MCSymbolRefExpr::Create(RegSymbolName, MCSymbolRefExpr::VK_None,
-                                   OutContext);
-    MCOp = MCOperand::CreateExpr(Expr);
+    if (MO.getReg() > 0) {
+      // Encode the register
+      EncodedReg = MFI->getEncodedRegister(MO.getReg());
+    } else {
+      EncodedReg = 0;
+    }
+    MCOp = MCOperand::CreateReg(EncodedReg);
     break;
   case MachineOperand::MO_Immediate:
     MCOp = MCOperand::CreateImm(MO.getImm());
@@ -594,4 +559,3 @@ extern "C" void LLVMInitializePTXAsmPrinter() {
   RegisterAsmPrinter<PTXAsmPrinter> X(ThePTX32Target);
   RegisterAsmPrinter<PTXAsmPrinter> Y(ThePTX64Target);
 }
-