[CUDA] Added rudimentary support for CUDA-9 and sm_70.

[tra]

Summary

For now CUDA-9 is not included in the list of CUDA versions clang
searches for, so the path to CUDA-9 must be explicitly passed
via --cuda-path=.

On LLVM side NVPTX added sm_70 GPU type which bumps required
PTX version to 6.0, but otherwise is equivalent to sm_62 at the moment.

[tra]

Added tests for sm_70 support.

[jlebar]

Looks great.

clang/include/clang/Basic/Cuda.h

 
 namespace clang {
 
 enum class CudaVersion {
   UNKNOWN,
   CUDA_70,
   CUDA_75,
   CUDA_80,
+  CUDA_90,
 };
 const char *CudaVersionToString(CudaVersion V);
 
 // No string -> CudaVersion conversion function because there's no canonical
 // spelling of the various CUDA versions.
 
 enum class CudaArch {
   UNKNOWN,
   SM_20,
   SM_21,
   SM_30,
   SM_32,
   SM_35,
   SM_37,
   SM_50,
   SM_52,
   SM_53,
   SM_60,
   SM_61,
   SM_62,
+  SM_70,
 };
 const char *CudaArchToString(CudaArch A);
 
 // The input should have the form "sm_20".
 CudaArch StringToCudaArch(llvm::StringRef S);
 
 enum class CudaVirtualArch {
   UNKNOWN,
   COMPUTE_20,
   COMPUTE_30,
   COMPUTE_32,
   COMPUTE_35,
   COMPUTE_37,
   COMPUTE_50,
   COMPUTE_52,
   COMPUTE_53,
   COMPUTE_60,
   COMPUTE_61,
   COMPUTE_62,
+  COMPUTE_70,
 };
 const char *CudaVirtualArchToString(CudaVirtualArch A);
 
 // The input should have the form "compute_20".
 CudaVirtualArch StringToCudaVirtualArch(llvm::StringRef S);
 
 /// Get the compute_xx corresponding to an sm_yy.
 CudaVirtualArch VirtualArchForCudaArch(CudaArch A);
 
 /// Get the earliest CudaVersion that supports the given CudaArch.
 CudaVersion MinVersionForCudaArch(CudaArch A);
 
 } // namespace clang
 
 #endif

clang/lib/Basic/Cuda.cpp

