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 
62  const char *arch_name;
63  const char *virtual_arch_name;
64 };
65 
66 #define SM2(sm, ca) \
67  { CudaArch::SM_##sm, "sm_" #sm, ca }
68 #define SM(sm) SM2(sm, "compute_" #sm)
69 #define GFX(gpu) \
70  { CudaArch::GFX##gpu, "gfx" #gpu, "compute_amdgcn" }
72  // clang-format off
73  {CudaArch::UNUSED, "", ""},
74  SM2(20, "compute_20"), SM2(21, "compute_20"), // Fermi
75  SM(30), SM(32), SM(35), SM(37), // Kepler
76  SM(50), SM(52), SM(53), // Maxwell
77  SM(60), SM(61), SM(62), // Pascal
78  SM(70), SM(72), // Volta
79  SM(75), // Turing
80  SM(80), SM(86), // Ampere
81  GFX(600), // gfx600
82  GFX(601), // gfx601
83  GFX(602), // gfx602
84  GFX(700), // gfx700
85  GFX(701), // gfx701
86  GFX(702), // gfx702
87  GFX(703), // gfx703
88  GFX(704), // gfx704
89  GFX(705), // gfx705
90  GFX(801), // gfx801
91  GFX(802), // gfx802
92  GFX(803), // gfx803
93  GFX(805), // gfx805
94  GFX(810), // gfx810
95  GFX(900), // gfx900
96  GFX(902), // gfx902
97  GFX(904), // gfx903
98  GFX(906), // gfx906
99  GFX(908), // gfx908
100  GFX(909), // gfx909
101  GFX(90a), // gfx90a
102  GFX(90c), // gfx90c
103  GFX(1010), // gfx1010
104  GFX(1011), // gfx1011
105  GFX(1012), // gfx1012
106  GFX(1030), // gfx1030
107  GFX(1031), // gfx1031
108  GFX(1032), // gfx1032
109  GFX(1033), // gfx1033
110  // clang-format on
111 };
112 #undef SM
113 #undef SM2
114 #undef GFX
115 
116 const char *CudaArchToString(CudaArch A) {
117  auto result = std::find_if(
118  std::begin(arch_names), std::end(arch_names),
119  [A](const CudaArchToStringMap &map) { return A == map.arch; });
120  if (result == std::end(arch_names))
121  return "unknown";
122  return result->arch_name;
123 }
124 
126  auto result = std::find_if(
127  std::begin(arch_names), std::end(arch_names),
128  [A](const CudaArchToStringMap &map) { return A == map.arch; });
129  if (result == std::end(arch_names))
130  return "unknown";
131  return result->virtual_arch_name;
132 }
133 
134 CudaArch StringToCudaArch(llvm::StringRef S) {
135  auto result = std::find_if(
136  std::begin(arch_names), std::end(arch_names),
137  [S](const CudaArchToStringMap &map) { return S == map.arch_name; });
138  if (result == std::end(arch_names))
139  return CudaArch::UNKNOWN;
140  return result->arch;
141 }
142 
144  if (A == CudaArch::UNKNOWN)
145  return CudaVersion::UNKNOWN;
146 
147  // AMD GPUs do not depend on CUDA versions.
148  if (IsAMDGpuArch(A))
149  return CudaVersion::CUDA_70;
150 
151  switch (A) {
152  case CudaArch::SM_20:
153  case CudaArch::SM_21:
154  case CudaArch::SM_30:
155  case CudaArch::SM_32:
156  case CudaArch::SM_35:
157  case CudaArch::SM_37:
158  case CudaArch::SM_50:
159  case CudaArch::SM_52:
160  case CudaArch::SM_53:
161  return CudaVersion::CUDA_70;
162  case CudaArch::SM_60:
163  case CudaArch::SM_61:
164  case CudaArch::SM_62:
165  return CudaVersion::CUDA_80;
166  case CudaArch::SM_70:
167  return CudaVersion::CUDA_90;
168  case CudaArch::SM_72:
169  return CudaVersion::CUDA_91;
170  case CudaArch::SM_75:
171  return CudaVersion::CUDA_100;
172  case CudaArch::SM_80:
173  return CudaVersion::CUDA_110;
174  case CudaArch::SM_86:
175  return CudaVersion::CUDA_111;
176  default:
177  llvm_unreachable("invalid enum");
178  }
179 }
180 
182  // AMD GPUs do not depend on CUDA versions.
183  if (IsAMDGpuArch(A))
184  return CudaVersion::LATEST;
185 
186  switch (A) {
187  case CudaArch::UNKNOWN:
188  return CudaVersion::UNKNOWN;
189  case CudaArch::SM_20:
190  case CudaArch::SM_21:
191  return CudaVersion::CUDA_80;
192  default:
193  return CudaVersion::LATEST;
194  }
195 }
196 
197 CudaVersion ToCudaVersion(llvm::VersionTuple Version) {
198  int IVer =
199  Version.getMajor() * 10 + Version.getMinor().getValueOr(0);
200  switch(IVer) {
201  case 70:
202  return CudaVersion::CUDA_70;
203  case 75:
204  return CudaVersion::CUDA_75;
205  case 80:
206  return CudaVersion::CUDA_80;
207  case 90:
208  return CudaVersion::CUDA_90;
209  case 91:
210  return CudaVersion::CUDA_91;
211  case 92:
212  return CudaVersion::CUDA_92;
213  case 100:
214  return CudaVersion::CUDA_100;
215  case 101:
216  return CudaVersion::CUDA_101;
217  case 102:
218  return CudaVersion::CUDA_102;
219  case 110:
220  return CudaVersion::CUDA_110;
221  case 111:
222  return CudaVersion::CUDA_111;
223  case 112:
224  return CudaVersion::CUDA_112;
225  default:
226  return CudaVersion::UNKNOWN;
227  }
228 }
229 
230 bool CudaFeatureEnabled(llvm::VersionTuple Version, CudaFeature Feature) {
231  return CudaFeatureEnabled(ToCudaVersion(Version), Feature);
232 }
233 
235  switch (Feature) {
237  return Version >= CudaVersion::CUDA_92;
239  return Version >= CudaVersion::CUDA_101;
240  }
241  llvm_unreachable("Unknown CUDA feature.");
242 }
243 } // 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:66
clang::CudaVersion::CUDA_100
@ CUDA_100
clang::CudaArchToStringMap::arch
CudaArch arch
Definition: Cuda.cpp:61
clang::CudaArchToVirtualArchString
const char * CudaArchToVirtualArchString(CudaArch A)
Definition: Cuda.cpp:125
clang::CudaFeature
CudaFeature
Definition: Cuda.h:114
clang::CudaVersion::CUDA_70
@ CUDA_70
clang::IsAMDGpuArch
static bool IsAMDGpuArch(CudaArch A)
Definition: Cuda.h:97
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:143
clang::StringToCudaArch
CudaArch StringToCudaArch(llvm::StringRef S)
Definition: Cuda.cpp:134
clang::CudaArch::SM_60
@ SM_60
clang::CudaArch::SM_20
@ SM_20
V
#define V(N, I)
Definition: ASTContext.h:3018
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
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:230
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
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:116
clang::ToCudaVersion
CudaVersion ToCudaVersion(llvm::VersionTuple)
Definition: Cuda.cpp:197
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:181
clang::arch_names
CudaArchToStringMap arch_names[]
Definition: Cuda.cpp:71
clang
Dataflow Directional Tag Classes.
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:69
clang::CudaArchToStringMap::arch_name
const char * arch_name
Definition: Cuda.cpp:62
clang::CudaArch
CudaArch
Definition: Cuda.h:41
clang::CudaArchToStringMap
Definition: Cuda.cpp:60
c
__device__ __2f16 float c
Definition: __clang_hip_libdevice_declares.h:303
clang::CudaArch::UNUSED
@ UNUSED
clang::CudaArch::SM_62
@ SM_62
SM
#define SM(sm)
Definition: Cuda.cpp:68
clang::CudaVersion::UNKNOWN
@ UNKNOWN
clang::CudaArch::SM_75
@ SM_75
clang::CudaArchToStringMap::virtual_arch_name
const char * virtual_arch_name
Definition: Cuda.cpp:63
clang::CudaVersion::CUDA_112
@ CUDA_112