Replace PTX back-end with NVPTX back-end in all places where Clang cares

NV_CONTRIB

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

10 files changed, 66 insertions, 176 deletions

diff --git a/include/clang/Basic/BuiltinsPTX.def b/include/clang/Basic/BuiltinsNVPTX.def
index f90a43f7f4..f90a43f7f4 100644
--- a/include/clang/Basic/BuiltinsPTX.def
+++ b/include/clang/Basic/BuiltinsNVPTX.def
diff --git a/include/clang/Basic/TargetBuiltins.h b/include/clang/Basic/TargetBuiltins.h
index 7c04bf7edf..e9b9f85521 100644
--- a/include/clang/Basic/TargetBuiltins.h
+++ b/include/clang/Basic/TargetBuiltins.h
@@ -35,12 +35,12 @@ namespace clang {
     };
   }
 
-  /// PTX builtins
-  namespace PTX {
+  /// NVPTX builtins
+  namespace NVPTX {
     enum {
         LastTIBuiltin = clang::Builtin::FirstTSBuiltin-1,
 #define BUILTIN(ID, TYPE, ATTRS) BI##ID,
-#include "clang/Basic/BuiltinsPTX.def"
+#include "clang/Basic/BuiltinsNVPTX.def"
         LastTSBuiltin
     };
   }
diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp
index e9752700f0..e38ca550c3 100644
--- a/lib/Basic/Targets.cpp
+++ b/lib/Basic/Targets.cpp
@@ -946,57 +946,39 @@ public:
 } // end anonymous namespace.
 
 namespace {
-  static const unsigned PTXAddrSpaceMap[] = {
-    0,    // opencl_global
-    4,    // opencl_local
-    1,    // opencl_constant
-    0,    // cuda_device
-    1,    // cuda_constant
-    4,    // cuda_shared
+  static const unsigned NVPTXAddrSpaceMap[] = {
+    1,    // opencl_global
+    3,    // opencl_local
+    4,    // opencl_constant
+    1,    // cuda_device
+    4,    // cuda_constant
+    3,    // cuda_shared
   };
-  class PTXTargetInfo : public TargetInfo {
+  class NVPTXTargetInfo : public TargetInfo {
     static const char * const GCCRegNames[];
     static const Builtin::Info BuiltinInfo[];
     std::vector<llvm::StringRef> AvailableFeatures;
   public:
-    PTXTargetInfo(const std::string& triple) : TargetInfo(triple) {
+    NVPTXTargetInfo(const std::string& triple) : TargetInfo(triple) {
       BigEndian = false;
       TLSSupported = false;
       LongWidth = LongAlign = 64;
-      AddrSpaceMap = &PTXAddrSpaceMap;
+      AddrSpaceMap = &NVPTXAddrSpaceMap;
       // Define available target features
-      // These must be defined in sorted order!      
-      AvailableFeatures.push_back("compute10");
-      AvailableFeatures.push_back("compute11");
-      AvailableFeatures.push_back("compute12");
-      AvailableFeatures.push_back("compute13");
-      AvailableFeatures.push_back("compute20");
-      AvailableFeatures.push_back("double");
-      AvailableFeatures.push_back("no-fma");
-      AvailableFeatures.push_back("ptx20");
-      AvailableFeatures.push_back("ptx21");
-      AvailableFeatures.push_back("ptx22");
-      AvailableFeatures.push_back("ptx23");
-      AvailableFeatures.push_back("sm10");
-      AvailableFeatures.push_back("sm11");
-      AvailableFeatures.push_back("sm12");
-      AvailableFeatures.push_back("sm13");
-      AvailableFeatures.push_back("sm20");
-      AvailableFeatures.push_back("sm21");
-      AvailableFeatures.push_back("sm22");
-      AvailableFeatures.push_back("sm23");
+      // These must be defined in sorted order!
     }
     virtual void getTargetDefines(const LangOptions &Opts,
                                   MacroBuilder &Builder) const {
       Builder.defineMacro("__PTX__");
+      Builder.defineMacro("__NVPTX__");
     }
     virtual void getTargetBuiltins(const Builtin::Info *&Records,
                                    unsigned &NumRecords) const {
       Records = BuiltinInfo;
-      NumRecords = clang::PTX::LastTSBuiltin-Builtin::FirstTSBuiltin;
+      NumRecords = clang::NVPTX::LastTSBuiltin-Builtin::FirstTSBuiltin;
     }
     virtual bool hasFeature(StringRef Feature) const {
-      return Feature == "ptx";
+      return Feature == "ptx" || Feature == "nvptx";
     }
     
     virtual void getGCCRegNames(const char * const *&Names,
@@ -1020,32 +1002,34 @@ namespace {
       // FIXME: implement
       return "typedef char* __builtin_va_list;";
     }
-
+    virtual bool setCPU(const std::string &Name) {
+      return Name == "sm_10" || Name == "sm_13" || Name == "sm_20";
+    }
     virtual bool setFeatureEnabled(llvm::StringMap<bool> &Features,
                                    StringRef Name,
                                    bool Enabled) const;
   };
 
-  const Builtin::Info PTXTargetInfo::BuiltinInfo[] = {
+  const Builtin::Info NVPTXTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS) { #ID, TYPE, ATTRS, 0, ALL_LANGUAGES },
 #define LIBBUILTIN(ID, TYPE, ATTRS, HEADER) { #ID, TYPE, ATTRS, HEADER,\
                                               ALL_LANGUAGES },
-#include "clang/Basic/BuiltinsPTX.def"
+#include "clang/Basic/BuiltinsNVPTX.def"
   };
 
-  const char * const PTXTargetInfo::GCCRegNames[] = {
+  const char * const NVPTXTargetInfo::GCCRegNames[] = {
     "r0"
   };
 
-  void PTXTargetInfo::getGCCRegNames(const char * const *&Names,
+  void NVPTXTargetInfo::getGCCRegNames(const char * const *&Names,
                                      unsigned &NumNames) const {
     Names = GCCRegNames;
     NumNames = llvm::array_lengthof(GCCRegNames);
   }
 
-  bool PTXTargetInfo::setFeatureEnabled(llvm::StringMap<bool> &Features,
-                                        StringRef Name,
-                                        bool Enabled) const {
+  bool NVPTXTargetInfo::setFeatureEnabled(llvm::StringMap<bool> &Features,
+                                          StringRef Name,
+                                          bool Enabled) const {
     if(std::binary_search(AvailableFeatures.begin(), AvailableFeatures.end(),
                           Name)) {
       Features[Name] = Enabled;
@@ -1055,117 +1039,28 @@ namespace {
     }
   }
 
-  class PTX32TargetInfo : public PTXTargetInfo {
-  public:
-  PTX32TargetInfo(const std::string& triple) : PTXTargetInfo(triple) {
-      PointerWidth = PointerAlign = 32;
-      SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedInt;
-      DescriptionString
-        = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64";
-    }
-  };
-
-  class PTX64TargetInfo : public PTXTargetInfo {
-  public:
-  PTX64TargetInfo(const std::string& triple) : PTXTargetInfo(triple) {
-      PointerWidth = PointerAlign = 64;
-      SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedLongLong;
-      DescriptionString
-        = "e-p:64:64-i64:64:64-f64:64:64-n1:8:16:32:64";
-    }
-  };
-}
-
-namespace {
-  static const unsigned NVPTXAddrSpaceMap[] = {
-    1,    // opencl_global
-    3,    // opencl_local
-    4,    // opencl_constant
-    1,    // cuda_device
-    4,    // cuda_constant
-    3,    // cuda_shared
-  };
-  class NVPTXTargetInfo : public TargetInfo {
-    static const char * const GCCRegNames[];
-  public:
-    NVPTXTargetInfo(const std::string& triple) : TargetInfo(triple) {
-      BigEndian = false;
-      TLSSupported = false;
-      LongWidth = LongAlign = 64;
-      AddrSpaceMap = &NVPTXAddrSpaceMap;
-    }
-    virtual void getTargetDefines(const LangOptions &Opts,
-                                  MacroBuilder &Builder) const {
-      Builder.defineMacro("__PTX__");
-    }
-    virtual void getTargetBuiltins(const Builtin::Info *&Records,
-                                   unsigned &NumRecords) const {
-      // FIXME: implement.
-      Records = 0;
-      NumRecords = 0;
-    }
-    virtual bool hasFeature(StringRef Feature) const {
-      return Feature == "nvptx";
-    }
-    
-    virtual void getGCCRegNames(const char * const *&Names,
-                                unsigned &NumNames) const;
-    virtual void getGCCRegAliases(const GCCRegAlias *&Aliases,
-                                  unsigned &NumAliases) const {
-      // No aliases.
-      Aliases = 0;
-      NumAliases = 0;
-    }
-    virtual bool validateAsmConstraint(const char *&Name,
-                                       TargetInfo::ConstraintInfo &info) const {
-      // FIXME: implement
-      return true;
-    }
-    virtual const char *getClobbers() const {
-      // FIXME: Is this really right?
-      return "";
-    }
-    virtual const char *getVAListDeclaration() const {
-      // FIXME: implement
-      return "typedef char* __builtin_va_list;";
-    }
-    virtual bool setCPU(const std::string &Name) {
-      return Name == "sm_10";
-    }
-  };
-
-  const char * const NVPTXTargetInfo::GCCRegNames[] = {
-    "r0"
-  };
-
-  void NVPTXTargetInfo::getGCCRegNames(const char * const *&Names,
-                                     unsigned &NumNames) const {
-    Names = GCCRegNames;
-    NumNames = llvm::array_lengthof(GCCRegNames);
-  }
-
   class NVPTX32TargetInfo : public NVPTXTargetInfo {
   public:
-  NVPTX32TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) {
+    NVPTX32TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) {
       PointerWidth = PointerAlign = 32;
-      SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedInt;
+      SizeType     = PtrDiffType = IntPtrType = TargetInfo::UnsignedInt;
       DescriptionString
         = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-"
           "f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-"
           "n16:32:64";
-    }
+  }
   };
 
   class NVPTX64TargetInfo : public NVPTXTargetInfo {
   public:
-  NVPTX64TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) {
+    NVPTX64TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) {
       PointerWidth = PointerAlign = 64;
-      SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedLongLong;
+      SizeType     = PtrDiffType = IntPtrType = TargetInfo::UnsignedLongLong;
       DescriptionString
         = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-"
           "f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-"
           "n16:32:64";
-    }
+  }
   };
 }
 
@@ -4139,11 +4034,6 @@ static TargetInfo *AllocateTarget(const std::string &T) {
       return new PPC64TargetInfo(T);
     }
 
-  case llvm::Triple::ptx32:
-    return new PTX32TargetInfo(T);
-  case llvm::Triple::ptx64:
-    return new PTX64TargetInfo(T);
-
   case llvm::Triple::nvptx:
     return new NVPTX32TargetInfo(T);
   case llvm::Triple::nvptx64:
diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp
index 97ca238130..357b3fe537 100644
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/lib/CodeGen/TargetInfo.cpp
@@ -2891,14 +2891,14 @@ llvm::Value *ARMABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty,
 }
 
 //===----------------------------------------------------------------------===//
-// PTX ABI Implementation
+// NVPTX ABI Implementation
 //===----------------------------------------------------------------------===//
 
 namespace {
 
-class PTXABIInfo : public ABIInfo {
+class NVPTXABIInfo : public ABIInfo {
 public:
-  PTXABIInfo(CodeGenTypes &CGT) : ABIInfo(CGT) {}
+  NVPTXABIInfo(CodeGenTypes &CGT) : ABIInfo(CGT) {}
 
   ABIArgInfo classifyReturnType(QualType RetTy) const;
   ABIArgInfo classifyArgumentType(QualType Ty) const;
@@ -2908,16 +2908,16 @@ public:
                                  CodeGenFunction &CFG) const;
 };
 
-class PTXTargetCodeGenInfo : public TargetCodeGenInfo {
+class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
 public:
-  PTXTargetCodeGenInfo(CodeGenTypes &CGT)
-    : TargetCodeGenInfo(new PTXABIInfo(CGT)) {}
+  NVPTXTargetCodeGenInfo(CodeGenTypes &CGT)
+    : TargetCodeGenInfo(new NVPTXABIInfo(CGT)) {}
     
   virtual void SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
                                    CodeGen::CodeGenModule &M) const;
 };
 
-ABIArgInfo PTXABIInfo::classifyReturnType(QualType RetTy) const {
+ABIArgInfo NVPTXABIInfo::classifyReturnType(QualType RetTy) const {
   if (RetTy->isVoidType())
     return ABIArgInfo::getIgnore();
   if (isAggregateTypeForABI(RetTy))
@@ -2925,14 +2925,14 @@ ABIArgInfo PTXABIInfo::classifyReturnType(QualType RetTy) const {
   return ABIArgInfo::getDirect();
 }
 
-ABIArgInfo PTXABIInfo::classifyArgumentType(QualType Ty) const {
+ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType Ty) const {
   if (isAggregateTypeForABI(Ty))
     return ABIArgInfo::getIndirect(0);
 
   return ABIArgInfo::getDirect();
 }
 
-void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
+void NVPTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
   FI.getReturnInfo() = classifyReturnType(FI.getReturnType());
   for (CGFunctionInfo::arg_iterator it = FI.arg_begin(), ie = FI.arg_end();
        it != ie; ++it)
@@ -2943,6 +2943,8 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
     return;
 
   // Calling convention as default by an ABI.
+  // We're still using the PTX_Kernel/PTX_Device calling conventions here,
+  // but we should switch to NVVM metadata later on.
   llvm::CallingConv::ID DefaultCC;
   const LangOptions &LangOpts = getContext().getLangOpts();
   if (LangOpts.OpenCL || LangOpts.CUDA) {
@@ -2961,14 +2963,14 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
    
 }
 
-llvm::Value *PTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty,
-                                   CodeGenFunction &CFG) const {
-  llvm_unreachable("PTX does not support varargs");
+llvm::Value *NVPTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty,
+                                     CodeGenFunction &CFG) const {
+  llvm_unreachable("NVPTX does not support varargs");
 }
 
-void PTXTargetCodeGenInfo::SetTargetAttributes(const Decl *D,
-                                               llvm::GlobalValue *GV,
-                                               CodeGen::CodeGenModule &M) const{
+void NVPTXTargetCodeGenInfo::
+SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
+                    CodeGen::CodeGenModule &M) const{
   const FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
   if (!FD) return;
 
@@ -3704,11 +3706,9 @@ const TargetCodeGenInfo &CodeGenModule::getTargetCodeGenInfo() {
   case llvm::Triple::ppc64:
     return *(TheTargetCodeGenInfo = new PPC64TargetCodeGenInfo(Types));
 
-  case llvm::Triple::ptx32:
-  case llvm::Triple::ptx64:
   case llvm::Triple::nvptx:
   case llvm::Triple::nvptx64:
-    return *(TheTargetCodeGenInfo = new PTXTargetCodeGenInfo(Types));
+    return *(TheTargetCodeGenInfo = new NVPTXTargetCodeGenInfo(Types));
 
   case llvm::Triple::mblaze:
     return *(TheTargetCodeGenInfo = new MBlazeTargetCodeGenInfo(Types));
diff --git a/test/CodeGen/builtins-ptx.c b/test/CodeGen/builtins-nvptx.c
index 6dd10188e9..4a094bbd7e 100644
--- a/test/CodeGen/builtins-ptx.c
+++ b/test/CodeGen/builtins-nvptx.c
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -triple ptx32-unknown-unknown -emit-llvm -o %t %s
-// RUN: %clang_cc1 -triple ptx64-unknown-unknown -emit-llvm -o %t %s
+// RUN: %clang_cc1 -triple nvptx-unknown-unknown -emit-llvm -o %t %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -emit-llvm -o %t %s
 
 
 int read_tid() {
diff --git a/test/CodeGen/ptx-cc.c b/test/CodeGen/nvptx-cc.c
index 2212d4260b..1c0d943f95 100644
--- a/test/CodeGen/ptx-cc.c
+++ b/test/CodeGen/nvptx-cc.c
@@ -1,7 +1,7 @@
-// RUN: %clang_cc1 -triple ptx32-unknown-unknown -O3 -S -o %t %s -emit-llvm
-// RUN: %clang_cc1 -triple ptx64-unknown-unknown -O3 -S -o %t %s -emit-llvm
+// RUN: %clang_cc1 -triple nvptx-unknown-unknown -O3 -S -o %t %s -emit-llvm
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -O3 -S -o %t %s -emit-llvm
 
-// Just make sure Clang uses the proper calling convention for the PTX back-end.
+// Just make sure Clang uses the proper calling convention for the NVPTX back-end.
 // If something is wrong, the back-end will fail.
 void foo(float* a,
          float* b) {
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index 2da61ec95a..61d4d6b6ba 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -1,24 +1,24 @@
-// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple ptx32-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
 
 #include "../SemaCUDA/cuda.h"
 
-// CHECK: @i = global
+// CHECK: @i = addrspace(1) global
 __device__ int i;
 
-// CHECK: @j = addrspace(1) global
+// CHECK: @j = addrspace(4) global
 __constant__ int j;
 
-// CHECK: @k = addrspace(4) global
+// CHECK: @k = addrspace(3) global
 __shared__ int k;
 
 __device__ void foo() {
-  // CHECK: load i32* @i
+  // CHECK: load i32* bitcast (i32 addrspace(1)* @i to i32*)
   i++;
 
-  // CHECK: load i32* bitcast (i32 addrspace(1)* @j to i32*)
+  // CHECK: load i32* bitcast (i32 addrspace(4)* @j to i32*)
   j++;
 
-  // CHECK: load i32* bitcast (i32 addrspace(4)* @k to i32*)
+  // CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*)
   k++;
 }
 
diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu
index ecca8519af..f0bf2952a1 100644
--- a/test/CodeGenCUDA/ptx-kernels.cu
+++ b/test/CodeGenCUDA/ptx-kernels.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
 
 #include "../SemaCUDA/cuda.h"
 
diff --git a/test/CodeGenOpenCL/ptx-calls.cl b/test/CodeGenOpenCL/ptx-calls.cl
index 6f336405c3..34a21c6c1d 100644
--- a/test/CodeGenOpenCL/ptx-calls.cl
+++ b/test/CodeGenOpenCL/ptx-calls.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 void device_function() {
 }
diff --git a/test/CodeGenOpenCL/ptx-kernels.cl b/test/CodeGenOpenCL/ptx-kernels.cl
index 4d6fa1084d..1d7e497b7c 100644
--- a/test/CodeGenOpenCL/ptx-kernels.cl
+++ b/test/CodeGenOpenCL/ptx-kernels.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -emit-llvm -o - | FileCheck %s
 
 void device_function() {
 }