const char *CudaVersionToString(CudaVersion V) {
   case CudaVersion::UNKNOWN:
     return "unknown";
   case CudaVersion::CUDA_70:
     return "7.0";
   case CudaVersion::CUDA_75:
     return "7.5";
   case CudaVersion::CUDA_80:
     return "8.0";
+  case CudaVersion::CUDA_90:
+    return "9.0";
   }
   llvm_unreachable("invalid enum");
 }
 
 const char *CudaArchToString(CudaArch A) {
   switch (A) {
   case CudaArch::UNKNOWN:
     return "unknown";
const char *CudaArchToString(CudaArch A) {
   case CudaArch::SM_53:
     return "sm_53";
   case CudaArch::SM_60:
     return "sm_60";
   case CudaArch::SM_61:
     return "sm_61";
   case CudaArch::SM_62:
     return "sm_62";
+  case CudaArch::SM_70:
+    return "sm_70";
   }
   llvm_unreachable("invalid enum");
 }
 
 CudaArch StringToCudaArch(llvm::StringRef S) {
   return llvm::StringSwitch<CudaArch>(S)
       .Case("sm_20", CudaArch::SM_20)
       .Case("sm_21", CudaArch::SM_21)
       .Case("sm_30", CudaArch::SM_30)
       .Case("sm_32", CudaArch::SM_32)
       .Case("sm_35", CudaArch::SM_35)
       .Case("sm_37", CudaArch::SM_37)
       .Case("sm_50", CudaArch::SM_50)
       .Case("sm_52", CudaArch::SM_52)
       .Case("sm_53", CudaArch::SM_53)
       .Case("sm_60", CudaArch::SM_60)
       .Case("sm_61", CudaArch::SM_61)
       .Case("sm_62", CudaArch::SM_62)
+      .Case("sm_70", CudaArch::SM_70)
       .Default(CudaArch::UNKNOWN);
 }
 
 const char *CudaVirtualArchToString(CudaVirtualArch A) {
   switch (A) {
   case CudaVirtualArch::UNKNOWN:
     return "unknown";
   case CudaVirtualArch::COMPUTE_20:
const char *CudaVirtualArchToString(CudaVirtualArch A) {
   case CudaVirtualArch::COMPUTE_53:
     return "compute_53";
   case CudaVirtualArch::COMPUTE_60:
     return "compute_60";
   case CudaVirtualArch::COMPUTE_61:
     return "compute_61";
   case CudaVirtualArch::COMPUTE_62:
     return "compute_62";
+  case CudaVirtualArch::COMPUTE_70:
+    return "compute_70";
   }
   llvm_unreachable("invalid enum");
 }
 
 CudaVirtualArch StringToCudaVirtualArch(llvm::StringRef S) {
   return llvm::StringSwitch<CudaVirtualArch>(S)
       .Case("compute_20", CudaVirtualArch::COMPUTE_20)
       .Case("compute_30", CudaVirtualArch::COMPUTE_30)
       .Case("compute_32", CudaVirtualArch::COMPUTE_32)
       .Case("compute_35", CudaVirtualArch::COMPUTE_35)
       .Case("compute_37", CudaVirtualArch::COMPUTE_37)
       .Case("compute_50", CudaVirtualArch::COMPUTE_50)
       .Case("compute_52", CudaVirtualArch::COMPUTE_52)
       .Case("compute_53", CudaVirtualArch::COMPUTE_53)
       .Case("compute_60", CudaVirtualArch::COMPUTE_60)
       .Case("compute_61", CudaVirtualArch::COMPUTE_61)
       .Case("compute_62", CudaVirtualArch::COMPUTE_62)
+      .Case("compute_70", CudaVirtualArch::COMPUTE_70)
       .Default(CudaVirtualArch::UNKNOWN);
 }
 
 CudaVirtualArch VirtualArchForCudaArch(CudaArch A) {
   switch (A) {
   case CudaArch::UNKNOWN:
     return CudaVirtualArch::UNKNOWN;
   case CudaArch::SM_20:
CudaVirtualArch VirtualArchForCudaArch(CudaArch A) {
   case CudaArch::SM_53:
     return CudaVirtualArch::COMPUTE_53;
   case CudaArch::SM_60:
     return CudaVirtualArch::COMPUTE_60;
   case CudaArch::SM_61:
     return CudaVirtualArch::COMPUTE_61;
   case CudaArch::SM_62:
     return CudaVirtualArch::COMPUTE_62;
+  case CudaArch::SM_70:
+    return CudaVirtualArch::COMPUTE_70;
   }
   llvm_unreachable("invalid enum");
 }
 
 CudaVersion MinVersionForCudaArch(CudaArch A) {
   switch (A) {
   case CudaArch::UNKNOWN:
     return CudaVersion::UNKNOWN;
   case CudaArch::SM_20:
   case CudaArch::SM_21:
   case CudaArch::SM_30:
   case CudaArch::SM_32:
   case CudaArch::SM_35:
   case CudaArch::SM_37:
   case CudaArch::SM_50:
   case CudaArch::SM_52:
   case CudaArch::SM_53:
     return CudaVersion::CUDA_70;
   case CudaArch::SM_60:
   case CudaArch::SM_61:
   case CudaArch::SM_62:
     return CudaVersion::CUDA_80;
+  case CudaArch::SM_70:
+    return CudaVersion::CUDA_90;
   }
   llvm_unreachable("invalid enum");
 }
 
 } // namespace clang

clang/lib/Basic/Targets/NVPTX.cpp

std::string CUDAArchCode = [this] {
       case CudaArch::SM_53:
         return "530";
       case CudaArch::SM_60:
         return "600";
       case CudaArch::SM_61:
         return "610";
       case CudaArch::SM_62:
         return "620";
+      case CudaArch::SM_70:
+        return "700";
       }
       llvm_unreachable("unhandled CudaArch");
     }();
     Builder.defineMacro("__CUDA_ARCH__", CUDAArchCode);
   }
 }
 
 ArrayRef<Builtin::Info> NVPTXTargetInfo::getTargetBuiltins() const {
   return llvm::makeArrayRef(BuiltinInfo, clang::NVPTX::LastTSBuiltin -
                                              Builtin::FirstTSBuiltin);
 }

