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