diff options
Diffstat (limited to 'clang/lib')
| -rw-r--r-- | clang/lib/Basic/Cuda.cpp | 12 | ||||
| -rw-r--r-- | clang/lib/Basic/Targets/NVPTX.cpp | 2 | ||||
| -rw-r--r-- | clang/lib/Driver/ToolChains/Cuda.cpp | 7 | ||||
| -rw-r--r-- | clang/lib/Headers/__clang_cuda_runtime_wrapper.h | 8 |
4 files changed, 27 insertions, 2 deletions
diff --git a/clang/lib/Basic/Cuda.cpp b/clang/lib/Basic/Cuda.cpp index dc7e61c02b2..43400c39a72 100644 --- a/clang/lib/Basic/Cuda.cpp +++ b/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: diff --git a/clang/lib/Basic/Targets/NVPTX.cpp b/clang/lib/Basic/Targets/NVPTX.cpp index fd4ee160606..2af28e34bd9 100644 --- a/clang/lib/Basic/Targets/NVPTX.cpp +++ b/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"); }(); diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp index 73e65011b45..a446c4d9d40 100644 --- a/clang/lib/Driver/ToolChains/Cuda.cpp +++ b/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"; diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index 09705a273a4..f05c0454a88 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/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__ |
