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