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