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