r342924 - [CUDA] Added basic support for compiling with CUDA-10.0
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 24 16:10:44 PDT 2018
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__
More information about the cfe-commits
mailing list