[CUDA] Added basic support for compiling with CUDA-10.0

llvm-svn: 342924

clang/include/clang/Basic/Cuda.h

@@ -24,7 +24,8 @@ enum class CudaVersion {
   CUDA_90,
   CUDA_91,
   CUDA_92,
-  LATEST = CUDA_92,
+  CUDA_100,
+  LATEST = CUDA_100,
 };
 const char *CudaVersionToString(CudaVersion V);
 
@@ -47,6 +48,7 @@ enum class CudaArch {
   SM_62,
   SM_70,
   SM_72,
+  SM_75,
   GFX600,
   GFX601,
   GFX700,
@@ -82,6 +84,7 @@ enum class CudaVirtualArch {
   COMPUTE_62,
   COMPUTE_70,
   COMPUTE_72,
+  COMPUTE_75,
   COMPUTE_AMDGCN,
 };
 const char *CudaVirtualArchToString(CudaVirtualArch A);

clang/lib/Basic/Cuda.cpp

@@ -22,6 +22,8 @@ const char *CudaVersionToString(CudaVersion V) {
     return "9.1";
   case CudaVersion::CUDA_92:
     return "9.2";
+  case CudaVersion::CUDA_100:
+    return "10.0";
   }
   llvm_unreachable("invalid enum");
 }
@@ -60,6 +62,8 @@ const char *CudaArchToString(CudaArch A) {
     return "sm_70";
   case CudaArch::SM_72:
     return "sm_72";
+  case CudaArch::SM_75:
+    return "sm_75";
   case CudaArch::GFX600: // tahiti
     return "gfx600";
   case CudaArch::GFX601: // pitcairn, verde, oland,hainan
@@ -106,6 +110,7 @@ CudaArch StringToCudaArch(llvm::StringRef S) {
       .Case("sm_62", CudaArch::SM_62)
       .Case("sm_70", CudaArch::SM_70)
       .Case("sm_72", CudaArch::SM_72)
+      .Case("sm_75", CudaArch::SM_75)
       .Case("gfx600", CudaArch::GFX600)
       .Case("gfx601", CudaArch::GFX601)
       .Case("gfx700", CudaArch::GFX700)
@@ -152,6 +157,8 @@ const char *CudaVirtualArchToString(CudaVirtualArch A) {
     return "compute_70";
   case CudaVirtualArch::COMPUTE_72:
     return "compute_72";
+  case CudaVirtualArch::COMPUTE_75:
+    return "compute_75";
   case CudaVirtualArch::COMPUTE_AMDGCN:
     return "compute_amdgcn";
   }
@@ -173,6 +180,7 @@ CudaVirtualArch StringToCudaVirtualArch(llvm::StringRef S) {
       .Case("compute_62", CudaVirtualArch::COMPUTE_62)
       .Case("compute_70", CudaVirtualArch::COMPUTE_70)
       .Case("compute_72", CudaVirtualArch::COMPUTE_72)
+      .Case("compute_75", CudaVirtualArch::COMPUTE_75)
       .Case("compute_amdgcn", CudaVirtualArch::COMPUTE_AMDGCN)
       .Default(CudaVirtualArch::UNKNOWN);
 }
@@ -210,6 +218,8 @@ CudaVirtualArch VirtualArchForCudaArch(CudaArch A) {
     return CudaVirtualArch::COMPUTE_70;
   case CudaArch::SM_72:
     return CudaVirtualArch::COMPUTE_72;
+  case CudaArch::SM_75:
+    return CudaVirtualArch::COMPUTE_75;
   case CudaArch::GFX600:
   case CudaArch::GFX601:
   case CudaArch::GFX700:
@@ -252,6 +262,8 @@ CudaVersion MinVersionForCudaArch(CudaArch A) {
     return CudaVersion::CUDA_90;
   case CudaArch::SM_72:
     return CudaVersion::CUDA_91;
+  case CudaArch::SM_75:
+    return CudaVersion::CUDA_100;
   case CudaArch::GFX600:
   case CudaArch::GFX601:
   case CudaArch::GFX700:

clang/lib/Basic/Targets/NVPTX.cpp

@@ -221,6 +221,8 @@ void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
         return "700";
       case CudaArch::SM_72:
         return "720";
+      case CudaArch::SM_75:
+        return "750";
       }
       llvm_unreachable("unhandled CudaArch");
     }();

clang/lib/Driver/ToolChains/Cuda.cpp

@@ -59,6 +59,8 @@ static CudaVersion ParseCudaVersionFile(llvm::StringRef V) {
     return CudaVersion::CUDA_91;
   if (Major == 9 && Minor == 2)
     return CudaVersion::CUDA_92;
+  if (Major == 10 && Minor == 0)
+    return CudaVersion::CUDA_100;
   return CudaVersion::UNKNOWN;
 }
 
@@ -165,7 +167,7 @@ CudaInstallationDetector::CudaInstallationDetector(
       if (FS.exists(FilePath)) {
         for (const char *GpuArchName :
              {"sm_30", "sm_32", "sm_35", "sm_37", "sm_50", "sm_52", "sm_53",
-              "sm_60", "sm_61", "sm_62", "sm_70", "sm_72"}) {
+              "sm_60", "sm_61", "sm_62", "sm_70", "sm_72", "sm_75"}) {
           const CudaArch GpuArch = StringToCudaArch(GpuArchName);
           if (Version >= MinVersionForCudaArch(GpuArch) &&
               Version <= MaxVersionForCudaArch(GpuArch))
@@ -628,6 +630,9 @@ void CudaToolChain::addClangTargetOptions(
   // defaults to. Use PTX4.2 by default, which is the PTX version that came with
   // CUDA-7.0.
   const char *PtxFeature = "+ptx42";
+  // TODO(tra): CUDA-10+ needs PTX 6.3 to support new features. However that
+  // requires fair amount of work on LLVM side. We'll keep using PTX 6.1 until
+  // all prerequisites are in place.
   if (CudaInstallation.version() >= CudaVersion::CUDA_91) {
     // CUDA-9.1 uses new instructions that are only available in PTX6.1+
     PtxFeature = "+ptx61";

clang/lib/Headers/__clang_cuda_runtime_wrapper.h

@@ -62,10 +62,15 @@
 #include "cuda.h"
 #if !defined(CUDA_VERSION)
 #error "cuda.h did not define CUDA_VERSION"
-#elif CUDA_VERSION < 7000 || CUDA_VERSION > 9020
+#elif CUDA_VERSION < 7000 || CUDA_VERSION > 10000
 #error "Unsupported CUDA version!"
 #endif
 
+#pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")
+#if CUDA_VERSION >= 10000
+#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__
+#endif
+
 // Make largest subset of device functions available during host
 // compilation -- SM_35 for the time being.
 #ifndef __CUDA_ARCH__
@@ -419,6 +424,7 @@ __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
 #pragma pop_macro("dim3")
 #pragma pop_macro("uint3")
 #pragma pop_macro("__USE_FAST_MATH__")
+#pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")
 
 #endif // __CUDA__
 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__

llvm/lib/Target/NVPTX/NVPTX.td

@@ -54,6 +54,8 @@ def SM70 : SubtargetFeature<"sm_70", "SmVersion", "70",
                              "Target SM 7.0">;
 def SM72 : SubtargetFeature<"sm_72", "SmVersion", "72",
                              "Target SM 7.2">;
+def SM75 : SubtargetFeature<"sm_75", "SmVersion", "75",
+                             "Target SM 7.5">;
 
 // PTX Versions
 def PTX32 : SubtargetFeature<"ptx32", "PTXVersion", "32",
@@ -72,6 +74,8 @@ def PTX60 : SubtargetFeature<"ptx60", "PTXVersion", "60",
                              "Use PTX version 6.0">;
 def PTX61 : SubtargetFeature<"ptx61", "PTXVersion", "61",
                              "Use PTX version 6.1">;
+def PTX63 : SubtargetFeature<"ptx63", "PTXVersion", "63",
+                             "Use PTX version 6.3">;
 
 //===----------------------------------------------------------------------===//
 // NVPTX supported processors.
@@ -94,6 +98,7 @@ def : Proc<"sm_61", [SM61, PTX50]>;
 def : Proc<"sm_62", [SM62, PTX50]>;
 def : Proc<"sm_70", [SM70, PTX60]>;
 def : Proc<"sm_72", [SM72, PTX61]>;
+def : Proc<"sm_75", [SM75, PTX63]>;
 
 def NVPTXInstrInfo : InstrInfo {
 }