clang  13.0.0git
Cuda.cpp
Go to the documentation of this file.
1 #include "clang/Basic/Cuda.h"
2 
3 #include "llvm/ADT/StringRef.h"
4 #include "llvm/ADT/StringSwitch.h"
5 #include "llvm/ADT/Twine.h"
6 #include "llvm/Support/ErrorHandling.h"
7 #include "llvm/Support/VersionTuple.h"
8 
9 namespace clang {
10 
12  switch (V) {
14  return "unknown";
16  return "7.0";
18  return "7.5";
20  return "8.0";
22  return "9.0";
24  return "9.1";
26  return "9.2";
28  return "10.0";
30  return "10.1";
32  return "10.2";
34  return "11.0";
36  return "11.1";
38  return "11.2";
39  }
40  llvm_unreachable("invalid enum");
41 }
42 
43 CudaVersion CudaStringToVersion(const llvm::Twine &S) {
44  return llvm::StringSwitch<CudaVersion>(S.str())
45  .Case("7.0", CudaVersion::CUDA_70)
46  .Case("7.5", CudaVersion::CUDA_75)
47  .Case("8.0", CudaVersion::CUDA_80)
48  .Case("9.0", CudaVersion::CUDA_90)
49  .Case("9.1", CudaVersion::CUDA_91)
50  .Case("9.2", CudaVersion::CUDA_92)
51  .Case("10.0", CudaVersion::CUDA_100)
52  .Case("10.1", CudaVersion::CUDA_101)
53  .Case("10.2", CudaVersion::CUDA_102)
54  .Case("11.0", CudaVersion::CUDA_110)
55  .Case("11.1", CudaVersion::CUDA_111)
56  .Case("11.2", CudaVersion::CUDA_112)
57  .Default(CudaVersion::UNKNOWN);
58 }
59 
60 namespace {
61 struct CudaArchToStringMap {
63  const char *arch_name;
64  const char *virtual_arch_name;
65 };
66 } // namespace
67 
68 #define SM2(sm, ca) \
69  { CudaArch::SM_##sm, "sm_" #sm, ca }
70 #define SM(sm) SM2(sm, "compute_" #sm)
71 #define GFX(gpu) \
72  { CudaArch::GFX##gpu, "gfx" #gpu, "compute_amdgcn" }
73 static const CudaArchToStringMap arch_names[] = {
74  // clang-format off
75  {CudaArch::UNUSED, "", ""},
76  SM2(20, "compute_20"), SM2(21, "compute_20"), // Fermi
77  SM(30), SM(32), SM(35), SM(37), // Kepler
78  SM(50), SM(52), SM(53), // Maxwell
79  SM(60), SM(61), SM(62), // Pascal
80  SM(70), SM(72), // Volta
81  SM(75), // Turing
82  SM(80), SM(86), // Ampere
83  GFX(600), // gfx600
84  GFX(601), // gfx601
85  GFX(602), // gfx602
86  GFX(700), // gfx700
87  GFX(701), // gfx701
88  GFX(702), // gfx702
89  GFX(703), // gfx703
90  GFX(704), // gfx704
91  GFX(705), // gfx705
92  GFX(801), // gfx801
93  GFX(802), // gfx802
94  GFX(803), // gfx803
95  GFX(805), // gfx805
96  GFX(810), // gfx810
97  GFX(900), // gfx900
98  GFX(902), // gfx902
99  GFX(904), // gfx903
100  GFX(906), // gfx906
101  GFX(908), // gfx908
102  GFX(909), // gfx909
103  GFX(90a), // gfx90a
104  GFX(90c), // gfx90c
105  GFX(1010), // gfx1010
106  GFX(1011), // gfx1011
107  GFX(1012), // gfx1012
108  GFX(1030), // gfx1030
109  GFX(1031), // gfx1031
110  GFX(1032), // gfx1032
111  GFX(1033), // gfx1033
112  GFX(1034), // gfx1034
113  // clang-format on
114 };
115 #undef SM
116 #undef SM2
117 #undef GFX
118 
119 const char *CudaArchToString(CudaArch A) {
120  auto result = std::find_if(
121  std::begin(arch_names), std::end(arch_names),
122  [A](const CudaArchToStringMap &map) { return A == map.arch; });
123  if (result == std::end(arch_names))
124  return "unknown";
125  return result->arch_name;
126 }
127 
129  auto result = std::find_if(
130  std::begin(arch_names), std::end(arch_names),
131  [A](const CudaArchToStringMap &map) { return A == map.arch; });
132  if (result == std::end(arch_names))
133  return "unknown";
134  return result->virtual_arch_name;
135 }
136 
137 CudaArch StringToCudaArch(llvm::StringRef S) {
138  auto result = std::find_if(
139  std::begin(arch_names), std::end(arch_names),
140  [S](const CudaArchToStringMap &map) { return S == map.arch_name; });
141  if (result == std::end(arch_names))
142  return CudaArch::UNKNOWN;
143  return result->arch;
144 }
145 
147  if (A == CudaArch::UNKNOWN)
148  return CudaVersion::UNKNOWN;
149 
150  // AMD GPUs do not depend on CUDA versions.
151  if (IsAMDGpuArch(A))
152  return CudaVersion::CUDA_70;
153 
154  switch (A) {
155  case CudaArch::SM_20:
156  case CudaArch::SM_21:
157  case CudaArch::SM_30:
158  case CudaArch::SM_32:
159  case CudaArch::SM_35:
160  case CudaArch::SM_37:
161  case CudaArch::SM_50:
162  case CudaArch::SM_52:
163  case CudaArch::SM_53:
164  return CudaVersion::CUDA_70;
165  case CudaArch::SM_60:
166  case CudaArch::SM_61:
167  case CudaArch::SM_62:
168  return CudaVersion::CUDA_80;
169  case CudaArch::SM_70:
170  return CudaVersion::CUDA_90;
171  case CudaArch::SM_72:
172  return CudaVersion::CUDA_91;
173  case CudaArch::SM_75:
174  return CudaVersion::CUDA_100;
175  case CudaArch::SM_80:
176  return CudaVersion::CUDA_110;
177  case CudaArch::SM_86:
178  return CudaVersion::CUDA_111;
179  default:
180  llvm_unreachable("invalid enum");
181  }
182 }
183 
185  // AMD GPUs do not depend on CUDA versions.
186  if (IsAMDGpuArch(A))
187  return CudaVersion::LATEST;
188 
189  switch (A) {
190  case CudaArch::UNKNOWN:
191  return CudaVersion::UNKNOWN;
192  case CudaArch::SM_20:
193  case CudaArch::SM_21:
194  return CudaVersion::CUDA_80;
195  default:
196  return CudaVersion::LATEST;
197  }
198 }
199 
200 CudaVersion ToCudaVersion(llvm::VersionTuple Version) {
201  int IVer =
202  Version.getMajor() * 10 + Version.getMinor().getValueOr(0);
203  switch(IVer) {
204  case 70:
205  return CudaVersion::CUDA_70;
206  case 75:
207  return CudaVersion::CUDA_75;
208  case 80:
209  return CudaVersion::CUDA_80;
210  case 90:
211  return CudaVersion::CUDA_90;
212  case 91:
213  return CudaVersion::CUDA_91;
214  case 92:
215  return CudaVersion::CUDA_92;
216  case 100:
217  return CudaVersion::CUDA_100;
218  case 101:
219  return CudaVersion::CUDA_101;
220  case 102:
221  return CudaVersion::CUDA_102;
222  case 110:
223  return CudaVersion::CUDA_110;
224  case 111:
225  return CudaVersion::CUDA_111;
226  case 112:
227  return CudaVersion::CUDA_112;
228  default:
229  return CudaVersion::UNKNOWN;
230  }
231 }
232 
233 bool CudaFeatureEnabled(llvm::VersionTuple Version, CudaFeature Feature) {
234  return CudaFeatureEnabled(ToCudaVersion(Version), Feature);
235 }
236 
238  switch (Feature) {
240  return Version >= CudaVersion::CUDA_92;
242  return Version >= CudaVersion::CUDA_101;
243  }
244  llvm_unreachable("Unknown CUDA feature.");
245 }
246 } // namespace clang
clang::CudaArch::SM_35
@ SM_35
clang::CudaVersion::CUDA_91
@ CUDA_91
clang::CudaArch::SM_70
@ SM_70
SM2
#define SM2(sm, ca)
Definition: Cuda.cpp:68
clang::CudaVersion::CUDA_100
@ CUDA_100
clang::CudaArchToVirtualArchString
const char * CudaArchToVirtualArchString(CudaArch A)
Definition: Cuda.cpp:128
clang::CudaFeature
CudaFeature
Definition: Cuda.h:115
clang::CudaVersion::CUDA_70
@ CUDA_70
clang::IsAMDGpuArch
static bool IsAMDGpuArch(CudaArch A)
Definition: Cuda.h:98
clang::arch_names
static const CudaArchToStringMap arch_names[]
Definition: Cuda.cpp:73
clang::CudaArch::SM_53
@ SM_53
clang::CudaVersion::CUDA_101
@ CUDA_101
clang::MinVersionForCudaArch
CudaVersion MinVersionForCudaArch(CudaArch A)
Get the earliest CudaVersion that supports the given CudaArch.
Definition: Cuda.cpp:146
clang::StringToCudaArch
CudaArch StringToCudaArch(llvm::StringRef S)
Definition: Cuda.cpp:137
clang::CudaArch::SM_60
@ SM_60
clang::CudaArch::SM_20
@ SM_20
arch
CudaArch arch
Definition: Cuda.cpp:62
V
#define V(N, I)
Definition: ASTContext.h:3039
clang::CudaVersion::CUDA_90
@ CUDA_90
clang::CudaArch::SM_86
@ SM_86
clang::CudaFeature::CUDA_USES_FATBIN_REGISTER_END
@ CUDA_USES_FATBIN_REGISTER_END
clang::CudaArch::SM_72
@ SM_72
clang::CudaVersion::CUDA_102
@ CUDA_102
virtual_arch_name
const char * virtual_arch_name
Definition: Cuda.cpp:64
clang::CudaVersion::CUDA_110
@ CUDA_110
clang::CudaVersion
CudaVersion
Definition: Cuda.h:20
clang::CudaArch::SM_21
@ SM_21
clang::CudaFeatureEnabled
bool CudaFeatureEnabled(llvm::VersionTuple, CudaFeature)
Definition: Cuda.cpp:233
clang::CudaFeature::CUDA_USES_NEW_LAUNCH
@ CUDA_USES_NEW_LAUNCH
clang::CudaVersion::CUDA_75
@ CUDA_75
clang::CudaArch::SM_52
@ SM_52
clang::CudaArch::SM_61
@ SM_61
clang::CudaVersion::CUDA_111
@ CUDA_111
clang::CudaVersion::LATEST
@ LATEST
arch_name
const char * arch_name
Definition: Cuda.cpp:63
clang::CudaArch::SM_80
@ SM_80
clang::CudaStringToVersion
CudaVersion CudaStringToVersion(const llvm::Twine &S)
Definition: Cuda.cpp:43
clang::CudaArch::SM_30
@ SM_30
clang::CudaArchToString
const char * CudaArchToString(CudaArch A)
Definition: Cuda.cpp:119
clang::ToCudaVersion
CudaVersion ToCudaVersion(llvm::VersionTuple)
Definition: Cuda.cpp:200
clang::CudaArch::UNKNOWN
@ UNKNOWN
clang::CudaVersion::CUDA_92
@ CUDA_92
clang::CudaArch::SM_32
@ SM_32
clang::CudaVersionToString
const char * CudaVersionToString(CudaVersion V)
Definition: Cuda.cpp:11
Cuda.h
clang::MaxVersionForCudaArch
CudaVersion MaxVersionForCudaArch(CudaArch A)
Get the latest CudaVersion that supports the given CudaArch.
Definition: Cuda.cpp:184
clang
Definition: CalledOnceCheck.h:17
clang::CudaVersion::CUDA_80
@ CUDA_80
clang::CudaArch::SM_37
@ SM_37
clang::CudaArch::SM_50
@ SM_50
GFX
#define GFX(gpu)
Definition: Cuda.cpp:71
clang::CudaArch
CudaArch
Definition: Cuda.h:41
c
__device__ __2f16 float c
Definition: __clang_hip_libdevice_declares.h:315
clang::CudaArch::UNUSED
@ UNUSED
clang::CudaArch::SM_62
@ SM_62
SM
#define SM(sm)
Definition: Cuda.cpp:70
clang::CudaVersion::UNKNOWN
@ UNKNOWN
clang::CudaArch::SM_75
@ SM_75
clang::CudaVersion::CUDA_112
@ CUDA_112