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