clang 23.0.0git
OffloadArch.cpp
Go to the documentation of this file.
1//===----------------------------------------------------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
9
10#include "llvm/ADT/STLExtras.h"
11#include "llvm/ADT/StringRef.h"
12#include "llvm/TargetParser/Triple.h"
13
14namespace clang {
15
16namespace {
17struct OffloadArchToStringMap {
18 OffloadArch Arch;
19 const char *ArchName;
20 const char *VirtualArchName;
21};
22} // namespace
23
24#define SM(sm) {OffloadArch::SM_##sm, "sm_" #sm, "compute_" #sm}
25#define GFX(gpu) {OffloadArch::GFX##gpu, "gfx" #gpu, "compute_amdgcn"}
26static const OffloadArchToStringMap ArchNames[] = {
27 // clang-format off
28 {OffloadArch::Unused, "", ""},
29 SM(20), {OffloadArch::SM_21, "sm_21", "compute_20"}, // Fermi
30 SM(30), {OffloadArch::SM_32_, "sm_32", "compute_32"}, SM(35), SM(37), // Kepler
31 SM(50), SM(52), SM(53), // Maxwell
32 SM(60), SM(61), SM(62), // Pascal
33 SM(70), SM(72), // Volta
34 SM(75), // Turing
35 SM(80), SM(86), // Ampere
36 SM(87), // Jetson/Drive AGX Orin
37 SM(88), // Ampere
38 SM(89), // Ada Lovelace
39 SM(90), // Hopper
40 SM(90a), // Hopper
41 SM(100), // Blackwell
42 SM(100a), // Blackwell
43 SM(101), // Blackwell
44 SM(101a), // Blackwell
45 SM(103), // Blackwell
46 SM(103a), // Blackwell
47 SM(110), // Blackwell
48 SM(110a), // Blackwell
49 SM(120), // Blackwell
50 SM(120a), // Blackwell
51 SM(121), // Blackwell
52 SM(121a), // Blackwell
53 GFX(600), // gfx600
54 GFX(601), // gfx601
55 GFX(602), // gfx602
56 GFX(700), // gfx700
57 GFX(701), // gfx701
58 GFX(702), // gfx702
59 GFX(703), // gfx703
60 GFX(704), // gfx704
61 GFX(705), // gfx705
62 GFX(801), // gfx801
63 GFX(802), // gfx802
64 GFX(803), // gfx803
65 GFX(805), // gfx805
66 GFX(810), // gfx810
67 {OffloadArch::GFX9_GENERIC, "gfx9-generic", "compute_amdgcn"},
68 GFX(900), // gfx900
69 GFX(902), // gfx902
70 GFX(904), // gfx903
71 GFX(906), // gfx906
72 GFX(908), // gfx908
73 GFX(909), // gfx909
74 GFX(90a), // gfx90a
75 GFX(90c), // gfx90c
76 {OffloadArch::GFX9_4_GENERIC, "gfx9-4-generic", "compute_amdgcn"},
77 GFX(942), // gfx942
78 GFX(950), // gfx950
79 {OffloadArch::GFX10_1_GENERIC, "gfx10-1-generic", "compute_amdgcn"},
80 GFX(1010), // gfx1010
81 GFX(1011), // gfx1011
82 GFX(1012), // gfx1012
83 GFX(1013), // gfx1013
84 {OffloadArch::GFX10_3_GENERIC, "gfx10-3-generic", "compute_amdgcn"},
85 GFX(1030), // gfx1030
86 GFX(1031), // gfx1031
87 GFX(1032), // gfx1032
88 GFX(1033), // gfx1033
89 GFX(1034), // gfx1034
90 GFX(1035), // gfx1035
91 GFX(1036), // gfx1036
92 {OffloadArch::GFX11_GENERIC, "gfx11-generic", "compute_amdgcn"},
93 GFX(1100), // gfx1100
94 GFX(1101), // gfx1101
95 GFX(1102), // gfx1102
96 GFX(1103), // gfx1103
97 GFX(1150), // gfx1150
98 GFX(1151), // gfx1151
99 GFX(1152), // gfx1152
100 GFX(1153), // gfx1153
101 GFX(1170), // gfx1170
102 GFX(1171), // gfx1171
103 GFX(1172), // gfx1172
104 {OffloadArch::GFX12_GENERIC, "gfx12-generic", "compute_amdgcn"},
105 GFX(1200), // gfx1200
106 GFX(1201), // gfx1201
107 GFX(1250), // gfx1250
108 GFX(1251), // gfx1251
109 {OffloadArch::GFX12_5_GENERIC, "gfx12-5-generic", "compute_amdgcn"},
110 GFX(1310), // gfx1310
111 {OffloadArch::AMDGCNSPIRV, "amdgcnspirv", "compute_amdgcn"},
112 // Intel CPUs
113 {OffloadArch::GRANITERAPIDS, "graniterapids", ""},
114 // Intel GPUS
115 {OffloadArch::BMG_G21, "bmg_g21", ""},
116 {OffloadArch::Generic, "generic", ""},
117 // clang-format on
118};
119#undef SM
120#undef GFX
121
123 auto Result =
124 llvm::find_if(ArchNames, [A](const OffloadArchToStringMap &Map) {
125 return A == Map.Arch;
126 });
127 if (Result == std::end(ArchNames))
128 return "unknown";
129 return Result->ArchName;
130}
131
133 auto Result =
134 llvm::find_if(ArchNames, [A](const OffloadArchToStringMap &Map) {
135 return A == Map.Arch;
136 });
137 if (Result == std::end(ArchNames))
138 return "unknown";
139 return Result->VirtualArchName;
140}
141
143 auto Result =
144 llvm::find_if(ArchNames, [S](const OffloadArchToStringMap &Map) {
145 return S == Map.ArchName;
146 });
147 if (Result == std::end(ArchNames))
149 return Result->Arch;
150}
151
152llvm::Triple OffloadArchToTriple(const llvm::Triple &DefaultToolchainTriple,
153 OffloadArch ID) {
154 if (ID == OffloadArch::AMDGCNSPIRV)
155 return llvm::Triple("spirv64-amd-amdhsa");
156
157 if (IsNVIDIAOffloadArch(ID))
158 return DefaultToolchainTriple.isArch64Bit()
159 ? llvm::Triple("nvptx64-nvidia-cuda")
160 : llvm::Triple("nvptx-nvidia-cuda");
161
162 if (IsAMDOffloadArch(ID))
163 return llvm::Triple("amdgcn-amd-amdhsa");
164
165 return {};
166}
167
168} // namespace clang
#define GFX(gpu)
#define SM(sm)
__device__ __2f16 float c
The JSON file list parser is used to communicate input to InstallAPI.
static const OffloadArchToStringMap ArchNames[]
static bool IsAMDOffloadArch(OffloadArch A)
@ Result
The result type of a method or function.
Definition TypeBase.h:905
static bool IsNVIDIAOffloadArch(OffloadArch A)
const char * OffloadArchToVirtualArchString(OffloadArch A)
OffloadArch StringToOffloadArch(llvm::StringRef S)
const char * OffloadArchToString(OffloadArch A)
llvm::Triple OffloadArchToTriple(const llvm::Triple &DefaultToolchainTriple, OffloadArch ID)