Simplify code. No functionality change.

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@144012 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Target/PTX/PTXAsmPrinter.cpp b/lib/Target/PTX/PTXAsmPrinter.cpp
index e329d5dcc7d006de9d818216cf08be42ed44408e..45a6afc8587d2c1d8246da005ad06e4ea53d7530 100644 (file)
--- a/lib/Target/PTX/PTXAsmPrinter.cpp
+++ b/lib/Target/PTX/PTXAsmPrinter.cpp
@@ -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"
@@ -139,15 +138,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();
@@ -179,68 +178,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";
-  }
+  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";
-  }
+  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";
-  }
+  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";
-  }
+  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";
-  }
+  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";
-  }
+  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 +227,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,32 +270,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();
-        
-    if (elementTy->isArrayTy())
-    {
+
+    if (elementTy->isArrayTy()) {
       assert(elementTy->isArrayTy() && "Only pointers to arrays are supported");
 
       ArrayType* arrayTy = dyn_cast<ArrayType>(elementTy);
@@ -329,7 +299,6 @@ void PTXAsmPrinter::EmitVariableDeclaration(const GlobalVariable *gv) {
       unsigned numElements = arrayTy->getNumElements();
 
       while (elementTy->isArrayTy()) {
-
         arrayTy = dyn_cast<ArrayType>(elementTy);
         elementTy = arrayTy->getElementType();
 
@@ -338,64 +307,46 @@ 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");
+
       // Find the size of the element in bits
       unsigned elementSize = elementTy->getPrimitiveSizeInBits();
 
-      decl += ".b";
-      decl += utostr(elementSize);
-      decl += " ";
-      decl += gvsym->getName();
-      decl += "[";
-      decl += utostr(numElements);
-      decl += "]";
-    }
-    else
-    {
-      decl += ".b8 ";
-      decl += gvsym->getName();
-      decl += "[]";
+      os << ".b" << elementSize << ' ' << gvsym->getName()
+         << '[' << numElements << ']';
+    } else {
+      os << ".b8" << gvsym->getName() << "[]";
     }
 
     // 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();
 }
 
@@ -414,43 +365,36 @@ void PTXAsmPrinter::EmitFunctionEntryLabel() {
   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(*i, MRI) << ' '
+           << MFI->getRegisterName(*i);
       }
     }
-    decl += ")";
+    os << ')';
   }
 
   // Print function name
-  decl += " ";
-  decl += CurrentFnSym->getName().str();
-
-  decl += " (";
+  os << ' ' << CurrentFnSym->getName() << " (";
 
   const Function *F = MF->getFunction();
 
@@ -458,64 +402,56 @@ 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(*i, MRI) << ' '
+         << MFI->getRegisterName(*i);
     }
   }
-  decl += ")";
+  os << ')';
 
-  OutStreamer.EmitRawText(Twine(decl));
+  OutStreamer.EmitRawText(os.str());
 }
 
 unsigned PTXAsmPrinter::GetOrCreateSourceID(StringRef FileName,