Author: tra Date: Mon Sep 24 16:10:44 2018 New Revision: 342924 URL: http://llvm.org/viewvc/llvm-project?rev=342924&view=rev Log: [CUDA] Added basic support for compiling with CUDA-10.0
Modified: cfe/trunk/include/clang/Basic/Cuda.h cfe/trunk/lib/Basic/Cuda.cpp cfe/trunk/lib/Basic/Targets/NVPTX.cpp cfe/trunk/lib/Driver/ToolChains/Cuda.cpp cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Modified: cfe/trunk/include/clang/Basic/Cuda.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Cuda.h?rev=342924&r1=342923&r2=342924&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/Cuda.h (original) +++ cfe/trunk/include/clang/Basic/Cuda.h Mon Sep 24 16:10:44 2018 @@ -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); Modified: cfe/trunk/lib/Basic/Cuda.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Cuda.cpp?rev=342924&r1=342923&r2=342924&view=diff ============================================================================== --- cfe/trunk/lib/Basic/Cuda.cpp (original) +++ cfe/trunk/lib/Basic/Cuda.cpp Mon Sep 24 16:10:44 2018 @@ -22,6 +22,8 @@ const char *CudaVersionToString(CudaVers 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::StringRe .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(Cuda 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( .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(C 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(CudaAr 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: Modified: cfe/trunk/lib/Basic/Targets/NVPTX.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/NVPTX.cpp?rev=342924&r1=342923&r2=342924&view=diff ============================================================================== --- cfe/trunk/lib/Basic/Targets/NVPTX.cpp (original) +++ cfe/trunk/lib/Basic/Targets/NVPTX.cpp Mon Sep 24 16:10:44 2018 @@ -221,6 +221,8 @@ void NVPTXTargetInfo::getTargetDefines(c return "700"; case CudaArch::SM_72: return "720"; + case CudaArch::SM_75: + return "750"; } llvm_unreachable("unhandled CudaArch"); }(); Modified: cfe/trunk/lib/Driver/ToolChains/Cuda.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/ToolChains/Cuda.cpp?rev=342924&r1=342923&r2=342924&view=diff ============================================================================== --- cfe/trunk/lib/Driver/ToolChains/Cuda.cpp (original) +++ cfe/trunk/lib/Driver/ToolChains/Cuda.cpp Mon Sep 24 16:10:44 2018 @@ -59,6 +59,8 @@ static CudaVersion ParseCudaVersionFile( 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::CudaInstallati 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::addClangTargetOption // 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"; Modified: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h?rev=342924&r1=342923&r2=342924&view=diff ============================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h (original) +++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Mon Sep 24 16:10:44 2018 @@ -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 #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__ _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits