1//===- SPIRVBuiltins.cpp - SPIR-V Built-in Functions ------------*- 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 implements lowering builtin function calls and types using their
10// demangled names and TableGen records.
11//
12//===----------------------------------------------------------------------===//
13
14#include "SPIRVBuiltins.h"
15#include "SPIRV.h"
16#include "SPIRVSubtarget.h"
17#include "SPIRVUtils.h"
18#include "llvm/ADT/StringExtras.h"
19#include "llvm/ADT/StringTable.h"
20#include "llvm/Analysis/ValueTracking.h"
21#include "llvm/IR/IntrinsicsSPIRV.h"
22#include <regex>
23#include <string>
24#include <tuple>
25
26#define DEBUG_TYPE "spirv-builtins"
27
28namespace llvm {
29namespace SPIRV {
30#define GET_BuiltinGroup_DECL
31#include "SPIRVGenTables.inc"
32
33struct DemangledBuiltin {
34 StringTable::Offset Name;
35 InstructionSet::InstructionSet Set;
36 BuiltinGroup Group;
37 uint8_t MinNumArgs;
38 uint8_t MaxNumArgs;
39
40 StringRef name() const;
41};
42
43#define GET_DemangledBuiltins_DECL
44#define GET_DemangledBuiltins_IMPL
45
46struct IncomingCall {
47 const std::string BuiltinName;
48 const DemangledBuiltin *Builtin;
49
50 const Register ReturnRegister;
51 const SPIRVTypeInst ReturnType;
52 const SmallVectorImpl<Register> &Arguments;
53
54 IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin,
55 const Register ReturnRegister, SPIRVTypeInst ReturnType,
56 const SmallVectorImpl<Register> &Arguments)
57 : BuiltinName(std::move(BuiltinName)), Builtin(Builtin),
58 ReturnRegister(ReturnRegister), ReturnType(ReturnType),
59 Arguments(Arguments) {}
60
61 bool isSpirvOp() const { return BuiltinName.rfind(s: "__spirv_", pos: 0) == 0; }
62};
63
64struct NativeBuiltin {
65 StringTable::Offset Name;
66 InstructionSet::InstructionSet Set;
67 uint32_t Opcode;
68};
69
70#define GET_NativeBuiltins_DECL
71#define GET_NativeBuiltins_IMPL
72
73struct GroupBuiltin {
74 StringTable::Offset Name;
75 uint32_t Opcode;
76 uint32_t GroupOperation;
77 bool IsElect;
78 bool IsAllOrAny;
79 bool IsAllEqual;
80 bool IsBallot;
81 bool IsInverseBallot;
82 bool IsBallotBitExtract;
83 bool IsBallotFindBit;
84 bool IsLogical;
85 bool NoGroupOperation;
86 bool HasBoolArg;
87};
88
89#define GET_GroupBuiltins_DECL
90#define GET_GroupBuiltins_IMPL
91
92struct IntelSubgroupsBuiltin {
93 StringTable::Offset Name;
94 uint32_t Opcode;
95 bool IsBlock;
96 bool IsWrite;
97 bool IsMedia;
98};
99
100#define GET_IntelSubgroupsBuiltins_DECL
101#define GET_IntelSubgroupsBuiltins_IMPL
102
103struct AtomicFloatingBuiltin {
104 StringTable::Offset Name;
105 uint32_t Opcode;
106};
107
108#define GET_AtomicFloatingBuiltins_DECL
109#define GET_AtomicFloatingBuiltins_IMPL
110struct GroupUniformBuiltin {
111 StringTable::Offset Name;
112 uint32_t Opcode;
113 bool IsLogical;
114};
115
116#define GET_GroupUniformBuiltins_DECL
117#define GET_GroupUniformBuiltins_IMPL
118
119struct GetBuiltin {
120 StringTable::Offset Name;
121 InstructionSet::InstructionSet Set;
122 BuiltIn::BuiltIn Value;
123};
124
125using namespace BuiltIn;
126#define GET_GetBuiltins_DECL
127#define GET_GetBuiltins_IMPL
128
129struct ImageQueryBuiltin {
130 StringTable::Offset Name;
131 InstructionSet::InstructionSet Set;
132 uint32_t Component;
133};
134
135#define GET_ImageQueryBuiltins_DECL
136#define GET_ImageQueryBuiltins_IMPL
137
138struct IntegerDotProductBuiltin {
139 StringTable::Offset Name;
140 uint32_t Opcode;
141 bool IsSwapReq;
142};
143
144#define GET_IntegerDotProductBuiltins_DECL
145#define GET_IntegerDotProductBuiltins_IMPL
146
147struct ConvertBuiltin {
148 StringTable::Offset Name;
149 InstructionSet::InstructionSet Set;
150 bool IsDestinationSigned;
151 bool IsSaturated;
152 bool IsRounded;
153 bool IsBfloat16;
154 bool IsTF32;
155 FPRoundingMode::FPRoundingMode RoundingMode;
156};
157
158struct VectorLoadStoreBuiltin {
159 StringTable::Offset Name;
160 InstructionSet::InstructionSet Set;
161 uint32_t Number;
162 uint32_t ElementCount;
163 bool IsRounded;
164 FPRoundingMode::FPRoundingMode RoundingMode;
165};
166
167using namespace FPRoundingMode;
168#define GET_ConvertBuiltins_DECL
169#define GET_ConvertBuiltins_IMPL
170
171using namespace InstructionSet;
172#define GET_VectorLoadStoreBuiltins_DECL
173#define GET_VectorLoadStoreBuiltins_IMPL
174
175#define GET_CLMemoryScope_DECL
176#define GET_CLSamplerAddressingMode_DECL
177#define GET_CLMemoryFenceFlags_DECL
178#define GET_ExtendedBuiltins_DECL
179#include "SPIRVGenTables.inc"
180
181// Defined here to reference declarations from tablegen.
182StringRef DemangledBuiltin::name() const {
183 return getDemangledBuiltinStr(Offset: Name);
184}
185} // namespace SPIRV
186
187//===----------------------------------------------------------------------===//
188// Misc functions for looking up builtins and veryfying requirements using
189// TableGen records
190//===----------------------------------------------------------------------===//
191
192namespace SPIRV {
193/// Parses the name part of the demangled builtin call.
194std::string lookupBuiltinNameHelper(StringRef DemangledCall,
195 FPDecorationId *DecorationId) {
196 StringRef PassPrefix = "(anonymous namespace)::";
197 StringRef SpvPrefix = "__spv::";
198 std::string BuiltinName = DemangledCall.str();
199
200 // Check if the extracted name contains type information between angle
201 // brackets. If so, the builtin is an instantiated template - needs to have
202 // the information after angle brackets and return type removed.
203 std::size_t Pos = BuiltinName.find(s: ">(");
204 if (Pos != std::string::npos) {
205 BuiltinName = BuiltinName.substr(pos: 0, n: BuiltinName.rfind(c: '<', pos: Pos));
206 } else {
207 Pos = BuiltinName.find(c: '(');
208 if (Pos != std::string::npos)
209 BuiltinName = BuiltinName.substr(pos: 0, n: Pos);
210 }
211 BuiltinName = BuiltinName.substr(pos: BuiltinName.find_last_of(c: ' ') + 1);
212
213 // Itanium Demangler result may have "(anonymous namespace)::" or "__spv::"
214 // prefix.
215 if (BuiltinName.find(svt: PassPrefix) == 0)
216 BuiltinName = BuiltinName.substr(pos: PassPrefix.size());
217 else if (BuiltinName.find(svt: SpvPrefix) == 0)
218 BuiltinName = BuiltinName.substr(pos: SpvPrefix.size());
219
220 // Account for possible "__spirv_ocl_" prefix in SPIR-V friendly LLVM IR
221 if (BuiltinName.rfind(s: "__spirv_ocl_", pos: 0) == 0)
222 BuiltinName = BuiltinName.substr(pos: 12);
223
224 // Check if the extracted name begins with:
225 // - "__spirv_ImageSampleExplicitLod"
226 // - "__spirv_ImageRead"
227 // - "__spirv_ImageWrite"
228 // - "__spirv_ImageQuerySizeLod"
229 // - "__spirv_UDotKHR"
230 // - "__spirv_SDotKHR"
231 // - "__spirv_SUDotKHR"
232 // - "__spirv_SDotAccSatKHR"
233 // - "__spirv_UDotAccSatKHR"
234 // - "__spirv_SUDotAccSatKHR"
235 // - "__spirv_ReadClockKHR"
236 // - "__spirv_SubgroupBlockReadINTEL"
237 // - "__spirv_SubgroupImageBlockReadINTEL"
238 // - "__spirv_SubgroupImageMediaBlockReadINTEL"
239 // - "__spirv_SubgroupImageMediaBlockWriteINTEL"
240 // - "__spirv_Convert"
241 // - "__spirv_Round"
242 // - "__spirv_UConvert"
243 // - "__spirv_SConvert"
244 // - "__spirv_FConvert"
245 // - "__spirv_SatConvert"
246 // and maybe contains return type information at the end "_R<type>".
247 // If so, extract the plain builtin name without the type information.
248 static const std::regex SpvWithR(
249 "(__spirv_(ImageSampleExplicitLod|ImageRead|ImageWrite|ImageQuerySizeLod|"
250 "UDotKHR|"
251 "SDotKHR|SUDotKHR|SDotAccSatKHR|UDotAccSatKHR|SUDotAccSatKHR|"
252 "ReadClockKHR|SubgroupBlockReadINTEL|SubgroupImageBlockReadINTEL|"
253 "SubgroupImageMediaBlockReadINTEL|SubgroupImageMediaBlockWriteINTEL|"
254 "Convert|Round|"
255 "UConvert|SConvert|FConvert|SatConvert)[^_]*)(_R[^_]*_?(\\w+)?.*)?");
256 std::smatch Match;
257 if (std::regex_match(s: BuiltinName, m&: Match, re: SpvWithR) && Match.size() > 1) {
258 std::ssub_match SubMatch;
259 if (DecorationId && Match.size() > 3) {
260 SubMatch = Match[4];
261 *DecorationId = demangledPostfixToDecorationId(S: SubMatch.str());
262 }
263 SubMatch = Match[1];
264 BuiltinName = SubMatch.str();
265 }
266
267 return BuiltinName;
268}
269} // namespace SPIRV
270
271/// Looks up the demangled builtin call in the SPIRVBuiltins.td records using
272/// the provided \p DemangledCall and specified \p Set.
273///
274/// The lookup follows the following algorithm, returning the first successful
275/// match:
276/// 1. Search with the plain demangled name (expecting a 1:1 match).
277/// 2. Search with the prefix before or suffix after the demangled name
278/// signyfying the type of the first argument.
279///
280/// \returns Wrapper around the demangled call and found builtin definition.
281static std::unique_ptr<const SPIRV::IncomingCall>
282lookupBuiltin(StringRef DemangledCall,
283 SPIRV::InstructionSet::InstructionSet Set,
284 Register ReturnRegister, SPIRVTypeInst ReturnType,
285 const SmallVectorImpl<Register> &Arguments) {
286 std::string BuiltinName = SPIRV::lookupBuiltinNameHelper(DemangledCall);
287
288 SmallVector<StringRef, 10> BuiltinArgumentTypes;
289 StringRef BuiltinArgs =
290 DemangledCall.slice(Start: DemangledCall.find(C: '(') + 1, End: DemangledCall.find(C: ')'));
291 BuiltinArgs.split(A&: BuiltinArgumentTypes, Separator: ',', MaxSplit: -1, KeepEmpty: false);
292
293 // Look up the builtin in the defined set. Start with the plain demangled
294 // name, expecting a 1:1 match in the defined builtin set.
295 const SPIRV::DemangledBuiltin *Builtin;
296 if ((Builtin = SPIRV::lookupBuiltin(Name: BuiltinName, Set)))
297 return std::make_unique<SPIRV::IncomingCall>(
298 args&: BuiltinName, args&: Builtin, args&: ReturnRegister, args&: ReturnType, args: Arguments);
299
300 // If the initial look up was unsuccessful and the demangled call takes at
301 // least 1 argument, add a prefix or suffix signifying the type of the first
302 // argument and repeat the search.
303 if (BuiltinArgumentTypes.size() >= 1) {
304 char FirstArgumentType = BuiltinArgumentTypes[0][0];
305 // Prefix to be added to the builtin's name for lookup.
306 // For example, OpenCL "abs" taking an unsigned value has a prefix "u_".
307 std::string Prefix;
308
309 switch (FirstArgumentType) {
310 // Unsigned:
311 case 'u':
312 if (Set == SPIRV::InstructionSet::OpenCL_std)
313 Prefix = "u_";
314 else if (Set == SPIRV::InstructionSet::GLSL_std_450)
315 Prefix = "u";
316 break;
317 // Signed:
318 case 'c':
319 case 's':
320 case 'i':
321 case 'l':
322 if (Set == SPIRV::InstructionSet::OpenCL_std)
323 Prefix = "s_";
324 else if (Set == SPIRV::InstructionSet::GLSL_std_450)
325 Prefix = "s";
326 break;
327 // Floating-point:
328 case 'f':
329 case 'd':
330 case 'h':
331 if (Set == SPIRV::InstructionSet::OpenCL_std ||
332 Set == SPIRV::InstructionSet::GLSL_std_450)
333 Prefix = "f";
334 break;
335 }
336
337 // If argument-type name prefix was added, look up the builtin again.
338 if (!Prefix.empty() &&
339 (Builtin = SPIRV::lookupBuiltin(Name: Prefix + BuiltinName, Set)))
340 return std::make_unique<SPIRV::IncomingCall>(
341 args&: BuiltinName, args&: Builtin, args&: ReturnRegister, args&: ReturnType, args: Arguments);
342
343 // If lookup with a prefix failed, find a suffix to be added to the
344 // builtin's name for lookup. For example, OpenCL "group_reduce_max" taking
345 // an unsigned value has a suffix "u".
346 std::string Suffix;
347
348 switch (FirstArgumentType) {
349 // Unsigned:
350 case 'u':
351 Suffix = "u";
352 break;
353 // Signed:
354 case 'c':
355 case 's':
356 case 'i':
357 case 'l':
358 Suffix = "s";
359 break;
360 // Floating-point:
361 case 'f':
362 case 'd':
363 case 'h':
364 Suffix = "f";
365 break;
366 }
367
368 // If argument-type name suffix was added, look up the builtin again.
369 if (!Suffix.empty() &&
370 (Builtin = SPIRV::lookupBuiltin(Name: BuiltinName + Suffix, Set)))
371 return std::make_unique<SPIRV::IncomingCall>(
372 args&: BuiltinName, args&: Builtin, args&: ReturnRegister, args&: ReturnType, args: Arguments);
373 }
374
375 // No builtin with such name was found in the set.
376 return nullptr;
377}
378
379static MachineInstr *getBlockStructInstr(Register ParamReg,
380 MachineRegisterInfo *MRI) {
381 // We expect ParamReg to be defined by G_ADDRSPACE_CAST with a source from
382 // G_GLOBAL_VALUE or spv_alloca. Returns the source instruction.
383 MachineInstr *MI = MRI->getUniqueVRegDef(Reg: ParamReg);
384 assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
385 MI->getOperand(1).isReg());
386 Register BitcastReg = MI->getOperand(i: 1).getReg();
387 MachineInstr *BitcastMI = MRI->getUniqueVRegDef(Reg: BitcastReg);
388 assert(BitcastMI && "Definition for source reg not found.");
389 if (BitcastMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
390 isSpvIntrinsic(MI: *BitcastMI, IntrinsicID: Intrinsic::spv_alloca))
391 return BitcastMI;
392 llvm_unreachable("getBlockStructInstr: unexpected instruction pattern");
393}
394
395// Return type of the instruction result from spv_assign_type intrinsic.
396// TODO: maybe unify with prelegalizer pass.
397static const Type *getMachineInstrType(MachineInstr *MI) {
398 MachineInstr *NextMI = MI->getNextNode();
399 if (!NextMI)
400 return nullptr;
401 if (isSpvIntrinsic(MI: *NextMI, IntrinsicID: Intrinsic::spv_assign_name))
402 if ((NextMI = NextMI->getNextNode()) == nullptr)
403 return nullptr;
404 Register ValueReg = MI->getOperand(i: 0).getReg();
405 if ((!isSpvIntrinsic(MI: *NextMI, IntrinsicID: Intrinsic::spv_assign_type) &&
406 !isSpvIntrinsic(MI: *NextMI, IntrinsicID: Intrinsic::spv_assign_ptr_type)) ||
407 NextMI->getOperand(i: 1).getReg() != ValueReg)
408 return nullptr;
409 Type *Ty = getMDOperandAsType(N: NextMI->getOperand(i: 2).getMetadata(), I: 0);
410 assert(Ty && "Type is expected");
411 return Ty;
412}
413
414static const Type *getBlockStructType(Register ParamReg,
415 MachineRegisterInfo *MRI) {
416 // In principle, this information should be passed to us from Clang via
417 // an elementtype attribute. However, said attribute requires that
418 // the function call be an intrinsic, which is not. Instead, we rely on being
419 // able to trace this to the declaration of a variable: OpenCL C specification
420 // section 6.12.5 should guarantee that we can do this.
421 MachineInstr *MI = getBlockStructInstr(ParamReg, MRI);
422 if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
423 return MI->getOperand(i: 1).getGlobal()->getValueType();
424 assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
425 "Blocks in OpenCL C must be traceable to allocation site");
426 return getMachineInstrType(MI);
427}
428
429//===----------------------------------------------------------------------===//
430// Helper functions for building misc instructions
431//===----------------------------------------------------------------------===//
432
433/// Helper function building either a resulting scalar or vector bool register
434/// depending on the expected \p ResultType.
435///
436/// \returns Tuple of the resulting register and its type.
437static std::tuple<Register, SPIRVTypeInst>
438buildBoolRegister(MachineIRBuilder &MIRBuilder, SPIRVTypeInst ResultType,
439 SPIRVGlobalRegistry *GR) {
440 LLT Type;
441 SPIRVTypeInst BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder, EmitIR: true);
442
443 if (ResultType->getOpcode() == SPIRV::OpTypeVector) {
444 unsigned VectorElements = GR->getScalarOrVectorComponentCount(Type: ResultType);
445 BoolType = GR->getOrCreateSPIRVVectorType(BaseType: BoolType, NumElements: VectorElements,
446 MIRBuilder, EmitIR: true);
447 const FixedVectorType *LLVMVectorType =
448 cast<FixedVectorType>(Val: GR->getTypeForSPIRVType(Ty: BoolType));
449 Type = LLT::vector(EC: LLVMVectorType->getElementCount(), ScalarSizeInBits: 1);
450 } else {
451 Type = LLT::scalar(SizeInBits: 1);
452 }
453
454 Register ResultRegister =
455 MIRBuilder.getMRI()->createGenericVirtualRegister(Ty: Type);
456 MIRBuilder.getMRI()->setRegClass(Reg: ResultRegister, RC: GR->getRegClass(SpvType: ResultType));
457 GR->assignSPIRVTypeToVReg(Type: BoolType, VReg: ResultRegister, MF: MIRBuilder.getMF());
458 return std::make_tuple(args&: ResultRegister, args&: BoolType);
459}
460
461/// Helper function for building either a vector or scalar select instruction
462/// depending on the expected \p ResultType.
463static bool buildSelectInst(MachineIRBuilder &MIRBuilder,
464 Register ReturnRegister, Register SourceRegister,
465 SPIRVTypeInst ReturnType, SPIRVGlobalRegistry *GR) {
466 Register TrueConst, FalseConst;
467
468 if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {
469 unsigned Bits = GR->getScalarOrVectorBitWidth(Type: ReturnType);
470 uint64_t AllOnes = APInt::getAllOnes(numBits: Bits).getZExtValue();
471 TrueConst =
472 GR->getOrCreateConsIntVector(Val: AllOnes, MIRBuilder, SpvType: ReturnType, EmitIR: true);
473 FalseConst = GR->getOrCreateConsIntVector(Val: 0, MIRBuilder, SpvType: ReturnType, EmitIR: true);
474 } else {
475 TrueConst = GR->buildConstantInt(Val: 1, MIRBuilder, SpvType: ReturnType, EmitIR: true);
476 FalseConst = GR->buildConstantInt(Val: 0, MIRBuilder, SpvType: ReturnType, EmitIR: true);
477 }
478
479 return MIRBuilder.buildSelect(Res: ReturnRegister, Tst: SourceRegister, Op0: TrueConst,
480 Op1: FalseConst);
481}
482
483/// Helper function for building a load instruction loading into the
484/// \p DestinationReg.
485static Register buildLoadInst(SPIRVTypeInst BaseType, Register PtrRegister,
486 MachineIRBuilder &MIRBuilder,
487 SPIRVGlobalRegistry *GR,
488 Register DestinationReg = Register(0)) {
489 if (!DestinationReg.isValid())
490 DestinationReg = createVirtualRegister(SpvType: BaseType, GR, MIRBuilder);
491 // TODO: consider using correct address space and alignment (p0 is canonical
492 // type for selection though).
493 MachinePointerInfo PtrInfo = MachinePointerInfo();
494 MIRBuilder.buildLoad(Res: DestinationReg, Addr: PtrRegister, PtrInfo, Alignment: Align());
495 return DestinationReg;
496}
497
498/// Helper function for building a load instruction for loading a builtin global
499/// variable of \p BuiltinValue value.
500static Register buildBuiltinVariableLoad(
501 MachineIRBuilder &MIRBuilder, SPIRVTypeInst VariableType,
502 SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, LLT LLType,
503 Register Reg = Register(0), bool isConst = true,
504 const std::optional<SPIRV::LinkageType::LinkageType> &LinkageTy = {
505 SPIRV::LinkageType::Import}) {
506 Register NewRegister =
507 MIRBuilder.getMRI()->createVirtualRegister(RegClass: &SPIRV::pIDRegClass);
508 MIRBuilder.getMRI()->setType(
509 VReg: NewRegister,
510 Ty: LLT::pointer(AddressSpace: storageClassToAddressSpace(SC: SPIRV::StorageClass::Function),
511 SizeInBits: GR->getPointerSize()));
512 SPIRVTypeInst PtrType = GR->getOrCreateSPIRVPointerType(
513 BaseType: VariableType, MIRBuilder, SC: SPIRV::StorageClass::Input);
514 GR->assignSPIRVTypeToVReg(Type: PtrType, VReg: NewRegister, MF: MIRBuilder.getMF());
515
516 // Set up the global OpVariable with the necessary builtin decorations.
517 Register Variable = GR->buildGlobalVariable(
518 Reg: NewRegister, BaseType: PtrType, Name: getLinkStringForBuiltIn(BuiltInValue: BuiltinValue), GV: nullptr,
519 Storage: SPIRV::StorageClass::Input, Init: nullptr, /* isConst= */ IsConst: isConst, LinkageType: LinkageTy,
520 MIRBuilder, IsInstSelector: false);
521
522 // Load the value from the global variable.
523 Register LoadedRegister =
524 buildLoadInst(BaseType: VariableType, PtrRegister: Variable, MIRBuilder, GR, DestinationReg: Reg);
525 MIRBuilder.getMRI()->setType(VReg: LoadedRegister, Ty: LLType);
526 return LoadedRegister;
527}
528
529/// Helper external function for assigning a SPIRV type to a register, ensuring
530/// the register class and type are set in MRI. Defined in
531/// SPIRVPreLegalizer.cpp.
532extern void updateRegType(Register Reg, Type *Ty, SPIRVTypeInst SpirvTy,
533 SPIRVGlobalRegistry *GR, MachineIRBuilder &MIB,
534 MachineRegisterInfo &MRI);
535
536// TODO: Move to TableGen.
537static SPIRV::MemorySemantics::MemorySemantics
538getSPIRVMemSemantics(std::memory_order MemOrder) {
539 switch (MemOrder) {
540 case std::memory_order_relaxed:
541 return SPIRV::MemorySemantics::None;
542 case std::memory_order_acquire:
543 return SPIRV::MemorySemantics::Acquire;
544 case std::memory_order_release:
545 return SPIRV::MemorySemantics::Release;
546 case std::memory_order_acq_rel:
547 return SPIRV::MemorySemantics::AcquireRelease;
548 case std::memory_order_seq_cst:
549 return SPIRV::MemorySemantics::SequentiallyConsistent;
550 default:
551 report_fatal_error(reason: "Unknown CL memory scope");
552 }
553}
554
555static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope) {
556 switch (ClScope) {
557 case SPIRV::CLMemoryScope::memory_scope_work_item:
558 return SPIRV::Scope::Invocation;
559 case SPIRV::CLMemoryScope::memory_scope_work_group:
560 return SPIRV::Scope::Workgroup;
561 case SPIRV::CLMemoryScope::memory_scope_device:
562 return SPIRV::Scope::Device;
563 case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:
564 return SPIRV::Scope::CrossDevice;
565 case SPIRV::CLMemoryScope::memory_scope_sub_group:
566 return SPIRV::Scope::Subgroup;
567 }
568 report_fatal_error(reason: "Unknown CL memory scope");
569}
570
571static Register buildConstantIntReg32(uint64_t Val,
572 MachineIRBuilder &MIRBuilder,
573 SPIRVGlobalRegistry *GR) {
574 return GR->buildConstantInt(
575 Val, MIRBuilder, SpvType: GR->getOrCreateSPIRVIntegerType(BitWidth: 32, MIRBuilder), EmitIR: true);
576}
577
578static Register buildScopeReg(Register CLScopeRegister,
579 SPIRV::Scope::Scope Scope,
580 MachineIRBuilder &MIRBuilder,
581 SPIRVGlobalRegistry *GR,
582 MachineRegisterInfo *MRI) {
583 if (CLScopeRegister.isValid()) {
584 auto CLScope =
585 static_cast<SPIRV::CLMemoryScope>(getIConstVal(ConstReg: CLScopeRegister, MRI));
586 Scope = getSPIRVScope(ClScope: CLScope);
587
588 if (CLScope == static_cast<unsigned>(Scope)) {
589 MRI->setRegClass(Reg: CLScopeRegister, RC: &SPIRV::iIDRegClass);
590 return CLScopeRegister;
591 }
592 }
593 return buildConstantIntReg32(Val: Scope, MIRBuilder, GR);
594}
595
596static void setRegClassIfNull(Register Reg, MachineRegisterInfo *MRI,
597 SPIRVGlobalRegistry *GR) {
598 if (MRI->getRegClassOrNull(Reg))
599 return;
600 SPIRVTypeInst SpvType = GR->getSPIRVTypeForVReg(VReg: Reg);
601 MRI->setRegClass(Reg,
602 RC: SpvType ? GR->getRegClass(SpvType) : &SPIRV::iIDRegClass);
603}
604
605static Register buildMemSemanticsReg(Register SemanticsRegister,
606 Register PtrRegister, unsigned &Semantics,
607 MachineIRBuilder &MIRBuilder,
608 SPIRVGlobalRegistry *GR) {
609 if (SemanticsRegister.isValid()) {
610 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
611 std::memory_order Order =
612 static_cast<std::memory_order>(getIConstVal(ConstReg: SemanticsRegister, MRI));
613 Semantics =
614 getSPIRVMemSemantics(MemOrder: Order) |
615 getMemSemanticsForStorageClass(SC: GR->getPointerStorageClass(VReg: PtrRegister));
616 if (static_cast<unsigned>(Order) == Semantics) {
617 MRI->setRegClass(Reg: SemanticsRegister, RC: &SPIRV::iIDRegClass);
618 return SemanticsRegister;
619 }
620 }
621 return buildConstantIntReg32(Val: Semantics, MIRBuilder, GR);
622}
623
624static bool buildOpFromWrapper(MachineIRBuilder &MIRBuilder, unsigned Opcode,
625 const SPIRV::IncomingCall *Call,
626 Register TypeReg,
627 ArrayRef<uint32_t> ImmArgs = {}) {
628 auto MIB = MIRBuilder.buildInstr(Opcode);
629 if (TypeReg.isValid())
630 MIB.addDef(RegNo: Call->ReturnRegister).addUse(RegNo: TypeReg);
631 unsigned Sz = Call->Arguments.size() - ImmArgs.size();
632 for (unsigned i = 0; i < Sz; ++i)
633 MIB.addUse(RegNo: Call->Arguments[i]);
634 for (uint32_t ImmArg : ImmArgs)
635 MIB.addImm(Val: ImmArg);
636 return true;
637}
638
639/// Helper function for translating atomic init to OpStore.
640static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call,
641 MachineIRBuilder &MIRBuilder) {
642 if (Call->isSpirvOp())
643 return buildOpFromWrapper(MIRBuilder, Opcode: SPIRV::OpStore, Call, TypeReg: Register(0));
644
645 assert(Call->Arguments.size() == 2 &&
646 "Need 2 arguments for atomic init translation");
647 MIRBuilder.buildInstr(Opcode: SPIRV::OpStore)
648 .addUse(RegNo: Call->Arguments[0])
649 .addUse(RegNo: Call->Arguments[1]);
650 return true;
651}
652
653/// Helper function for building an atomic load instruction.
654static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call,
655 MachineIRBuilder &MIRBuilder,
656 SPIRVGlobalRegistry *GR) {
657 Register TypeReg = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
658 if (Call->isSpirvOp())
659 return buildOpFromWrapper(MIRBuilder, Opcode: SPIRV::OpAtomicLoad, Call, TypeReg);
660
661 Register PtrRegister = Call->Arguments[0];
662 // TODO: if true insert call to __translate_ocl_memory_sccope before
663 // OpAtomicLoad and the function implementation. We can use Translator's
664 // output for transcoding/atomic_explicit_arguments.cl as an example.
665 Register ScopeRegister =
666 Call->Arguments.size() > 1
667 ? Call->Arguments[1]
668 : buildConstantIntReg32(Val: SPIRV::Scope::Device, MIRBuilder, GR);
669 Register MemSemanticsReg;
670 if (Call->Arguments.size() > 2) {
671 // TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
672 MemSemanticsReg = Call->Arguments[2];
673 } else {
674 int Semantics =
675 SPIRV::MemorySemantics::SequentiallyConsistent |
676 getMemSemanticsForStorageClass(SC: GR->getPointerStorageClass(VReg: PtrRegister));
677 MemSemanticsReg = buildConstantIntReg32(Val: Semantics, MIRBuilder, GR);
678 }
679
680 MIRBuilder.buildInstr(Opcode: SPIRV::OpAtomicLoad)
681 .addDef(RegNo: Call->ReturnRegister)
682 .addUse(RegNo: TypeReg)
683 .addUse(RegNo: PtrRegister)
684 .addUse(RegNo: ScopeRegister)
685 .addUse(RegNo: MemSemanticsReg);
686 return true;
687}
688
689/// Helper function for building an atomic store instruction.
690static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call,
691 MachineIRBuilder &MIRBuilder,
692 SPIRVGlobalRegistry *GR) {
693 if (Call->isSpirvOp())
694 return buildOpFromWrapper(MIRBuilder, Opcode: SPIRV::OpAtomicStore, Call,
695 TypeReg: Register(0));
696
697 Register ScopeRegister =
698 buildConstantIntReg32(Val: SPIRV::Scope::Device, MIRBuilder, GR);
699 Register PtrRegister = Call->Arguments[0];
700 int Semantics =
701 SPIRV::MemorySemantics::SequentiallyConsistent |
702 getMemSemanticsForStorageClass(SC: GR->getPointerStorageClass(VReg: PtrRegister));
703 Register MemSemanticsReg = buildConstantIntReg32(Val: Semantics, MIRBuilder, GR);
704 MIRBuilder.buildInstr(Opcode: SPIRV::OpAtomicStore)
705 .addUse(RegNo: PtrRegister)
706 .addUse(RegNo: ScopeRegister)
707 .addUse(RegNo: MemSemanticsReg)
708 .addUse(RegNo: Call->Arguments[1]);
709 return true;
710}
711
712/// Helper function for building an atomic compare-exchange instruction.
713static bool buildAtomicCompareExchangeInst(
714 const SPIRV::IncomingCall *Call, const SPIRV::DemangledBuiltin *Builtin,
715 unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
716 if (Call->isSpirvOp())
717 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
718 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
719
720 bool IsCmpxchg = Call->Builtin->name().contains(Other: "cmpxchg");
721 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
722
723 Register ObjectPtr = Call->Arguments[0]; // Pointer (volatile A *object.)
724 Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
725 Register Desired = Call->Arguments[2]; // Value (C Desired).
726 SPIRVTypeInst SpvDesiredTy = GR->getSPIRVTypeForVReg(VReg: Desired);
727 LLT DesiredLLT = MRI->getType(Reg: Desired);
728
729 assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
730 SPIRV::OpTypePointer);
731 unsigned ExpectedType = GR->getSPIRVTypeForVReg(VReg: ExpectedArg)->getOpcode();
732 (void)ExpectedType;
733 assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
734 : ExpectedType == SPIRV::OpTypePointer);
735 assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
736
737 SPIRVTypeInst SpvObjectPtrTy = GR->getSPIRVTypeForVReg(VReg: ObjectPtr);
738 assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
739 auto StorageClass = static_cast<SPIRV::StorageClass::StorageClass>(
740 SpvObjectPtrTy->getOperand(i: 1).getImm());
741 auto MemSemStorage = getMemSemanticsForStorageClass(SC: StorageClass);
742
743 Register MemSemEqualReg;
744 Register MemSemUnequalReg;
745 uint64_t MemSemEqual =
746 IsCmpxchg
747 ? SPIRV::MemorySemantics::None
748 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
749 uint64_t MemSemUnequal =
750 IsCmpxchg
751 ? SPIRV::MemorySemantics::None
752 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
753 if (Call->Arguments.size() >= 4) {
754 assert(Call->Arguments.size() >= 5 &&
755 "Need 5+ args for explicit atomic cmpxchg");
756 auto MemOrdEq =
757 static_cast<std::memory_order>(getIConstVal(ConstReg: Call->Arguments[3], MRI));
758 auto MemOrdNeq =
759 static_cast<std::memory_order>(getIConstVal(ConstReg: Call->Arguments[4], MRI));
760 MemSemEqual = getSPIRVMemSemantics(MemOrder: MemOrdEq) | MemSemStorage;
761 MemSemUnequal = getSPIRVMemSemantics(MemOrder: MemOrdNeq) | MemSemStorage;
762 if (static_cast<unsigned>(MemOrdEq) == MemSemEqual)
763 MemSemEqualReg = Call->Arguments[3];
764 if (static_cast<unsigned>(MemOrdNeq) == MemSemUnequal)
765 MemSemUnequalReg = Call->Arguments[4];
766 }
767 if (!MemSemEqualReg.isValid())
768 MemSemEqualReg = buildConstantIntReg32(Val: MemSemEqual, MIRBuilder, GR);
769 if (!MemSemUnequalReg.isValid())
770 MemSemUnequalReg = buildConstantIntReg32(Val: MemSemUnequal, MIRBuilder, GR);
771
772 Register ScopeReg;
773 auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
774 if (Call->Arguments.size() >= 6) {
775 assert(Call->Arguments.size() == 6 &&
776 "Extra args for explicit atomic cmpxchg");
777 auto ClScope = static_cast<SPIRV::CLMemoryScope>(
778 getIConstVal(ConstReg: Call->Arguments[5], MRI));
779 Scope = getSPIRVScope(ClScope);
780 if (ClScope == static_cast<unsigned>(Scope))
781 ScopeReg = Call->Arguments[5];
782 }
783 if (!ScopeReg.isValid())
784 ScopeReg = buildConstantIntReg32(Val: Scope, MIRBuilder, GR);
785
786 Register Expected =
787 IsCmpxchg ? ExpectedArg
788 : buildLoadInst(BaseType: SpvDesiredTy, PtrRegister: ExpectedArg, MIRBuilder, GR);
789 MRI->setType(VReg: Expected, Ty: DesiredLLT);
790 Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(Ty: DesiredLLT)
791 : Call->ReturnRegister;
792 if (!MRI->getRegClassOrNull(Reg: Tmp))
793 MRI->setRegClass(Reg: Tmp, RC: GR->getRegClass(SpvType: SpvDesiredTy));
794 GR->assignSPIRVTypeToVReg(Type: SpvDesiredTy, VReg: Tmp, MF: MIRBuilder.getMF());
795
796 MIRBuilder.buildInstr(Opcode)
797 .addDef(RegNo: Tmp)
798 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: SpvDesiredTy))
799 .addUse(RegNo: ObjectPtr)
800 .addUse(RegNo: ScopeReg)
801 .addUse(RegNo: MemSemEqualReg)
802 .addUse(RegNo: MemSemUnequalReg)
803 .addUse(RegNo: Desired)
804 .addUse(RegNo: Expected);
805 if (!IsCmpxchg) {
806 MIRBuilder.buildInstr(Opcode: SPIRV::OpStore).addUse(RegNo: ExpectedArg).addUse(RegNo: Tmp);
807 MIRBuilder.buildICmp(Pred: CmpInst::ICMP_EQ, Res: Call->ReturnRegister, Op0: Tmp, Op1: Expected);
808 }
809 return true;
810}
811
812/// Helper function for building atomic instructions.
813static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
814 MachineIRBuilder &MIRBuilder,
815 SPIRVGlobalRegistry *GR) {
816 if (Call->isSpirvOp())
817 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
818 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
819
820 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
821 Register ScopeRegister =
822 Call->Arguments.size() >= 4 ? Call->Arguments[3] : Register();
823
824 assert(Call->Arguments.size() <= 4 &&
825 "Too many args for explicit atomic RMW");
826 ScopeRegister = buildScopeReg(CLScopeRegister: ScopeRegister, Scope: SPIRV::Scope::Workgroup,
827 MIRBuilder, GR, MRI);
828
829 Register PtrRegister = Call->Arguments[0];
830 unsigned Semantics = SPIRV::MemorySemantics::None;
831 Register MemSemanticsReg =
832 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
833 MemSemanticsReg = buildMemSemanticsReg(SemanticsRegister: MemSemanticsReg, PtrRegister,
834 Semantics, MIRBuilder, GR);
835 Register ValueReg = Call->Arguments[1];
836 Register ValueTypeReg = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
837 // support cl_ext_float_atomics
838 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeFloat) {
839 if (Opcode == SPIRV::OpAtomicIAdd) {
840 Opcode = SPIRV::OpAtomicFAddEXT;
841 } else if (Opcode == SPIRV::OpAtomicISub) {
842 // Translate OpAtomicISub applied to a floating type argument to
843 // OpAtomicFAddEXT with the negative value operand
844 Opcode = SPIRV::OpAtomicFAddEXT;
845 Register NegValueReg =
846 MRI->createGenericVirtualRegister(Ty: MRI->getType(Reg: ValueReg));
847 MRI->setRegClass(Reg: NegValueReg, RC: GR->getRegClass(SpvType: Call->ReturnType));
848 GR->assignSPIRVTypeToVReg(Type: Call->ReturnType, VReg: NegValueReg,
849 MF: MIRBuilder.getMF());
850 MIRBuilder.buildInstr(Opcode: TargetOpcode::G_FNEG)
851 .addDef(RegNo: NegValueReg)
852 .addUse(RegNo: ValueReg);
853 updateRegType(Reg: NegValueReg, Ty: nullptr, SpirvTy: Call->ReturnType, GR, MIB&: MIRBuilder,
854 MRI&: MIRBuilder.getMF().getRegInfo());
855 ValueReg = NegValueReg;
856 }
857 }
858 MIRBuilder.buildInstr(Opcode)
859 .addDef(RegNo: Call->ReturnRegister)
860 .addUse(RegNo: ValueTypeReg)
861 .addUse(RegNo: PtrRegister)
862 .addUse(RegNo: ScopeRegister)
863 .addUse(RegNo: MemSemanticsReg)
864 .addUse(RegNo: ValueReg);
865 return true;
866}
867
868/// Helper function for building an atomic floating-type instruction.
869static bool buildAtomicFloatingRMWInst(const SPIRV::IncomingCall *Call,
870 unsigned Opcode,
871 MachineIRBuilder &MIRBuilder,
872 SPIRVGlobalRegistry *GR) {
873 assert(Call->Arguments.size() == 4 &&
874 "Wrong number of atomic floating-type builtin");
875 Register PtrReg = Call->Arguments[0];
876 Register ScopeReg = Call->Arguments[1];
877 Register MemSemanticsReg = Call->Arguments[2];
878 Register ValueReg = Call->Arguments[3];
879 MIRBuilder.buildInstr(Opcode)
880 .addDef(RegNo: Call->ReturnRegister)
881 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
882 .addUse(RegNo: PtrReg)
883 .addUse(RegNo: ScopeReg)
884 .addUse(RegNo: MemSemanticsReg)
885 .addUse(RegNo: ValueReg);
886 return true;
887}
888
889/// Helper function for building atomic flag instructions (e.g.
890/// OpAtomicFlagTestAndSet).
891static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call,
892 unsigned Opcode, MachineIRBuilder &MIRBuilder,
893 SPIRVGlobalRegistry *GR) {
894 bool IsSet = Opcode == SPIRV::OpAtomicFlagTestAndSet;
895 Register TypeReg = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
896 if (Call->isSpirvOp())
897 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
898 TypeReg: IsSet ? TypeReg : Register(0));
899
900 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
901 Register PtrRegister = Call->Arguments[0];
902 unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;
903 Register MemSemanticsReg =
904 Call->Arguments.size() >= 2 ? Call->Arguments[1] : Register();
905 MemSemanticsReg = buildMemSemanticsReg(SemanticsRegister: MemSemanticsReg, PtrRegister,
906 Semantics, MIRBuilder, GR);
907
908 assert((Opcode != SPIRV::OpAtomicFlagClear ||
909 (Semantics != SPIRV::MemorySemantics::Acquire &&
910 Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&
911 "Invalid memory order argument!");
912
913 Register ScopeRegister =
914 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
915 ScopeRegister =
916 buildScopeReg(CLScopeRegister: ScopeRegister, Scope: SPIRV::Scope::Device, MIRBuilder, GR, MRI);
917
918 auto MIB = MIRBuilder.buildInstr(Opcode);
919 if (IsSet)
920 MIB.addDef(RegNo: Call->ReturnRegister).addUse(RegNo: TypeReg);
921
922 MIB.addUse(RegNo: PtrRegister).addUse(RegNo: ScopeRegister).addUse(RegNo: MemSemanticsReg);
923 return true;
924}
925
926/// Helper function for building barriers, i.e., memory/control ordering
927/// operations.
928static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
929 MachineIRBuilder &MIRBuilder,
930 SPIRVGlobalRegistry *GR) {
931 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
932 const auto *ST =
933 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
934 if ((Opcode == SPIRV::OpControlBarrierArriveINTEL ||
935 Opcode == SPIRV::OpControlBarrierWaitINTEL) &&
936 !ST->canUseExtension(E: SPIRV::Extension::SPV_INTEL_split_barrier)) {
937 std::string DiagMsg = std::string(Builtin->name()) +
938 ": the builtin requires the following SPIR-V "
939 "extension: SPV_INTEL_split_barrier";
940 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
941 }
942
943 if (Call->isSpirvOp())
944 return buildOpFromWrapper(MIRBuilder, Opcode, Call, TypeReg: Register(0));
945
946 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
947 unsigned MemFlags = getIConstVal(ConstReg: Call->Arguments[0], MRI);
948 unsigned MemSemantics = SPIRV::MemorySemantics::None;
949
950 if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
951 MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
952
953 if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
954 MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;
955
956 if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)
957 MemSemantics |= SPIRV::MemorySemantics::ImageMemory;
958
959 if (Opcode == SPIRV::OpMemoryBarrier)
960 MemSemantics = getSPIRVMemSemantics(MemOrder: static_cast<std::memory_order>(
961 getIConstVal(ConstReg: Call->Arguments[1], MRI))) |
962 MemSemantics;
963 else if (Opcode == SPIRV::OpControlBarrierArriveINTEL)
964 MemSemantics |= SPIRV::MemorySemantics::Release;
965 else if (Opcode == SPIRV::OpControlBarrierWaitINTEL)
966 MemSemantics |= SPIRV::MemorySemantics::Acquire;
967 else
968 MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
969
970 Register MemSemanticsReg =
971 MemFlags == MemSemantics
972 ? Call->Arguments[0]
973 : buildConstantIntReg32(Val: MemSemantics, MIRBuilder, GR);
974 Register ScopeReg;
975 SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
976 SPIRV::Scope::Scope MemScope = Scope;
977 if (Call->Arguments.size() >= 2) {
978 assert(
979 ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
980 (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
981 "Extra args for explicitly scoped barrier");
982 Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]
983 : Call->Arguments[1];
984 SPIRV::CLMemoryScope CLScope =
985 static_cast<SPIRV::CLMemoryScope>(getIConstVal(ConstReg: ScopeArg, MRI));
986 MemScope = getSPIRVScope(ClScope: CLScope);
987 if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
988 (Opcode == SPIRV::OpMemoryBarrier))
989 Scope = MemScope;
990 if (CLScope == static_cast<unsigned>(Scope))
991 ScopeReg = Call->Arguments[1];
992 }
993
994 if (!ScopeReg.isValid())
995 ScopeReg = buildConstantIntReg32(Val: Scope, MIRBuilder, GR);
996
997 auto MIB = MIRBuilder.buildInstr(Opcode).addUse(RegNo: ScopeReg);
998 if (Opcode != SPIRV::OpMemoryBarrier)
999 MIB.addUse(RegNo: buildConstantIntReg32(Val: MemScope, MIRBuilder, GR));
1000 MIB.addUse(RegNo: MemSemanticsReg);
1001 return true;
1002}
1003
1004/// Helper function for building extended bit operations.
1005static bool buildExtendedBitOpsInst(const SPIRV::IncomingCall *Call,
1006 unsigned Opcode,
1007 MachineIRBuilder &MIRBuilder,
1008 SPIRVGlobalRegistry *GR) {
1009 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1010 const auto *ST =
1011 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
1012 if ((Opcode == SPIRV::OpBitFieldInsert ||
1013 Opcode == SPIRV::OpBitFieldSExtract ||
1014 Opcode == SPIRV::OpBitFieldUExtract || Opcode == SPIRV::OpBitReverse) &&
1015 !ST->canUseExtension(E: SPIRV::Extension::SPV_KHR_bit_instructions)) {
1016 std::string DiagMsg = std::string(Builtin->name()) +
1017 ": the builtin requires the following SPIR-V "
1018 "extension: SPV_KHR_bit_instructions";
1019 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
1020 }
1021
1022 // Generate SPIRV instruction accordingly.
1023 if (Call->isSpirvOp())
1024 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1025 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1026
1027 auto MIB = MIRBuilder.buildInstr(Opcode)
1028 .addDef(RegNo: Call->ReturnRegister)
1029 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1030 for (unsigned i = 0; i < Call->Arguments.size(); ++i)
1031 MIB.addUse(RegNo: Call->Arguments[i]);
1032
1033 return true;
1034}
1035
1036/// Helper function for building Intel's bindless image instructions.
1037static bool buildBindlessImageINTELInst(const SPIRV::IncomingCall *Call,
1038 unsigned Opcode,
1039 MachineIRBuilder &MIRBuilder,
1040 SPIRVGlobalRegistry *GR) {
1041 // Generate SPIRV instruction accordingly.
1042 if (Call->isSpirvOp())
1043 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1044 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1045
1046 MIRBuilder.buildInstr(Opcode)
1047 .addDef(RegNo: Call->ReturnRegister)
1048 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
1049 .addUse(RegNo: Call->Arguments[0]);
1050
1051 return true;
1052}
1053
1054/// Helper function for building Intel's OpBitwiseFunctionINTEL instruction.
1055static bool buildTernaryBitwiseFunctionINTELInst(
1056 const SPIRV::IncomingCall *Call, unsigned Opcode,
1057 MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
1058 // Generate SPIRV instruction accordingly.
1059 if (Call->isSpirvOp())
1060 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1061 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1062
1063 auto MIB = MIRBuilder.buildInstr(Opcode)
1064 .addDef(RegNo: Call->ReturnRegister)
1065 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1066 for (unsigned i = 0; i < Call->Arguments.size(); ++i)
1067 MIB.addUse(RegNo: Call->Arguments[i]);
1068
1069 return true;
1070}
1071
1072static bool buildImageChannelDataTypeInst(const SPIRV::IncomingCall *Call,
1073 unsigned Opcode,
1074 MachineIRBuilder &MIRBuilder,
1075 SPIRVGlobalRegistry *GR) {
1076 if (Call->isSpirvOp())
1077 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1078 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1079
1080 auto MIB = MIRBuilder.buildInstr(Opcode)
1081 .addDef(RegNo: Call->ReturnRegister)
1082 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1083 for (unsigned i = 0; i < Call->Arguments.size(); ++i)
1084 MIB.addUse(RegNo: Call->Arguments[i]);
1085
1086 return true;
1087}
1088
1089/// Helper function for building Intel's 2d block io instructions.
1090static bool build2DBlockIOINTELInst(const SPIRV::IncomingCall *Call,
1091 unsigned Opcode,
1092 MachineIRBuilder &MIRBuilder,
1093 SPIRVGlobalRegistry *GR) {
1094 // Generate SPIRV instruction accordingly.
1095 if (Call->isSpirvOp())
1096 return buildOpFromWrapper(MIRBuilder, Opcode, Call, TypeReg: Register(0));
1097
1098 auto MIB = MIRBuilder.buildInstr(Opcode)
1099 .addDef(RegNo: Call->ReturnRegister)
1100 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1101 for (unsigned i = 0; i < Call->Arguments.size(); ++i)
1102 MIB.addUse(RegNo: Call->Arguments[i]);
1103
1104 return true;
1105}
1106
1107static bool buildPipeInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
1108 unsigned Scope, MachineIRBuilder &MIRBuilder,
1109 SPIRVGlobalRegistry *GR) {
1110 switch (Opcode) {
1111 case SPIRV::OpCommitReadPipe:
1112 case SPIRV::OpCommitWritePipe:
1113 return buildOpFromWrapper(MIRBuilder, Opcode, Call, TypeReg: Register(0));
1114 case SPIRV::OpGroupCommitReadPipe:
1115 case SPIRV::OpGroupCommitWritePipe:
1116 case SPIRV::OpGroupReserveReadPipePackets:
1117 case SPIRV::OpGroupReserveWritePipePackets: {
1118 Register ScopeConstReg =
1119 MIRBuilder.buildConstant(Res: LLT::scalar(SizeInBits: 32), Val: Scope).getReg(Idx: 0);
1120 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1121 MRI->setRegClass(Reg: ScopeConstReg, RC: &SPIRV::iIDRegClass);
1122 MachineInstrBuilder MIB;
1123 MIB = MIRBuilder.buildInstr(Opcode);
1124 // Add Return register and type.
1125 if (Opcode == SPIRV::OpGroupReserveReadPipePackets ||
1126 Opcode == SPIRV::OpGroupReserveWritePipePackets)
1127 MIB.addDef(RegNo: Call->ReturnRegister)
1128 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1129
1130 MIB.addUse(RegNo: ScopeConstReg);
1131 for (unsigned int i = 0; i < Call->Arguments.size(); ++i)
1132 MIB.addUse(RegNo: Call->Arguments[i]);
1133
1134 return true;
1135 }
1136 default:
1137 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1138 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1139 }
1140}
1141
1142static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
1143 switch (dim) {
1144 case SPIRV::Dim::DIM_1D:
1145 case SPIRV::Dim::DIM_Buffer:
1146 return 1;
1147 case SPIRV::Dim::DIM_2D:
1148 case SPIRV::Dim::DIM_Cube:
1149 case SPIRV::Dim::DIM_Rect:
1150 return 2;
1151 case SPIRV::Dim::DIM_3D:
1152 return 3;
1153 default:
1154 report_fatal_error(reason: "Cannot get num components for given Dim");
1155 }
1156}
1157
1158/// Helper function for obtaining the number of size components.
1159static unsigned getNumSizeComponents(SPIRVTypeInst imgType) {
1160 assert(imgType->getOpcode() == SPIRV::OpTypeImage);
1161 auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(i: 2).getImm());
1162 unsigned numComps = getNumComponentsForDim(dim);
1163 bool arrayed = imgType->getOperand(i: 4).getImm() == 1;
1164 return arrayed ? numComps + 1 : numComps;
1165}
1166
1167static bool builtinMayNeedPromotionToVec(uint32_t BuiltinNumber) {
1168 switch (BuiltinNumber) {
1169 case SPIRV::OpenCLExtInst::s_min:
1170 case SPIRV::OpenCLExtInst::u_min:
1171 case SPIRV::OpenCLExtInst::s_max:
1172 case SPIRV::OpenCLExtInst::u_max:
1173 case SPIRV::OpenCLExtInst::fmax:
1174 case SPIRV::OpenCLExtInst::fmin:
1175 case SPIRV::OpenCLExtInst::fmax_common:
1176 case SPIRV::OpenCLExtInst::fmin_common:
1177 case SPIRV::OpenCLExtInst::s_clamp:
1178 case SPIRV::OpenCLExtInst::fclamp:
1179 case SPIRV::OpenCLExtInst::u_clamp:
1180 case SPIRV::OpenCLExtInst::mix:
1181 case SPIRV::OpenCLExtInst::step:
1182 case SPIRV::OpenCLExtInst::smoothstep:
1183 case SPIRV::OpenCLExtInst::ldexp:
1184 case SPIRV::OpenCLExtInst::pown:
1185 case SPIRV::OpenCLExtInst::rootn:
1186 return true;
1187 default:
1188 break;
1189 }
1190 return false;
1191}
1192
1193//===----------------------------------------------------------------------===//
1194// Implementation functions for each builtin group
1195//===----------------------------------------------------------------------===//
1196
1197static SmallVector<Register>
1198getBuiltinCallArguments(const SPIRV::IncomingCall *Call, uint32_t BuiltinNumber,
1199 MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
1200
1201 Register ReturnTypeId = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
1202 unsigned ResultElementCount =
1203 GR->getScalarOrVectorComponentCount(VReg: ReturnTypeId);
1204 bool MayNeedPromotionToVec =
1205 builtinMayNeedPromotionToVec(BuiltinNumber) && ResultElementCount > 1;
1206
1207 if (!MayNeedPromotionToVec)
1208 return {Call->Arguments.begin(), Call->Arguments.end()};
1209
1210 SmallVector<Register> Arguments;
1211 for (Register Argument : Call->Arguments) {
1212 Register VecArg = Argument;
1213 SPIRVTypeInst ArgumentType = GR->getSPIRVTypeForVReg(VReg: Argument);
1214 if (GR->getScalarOrVectorComponentCount(Type: ArgumentType) == 1 &&
1215 ArgumentType != Call->ReturnType) {
1216 SPIRVTypeInst VecType = GR->getOrCreateSPIRVVectorType(
1217 BaseType: ArgumentType, NumElements: ResultElementCount, MIRBuilder, /*EmitIR=*/true);
1218 VecArg = createVirtualRegister(SpvType: VecType, GR, MIRBuilder);
1219 Register VecTypeId = GR->getSPIRVTypeID(SpirvType: VecType);
1220 auto VecSplat = MIRBuilder.buildInstr(Opcode: SPIRV::OpCompositeConstruct)
1221 .addDef(RegNo: VecArg)
1222 .addUse(RegNo: VecTypeId);
1223 for (unsigned I = 0; I != ResultElementCount; ++I)
1224 VecSplat.addUse(RegNo: Argument);
1225 }
1226 Arguments.push_back(Elt: VecArg);
1227 }
1228 return Arguments;
1229}
1230
1231static bool generateExtInst(const SPIRV::IncomingCall *Call,
1232 MachineIRBuilder &MIRBuilder,
1233 SPIRVGlobalRegistry *GR, const CallBase &CB) {
1234 // Lookup the extended instruction number in the TableGen records.
1235 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1236 uint32_t Number =
1237 SPIRV::lookupExtendedBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Number;
1238 // fmin_common and fmax_common are now deprecated, and we should use fmin and
1239 // fmax with NotInf and NotNaN flags instead. Keep original number to add
1240 // later the NoNans and NoInfs flags.
1241 uint32_t OrigNumber = Number;
1242 const SPIRVSubtarget &ST =
1243 cast<SPIRVSubtarget>(Val: MIRBuilder.getMF().getSubtarget());
1244 if (ST.canUseExtension(E: SPIRV::Extension::SPV_KHR_float_controls2) &&
1245 (Number == SPIRV::OpenCLExtInst::fmin_common ||
1246 Number == SPIRV::OpenCLExtInst::fmax_common)) {
1247 Number = (Number == SPIRV::OpenCLExtInst::fmin_common)
1248 ? SPIRV::OpenCLExtInst::fmin
1249 : SPIRV::OpenCLExtInst::fmax;
1250 }
1251
1252 Register ReturnTypeId = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
1253 SmallVector<Register> Arguments =
1254 getBuiltinCallArguments(Call, BuiltinNumber: Number, MIRBuilder, GR);
1255
1256 MachineInstrBuilder MIB;
1257 if (ST.canUseExtension(E: SPIRV::Extension::SPV_KHR_fma) &&
1258 Number == SPIRV::OpenCLExtInst::fma) {
1259 // Use the SPIR-V fma instruction instead of the OpenCL extended
1260 // instruction if the extension is available.
1261 MIB = MIRBuilder.buildInstr(Opcode: SPIRV::OpFmaKHR)
1262 .addDef(RegNo: Call->ReturnRegister)
1263 .addUse(RegNo: ReturnTypeId);
1264 } else {
1265 // Build extended instruction.
1266 MIB = MIRBuilder.buildInstr(Opcode: SPIRV::OpExtInst)
1267 .addDef(RegNo: Call->ReturnRegister)
1268 .addUse(RegNo: ReturnTypeId)
1269 .addImm(Val: static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
1270 .addImm(Val: Number);
1271 }
1272
1273 for (Register Argument : Arguments)
1274 MIB.addUse(RegNo: Argument);
1275
1276 MIB.getInstr()->copyIRFlags(I: CB);
1277 if (OrigNumber == SPIRV::OpenCLExtInst::fmin_common ||
1278 OrigNumber == SPIRV::OpenCLExtInst::fmax_common) {
1279 // Add NoNans and NoInfs flags to fmin/fmax instruction.
1280 MIB.getInstr()->setFlag(MachineInstr::MIFlag::FmNoNans);
1281 MIB.getInstr()->setFlag(MachineInstr::MIFlag::FmNoInfs);
1282 }
1283
1284 // Derive fast-math flags from nofpclass attributes on the called function.
1285 // FPFastMathMode decoration is valid on ExtInst in Kernel environments
1286 // (SPIR-V core) or with SPV_KHR_float_controls2 for any environment.
1287 if (ST.isKernel() ||
1288 ST.canUseExtension(E: SPIRV::Extension::SPV_KHR_float_controls2)) {
1289 if (const Function *F = CB.getCalledFunction()) {
1290 bool AddNoNan = CB.getRetNoFPClass() & fcNan;
1291 bool AddNoInf = CB.getRetNoFPClass() & fcInf;
1292 FunctionType *FTy = F->getFunctionType();
1293 for (unsigned I = 0, E = FTy->getNumParams();
1294 I != E && (AddNoNan || AddNoInf); ++I) {
1295 if (!FTy->getParamType(i: I)->isFloatingPointTy())
1296 continue;
1297 FPClassTest ArgTest = CB.getParamNoFPClass(i: I);
1298 AddNoNan = AddNoNan && ArgTest & fcNan;
1299 AddNoInf = AddNoInf && ArgTest & fcInf;
1300 }
1301 if (AddNoNan)
1302 MIB.getInstr()->setFlag(MachineInstr::MIFlag::FmNoNans);
1303 if (AddNoInf)
1304 MIB.getInstr()->setFlag(MachineInstr::MIFlag::FmNoInfs);
1305 }
1306 }
1307
1308 return true;
1309}
1310
1311static bool generateRelationalInst(const SPIRV::IncomingCall *Call,
1312 MachineIRBuilder &MIRBuilder,
1313 SPIRVGlobalRegistry *GR) {
1314 // Lookup the instruction opcode in the TableGen records.
1315 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1316 unsigned Opcode =
1317 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
1318
1319 Register CompareRegister;
1320 SPIRVTypeInst RelationType = nullptr;
1321 std::tie(args&: CompareRegister, args&: RelationType) =
1322 buildBoolRegister(MIRBuilder, ResultType: Call->ReturnType, GR);
1323
1324 // OpAny/OpAll require a boolean vector input, but OpenCL any()/all()
1325 // builtins receive integer vectors. Convert via OpINotEqual against zero.
1326 SmallVector<Register> Arguments(Call->Arguments.begin(),
1327 Call->Arguments.end());
1328 if ((Opcode == SPIRV::OpAny || Opcode == SPIRV::OpAll) &&
1329 !GR->isScalarOrVectorOfType(VReg: Arguments[0], TypeOpcode: SPIRV::OpTypeBool)) {
1330 SPIRVTypeInst ArgType = GR->getSPIRVTypeForVReg(VReg: Arguments[0]);
1331 unsigned NumElts = GR->getScalarOrVectorComponentCount(Type: ArgType);
1332 SPIRVTypeInst BoolVecTy = GR->getOrCreateSPIRVVectorType(
1333 BaseType: GR->getOrCreateSPIRVBoolType(MIRBuilder, /*EmitIR=*/true), NumElements: NumElts,
1334 MIRBuilder, /*EmitIR=*/true);
1335 Register ZeroReg =
1336 GR->getOrCreateConsIntVector(Val: uint64_t(0), MIRBuilder, SpvType: ArgType,
1337 /*EmitIR=*/true);
1338 Register BoolVecReg = createVirtualRegister(SpvType: BoolVecTy, GR, MIRBuilder);
1339 MIRBuilder.buildInstr(Opcode: SPIRV::OpINotEqual)
1340 .addDef(RegNo: BoolVecReg)
1341 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: BoolVecTy))
1342 .addUse(RegNo: Arguments[0])
1343 .addUse(RegNo: ZeroReg);
1344 Arguments[0] = BoolVecReg;
1345 }
1346
1347 // Build relational instruction.
1348 auto MIB = MIRBuilder.buildInstr(Opcode)
1349 .addDef(RegNo: CompareRegister)
1350 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: RelationType));
1351
1352 for (auto Argument : Arguments)
1353 MIB.addUse(RegNo: Argument);
1354
1355 // Build select instruction.
1356 return buildSelectInst(MIRBuilder, ReturnRegister: Call->ReturnRegister, SourceRegister: CompareRegister,
1357 ReturnType: Call->ReturnType, GR);
1358}
1359
1360static bool generateGroupInst(const SPIRV::IncomingCall *Call,
1361 MachineIRBuilder &MIRBuilder,
1362 SPIRVGlobalRegistry *GR) {
1363 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1364 const SPIRV::GroupBuiltin *GroupBuiltin =
1365 SPIRV::lookupGroupBuiltin(Name: Builtin->name());
1366
1367 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1368 if (Call->isSpirvOp()) {
1369 if (GroupBuiltin->NoGroupOperation) {
1370 SmallVector<uint32_t, 1> ImmArgs;
1371 if (GroupBuiltin->Opcode ==
1372 SPIRV::OpSubgroupMatrixMultiplyAccumulateINTEL &&
1373 Call->Arguments.size() > 4)
1374 ImmArgs.push_back(Elt: getIConstVal(ConstReg: Call->Arguments[4], MRI));
1375 return buildOpFromWrapper(MIRBuilder, Opcode: GroupBuiltin->Opcode, Call,
1376 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType), ImmArgs);
1377 }
1378
1379 // Group Operation is a literal
1380 Register GroupOpReg = Call->Arguments[1];
1381 const MachineInstr *MI = getDefInstrMaybeConstant(ConstReg&: GroupOpReg, MRI);
1382 if (!MI || MI->getOpcode() != TargetOpcode::G_CONSTANT)
1383 report_fatal_error(
1384 reason: "Group Operation parameter must be an integer constant");
1385 uint64_t GrpOp = MI->getOperand(i: 1).getCImm()->getValue().getZExtValue();
1386 Register ScopeReg = Call->Arguments[0];
1387 auto MIB = MIRBuilder.buildInstr(Opcode: GroupBuiltin->Opcode)
1388 .addDef(RegNo: Call->ReturnRegister)
1389 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
1390 .addUse(RegNo: ScopeReg)
1391 .addImm(Val: GrpOp);
1392 for (unsigned i = 2; i < Call->Arguments.size(); ++i)
1393 MIB.addUse(RegNo: Call->Arguments[i]);
1394 return true;
1395 }
1396
1397 Register Arg0;
1398 if (GroupBuiltin->HasBoolArg) {
1399 SPIRVTypeInst BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder, EmitIR: true);
1400 Register BoolReg = Call->Arguments[0];
1401 SPIRVTypeInst BoolRegType = GR->getSPIRVTypeForVReg(VReg: BoolReg);
1402 if (!BoolRegType)
1403 report_fatal_error(reason: "Can't find a register's type definition");
1404 MachineInstr *ArgInstruction = getDefInstrMaybeConstant(ConstReg&: BoolReg, MRI);
1405 if (ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT) {
1406 if (BoolRegType->getOpcode() != SPIRV::OpTypeBool)
1407 Arg0 = GR->buildConstantInt(Val: getIConstVal(ConstReg: BoolReg, MRI), MIRBuilder,
1408 SpvType: BoolType, EmitIR: true);
1409 } else {
1410 if (BoolRegType->getOpcode() == SPIRV::OpTypeInt) {
1411 Arg0 = MRI->createGenericVirtualRegister(Ty: LLT::scalar(SizeInBits: 1));
1412 MRI->setRegClass(Reg: Arg0, RC: &SPIRV::iIDRegClass);
1413 GR->assignSPIRVTypeToVReg(Type: BoolType, VReg: Arg0, MF: MIRBuilder.getMF());
1414 MIRBuilder.buildICmp(
1415 Pred: CmpInst::ICMP_NE, Res: Arg0, Op0: BoolReg,
1416 Op1: GR->buildConstantInt(Val: 0, MIRBuilder, SpvType: BoolRegType, EmitIR: true));
1417 updateRegType(Reg: Arg0, Ty: nullptr, SpirvTy: BoolType, GR, MIB&: MIRBuilder,
1418 MRI&: MIRBuilder.getMF().getRegInfo());
1419 } else if (BoolRegType->getOpcode() != SPIRV::OpTypeBool) {
1420 report_fatal_error(reason: "Expect a boolean argument");
1421 }
1422 // if BoolReg is a boolean register, we don't need to do anything
1423 }
1424 }
1425
1426 Register GroupResultRegister = Call->ReturnRegister;
1427 SPIRVTypeInst GroupResultType = Call->ReturnType;
1428
1429 // TODO: maybe we need to check whether the result type is already boolean
1430 // and in this case do not insert select instruction.
1431 const bool HasBoolReturnTy =
1432 GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
1433 GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
1434 GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
1435
1436 if (HasBoolReturnTy)
1437 std::tie(args&: GroupResultRegister, args&: GroupResultType) =
1438 buildBoolRegister(MIRBuilder, ResultType: Call->ReturnType, GR);
1439
1440 auto Scope = Builtin->name().starts_with(Prefix: "sub_group")
1441 ? SPIRV::Scope::Subgroup
1442 : SPIRV::Scope::Workgroup;
1443 Register ScopeRegister = buildConstantIntReg32(Val: Scope, MIRBuilder, GR);
1444
1445 Register VecReg;
1446 if (GroupBuiltin->Opcode == SPIRV::OpGroupBroadcast &&
1447 Call->Arguments.size() > 2) {
1448 // For OpGroupBroadcast "LocalId must be an integer datatype. It must be a
1449 // scalar, a vector with 2 components, or a vector with 3 components.",
1450 // meaning that we must create a vector from the function arguments if
1451 // it's a work_group_broadcast(val, local_id_x, local_id_y) or
1452 // work_group_broadcast(val, local_id_x, local_id_y, local_id_z) call.
1453 Register ElemReg = Call->Arguments[1];
1454 SPIRVTypeInst ElemType = GR->getSPIRVTypeForVReg(VReg: ElemReg);
1455 if (!ElemType || ElemType->getOpcode() != SPIRV::OpTypeInt)
1456 report_fatal_error(reason: "Expect an integer <LocalId> argument");
1457 unsigned VecLen = Call->Arguments.size() - 1;
1458 VecReg = MRI->createGenericVirtualRegister(
1459 Ty: LLT::fixed_vector(NumElements: VecLen, ScalarTy: MRI->getType(Reg: ElemReg)));
1460 MRI->setRegClass(Reg: VecReg, RC: &SPIRV::viIDRegClass);
1461 SPIRVTypeInst VecType =
1462 GR->getOrCreateSPIRVVectorType(BaseType: ElemType, NumElements: VecLen, MIRBuilder, EmitIR: true);
1463 GR->assignSPIRVTypeToVReg(Type: VecType, VReg: VecReg, MF: MIRBuilder.getMF());
1464 auto MIB =
1465 MIRBuilder.buildInstr(Opcode: TargetOpcode::G_BUILD_VECTOR).addDef(RegNo: VecReg);
1466 for (unsigned i = 1; i < Call->Arguments.size(); i++) {
1467 MIB.addUse(RegNo: Call->Arguments[i]);
1468 setRegClassIfNull(Reg: Call->Arguments[i], MRI, GR);
1469 }
1470 updateRegType(Reg: VecReg, Ty: nullptr, SpirvTy: VecType, GR, MIB&: MIRBuilder,
1471 MRI&: MIRBuilder.getMF().getRegInfo());
1472 }
1473
1474 // Build work/sub group instruction.
1475 auto MIB = MIRBuilder.buildInstr(Opcode: GroupBuiltin->Opcode)
1476 .addDef(RegNo: GroupResultRegister)
1477 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: GroupResultType))
1478 .addUse(RegNo: ScopeRegister);
1479
1480 if (!GroupBuiltin->NoGroupOperation)
1481 MIB.addImm(Val: GroupBuiltin->GroupOperation);
1482 if (Call->Arguments.size() > 0) {
1483 MIB.addUse(RegNo: Arg0.isValid() ? Arg0 : Call->Arguments[0]);
1484 setRegClassIfNull(Reg: Call->Arguments[0], MRI, GR);
1485 if (VecReg.isValid())
1486 MIB.addUse(RegNo: VecReg);
1487 else
1488 for (unsigned i = 1; i < Call->Arguments.size(); i++)
1489 MIB.addUse(RegNo: Call->Arguments[i]);
1490 }
1491
1492 // Build select instruction.
1493 if (HasBoolReturnTy)
1494 buildSelectInst(MIRBuilder, ReturnRegister: Call->ReturnRegister, SourceRegister: GroupResultRegister,
1495 ReturnType: Call->ReturnType, GR);
1496 return true;
1497}
1498
1499static bool generateIntelSubgroupsInst(const SPIRV::IncomingCall *Call,
1500 MachineIRBuilder &MIRBuilder,
1501 SPIRVGlobalRegistry *GR) {
1502 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1503 MachineFunction &MF = MIRBuilder.getMF();
1504 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1505 const SPIRV::IntelSubgroupsBuiltin *IntelSubgroups =
1506 SPIRV::lookupIntelSubgroupsBuiltin(Name: Builtin->name());
1507
1508 if (IntelSubgroups->IsMedia &&
1509 !ST->canUseExtension(E: SPIRV::Extension::SPV_INTEL_media_block_io)) {
1510 std::string DiagMsg = std::string(Builtin->name()) +
1511 ": the builtin requires the following SPIR-V "
1512 "extension: SPV_INTEL_media_block_io";
1513 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
1514 } else if (!IntelSubgroups->IsMedia &&
1515 !ST->canUseExtension(E: SPIRV::Extension::SPV_INTEL_subgroups)) {
1516 std::string DiagMsg = std::string(Builtin->name()) +
1517 ": the builtin requires the following SPIR-V "
1518 "extension: SPV_INTEL_subgroups";
1519 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
1520 }
1521
1522 uint32_t OpCode = IntelSubgroups->Opcode;
1523 if (Call->isSpirvOp()) {
1524 bool IsSet = OpCode != SPIRV::OpSubgroupBlockWriteINTEL &&
1525 OpCode != SPIRV::OpSubgroupImageBlockWriteINTEL &&
1526 OpCode != SPIRV::OpSubgroupImageMediaBlockWriteINTEL;
1527 return buildOpFromWrapper(MIRBuilder, Opcode: OpCode, Call,
1528 TypeReg: IsSet ? GR->getSPIRVTypeID(SpirvType: Call->ReturnType)
1529 : Register(0));
1530 }
1531
1532 if (IntelSubgroups->IsBlock) {
1533 // Minimal number or arguments set in TableGen records is 1
1534 if (SPIRVTypeInst Arg0Type = GR->getSPIRVTypeForVReg(VReg: Call->Arguments[0])) {
1535 if (Arg0Type->getOpcode() == SPIRV::OpTypeImage) {
1536 // TODO: add required validation from the specification:
1537 // "'Image' must be an object whose type is OpTypeImage with a 'Sampled'
1538 // operand of 0 or 2. If the 'Sampled' operand is 2, then some
1539 // dimensions require a capability."
1540 switch (OpCode) {
1541 case SPIRV::OpSubgroupBlockReadINTEL:
1542 OpCode = SPIRV::OpSubgroupImageBlockReadINTEL;
1543 break;
1544 case SPIRV::OpSubgroupBlockWriteINTEL:
1545 OpCode = SPIRV::OpSubgroupImageBlockWriteINTEL;
1546 break;
1547 }
1548 }
1549 }
1550 }
1551
1552 // TODO: opaque pointers types should be eventually resolved in such a way
1553 // that validation of block read is enabled with respect to the following
1554 // specification requirement:
1555 // "'Result Type' may be a scalar or vector type, and its component type must
1556 // be equal to the type pointed to by 'Ptr'."
1557 // For example, function parameter type should not be default i8 pointer, but
1558 // depend on the result type of the instruction where it is used as a pointer
1559 // argument of OpSubgroupBlockReadINTEL
1560
1561 // Build Intel subgroups instruction
1562 MachineInstrBuilder MIB =
1563 IntelSubgroups->IsWrite
1564 ? MIRBuilder.buildInstr(Opcode: OpCode)
1565 : MIRBuilder.buildInstr(Opcode: OpCode)
1566 .addDef(RegNo: Call->ReturnRegister)
1567 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1568 for (size_t i = 0; i < Call->Arguments.size(); ++i)
1569 MIB.addUse(RegNo: Call->Arguments[i]);
1570 return true;
1571}
1572
1573static bool generateGroupUniformInst(const SPIRV::IncomingCall *Call,
1574 MachineIRBuilder &MIRBuilder,
1575 SPIRVGlobalRegistry *GR) {
1576 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1577 MachineFunction &MF = MIRBuilder.getMF();
1578 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1579 if (!ST->canUseExtension(
1580 E: SPIRV::Extension::SPV_KHR_uniform_group_instructions)) {
1581 std::string DiagMsg = std::string(Builtin->name()) +
1582 ": the builtin requires the following SPIR-V "
1583 "extension: SPV_KHR_uniform_group_instructions";
1584 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
1585 }
1586 const SPIRV::GroupUniformBuiltin *GroupUniform =
1587 SPIRV::lookupGroupUniformBuiltin(Name: Builtin->name());
1588 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1589
1590 Register GroupResultReg = Call->ReturnRegister;
1591 Register ScopeReg = Call->Arguments[0];
1592 Register ValueReg = Call->Arguments[2];
1593
1594 // Group Operation
1595 Register ConstGroupOpReg = Call->Arguments[1];
1596 const MachineInstr *Const = getDefInstrMaybeConstant(ConstReg&: ConstGroupOpReg, MRI);
1597 if (!Const || Const->getOpcode() != TargetOpcode::G_CONSTANT)
1598 report_fatal_error(
1599 reason: "expect a constant group operation for a uniform group instruction",
1600 gen_crash_diag: false);
1601 const MachineOperand &ConstOperand = Const->getOperand(i: 1);
1602 if (!ConstOperand.isCImm())
1603 report_fatal_error(reason: "uniform group instructions: group operation must be an "
1604 "integer constant",
1605 gen_crash_diag: false);
1606
1607 auto MIB = MIRBuilder.buildInstr(Opcode: GroupUniform->Opcode)
1608 .addDef(RegNo: GroupResultReg)
1609 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
1610 .addUse(RegNo: ScopeReg);
1611 addNumImm(Imm: ConstOperand.getCImm()->getValue(), MIB);
1612 MIB.addUse(RegNo: ValueReg);
1613
1614 return true;
1615}
1616
1617static bool generateKernelClockInst(const SPIRV::IncomingCall *Call,
1618 MachineIRBuilder &MIRBuilder,
1619 SPIRVGlobalRegistry *GR) {
1620 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1621 MachineFunction &MF = MIRBuilder.getMF();
1622 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1623 if (!ST->canUseExtension(E: SPIRV::Extension::SPV_KHR_shader_clock)) {
1624 std::string DiagMsg = std::string(Builtin->name()) +
1625 ": the builtin requires the following SPIR-V "
1626 "extension: SPV_KHR_shader_clock";
1627 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
1628 }
1629
1630 Register ResultReg = Call->ReturnRegister;
1631
1632 if (Builtin->name() == "__spirv_ReadClockKHR") {
1633 MIRBuilder.buildInstr(Opcode: SPIRV::OpReadClockKHR)
1634 .addDef(RegNo: ResultReg)
1635 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
1636 .addUse(RegNo: Call->Arguments[0]);
1637 } else {
1638 // Deduce the `Scope` operand from the builtin function name.
1639 SPIRV::Scope::Scope ScopeArg =
1640 StringSwitch<SPIRV::Scope::Scope>(Builtin->name())
1641 .EndsWith(S: "device", Value: SPIRV::Scope::Scope::Device)
1642 .EndsWith(S: "work_group", Value: SPIRV::Scope::Scope::Workgroup)
1643 .EndsWith(S: "sub_group", Value: SPIRV::Scope::Scope::Subgroup);
1644 Register ScopeReg = buildConstantIntReg32(Val: ScopeArg, MIRBuilder, GR);
1645
1646 MIRBuilder.buildInstr(Opcode: SPIRV::OpReadClockKHR)
1647 .addDef(RegNo: ResultReg)
1648 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
1649 .addUse(RegNo: ScopeReg);
1650 }
1651
1652 return true;
1653}
1654
1655// These queries ask for a single size_t result for a given dimension index,
1656// e.g. size_t get_global_id(uint dimindex). In SPIR-V, the builtins
1657// corresponding to these values are all vec3 types, so we need to extract the
1658// correct index or return DefaultValue (0 or 1 depending on the query). We also
1659// handle extending or truncating in case size_t does not match the expected
1660// result type's bitwidth.
1661//
1662// For a constant index >= 3 we generate:
1663// %res = OpConstant %SizeT DefaultValue
1664//
1665// For other indices we generate:
1666// %g = OpVariable %ptr_V3_SizeT Input
1667// OpDecorate %g BuiltIn XXX
1668// OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"
1669// OpDecorate %g Constant
1670// %loadedVec = OpLoad %V3_SizeT %g
1671//
1672// Then, if the index is constant < 3, we generate:
1673// %res = OpCompositeExtract %SizeT %loadedVec idx
1674// If the index is dynamic, we generate:
1675// %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
1676// %cmp = OpULessThan %bool %idx %const_3
1677// %res = OpSelect %SizeT %cmp %tmp %const_<DefaultValue>
1678//
1679// If the bitwidth of %res does not match the expected return type, we add an
1680// extend or truncate.
1681static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call,
1682 MachineIRBuilder &MIRBuilder,
1683 SPIRVGlobalRegistry *GR,
1684 SPIRV::BuiltIn::BuiltIn BuiltinValue,
1685 uint64_t DefaultValue) {
1686 Register IndexRegister = Call->Arguments[0];
1687 const unsigned ResultWidth = Call->ReturnType->getOperand(i: 1).getImm();
1688 const unsigned PointerSize = GR->getPointerSize();
1689 const SPIRVTypeInst PointerSizeType =
1690 GR->getOrCreateSPIRVIntegerType(BitWidth: PointerSize, MIRBuilder);
1691 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1692 auto IndexInstruction = getDefInstrMaybeConstant(ConstReg&: IndexRegister, MRI);
1693
1694 // Set up the final register to do truncation or extension on at the end.
1695 Register ToTruncate = Call->ReturnRegister;
1696
1697 // If the index is constant, we can statically determine if it is in range.
1698 bool IsConstantIndex =
1699 IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
1700
1701 // If it's out of range (max dimension is 3), we can just return the constant
1702 // default value (0 or 1 depending on which query function).
1703 if (IsConstantIndex && getIConstVal(ConstReg: IndexRegister, MRI) >= 3) {
1704 Register DefaultReg = Call->ReturnRegister;
1705 if (PointerSize != ResultWidth) {
1706 DefaultReg = MRI->createGenericVirtualRegister(Ty: LLT::scalar(SizeInBits: PointerSize));
1707 MRI->setRegClass(Reg: DefaultReg, RC: &SPIRV::iIDRegClass);
1708 GR->assignSPIRVTypeToVReg(Type: PointerSizeType, VReg: DefaultReg,
1709 MF: MIRBuilder.getMF());
1710 ToTruncate = DefaultReg;
1711 }
1712 auto NewRegister =
1713 GR->buildConstantInt(Val: DefaultValue, MIRBuilder, SpvType: PointerSizeType, EmitIR: true);
1714 MIRBuilder.buildCopy(Res: DefaultReg, Op: NewRegister);
1715 } else { // If it could be in range, we need to load from the given builtin.
1716 auto Vec3Ty =
1717 GR->getOrCreateSPIRVVectorType(BaseType: PointerSizeType, NumElements: 3, MIRBuilder, EmitIR: true);
1718 Register LoadedVector =
1719 buildBuiltinVariableLoad(MIRBuilder, VariableType: Vec3Ty, GR, BuiltinValue,
1720 LLType: LLT::fixed_vector(NumElements: 3, ScalarSizeInBits: PointerSize));
1721 // Set up the vreg to extract the result to (possibly a new temporary one).
1722 Register Extracted = Call->ReturnRegister;
1723 if (!IsConstantIndex || PointerSize != ResultWidth) {
1724 Extracted = MRI->createGenericVirtualRegister(Ty: LLT::scalar(SizeInBits: PointerSize));
1725 MRI->setRegClass(Reg: Extracted, RC: &SPIRV::iIDRegClass);
1726 GR->assignSPIRVTypeToVReg(Type: PointerSizeType, VReg: Extracted, MF: MIRBuilder.getMF());
1727 }
1728 // Use Intrinsic::spv_extractelt so dynamic vs static extraction is
1729 // handled later: extr = spv_extractelt LoadedVector, IndexRegister.
1730 MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
1731 ID: Intrinsic::spv_extractelt, Res: ArrayRef<Register>{Extracted}, HasSideEffects: true, isConvergent: false);
1732 ExtractInst.addUse(RegNo: LoadedVector).addUse(RegNo: IndexRegister);
1733
1734 // If the index is dynamic, need check if it's < 3, and then use a select.
1735 if (!IsConstantIndex) {
1736 updateRegType(Reg: Extracted, Ty: nullptr, SpirvTy: PointerSizeType, GR, MIB&: MIRBuilder, MRI&: *MRI);
1737
1738 auto IndexType = GR->getSPIRVTypeForVReg(VReg: IndexRegister);
1739 auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder, EmitIR: true);
1740
1741 Register CompareRegister =
1742 MRI->createGenericVirtualRegister(Ty: LLT::scalar(SizeInBits: 1));
1743 MRI->setRegClass(Reg: CompareRegister, RC: &SPIRV::iIDRegClass);
1744 GR->assignSPIRVTypeToVReg(Type: BoolType, VReg: CompareRegister, MF: MIRBuilder.getMF());
1745
1746 // Use G_ICMP to check if idxVReg < 3.
1747 MIRBuilder.buildICmp(
1748 Pred: CmpInst::ICMP_ULT, Res: CompareRegister, Op0: IndexRegister,
1749 Op1: GR->buildConstantInt(Val: 3, MIRBuilder, SpvType: IndexType, EmitIR: true));
1750
1751 // Get constant for the default value (0 or 1 depending on which
1752 // function).
1753 Register DefaultRegister =
1754 GR->buildConstantInt(Val: DefaultValue, MIRBuilder, SpvType: PointerSizeType, EmitIR: true);
1755
1756 // Get a register for the selection result (possibly a new temporary one).
1757 Register SelectionResult = Call->ReturnRegister;
1758 if (PointerSize != ResultWidth) {
1759 SelectionResult =
1760 MRI->createGenericVirtualRegister(Ty: LLT::scalar(SizeInBits: PointerSize));
1761 MRI->setRegClass(Reg: SelectionResult, RC: &SPIRV::iIDRegClass);
1762 GR->assignSPIRVTypeToVReg(Type: PointerSizeType, VReg: SelectionResult,
1763 MF: MIRBuilder.getMF());
1764 }
1765 // Create the final G_SELECT to return the extracted value or the default.
1766 MIRBuilder.buildSelect(Res: SelectionResult, Tst: CompareRegister, Op0: Extracted,
1767 Op1: DefaultRegister);
1768 ToTruncate = SelectionResult;
1769 } else {
1770 ToTruncate = Extracted;
1771 }
1772 }
1773 // Alter the result's bitwidth if it does not match the SizeT value extracted.
1774 if (PointerSize != ResultWidth)
1775 MIRBuilder.buildZExtOrTrunc(Res: Call->ReturnRegister, Op: ToTruncate);
1776 return true;
1777}
1778
1779static bool generateBuiltinVar(const SPIRV::IncomingCall *Call,
1780 MachineIRBuilder &MIRBuilder,
1781 SPIRVGlobalRegistry *GR) {
1782 // Lookup the builtin variable record.
1783 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1784 SPIRV::BuiltIn::BuiltIn Value =
1785 SPIRV::lookupGetBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Value;
1786
1787 if (Value == SPIRV::BuiltIn::GlobalInvocationId)
1788 return genWorkgroupQuery(Call, MIRBuilder, GR, BuiltinValue: Value, DefaultValue: 0);
1789
1790 // Build a load instruction for the builtin variable.
1791 unsigned BitWidth = GR->getScalarOrVectorBitWidth(Type: Call->ReturnType);
1792 LLT LLType;
1793 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
1794 LLType = LLT::fixed_vector(
1795 NumElements: GR->getScalarOrVectorComponentCount(Type: Call->ReturnType), ScalarSizeInBits: BitWidth);
1796 else
1797 LLType = LLT::scalar(SizeInBits: BitWidth);
1798
1799 return buildBuiltinVariableLoad(MIRBuilder, VariableType: Call->ReturnType, GR, BuiltinValue: Value,
1800 LLType, Reg: Call->ReturnRegister);
1801}
1802
1803static bool generateAtomicInst(const SPIRV::IncomingCall *Call,
1804 MachineIRBuilder &MIRBuilder,
1805 SPIRVGlobalRegistry *GR) {
1806 // Lookup the instruction opcode in the TableGen records.
1807 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1808 unsigned Opcode =
1809 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
1810
1811 switch (Opcode) {
1812 case SPIRV::OpStore:
1813 return buildAtomicInitInst(Call, MIRBuilder);
1814 case SPIRV::OpAtomicLoad:
1815 return buildAtomicLoadInst(Call, MIRBuilder, GR);
1816 case SPIRV::OpAtomicStore:
1817 return buildAtomicStoreInst(Call, MIRBuilder, GR);
1818 case SPIRV::OpAtomicCompareExchange:
1819 case SPIRV::OpAtomicCompareExchangeWeak:
1820 return buildAtomicCompareExchangeInst(Call, Builtin, Opcode, MIRBuilder,
1821 GR);
1822 case SPIRV::OpAtomicIAdd:
1823 case SPIRV::OpAtomicISub:
1824 case SPIRV::OpAtomicOr:
1825 case SPIRV::OpAtomicXor:
1826 case SPIRV::OpAtomicAnd:
1827 case SPIRV::OpAtomicExchange:
1828 case SPIRV::OpAtomicSMax:
1829 case SPIRV::OpAtomicSMin:
1830 case SPIRV::OpAtomicUMax:
1831 case SPIRV::OpAtomicUMin:
1832 return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
1833 case SPIRV::OpMemoryBarrier:
1834 return buildBarrierInst(Call, Opcode: SPIRV::OpMemoryBarrier, MIRBuilder, GR);
1835 case SPIRV::OpAtomicFlagTestAndSet:
1836 case SPIRV::OpAtomicFlagClear:
1837 return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);
1838 default:
1839 if (Call->isSpirvOp())
1840 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1841 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1842 return false;
1843 }
1844}
1845
1846static bool generateAtomicFloatingInst(const SPIRV::IncomingCall *Call,
1847 MachineIRBuilder &MIRBuilder,
1848 SPIRVGlobalRegistry *GR) {
1849 // Lookup the instruction opcode in the TableGen records.
1850 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1851 unsigned Opcode = SPIRV::lookupAtomicFloatingBuiltin(Name: Builtin->name())->Opcode;
1852
1853 switch (Opcode) {
1854 case SPIRV::OpAtomicFAddEXT:
1855 case SPIRV::OpAtomicFMinEXT:
1856 case SPIRV::OpAtomicFMaxEXT:
1857 return buildAtomicFloatingRMWInst(Call, Opcode, MIRBuilder, GR);
1858 default:
1859 return false;
1860 }
1861}
1862
1863static bool generateBarrierInst(const SPIRV::IncomingCall *Call,
1864 MachineIRBuilder &MIRBuilder,
1865 SPIRVGlobalRegistry *GR) {
1866 // Lookup the instruction opcode in the TableGen records.
1867 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1868 unsigned Opcode =
1869 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
1870
1871 return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
1872}
1873
1874static bool generateCastToPtrInst(const SPIRV::IncomingCall *Call,
1875 MachineIRBuilder &MIRBuilder,
1876 SPIRVGlobalRegistry *GR) {
1877 // Lookup the instruction opcode in the TableGen records.
1878 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1879 unsigned Opcode =
1880 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
1881
1882 if (Opcode == SPIRV::OpGenericCastToPtrExplicit) {
1883 SPIRV::StorageClass::StorageClass ResSC =
1884 GR->getPointerStorageClass(VReg: Call->ReturnRegister);
1885 if (!isGenericCastablePtr(SC: ResSC))
1886 return false;
1887
1888 MIRBuilder.buildInstr(Opcode)
1889 .addDef(RegNo: Call->ReturnRegister)
1890 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
1891 .addUse(RegNo: Call->Arguments[0])
1892 .addImm(Val: ResSC);
1893 } else {
1894 MIRBuilder.buildInstr(Opcode: TargetOpcode::G_ADDRSPACE_CAST)
1895 .addDef(RegNo: Call->ReturnRegister)
1896 .addUse(RegNo: Call->Arguments[0]);
1897 }
1898 return true;
1899}
1900
1901static bool generateDotOrFMulInst(StringRef DemangledCall,
1902 const SPIRV::IncomingCall *Call,
1903 MachineIRBuilder &MIRBuilder,
1904 SPIRVGlobalRegistry *GR) {
1905 if (Call->isSpirvOp())
1906 return buildOpFromWrapper(MIRBuilder, Opcode: SPIRV::OpDot, Call,
1907 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1908
1909 bool IsVec = GR->getSPIRVTypeForVReg(VReg: Call->Arguments[0])->getOpcode() ==
1910 SPIRV::OpTypeVector;
1911 // Use OpDot only in case of vector args and OpFMul in case of scalar args.
1912 uint32_t OC = IsVec ? SPIRV::OpDot : SPIRV::OpFMulS;
1913 bool IsSwapReq = false;
1914
1915 const auto *ST =
1916 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
1917 if (GR->isScalarOrVectorOfType(VReg: Call->ReturnRegister, TypeOpcode: SPIRV::OpTypeInt) &&
1918 (ST->canUseExtension(E: SPIRV::Extension::SPV_KHR_integer_dot_product) ||
1919 ST->isAtLeastSPIRVVer(VerToCompareTo: VersionTuple(1, 6)))) {
1920 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1921 const SPIRV::IntegerDotProductBuiltin *IntDot =
1922 SPIRV::lookupIntegerDotProductBuiltin(Name: Builtin->name());
1923 if (IntDot) {
1924 OC = IntDot->Opcode;
1925 IsSwapReq = IntDot->IsSwapReq;
1926 } else if (IsVec) {
1927 // Handling "dot" and "dot_acc_sat" builtins which use vectors of
1928 // integers.
1929 LLVMContext &Ctx = MIRBuilder.getContext();
1930 SmallVector<StringRef, 10> TypeStrs;
1931 SPIRV::parseBuiltinTypeStr(BuiltinArgsTypeStrs&: TypeStrs, DemangledCall, Ctx);
1932 bool IsFirstSigned = TypeStrs[0].trim()[0] != 'u';
1933 bool IsSecondSigned = TypeStrs[1].trim()[0] != 'u';
1934
1935 if (Call->BuiltinName == "dot") {
1936 if (IsFirstSigned && IsSecondSigned)
1937 OC = SPIRV::OpSDot;
1938 else if (!IsFirstSigned && !IsSecondSigned)
1939 OC = SPIRV::OpUDot;
1940 else {
1941 OC = SPIRV::OpSUDot;
1942 if (!IsFirstSigned)
1943 IsSwapReq = true;
1944 }
1945 } else if (Call->BuiltinName == "dot_acc_sat") {
1946 if (IsFirstSigned && IsSecondSigned)
1947 OC = SPIRV::OpSDotAccSat;
1948 else if (!IsFirstSigned && !IsSecondSigned)
1949 OC = SPIRV::OpUDotAccSat;
1950 else {
1951 OC = SPIRV::OpSUDotAccSat;
1952 if (!IsFirstSigned)
1953 IsSwapReq = true;
1954 }
1955 }
1956 }
1957 }
1958
1959 MachineInstrBuilder MIB = MIRBuilder.buildInstr(Opcode: OC)
1960 .addDef(RegNo: Call->ReturnRegister)
1961 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
1962
1963 if (IsSwapReq) {
1964 MIB.addUse(RegNo: Call->Arguments[1]);
1965 MIB.addUse(RegNo: Call->Arguments[0]);
1966 // needed for dot_acc_sat* builtins
1967 for (size_t i = 2; i < Call->Arguments.size(); ++i)
1968 MIB.addUse(RegNo: Call->Arguments[i]);
1969 } else {
1970 for (size_t i = 0; i < Call->Arguments.size(); ++i)
1971 MIB.addUse(RegNo: Call->Arguments[i]);
1972 }
1973
1974 // Add Packed Vector Format for Integer dot product builtins if arguments are
1975 // scalar
1976 if (!IsVec && OC != SPIRV::OpFMulS)
1977 MIB.addImm(Val: SPIRV::PackedVectorFormat4x8Bit);
1978
1979 return true;
1980}
1981
1982static bool generateWaveInst(const SPIRV::IncomingCall *Call,
1983 MachineIRBuilder &MIRBuilder,
1984 SPIRVGlobalRegistry *GR) {
1985 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1986 SPIRV::BuiltIn::BuiltIn Value =
1987 SPIRV::lookupGetBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Value;
1988
1989 // For now, we only support a single Wave intrinsic with a single return type.
1990 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt);
1991 LLT LLType = LLT::scalar(SizeInBits: GR->getScalarOrVectorBitWidth(Type: Call->ReturnType));
1992
1993 return buildBuiltinVariableLoad(
1994 MIRBuilder, VariableType: Call->ReturnType, GR, BuiltinValue: Value, LLType, Reg: Call->ReturnRegister,
1995 /* isConst= */ false, /* LinkageType= */ LinkageTy: std::nullopt);
1996}
1997
1998// Build a SPIR-V instruction with struct return via sret pointer:
1999// Res = Opcode RetType Op1 Op2
2000// OpStore SRetReg Res
2001static void buildSRetInst(unsigned Opcode, Register SRetReg, Register Op1Reg,
2002 Register Op2Reg, SPIRVTypeInst RetType,
2003 MachineIRBuilder &MIRBuilder,
2004 SPIRVGlobalRegistry *GR) {
2005 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2006 Register ResReg = MRI->createVirtualRegister(RegClass: &SPIRV::iIDRegClass);
2007 if (const TargetRegisterClass *DstRC = MRI->getRegClassOrNull(Reg: Op1Reg)) {
2008 MRI->setRegClass(Reg: ResReg, RC: DstRC);
2009 MRI->setType(VReg: ResReg, Ty: MRI->getType(Reg: Op1Reg));
2010 }
2011 GR->assignSPIRVTypeToVReg(Type: RetType, VReg: ResReg, MF: MIRBuilder.getMF());
2012 MIRBuilder.buildInstr(Opcode)
2013 .addDef(RegNo: ResReg)
2014 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: RetType))
2015 .addUse(RegNo: Op1Reg)
2016 .addUse(RegNo: Op2Reg);
2017 MIRBuilder.buildInstr(Opcode: SPIRV::OpStore).addUse(RegNo: SRetReg).addUse(RegNo: ResReg);
2018}
2019
2020// We expect a builtin
2021// Name(ptr sret([RetType]) %result, Type %operand1, Type %operand1)
2022// where %result is a pointer to where the result of the builtin execution
2023// is to be stored, and generate the following instructions:
2024// Res = Opcode RetType Operand1 Operand1
2025// OpStore RetVariable Res
2026static bool generateICarryBorrowInst(const SPIRV::IncomingCall *Call,
2027 MachineIRBuilder &MIRBuilder,
2028 SPIRVGlobalRegistry *GR) {
2029 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2030 unsigned Opcode =
2031 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2032
2033 Register SRetReg = Call->Arguments[0];
2034 SPIRVTypeInst PtrRetType = GR->getSPIRVTypeForVReg(VReg: SRetReg);
2035 SPIRVTypeInst RetType = GR->getPointeeType(PtrType: PtrRetType);
2036 if (!RetType)
2037 report_fatal_error(reason: "The first parameter must be a pointer");
2038 if (RetType->getOpcode() != SPIRV::OpTypeStruct)
2039 report_fatal_error(reason: "Expected struct type result for the arithmetic with "
2040 "overflow builtins");
2041
2042 SPIRVTypeInst OpType1 = GR->getSPIRVTypeForVReg(VReg: Call->Arguments[1]);
2043 SPIRVTypeInst OpType2 = GR->getSPIRVTypeForVReg(VReg: Call->Arguments[2]);
2044 if (!OpType1 || !OpType2 || OpType1 != OpType2)
2045 report_fatal_error(reason: "Operands must have the same type");
2046 if (OpType1->getOpcode() == SPIRV::OpTypeVector)
2047 switch (Opcode) {
2048 case SPIRV::OpIAddCarryS:
2049 Opcode = SPIRV::OpIAddCarryV;
2050 break;
2051 case SPIRV::OpISubBorrowS:
2052 Opcode = SPIRV::OpISubBorrowV;
2053 break;
2054 }
2055
2056 buildSRetInst(Opcode, SRetReg, Op1Reg: Call->Arguments[1], Op2Reg: Call->Arguments[2],
2057 RetType, MIRBuilder, GR);
2058 return true;
2059}
2060
2061// We expect a builtin in one of two forms:
2062//
2063// (1) sret convention (3 arguments):
2064// void Name(ptr sret([RetType]) %result, Type %operand1, Type %operand2)
2065// => Res = Opcode RetType Operand1 Operand2
2066// OpStore %result Res
2067//
2068// (2) direct return convention (2 arguments):
2069// RetType Name(Type %operand1, Type %operand2)
2070// => Res = Opcode RetType Operand1 Operand2
2071//
2072// RetType is a struct with two members of the same type as the operands.
2073static bool generateMulExtendedInst(const SPIRV::IncomingCall *Call,
2074 MachineIRBuilder &MIRBuilder,
2075 SPIRVGlobalRegistry *GR) {
2076 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2077 unsigned Opcode =
2078 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2079 assert((Opcode == SPIRV::OpUMulExtended || Opcode == SPIRV::OpSMulExtended) &&
2080 "Expected OpUMulExtended or OpSMulExtended");
2081
2082 const bool IsSret =
2083 !Call->ReturnType || Call->ReturnType->getOpcode() == SPIRV::OpTypeVoid;
2084 Register Op1Reg = IsSret ? Call->Arguments[1] : Call->Arguments[0];
2085 Register Op2Reg = IsSret ? Call->Arguments[2] : Call->Arguments[1];
2086
2087 SPIRVTypeInst RetType = nullptr;
2088 if (IsSret) {
2089 Register SRetReg = Call->Arguments[0];
2090 SPIRVTypeInst PtrRetType = GR->getSPIRVTypeForVReg(VReg: SRetReg);
2091 RetType = GR->getPointeeType(PtrType: PtrRetType);
2092 if (!RetType)
2093 report_fatal_error(reason: "The first parameter must be a pointer");
2094 } else {
2095 RetType = Call->ReturnType;
2096 }
2097
2098 if (!RetType || RetType->getOpcode() != SPIRV::OpTypeStruct)
2099 report_fatal_error(reason: "Expected struct type result for the extended "
2100 "multiplication builtins");
2101 if (RetType->getNumOperands() != 3)
2102 report_fatal_error(reason: "Expected struct with exactly two members for the "
2103 "extended multiplication builtins");
2104 SPIRVTypeInst Member0Type =
2105 GR->getSPIRVTypeForVReg(VReg: RetType->getOperand(i: 1).getReg());
2106 SPIRVTypeInst Member1Type =
2107 GR->getSPIRVTypeForVReg(VReg: RetType->getOperand(i: 2).getReg());
2108 if (!Member0Type || !Member1Type || Member0Type != Member1Type)
2109 report_fatal_error(reason: "Both struct members must be the same type");
2110
2111 SPIRVTypeInst OpType1 = GR->getSPIRVTypeForVReg(VReg: Op1Reg);
2112 SPIRVTypeInst OpType2 = GR->getSPIRVTypeForVReg(VReg: Op2Reg);
2113 if (!OpType1 || !OpType2 || OpType1 != OpType2)
2114 report_fatal_error(reason: "Operands must have the same type");
2115 if (OpType1 != Member0Type)
2116 report_fatal_error(reason: "Operand type must match the struct member type");
2117
2118 if (IsSret) {
2119 buildSRetInst(Opcode, SRetReg: Call->Arguments[0], Op1Reg, Op2Reg, RetType,
2120 MIRBuilder, GR);
2121 } else {
2122 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2123 Register ResReg = Call->ReturnRegister;
2124 if (const TargetRegisterClass *DstRC = MRI->getRegClassOrNull(Reg: Op1Reg)) {
2125 MRI->setRegClass(Reg: ResReg, RC: DstRC);
2126 }
2127 GR->assignSPIRVTypeToVReg(Type: RetType, VReg: ResReg, MF: MIRBuilder.getMF());
2128 MIRBuilder.buildInstr(Opcode)
2129 .addDef(RegNo: ResReg)
2130 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: RetType))
2131 .addUse(RegNo: Op1Reg)
2132 .addUse(RegNo: Op2Reg);
2133 }
2134 return true;
2135}
2136
2137static bool generateArithmeticInst(const SPIRV::IncomingCall *Call,
2138 MachineIRBuilder &MIRBuilder,
2139 SPIRVGlobalRegistry *GR) {
2140 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2141 unsigned Opcode =
2142 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2143
2144 auto MIB = MIRBuilder.buildInstr(Opcode)
2145 .addDef(RegNo: Call->ReturnRegister)
2146 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
2147 for (Register Arg : Call->Arguments)
2148 MIB.addUse(RegNo: Arg);
2149 return true;
2150}
2151
2152static bool generateGetQueryInst(const SPIRV::IncomingCall *Call,
2153 MachineIRBuilder &MIRBuilder,
2154 SPIRVGlobalRegistry *GR) {
2155 // Lookup the builtin record.
2156 SPIRV::BuiltIn::BuiltIn Value =
2157 SPIRV::lookupGetBuiltin(Name: Call->Builtin->name(), Set: Call->Builtin->Set)->Value;
2158 const bool IsDefaultOne = (Value == SPIRV::BuiltIn::GlobalSize ||
2159 Value == SPIRV::BuiltIn::NumWorkgroups ||
2160 Value == SPIRV::BuiltIn::WorkgroupSize ||
2161 Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
2162 return genWorkgroupQuery(Call, MIRBuilder, GR, BuiltinValue: Value, DefaultValue: IsDefaultOne ? 1 : 0);
2163}
2164
2165static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call,
2166 MachineIRBuilder &MIRBuilder,
2167 SPIRVGlobalRegistry *GR) {
2168 // Lookup the image size query component number in the TableGen records.
2169 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2170 uint32_t Component =
2171 SPIRV::lookupImageQueryBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Component;
2172 // Query result may either be a vector or a scalar. If return type is not a
2173 // vector, expect only a single size component. Otherwise get the number of
2174 // expected components.
2175 unsigned NumExpectedRetComponents =
2176 GR->getScalarOrVectorComponentCount(Type: Call->ReturnType);
2177 // Get the actual number of query result/size components.
2178 SPIRVTypeInst ImgType = GR->getSPIRVTypeForVReg(VReg: Call->Arguments[0]);
2179 unsigned NumActualRetComponents = getNumSizeComponents(imgType: ImgType);
2180 Register QueryResult = Call->ReturnRegister;
2181 SPIRVTypeInst QueryResultType = Call->ReturnType;
2182 if (NumExpectedRetComponents != NumActualRetComponents) {
2183 unsigned Bitwidth = Call->ReturnType->getOpcode() == SPIRV::OpTypeInt
2184 ? Call->ReturnType->getOperand(i: 1).getImm()
2185 : 32;
2186 QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
2187 Ty: LLT::fixed_vector(NumElements: NumActualRetComponents, ScalarSizeInBits: Bitwidth));
2188 MIRBuilder.getMRI()->setRegClass(Reg: QueryResult, RC: &SPIRV::viIDRegClass);
2189 SPIRVTypeInst IntTy = GR->getOrCreateSPIRVIntegerType(BitWidth: Bitwidth, MIRBuilder);
2190 QueryResultType = GR->getOrCreateSPIRVVectorType(
2191 BaseType: IntTy, NumElements: NumActualRetComponents, MIRBuilder, EmitIR: true);
2192 GR->assignSPIRVTypeToVReg(Type: QueryResultType, VReg: QueryResult, MF: MIRBuilder.getMF());
2193 }
2194 bool IsDimBuf = ImgType->getOperand(i: 2).getImm() == SPIRV::Dim::DIM_Buffer;
2195 bool IsMultisampled = ImgType->getOperand(i: 5).getImm() != 0;
2196 bool UseQuerySize = IsDimBuf || IsMultisampled;
2197 unsigned Opcode =
2198 UseQuerySize ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
2199 auto MIB = MIRBuilder.buildInstr(Opcode)
2200 .addDef(RegNo: QueryResult)
2201 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: QueryResultType))
2202 .addUse(RegNo: Call->Arguments[0]);
2203 if (!UseQuerySize)
2204 MIB.addUse(RegNo: buildConstantIntReg32(Val: 0, MIRBuilder, GR)); // Lod id.
2205 if (NumExpectedRetComponents == NumActualRetComponents)
2206 return true;
2207 if (NumExpectedRetComponents == 1) {
2208 // Only 1 component is expected, build OpCompositeExtract instruction.
2209 unsigned ExtractedComposite =
2210 Component == 3 ? NumActualRetComponents - 1 : Component;
2211 assert(ExtractedComposite < NumActualRetComponents &&
2212 "Invalid composite index!");
2213 Register TypeReg = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
2214 SPIRVTypeInst NewType = nullptr;
2215 if (QueryResultType->getOpcode() == SPIRV::OpTypeVector) {
2216 NewType = GR->getScalarOrVectorComponentType(Type: QueryResultType);
2217 Register NewTypeReg = GR->getSPIRVTypeID(SpirvType: NewType);
2218 if (TypeReg != NewTypeReg)
2219 TypeReg = NewTypeReg;
2220 else
2221 NewType = nullptr;
2222 }
2223 MIRBuilder.buildInstr(Opcode: SPIRV::OpCompositeExtract)
2224 .addDef(RegNo: Call->ReturnRegister)
2225 .addUse(RegNo: TypeReg)
2226 .addUse(RegNo: QueryResult)
2227 .addImm(Val: ExtractedComposite);
2228 if (NewType)
2229 updateRegType(Reg: Call->ReturnRegister, Ty: nullptr, SpirvTy: NewType, GR, MIB&: MIRBuilder,
2230 MRI&: MIRBuilder.getMF().getRegInfo());
2231 } else {
2232 // More than 1 component is expected, fill a new vector.
2233 auto MIB = MIRBuilder.buildInstr(Opcode: SPIRV::OpVectorShuffle)
2234 .addDef(RegNo: Call->ReturnRegister)
2235 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
2236 .addUse(RegNo: QueryResult)
2237 .addUse(RegNo: QueryResult);
2238 for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
2239 MIB.addImm(Val: i < NumActualRetComponents ? i : 0xffffffff);
2240 }
2241 return true;
2242}
2243
2244static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call,
2245 MachineIRBuilder &MIRBuilder,
2246 SPIRVGlobalRegistry *GR) {
2247 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
2248 "Image samples query result must be of int type!");
2249
2250 // Lookup the instruction opcode in the TableGen records.
2251 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2252 unsigned Opcode =
2253 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2254
2255 Register Image = Call->Arguments[0];
2256 SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
2257 GR->getSPIRVTypeForVReg(VReg: Image)->getOperand(i: 2).getImm());
2258 (void)ImageDimensionality;
2259
2260 switch (Opcode) {
2261 case SPIRV::OpImageQuerySamples:
2262 assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
2263 "Image must be of 2D dimensionality");
2264 break;
2265 case SPIRV::OpImageQueryLevels:
2266 assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
2267 ImageDimensionality == SPIRV::Dim::DIM_2D ||
2268 ImageDimensionality == SPIRV::Dim::DIM_3D ||
2269 ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
2270 "Image must be of 1D/2D/3D/Cube dimensionality");
2271 break;
2272 }
2273
2274 MIRBuilder.buildInstr(Opcode)
2275 .addDef(RegNo: Call->ReturnRegister)
2276 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
2277 .addUse(RegNo: Image);
2278 return true;
2279}
2280
2281// TODO: Move to TableGen.
2282static SPIRV::SamplerAddressingMode::SamplerAddressingMode
2283getSamplerAddressingModeFromBitmask(unsigned Bitmask) {
2284 switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
2285 case SPIRV::CLK_ADDRESS_CLAMP:
2286 return SPIRV::SamplerAddressingMode::Clamp;
2287 case SPIRV::CLK_ADDRESS_CLAMP_TO_EDGE:
2288 return SPIRV::SamplerAddressingMode::ClampToEdge;
2289 case SPIRV::CLK_ADDRESS_REPEAT:
2290 return SPIRV::SamplerAddressingMode::Repeat;
2291 case SPIRV::CLK_ADDRESS_MIRRORED_REPEAT:
2292 return SPIRV::SamplerAddressingMode::RepeatMirrored;
2293 case SPIRV::CLK_ADDRESS_NONE:
2294 return SPIRV::SamplerAddressingMode::None;
2295 default:
2296 report_fatal_error(reason: "Unknown CL address mode");
2297 }
2298}
2299
2300static unsigned getSamplerParamFromBitmask(unsigned Bitmask) {
2301 return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;
2302}
2303
2304static SPIRV::SamplerFilterMode::SamplerFilterMode
2305getSamplerFilterModeFromBitmask(unsigned Bitmask) {
2306 if (Bitmask & SPIRV::CLK_FILTER_LINEAR)
2307 return SPIRV::SamplerFilterMode::Linear;
2308 if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
2309 return SPIRV::SamplerFilterMode::Nearest;
2310 return SPIRV::SamplerFilterMode::Nearest;
2311}
2312
2313static bool generateReadImageInst(StringRef DemangledCall,
2314 const SPIRV::IncomingCall *Call,
2315 MachineIRBuilder &MIRBuilder,
2316 SPIRVGlobalRegistry *GR) {
2317 if (Call->isSpirvOp())
2318 return buildOpFromWrapper(MIRBuilder, Opcode: SPIRV::OpImageRead, Call,
2319 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
2320 Register Image = Call->Arguments[0];
2321 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2322 bool HasOclSampler = DemangledCall.contains_insensitive(Other: "ocl_sampler");
2323 bool HasMsaa = DemangledCall.contains_insensitive(Other: "msaa");
2324 if (HasOclSampler) {
2325 Register Sampler = Call->Arguments[1];
2326
2327 if (!GR->isScalarOfType(VReg: Sampler, TypeOpcode: SPIRV::OpTypeSampler) &&
2328 getDefInstrMaybeConstant(ConstReg&: Sampler, MRI)->getOperand(i: 1).isCImm()) {
2329 uint64_t SamplerMask = getIConstVal(ConstReg: Sampler, MRI);
2330 Sampler = GR->buildConstantSampler(
2331 Res: Register(), AddrMode: getSamplerAddressingModeFromBitmask(Bitmask: SamplerMask),
2332 Param: getSamplerParamFromBitmask(Bitmask: SamplerMask),
2333 FilerMode: getSamplerFilterModeFromBitmask(Bitmask: SamplerMask), MIRBuilder);
2334 }
2335 SPIRVTypeInst ImageType = GR->getSPIRVTypeForVReg(VReg: Image);
2336 SPIRVTypeInst SampledImageType =
2337 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
2338 Register SampledImage = MRI->createVirtualRegister(RegClass: &SPIRV::iIDRegClass);
2339
2340 MIRBuilder.buildInstr(Opcode: SPIRV::OpSampledImage)
2341 .addDef(RegNo: SampledImage)
2342 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: SampledImageType))
2343 .addUse(RegNo: Image)
2344 .addUse(RegNo: Sampler);
2345
2346 Register Lod = GR->buildConstantFP(Val: APFloat::getZero(Sem: APFloat::IEEEsingle()),
2347 MIRBuilder);
2348
2349 if (Call->ReturnType->getOpcode() != SPIRV::OpTypeVector) {
2350 SPIRVTypeInst TempType =
2351 GR->getOrCreateSPIRVVectorType(BaseType: Call->ReturnType, NumElements: 4, MIRBuilder, EmitIR: true);
2352 Register TempRegister =
2353 MRI->createGenericVirtualRegister(Ty: GR->getRegType(SpvType: TempType));
2354 MRI->setRegClass(Reg: TempRegister, RC: GR->getRegClass(SpvType: TempType));
2355 GR->assignSPIRVTypeToVReg(Type: TempType, VReg: TempRegister, MF: MIRBuilder.getMF());
2356 MIRBuilder.buildInstr(Opcode: SPIRV::OpImageSampleExplicitLod)
2357 .addDef(RegNo: TempRegister)
2358 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: TempType))
2359 .addUse(RegNo: SampledImage)
2360 .addUse(RegNo: Call->Arguments[2]) // Coordinate.
2361 .addImm(Val: SPIRV::ImageOperand::Lod)
2362 .addUse(RegNo: Lod);
2363 MIRBuilder.buildInstr(Opcode: SPIRV::OpCompositeExtract)
2364 .addDef(RegNo: Call->ReturnRegister)
2365 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
2366 .addUse(RegNo: TempRegister)
2367 .addImm(Val: 0);
2368 } else {
2369 MIRBuilder.buildInstr(Opcode: SPIRV::OpImageSampleExplicitLod)
2370 .addDef(RegNo: Call->ReturnRegister)
2371 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
2372 .addUse(RegNo: SampledImage)
2373 .addUse(RegNo: Call->Arguments[2]) // Coordinate.
2374 .addImm(Val: SPIRV::ImageOperand::Lod)
2375 .addUse(RegNo: Lod);
2376 }
2377 } else if (HasMsaa) {
2378 MIRBuilder.buildInstr(Opcode: SPIRV::OpImageRead)
2379 .addDef(RegNo: Call->ReturnRegister)
2380 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
2381 .addUse(RegNo: Image)
2382 .addUse(RegNo: Call->Arguments[1]) // Coordinate.
2383 .addImm(Val: SPIRV::ImageOperand::Sample)
2384 .addUse(RegNo: Call->Arguments[2]);
2385 } else {
2386 MIRBuilder.buildInstr(Opcode: SPIRV::OpImageRead)
2387 .addDef(RegNo: Call->ReturnRegister)
2388 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
2389 .addUse(RegNo: Image)
2390 .addUse(RegNo: Call->Arguments[1]); // Coordinate.
2391 }
2392 return true;
2393}
2394
2395static bool generateWriteImageInst(const SPIRV::IncomingCall *Call,
2396 MachineIRBuilder &MIRBuilder,
2397 SPIRVGlobalRegistry *GR) {
2398 if (Call->isSpirvOp())
2399 return buildOpFromWrapper(MIRBuilder, Opcode: SPIRV::OpImageWrite, Call,
2400 TypeReg: Register(0));
2401 MIRBuilder.buildInstr(Opcode: SPIRV::OpImageWrite)
2402 .addUse(RegNo: Call->Arguments[0]) // Image.
2403 .addUse(RegNo: Call->Arguments[1]) // Coordinate.
2404 .addUse(RegNo: Call->Arguments[2]); // Texel.
2405 return true;
2406}
2407
2408static bool generateSampleImageInst(StringRef DemangledCall,
2409 const SPIRV::IncomingCall *Call,
2410 MachineIRBuilder &MIRBuilder,
2411 SPIRVGlobalRegistry *GR) {
2412 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2413 if (Call->Builtin->name().contains_insensitive(
2414 Other: "__translate_sampler_initializer")) {
2415 // Build sampler literal.
2416 uint64_t Bitmask = getIConstVal(ConstReg: Call->Arguments[0], MRI);
2417 Register Sampler = GR->buildConstantSampler(
2418 Res: Call->ReturnRegister, AddrMode: getSamplerAddressingModeFromBitmask(Bitmask),
2419 Param: getSamplerParamFromBitmask(Bitmask),
2420 FilerMode: getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder);
2421 return Sampler.isValid();
2422 } else if (Call->Builtin->name().contains_insensitive(
2423 Other: "__spirv_SampledImage")) {
2424 // Create OpSampledImage.
2425 Register Image = Call->Arguments[0];
2426 SPIRVTypeInst ImageType = GR->getSPIRVTypeForVReg(VReg: Image);
2427 SPIRVTypeInst SampledImageType =
2428 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
2429 Register SampledImage =
2430 Call->ReturnRegister.isValid()
2431 ? Call->ReturnRegister
2432 : MRI->createVirtualRegister(RegClass: &SPIRV::iIDRegClass);
2433 MIRBuilder.buildInstr(Opcode: SPIRV::OpSampledImage)
2434 .addDef(RegNo: SampledImage)
2435 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: SampledImageType))
2436 .addUse(RegNo: Image)
2437 .addUse(RegNo: Call->Arguments[1]); // Sampler.
2438 return true;
2439 } else if (Call->Builtin->name().contains_insensitive(
2440 Other: "__spirv_ImageSampleExplicitLod")) {
2441 // Sample an image using an explicit level of detail.
2442 std::string ReturnType = DemangledCall.str();
2443 if (DemangledCall.contains(Other: "_R")) {
2444 ReturnType = ReturnType.substr(pos: ReturnType.find(s: "_R") + 2);
2445 ReturnType = ReturnType.substr(pos: 0, n: ReturnType.find(c: '('));
2446 }
2447 SPIRVTypeInst Type = Call->ReturnType
2448 ? Call->ReturnType
2449 : SPIRVTypeInst(GR->getOrCreateSPIRVTypeByName(
2450 TypeStr: ReturnType, MIRBuilder, EmitIR: true));
2451 if (!Type) {
2452 std::string DiagMsg =
2453 "Unable to recognize SPIRV type name: " + ReturnType;
2454 report_fatal_error(reason: DiagMsg.c_str());
2455 }
2456 MIRBuilder.buildInstr(Opcode: SPIRV::OpImageSampleExplicitLod)
2457 .addDef(RegNo: Call->ReturnRegister)
2458 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Type))
2459 .addUse(RegNo: Call->Arguments[0]) // Image.
2460 .addUse(RegNo: Call->Arguments[1]) // Coordinate.
2461 .addImm(Val: SPIRV::ImageOperand::Lod)
2462 .addUse(RegNo: Call->Arguments[3]);
2463 return true;
2464 }
2465 return false;
2466}
2467
2468static bool generateSelectInst(const SPIRV::IncomingCall *Call,
2469 MachineIRBuilder &MIRBuilder) {
2470 const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2471 LLT ResTy = MRI->getType(Reg: Call->ReturnRegister);
2472 LLT CondTy = MRI->getType(Reg: Call->Arguments[0]);
2473 if (!ResTy.isVector() && CondTy.isVector())
2474 report_fatal_error(reason: "OpSelect with a scalar result requires a scalar "
2475 "boolean condition");
2476 MIRBuilder.buildSelect(Res: Call->ReturnRegister, Tst: Call->Arguments[0],
2477 Op0: Call->Arguments[1], Op1: Call->Arguments[2]);
2478 return true;
2479}
2480
2481static bool generateConstructInst(const SPIRV::IncomingCall *Call,
2482 MachineIRBuilder &MIRBuilder,
2483 SPIRVGlobalRegistry *GR) {
2484 createContinuedInstructions(MIRBuilder, Opcode: SPIRV::OpCompositeConstruct, MinWC: 3,
2485 ContinuedOpcode: SPIRV::OpCompositeConstructContinuedINTEL,
2486 Args: Call->Arguments, ReturnRegister: Call->ReturnRegister,
2487 TypeID: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
2488 return true;
2489}
2490
2491static bool generateCoopMatrInst(const SPIRV::IncomingCall *Call,
2492 MachineIRBuilder &MIRBuilder,
2493 SPIRVGlobalRegistry *GR) {
2494 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2495 unsigned Opcode =
2496 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2497 bool IsSet = Opcode != SPIRV::OpCooperativeMatrixStoreKHR &&
2498 Opcode != SPIRV::OpCooperativeMatrixStoreCheckedINTEL &&
2499 Opcode != SPIRV::OpCooperativeMatrixPrefetchINTEL;
2500 unsigned ArgSz = Call->Arguments.size();
2501 unsigned LiteralIdx = 0;
2502 switch (Opcode) {
2503 // Memory operand is optional and is literal.
2504 case SPIRV::OpCooperativeMatrixLoadKHR:
2505 LiteralIdx = ArgSz > 3 ? 3 : 0;
2506 break;
2507 case SPIRV::OpCooperativeMatrixStoreKHR:
2508 LiteralIdx = ArgSz > 4 ? 4 : 0;
2509 break;
2510 case SPIRV::OpCooperativeMatrixLoadCheckedINTEL:
2511 LiteralIdx = ArgSz > 7 ? 7 : 0;
2512 break;
2513 case SPIRV::OpCooperativeMatrixStoreCheckedINTEL:
2514 LiteralIdx = ArgSz > 8 ? 8 : 0;
2515 break;
2516 // Cooperative Matrix Operands operand is optional and is literal.
2517 case SPIRV::OpCooperativeMatrixMulAddKHR:
2518 LiteralIdx = ArgSz > 3 ? 3 : 0;
2519 break;
2520 };
2521
2522 SmallVector<uint32_t, 1> ImmArgs;
2523 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2524 if (Opcode == SPIRV::OpCooperativeMatrixPrefetchINTEL) {
2525 const uint32_t CacheLevel = getIConstVal(ConstReg: Call->Arguments[3], MRI);
2526 auto MIB = MIRBuilder.buildInstr(Opcode: SPIRV::OpCooperativeMatrixPrefetchINTEL)
2527 .addUse(RegNo: Call->Arguments[0]) // pointer
2528 .addUse(RegNo: Call->Arguments[1]) // rows
2529 .addUse(RegNo: Call->Arguments[2]) // columns
2530 .addImm(Val: CacheLevel) // cache level
2531 .addUse(RegNo: Call->Arguments[4]); // memory layout
2532 if (ArgSz > 5)
2533 MIB.addUse(RegNo: Call->Arguments[5]); // stride
2534 if (ArgSz > 6) {
2535 const uint32_t MemOp = getIConstVal(ConstReg: Call->Arguments[6], MRI);
2536 MIB.addImm(Val: MemOp); // memory operand
2537 }
2538 return true;
2539 }
2540 if (LiteralIdx > 0)
2541 ImmArgs.push_back(Elt: getIConstVal(ConstReg: Call->Arguments[LiteralIdx], MRI));
2542 Register TypeReg = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
2543 if (Opcode == SPIRV::OpCooperativeMatrixLengthKHR) {
2544 SPIRVTypeInst CoopMatrType = GR->getSPIRVTypeForVReg(VReg: Call->Arguments[0]);
2545 if (!CoopMatrType)
2546 report_fatal_error(reason: "Can't find a register's type definition");
2547 MIRBuilder.buildInstr(Opcode)
2548 .addDef(RegNo: Call->ReturnRegister)
2549 .addUse(RegNo: TypeReg)
2550 .addUse(RegNo: CoopMatrType->getOperand(i: 0).getReg());
2551 return true;
2552 }
2553 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2554 TypeReg: IsSet ? TypeReg : Register(0), ImmArgs);
2555}
2556
2557static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call,
2558 MachineIRBuilder &MIRBuilder,
2559 SPIRVGlobalRegistry *GR) {
2560 // Lookup the instruction opcode in the TableGen records.
2561 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2562 unsigned Opcode =
2563 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2564 const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2565
2566 switch (Opcode) {
2567 case SPIRV::OpSpecConstant: {
2568 // Determine the constant MI.
2569 Register ConstRegister = Call->Arguments[1];
2570 const MachineInstr *Const = getDefInstrMaybeConstant(ConstReg&: ConstRegister, MRI);
2571 assert(Const &&
2572 (Const->getOpcode() == TargetOpcode::G_CONSTANT ||
2573 Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&
2574 "Argument should be either an int or floating-point constant");
2575 // Determine the opcode and built the OpSpec MI.
2576 const MachineOperand &ConstOperand = Const->getOperand(i: 1);
2577 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
2578 assert(ConstOperand.isCImm() && "Int constant operand is expected");
2579 Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
2580 ? SPIRV::OpSpecConstantTrue
2581 : SPIRV::OpSpecConstantFalse;
2582 }
2583 auto MIB = MIRBuilder.buildInstr(Opcode)
2584 .addDef(RegNo: Call->ReturnRegister)
2585 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
2586
2587 if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
2588 if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
2589 addNumImm(Imm: ConstOperand.getCImm()->getValue(), MIB);
2590 else
2591 addNumImm(Imm: ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
2592 }
2593 // Build the SpecID decoration.
2594 unsigned SpecId =
2595 static_cast<unsigned>(getIConstVal(ConstReg: Call->Arguments[0], MRI));
2596 buildOpDecorate(Reg: Call->ReturnRegister, MIRBuilder, Dec: SPIRV::Decoration::SpecId,
2597 DecArgs: {SpecId});
2598 return true;
2599 }
2600 case SPIRV::OpSpecConstantComposite: {
2601 createContinuedInstructions(MIRBuilder, Opcode, MinWC: 3,
2602 ContinuedOpcode: SPIRV::OpSpecConstantCompositeContinuedINTEL,
2603 Args: Call->Arguments, ReturnRegister: Call->ReturnRegister,
2604 TypeID: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
2605 return true;
2606 }
2607 default:
2608 return false;
2609 }
2610}
2611
2612static bool generateExtendedBitOpsInst(const SPIRV::IncomingCall *Call,
2613 MachineIRBuilder &MIRBuilder,
2614 SPIRVGlobalRegistry *GR) {
2615 // Lookup the instruction opcode in the TableGen records.
2616 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2617 unsigned Opcode =
2618 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2619
2620 return buildExtendedBitOpsInst(Call, Opcode, MIRBuilder, GR);
2621}
2622
2623static bool generateBindlessImageINTELInst(const SPIRV::IncomingCall *Call,
2624 MachineIRBuilder &MIRBuilder,
2625 SPIRVGlobalRegistry *GR) {
2626 // Lookup the instruction opcode in the TableGen records.
2627 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2628 unsigned Opcode =
2629 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2630
2631 return buildBindlessImageINTELInst(Call, Opcode, MIRBuilder, GR);
2632}
2633
2634static bool generateBlockingPipesInst(const SPIRV::IncomingCall *Call,
2635 MachineIRBuilder &MIRBuilder,
2636 SPIRVGlobalRegistry *GR) {
2637 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2638 unsigned Opcode =
2639 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2640 return buildOpFromWrapper(MIRBuilder, Opcode, Call, TypeReg: Register(0));
2641}
2642
2643static bool buildAPFixedPointInst(const SPIRV::IncomingCall *Call,
2644 unsigned Opcode, MachineIRBuilder &MIRBuilder,
2645 SPIRVGlobalRegistry *GR) {
2646 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2647 SmallVector<uint32_t, 1> ImmArgs;
2648 Register InputReg = Call->Arguments[0];
2649 const Type *RetTy = GR->getTypeForSPIRVType(Ty: Call->ReturnType);
2650 bool IsSRet = RetTy->isVoidTy();
2651
2652 if (IsSRet) {
2653 const LLT ValTy = MRI->getType(Reg: InputReg);
2654 Register ActualRetValReg = MRI->createGenericVirtualRegister(Ty: ValTy);
2655 SPIRVTypeInst InstructionType =
2656 GR->getPointeeType(PtrType: GR->getSPIRVTypeForVReg(VReg: InputReg));
2657 InputReg = Call->Arguments[1];
2658 auto InputType = GR->getTypeForSPIRVType(Ty: GR->getSPIRVTypeForVReg(VReg: InputReg));
2659 Register PtrInputReg;
2660 if (InputType->getTypeID() == llvm::Type::TypeID::TypedPointerTyID) {
2661 LLT InputLLT = MRI->getType(Reg: InputReg);
2662 PtrInputReg = MRI->createGenericVirtualRegister(Ty: InputLLT);
2663 SPIRVTypeInst PtrType =
2664 GR->getPointeeType(PtrType: GR->getSPIRVTypeForVReg(VReg: InputReg));
2665 MachineMemOperand *MMO1 = MIRBuilder.getMF().getMachineMemOperand(
2666 PtrInfo: MachinePointerInfo(), F: MachineMemOperand::MOLoad,
2667 Size: InputLLT.getSizeInBytes(), BaseAlignment: Align(4));
2668 MIRBuilder.buildLoad(Res: PtrInputReg, Addr: InputReg, MMO&: *MMO1);
2669 MRI->setRegClass(Reg: PtrInputReg, RC: &SPIRV::iIDRegClass);
2670 GR->assignSPIRVTypeToVReg(Type: PtrType, VReg: PtrInputReg, MF: MIRBuilder.getMF());
2671 }
2672
2673 for (unsigned index = 2; index < 7; index++) {
2674 ImmArgs.push_back(Elt: getIConstVal(ConstReg: Call->Arguments[index], MRI));
2675 }
2676
2677 // Emit the instruction
2678 auto MIB = MIRBuilder.buildInstr(Opcode)
2679 .addDef(RegNo: ActualRetValReg)
2680 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: InstructionType));
2681 if (PtrInputReg)
2682 MIB.addUse(RegNo: PtrInputReg);
2683 else
2684 MIB.addUse(RegNo: InputReg);
2685
2686 for (uint32_t Imm : ImmArgs)
2687 MIB.addImm(Val: Imm);
2688 unsigned Size = ValTy.getSizeInBytes();
2689 // Store result to the pointer passed in Arg[0]
2690 MachineMemOperand *MMO = MIRBuilder.getMF().getMachineMemOperand(
2691 PtrInfo: MachinePointerInfo(), F: MachineMemOperand::MOStore, Size, BaseAlignment: Align(4));
2692 MRI->setRegClass(Reg: ActualRetValReg, RC: &SPIRV::pIDRegClass);
2693 MIRBuilder.buildStore(Val: ActualRetValReg, Addr: Call->Arguments[0], MMO&: *MMO);
2694 return true;
2695 } else {
2696 for (unsigned index = 1; index < 6; index++)
2697 ImmArgs.push_back(Elt: getIConstVal(ConstReg: Call->Arguments[index], MRI));
2698
2699 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2700 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType), ImmArgs);
2701 }
2702}
2703
2704static bool generateAPFixedPointInst(const SPIRV::IncomingCall *Call,
2705 MachineIRBuilder &MIRBuilder,
2706 SPIRVGlobalRegistry *GR) {
2707 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2708 unsigned Opcode =
2709 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2710
2711 return buildAPFixedPointInst(Call, Opcode, MIRBuilder, GR);
2712}
2713
2714static bool
2715generateTernaryBitwiseFunctionINTELInst(const SPIRV::IncomingCall *Call,
2716 MachineIRBuilder &MIRBuilder,
2717 SPIRVGlobalRegistry *GR) {
2718 // Lookup the instruction opcode in the TableGen records.
2719 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2720 unsigned Opcode =
2721 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2722
2723 return buildTernaryBitwiseFunctionINTELInst(Call, Opcode, MIRBuilder, GR);
2724}
2725
2726static bool generateImageChannelDataTypeInst(const SPIRV::IncomingCall *Call,
2727 MachineIRBuilder &MIRBuilder,
2728 SPIRVGlobalRegistry *GR) {
2729 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2730 unsigned Opcode =
2731 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2732
2733 return buildImageChannelDataTypeInst(Call, Opcode, MIRBuilder, GR);
2734}
2735
2736static bool generate2DBlockIOINTELInst(const SPIRV::IncomingCall *Call,
2737 MachineIRBuilder &MIRBuilder,
2738 SPIRVGlobalRegistry *GR) {
2739 // Lookup the instruction opcode in the TableGen records.
2740 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2741 unsigned Opcode =
2742 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2743
2744 return build2DBlockIOINTELInst(Call, Opcode, MIRBuilder, GR);
2745}
2746
2747static bool generatePipeInst(const SPIRV::IncomingCall *Call,
2748 MachineIRBuilder &MIRBuilder,
2749 SPIRVGlobalRegistry *GR) {
2750 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2751 unsigned Opcode =
2752 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2753
2754 unsigned Scope = SPIRV::Scope::Workgroup;
2755 if (Builtin->name().contains(Other: "sub_group"))
2756 Scope = SPIRV::Scope::Subgroup;
2757
2758 return buildPipeInst(Call, Opcode, Scope, MIRBuilder, GR);
2759}
2760
2761static bool generatePredicatedLoadStoreInst(const SPIRV::IncomingCall *Call,
2762 MachineIRBuilder &MIRBuilder,
2763 SPIRVGlobalRegistry *GR) {
2764 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2765 unsigned Opcode =
2766 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
2767
2768 bool IsSet = Opcode != SPIRV::OpPredicatedStoreINTEL;
2769 unsigned ArgSz = Call->Arguments.size();
2770 SmallVector<uint32_t, 1> ImmArgs;
2771 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2772 // Memory operand is optional and is literal.
2773 if (ArgSz > 3)
2774 ImmArgs.push_back(Elt: getIConstVal(ConstReg: Call->Arguments[/*Literal index*/ 3], MRI));
2775
2776 Register TypeReg = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
2777 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2778 TypeReg: IsSet ? TypeReg : Register(0), ImmArgs);
2779}
2780
2781static bool buildNDRange(const SPIRV::IncomingCall *Call,
2782 MachineIRBuilder &MIRBuilder,
2783 SPIRVGlobalRegistry *GR) {
2784 // The OpenCL ndrange_*D functions are overloaded and support 1D, 2D, and 3D
2785 // variants, accepting 1 to 3 arguments:
2786 // (global_work_size)
2787 // (global_work_size, local_work_size)
2788 // (global_work_offset, global_work_size, local_work_size)
2789 // Note: When all three arguments are provided, they are reordered compared
2790 // to the one- or two-argument form.
2791 //
2792 // The function may return data through an sret argument at position 0 (with
2793 // a void function return type). When present, all other argument indices are
2794 // adjusted accordingly.
2795 //
2796 // SPIR-V's OpBuildNDRange requires all three arguments (GlobalWorkSize,
2797 // LocalWorkSize, GlobalWorkOffset). For 1D kernels, the values are scalars;
2798 // for 2D/3D kernels, they are arrays of 2 or 3 elements. Missing arguments
2799 // default to zero.
2800 //
2801 // Calculate argument indices based on the number of arguments and presence
2802 // of sret:
2803 const unsigned NumCallArgs = Call->Arguments.size();
2804 const unsigned MaxCallArgs = Call->Builtin->MaxNumArgs;
2805 const unsigned IncorrectArgIdx = MaxCallArgs + 1;
2806
2807 const Type *RetTy = GR->getTypeForSPIRVType(Ty: Call->ReturnType);
2808 bool HasSRetArg = RetTy->isVoidTy();
2809
2810 const unsigned SRetArgIdx = HasSRetArg ? 0 : IncorrectArgIdx;
2811 const unsigned ArgBase = HasSRetArg ? 1 : 0;
2812 const unsigned MaxNDRangeArgs = 3;
2813 const unsigned NumNDRangeArgs = NumCallArgs - ArgBase;
2814
2815 const unsigned GlobalWorkSizeArgIdx =
2816 NumNDRangeArgs < MaxNDRangeArgs ? ArgBase : ArgBase + 1;
2817 const unsigned LocalWorkSizeArgIdx =
2818 (NumNDRangeArgs == 1)
2819 ? IncorrectArgIdx
2820 : (NumNDRangeArgs == MaxNDRangeArgs ? ArgBase + 2 : ArgBase + 1);
2821 const unsigned GlobalWorkOffsetArgIdx =
2822 NumNDRangeArgs == MaxNDRangeArgs ? ArgBase : IncorrectArgIdx;
2823
2824 // Each nd_range field is an array of <Dimension> integers matching the
2825 // address model width (32 or 64 bits).
2826 const unsigned AddressModelBits = GR->getPointerSize();
2827 assert(AddressModelBits == 64 || AddressModelBits == 32);
2828
2829 // The dimension is encoded in the function name as "ndrange_XD" where X is
2830 // 1, 2, or 3.
2831 unsigned Dimension = 0;
2832 Call->Builtin->name().substr(Start: 8, N: 1).getAsInteger(Radix: 10, Result&: Dimension);
2833 assert(Dimension <= 3 && Dimension >= 1);
2834
2835 // Determine the work size type based on the dimension. For missing arguments,
2836 // create a zero constant of the appropriate type.
2837 MachineFunction &MF = MIRBuilder.getMF();
2838 SPIRVTypeInst SpvFieldTy;
2839 Register ConstZero;
2840 if (Dimension == 1) {
2841 SpvFieldTy = GR->getSPIRVTypeForVReg(VReg: Call->Arguments[GlobalWorkSizeArgIdx]);
2842 assert(SpvFieldTy && SpvFieldTy->getOpcode() == SPIRV::OpTypeInt &&
2843 "Expected scalar integer type");
2844
2845 if (NumNDRangeArgs < MaxNDRangeArgs)
2846 ConstZero = GR->buildConstantInt(Val: 0, MIRBuilder, SpvType: SpvFieldTy, EmitIR: true);
2847 } else {
2848 Type *BaseTy =
2849 IntegerType::get(C&: MF.getFunction().getContext(), NumBits: AddressModelBits);
2850 Type *FieldTy = ArrayType::get(ElementType: BaseTy, NumElements: Dimension);
2851 SpvFieldTy = GR->getOrCreateSPIRVType(
2852 Type: FieldTy, MIRBuilder, AQ: SPIRV::AccessQualifier::ReadOnly, EmitIR: true);
2853
2854 if (NumNDRangeArgs < MaxNDRangeArgs) {
2855 auto InsertIt = MIRBuilder.getInsertPt();
2856 MachineBasicBlock &MBB = MIRBuilder.getMBB();
2857 MachineInstr &InsertMI = (InsertIt != MBB.end()) ? *InsertIt : MBB.back();
2858 const SPIRVSubtarget &ST = cast<SPIRVSubtarget>(Val: MF.getSubtarget());
2859 ConstZero = GR->getOrCreateConstIntArray(Val: 0, Num: Dimension, I&: InsertMI,
2860 SpvType: SpvFieldTy, TII: *ST.getInstrInfo());
2861 }
2862 }
2863
2864 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2865
2866 auto CreateDataRegister = [&](unsigned Idx) -> Register {
2867 Register Reg = (Idx == IncorrectArgIdx) ? ConstZero : Call->Arguments[Idx];
2868
2869 if (GR->getSPIRVTypeForVReg(VReg: Reg) == SpvFieldTy) {
2870 // Already has the correct type.
2871 return Reg;
2872 }
2873
2874 assert(GR->getSPIRVTypeForVReg(Reg)->getOpcode() == SPIRV::OpTypePointer &&
2875 "Only pointer types are supported for loading values");
2876
2877 Register Ptr = Reg;
2878
2879 Reg = MRI->createVirtualRegister(RegClass: &SPIRV::iIDRegClass);
2880 GR->assignSPIRVTypeToVReg(Type: SpvFieldTy, VReg: Reg, MF);
2881
2882 MIRBuilder.buildInstr(Opcode: SPIRV::OpLoad)
2883 .addDef(RegNo: Reg)
2884 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: SpvFieldTy))
2885 .addUse(RegNo: Ptr);
2886 return Reg;
2887 };
2888
2889 Register GlobalWorkSize = CreateDataRegister(GlobalWorkSizeArgIdx);
2890 Register LocalWorkSize = CreateDataRegister(LocalWorkSizeArgIdx);
2891 Register GlobalWorkOffset = CreateDataRegister(GlobalWorkOffsetArgIdx);
2892
2893 if (!HasSRetArg) {
2894 return MIRBuilder.buildInstr(Opcode: SPIRV::OpBuildNDRange)
2895 .addDef(RegNo: Call->ReturnRegister)
2896 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
2897 .addUse(RegNo: GlobalWorkSize)
2898 .addUse(RegNo: LocalWorkSize)
2899 .addUse(RegNo: GlobalWorkOffset);
2900 }
2901
2902 // When sret is used, store nd_range struct through the pointer in the first
2903 // argument.
2904 Register SRetReg = Call->Arguments[SRetArgIdx];
2905 SPIRVTypeInst SRetPtrType = GR->getSPIRVTypeForVReg(VReg: SRetReg);
2906 SPIRVTypeInst SRetType = GR->getPointeeType(PtrType: SRetPtrType);
2907
2908 Register TmpReg = MRI->createVirtualRegister(RegClass: &SPIRV::iIDRegClass);
2909 GR->assignSPIRVTypeToVReg(Type: SRetType, VReg: TmpReg, MF);
2910
2911 MIRBuilder.buildInstr(Opcode: SPIRV::OpBuildNDRange)
2912 .addDef(RegNo: TmpReg)
2913 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: SRetType))
2914 .addUse(RegNo: GlobalWorkSize)
2915 .addUse(RegNo: LocalWorkSize)
2916 .addUse(RegNo: GlobalWorkOffset);
2917 return MIRBuilder.buildInstr(Opcode: SPIRV::OpStore)
2918 .addUse(RegNo: Call->Arguments[SRetArgIdx])
2919 .addUse(RegNo: TmpReg);
2920}
2921
2922static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
2923 MachineIRBuilder &MIRBuilder,
2924 SPIRVGlobalRegistry *GR) {
2925 // In this function there are three stages:
2926 // 1. prepare call indexes in order we expect them.
2927 // 2. process all arguments which requered preparation.
2928 // 3. create a SPIRV operator with arguments.
2929
2930 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2931 const DataLayout &DL = MIRBuilder.getDataLayout();
2932 const SPIRVTypeInst Int32Ty = GR->getOrCreateSPIRVIntegerType(BitWidth: 32, MIRBuilder);
2933
2934 // 1. prepare call indexes in order we expect them.
2935 // Based on clang sources, clang/lib/CodeGen/CGBuiltin.cpp, BIenqueue_kernel,
2936 // We expect 4 different layouts of call arguments:
2937 // 1) No events, no vargs: {Queue, Flags, Range, Kernel, Block};
2938 // 2) No events, varargs: {Queue, Flags, Range, Kernel, Block, NumElem,
2939 // ElemPtr};
2940 // 3) events, no varargs: {Queue, Flags, Range, NumEvents,
2941 // EventWaitList, EventRet, Kernel, Block};
2942 // 4) events, varargs: {Queue,
2943 // Flags, Range, NumEvents, EventWaitList, EventRet, Kernel, Block,
2944 // NumElem, ElemPtr};
2945 //
2946 // We also may expect __spirv_EnqueueKernel
2947
2948 bool IsSpirvOp = Call->isSpirvOp();
2949 bool HasEvents = Call->Builtin->name().contains(Other: "_events") || IsSpirvOp;
2950 bool HasVarArgs = Call->Builtin->name().contains(Other: "_varargs") || IsSpirvOp;
2951
2952 const unsigned NumArgs = Call->Arguments.size();
2953 const unsigned BaseArgIdx = 0;
2954 const unsigned IncorrectIdx = NumArgs + 1;
2955
2956 const unsigned QueueIdx = BaseArgIdx;
2957 const unsigned FlagsIdx = BaseArgIdx + 1;
2958 const unsigned NDRangeIdx = BaseArgIdx + 2;
2959 const unsigned NumEventsIdx = HasEvents ? BaseArgIdx + 3 : IncorrectIdx;
2960 const unsigned WaitEventsIdx = HasEvents ? BaseArgIdx + 4 : IncorrectIdx;
2961 const unsigned RetEventIdx = HasEvents ? BaseArgIdx + 5 : IncorrectIdx;
2962 const unsigned InvokeIdx = BaseArgIdx + 3 + (HasEvents ? 3 : 0);
2963 const unsigned ParamIdx = BaseArgIdx + 4 + (HasEvents ? 3 : 0);
2964 const unsigned LocalSizeNumElemIdx =
2965 HasVarArgs ? (BaseArgIdx + 5 + (HasEvents ? 3 : 0)) : IncorrectIdx;
2966 const unsigned LocalSizeElemPtrIdx =
2967 HasVarArgs ? (BaseArgIdx + 6 + (HasEvents ? 3 : 0)) : IncorrectIdx;
2968
2969 [[maybe_unused]] const unsigned LastArgIdx =
2970 (BaseArgIdx + 4 + (HasEvents ? 3 : 0) + (HasVarArgs ? 2 : 0));
2971 assert(LastArgIdx < NumArgs && "Incorrect number arguments");
2972
2973 // 2. Process all arguments which requered preparation.
2974 // 2.1 Events - use Call arguments, or use dummy nulls in case of absence of
2975 // events
2976
2977 auto BuildDeviceEventNullPtr = [&]() {
2978 LLVMContext &Ctx = MIRBuilder.getMF().getFunction().getContext();
2979 Type *DeviceEventTy = TargetExtType::get(Context&: Ctx, Name: "spirv.DeviceEvent");
2980 SPIRVTypeInst DeviceEventPtrTy = GR->getOrCreateSPIRVPointerType(
2981 BaseType: DeviceEventTy, MIRBuilder, SC: SPIRV::StorageClass::Generic);
2982 return GR->getOrCreateConstNullPtr(MIRBuilder, SpvType: DeviceEventPtrTy);
2983 };
2984
2985 Register NumEventsReg;
2986 Register WaitEventsReg;
2987 Register RetEventReg;
2988 if (HasEvents) {
2989 auto IsNullEvent = [&](Register R) {
2990 MachineInstr *Def = getDefInstrMaybeConstant(ConstReg&: R, MRI);
2991 return Def->getOpcode() == TargetOpcode::G_CONSTANT &&
2992 Def->getOperand(i: 1).getCImm()->isZero();
2993 };
2994
2995 NumEventsReg = Call->Arguments[NumEventsIdx];
2996 WaitEventsReg = Call->Arguments[WaitEventsIdx];
2997 RetEventReg = Call->Arguments[RetEventIdx];
2998 if (IsNullEvent(WaitEventsReg))
2999 WaitEventsReg = BuildDeviceEventNullPtr();
3000 if (IsNullEvent(RetEventReg))
3001 RetEventReg = BuildDeviceEventNullPtr();
3002 } else {
3003 NumEventsReg = buildConstantIntReg32(Val: 0, MIRBuilder, GR);
3004 Register NullPtr = BuildDeviceEventNullPtr();
3005 WaitEventsReg = NullPtr;
3006 RetEventReg = NullPtr;
3007 }
3008
3009 // 2.2 Invoke (Kernel)
3010 // The Invoke operand of OpEnqueueKernel must be the function's <id>
3011 // (per SPIR-V spec). The frontend hands us the result of an
3012 // addrspacecast of @block_invoke_kernel; bypass that cast so the
3013 // operand references the underlying G_GLOBAL_VALUE register, which
3014 // selectGlobalValue lowers to a placeholder later rewritten by
3015 // SPIRVModuleAnalysis to the OpFunction <id>.
3016 MachineInstr *InvokeGlobalMI =
3017 getBlockStructInstr(ParamReg: Call->Arguments[InvokeIdx], MRI);
3018 assert(InvokeGlobalMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
3019 Register InvokeReg = InvokeGlobalMI->getOperand(i: 0).getReg();
3020 // OpEnqueueKernel's Invoke operand uses the pID register class.
3021 MRI->setRegClass(Reg: InvokeReg, RC: &SPIRV::pIDRegClass);
3022
3023 // 2.3 Param, Param Size, Param Align
3024 Register BlockLiteralReg = Call->Arguments[ParamIdx];
3025 const SPIRVTypeInst Int8Ty = GR->getOrCreateSPIRVIntegerType(BitWidth: 8, MIRBuilder);
3026 const SPIRVTypeInst Int8PtrGen = GR->getOrCreateSPIRVPointerType(
3027 BaseType: Int8Ty, MIRBuilder, SC: SPIRV::StorageClass::Generic);
3028 Type *PType = const_cast<Type *>(getBlockStructType(ParamReg: BlockLiteralReg, MRI));
3029
3030 Register ParamReg = createVirtualRegister(SpvType: Int8PtrGen, GR, MIRBuilder);
3031 MIRBuilder.buildInstr(Opcode: SPIRV::OpBitcast)
3032 .addDef(RegNo: ParamReg)
3033 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Int8PtrGen))
3034 .addUse(RegNo: BlockLiteralReg);
3035 // TODO: these numbers should be obtained from block literal structure.
3036 Register ParamSizeReg =
3037 buildConstantIntReg32(Val: DL.getTypeStoreSize(Ty: PType), MIRBuilder, GR);
3038 Register ParamAlignReg =
3039 buildConstantIntReg32(Val: DL.getPrefTypeAlign(Ty: PType).value(), MIRBuilder, GR);
3040
3041 // 2.4 Local Size Array
3042 SmallVector<Register, 16> LocalSizes;
3043 if (HasVarArgs) {
3044 Register LocalSizeNumElem = Call->Arguments[LocalSizeNumElemIdx];
3045 MachineInstr *LocalSizeNumElemMI = MRI->getUniqueVRegDef(Reg: LocalSizeNumElem);
3046 const MachineOperand &ConstOp = LocalSizeNumElemMI->getOperand(i: 1);
3047 assert(LocalSizeNumElemMI->getOpcode() == TargetOpcode::G_CONSTANT &&
3048 ConstOp.isCImm() && "Expected constant immediate");
3049 uint64_t NumElem = ConstOp.getCImm()->getValue().getZExtValue();
3050
3051 Register LocalSizeArrayReg = Call->Arguments[LocalSizeElemPtrIdx];
3052
3053 for (unsigned i = 0; i < NumElem; ++i) {
3054 Register Reg = MRI->createVirtualRegister(RegClass: &SPIRV::pIDRegClass);
3055 auto GEPInst = MIRBuilder.buildIntrinsic(
3056 ID: Intrinsic::spv_gep, Res: ArrayRef<Register>{Reg}, HasSideEffects: true, isConvergent: false);
3057 GEPInst
3058 .addImm(Val: 0) // In bound.
3059 .addUse(RegNo: LocalSizeArrayReg) // Base pointer.
3060 .addUse(RegNo: buildConstantIntReg32(Val: 0, MIRBuilder, GR)) // Indices.
3061 .addUse(RegNo: buildConstantIntReg32(Val: i, MIRBuilder, GR));
3062 LocalSizes.push_back(Elt: Reg);
3063 }
3064 }
3065
3066 // 3. create a SPIRV operator with arguments.
3067 auto MIB = MIRBuilder.buildInstr(Opcode: SPIRV::OpEnqueueKernel)
3068 .addDef(RegNo: Call->ReturnRegister)
3069 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Int32Ty))
3070 .addUse(RegNo: Call->Arguments[QueueIdx])
3071 .addUse(RegNo: Call->Arguments[FlagsIdx])
3072 .addUse(RegNo: Call->Arguments[NDRangeIdx])
3073 .addUse(RegNo: NumEventsReg)
3074 .addUse(RegNo: WaitEventsReg)
3075 .addUse(RegNo: RetEventReg)
3076 .addUse(RegNo: InvokeReg)
3077 .addUse(RegNo: ParamReg)
3078 .addUse(RegNo: ParamSizeReg)
3079 .addUse(RegNo: ParamAlignReg);
3080 for (auto &LocalSize : LocalSizes)
3081 MIB.addUse(RegNo: LocalSize);
3082
3083 return true;
3084}
3085
3086static bool generateEnqueueInst(const SPIRV::IncomingCall *Call,
3087 MachineIRBuilder &MIRBuilder,
3088 SPIRVGlobalRegistry *GR) {
3089 // Lookup the instruction opcode in the TableGen records.
3090 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
3091 unsigned Opcode =
3092 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
3093
3094 switch (Opcode) {
3095 case SPIRV::OpRetainEvent:
3096 case SPIRV::OpReleaseEvent:
3097 return MIRBuilder.buildInstr(Opcode).addUse(RegNo: Call->Arguments[0]);
3098 case SPIRV::OpCreateUserEvent:
3099 case SPIRV::OpGetDefaultQueue:
3100 return MIRBuilder.buildInstr(Opcode)
3101 .addDef(RegNo: Call->ReturnRegister)
3102 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
3103 case SPIRV::OpIsValidEvent:
3104 return MIRBuilder.buildInstr(Opcode)
3105 .addDef(RegNo: Call->ReturnRegister)
3106 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
3107 .addUse(RegNo: Call->Arguments[0]);
3108 case SPIRV::OpSetUserEventStatus:
3109 return MIRBuilder.buildInstr(Opcode)
3110 .addUse(RegNo: Call->Arguments[0])
3111 .addUse(RegNo: Call->Arguments[1]);
3112 case SPIRV::OpCaptureEventProfilingInfo:
3113 return MIRBuilder.buildInstr(Opcode)
3114 .addUse(RegNo: Call->Arguments[0])
3115 .addUse(RegNo: Call->Arguments[1])
3116 .addUse(RegNo: Call->Arguments[2]);
3117 case SPIRV::OpBuildNDRange:
3118 return buildNDRange(Call, MIRBuilder, GR);
3119 case SPIRV::OpEnqueueKernel:
3120 return buildEnqueueKernel(Call, MIRBuilder, GR);
3121 default:
3122 return false;
3123 }
3124}
3125
3126static bool generateAsyncCopy(const SPIRV::IncomingCall *Call,
3127 MachineIRBuilder &MIRBuilder,
3128 SPIRVGlobalRegistry *GR) {
3129 // Lookup the instruction opcode in the TableGen records.
3130 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
3131 unsigned Opcode =
3132 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
3133
3134 bool IsSet = Opcode == SPIRV::OpGroupAsyncCopy;
3135 Register TypeReg = GR->getSPIRVTypeID(SpirvType: Call->ReturnType);
3136 if (Call->isSpirvOp())
3137 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
3138 TypeReg: IsSet ? TypeReg : Register(0));
3139
3140 auto Scope = buildConstantIntReg32(Val: SPIRV::Scope::Workgroup, MIRBuilder, GR);
3141
3142 switch (Opcode) {
3143 case SPIRV::OpGroupAsyncCopy: {
3144 SPIRVTypeInst NewType =
3145 Call->ReturnType->getOpcode() == SPIRV::OpTypeEvent
3146 ? nullptr
3147 : GR->getOrCreateSPIRVTypeByName(TypeStr: "spirv.Event", MIRBuilder, EmitIR: true);
3148 Register TypeReg = GR->getSPIRVTypeID(SpirvType: NewType ? NewType : Call->ReturnType);
3149 unsigned NumArgs = Call->Arguments.size();
3150 Register EventReg = Call->Arguments[NumArgs - 1];
3151 bool Res = MIRBuilder.buildInstr(Opcode)
3152 .addDef(RegNo: Call->ReturnRegister)
3153 .addUse(RegNo: TypeReg)
3154 .addUse(RegNo: Scope)
3155 .addUse(RegNo: Call->Arguments[0])
3156 .addUse(RegNo: Call->Arguments[1])
3157 .addUse(RegNo: Call->Arguments[2])
3158 .addUse(RegNo: Call->Arguments.size() > 4
3159 ? Call->Arguments[3]
3160 : buildConstantIntReg32(Val: 1, MIRBuilder, GR))
3161 .addUse(RegNo: EventReg);
3162 if (NewType)
3163 updateRegType(Reg: Call->ReturnRegister, Ty: nullptr, SpirvTy: NewType, GR, MIB&: MIRBuilder,
3164 MRI&: MIRBuilder.getMF().getRegInfo());
3165 return Res;
3166 }
3167 case SPIRV::OpGroupWaitEvents:
3168 return MIRBuilder.buildInstr(Opcode)
3169 .addUse(RegNo: Scope)
3170 .addUse(RegNo: Call->Arguments[0])
3171 .addUse(RegNo: Call->Arguments[1]);
3172 default:
3173 return false;
3174 }
3175}
3176
3177static bool generateConvertInst(StringRef DemangledCall,
3178 const SPIRV::IncomingCall *Call,
3179 MachineIRBuilder &MIRBuilder,
3180 SPIRVGlobalRegistry *GR) {
3181 // Lookup the conversion builtin in the TableGen records.
3182 const SPIRV::ConvertBuiltin *Builtin =
3183 SPIRV::lookupConvertBuiltin(Name: Call->Builtin->name(), Set: Call->Builtin->Set);
3184
3185 if (!Builtin && Call->isSpirvOp()) {
3186 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
3187 unsigned Opcode =
3188 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
3189 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
3190 TypeReg: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
3191 }
3192
3193 assert(Builtin && "Conversion builtin not found.");
3194 if (Builtin->IsSaturated)
3195 buildOpDecorate(Reg: Call->ReturnRegister, MIRBuilder,
3196 Dec: SPIRV::Decoration::SaturatedConversion, DecArgs: {});
3197
3198 if (Builtin->IsRounded) {
3199 bool AnyTypeIsFloat =
3200 GR->isScalarOrVectorOfType(VReg: Call->ReturnRegister, TypeOpcode: SPIRV::OpTypeFloat) ||
3201 GR->isScalarOrVectorOfType(VReg: Call->Arguments[0], TypeOpcode: SPIRV::OpTypeFloat);
3202
3203 // Rounding mode decorations are only valid for floating point types.
3204 // Conversion builtins from integer to integer are equivalent to their
3205 // non-rounded counterparts.
3206 if (AnyTypeIsFloat) {
3207 buildOpDecorate(Reg: Call->ReturnRegister, MIRBuilder,
3208 Dec: SPIRV::Decoration::FPRoundingMode,
3209 DecArgs: {(unsigned)Builtin->RoundingMode});
3210 }
3211 }
3212
3213 std::string NeedExtMsg; // no errors if empty
3214 bool IsRightComponentsNumber = true; // check if input/output accepts vectors
3215 unsigned Opcode = SPIRV::OpNop;
3216 if (GR->isScalarOrVectorOfType(VReg: Call->Arguments[0], TypeOpcode: SPIRV::OpTypeInt)) {
3217 // Int -> ...
3218 bool IsSourceSigned =
3219 DemangledCall[DemangledCall.find_first_of(C: '(') + 1] != 'u';
3220 if (GR->isScalarOrVectorOfType(VReg: Call->ReturnRegister, TypeOpcode: SPIRV::OpTypeInt)) {
3221 // Int -> Int
3222 if (Builtin->IsSaturated)
3223 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
3224 : SPIRV::OpSatConvertSToU;
3225 else
3226 Opcode = IsSourceSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3227 } else if (GR->isScalarOrVectorOfType(VReg: Call->ReturnRegister,
3228 TypeOpcode: SPIRV::OpTypeFloat)) {
3229 // Int -> Float
3230 if (Builtin->IsBfloat16) {
3231 const auto *ST = static_cast<const SPIRVSubtarget *>(
3232 &MIRBuilder.getMF().getSubtarget());
3233 if (!ST->canUseExtension(
3234 E: SPIRV::Extension::SPV_INTEL_bfloat16_conversion))
3235 NeedExtMsg = "SPV_INTEL_bfloat16_conversion";
3236 IsRightComponentsNumber =
3237 GR->getScalarOrVectorComponentCount(VReg: Call->Arguments[0]) ==
3238 GR->getScalarOrVectorComponentCount(VReg: Call->ReturnRegister);
3239 Opcode = SPIRV::OpConvertBF16ToFINTEL;
3240 } else {
3241 Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
3242 }
3243 }
3244 } else if (GR->isScalarOrVectorOfType(VReg: Call->Arguments[0],
3245 TypeOpcode: SPIRV::OpTypeFloat)) {
3246 // Float -> ...
3247 if (GR->isScalarOrVectorOfType(VReg: Call->ReturnRegister, TypeOpcode: SPIRV::OpTypeInt)) {
3248 // Float -> Int
3249 if (Builtin->IsBfloat16) {
3250 const auto *ST = static_cast<const SPIRVSubtarget *>(
3251 &MIRBuilder.getMF().getSubtarget());
3252 if (!ST->canUseExtension(
3253 E: SPIRV::Extension::SPV_INTEL_bfloat16_conversion))
3254 NeedExtMsg = "SPV_INTEL_bfloat16_conversion";
3255 IsRightComponentsNumber =
3256 GR->getScalarOrVectorComponentCount(VReg: Call->Arguments[0]) ==
3257 GR->getScalarOrVectorComponentCount(VReg: Call->ReturnRegister);
3258 Opcode = SPIRV::OpConvertFToBF16INTEL;
3259 } else {
3260 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
3261 : SPIRV::OpConvertFToU;
3262 }
3263 } else if (GR->isScalarOrVectorOfType(VReg: Call->ReturnRegister,
3264 TypeOpcode: SPIRV::OpTypeFloat)) {
3265 if (Builtin->IsTF32) {
3266 const auto *ST = static_cast<const SPIRVSubtarget *>(
3267 &MIRBuilder.getMF().getSubtarget());
3268 if (!ST->canUseExtension(
3269 E: SPIRV::Extension::SPV_INTEL_tensor_float32_conversion))
3270 NeedExtMsg = "SPV_INTEL_tensor_float32_conversion";
3271 IsRightComponentsNumber =
3272 GR->getScalarOrVectorComponentCount(VReg: Call->Arguments[0]) ==
3273 GR->getScalarOrVectorComponentCount(VReg: Call->ReturnRegister);
3274 Opcode = SPIRV::OpRoundFToTF32INTEL;
3275 } else {
3276 // Float -> Float
3277 Opcode = SPIRV::OpFConvert;
3278 }
3279 }
3280 }
3281
3282 StringRef BuiltinName = SPIRV::getConvertBuiltinStr(Offset: Builtin->Name);
3283 if (!NeedExtMsg.empty()) {
3284 std::string DiagMsg = std::string(BuiltinName) +
3285 ": the builtin requires the following SPIR-V "
3286 "extension: " +
3287 NeedExtMsg;
3288 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
3289 }
3290 if (!IsRightComponentsNumber) {
3291 std::string DiagMsg =
3292 std::string(BuiltinName) +
3293 ": result and argument must have the same number of components";
3294 report_fatal_error(reason: DiagMsg.c_str(), gen_crash_diag: false);
3295 }
3296 assert(Opcode != SPIRV::OpNop &&
3297 "Conversion between the types not implemented!");
3298
3299 MIRBuilder.buildInstr(Opcode)
3300 .addDef(RegNo: Call->ReturnRegister)
3301 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
3302 .addUse(RegNo: Call->Arguments[0]);
3303 return true;
3304}
3305
3306static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call,
3307 MachineIRBuilder &MIRBuilder,
3308 SPIRVGlobalRegistry *GR) {
3309 // Lookup the vector load/store builtin in the TableGen records.
3310 const SPIRV::VectorLoadStoreBuiltin *Builtin =
3311 SPIRV::lookupVectorLoadStoreBuiltin(Name: Call->Builtin->name(),
3312 Set: Call->Builtin->Set);
3313 // Build extended instruction.
3314 auto MIB =
3315 MIRBuilder.buildInstr(Opcode: SPIRV::OpExtInst)
3316 .addDef(RegNo: Call->ReturnRegister)
3317 .addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType))
3318 .addImm(Val: static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
3319 .addImm(Val: Builtin->Number);
3320 for (auto Argument : Call->Arguments)
3321 MIB.addUse(RegNo: Argument);
3322 StringRef BuiltinName = SPIRV::getVectorLoadStoreBuiltinStr(Offset: Builtin->Name);
3323 if (BuiltinName.contains(Other: "load") && Builtin->ElementCount > 1)
3324 MIB.addImm(Val: Builtin->ElementCount);
3325
3326 // Rounding mode should be passed as a last argument in the MI for builtins
3327 // like "vstorea_halfn_r".
3328 if (Builtin->IsRounded)
3329 MIB.addImm(Val: static_cast<uint32_t>(Builtin->RoundingMode));
3330 return true;
3331}
3332
3333static bool generateAFPInst(const SPIRV::IncomingCall *Call,
3334 MachineIRBuilder &MIRBuilder,
3335 SPIRVGlobalRegistry *GR) {
3336 const auto *Builtin = Call->Builtin;
3337 auto *MRI = MIRBuilder.getMRI();
3338 unsigned Opcode =
3339 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
3340 const Type *RetTy = GR->getTypeForSPIRVType(Ty: Call->ReturnType);
3341 bool IsVoid = RetTy->isVoidTy();
3342 auto MIB = MIRBuilder.buildInstr(Opcode);
3343 Register DestReg;
3344 if (IsVoid) {
3345 LLT PtrTy = MRI->getType(Reg: Call->Arguments[0]);
3346 DestReg = MRI->createGenericVirtualRegister(Ty: PtrTy);
3347 MRI->setRegClass(Reg: DestReg, RC: &SPIRV::pIDRegClass);
3348 SPIRVTypeInst PointeeTy =
3349 GR->getPointeeType(PtrType: GR->getSPIRVTypeForVReg(VReg: Call->Arguments[0]));
3350 MIB.addDef(RegNo: DestReg);
3351 MIB.addUse(RegNo: GR->getSPIRVTypeID(SpirvType: PointeeTy));
3352 } else {
3353 MIB.addDef(RegNo: Call->ReturnRegister);
3354 MIB.addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
3355 }
3356 for (unsigned i = IsVoid ? 1 : 0; i < Call->Arguments.size(); ++i) {
3357 Register Arg = Call->Arguments[i];
3358 MachineInstr *DefMI = MRI->getUniqueVRegDef(Reg: Arg);
3359 if (DefMI->getOpcode() == TargetOpcode::G_CONSTANT &&
3360 DefMI->getOperand(i: 1).isCImm()) {
3361 MIB.addImm(Val: getIConstVal(ConstReg: Arg, MRI));
3362 } else {
3363 MIB.addUse(RegNo: Arg);
3364 }
3365 }
3366 if (IsVoid) {
3367 LLT PtrTy = MRI->getType(Reg: Call->Arguments[0]);
3368 MachineMemOperand *MMO = MIRBuilder.getMF().getMachineMemOperand(
3369 PtrInfo: MachinePointerInfo(), F: MachineMemOperand::MOStore,
3370 Size: PtrTy.getSizeInBytes(), BaseAlignment: Align(4));
3371 MIRBuilder.buildStore(Val: DestReg, Addr: Call->Arguments[0], MMO&: *MMO);
3372 }
3373 return true;
3374}
3375
3376static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call,
3377 MachineIRBuilder &MIRBuilder,
3378 SPIRVGlobalRegistry *GR) {
3379 // Lookup the instruction opcode in the TableGen records.
3380 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
3381 unsigned Opcode =
3382 SPIRV::lookupNativeBuiltin(Name: Builtin->name(), Set: Builtin->Set)->Opcode;
3383 bool IsLoad = Opcode == SPIRV::OpLoad;
3384 // Build the instruction.
3385 auto MIB = MIRBuilder.buildInstr(Opcode);
3386 if (IsLoad) {
3387 MIB.addDef(RegNo: Call->ReturnRegister);
3388 MIB.addUse(RegNo: GR->getSPIRVTypeID(SpirvType: Call->ReturnType));
3389 }
3390 // Add a pointer to the value to load/store.
3391 MIB.addUse(RegNo: Call->Arguments[0]);
3392 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3393 // Add a value to store.
3394 if (!IsLoad)
3395 MIB.addUse(RegNo: Call->Arguments[1]);
3396 // Add optional memory attributes and an alignment.
3397 unsigned NumArgs = Call->Arguments.size();
3398 if ((IsLoad && NumArgs >= 2) || NumArgs >= 3)
3399 MIB.addImm(Val: getIConstVal(ConstReg: Call->Arguments[IsLoad ? 1 : 2], MRI));
3400 if ((IsLoad && NumArgs >= 3) || NumArgs >= 4)
3401 MIB.addImm(Val: getIConstVal(ConstReg: Call->Arguments[IsLoad ? 2 : 3], MRI));
3402 return true;
3403}
3404
3405namespace SPIRV {
3406// Try to find a builtin function attributes by a demangled function name and
3407// return a tuple <builtin group, op code, ext instruction number>, or a special
3408// tuple value <-1, 0, 0> if the builtin function is not found.
3409// Not all builtin functions are supported, only those with a ready-to-use op
3410// code or instruction number defined in TableGen.
3411// TODO: consider a major rework of mapping demangled calls into a builtin
3412// functions to unify search and decrease number of individual cases.
3413std::tuple<int, unsigned, unsigned>
3414mapBuiltinToOpcode(StringRef DemangledCall,
3415 SPIRV::InstructionSet::InstructionSet Set) {
3416 Register Reg;
3417 SmallVector<Register> Args;
3418 std::unique_ptr<const IncomingCall> Call =
3419 lookupBuiltin(DemangledCall, Set, ReturnRegister: Reg, ReturnType: nullptr, Arguments: Args);
3420 if (!Call)
3421 return std::make_tuple(args: -1, args: 0, args: 0);
3422
3423 switch (Call->Builtin->Group) {
3424 case SPIRV::Relational:
3425 case SPIRV::Atomic:
3426 case SPIRV::Barrier:
3427 case SPIRV::CastToPtr:
3428 case SPIRV::ImageMiscQuery:
3429 case SPIRV::SpecConstant:
3430 case SPIRV::Enqueue:
3431 case SPIRV::AsyncCopy:
3432 case SPIRV::LoadStore:
3433 case SPIRV::CoopMatr:
3434 case SPIRV::Arithmetic:
3435 if (const auto *R = SPIRV::lookupNativeBuiltin(Name: Call->Builtin->name(),
3436 Set: Call->Builtin->Set))
3437 return std::make_tuple(args: Call->Builtin->Group, args: R->Opcode, args: 0);
3438 break;
3439 case SPIRV::Extended:
3440 if (const auto *R = SPIRV::lookupExtendedBuiltin(Name: Call->Builtin->name(),
3441 Set: Call->Builtin->Set))
3442 return std::make_tuple(args: Call->Builtin->Group, args: 0, args: R->Number);
3443 break;
3444 case SPIRV::VectorLoadStore:
3445 if (const auto *R = SPIRV::lookupVectorLoadStoreBuiltin(
3446 Name: Call->Builtin->name(), Set: Call->Builtin->Set))
3447 return std::make_tuple(args: SPIRV::Extended, args: 0, args: R->Number);
3448 break;
3449 case SPIRV::Group:
3450 if (const auto *R = SPIRV::lookupGroupBuiltin(Name: Call->Builtin->name()))
3451 return std::make_tuple(args: Call->Builtin->Group, args: R->Opcode, args: 0);
3452 break;
3453 case SPIRV::AtomicFloating:
3454 if (const auto *R =
3455 SPIRV::lookupAtomicFloatingBuiltin(Name: Call->Builtin->name()))
3456 return std::make_tuple(args: Call->Builtin->Group, args: R->Opcode, args: 0);
3457 break;
3458 case SPIRV::IntelSubgroups:
3459 if (const auto *R =
3460 SPIRV::lookupIntelSubgroupsBuiltin(Name: Call->Builtin->name()))
3461 return std::make_tuple(args: Call->Builtin->Group, args: R->Opcode, args: 0);
3462 break;
3463 case SPIRV::GroupUniform:
3464 if (const auto *R = SPIRV::lookupGroupUniformBuiltin(Name: Call->Builtin->name()))
3465 return std::make_tuple(args: Call->Builtin->Group, args: R->Opcode, args: 0);
3466 break;
3467 case SPIRV::IntegerDot:
3468 if (const auto *R =
3469 SPIRV::lookupIntegerDotProductBuiltin(Name: Call->Builtin->name()))
3470 return std::make_tuple(args: Call->Builtin->Group, args: R->Opcode, args: 0);
3471 break;
3472 case SPIRV::WriteImage:
3473 return std::make_tuple(args: Call->Builtin->Group, args: SPIRV::OpImageWrite, args: 0);
3474 case SPIRV::Select:
3475 return std::make_tuple(args: Call->Builtin->Group, args: TargetOpcode::G_SELECT, args: 0);
3476 case SPIRV::Construct:
3477 return std::make_tuple(args: Call->Builtin->Group, args: SPIRV::OpCompositeConstruct,
3478 args: 0);
3479 case SPIRV::KernelClock:
3480 return std::make_tuple(args: Call->Builtin->Group, args: SPIRV::OpReadClockKHR, args: 0);
3481 default:
3482 return std::make_tuple(args: -1, args: 0, args: 0);
3483 }
3484 return std::make_tuple(args: -1, args: 0, args: 0);
3485}
3486
3487std::optional<bool> lowerBuiltin(StringRef DemangledCall,
3488 SPIRV::InstructionSet::InstructionSet Set,
3489 MachineIRBuilder &MIRBuilder,
3490 const Register OrigRet, const Type *OrigRetTy,
3491 const SmallVectorImpl<Register> &Args,
3492 SPIRVGlobalRegistry *GR, const CallBase &CB) {
3493 LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
3494
3495 // Lookup the builtin in the TableGen records.
3496 SPIRVTypeInst SpvType = GR->getSPIRVTypeForVReg(VReg: OrigRet);
3497 assert(SpvType && "Inconsistent return register: expected valid type info");
3498 std::unique_ptr<const IncomingCall> Call =
3499 lookupBuiltin(DemangledCall, Set, ReturnRegister: OrigRet, ReturnType: SpvType, Arguments: Args);
3500
3501 if (!Call) {
3502 LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
3503 return std::nullopt;
3504 }
3505
3506 // Check if the provided args meet the builtin requirements. If not, treat
3507 // the call as a regular function call rather than crashing.
3508 if (Args.size() < Call->Builtin->MinNumArgs) {
3509 LLVM_DEBUG(dbgs() << "Too few arguments for builtin " << DemangledCall
3510 << ": expected at least " << Call->Builtin->MinNumArgs
3511 << ", got " << Args.size()
3512 << "; treating as a normal function\n");
3513 return std::nullopt;
3514 }
3515 if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs) {
3516 LLVM_DEBUG(dbgs() << "Too many arguments for builtin " << DemangledCall
3517 << ": expected at most " << Call->Builtin->MaxNumArgs
3518 << ", got " << Args.size()
3519 << "; treating as a normal function\n");
3520 return std::nullopt;
3521 }
3522
3523 // Match the builtin with implementation based on the grouping.
3524 switch (Call->Builtin->Group) {
3525 case SPIRV::Extended:
3526 return generateExtInst(Call: Call.get(), MIRBuilder, GR, CB);
3527 case SPIRV::Relational:
3528 return generateRelationalInst(Call: Call.get(), MIRBuilder, GR);
3529 case SPIRV::Group:
3530 return generateGroupInst(Call: Call.get(), MIRBuilder, GR);
3531 case SPIRV::Variable:
3532 return generateBuiltinVar(Call: Call.get(), MIRBuilder, GR);
3533 case SPIRV::Atomic:
3534 return generateAtomicInst(Call: Call.get(), MIRBuilder, GR);
3535 case SPIRV::AtomicFloating:
3536 return generateAtomicFloatingInst(Call: Call.get(), MIRBuilder, GR);
3537 case SPIRV::Barrier:
3538 return generateBarrierInst(Call: Call.get(), MIRBuilder, GR);
3539 case SPIRV::CastToPtr:
3540 return generateCastToPtrInst(Call: Call.get(), MIRBuilder, GR);
3541 case SPIRV::Dot:
3542 case SPIRV::IntegerDot:
3543 return generateDotOrFMulInst(DemangledCall, Call: Call.get(), MIRBuilder, GR);
3544 case SPIRV::Wave:
3545 return generateWaveInst(Call: Call.get(), MIRBuilder, GR);
3546 case SPIRV::ICarryBorrow:
3547 return generateICarryBorrowInst(Call: Call.get(), MIRBuilder, GR);
3548 case SPIRV::MulExtended:
3549 return generateMulExtendedInst(Call: Call.get(), MIRBuilder, GR);
3550 case SPIRV::Arithmetic:
3551 return generateArithmeticInst(Call: Call.get(), MIRBuilder, GR);
3552 case SPIRV::GetQuery:
3553 return generateGetQueryInst(Call: Call.get(), MIRBuilder, GR);
3554 case SPIRV::ImageSizeQuery:
3555 return generateImageSizeQueryInst(Call: Call.get(), MIRBuilder, GR);
3556 case SPIRV::ImageMiscQuery:
3557 return generateImageMiscQueryInst(Call: Call.get(), MIRBuilder, GR);
3558 case SPIRV::ReadImage:
3559 return generateReadImageInst(DemangledCall, Call: Call.get(), MIRBuilder, GR);
3560 case SPIRV::WriteImage:
3561 return generateWriteImageInst(Call: Call.get(), MIRBuilder, GR);
3562 case SPIRV::SampleImage:
3563 return generateSampleImageInst(DemangledCall, Call: Call.get(), MIRBuilder, GR);
3564 case SPIRV::Select:
3565 return generateSelectInst(Call: Call.get(), MIRBuilder);
3566 case SPIRV::Construct:
3567 return generateConstructInst(Call: Call.get(), MIRBuilder, GR);
3568 case SPIRV::SpecConstant:
3569 return generateSpecConstantInst(Call: Call.get(), MIRBuilder, GR);
3570 case SPIRV::Enqueue:
3571 return generateEnqueueInst(Call: Call.get(), MIRBuilder, GR);
3572 case SPIRV::AsyncCopy:
3573 return generateAsyncCopy(Call: Call.get(), MIRBuilder, GR);
3574 case SPIRV::Convert:
3575 return generateConvertInst(DemangledCall, Call: Call.get(), MIRBuilder, GR);
3576 case SPIRV::VectorLoadStore:
3577 return generateVectorLoadStoreInst(Call: Call.get(), MIRBuilder, GR);
3578 case SPIRV::LoadStore:
3579 return generateLoadStoreInst(Call: Call.get(), MIRBuilder, GR);
3580 case SPIRV::IntelSubgroups:
3581 return generateIntelSubgroupsInst(Call: Call.get(), MIRBuilder, GR);
3582 case SPIRV::GroupUniform:
3583 return generateGroupUniformInst(Call: Call.get(), MIRBuilder, GR);
3584 case SPIRV::KernelClock:
3585 return generateKernelClockInst(Call: Call.get(), MIRBuilder, GR);
3586 case SPIRV::CoopMatr:
3587 return generateCoopMatrInst(Call: Call.get(), MIRBuilder, GR);
3588 case SPIRV::ExtendedBitOps:
3589 return generateExtendedBitOpsInst(Call: Call.get(), MIRBuilder, GR);
3590 case SPIRV::BindlessINTEL:
3591 return generateBindlessImageINTELInst(Call: Call.get(), MIRBuilder, GR);
3592 case SPIRV::TernaryBitwiseINTEL:
3593 return generateTernaryBitwiseFunctionINTELInst(Call: Call.get(), MIRBuilder, GR);
3594 case SPIRV::Block2DLoadStore:
3595 return generate2DBlockIOINTELInst(Call: Call.get(), MIRBuilder, GR);
3596 case SPIRV::Pipe:
3597 return generatePipeInst(Call: Call.get(), MIRBuilder, GR);
3598 case SPIRV::PredicatedLoadStore:
3599 return generatePredicatedLoadStoreInst(Call: Call.get(), MIRBuilder, GR);
3600 case SPIRV::BlockingPipes:
3601 return generateBlockingPipesInst(Call: Call.get(), MIRBuilder, GR);
3602 case SPIRV::ArbitraryPrecisionFixedPoint:
3603 return generateAPFixedPointInst(Call: Call.get(), MIRBuilder, GR);
3604 case SPIRV::ImageChannelDataTypes:
3605 return generateImageChannelDataTypeInst(Call: Call.get(), MIRBuilder, GR);
3606 case SPIRV::ArbitraryFloatingPoint:
3607 return generateAFPInst(Call: Call.get(), MIRBuilder, GR);
3608 }
3609 return false;
3610}
3611
3612Type *parseBuiltinCallArgumentType(StringRef TypeStr, LLVMContext &Ctx) {
3613 // Parse strings representing OpenCL builtin types.
3614 if (hasBuiltinTypePrefix(Name: TypeStr)) {
3615 // OpenCL builtin types in demangled call strings have the following format:
3616 // e.g. ocl_image2d_ro
3617 [[maybe_unused]] bool IsOCLBuiltinType = TypeStr.consume_front(Prefix: "ocl_");
3618 assert(IsOCLBuiltinType && "Invalid OpenCL builtin prefix");
3619
3620 // Check if this is pointer to a builtin type and not just pointer
3621 // representing a builtin type. In case it is a pointer to builtin type,
3622 // this will require additional handling in the method calling
3623 // parseBuiltinCallArgumentBaseType(...) as this function only retrieves the
3624 // base types.
3625 if (TypeStr.ends_with(Suffix: "*"))
3626 TypeStr = TypeStr.slice(Start: 0, End: TypeStr.find_first_of(Chars: " *"));
3627
3628 return parseBuiltinTypeNameToTargetExtType(TypeName: "opencl." + TypeStr.str() + "_t",
3629 Context&: Ctx);
3630 }
3631
3632 // Parse type name in either "typeN" or "type vector[N]" format, where
3633 // N is the number of elements of the vector.
3634 Type *BaseType;
3635 unsigned VecElts = 0;
3636
3637 BaseType = parseBasicTypeName(TypeName&: TypeStr, Ctx);
3638 if (!BaseType)
3639 // Unable to recognize SPIRV type name.
3640 return nullptr;
3641
3642 // Handle "typeN*" or "type vector[N]*".
3643 TypeStr.consume_back(Suffix: "*");
3644
3645 if (TypeStr.consume_front(Prefix: " vector["))
3646 TypeStr = TypeStr.substr(Start: 0, N: TypeStr.find(C: ']'));
3647
3648 TypeStr.getAsInteger(Radix: 10, Result&: VecElts);
3649 if (VecElts > 0)
3650 BaseType = VectorType::get(
3651 ElementType: BaseType->isVoidTy() ? Type::getInt8Ty(C&: Ctx) : BaseType, NumElements: VecElts, Scalable: false);
3652
3653 return BaseType;
3654}
3655
3656bool parseBuiltinTypeStr(SmallVector<StringRef, 10> &BuiltinArgsTypeStrs,
3657 StringRef DemangledCall, LLVMContext &Ctx) {
3658 auto Pos1 = DemangledCall.find(C: '(');
3659 if (Pos1 == StringRef::npos)
3660 return false;
3661 auto Pos2 = DemangledCall.find(C: ')');
3662 if (Pos2 == StringRef::npos || Pos1 > Pos2)
3663 return false;
3664 DemangledCall.slice(Start: Pos1 + 1, End: Pos2)
3665 .split(A&: BuiltinArgsTypeStrs, Separator: ',', MaxSplit: -1, KeepEmpty: false);
3666 return true;
3667}
3668
3669Type *parseBuiltinCallArgumentBaseType(StringRef DemangledCall, unsigned ArgIdx,
3670 LLVMContext &Ctx) {
3671 SmallVector<StringRef, 10> BuiltinArgsTypeStrs;
3672 parseBuiltinTypeStr(BuiltinArgsTypeStrs, DemangledCall, Ctx);
3673 if (ArgIdx >= BuiltinArgsTypeStrs.size())
3674 return nullptr;
3675 StringRef TypeStr = BuiltinArgsTypeStrs[ArgIdx].trim();
3676 return parseBuiltinCallArgumentType(TypeStr, Ctx);
3677}
3678
3679struct BuiltinType {
3680 StringTable::Offset Name;
3681 uint32_t Opcode;
3682};
3683
3684#define GET_BuiltinTypes_DECL
3685#define GET_BuiltinTypes_IMPL
3686
3687struct OpenCLType {
3688 StringTable::Offset Name;
3689 StringTable::Offset SpirvTypeLiteral;
3690};
3691
3692#define GET_OpenCLTypes_DECL
3693#define GET_OpenCLTypes_IMPL
3694
3695#include "SPIRVGenTables.inc"
3696} // namespace SPIRV
3697
3698//===----------------------------------------------------------------------===//
3699// Misc functions for parsing builtin types.
3700//===----------------------------------------------------------------------===//
3701
3702static Type *parseTypeString(StringRef Name, LLVMContext &Context) {
3703 if (Name.starts_with(Prefix: "void"))
3704 return Type::getVoidTy(C&: Context);
3705 else if (Name.starts_with(Prefix: "int") || Name.starts_with(Prefix: "uint"))
3706 return Type::getInt32Ty(C&: Context);
3707 else if (Name.starts_with(Prefix: "bfloat"))
3708 return Type::getBFloatTy(C&: Context);
3709 else if (Name.starts_with(Prefix: "float"))
3710 return Type::getFloatTy(C&: Context);
3711 else if (Name.starts_with(Prefix: "half"))
3712 return Type::getHalfTy(C&: Context);
3713 else if (Name.starts_with(Prefix: "double"))
3714 return Type::getDoubleTy(C&: Context);
3715 report_fatal_error(reason: "Unable to recognize type!");
3716}
3717
3718//===----------------------------------------------------------------------===//
3719// Implementation functions for builtin types.
3720//===----------------------------------------------------------------------===//
3721
3722static SPIRVTypeInst
3723getNonParameterizedType(const TargetExtType *ExtensionType,
3724 const SPIRV::BuiltinType *TypeRecord,
3725 MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
3726 unsigned Opcode = TypeRecord->Opcode;
3727 // Create or get an existing type from GlobalRegistry.
3728 return GR->getOrCreateOpTypeByOpcode(Ty: ExtensionType, MIRBuilder, Opcode);
3729}
3730
3731static SPIRVTypeInst getSamplerType(MachineIRBuilder &MIRBuilder,
3732 SPIRVGlobalRegistry *GR) {
3733 // Create or get an existing type from GlobalRegistry.
3734 return GR->getOrCreateOpTypeSampler(MIRBuilder);
3735}
3736
3737static SPIRVTypeInst getPipeType(const TargetExtType *ExtensionType,
3738 MachineIRBuilder &MIRBuilder,
3739 SPIRVGlobalRegistry *GR) {
3740 assert(ExtensionType->getNumIntParameters() == 1 &&
3741 "Invalid number of parameters for SPIR-V pipe builtin!");
3742 // Create or get an existing type from GlobalRegistry.
3743 return GR->getOrCreateOpTypePipe(MIRBuilder,
3744 AccQual: SPIRV::AccessQualifier::AccessQualifier(
3745 ExtensionType->getIntParameter(i: 0)));
3746}
3747
3748static SPIRVTypeInst getCoopMatrType(const TargetExtType *ExtensionType,
3749 MachineIRBuilder &MIRBuilder,
3750 SPIRVGlobalRegistry *GR) {
3751 assert(ExtensionType->getNumIntParameters() == 4 &&
3752 "Invalid number of parameters for SPIR-V coop matrices builtin!");
3753 assert(ExtensionType->getNumTypeParameters() == 1 &&
3754 "SPIR-V coop matrices builtin type must have a type parameter!");
3755 SPIRVTypeInst ElemType =
3756 GR->getOrCreateSPIRVType(Type: ExtensionType->getTypeParameter(i: 0), MIRBuilder,
3757 AQ: SPIRV::AccessQualifier::ReadWrite, EmitIR: true);
3758 // Create or get an existing type from GlobalRegistry.
3759 return GR->getOrCreateOpTypeCoopMatr(
3760 MIRBuilder, ExtensionType, ElemType, Scope: ExtensionType->getIntParameter(i: 0),
3761 Rows: ExtensionType->getIntParameter(i: 1), Columns: ExtensionType->getIntParameter(i: 2),
3762 Use: ExtensionType->getIntParameter(i: 3), EmitIR: true);
3763}
3764
3765static SPIRVTypeInst getSampledImageType(const TargetExtType *OpaqueType,
3766 MachineIRBuilder &MIRBuilder,
3767 SPIRVGlobalRegistry *GR) {
3768 SPIRVTypeInst OpaqueImageType = GR->getImageType(
3769 ExtensionType: OpaqueType, Qualifier: SPIRV::AccessQualifier::ReadOnly, MIRBuilder);
3770 // Create or get an existing type from GlobalRegistry.
3771 return GR->getOrCreateOpTypeSampledImage(ImageType: OpaqueImageType, MIRBuilder);
3772}
3773
3774static SPIRVTypeInst getInlineSpirvType(const TargetExtType *ExtensionType,
3775 MachineIRBuilder &MIRBuilder,
3776 SPIRVGlobalRegistry *GR) {
3777 assert(ExtensionType->getNumIntParameters() == 3 &&
3778 "Inline SPIR-V type builtin takes an opcode, size, and alignment "
3779 "parameter");
3780 auto Opcode = ExtensionType->getIntParameter(i: 0);
3781
3782 SmallVector<MCOperand> Operands;
3783 for (Type *Param : ExtensionType->type_params()) {
3784 if (const TargetExtType *ParamEType = dyn_cast<TargetExtType>(Val: Param)) {
3785 if (ParamEType->getName() == "spirv.IntegralConstant") {
3786 assert(ParamEType->getNumTypeParameters() == 1 &&
3787 "Inline SPIR-V integral constant builtin must have a type "
3788 "parameter");
3789 assert(ParamEType->getNumIntParameters() == 1 &&
3790 "Inline SPIR-V integral constant builtin must have a "
3791 "value parameter");
3792
3793 auto OperandValue = ParamEType->getIntParameter(i: 0);
3794 auto *OperandType = ParamEType->getTypeParameter(i: 0);
3795
3796 SPIRVTypeInst OperandSPIRVType = GR->getOrCreateSPIRVType(
3797 Type: OperandType, MIRBuilder, AQ: SPIRV::AccessQualifier::ReadWrite, EmitIR: true);
3798
3799 Operands.push_back(Elt: MCOperand::createReg(Reg: GR->buildConstantInt(
3800 Val: OperandValue, MIRBuilder, SpvType: OperandSPIRVType, EmitIR: true)));
3801 continue;
3802 } else if (ParamEType->getName() == "spirv.Literal") {
3803 assert(ParamEType->getNumTypeParameters() == 0 &&
3804 "Inline SPIR-V literal builtin does not take type "
3805 "parameters");
3806 assert(ParamEType->getNumIntParameters() == 1 &&
3807 "Inline SPIR-V literal builtin must have an integer "
3808 "parameter");
3809
3810 auto OperandValue = ParamEType->getIntParameter(i: 0);
3811
3812 Operands.push_back(Elt: MCOperand::createImm(Val: OperandValue));
3813 continue;
3814 }
3815 }
3816 SPIRVTypeInst TypeOperand = GR->getOrCreateSPIRVType(
3817 Type: Param, MIRBuilder, AQ: SPIRV::AccessQualifier::ReadWrite, EmitIR: true);
3818 Operands.push_back(Elt: MCOperand::createReg(Reg: GR->getSPIRVTypeID(SpirvType: TypeOperand)));
3819 }
3820
3821 return GR->getOrCreateUnknownType(Ty: ExtensionType, MIRBuilder, Opcode,
3822 Operands);
3823}
3824
3825static SPIRVTypeInst getVulkanBufferType(const TargetExtType *ExtensionType,
3826 MachineIRBuilder &MIRBuilder,
3827 SPIRVGlobalRegistry *GR) {
3828 assert(ExtensionType->getNumTypeParameters() == 1 &&
3829 "Vulkan buffers have exactly one type for the type of the buffer.");
3830 assert(ExtensionType->getNumIntParameters() == 2 &&
3831 "Vulkan buffer have 2 integer parameters: storage class and is "
3832 "writable.");
3833
3834 auto *T = ExtensionType->getTypeParameter(i: 0);
3835 auto SC = static_cast<SPIRV::StorageClass::StorageClass>(
3836 ExtensionType->getIntParameter(i: 0));
3837 bool IsWritable = ExtensionType->getIntParameter(i: 1);
3838 return GR->getOrCreateVulkanBufferType(MIRBuilder, ElemType: T, SC, IsWritable);
3839}
3840
3841static SPIRVTypeInst
3842getVulkanPushConstantType(const TargetExtType *ExtensionType,
3843 MachineIRBuilder &MIRBuilder,
3844 SPIRVGlobalRegistry *GR) {
3845 assert(ExtensionType->getNumTypeParameters() == 1 &&
3846 "Vulkan push constants have exactly one type as argument.");
3847 auto *T = ExtensionType->getTypeParameter(i: 0);
3848 return GR->getOrCreateVulkanPushConstantType(MIRBuilder, ElemType: T);
3849}
3850
3851static SPIRVTypeInst getLayoutType(const TargetExtType *ExtensionType,
3852 MachineIRBuilder &MIRBuilder,
3853 SPIRVGlobalRegistry *GR) {
3854 return GR->getOrCreateLayoutType(MIRBuilder, T: ExtensionType);
3855}
3856
3857namespace SPIRV {
3858TargetExtType *parseBuiltinTypeNameToTargetExtType(std::string TypeName,
3859 LLVMContext &Context) {
3860 StringRef NameWithParameters = TypeName;
3861
3862 // Pointers-to-opaque-structs representing OpenCL types are first translated
3863 // to equivalent SPIR-V types. OpenCL builtin type names should have the
3864 // following format: e.g. %opencl.event_t
3865 if (NameWithParameters.starts_with(Prefix: "opencl.")) {
3866 const SPIRV::OpenCLType *OCLTypeRecord =
3867 SPIRV::lookupOpenCLType(Name: NameWithParameters);
3868 if (!OCLTypeRecord)
3869 report_fatal_error(reason: "Missing TableGen record for OpenCL type: " +
3870 NameWithParameters);
3871 NameWithParameters =
3872 SPIRV::getOpenCLTypeStr(Offset: OCLTypeRecord->SpirvTypeLiteral);
3873 // Continue with the SPIR-V builtin type...
3874 }
3875
3876 // Names of the opaque structs representing a SPIR-V builtins without
3877 // parameters should have the following format: e.g. %spirv.Event
3878 assert(NameWithParameters.starts_with("spirv.") &&
3879 "Unknown builtin opaque type!");
3880
3881 // Parameterized SPIR-V builtins names follow this format:
3882 // e.g. %spirv.Image._void_1_0_0_0_0_0_0, %spirv.Pipe._0
3883 if (!NameWithParameters.contains(C: '_'))
3884 return TargetExtType::get(Context, Name: NameWithParameters);
3885
3886 SmallVector<StringRef> Parameters;
3887 unsigned BaseNameLength = NameWithParameters.find(C: '_') - 1;
3888 SplitString(Source: NameWithParameters.substr(Start: BaseNameLength + 1), OutFragments&: Parameters, Delimiters: "_");
3889
3890 SmallVector<Type *, 1> TypeParameters;
3891 bool HasTypeParameter = !isDigit(C: Parameters[0][0]);
3892 if (HasTypeParameter)
3893 TypeParameters.push_back(Elt: parseTypeString(Name: Parameters[0], Context));
3894 SmallVector<unsigned> IntParameters;
3895 for (unsigned i = HasTypeParameter ? 1 : 0; i < Parameters.size(); i++) {
3896 unsigned IntParameter = 0;
3897 bool ValidLiteral = !Parameters[i].getAsInteger(Radix: 10, Result&: IntParameter);
3898 (void)ValidLiteral;
3899 assert(ValidLiteral &&
3900 "Invalid format of SPIR-V builtin parameter literal!");
3901 IntParameters.push_back(Elt: IntParameter);
3902 }
3903 return TargetExtType::get(Context,
3904 Name: NameWithParameters.substr(Start: 0, N: BaseNameLength),
3905 Types: TypeParameters, Ints: IntParameters);
3906}
3907
3908SPIRVTypeInst
3909lowerBuiltinType(const Type *OpaqueType,
3910 SPIRV::AccessQualifier::AccessQualifier AccessQual,
3911 MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
3912 // In LLVM IR, SPIR-V and OpenCL builtin types are represented as either
3913 // target(...) target extension types or pointers-to-opaque-structs. The
3914 // approach relying on structs is deprecated and works only in the non-opaque
3915 // pointer mode (-opaque-pointers=0).
3916 // In order to maintain compatibility with LLVM IR generated by older versions
3917 // of Clang and LLVM/SPIR-V Translator, the pointers-to-opaque-structs are
3918 // "translated" to target extension types. This translation is temporary and
3919 // will be removed in the future release of LLVM.
3920 const TargetExtType *BuiltinType = dyn_cast<TargetExtType>(Val: OpaqueType);
3921 if (!BuiltinType)
3922 BuiltinType = parseBuiltinTypeNameToTargetExtType(
3923 TypeName: OpaqueType->getStructName().str(), Context&: MIRBuilder.getContext());
3924
3925 unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();
3926
3927 StringRef Name = BuiltinType->getName();
3928 LLVM_DEBUG(dbgs() << "Lowering builtin type: " << Name << "\n");
3929
3930 SPIRVTypeInst TargetType = nullptr;
3931 if (Name == "spirv.Type") {
3932 TargetType = getInlineSpirvType(ExtensionType: BuiltinType, MIRBuilder, GR);
3933 } else if (Name == "spirv.VulkanBuffer") {
3934 TargetType = getVulkanBufferType(ExtensionType: BuiltinType, MIRBuilder, GR);
3935 } else if (Name == "spirv.Padding") {
3936 TargetType = GR->getOrCreatePaddingType(MIRBuilder);
3937 } else if (Name == "spirv.PushConstant") {
3938 TargetType = getVulkanPushConstantType(ExtensionType: BuiltinType, MIRBuilder, GR);
3939 } else if (Name == "spirv.Layout") {
3940 TargetType = getLayoutType(ExtensionType: BuiltinType, MIRBuilder, GR);
3941 } else {
3942 // Lookup the demangled builtin type in the TableGen records.
3943 const SPIRV::BuiltinType *TypeRecord = SPIRV::lookupBuiltinType(Name);
3944 if (!TypeRecord)
3945 report_fatal_error(reason: "Missing TableGen record for builtin type: " + Name);
3946
3947 // "Lower" the BuiltinType into TargetType. The following get<...>Type
3948 // methods use the implementation details from TableGen records or
3949 // TargetExtType parameters to either create a new OpType<...> machine
3950 // instruction or get an existing equivalent SPIRV type from
3951 // GlobalRegistry.
3952
3953 switch (TypeRecord->Opcode) {
3954 case SPIRV::OpTypeImage:
3955 TargetType = GR->getImageType(ExtensionType: BuiltinType, Qualifier: AccessQual, MIRBuilder);
3956 break;
3957 case SPIRV::OpTypePipe:
3958 TargetType = getPipeType(ExtensionType: BuiltinType, MIRBuilder, GR);
3959 break;
3960 case SPIRV::OpTypeDeviceEvent:
3961 TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
3962 break;
3963 case SPIRV::OpTypeSampler:
3964 TargetType = getSamplerType(MIRBuilder, GR);
3965 break;
3966 case SPIRV::OpTypeSampledImage:
3967 TargetType = getSampledImageType(OpaqueType: BuiltinType, MIRBuilder, GR);
3968 break;
3969 case SPIRV::OpTypeCooperativeMatrixKHR:
3970 TargetType = getCoopMatrType(ExtensionType: BuiltinType, MIRBuilder, GR);
3971 break;
3972 default:
3973 TargetType =
3974 getNonParameterizedType(ExtensionType: BuiltinType, TypeRecord, MIRBuilder, GR);
3975 break;
3976 }
3977 }
3978
3979 // Emit OpName instruction if a new OpType<...> instruction was added
3980 // (equivalent type was not found in GlobalRegistry).
3981 if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
3982 buildOpName(Target: GR->getSPIRVTypeID(SpirvType: TargetType), Name, MIRBuilder);
3983
3984 return TargetType;
3985}
3986
3987bool isPipeOrAddressSpaceCastBuiltin(StringRef Name) {
3988 const DemangledBuiltin *Builtin = lookupBuiltin(Name, Set: OpenCL_std);
3989 if (!Builtin)
3990 return false;
3991 return Builtin->Group == Pipe || Builtin->Group == CastToPtr ||
3992 Builtin->Group == BlockingPipes;
3993}
3994} // namespace SPIRV
3995} // namespace llvm
3996