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