clang  16.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  const char *Name;
14  llvm::VersionTuple TVersion;
15 };
16 #define CUDA_ENTRY(major, minor) \
17  { \
18 #major "." #minor, CudaVersion::CUDA_##major##minor, \
19  llvm::VersionTuple(major, minor) \
20  }
21 
23  CUDA_ENTRY(7, 0),
24  CUDA_ENTRY(7, 5),
25  CUDA_ENTRY(8, 0),
26  CUDA_ENTRY(9, 0),
27  CUDA_ENTRY(9, 1),
28  CUDA_ENTRY(9, 2),
29  CUDA_ENTRY(10, 0),
30  CUDA_ENTRY(10, 1),
31  CUDA_ENTRY(10, 2),
32  CUDA_ENTRY(11, 0),
33  CUDA_ENTRY(11, 1),
34  CUDA_ENTRY(11, 2),
35  CUDA_ENTRY(11, 3),
36  CUDA_ENTRY(11, 4),
37  CUDA_ENTRY(11, 5),
38  CUDA_ENTRY(11, 6),
39  CUDA_ENTRY(11, 7),
40  CUDA_ENTRY(11, 8),
41  {"", CudaVersion::NEW, llvm::VersionTuple(std::numeric_limits<int>::max())},
42  {"unknown", CudaVersion::UNKNOWN, {}} // End of list tombstone.
43 };
44 #undef CUDA_ENTRY
45 
47  for (auto *I = CudaNameVersionMap; I->Version != CudaVersion::UNKNOWN; ++I)
48  if (I->Version == V)
49  return I->Name;
50 
52 }
53 
54 CudaVersion CudaStringToVersion(const llvm::Twine &S) {
55  std::string VS = S.str();
56  for (auto *I = CudaNameVersionMap; I->Version != CudaVersion::UNKNOWN; ++I)
57  if (I->Name == VS)
58  return I->Version;
59  return CudaVersion::UNKNOWN;
60 }
61 
62 CudaVersion ToCudaVersion(llvm::VersionTuple Version) {
63  for (auto *I = CudaNameVersionMap; I->Version != CudaVersion::UNKNOWN; ++I)
64  if (I->TVersion == Version)
65  return I->Version;
66  return CudaVersion::UNKNOWN;
67 }
68 
69 namespace {
70 struct CudaArchToStringMap {
72  const char *arch_name;
73  const char *virtual_arch_name;
74 };
75 } // namespace
76 
77 #define SM2(sm, ca) \
78  { CudaArch::SM_##sm, "sm_" #sm, ca }
79 #define SM(sm) SM2(sm, "compute_" #sm)
80 #define GFX(gpu) \
81  { CudaArch::GFX##gpu, "gfx" #gpu, "compute_amdgcn" }
82 static const CudaArchToStringMap arch_names[] = {
83  // clang-format off
84  {CudaArch::UNUSED, "", ""},
85  SM2(20, "compute_20"), SM2(21, "compute_20"), // Fermi
86  SM(30), SM(32), SM(35), SM(37), // Kepler
87  SM(50), SM(52), SM(53), // Maxwell
88  SM(60), SM(61), SM(62), // Pascal
89  SM(70), SM(72), // Volta
90  SM(75), // Turing
91  SM(80), SM(86), // Ampere
92  SM(87), // Jetson/Drive AGX Orin
93  SM(89), // Ada Lovelace
94  SM(90), // Hopper
95  GFX(600), // gfx600
96  GFX(601), // gfx601
97  GFX(602), // gfx602
98  GFX(700), // gfx700
99  GFX(701), // gfx701
100  GFX(702), // gfx702
101  GFX(703), // gfx703
102  GFX(704), // gfx704
103  GFX(705), // gfx705
104  GFX(801), // gfx801
105  GFX(802), // gfx802
106  GFX(803), // gfx803
107  GFX(805), // gfx805
108  GFX(810), // gfx810
109  GFX(900), // gfx900
110  GFX(902), // gfx902
111  GFX(904), // gfx903
112  GFX(906), // gfx906
113  GFX(908), // gfx908
114  GFX(909), // gfx909
115  GFX(90a), // gfx90a
116  GFX(90c), // gfx90c
117  GFX(940), // gfx940
118  GFX(1010), // gfx1010
119  GFX(1011), // gfx1011
120  GFX(1012), // gfx1012
121  GFX(1013), // gfx1013
122  GFX(1030), // gfx1030
123  GFX(1031), // gfx1031
124  GFX(1032), // gfx1032
125  GFX(1033), // gfx1033
126  GFX(1034), // gfx1034
127  GFX(1035), // gfx1035
128  GFX(1036), // gfx1036
129  GFX(1100), // gfx1100
130  GFX(1101), // gfx1101
131  GFX(1102), // gfx1102
132  GFX(1103), // gfx1103
133  {CudaArch::Generic, "generic", ""},
134  // clang-format on
135 };
136 #undef SM
137 #undef SM2
138 #undef GFX
139 
140 const char *CudaArchToString(CudaArch A) {
141  auto result = std::find_if(
142  std::begin(arch_names), std::end(arch_names),
143  [A](const CudaArchToStringMap &map) { return A == map.arch; });
144  if (result == std::end(arch_names))
145  return "unknown";
146  return result->arch_name;
147 }
148 
150  auto result = std::find_if(
151  std::begin(arch_names), std::end(arch_names),
152  [A](const CudaArchToStringMap &map) { return A == map.arch; });
153  if (result == std::end(arch_names))
154  return "unknown";
155  return result->virtual_arch_name;
156 }
157 
158 CudaArch StringToCudaArch(llvm::StringRef S) {
159  auto result = std::find_if(
160  std::begin(arch_names), std::end(arch_names),
161  [S](const CudaArchToStringMap &map) { return S == map.arch_name; });
162  if (result == std::end(arch_names))
163  return CudaArch::UNKNOWN;
164  return result->arch;
165 }
166 
168  if (A == CudaArch::UNKNOWN)
169  return CudaVersion::UNKNOWN;
170 
171  // AMD GPUs do not depend on CUDA versions.
172  if (IsAMDGpuArch(A))
173  return CudaVersion::CUDA_70;
174 
175  switch (A) {
176  case CudaArch::SM_20:
177  case CudaArch::SM_21:
178  case CudaArch::SM_30:
179  case CudaArch::SM_32:
180  case CudaArch::SM_35:
181  case CudaArch::SM_37:
182  case CudaArch::SM_50:
183  case CudaArch::SM_52:
184  case CudaArch::SM_53:
185  return CudaVersion::CUDA_70;
186  case CudaArch::SM_60:
187  case CudaArch::SM_61:
188  case CudaArch::SM_62:
189  return CudaVersion::CUDA_80;
190  case CudaArch::SM_70:
191  return CudaVersion::CUDA_90;
192  case CudaArch::SM_72:
193  return CudaVersion::CUDA_91;
194  case CudaArch::SM_75:
195  return CudaVersion::CUDA_100;
196  case CudaArch::SM_80:
197  return CudaVersion::CUDA_110;
198  case CudaArch::SM_86:
199  return CudaVersion::CUDA_111;
200  case CudaArch::SM_87:
201  return CudaVersion::CUDA_114;
202  case CudaArch::SM_89:
203  case CudaArch::SM_90:
204  return CudaVersion::CUDA_118;
205  default:
206  llvm_unreachable("invalid enum");
207  }
208 }
209 
211  // AMD GPUs do not depend on CUDA versions.
212  if (IsAMDGpuArch(A))
213  return CudaVersion::NEW;
214 
215  switch (A) {
216  case CudaArch::UNKNOWN:
217  return CudaVersion::UNKNOWN;
218  case CudaArch::SM_20:
219  case CudaArch::SM_21:
220  return CudaVersion::CUDA_80;
221  case CudaArch::SM_30:
222  return CudaVersion::CUDA_110;
223  default:
224  return CudaVersion::NEW;
225  }
226 }
227 
228 bool CudaFeatureEnabled(llvm::VersionTuple Version, CudaFeature Feature) {
229  return CudaFeatureEnabled(ToCudaVersion(Version), Feature);
230 }
231 
233  switch (Feature) {
235  return Version >= CudaVersion::CUDA_92;
237  return Version >= CudaVersion::CUDA_101;
238  }
239  llvm_unreachable("Unknown CUDA feature.");
240 }
241 } // 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
max
__DEVICE__ int max(int __a, int __b)
Definition: __clang_cuda_math.h:196
SM2
#define SM2(sm, ca)
Definition: Cuda.cpp:77
clang::CudaVersion::CUDA_100
@ CUDA_100
clang::CudaArchToVirtualArchString
const char * CudaArchToVirtualArchString(CudaArch A)
Definition: Cuda.cpp:149
clang::CudaVersion::NEW
@ NEW
clang::CudaFeature
CudaFeature
Definition: Cuda.h:140
string
string(SUBSTRING ${CMAKE_CURRENT_BINARY_DIR} 0 ${PATH_LIB_START} PATH_HEAD) string(SUBSTRING $
Definition: CMakeLists.txt:22
clang::CudaVersion::CUDA_70
@ CUDA_70
clang::IsAMDGpuArch
static bool IsAMDGpuArch(CudaArch A)
Definition: Cuda.h:122
clang::CudaVersionMapEntry
Definition: Cuda.cpp:11
clang::CudaArch::SM_90
@ SM_90
clang::arch_names
static const CudaArchToStringMap arch_names[]
Definition: Cuda.cpp:82
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:167
clang::StringToCudaArch
CudaArch StringToCudaArch(llvm::StringRef S)
Definition: Cuda.cpp:158
clang::CudaArch::SM_60
@ SM_60
clang::CudaArch::SM_20
@ SM_20
arch
CudaArch arch
Definition: Cuda.cpp:71
V
#define V(N, I)
Definition: ASTContext.h:3237
clang::CudaVersion::CUDA_90
@ CUDA_90
clang::CudaArch::SM_86
@ SM_86
clang::CudaVersionMapEntry::Version
CudaVersion Version
Definition: Cuda.cpp:13
clang::CudaFeature::CUDA_USES_FATBIN_REGISTER_END
@ CUDA_USES_FATBIN_REGISTER_END
clang::CudaArch::SM_72
@ SM_72
clang::CudaArch::SM_87
@ SM_87
virtual_arch_name
const char * virtual_arch_name
Definition: Cuda.cpp:73
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:228
clang::CudaFeature::CUDA_USES_NEW_LAUNCH
@ CUDA_USES_NEW_LAUNCH
clang::CudaArch::SM_52
@ SM_52
CUDA_ENTRY
#define CUDA_ENTRY(major, minor)
Definition: Cuda.cpp:16
clang::CudaArch::SM_61
@ SM_61
clang::CudaVersion::CUDA_118
@ CUDA_118
clang::CudaVersionMapEntry::TVersion
llvm::VersionTuple TVersion
Definition: Cuda.cpp:14
clang::CudaVersion::CUDA_111
@ CUDA_111
clang::CudaNameVersionMap
static const CudaVersionMapEntry CudaNameVersionMap[]
Definition: Cuda.cpp:22
arch_name
const char * arch_name
Definition: Cuda.cpp:72
clang::CudaArch::SM_80
@ SM_80
clang::CudaStringToVersion
CudaVersion CudaStringToVersion(const llvm::Twine &S)
Definition: Cuda.cpp:54
clang::CudaArch::SM_30
@ SM_30
clang::CudaArchToString
const char * CudaArchToString(CudaArch A)
Definition: Cuda.cpp:140
clang::ToCudaVersion
CudaVersion ToCudaVersion(llvm::VersionTuple)
Definition: Cuda.cpp:62
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:46
Cuda.h
clang::MaxVersionForCudaArch
CudaVersion MaxVersionForCudaArch(CudaArch A)
Get the latest CudaVersion that supports the given CudaArch.
Definition: Cuda.cpp:210
clang
Definition: CalledOnceCheck.h:17
clang::CudaVersion::CUDA_80
@ CUDA_80
clang::CudaVersionMapEntry::Name
const char * Name
Definition: Cuda.cpp:12
clang::CudaArch::SM_37
@ SM_37
clang::CudaArch::SM_50
@ SM_50
GFX
#define GFX(gpu)
Definition: Cuda.cpp:80
clang::CudaArch
CudaArch
Definition: Cuda.h:49
clang::CudaArch::Generic
@ Generic
c
__device__ __2f16 float c
Definition: __clang_hip_libdevice_declares.h:320
clang::CudaArch::UNUSED
@ UNUSED
clang::CudaArch::SM_62
@ SM_62
SM
#define SM(sm)
Definition: Cuda.cpp:79
clang::CudaVersion::UNKNOWN
@ UNKNOWN
clang::CudaArch::SM_75
@ SM_75
clang::CudaArch::SM_89
@ SM_89