From: Artem Belevich Date: Mon, 24 Sep 2018 23:10:44 +0000 (+0000) Subject: [CUDA] Added basic support for compiling with CUDA-10.0 X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=252ecba16a1694e520ca44a2cb065e7899786fbf;p=clang [CUDA] Added basic support for compiling with CUDA-10.0 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@342924 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/Basic/Cuda.h b/include/clang/Basic/Cuda.h index 24159e160f..f9fcbd95d8 100644 --- a/include/clang/Basic/Cuda.h +++ b/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); diff --git a/lib/Basic/Cuda.cpp b/lib/Basic/Cuda.cpp index dc7e61c02b..43400c39a7 100644 --- a/lib/Basic/Cuda.cpp +++ b/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/lib/Basic/Targets/NVPTX.cpp b/lib/Basic/Targets/NVPTX.cpp index fd4ee16060..2af28e34bd 100644 --- a/lib/Basic/Targets/NVPTX.cpp +++ b/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/lib/Driver/ToolChains/Cuda.cpp b/lib/Driver/ToolChains/Cuda.cpp index 73e65011b4..a446c4d9d4 100644 --- a/lib/Driver/ToolChains/Cuda.cpp +++ b/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/lib/Headers/__clang_cuda_runtime_wrapper.h b/lib/Headers/__clang_cuda_runtime_wrapper.h index 09705a273a..f05c0454a8 100644 --- a/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/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__