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