1//===--- SPIR.h - Declare SPIR and SPIR-V target feature support *- C++ -*-===//
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//===----------------------------------------------------------------------===//
8//
9// This file declares SPIR and SPIR-V TargetInfo objects.
10//
11//===----------------------------------------------------------------------===//
12
13#ifndef LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
14#define LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
15
16#include "Targets.h"
17#include "clang/Basic/AddressSpaces.h"
18#include "clang/Basic/Diagnostic.h"
19#include "clang/Basic/DiagnosticFrontend.h"
20#include "clang/Basic/TargetInfo.h"
21#include "clang/Basic/TargetOptions.h"
22#include "llvm/Support/Compiler.h"
23#include "llvm/Support/VersionTuple.h"
24#include "llvm/TargetParser/Triple.h"
25#include <optional>
26
27namespace clang {
28namespace targets {
29
30// Used by both the SPIR and SPIR-V targets.
31static constexpr LangASMap SPIRDefIsPrivMap = {
32 {LangAS::opencl_global, 1}, {LangAS::opencl_local, 3},
33 {LangAS::opencl_constant, 2}, {LangAS::opencl_generic, 4},
34 {LangAS::opencl_global_device, 5}, {LangAS::opencl_global_host, 6},
35 {LangAS::hlsl_groupshared, 3}, {LangAS::hlsl_constant, 12},
36 {LangAS::hlsl_private, 10}, {LangAS::hlsl_device, 11},
37 {LangAS::hlsl_input, 7}, {LangAS::hlsl_output, 8},
38 {LangAS::hlsl_push_constant, 13},
39};
40
41// Used by both the SPIR and SPIR-V targets.
42static constexpr LangASMap SPIRDefIsGenMap = {
43 {LangAS::Default, 4},
44 {LangAS::opencl_global, 1},
45 {LangAS::opencl_local, 3},
46 {LangAS::opencl_constant, 2},
47 {LangAS::opencl_generic, 4},
48 {LangAS::opencl_global_device, 5},
49 {LangAS::opencl_global_host, 6},
50 // cuda_* address space mapping is intended for HIPSPV (HIP to SPIR-V
51 // translation). This mapping is enabled when the language mode is HIP.
52 {LangAS::cuda_device, 1},
53 // cuda_constant pointer can be casted to default/"flat" pointer, but in
54 // SPIR-V casts between constant and generic pointers are not allowed. For
55 // this reason cuda_constant is mapped to SPIR-V CrossWorkgroup.
56 {LangAS::cuda_constant, 1},
57 {LangAS::cuda_shared, 3},
58 {LangAS::sycl_global, 1},
59 {LangAS::sycl_global_device, 5},
60 {LangAS::sycl_global_host, 6},
61 {LangAS::sycl_local, 3},
62 {LangAS::hlsl_groupshared, 3},
63 {LangAS::hlsl_private, 10},
64 {LangAS::hlsl_device, 11},
65 {LangAS::hlsl_input, 7},
66 {LangAS::hlsl_output, 8},
67 {LangAS::hlsl_push_constant, 13},
68};
69
70// Base class for SPIR and SPIR-V target info.
71class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public TargetInfo {
72 std::unique_ptr<TargetInfo> HostTarget;
73
74protected:
75 BaseSPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
76 : TargetInfo(Triple) {
77 assert((Triple.isSPIR() || Triple.isSPIRV()) &&
78 "Invalid architecture for SPIR or SPIR-V.");
79 TLSSupported = false;
80 VLASupported = false;
81 LongWidth = LongAlign = 64;
82 AddrSpaceMap = &SPIRDefIsPrivMap;
83 UseAddrSpaceMapMangling = true;
84 HasFastHalfType = true;
85 HasFloat16 = true;
86 HasBFloat16 = true;
87 HasFullBFloat16 = true;
88 BFloat16Width = BFloat16Align = 16;
89 BFloat16Format = &llvm::APFloat::BFloat();
90 // Define available target features
91 // These must be defined in sorted order!
92 NoAsmVariants = true;
93
94 llvm::Triple HostTriple(Opts.HostTriple);
95 if (!HostTriple.isSPIR() && !HostTriple.isSPIRV() &&
96 HostTriple.getArch() != llvm::Triple::UnknownArch) {
97 HostTarget = AllocateTarget(Triple: llvm::Triple(Opts.HostTriple), Opts);
98
99 // Copy properties from host target.
100 BoolWidth = HostTarget->getBoolWidth();
101 BoolAlign = HostTarget->getBoolAlign();
102 IntWidth = HostTarget->getIntWidth();
103 IntAlign = HostTarget->getIntAlign();
104 HalfWidth = HostTarget->getHalfWidth();
105 HalfAlign = HostTarget->getHalfAlign();
106 FloatWidth = HostTarget->getFloatWidth();
107 FloatAlign = HostTarget->getFloatAlign();
108 DoubleWidth = HostTarget->getDoubleWidth();
109 DoubleAlign = HostTarget->getDoubleAlign();
110 LongWidth = HostTarget->getLongWidth();
111 LongAlign = HostTarget->getLongAlign();
112 LongLongWidth = HostTarget->getLongLongWidth();
113 LongLongAlign = HostTarget->getLongLongAlign();
114 MinGlobalAlign =
115 HostTarget->getMinGlobalAlign(/* TypeSize = */ Size: 0,
116 /* HasNonWeakDef = */ HasNonWeakDef: true);
117 NewAlign = HostTarget->getNewAlign();
118 DefaultAlignForAttributeAligned =
119 HostTarget->getDefaultAlignForAttributeAligned();
120 IntMaxType = HostTarget->getIntMaxType();
121 WCharType = HostTarget->getWCharType();
122 WIntType = HostTarget->getWIntType();
123 Char16Type = HostTarget->getChar16Type();
124 Char32Type = HostTarget->getChar32Type();
125 Int64Type = HostTarget->getInt64Type();
126 SigAtomicType = HostTarget->getSigAtomicType();
127 ProcessIDType = HostTarget->getProcessIDType();
128
129 UseBitFieldTypeAlignment = HostTarget->useBitFieldTypeAlignment();
130 UseZeroLengthBitfieldAlignment =
131 HostTarget->useZeroLengthBitfieldAlignment();
132 UseExplicitBitFieldAlignment = HostTarget->useExplicitBitFieldAlignment();
133 ZeroLengthBitfieldBoundary = HostTarget->getZeroLengthBitfieldBoundary();
134
135 // This is a bit of a lie, but it controls __GCC_ATOMIC_XXX_LOCK_FREE, and
136 // we need those macros to be identical on host and device, because (among
137 // other things) they affect which standard library classes are defined,
138 // and we need all classes to be defined on both the host and device.
139 MaxAtomicInlineWidth = HostTarget->getMaxAtomicInlineWidth();
140 }
141 }
142
143public:
144 llvm::SmallVector<Builtin::InfosShard> getTargetBuiltins() const override {
145 return {};
146 }
147
148 std::string_view getClobbers() const override { return ""; }
149
150 ArrayRef<const char *> getGCCRegNames() const override { return {}; }
151
152 bool validateAsmConstraint(const char *&Name,
153 TargetInfo::ConstraintInfo &info) const override {
154 return true;
155 }
156
157 ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {
158 return {};
159 }
160
161 BuiltinVaListKind getBuiltinVaListKind() const override {
162 return TargetInfo::VoidPtrBuiltinVaList;
163 }
164
165 std::optional<unsigned>
166 getDWARFAddressSpace(unsigned AddressSpace) const override {
167 return AddressSpace;
168 }
169
170 CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {
171 return (CC == CC_SpirFunction || CC == CC_DeviceKernel) ? CCCR_OK
172 : CCCR_Warning;
173 }
174
175 CallingConv getDefaultCallingConv() const override {
176 return CC_SpirFunction;
177 }
178
179 void setAddressSpaceMap(bool DefaultIsGeneric) {
180 AddrSpaceMap = DefaultIsGeneric ? &SPIRDefIsGenMap : &SPIRDefIsPrivMap;
181 }
182
183 void adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
184 const TargetInfo *Aux) override {
185 TargetInfo::adjust(Diags, Opts, Aux);
186 // FIXME: SYCL specification considers unannotated pointers and references
187 // to be pointing to the generic address space. See section 5.9.3 of
188 // SYCL 2020 specification.
189 // Currently, there is no way of representing SYCL's and HIP/CUDA's default
190 // address space language semantic along with the semantics of embedded C's
191 // default address space in the same address space map. Hence the map needs
192 // to be reset to allow mapping to the desired value of 'Default' entry for
193 // SYCL and HIP/CUDA.
194 setAddressSpaceMap(
195 /*DefaultIsGeneric=*/Opts.SYCLIsDevice ||
196 // The address mapping from HIP/CUDA language for device code is only
197 // defined for SPIR-V, and all Intel SPIR-V code should have the default
198 // AS as generic.
199 (getTriple().isSPIRV() &&
200 (Opts.CUDAIsDevice ||
201 getTriple().getVendor() == llvm::Triple::Intel)));
202 }
203
204 void setSupportedOpenCLOpts() override {
205 // Assume all OpenCL extensions and optional core features are supported
206 // for SPIR and SPIR-V since they are generic targets.
207 supportAllOpenCLOpts();
208 }
209
210 bool hasBitIntType() const override { return true; }
211
212 bool hasInt128Type() const override { return false; }
213};
214
215class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public BaseSPIRTargetInfo {
216public:
217 SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
218 : BaseSPIRTargetInfo(Triple, Opts) {
219 assert(Triple.isSPIR() && "Invalid architecture for SPIR.");
220 assert(getTriple().getOS() == llvm::Triple::UnknownOS &&
221 "SPIR target must use unknown OS");
222 assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
223 "SPIR target must use unknown environment type");
224 }
225
226 void getTargetDefines(const LangOptions &Opts,
227 MacroBuilder &Builder) const override;
228
229 bool hasFeature(StringRef Feature) const override {
230 return Feature == "spir";
231 }
232
233 bool checkArithmeticFenceSupported() const override { return true; }
234};
235
236class LLVM_LIBRARY_VISIBILITY SPIR32TargetInfo : public SPIRTargetInfo {
237public:
238 SPIR32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
239 : SPIRTargetInfo(Triple, Opts) {
240 assert(Triple.getArch() == llvm::Triple::spir &&
241 "Invalid architecture for 32-bit SPIR.");
242 PointerWidth = PointerAlign = 32;
243 SizeType = TargetInfo::UnsignedInt;
244 PtrDiffType = IntPtrType = TargetInfo::SignedInt;
245 // SPIR32 has support for atomic ops if atomic extension is enabled.
246 // Take the maximum because it's possible the Host supports wider types.
247 MaxAtomicInlineWidth = std::max<unsigned char>(a: MaxAtomicInlineWidth, b: 64);
248 resetDataLayout(DL: "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-"
249 "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
250 }
251
252 void getTargetDefines(const LangOptions &Opts,
253 MacroBuilder &Builder) const override;
254};
255
256class LLVM_LIBRARY_VISIBILITY SPIR64TargetInfo : public SPIRTargetInfo {
257public:
258 SPIR64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
259 : SPIRTargetInfo(Triple, Opts) {
260 assert(Triple.getArch() == llvm::Triple::spir64 &&
261 "Invalid architecture for 64-bit SPIR.");
262 PointerWidth = PointerAlign = 64;
263 SizeType = TargetInfo::UnsignedLong;
264 PtrDiffType = IntPtrType = TargetInfo::SignedLong;
265 // SPIR64 has support for atomic ops if atomic extension is enabled.
266 // Take the maximum because it's possible the Host supports wider types.
267 MaxAtomicInlineWidth = std::max<unsigned char>(a: MaxAtomicInlineWidth, b: 64);
268 resetDataLayout(DL: "e-i64:64-v16:16-v24:32-v32:32-v48:64-"
269 "v96:128-v192:256-v256:256-v512:512-v1024:1024-G1");
270 }
271
272 void getTargetDefines(const LangOptions &Opts,
273 MacroBuilder &Builder) const override;
274};
275
276class LLVM_LIBRARY_VISIBILITY BaseSPIRVTargetInfo : public BaseSPIRTargetInfo {
277public:
278 BaseSPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
279 : BaseSPIRTargetInfo(Triple, Opts) {
280 assert(Triple.isSPIRV() && "Invalid architecture for SPIR-V.");
281 }
282
283 llvm::SmallVector<Builtin::InfosShard> getTargetBuiltins() const override;
284
285 bool hasFeature(StringRef Feature) const override {
286 return Feature == "spirv";
287 }
288
289 virtual bool isAddressSpaceSupersetOf(LangAS A, LangAS B) const override {
290 // The generic space AS(4) is a superset of all the other address
291 // spaces used by the backend target except constant address space.
292 return A == B || ((A == LangAS::Default ||
293 (isTargetAddressSpace(AS: A) &&
294 toTargetAddressSpace(AS: A) == /*Generic=*/4)) &&
295 isTargetAddressSpace(AS: B) &&
296 (toTargetAddressSpace(AS: B) <= /*Generic=*/4 &&
297 toTargetAddressSpace(AS: B) != /*Constant=*/2));
298 }
299
300 void getTargetDefines(const LangOptions &Opts,
301 MacroBuilder &Builder) const override;
302};
303
304class LLVM_LIBRARY_VISIBILITY SPIRVTargetInfo : public BaseSPIRVTargetInfo {
305public:
306 SPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
307 : BaseSPIRVTargetInfo(Triple, Opts) {
308 assert(Triple.getArch() == llvm::Triple::spirv &&
309 "Invalid architecture for Logical SPIR-V.");
310 PointerWidth = PointerAlign = 64;
311
312 // SPIR-V IDs are represented with a single 32-bit word.
313 SizeType = TargetInfo::UnsignedInt;
314 VectorsAreElementAligned = true;
315 resetDataLayout();
316 }
317
318 // SPIR-V targeting requires a fully specified Vulkan environment.
319 // SPIR-V requires the enviornment to be in a valid shader stage as well.
320 // Validate here before CreateTargetInfo() to emit a proper diagnostic.
321 bool validateTarget(DiagnosticsEngine &Diags) const override {
322 if (getTriple().getOS() != llvm::Triple::Vulkan ||
323 getTriple().getVulkanVersion() == llvm::VersionTuple(0)) {
324 Diags.Report(DiagID: diag::err_target_spirv_requires_vulkan);
325 return false;
326 }
327 if (getTriple().getEnvironment() != llvm::Triple::UnknownEnvironment &&
328 (getTriple().getEnvironment() < llvm::Triple::Pixel ||
329 getTriple().getEnvironment() > llvm::Triple::Amplification)) {
330 Diags.Report(DiagID: diag::err_target_spirv_invalid_shader_stage);
331 return false;
332 }
333 return true;
334 }
335
336 void getTargetDefines(const LangOptions &Opts,
337 MacroBuilder &Builder) const override;
338};
339
340class LLVM_LIBRARY_VISIBILITY SPIRV32TargetInfo : public BaseSPIRVTargetInfo {
341public:
342 SPIRV32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
343 : BaseSPIRVTargetInfo(Triple, Opts) {
344 assert(Triple.getArch() == llvm::Triple::spirv32 &&
345 "Invalid architecture for 32-bit SPIR-V.");
346 assert((getTriple().getOS() == llvm::Triple::UnknownOS ||
347 getTriple().getOS() == llvm::Triple::ChipStar ||
348 getTriple().getOS() == llvm::Triple::Vulkan) &&
349 "32-bit SPIR-V target must use unknown, chipstar, or vulkan OS");
350 assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
351 "32-bit SPIR-V target must use unknown environment type");
352 PointerWidth = PointerAlign = 32;
353 SizeType = TargetInfo::UnsignedInt;
354 PtrDiffType = IntPtrType = TargetInfo::SignedInt;
355 // SPIR-V has core support for atomic ops, and Int32 is always available;
356 // we take the maximum because it's possible the Host supports wider types.
357 MaxAtomicInlineWidth = std::max<unsigned char>(a: MaxAtomicInlineWidth, b: 64);
358 resetDataLayout();
359 }
360
361 void getTargetDefines(const LangOptions &Opts,
362 MacroBuilder &Builder) const override;
363};
364
365class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public BaseSPIRVTargetInfo {
366public:
367 SPIRV64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
368 : BaseSPIRVTargetInfo(Triple, Opts) {
369 assert(Triple.getArch() == llvm::Triple::spirv64 &&
370 "Invalid architecture for 64-bit SPIR-V.");
371 assert((getTriple().getOS() == llvm::Triple::UnknownOS ||
372 getTriple().getOS() == llvm::Triple::ChipStar ||
373 getTriple().getOS() == llvm::Triple::Vulkan) &&
374 "64-bit SPIR-V target must use unknown, chipstar, or vulkan OS");
375 assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
376 "64-bit SPIR-V target must use unknown environment type");
377 PointerWidth = PointerAlign = 64;
378 SizeType = TargetInfo::UnsignedLong;
379 PtrDiffType = IntPtrType = TargetInfo::SignedLong;
380 // SPIR-V has core support for atomic ops, and Int64 is always available;
381 // we take the maximum because it's possible the Host supports wider types.
382 MaxAtomicInlineWidth = std::max<unsigned char>(a: MaxAtomicInlineWidth, b: 64);
383 resetDataLayout();
384 }
385
386 void getTargetDefines(const LangOptions &Opts,
387 MacroBuilder &Builder) const override;
388
389 const llvm::omp::GV &getGridValue() const override {
390 return llvm::omp::SPIRVGridValues;
391 }
392
393 std::optional<LangAS> getConstantAddressSpace() const override {
394 return ConstantAS;
395 }
396 void adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
397 const TargetInfo *Aux) override {
398 BaseSPIRVTargetInfo::adjust(Diags, Opts, Aux);
399 // opencl_constant will map to UniformConstant in SPIR-V
400 if (Opts.OpenCL)
401 ConstantAS = LangAS::opencl_constant;
402 }
403
404private:
405 // opencl_global will map to CrossWorkgroup in SPIR-V
406 LangAS ConstantAS = LangAS::opencl_global;
407};
408
409class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
410 : public BaseSPIRVTargetInfo {
411public:
412 SPIRV64AMDGCNTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
413 : BaseSPIRVTargetInfo(Triple, Opts) {
414 assert(Triple.getArch() == llvm::Triple::spirv64 &&
415 "Invalid architecture for 64-bit AMDGCN SPIR-V.");
416 assert(Triple.getVendor() == llvm::Triple::VendorType::AMD &&
417 "64-bit AMDGCN SPIR-V target must use AMD vendor");
418 assert(getTriple().getOS() == llvm::Triple::OSType::AMDHSA &&
419 "64-bit AMDGCN SPIR-V target must use AMDHSA OS");
420 assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
421 "64-bit SPIR-V target must use unknown environment type");
422 PointerWidth = PointerAlign = 64;
423 SizeType = TargetInfo::UnsignedLong;
424 PtrDiffType = IntPtrType = TargetInfo::SignedLong;
425 AddrSpaceMap = &SPIRDefIsGenMap;
426
427 resetDataLayout();
428
429 HasFastHalfType = true;
430 HasFloat16 = true;
431 HalfArgsAndReturns = true;
432
433 MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
434 }
435
436 ArrayRef<const char *> getGCCRegNames() const override;
437
438 BuiltinVaListKind getBuiltinVaListKind() const override {
439 return TargetInfo::CharPtrBuiltinVaList;
440 }
441
442 bool initFeatureMap(llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags,
443 StringRef,
444 const std::vector<std::string> &) const override;
445
446 bool validateAsmConstraint(const char *&Name,
447 TargetInfo::ConstraintInfo &Info) const override;
448
449 std::string convertConstraint(const char *&Constraint) const override;
450
451 llvm::SmallVector<Builtin::InfosShard> getTargetBuiltins() const override;
452
453 void getTargetDefines(const LangOptions &Opts,
454 MacroBuilder &Builder) const override;
455
456 void setAuxTarget(const TargetInfo *Aux) override;
457
458 void adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
459 const TargetInfo *Aux) override {
460 TargetInfo::adjust(Diags, Opts, Aux);
461
462 AtomicOpts = AtomicOptions(Opts);
463 }
464
465 bool hasInt128Type() const override { return TargetInfo::hasInt128Type(); }
466
467 // This is only needed for validating arguments passed to
468 // __builtin_amdgcn_processor_is
469 bool isValidCPUName(StringRef Name) const override;
470 void fillValidCPUList(SmallVectorImpl<StringRef> &Values) const override;
471};
472
473class LLVM_LIBRARY_VISIBILITY SPIRV64IntelTargetInfo final
474 : public SPIRV64TargetInfo {
475public:
476 SPIRV64IntelTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
477 : SPIRV64TargetInfo(Triple, Opts) {
478 assert(Triple.getVendor() == llvm::Triple::VendorType::Intel &&
479 "64-bit Intel SPIR-V target must use Intel vendor");
480 resetDataLayout();
481 }
482};
483} // namespace targets
484} // namespace clang
485#endif // LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
486