clang/lib/Driver/ToolChains/Cuda.cpp

if (Major == 7 && Minor == 0) {
     // This doesn't appear to ever happen -- version.txt doesn't exist in the
     // CUDA 7 installs I've seen.  But no harm in checking.
     return CudaVersion::CUDA_70;
   }
   if (Major == 7 && Minor == 5)
     return CudaVersion::CUDA_75;
   if (Major == 8 && Minor == 0)
     return CudaVersion::CUDA_80;
+  if (Major == 9 && Minor == 0)
+    return CudaVersion::CUDA_90;
   return CudaVersion::UNKNOWN;
 }
 
 CudaInstallationDetector::CudaInstallationDetector(
     const Driver &D, const llvm::Triple &HostTriple,
     const llvm::opt::ArgList &Args)
     : D(D) {
   SmallVector<std::string, 4> CudaPathCandidates;
for (const auto &CudaPath : CudaPathCandidates) {
     if (!VersionFile) {
       // CUDA 7.0 doesn't have a version.txt, so guess that's our version if
       // version.txt isn't present.
       Version = CudaVersion::CUDA_70;
     } else {
       Version = ParseCudaVersionFile((*VersionFile)->getBuffer());
     }
 
+    if (Version == CudaVersion::CUDA_90) {
+      // CUDA-9 uses single libdevice file for all GPU variants.
+      std::string FilePath = LibDevicePath + "/libdevice.10.bc";
+      if (FS.exists(FilePath)) {
+        for (const char *GpuArch :
+             {"sm_20", "sm_30", "sm_32", "sm_35", "sm_50", "sm_52", "sm_53",
+              "sm_60", "sm_61", "sm_62", "sm_70"})
+          LibDeviceMap[GpuArch] = FilePath;
+      }
+    } else {
       std::error_code EC;
       for (llvm::sys::fs::directory_iterator LI(LibDevicePath, EC), LE;
            !EC && LI != LE; LI = LI.increment(EC)) {
         StringRef FilePath = LI->path();
         StringRef FileName = llvm::sys::path::filename(FilePath);
-      // Process all bitcode filenames that look like libdevice.compute_XX.YY.bc
+        // Process all bitcode filenames that look like
+        // libdevice.compute_XX.YY.bc
         const StringRef LibDeviceName = "libdevice.";
         if (!(FileName.startswith(LibDeviceName) && FileName.endswith(".bc")))
           continue;
         StringRef GpuArch = FileName.slice(
             LibDeviceName.size(), FileName.find('.', LibDeviceName.size()));
         LibDeviceMap[GpuArch] = FilePath.str();
         // Insert map entries for specifc devices with this compute
         // capability. NVCC's choice of the libdevice library version is
         // rather peculiar and depends on the CUDA version.
         if (GpuArch == "compute_20") {
           LibDeviceMap["sm_20"] = FilePath;
           LibDeviceMap["sm_21"] = FilePath;
           LibDeviceMap["sm_32"] = FilePath;
         } else if (GpuArch == "compute_30") {
           LibDeviceMap["sm_30"] = FilePath;
           if (Version < CudaVersion::CUDA_80) {
             LibDeviceMap["sm_50"] = FilePath;
             LibDeviceMap["sm_52"] = FilePath;
             LibDeviceMap["sm_53"] = FilePath;
           }
           LibDeviceMap["sm_60"] = FilePath;
           LibDeviceMap["sm_61"] = FilePath;
           LibDeviceMap["sm_62"] = FilePath;
         } else if (GpuArch == "compute_35") {
           LibDeviceMap["sm_35"] = FilePath;
           LibDeviceMap["sm_37"] = FilePath;
         } else if (GpuArch == "compute_50") {
           if (Version >= CudaVersion::CUDA_80) {
             LibDeviceMap["sm_50"] = FilePath;
             LibDeviceMap["sm_52"] = FilePath;
             LibDeviceMap["sm_53"] = FilePath;
           }
         }
       }
+    }
 
     // This code prevents IsValid from being set when
     // no libdevice has been found.
     bool allEmpty = true;
     std::string LibDeviceFile;
     for (auto key : LibDeviceMap.keys()) {
       LibDeviceFile = LibDeviceMap.lookup(key);
       if (!LibDeviceFile.empty())

clang/lib/Headers/__clang_cuda_runtime_wrapper.h

 #pragma push_macro("__CUDA_ARCH__")
 
 // WARNING: Preprocessor hacks below are based on specific details of
 // CUDA-7.x headers and are not expected to work with any other
 // version of CUDA headers.
 #include "cuda.h"
 #if !defined(CUDA_VERSION)
 #error "cuda.h did not define CUDA_VERSION"
-#elif CUDA_VERSION < 7000 || CUDA_VERSION > 8000
+#elif CUDA_VERSION < 7000 || CUDA_VERSION > 9000
 #error "Unsupported CUDA version!"
 #endif
 
 // Make largest subset of device functions available during host
 // compilation -- SM_35 for the time being.
 #ifndef __CUDA_ARCH__
 #define __CUDA_ARCH__ 350
 #endif
 
 #include "__clang_cuda_builtin_vars.h"
 
 // No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above
 // has taken care of builtin variables declared in the file.
 #define __DEVICE_LAUNCH_PARAMETERS_H__
 
 // {math,device}_functions.h only have declarations of the
 // functions. We don't need them as we're going to pull in their
 // definitions from .hpp files.
 #define __DEVICE_FUNCTIONS_H__
 #define __MATH_FUNCTIONS_H__
 #define __COMMON_FUNCTIONS_H__
 
 #undef __CUDACC__
+#if CUDA_VERSION < 9000
 #define __CUDABE__
+#else
+#define __CUDA_LIBDEVICE__
+#endif
 // Disables definitions of device-side runtime support stubs in
 // cuda_device_runtime_api.h
 #include "driver_types.h"
 #include "host_config.h"
 #include "host_defines.h"
 
 #undef __CUDABE__
+#undef __CUDA_LIBDEVICE__
 #define __CUDACC__
 #include "cuda_runtime.h"
 
 #undef __CUDACC__
 #define __CUDABE__
 
 // CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does
 // not have at the moment. Emulate them with a builtin memcpy/memset.
 #define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n)
 #define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n)
 
+#if CUDA_VERSION < 9000
 #include "crt/device_runtime.h"
+#endif
 #include "crt/host_runtime.h"
 // device_runtime.h defines __cxa_* macros that will conflict with
 // cxxabi.h.
 // FIXME: redefine these as __device__ functions.
 #undef __cxa_vec_ctor
 #undef __cxa_vec_cctor
 #undef __cxa_vec_dtor
 #undef __cxa_vec_new

clang/test/Driver/cuda-arch-translation.cu

 // Tests that "sm_XX" gets correctly converted to "compute_YY" when we invoke
 // fatbinary.
 //
 // REQUIRES: clang-driver
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
-// CHECK:fatbinary
-
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_20 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM20 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM20 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_21 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM21 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM21 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_30 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM30 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM30 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_32 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM32 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM32 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_35 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM35 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM35 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_37 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM37 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM37 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_50 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM50 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM50 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_52 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM52 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM52 %s
 // RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_53 %s 2>&1 \
-// RUN: | FileCheck -check-prefix ARCH64 -check-prefix SM53 %s
+// RUN: | FileCheck -check-prefixes=COMMON,SM53 %s
+// RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_60 %s 2>&1 \
+// RUN: | FileCheck -check-prefixes=COMMON,SM60 %s
+// RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_61 %s 2>&1 \
+// RUN: | FileCheck -check-prefixes=COMMON,SM61 %s
+// RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_62 %s 2>&1 \
+// RUN: | FileCheck -check-prefixes=COMMON,SM62 %s
+// RUN: %clang -### -target x86_64-linux-gnu -c --cuda-gpu-arch=sm_70 %s 2>&1 \
+// RUN: | FileCheck -check-prefixes=COMMON,SM70 %s
+
+// COMMON: ptxas
+// COMMON-SAME: -m64
+// COMMON: fatbinary
 
 // SM20:--image=profile=sm_20{{.*}}--image=profile=compute_20
 // SM21:--image=profile=sm_21{{.*}}--image=profile=compute_20
 // SM30:--image=profile=sm_30{{.*}}--image=profile=compute_30
 // SM32:--image=profile=sm_32{{.*}}--image=profile=compute_32
 // SM35:--image=profile=sm_35{{.*}}--image=profile=compute_35
 // SM37:--image=profile=sm_37{{.*}}--image=profile=compute_37
 // SM50:--image=profile=sm_50{{.*}}--image=profile=compute_50
 // SM52:--image=profile=sm_52{{.*}}--image=profile=compute_52
 // SM53:--image=profile=sm_53{{.*}}--image=profile=compute_53
+// SM60:--image=profile=sm_60{{.*}}--image=profile=compute_60
+// SM61:--image=profile=sm_61{{.*}}--image=profile=compute_61
+// SM62:--image=profile=sm_62{{.*}}--image=profile=compute_62
+// SM70:--image=profile=sm_70{{.*}}--image=profile=compute_70

llvm/lib/Target/NVPTX/NVPTX.td

 def SM53 : SubtargetFeature<"sm_53", "SmVersion", "53",
                             "Target SM 5.3">;
 def SM60 : SubtargetFeature<"sm_60", "SmVersion", "60",
                              "Target SM 6.0">;
 def SM61 : SubtargetFeature<"sm_61", "SmVersion", "61",
                              "Target SM 6.1">;
 def SM62 : SubtargetFeature<"sm_62", "SmVersion", "62",
                              "Target SM 6.2">;
+def SM70 : SubtargetFeature<"sm_70", "SmVersion", "70",
+                             "Target SM 7.0">;
 
 def SATOM : SubtargetFeature<"satom", "HasAtomScope", "true",
                              "Atomic operations with scope">;
 
 // PTX Versions
 def PTX32 : SubtargetFeature<"ptx32", "PTXVersion", "32",
                              "Use PTX version 3.2">;
 def PTX40 : SubtargetFeature<"ptx40", "PTXVersion", "40",
                              "Use PTX version 4.0">;
 def PTX41 : SubtargetFeature<"ptx41", "PTXVersion", "41",
                              "Use PTX version 4.1">;
 def PTX42 : SubtargetFeature<"ptx42", "PTXVersion", "42",
                              "Use PTX version 4.2">;
 def PTX43 : SubtargetFeature<"ptx43", "PTXVersion", "43",
                              "Use PTX version 4.3">;
 def PTX50 : SubtargetFeature<"ptx50", "PTXVersion", "50",
                              "Use PTX version 5.0">;
+def PTX60 : SubtargetFeature<"ptx60", "PTXVersion", "60",
+                             "Use PTX version 6.0">;
 
 //===----------------------------------------------------------------------===//
 // NVPTX supported processors.
 //===----------------------------------------------------------------------===//
 
 class Proc<string Name, list<SubtargetFeature> Features>
  : Processor<Name, NoItineraries, Features>;
 
 def : Proc<"sm_20", [SM20]>;
 def : Proc<"sm_21", [SM21]>;
 def : Proc<"sm_30", [SM30]>;
 def : Proc<"sm_32", [SM32, PTX40]>;
 def : Proc<"sm_35", [SM35]>;
 def : Proc<"sm_37", [SM37, PTX41]>;
 def : Proc<"sm_50", [SM50, PTX40]>;
 def : Proc<"sm_52", [SM52, PTX41]>;
 def : Proc<"sm_53", [SM53, PTX42]>;
 def : Proc<"sm_60", [SM60, PTX50, SATOM]>;
 def : Proc<"sm_61", [SM61, PTX50, SATOM]>;
 def : Proc<"sm_62", [SM62, PTX50, SATOM]>;
+def : Proc<"sm_70", [SM70, PTX60, SATOM]>;
 
 def NVPTXInstrInfo : InstrInfo {
 }
 
 def NVPTX : Target {
   let InstructionSet = NVPTXInstrInfo;
 }

llvm/test/CodeGen/NVPTX/sm-version-70.ll

+; RUN: llc < %s -march=nvptx -mcpu=sm_70 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 | FileCheck %s
+
+; CHECK: .version 6.0
+; CHECK: .target sm_70