clang  15.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";
40  return "11.3";
42  return "11.4";
44  return "11.5";
45  case CudaVersion::NEW:
46  return "";
47  }
48  llvm_unreachable("invalid enum");
49 }
50 
51 CudaVersion CudaStringToVersion(const llvm::Twine &S) {
52  return llvm::StringSwitch<CudaVersion>(S.str())
53  .Case("7.0", CudaVersion::CUDA_70)
54  .Case("7.5", CudaVersion::CUDA_75)
55  .Case("8.0", CudaVersion::CUDA_80)
56  .Case("9.0", CudaVersion::CUDA_90)
57  .Case("9.1", CudaVersion::CUDA_91)
58  .Case("9.2", CudaVersion::CUDA_92)
59  .Case("10.0", CudaVersion::CUDA_100)
60  .Case("10.1", CudaVersion::CUDA_101)
61  .Case("10.2", CudaVersion::CUDA_102)
62  .Case("11.0", CudaVersion::CUDA_110)
63  .Case("11.1", CudaVersion::CUDA_111)
64  .Case("11.2", CudaVersion::CUDA_112)
65  .Case("11.3", CudaVersion::CUDA_113)
66  .Case("11.4", CudaVersion::CUDA_114)
67  .Case("11.5", CudaVersion::CUDA_115)
68  .Default(CudaVersion::UNKNOWN);
69 }
70 
71 namespace {
72 struct CudaArchToStringMap {
74  const char *arch_name;
75  const char *virtual_arch_name;
76 };
77 } // namespace
78 
79 #define SM2(sm, ca) \
80  { CudaArch::SM_##sm, "sm_" #sm, ca }
81 #define SM(sm) SM2(sm, "compute_" #sm)
82 #define GFX(gpu) \
83  { CudaArch::GFX##gpu, "gfx" #gpu, "compute_amdgcn" }
84 static const CudaArchToStringMap arch_names[] = {
85  // clang-format off
86  {CudaArch::UNUSED, "", ""},
87  SM2(20, "compute_20"), SM2(21, "compute_20"), // Fermi
88  SM(30), SM(32), SM(35), SM(37), // Kepler
89  SM(50), SM(52), SM(53), // Maxwell
90  SM(60), SM(61), SM(62), // Pascal
91  SM(70), SM(72), // Volta
92  SM(75), // Turing
93  SM(80), SM(86), // Ampere
94  GFX(600), // gfx600
95  GFX(601), // gfx601
96  GFX(602), // gfx602
97  GFX(700), // gfx700
98  GFX(701), // gfx701
99  GFX(702), // gfx702
100  GFX(703), // gfx703
101  GFX(704), // gfx704
102  GFX(705), // gfx705
103  GFX(801), // gfx801
104  GFX(802), // gfx802
105  GFX(803), // gfx803
106  GFX(805), // gfx805
107  GFX(810), // gfx810
108  GFX(900), // gfx900
109  GFX(902), // gfx902
110  GFX(904), // gfx903
111  GFX(906), // gfx906
112  GFX(908), // gfx908
113  GFX(909), // gfx909
114  GFX(90a), // gfx90a
115  GFX(90c), // gfx90c
116  GFX(940), // gfx940
117  GFX(1010), // gfx1010
118  GFX(1011), // gfx1011
119  GFX(1012), // gfx1012
120  GFX(1013), // gfx1013
121  GFX(1030), // gfx1030
122  GFX(1031), // gfx1031
123  GFX(1032), // gfx1032
124  GFX(1033), // gfx1033
125  GFX(1034), // gfx1034
126  GFX(1035), // gfx1035
127  GFX(1036), // gfx1036
128  GFX(1100), // gfx1100
129  GFX(1101), // gfx1101
130  GFX(1102), // gfx1102
131  GFX(1103), // gfx1103
132  {CudaArch::Generic, "generic", ""},
133  // clang-format on
134 };
135 #undef SM
136 #undef SM2
137 #undef GFX
138 
139 const char *CudaArchToString(CudaArch A) {
140  auto result = std::find_if(
141  std::begin(arch_names), std::end(arch_names),
142  [A](const CudaArchToStringMap &map) { return A == map.arch; });
143  if (result == std::end(arch_names))
144  return "unknown";
145  return result->arch_name;
146 }
147 
149  auto result = std::find_if(
150  std::begin(arch_names), std::end(arch_names),
151  [A](const CudaArchToStringMap &map) { return A == map.arch; });
152  if (result == std::end(arch_names))
153  return "unknown";
154  return result->virtual_arch_name;
155 }
156 
157 CudaArch StringToCudaArch(llvm::StringRef S) {
158  auto result = std::find_if(
159  std::begin(arch_names), std::end(arch_names),
160  [S](const CudaArchToStringMap &map) { return S == map.arch_name; });
161  if (result == std::end(arch_names))
162  return CudaArch::UNKNOWN;
163  return result->arch;
164 }
165 
167  if (A == CudaArch::UNKNOWN)
168  return CudaVersion::UNKNOWN;
169 
170  // AMD GPUs do not depend on CUDA versions.
171  if (IsAMDGpuArch(A))
172  return CudaVersion::CUDA_70;
173 
174  switch (A) {
175  case CudaArch::SM_20:
176  case CudaArch::SM_21:
177  case CudaArch::SM_30:
178  case CudaArch::SM_32:
179  case CudaArch::SM_35:
180  case CudaArch::SM_37:
181  case CudaArch::SM_50:
182  case CudaArch::SM_52:
183  case CudaArch::SM_53:
184  return CudaVersion::CUDA_70;
185  case CudaArch::SM_60:
186  case CudaArch::SM_61:
187  case CudaArch::SM_62:
188  return CudaVersion::CUDA_80;
189  case CudaArch::SM_70:
190  return CudaVersion::CUDA_90;
191  case CudaArch::SM_72:
192  return CudaVersion::CUDA_91;
193  case CudaArch::SM_75:
194  return CudaVersion::CUDA_100;
195  case CudaArch::SM_80:
196  return CudaVersion::CUDA_110;
197  case CudaArch::SM_86:
198  return CudaVersion::CUDA_111;
199  default:
200  llvm_unreachable("invalid enum");
201  }
202 }
203 
205  // AMD GPUs do not depend on CUDA versions.
206  if (IsAMDGpuArch(A))
207  return CudaVersion::NEW;
208 
209  switch (A) {
210  case CudaArch::UNKNOWN:
211  return CudaVersion::UNKNOWN;
212  case CudaArch::SM_20:
213  case CudaArch::SM_21:
214  return CudaVersion::CUDA_80;
215  case CudaArch::SM_30:
216  return CudaVersion::CUDA_110;
217  default:
218  return CudaVersion::NEW;
219  }
220 }
221 
222 CudaVersion ToCudaVersion(llvm::VersionTuple Version) {
223  int IVer = Version.getMajor() * 10 + Version.getMinor().value_or(0);
224  switch(IVer) {
225  case 70:
226  return CudaVersion::CUDA_70;
227  case 75:
228  return CudaVersion::CUDA_75;
229  case 80:
230  return CudaVersion::CUDA_80;
231  case 90:
232  return CudaVersion::CUDA_90;
233  case 91:
234  return CudaVersion::CUDA_91;
235  case 92:
236  return CudaVersion::CUDA_92;
237  case 100:
238  return CudaVersion::CUDA_100;
239  case 101:
240  return CudaVersion::CUDA_101;
241  case 102:
242  return CudaVersion::CUDA_102;
243  case 110:
244  return CudaVersion::CUDA_110;
245  case 111:
246  return CudaVersion::CUDA_111;
247  case 112:
248  return CudaVersion::CUDA_112;
249  case 113:
250  return CudaVersion::CUDA_113;
251  case 114:
252  return CudaVersion::CUDA_114;
253  case 115:
254  return CudaVersion::CUDA_115;
255  default:
256  return CudaVersion::UNKNOWN;
257  }
258 }
259 
260 bool CudaFeatureEnabled(llvm::VersionTuple Version, CudaFeature Feature) {
261  return CudaFeatureEnabled(ToCudaVersion(Version), Feature);
262 }
263 
265  switch (Feature) {
267  return Version >= CudaVersion::CUDA_92;
269  return Version >= CudaVersion::CUDA_101;
270  }
271  llvm_unreachable("Unknown CUDA feature.");
272 }
273 } // namespace clang
clang::CudaArch::SM_35
@ SM_35
clang::CudaVersion::CUDA_91
@ CUDA_91
clang::CudaArch::SM_70
@ SM_70
clang::CudaVersion::CUDA_114
@ CUDA_114
SM2
#define SM2(sm, ca)
Definition: Cuda.cpp:79
clang::CudaVersion::CUDA_100
@ CUDA_100
clang::CudaArchToVirtualArchString
const char * CudaArchToVirtualArchString(CudaArch A)
Definition: Cuda.cpp:148
clang::CudaVersion::NEW
@ NEW
clang::CudaFeature
CudaFeature
Definition: Cuda.h:134
clang::CudaVersion::CUDA_70
@ CUDA_70
clang::IsAMDGpuArch
static bool IsAMDGpuArch(CudaArch A)
Definition: Cuda.h:116
clang::arch_names
static const CudaArchToStringMap arch_names[]
Definition: Cuda.cpp:84
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:166
clang::StringToCudaArch
CudaArch StringToCudaArch(llvm::StringRef S)
Definition: Cuda.cpp:157
clang::CudaArch::SM_60
@ SM_60
clang::CudaArch::SM_20
@ SM_20
arch
CudaArch arch
Definition: Cuda.cpp:73
V
#define V(N, I)
Definition: ASTContext.h:3176
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:75
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:260
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
arch_name
const char * arch_name
Definition: Cuda.cpp:74
clang::CudaArch::SM_80
@ SM_80
clang::CudaStringToVersion
CudaVersion CudaStringToVersion(const llvm::Twine &S)
Definition: Cuda.cpp:51
clang::CudaArch::SM_30
@ SM_30
clang::CudaArchToString
const char * CudaArchToString(CudaArch A)
Definition: Cuda.cpp:139
clang::ToCudaVersion
CudaVersion ToCudaVersion(llvm::VersionTuple)
Definition: Cuda.cpp:222
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:204
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:82
clang::CudaArch
CudaArch
Definition: Cuda.h:46
clang::CudaArch::Generic
@ Generic
c
__device__ __2f16 float c
Definition: __clang_hip_libdevice_declares.h:315
clang::CudaArch::UNUSED
@ UNUSED
clang::CudaVersion::CUDA_115
@ CUDA_115
clang::CudaArch::SM_62
@ SM_62
clang::CudaVersion::CUDA_113
@ CUDA_113
SM
#define SM(sm)
Definition: Cuda.cpp:81
clang::CudaVersion::UNKNOWN
@ UNKNOWN
clang::CudaArch::SM_75
@ SM_75
clang::CudaVersion::CUDA_112
@ CUDA_112