LLVM 19.0.0git
SPIRVBuiltins.cpp
Go to the documentation of this file.
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"
20#include "llvm/IR/IntrinsicsSPIRV.h"
21#include <string>
22#include <tuple>
23
24#define DEBUG_TYPE "spirv-builtins"
25
26namespace llvm {
27namespace SPIRV {
28#define GET_BuiltinGroup_DECL
29#include "SPIRVGenTables.inc"
30
33 InstructionSet::InstructionSet Set;
34 BuiltinGroup Group;
35 uint8_t MinNumArgs;
36 uint8_t MaxNumArgs;
37};
38
39#define GET_DemangledBuiltins_DECL
40#define GET_DemangledBuiltins_IMPL
41
43 const std::string BuiltinName;
45
49
56
57 bool isSpirvOp() const { return BuiltinName.rfind("__spirv_", 0) == 0; }
58};
59
62 InstructionSet::InstructionSet Set;
64};
65
66#define GET_NativeBuiltins_DECL
67#define GET_NativeBuiltins_IMPL
68
73 bool IsElect;
83};
84
85#define GET_GroupBuiltins_DECL
86#define GET_GroupBuiltins_IMPL
87
91 bool IsBlock;
92 bool IsWrite;
93};
94
95#define GET_IntelSubgroupsBuiltins_DECL
96#define GET_IntelSubgroupsBuiltins_IMPL
97
101};
102
103#define GET_AtomicFloatingBuiltins_DECL
104#define GET_AtomicFloatingBuiltins_IMPL
109};
110
111#define GET_GroupUniformBuiltins_DECL
112#define GET_GroupUniformBuiltins_IMPL
113
116 InstructionSet::InstructionSet Set;
117 BuiltIn::BuiltIn Value;
118};
119
120using namespace BuiltIn;
121#define GET_GetBuiltins_DECL
122#define GET_GetBuiltins_IMPL
123
126 InstructionSet::InstructionSet Set;
128};
129
130#define GET_ImageQueryBuiltins_DECL
131#define GET_ImageQueryBuiltins_IMPL
132
135 InstructionSet::InstructionSet Set;
140 FPRoundingMode::FPRoundingMode RoundingMode;
141};
142
145 InstructionSet::InstructionSet Set;
149 FPRoundingMode::FPRoundingMode RoundingMode;
150};
151
152using namespace FPRoundingMode;
153#define GET_ConvertBuiltins_DECL
154#define GET_ConvertBuiltins_IMPL
155
156using namespace InstructionSet;
157#define GET_VectorLoadStoreBuiltins_DECL
158#define GET_VectorLoadStoreBuiltins_IMPL
159
160#define GET_CLMemoryScope_DECL
161#define GET_CLSamplerAddressingMode_DECL
162#define GET_CLMemoryFenceFlags_DECL
163#define GET_ExtendedBuiltins_DECL
164#include "SPIRVGenTables.inc"
165} // namespace SPIRV
166
167//===----------------------------------------------------------------------===//
168// Misc functions for looking up builtins and veryfying requirements using
169// TableGen records
170//===----------------------------------------------------------------------===//
171
172/// Looks up the demangled builtin call in the SPIRVBuiltins.td records using
173/// the provided \p DemangledCall and specified \p Set.
174///
175/// The lookup follows the following algorithm, returning the first successful
176/// match:
177/// 1. Search with the plain demangled name (expecting a 1:1 match).
178/// 2. Search with the prefix before or suffix after the demangled name
179/// signyfying the type of the first argument.
180///
181/// \returns Wrapper around the demangled call and found builtin definition.
182static std::unique_ptr<const SPIRV::IncomingCall>
184 SPIRV::InstructionSet::InstructionSet Set,
185 Register ReturnRegister, const SPIRVType *ReturnType,
187 // Extract the builtin function name and types of arguments from the call
188 // skeleton.
189 std::string BuiltinName =
190 DemangledCall.substr(0, DemangledCall.find('(')).str();
191
192 // Account for possible "__spirv_ocl_" prefix in SPIR-V friendly LLVM IR
193 if (BuiltinName.rfind("__spirv_ocl_", 0) == 0)
194 BuiltinName = BuiltinName.substr(12);
195
196 // Check if the extracted name contains type information between angle
197 // brackets. If so, the builtin is an instantiated template - needs to have
198 // the information after angle brackets and return type removed.
199 if (BuiltinName.find('<') && BuiltinName.back() == '>') {
200 BuiltinName = BuiltinName.substr(0, BuiltinName.find('<'));
201 BuiltinName = BuiltinName.substr(BuiltinName.find_last_of(' ') + 1);
202 }
203
204 // Check if the extracted name begins with "__spirv_ImageSampleExplicitLod"
205 // contains return type information at the end "_R<type>", if so extract the
206 // plain builtin name without the type information.
207 if (StringRef(BuiltinName).contains("__spirv_ImageSampleExplicitLod") &&
208 StringRef(BuiltinName).contains("_R")) {
209 BuiltinName = BuiltinName.substr(0, BuiltinName.find("_R"));
210 }
211
212 SmallVector<StringRef, 10> BuiltinArgumentTypes;
213 StringRef BuiltinArgs =
214 DemangledCall.slice(DemangledCall.find('(') + 1, DemangledCall.find(')'));
215 BuiltinArgs.split(BuiltinArgumentTypes, ',', -1, false);
216
217 // Look up the builtin in the defined set. Start with the plain demangled
218 // name, expecting a 1:1 match in the defined builtin set.
219 const SPIRV::DemangledBuiltin *Builtin;
220 if ((Builtin = SPIRV::lookupBuiltin(BuiltinName, Set)))
221 return std::make_unique<SPIRV::IncomingCall>(
222 BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
223
224 // If the initial look up was unsuccessful and the demangled call takes at
225 // least 1 argument, add a prefix or suffix signifying the type of the first
226 // argument and repeat the search.
227 if (BuiltinArgumentTypes.size() >= 1) {
228 char FirstArgumentType = BuiltinArgumentTypes[0][0];
229 // Prefix to be added to the builtin's name for lookup.
230 // For example, OpenCL "abs" taking an unsigned value has a prefix "u_".
231 std::string Prefix;
232
233 switch (FirstArgumentType) {
234 // Unsigned:
235 case 'u':
236 if (Set == SPIRV::InstructionSet::OpenCL_std)
237 Prefix = "u_";
238 else if (Set == SPIRV::InstructionSet::GLSL_std_450)
239 Prefix = "u";
240 break;
241 // Signed:
242 case 'c':
243 case 's':
244 case 'i':
245 case 'l':
246 if (Set == SPIRV::InstructionSet::OpenCL_std)
247 Prefix = "s_";
248 else if (Set == SPIRV::InstructionSet::GLSL_std_450)
249 Prefix = "s";
250 break;
251 // Floating-point:
252 case 'f':
253 case 'd':
254 case 'h':
255 if (Set == SPIRV::InstructionSet::OpenCL_std ||
256 Set == SPIRV::InstructionSet::GLSL_std_450)
257 Prefix = "f";
258 break;
259 }
260
261 // If argument-type name prefix was added, look up the builtin again.
262 if (!Prefix.empty() &&
263 (Builtin = SPIRV::lookupBuiltin(Prefix + BuiltinName, Set)))
264 return std::make_unique<SPIRV::IncomingCall>(
265 BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
266
267 // If lookup with a prefix failed, find a suffix to be added to the
268 // builtin's name for lookup. For example, OpenCL "group_reduce_max" taking
269 // an unsigned value has a suffix "u".
270 std::string Suffix;
271
272 switch (FirstArgumentType) {
273 // Unsigned:
274 case 'u':
275 Suffix = "u";
276 break;
277 // Signed:
278 case 'c':
279 case 's':
280 case 'i':
281 case 'l':
282 Suffix = "s";
283 break;
284 // Floating-point:
285 case 'f':
286 case 'd':
287 case 'h':
288 Suffix = "f";
289 break;
290 }
291
292 // If argument-type name suffix was added, look up the builtin again.
293 if (!Suffix.empty() &&
294 (Builtin = SPIRV::lookupBuiltin(BuiltinName + Suffix, Set)))
295 return std::make_unique<SPIRV::IncomingCall>(
296 BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
297 }
298
299 // No builtin with such name was found in the set.
300 return nullptr;
301}
302
305 // We expect the following sequence of instructions:
306 // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
307 // or = G_GLOBAL_VALUE @block_literal_global
308 // %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0
309 // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
310 MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
311 assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
312 MI->getOperand(1).isReg());
313 Register BitcastReg = MI->getOperand(1).getReg();
314 MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg);
315 assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&
316 BitcastMI->getOperand(2).isReg());
317 Register ValueReg = BitcastMI->getOperand(2).getReg();
318 MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
319 return ValueMI;
320}
321
322// Return an integer constant corresponding to the given register and
323// defined in spv_track_constant.
324// TODO: maybe unify with prelegalizer pass.
326 MachineInstr *DefMI = MRI->getUniqueVRegDef(Reg);
327 assert(isSpvIntrinsic(*DefMI, Intrinsic::spv_track_constant) &&
328 DefMI->getOperand(2).isReg());
329 MachineInstr *DefMI2 = MRI->getUniqueVRegDef(DefMI->getOperand(2).getReg());
330 assert(DefMI2->getOpcode() == TargetOpcode::G_CONSTANT &&
331 DefMI2->getOperand(1).isCImm());
332 return DefMI2->getOperand(1).getCImm()->getValue().getZExtValue();
333}
334
335// Return type of the instruction result from spv_assign_type intrinsic.
336// TODO: maybe unify with prelegalizer pass.
338 MachineInstr *NextMI = MI->getNextNode();
339 if (!NextMI)
340 return nullptr;
341 if (isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_name))
342 if ((NextMI = NextMI->getNextNode()) == nullptr)
343 return nullptr;
344 Register ValueReg = MI->getOperand(0).getReg();
345 if ((!isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_type) &&
346 !isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_ptr_type)) ||
347 NextMI->getOperand(1).getReg() != ValueReg)
348 return nullptr;
349 Type *Ty = getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0);
350 assert(Ty && "Type is expected");
351 return Ty;
352}
353
354static const Type *getBlockStructType(Register ParamReg,
356 // In principle, this information should be passed to us from Clang via
357 // an elementtype attribute. However, said attribute requires that
358 // the function call be an intrinsic, which is not. Instead, we rely on being
359 // able to trace this to the declaration of a variable: OpenCL C specification
360 // section 6.12.5 should guarantee that we can do this.
362 if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
363 return MI->getOperand(1).getGlobal()->getType();
364 assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
365 "Blocks in OpenCL C must be traceable to allocation site");
366 return getMachineInstrType(MI);
367}
368
369//===----------------------------------------------------------------------===//
370// Helper functions for building misc instructions
371//===----------------------------------------------------------------------===//
372
373/// Helper function building either a resulting scalar or vector bool register
374/// depending on the expected \p ResultType.
375///
376/// \returns Tuple of the resulting register and its type.
377static std::tuple<Register, SPIRVType *>
378buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType,
380 LLT Type;
381 SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
382
383 if (ResultType->getOpcode() == SPIRV::OpTypeVector) {
384 unsigned VectorElements = ResultType->getOperand(2).getImm();
385 BoolType =
386 GR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder);
388 cast<FixedVectorType>(GR->getTypeForSPIRVType(BoolType));
389 Type = LLT::vector(LLVMVectorType->getElementCount(), 1);
390 } else {
391 Type = LLT::scalar(1);
392 }
393
394 Register ResultRegister =
396 MIRBuilder.getMRI()->setRegClass(ResultRegister, &SPIRV::IDRegClass);
397 GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());
398 return std::make_tuple(ResultRegister, BoolType);
399}
400
401/// Helper function for building either a vector or scalar select instruction
402/// depending on the expected \p ResultType.
403static bool buildSelectInst(MachineIRBuilder &MIRBuilder,
404 Register ReturnRegister, Register SourceRegister,
405 const SPIRVType *ReturnType,
407 Register TrueConst, FalseConst;
408
409 if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {
410 unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType);
412 TrueConst = GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType);
413 FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType);
414 } else {
415 TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType);
416 FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType);
417 }
418 return MIRBuilder.buildSelect(ReturnRegister, SourceRegister, TrueConst,
419 FalseConst);
420}
421
422/// Helper function for building a load instruction loading into the
423/// \p DestinationReg.
425 MachineIRBuilder &MIRBuilder,
426 SPIRVGlobalRegistry *GR, LLT LowLevelType,
427 Register DestinationReg = Register(0)) {
428 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
429 if (!DestinationReg.isValid()) {
430 DestinationReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
431 MRI->setType(DestinationReg, LLT::scalar(32));
432 GR->assignSPIRVTypeToVReg(BaseType, DestinationReg, MIRBuilder.getMF());
433 }
434 // TODO: consider using correct address space and alignment (p0 is canonical
435 // type for selection though).
437 MIRBuilder.buildLoad(DestinationReg, PtrRegister, PtrInfo, Align());
438 return DestinationReg;
439}
440
441/// Helper function for building a load instruction for loading a builtin global
442/// variable of \p BuiltinValue value.
444 MachineIRBuilder &MIRBuilder, SPIRVType *VariableType,
445 SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, LLT LLType,
446 Register Reg = Register(0), bool isConst = true, bool hasLinkageTy = true) {
447 Register NewRegister =
448 MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
449 MIRBuilder.getMRI()->setType(NewRegister,
450 LLT::pointer(0, GR->getPointerSize()));
452 VariableType, MIRBuilder, SPIRV::StorageClass::Input);
453 GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
454
455 // Set up the global OpVariable with the necessary builtin decorations.
456 Register Variable = GR->buildGlobalVariable(
457 NewRegister, PtrType, getLinkStringForBuiltIn(BuiltinValue), nullptr,
458 SPIRV::StorageClass::Input, nullptr, /* isConst= */ isConst,
459 /* HasLinkageTy */ hasLinkageTy, SPIRV::LinkageType::Import, MIRBuilder,
460 false);
461
462 // Load the value from the global variable.
463 Register LoadedRegister =
464 buildLoadInst(VariableType, Variable, MIRBuilder, GR, LLType, Reg);
465 MIRBuilder.getMRI()->setType(LoadedRegister, LLType);
466 return LoadedRegister;
467}
468
469/// Helper external function for inserting ASSIGN_TYPE instuction between \p Reg
470/// and its definition, set the new register as a destination of the definition,
471/// assign SPIRVType to both registers. If SpirvTy is provided, use it as
472/// SPIRVType in ASSIGN_TYPE, otherwise create it from \p Ty. Defined in
473/// SPIRVPreLegalizer.cpp.
474extern Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy,
475 SPIRVGlobalRegistry *GR,
476 MachineIRBuilder &MIB,
477 MachineRegisterInfo &MRI);
478
479// TODO: Move to TableGen.
480static SPIRV::MemorySemantics::MemorySemantics
481getSPIRVMemSemantics(std::memory_order MemOrder) {
482 switch (MemOrder) {
483 case std::memory_order::memory_order_relaxed:
484 return SPIRV::MemorySemantics::None;
485 case std::memory_order::memory_order_acquire:
486 return SPIRV::MemorySemantics::Acquire;
487 case std::memory_order::memory_order_release:
488 return SPIRV::MemorySemantics::Release;
489 case std::memory_order::memory_order_acq_rel:
490 return SPIRV::MemorySemantics::AcquireRelease;
491 case std::memory_order::memory_order_seq_cst:
492 return SPIRV::MemorySemantics::SequentiallyConsistent;
493 default:
494 report_fatal_error("Unknown CL memory scope");
495 }
496}
497
498static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope) {
499 switch (ClScope) {
500 case SPIRV::CLMemoryScope::memory_scope_work_item:
501 return SPIRV::Scope::Invocation;
502 case SPIRV::CLMemoryScope::memory_scope_work_group:
503 return SPIRV::Scope::Workgroup;
504 case SPIRV::CLMemoryScope::memory_scope_device:
505 return SPIRV::Scope::Device;
506 case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:
507 return SPIRV::Scope::CrossDevice;
508 case SPIRV::CLMemoryScope::memory_scope_sub_group:
509 return SPIRV::Scope::Subgroup;
510 }
511 report_fatal_error("Unknown CL memory scope");
512}
513
516 unsigned BitWidth = 32) {
517 SPIRVType *IntType = GR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
518 return GR->buildConstantInt(Val, MIRBuilder, IntType);
519}
520
521static Register buildScopeReg(Register CLScopeRegister,
522 SPIRV::Scope::Scope Scope,
523 MachineIRBuilder &MIRBuilder,
526 if (CLScopeRegister.isValid()) {
527 auto CLScope =
528 static_cast<SPIRV::CLMemoryScope>(getIConstVal(CLScopeRegister, MRI));
529 Scope = getSPIRVScope(CLScope);
530
531 if (CLScope == static_cast<unsigned>(Scope)) {
532 MRI->setRegClass(CLScopeRegister, &SPIRV::IDRegClass);
533 return CLScopeRegister;
534 }
535 }
536 return buildConstantIntReg(Scope, MIRBuilder, GR);
537}
538
539static Register buildMemSemanticsReg(Register SemanticsRegister,
540 Register PtrRegister, unsigned &Semantics,
541 MachineIRBuilder &MIRBuilder,
543 if (SemanticsRegister.isValid()) {
544 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
545 std::memory_order Order =
546 static_cast<std::memory_order>(getIConstVal(SemanticsRegister, MRI));
547 Semantics =
548 getSPIRVMemSemantics(Order) |
550
551 if (Order == Semantics) {
552 MRI->setRegClass(SemanticsRegister, &SPIRV::IDRegClass);
553 return SemanticsRegister;
554 }
555 }
556 return buildConstantIntReg(Semantics, MIRBuilder, GR);
557}
558
559static bool buildOpFromWrapper(MachineIRBuilder &MIRBuilder, unsigned Opcode,
560 const SPIRV::IncomingCall *Call,
561 Register TypeReg = Register(0)) {
562 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
563 auto MIB = MIRBuilder.buildInstr(Opcode);
564 if (TypeReg.isValid())
565 MIB.addDef(Call->ReturnRegister).addUse(TypeReg);
566 for (Register ArgReg : Call->Arguments) {
567 if (!MRI->getRegClassOrNull(ArgReg))
568 MRI->setRegClass(ArgReg, &SPIRV::IDRegClass);
569 MIB.addUse(ArgReg);
570 }
571 return true;
572}
573
574/// Helper function for translating atomic init to OpStore.
576 MachineIRBuilder &MIRBuilder) {
577 if (Call->isSpirvOp())
578 return buildOpFromWrapper(MIRBuilder, SPIRV::OpStore, Call);
579
580 assert(Call->Arguments.size() == 2 &&
581 "Need 2 arguments for atomic init translation");
582 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
583 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
584 MIRBuilder.buildInstr(SPIRV::OpStore)
585 .addUse(Call->Arguments[0])
586 .addUse(Call->Arguments[1]);
587 return true;
588}
589
590/// Helper function for building an atomic load instruction.
592 MachineIRBuilder &MIRBuilder,
594 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
595 if (Call->isSpirvOp())
596 return buildOpFromWrapper(MIRBuilder, SPIRV::OpAtomicLoad, Call, TypeReg);
597
598 Register PtrRegister = Call->Arguments[0];
599 MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass);
600 // TODO: if true insert call to __translate_ocl_memory_sccope before
601 // OpAtomicLoad and the function implementation. We can use Translator's
602 // output for transcoding/atomic_explicit_arguments.cl as an example.
603 Register ScopeRegister;
604 if (Call->Arguments.size() > 1) {
605 ScopeRegister = Call->Arguments[1];
606 MIRBuilder.getMRI()->setRegClass(ScopeRegister, &SPIRV::IDRegClass);
607 } else
608 ScopeRegister = buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
609
610 Register MemSemanticsReg;
611 if (Call->Arguments.size() > 2) {
612 // TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
613 MemSemanticsReg = Call->Arguments[2];
614 MIRBuilder.getMRI()->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass);
615 } else {
616 int Semantics =
617 SPIRV::MemorySemantics::SequentiallyConsistent |
619 MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
620 }
621
622 MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
623 .addDef(Call->ReturnRegister)
624 .addUse(TypeReg)
625 .addUse(PtrRegister)
626 .addUse(ScopeRegister)
627 .addUse(MemSemanticsReg);
628 return true;
629}
630
631/// Helper function for building an atomic store instruction.
633 MachineIRBuilder &MIRBuilder,
635 if (Call->isSpirvOp())
636 return buildOpFromWrapper(MIRBuilder, SPIRV::OpAtomicStore, Call);
637
638 Register ScopeRegister =
639 buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
640 Register PtrRegister = Call->Arguments[0];
641 MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass);
642 int Semantics =
643 SPIRV::MemorySemantics::SequentiallyConsistent |
645 Register MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
646 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
647 MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
648 .addUse(PtrRegister)
649 .addUse(ScopeRegister)
650 .addUse(MemSemanticsReg)
651 .addUse(Call->Arguments[1]);
652 return true;
653}
654
655/// Helper function for building an atomic compare-exchange instruction.
657 const SPIRV::IncomingCall *Call, const SPIRV::DemangledBuiltin *Builtin,
658 unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
659 if (Call->isSpirvOp())
660 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
661 GR->getSPIRVTypeID(Call->ReturnType));
662
663 bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");
664 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
665
666 Register ObjectPtr = Call->Arguments[0]; // Pointer (volatile A *object.)
667 Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
668 Register Desired = Call->Arguments[2]; // Value (C Desired).
669 MRI->setRegClass(ObjectPtr, &SPIRV::IDRegClass);
670 MRI->setRegClass(ExpectedArg, &SPIRV::IDRegClass);
671 MRI->setRegClass(Desired, &SPIRV::IDRegClass);
672 SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);
673 LLT DesiredLLT = MRI->getType(Desired);
674
675 assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
676 SPIRV::OpTypePointer);
677 unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
678 (void)ExpectedType;
679 assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
680 : ExpectedType == SPIRV::OpTypePointer);
681 assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
682
683 SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);
684 assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
685 auto StorageClass = static_cast<SPIRV::StorageClass::StorageClass>(
686 SpvObjectPtrTy->getOperand(1).getImm());
687 auto MemSemStorage = getMemSemanticsForStorageClass(StorageClass);
688
689 Register MemSemEqualReg;
690 Register MemSemUnequalReg;
691 uint64_t MemSemEqual =
692 IsCmpxchg
693 ? SPIRV::MemorySemantics::None
694 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
695 uint64_t MemSemUnequal =
696 IsCmpxchg
697 ? SPIRV::MemorySemantics::None
698 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
699 if (Call->Arguments.size() >= 4) {
700 assert(Call->Arguments.size() >= 5 &&
701 "Need 5+ args for explicit atomic cmpxchg");
702 auto MemOrdEq =
703 static_cast<std::memory_order>(getIConstVal(Call->Arguments[3], MRI));
704 auto MemOrdNeq =
705 static_cast<std::memory_order>(getIConstVal(Call->Arguments[4], MRI));
706 MemSemEqual = getSPIRVMemSemantics(MemOrdEq) | MemSemStorage;
707 MemSemUnequal = getSPIRVMemSemantics(MemOrdNeq) | MemSemStorage;
708 if (MemOrdEq == MemSemEqual)
709 MemSemEqualReg = Call->Arguments[3];
710 if (MemOrdNeq == MemSemEqual)
711 MemSemUnequalReg = Call->Arguments[4];
712 MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass);
713 MRI->setRegClass(Call->Arguments[4], &SPIRV::IDRegClass);
714 }
715 if (!MemSemEqualReg.isValid())
716 MemSemEqualReg = buildConstantIntReg(MemSemEqual, MIRBuilder, GR);
717 if (!MemSemUnequalReg.isValid())
718 MemSemUnequalReg = buildConstantIntReg(MemSemUnequal, MIRBuilder, GR);
719
720 Register ScopeReg;
721 auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
722 if (Call->Arguments.size() >= 6) {
723 assert(Call->Arguments.size() == 6 &&
724 "Extra args for explicit atomic cmpxchg");
725 auto ClScope = static_cast<SPIRV::CLMemoryScope>(
726 getIConstVal(Call->Arguments[5], MRI));
727 Scope = getSPIRVScope(ClScope);
728 if (ClScope == static_cast<unsigned>(Scope))
729 ScopeReg = Call->Arguments[5];
730 MRI->setRegClass(Call->Arguments[5], &SPIRV::IDRegClass);
731 }
732 if (!ScopeReg.isValid())
733 ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
734
735 Register Expected = IsCmpxchg
736 ? ExpectedArg
737 : buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,
738 GR, LLT::scalar(32));
739 MRI->setType(Expected, DesiredLLT);
740 Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT)
741 : Call->ReturnRegister;
742 if (!MRI->getRegClassOrNull(Tmp))
743 MRI->setRegClass(Tmp, &SPIRV::IDRegClass);
744 GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
745
746 SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
747 MIRBuilder.buildInstr(Opcode)
748 .addDef(Tmp)
749 .addUse(GR->getSPIRVTypeID(IntTy))
750 .addUse(ObjectPtr)
751 .addUse(ScopeReg)
752 .addUse(MemSemEqualReg)
753 .addUse(MemSemUnequalReg)
754 .addUse(Desired)
756 if (!IsCmpxchg) {
757 MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);
758 MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected);
759 }
760 return true;
761}
762
763/// Helper function for building an atomic load instruction.
764static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
765 MachineIRBuilder &MIRBuilder,
767 if (Call->isSpirvOp())
768 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
769 GR->getSPIRVTypeID(Call->ReturnType));
770
771 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
772 Register ScopeRegister =
773 Call->Arguments.size() >= 4 ? Call->Arguments[3] : Register();
774
775 assert(Call->Arguments.size() <= 4 &&
776 "Too many args for explicit atomic RMW");
777 ScopeRegister = buildScopeReg(ScopeRegister, SPIRV::Scope::Workgroup,
778 MIRBuilder, GR, MRI);
779
780 Register PtrRegister = Call->Arguments[0];
781 unsigned Semantics = SPIRV::MemorySemantics::None;
782 MRI->setRegClass(PtrRegister, &SPIRV::IDRegClass);
783 Register MemSemanticsReg =
784 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
785 MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
786 Semantics, MIRBuilder, GR);
787 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
788 MIRBuilder.buildInstr(Opcode)
789 .addDef(Call->ReturnRegister)
790 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
791 .addUse(PtrRegister)
792 .addUse(ScopeRegister)
793 .addUse(MemSemanticsReg)
794 .addUse(Call->Arguments[1]);
795 return true;
796}
797
798/// Helper function for building an atomic floating-type instruction.
800 unsigned Opcode,
801 MachineIRBuilder &MIRBuilder,
803 assert(Call->Arguments.size() == 4 &&
804 "Wrong number of atomic floating-type builtin");
805
806 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
807
808 Register PtrReg = Call->Arguments[0];
809 MRI->setRegClass(PtrReg, &SPIRV::IDRegClass);
810
811 Register ScopeReg = Call->Arguments[1];
812 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass);
813
814 Register MemSemanticsReg = Call->Arguments[2];
815 MRI->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass);
816
817 Register ValueReg = Call->Arguments[3];
818 MRI->setRegClass(ValueReg, &SPIRV::IDRegClass);
819
820 MIRBuilder.buildInstr(Opcode)
821 .addDef(Call->ReturnRegister)
822 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
823 .addUse(PtrReg)
824 .addUse(ScopeReg)
825 .addUse(MemSemanticsReg)
826 .addUse(ValueReg);
827 return true;
828}
829
830/// Helper function for building atomic flag instructions (e.g.
831/// OpAtomicFlagTestAndSet).
833 unsigned Opcode, MachineIRBuilder &MIRBuilder,
835 bool IsSet = Opcode == SPIRV::OpAtomicFlagTestAndSet;
836 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
837 if (Call->isSpirvOp())
838 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
839 IsSet ? TypeReg : Register(0));
840
841 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
842 Register PtrRegister = Call->Arguments[0];
843 unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;
844 Register MemSemanticsReg =
845 Call->Arguments.size() >= 2 ? Call->Arguments[1] : Register();
846 MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
847 Semantics, MIRBuilder, GR);
848
849 assert((Opcode != SPIRV::OpAtomicFlagClear ||
850 (Semantics != SPIRV::MemorySemantics::Acquire &&
851 Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&
852 "Invalid memory order argument!");
853
854 Register ScopeRegister =
855 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
856 ScopeRegister =
857 buildScopeReg(ScopeRegister, SPIRV::Scope::Device, MIRBuilder, GR, MRI);
858
859 auto MIB = MIRBuilder.buildInstr(Opcode);
860 if (IsSet)
861 MIB.addDef(Call->ReturnRegister).addUse(TypeReg);
862
863 MIB.addUse(PtrRegister).addUse(ScopeRegister).addUse(MemSemanticsReg);
864 return true;
865}
866
867/// Helper function for building barriers, i.e., memory/control ordering
868/// operations.
869static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
870 MachineIRBuilder &MIRBuilder,
872 if (Call->isSpirvOp())
873 return buildOpFromWrapper(MIRBuilder, Opcode, Call);
874
875 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
876 unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI);
877 unsigned MemSemantics = SPIRV::MemorySemantics::None;
878
879 if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
880 MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
881
882 if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
883 MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;
884
885 if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)
886 MemSemantics |= SPIRV::MemorySemantics::ImageMemory;
887
888 if (Opcode == SPIRV::OpMemoryBarrier) {
889 std::memory_order MemOrder =
890 static_cast<std::memory_order>(getIConstVal(Call->Arguments[1], MRI));
891 MemSemantics = getSPIRVMemSemantics(MemOrder) | MemSemantics;
892 } else {
893 MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
894 }
895
896 Register MemSemanticsReg;
897 if (MemFlags == MemSemantics) {
898 MemSemanticsReg = Call->Arguments[0];
899 MRI->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass);
900 } else
901 MemSemanticsReg = buildConstantIntReg(MemSemantics, MIRBuilder, GR);
902
903 Register ScopeReg;
904 SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
905 SPIRV::Scope::Scope MemScope = Scope;
906 if (Call->Arguments.size() >= 2) {
907 assert(
908 ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
909 (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
910 "Extra args for explicitly scoped barrier");
911 Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]
912 : Call->Arguments[1];
913 SPIRV::CLMemoryScope CLScope =
914 static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg, MRI));
915 MemScope = getSPIRVScope(CLScope);
916 if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
917 (Opcode == SPIRV::OpMemoryBarrier))
918 Scope = MemScope;
919
920 if (CLScope == static_cast<unsigned>(Scope)) {
921 ScopeReg = Call->Arguments[1];
922 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass);
923 }
924 }
925
926 if (!ScopeReg.isValid())
927 ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
928
929 auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);
930 if (Opcode != SPIRV::OpMemoryBarrier)
931 MIB.addUse(buildConstantIntReg(MemScope, MIRBuilder, GR));
932 MIB.addUse(MemSemanticsReg);
933 return true;
934}
935
936static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
937 switch (dim) {
938 case SPIRV::Dim::DIM_1D:
939 case SPIRV::Dim::DIM_Buffer:
940 return 1;
941 case SPIRV::Dim::DIM_2D:
942 case SPIRV::Dim::DIM_Cube:
943 case SPIRV::Dim::DIM_Rect:
944 return 2;
945 case SPIRV::Dim::DIM_3D:
946 return 3;
947 default:
948 report_fatal_error("Cannot get num components for given Dim");
949 }
950}
951
952/// Helper function for obtaining the number of size components.
953static unsigned getNumSizeComponents(SPIRVType *imgType) {
954 assert(imgType->getOpcode() == SPIRV::OpTypeImage);
955 auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm());
956 unsigned numComps = getNumComponentsForDim(dim);
957 bool arrayed = imgType->getOperand(4).getImm() == 1;
958 return arrayed ? numComps + 1 : numComps;
959}
960
961//===----------------------------------------------------------------------===//
962// Implementation functions for each builtin group
963//===----------------------------------------------------------------------===//
964
965static bool generateExtInst(const SPIRV::IncomingCall *Call,
966 MachineIRBuilder &MIRBuilder,
968 // Lookup the extended instruction number in the TableGen records.
969 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
971 SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;
972
973 // Build extended instruction.
974 auto MIB =
975 MIRBuilder.buildInstr(SPIRV::OpExtInst)
976 .addDef(Call->ReturnRegister)
977 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
978 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
979 .addImm(Number);
980
981 for (auto Argument : Call->Arguments)
982 MIB.addUse(Argument);
983 return true;
984}
985
987 MachineIRBuilder &MIRBuilder,
989 // Lookup the instruction opcode in the TableGen records.
990 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
991 unsigned Opcode =
992 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
993
994 Register CompareRegister;
995 SPIRVType *RelationType;
996 std::tie(CompareRegister, RelationType) =
997 buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
998
999 // Build relational instruction.
1000 auto MIB = MIRBuilder.buildInstr(Opcode)
1001 .addDef(CompareRegister)
1002 .addUse(GR->getSPIRVTypeID(RelationType));
1003
1004 for (auto Argument : Call->Arguments)
1005 MIB.addUse(Argument);
1006
1007 // Build select instruction.
1008 return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,
1009 Call->ReturnType, GR);
1010}
1011
1013 MachineIRBuilder &MIRBuilder,
1014 SPIRVGlobalRegistry *GR) {
1015 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1016 const SPIRV::GroupBuiltin *GroupBuiltin =
1017 SPIRV::lookupGroupBuiltin(Builtin->Name);
1018
1019 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1020 if (Call->isSpirvOp()) {
1021 if (GroupBuiltin->NoGroupOperation)
1022 return buildOpFromWrapper(MIRBuilder, GroupBuiltin->Opcode, Call,
1023 GR->getSPIRVTypeID(Call->ReturnType));
1024
1025 // Group Operation is a literal
1026 Register GroupOpReg = Call->Arguments[1];
1027 const MachineInstr *MI = getDefInstrMaybeConstant(GroupOpReg, MRI);
1028 if (!MI || MI->getOpcode() != TargetOpcode::G_CONSTANT)
1030 "Group Operation parameter must be an integer constant");
1031 uint64_t GrpOp = MI->getOperand(1).getCImm()->getValue().getZExtValue();
1032 Register ScopeReg = Call->Arguments[0];
1033 if (!MRI->getRegClassOrNull(ScopeReg))
1034 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass);
1035 Register ValueReg = Call->Arguments[2];
1036 if (!MRI->getRegClassOrNull(ValueReg))
1037 MRI->setRegClass(ValueReg, &SPIRV::IDRegClass);
1038 MIRBuilder.buildInstr(GroupBuiltin->Opcode)
1039 .addDef(Call->ReturnRegister)
1040 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1041 .addUse(ScopeReg)
1042 .addImm(GrpOp)
1043 .addUse(ValueReg);
1044 return true;
1045 }
1046
1047 Register Arg0;
1048 if (GroupBuiltin->HasBoolArg) {
1049 Register ConstRegister = Call->Arguments[0];
1050 auto ArgInstruction = getDefInstrMaybeConstant(ConstRegister, MRI);
1051 (void)ArgInstruction;
1052 // TODO: support non-constant bool values.
1053 assert(ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT &&
1054 "Only constant bool value args are supported");
1055 if (GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() !=
1056 SPIRV::OpTypeBool)
1057 Arg0 = GR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder,
1058 GR->getOrCreateSPIRVBoolType(MIRBuilder));
1059 }
1060
1061 Register GroupResultRegister = Call->ReturnRegister;
1062 SPIRVType *GroupResultType = Call->ReturnType;
1063
1064 // TODO: maybe we need to check whether the result type is already boolean
1065 // and in this case do not insert select instruction.
1066 const bool HasBoolReturnTy =
1067 GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
1068 GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
1069 GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
1070
1071 if (HasBoolReturnTy)
1072 std::tie(GroupResultRegister, GroupResultType) =
1073 buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
1074
1075 auto Scope = Builtin->Name.starts_with("sub_group") ? SPIRV::Scope::Subgroup
1076 : SPIRV::Scope::Workgroup;
1077 Register ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
1078
1079 // Build work/sub group instruction.
1080 auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
1081 .addDef(GroupResultRegister)
1082 .addUse(GR->getSPIRVTypeID(GroupResultType))
1083 .addUse(ScopeRegister);
1084
1085 if (!GroupBuiltin->NoGroupOperation)
1086 MIB.addImm(GroupBuiltin->GroupOperation);
1087 if (Call->Arguments.size() > 0) {
1088 MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);
1089 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1090 for (unsigned i = 1; i < Call->Arguments.size(); i++) {
1091 MIB.addUse(Call->Arguments[i]);
1092 MRI->setRegClass(Call->Arguments[i], &SPIRV::IDRegClass);
1093 }
1094 }
1095
1096 // Build select instruction.
1097 if (HasBoolReturnTy)
1098 buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
1099 Call->ReturnType, GR);
1100 return true;
1101}
1102
1104 MachineIRBuilder &MIRBuilder,
1105 SPIRVGlobalRegistry *GR) {
1106 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1107 MachineFunction &MF = MIRBuilder.getMF();
1108 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1109 if (!ST->canUseExtension(SPIRV::Extension::SPV_INTEL_subgroups)) {
1110 std::string DiagMsg = std::string(Builtin->Name) +
1111 ": the builtin requires the following SPIR-V "
1112 "extension: SPV_INTEL_subgroups";
1113 report_fatal_error(DiagMsg.c_str(), false);
1114 }
1115 const SPIRV::IntelSubgroupsBuiltin *IntelSubgroups =
1116 SPIRV::lookupIntelSubgroupsBuiltin(Builtin->Name);
1117
1118 uint32_t OpCode = IntelSubgroups->Opcode;
1119 if (Call->isSpirvOp()) {
1120 bool IsSet = OpCode != SPIRV::OpSubgroupBlockWriteINTEL &&
1121 OpCode != SPIRV::OpSubgroupImageBlockWriteINTEL;
1122 return buildOpFromWrapper(MIRBuilder, OpCode, Call,
1123 IsSet ? GR->getSPIRVTypeID(Call->ReturnType)
1124 : Register(0));
1125 }
1126
1127 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1128 if (IntelSubgroups->IsBlock) {
1129 // Minimal number or arguments set in TableGen records is 1
1130 if (SPIRVType *Arg0Type = GR->getSPIRVTypeForVReg(Call->Arguments[0])) {
1131 if (Arg0Type->getOpcode() == SPIRV::OpTypeImage) {
1132 // TODO: add required validation from the specification:
1133 // "'Image' must be an object whose type is OpTypeImage with a 'Sampled'
1134 // operand of 0 or 2. If the 'Sampled' operand is 2, then some
1135 // dimensions require a capability."
1136 switch (OpCode) {
1137 case SPIRV::OpSubgroupBlockReadINTEL:
1138 OpCode = SPIRV::OpSubgroupImageBlockReadINTEL;
1139 break;
1140 case SPIRV::OpSubgroupBlockWriteINTEL:
1141 OpCode = SPIRV::OpSubgroupImageBlockWriteINTEL;
1142 break;
1143 }
1144 }
1145 }
1146 }
1147
1148 // TODO: opaque pointers types should be eventually resolved in such a way
1149 // that validation of block read is enabled with respect to the following
1150 // specification requirement:
1151 // "'Result Type' may be a scalar or vector type, and its component type must
1152 // be equal to the type pointed to by 'Ptr'."
1153 // For example, function parameter type should not be default i8 pointer, but
1154 // depend on the result type of the instruction where it is used as a pointer
1155 // argument of OpSubgroupBlockReadINTEL
1156
1157 // Build Intel subgroups instruction
1159 IntelSubgroups->IsWrite
1160 ? MIRBuilder.buildInstr(OpCode)
1161 : MIRBuilder.buildInstr(OpCode)
1162 .addDef(Call->ReturnRegister)
1163 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1164 for (size_t i = 0; i < Call->Arguments.size(); ++i) {
1165 MIB.addUse(Call->Arguments[i]);
1166 MRI->setRegClass(Call->Arguments[i], &SPIRV::IDRegClass);
1167 }
1168
1169 return true;
1170}
1171
1173 MachineIRBuilder &MIRBuilder,
1174 SPIRVGlobalRegistry *GR) {
1175 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1176 MachineFunction &MF = MIRBuilder.getMF();
1177 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1178 if (!ST->canUseExtension(
1179 SPIRV::Extension::SPV_KHR_uniform_group_instructions)) {
1180 std::string DiagMsg = std::string(Builtin->Name) +
1181 ": the builtin requires the following SPIR-V "
1182 "extension: SPV_KHR_uniform_group_instructions";
1183 report_fatal_error(DiagMsg.c_str(), false);
1184 }
1185 const SPIRV::GroupUniformBuiltin *GroupUniform =
1186 SPIRV::lookupGroupUniformBuiltin(Builtin->Name);
1187 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1188
1189 Register GroupResultReg = Call->ReturnRegister;
1190 MRI->setRegClass(GroupResultReg, &SPIRV::IDRegClass);
1191
1192 // Scope
1193 Register ScopeReg = Call->Arguments[0];
1194 MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass);
1195
1196 // Group Operation
1197 Register ConstGroupOpReg = Call->Arguments[1];
1198 const MachineInstr *Const = getDefInstrMaybeConstant(ConstGroupOpReg, MRI);
1199 if (!Const || Const->getOpcode() != TargetOpcode::G_CONSTANT)
1201 "expect a constant group operation for a uniform group instruction",
1202 false);
1203 const MachineOperand &ConstOperand = Const->getOperand(1);
1204 if (!ConstOperand.isCImm())
1205 report_fatal_error("uniform group instructions: group operation must be an "
1206 "integer constant",
1207 false);
1208
1209 // Value
1210 Register ValueReg = Call->Arguments[2];
1211 MRI->setRegClass(ValueReg, &SPIRV::IDRegClass);
1212
1213 auto MIB = MIRBuilder.buildInstr(GroupUniform->Opcode)
1214 .addDef(GroupResultReg)
1215 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1216 .addUse(ScopeReg);
1217 addNumImm(ConstOperand.getCImm()->getValue(), MIB);
1218 MIB.addUse(ValueReg);
1219
1220 return true;
1221}
1222
1224 MachineIRBuilder &MIRBuilder,
1225 SPIRVGlobalRegistry *GR) {
1226 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1227 MachineFunction &MF = MIRBuilder.getMF();
1228 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1229 if (!ST->canUseExtension(SPIRV::Extension::SPV_KHR_shader_clock)) {
1230 std::string DiagMsg = std::string(Builtin->Name) +
1231 ": the builtin requires the following SPIR-V "
1232 "extension: SPV_KHR_shader_clock";
1233 report_fatal_error(DiagMsg.c_str(), false);
1234 }
1235
1236 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1237 Register ResultReg = Call->ReturnRegister;
1238 MRI->setRegClass(ResultReg, &SPIRV::IDRegClass);
1239
1240 // Deduce the `Scope` operand from the builtin function name.
1241 SPIRV::Scope::Scope ScopeArg =
1243 .EndsWith("device", SPIRV::Scope::Scope::Device)
1244 .EndsWith("work_group", SPIRV::Scope::Scope::Workgroup)
1245 .EndsWith("sub_group", SPIRV::Scope::Scope::Subgroup);
1246 Register ScopeReg = buildConstantIntReg(ScopeArg, MIRBuilder, GR);
1247
1248 MIRBuilder.buildInstr(SPIRV::OpReadClockKHR)
1249 .addDef(ResultReg)
1250 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1251 .addUse(ScopeReg);
1252
1253 return true;
1254}
1255
1256// These queries ask for a single size_t result for a given dimension index, e.g
1257// size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
1258// these values are all vec3 types, so we need to extract the correct index or
1259// return defaultVal (0 or 1 depending on the query). We also handle extending
1260// or tuncating in case size_t does not match the expected result type's
1261// bitwidth.
1262//
1263// For a constant index >= 3 we generate:
1264// %res = OpConstant %SizeT 0
1265//
1266// For other indices we generate:
1267// %g = OpVariable %ptr_V3_SizeT Input
1268// OpDecorate %g BuiltIn XXX
1269// OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"
1270// OpDecorate %g Constant
1271// %loadedVec = OpLoad %V3_SizeT %g
1272//
1273// Then, if the index is constant < 3, we generate:
1274// %res = OpCompositeExtract %SizeT %loadedVec idx
1275// If the index is dynamic, we generate:
1276// %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
1277// %cmp = OpULessThan %bool %idx %const_3
1278// %res = OpSelect %SizeT %cmp %tmp %const_0
1279//
1280// If the bitwidth of %res does not match the expected return type, we add an
1281// extend or truncate.
1283 MachineIRBuilder &MIRBuilder,
1285 SPIRV::BuiltIn::BuiltIn BuiltinValue,
1286 uint64_t DefaultValue) {
1287 Register IndexRegister = Call->Arguments[0];
1288 const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();
1289 const unsigned PointerSize = GR->getPointerSize();
1290 const SPIRVType *PointerSizeType =
1291 GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
1292 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1293 auto IndexInstruction = getDefInstrMaybeConstant(IndexRegister, MRI);
1294
1295 // Set up the final register to do truncation or extension on at the end.
1296 Register ToTruncate = Call->ReturnRegister;
1297
1298 // If the index is constant, we can statically determine if it is in range.
1299 bool IsConstantIndex =
1300 IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
1301
1302 // If it's out of range (max dimension is 3), we can just return the constant
1303 // default value (0 or 1 depending on which query function).
1304 if (IsConstantIndex && getIConstVal(IndexRegister, MRI) >= 3) {
1305 Register DefaultReg = Call->ReturnRegister;
1306 if (PointerSize != ResultWidth) {
1307 DefaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
1308 MRI->setRegClass(DefaultReg, &SPIRV::IDRegClass);
1309 GR->assignSPIRVTypeToVReg(PointerSizeType, DefaultReg,
1310 MIRBuilder.getMF());
1311 ToTruncate = DefaultReg;
1312 }
1313 auto NewRegister =
1314 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
1315 MIRBuilder.buildCopy(DefaultReg, NewRegister);
1316 } else { // If it could be in range, we need to load from the given builtin.
1317 auto Vec3Ty =
1318 GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);
1319 Register LoadedVector =
1320 buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,
1321 LLT::fixed_vector(3, PointerSize));
1322 // Set up the vreg to extract the result to (possibly a new temporary one).
1323 Register Extracted = Call->ReturnRegister;
1324 if (!IsConstantIndex || PointerSize != ResultWidth) {
1325 Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
1326 MRI->setRegClass(Extracted, &SPIRV::IDRegClass);
1327 GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());
1328 }
1329 // Use Intrinsic::spv_extractelt so dynamic vs static extraction is
1330 // handled later: extr = spv_extractelt LoadedVector, IndexRegister.
1331 MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
1332 Intrinsic::spv_extractelt, ArrayRef<Register>{Extracted}, true, false);
1333 ExtractInst.addUse(LoadedVector).addUse(IndexRegister);
1334
1335 // If the index is dynamic, need check if it's < 3, and then use a select.
1336 if (!IsConstantIndex) {
1337 insertAssignInstr(Extracted, nullptr, PointerSizeType, GR, MIRBuilder,
1338 *MRI);
1339
1340 auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);
1341 auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
1342
1343 Register CompareRegister =
1344 MRI->createGenericVirtualRegister(LLT::scalar(1));
1345 MRI->setRegClass(CompareRegister, &SPIRV::IDRegClass);
1346 GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
1347
1348 // Use G_ICMP to check if idxVReg < 3.
1349 MIRBuilder.buildICmp(CmpInst::ICMP_ULT, CompareRegister, IndexRegister,
1350 GR->buildConstantInt(3, MIRBuilder, IndexType));
1351
1352 // Get constant for the default value (0 or 1 depending on which
1353 // function).
1354 Register DefaultRegister =
1355 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
1356
1357 // Get a register for the selection result (possibly a new temporary one).
1358 Register SelectionResult = Call->ReturnRegister;
1359 if (PointerSize != ResultWidth) {
1360 SelectionResult =
1361 MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
1362 MRI->setRegClass(SelectionResult, &SPIRV::IDRegClass);
1363 GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
1364 MIRBuilder.getMF());
1365 }
1366 // Create the final G_SELECT to return the extracted value or the default.
1367 MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,
1368 DefaultRegister);
1369 ToTruncate = SelectionResult;
1370 } else {
1371 ToTruncate = Extracted;
1372 }
1373 }
1374 // Alter the result's bitwidth if it does not match the SizeT value extracted.
1375 if (PointerSize != ResultWidth)
1376 MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);
1377 return true;
1378}
1379
1381 MachineIRBuilder &MIRBuilder,
1382 SPIRVGlobalRegistry *GR) {
1383 // Lookup the builtin variable record.
1384 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1385 SPIRV::BuiltIn::BuiltIn Value =
1386 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
1387
1388 if (Value == SPIRV::BuiltIn::GlobalInvocationId)
1389 return genWorkgroupQuery(Call, MIRBuilder, GR, Value, 0);
1390
1391 // Build a load instruction for the builtin variable.
1392 unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);
1393 LLT LLType;
1394 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
1395 LLType =
1396 LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth);
1397 else
1398 LLType = LLT::scalar(BitWidth);
1399
1400 return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value,
1401 LLType, Call->ReturnRegister);
1402}
1403
1405 MachineIRBuilder &MIRBuilder,
1406 SPIRVGlobalRegistry *GR) {
1407 // Lookup the instruction opcode in the TableGen records.
1408 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1409 unsigned Opcode =
1410 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1411
1412 switch (Opcode) {
1413 case SPIRV::OpStore:
1414 return buildAtomicInitInst(Call, MIRBuilder);
1415 case SPIRV::OpAtomicLoad:
1416 return buildAtomicLoadInst(Call, MIRBuilder, GR);
1417 case SPIRV::OpAtomicStore:
1418 return buildAtomicStoreInst(Call, MIRBuilder, GR);
1419 case SPIRV::OpAtomicCompareExchange:
1420 case SPIRV::OpAtomicCompareExchangeWeak:
1421 return buildAtomicCompareExchangeInst(Call, Builtin, Opcode, MIRBuilder,
1422 GR);
1423 case SPIRV::OpAtomicIAdd:
1424 case SPIRV::OpAtomicISub:
1425 case SPIRV::OpAtomicOr:
1426 case SPIRV::OpAtomicXor:
1427 case SPIRV::OpAtomicAnd:
1428 case SPIRV::OpAtomicExchange:
1429 return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
1430 case SPIRV::OpMemoryBarrier:
1431 return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
1432 case SPIRV::OpAtomicFlagTestAndSet:
1433 case SPIRV::OpAtomicFlagClear:
1434 return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);
1435 default:
1436 return false;
1437 }
1438}
1439
1441 MachineIRBuilder &MIRBuilder,
1442 SPIRVGlobalRegistry *GR) {
1443 // Lookup the instruction opcode in the TableGen records.
1444 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1445 unsigned Opcode = SPIRV::lookupAtomicFloatingBuiltin(Builtin->Name)->Opcode;
1446
1447 switch (Opcode) {
1448 case SPIRV::OpAtomicFAddEXT:
1449 case SPIRV::OpAtomicFMinEXT:
1450 case SPIRV::OpAtomicFMaxEXT:
1451 return buildAtomicFloatingRMWInst(Call, Opcode, MIRBuilder, GR);
1452 default:
1453 return false;
1454 }
1455}
1456
1458 MachineIRBuilder &MIRBuilder,
1459 SPIRVGlobalRegistry *GR) {
1460 // Lookup the instruction opcode in the TableGen records.
1461 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1462 unsigned Opcode =
1463 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1464
1465 return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
1466}
1467
1469 MachineIRBuilder &MIRBuilder) {
1470 MIRBuilder.buildInstr(TargetOpcode::G_ADDRSPACE_CAST)
1471 .addDef(Call->ReturnRegister)
1472 .addUse(Call->Arguments[0]);
1473 return true;
1474}
1475
1477 MachineIRBuilder &MIRBuilder,
1478 SPIRVGlobalRegistry *GR) {
1479 unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();
1480 bool IsVec = Opcode == SPIRV::OpTypeVector;
1481 // Use OpDot only in case of vector args and OpFMul in case of scalar args.
1482 MIRBuilder.buildInstr(IsVec ? SPIRV::OpDot : SPIRV::OpFMulS)
1483 .addDef(Call->ReturnRegister)
1484 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1485 .addUse(Call->Arguments[0])
1486 .addUse(Call->Arguments[1]);
1487 return true;
1488}
1489
1491 MachineIRBuilder &MIRBuilder,
1492 SPIRVGlobalRegistry *GR) {
1493 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1494 SPIRV::BuiltIn::BuiltIn Value =
1495 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
1496
1497 // For now, we only support a single Wave intrinsic with a single return type.
1498 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt);
1499 LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(Call->ReturnType));
1500
1502 MIRBuilder, Call->ReturnType, GR, Value, LLType, Call->ReturnRegister,
1503 /* isConst= */ false, /* hasLinkageTy= */ false);
1504}
1505
1507 MachineIRBuilder &MIRBuilder,
1508 SPIRVGlobalRegistry *GR) {
1509 // Lookup the builtin record.
1510 SPIRV::BuiltIn::BuiltIn Value =
1511 SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;
1512 uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||
1513 Value == SPIRV::BuiltIn::WorkgroupSize ||
1514 Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
1515 return genWorkgroupQuery(Call, MIRBuilder, GR, Value, IsDefault ? 1 : 0);
1516}
1517
1519 MachineIRBuilder &MIRBuilder,
1520 SPIRVGlobalRegistry *GR) {
1521 // Lookup the image size query component number in the TableGen records.
1522 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1523 uint32_t Component =
1524 SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;
1525 // Query result may either be a vector or a scalar. If return type is not a
1526 // vector, expect only a single size component. Otherwise get the number of
1527 // expected components.
1528 SPIRVType *RetTy = Call->ReturnType;
1529 unsigned NumExpectedRetComponents = RetTy->getOpcode() == SPIRV::OpTypeVector
1530 ? RetTy->getOperand(2).getImm()
1531 : 1;
1532 // Get the actual number of query result/size components.
1533 SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1534 unsigned NumActualRetComponents = getNumSizeComponents(ImgType);
1535 Register QueryResult = Call->ReturnRegister;
1536 SPIRVType *QueryResultType = Call->ReturnType;
1537 if (NumExpectedRetComponents != NumActualRetComponents) {
1538 QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
1539 LLT::fixed_vector(NumActualRetComponents, 32));
1540 MIRBuilder.getMRI()->setRegClass(QueryResult, &SPIRV::IDRegClass);
1541 SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1542 QueryResultType = GR->getOrCreateSPIRVVectorType(
1543 IntTy, NumActualRetComponents, MIRBuilder);
1544 GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());
1545 }
1546 bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;
1547 unsigned Opcode =
1548 IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
1549 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1550 auto MIB = MIRBuilder.buildInstr(Opcode)
1551 .addDef(QueryResult)
1552 .addUse(GR->getSPIRVTypeID(QueryResultType))
1553 .addUse(Call->Arguments[0]);
1554 if (!IsDimBuf)
1555 MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Lod id.
1556 if (NumExpectedRetComponents == NumActualRetComponents)
1557 return true;
1558 if (NumExpectedRetComponents == 1) {
1559 // Only 1 component is expected, build OpCompositeExtract instruction.
1560 unsigned ExtractedComposite =
1561 Component == 3 ? NumActualRetComponents - 1 : Component;
1562 assert(ExtractedComposite < NumActualRetComponents &&
1563 "Invalid composite index!");
1564 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
1565 SPIRVType *NewType = nullptr;
1566 if (QueryResultType->getOpcode() == SPIRV::OpTypeVector) {
1567 Register NewTypeReg = QueryResultType->getOperand(1).getReg();
1568 if (TypeReg != NewTypeReg &&
1569 (NewType = GR->getSPIRVTypeForVReg(NewTypeReg)) != nullptr)
1570 TypeReg = NewTypeReg;
1571 }
1572 MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1573 .addDef(Call->ReturnRegister)
1574 .addUse(TypeReg)
1575 .addUse(QueryResult)
1576 .addImm(ExtractedComposite);
1577 if (NewType != nullptr)
1578 insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder,
1579 MIRBuilder.getMF().getRegInfo());
1580 } else {
1581 // More than 1 component is expected, fill a new vector.
1582 auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)
1583 .addDef(Call->ReturnRegister)
1584 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1585 .addUse(QueryResult)
1586 .addUse(QueryResult);
1587 for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
1588 MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);
1589 }
1590 return true;
1591}
1592
1594 MachineIRBuilder &MIRBuilder,
1595 SPIRVGlobalRegistry *GR) {
1596 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
1597 "Image samples query result must be of int type!");
1598
1599 // Lookup the instruction opcode in the TableGen records.
1600 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1601 unsigned Opcode =
1602 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1603
1604 Register Image = Call->Arguments[0];
1605 MIRBuilder.getMRI()->setRegClass(Image, &SPIRV::IDRegClass);
1606 SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
1607 GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
1608 (void)ImageDimensionality;
1609
1610 switch (Opcode) {
1611 case SPIRV::OpImageQuerySamples:
1612 assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
1613 "Image must be of 2D dimensionality");
1614 break;
1615 case SPIRV::OpImageQueryLevels:
1616 assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
1617 ImageDimensionality == SPIRV::Dim::DIM_2D ||
1618 ImageDimensionality == SPIRV::Dim::DIM_3D ||
1619 ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
1620 "Image must be of 1D/2D/3D/Cube dimensionality");
1621 break;
1622 }
1623
1624 MIRBuilder.buildInstr(Opcode)
1625 .addDef(Call->ReturnRegister)
1626 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1627 .addUse(Image);
1628 return true;
1629}
1630
1631// TODO: Move to TableGen.
1632static SPIRV::SamplerAddressingMode::SamplerAddressingMode
1634 switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
1635 case SPIRV::CLK_ADDRESS_CLAMP:
1636 return SPIRV::SamplerAddressingMode::Clamp;
1637 case SPIRV::CLK_ADDRESS_CLAMP_TO_EDGE:
1638 return SPIRV::SamplerAddressingMode::ClampToEdge;
1639 case SPIRV::CLK_ADDRESS_REPEAT:
1640 return SPIRV::SamplerAddressingMode::Repeat;
1641 case SPIRV::CLK_ADDRESS_MIRRORED_REPEAT:
1642 return SPIRV::SamplerAddressingMode::RepeatMirrored;
1643 case SPIRV::CLK_ADDRESS_NONE:
1644 return SPIRV::SamplerAddressingMode::None;
1645 default:
1646 report_fatal_error("Unknown CL address mode");
1647 }
1648}
1649
1650static unsigned getSamplerParamFromBitmask(unsigned Bitmask) {
1651 return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;
1652}
1653
1654static SPIRV::SamplerFilterMode::SamplerFilterMode
1656 if (Bitmask & SPIRV::CLK_FILTER_LINEAR)
1657 return SPIRV::SamplerFilterMode::Linear;
1658 if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
1659 return SPIRV::SamplerFilterMode::Nearest;
1660 return SPIRV::SamplerFilterMode::Nearest;
1661}
1662
1663static bool generateReadImageInst(const StringRef DemangledCall,
1664 const SPIRV::IncomingCall *Call,
1665 MachineIRBuilder &MIRBuilder,
1666 SPIRVGlobalRegistry *GR) {
1667 Register Image = Call->Arguments[0];
1668 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1669 MRI->setRegClass(Image, &SPIRV::IDRegClass);
1670 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
1671 bool HasOclSampler = DemangledCall.contains_insensitive("ocl_sampler");
1672 bool HasMsaa = DemangledCall.contains_insensitive("msaa");
1673 if (HasOclSampler || HasMsaa)
1674 MRI->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
1675 if (HasOclSampler) {
1676 Register Sampler = Call->Arguments[1];
1677
1678 if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
1679 getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) {
1680 uint64_t SamplerMask = getIConstVal(Sampler, MRI);
1681 Sampler = GR->buildConstantSampler(
1683 getSamplerParamFromBitmask(SamplerMask),
1684 getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder,
1685 GR->getSPIRVTypeForVReg(Sampler));
1686 }
1687 SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1688 SPIRVType *SampledImageType =
1689 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1690 Register SampledImage = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1691
1692 MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1693 .addDef(SampledImage)
1694 .addUse(GR->getSPIRVTypeID(SampledImageType))
1695 .addUse(Image)
1696 .addUse(Sampler);
1697
1699 MIRBuilder);
1700 SPIRVType *TempType = Call->ReturnType;
1701 bool NeedsExtraction = false;
1702 if (TempType->getOpcode() != SPIRV::OpTypeVector) {
1703 TempType =
1704 GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);
1705 NeedsExtraction = true;
1706 }
1707 LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(TempType));
1708 Register TempRegister = MRI->createGenericVirtualRegister(LLType);
1709 MRI->setRegClass(TempRegister, &SPIRV::IDRegClass);
1710 GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
1711
1712 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1713 .addDef(NeedsExtraction ? TempRegister : Call->ReturnRegister)
1714 .addUse(GR->getSPIRVTypeID(TempType))
1715 .addUse(SampledImage)
1716 .addUse(Call->Arguments[2]) // Coordinate.
1717 .addImm(SPIRV::ImageOperand::Lod)
1718 .addUse(Lod);
1719
1720 if (NeedsExtraction)
1721 MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1722 .addDef(Call->ReturnRegister)
1723 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1724 .addUse(TempRegister)
1725 .addImm(0);
1726 } else if (HasMsaa) {
1727 MIRBuilder.buildInstr(SPIRV::OpImageRead)
1728 .addDef(Call->ReturnRegister)
1729 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1730 .addUse(Image)
1731 .addUse(Call->Arguments[1]) // Coordinate.
1732 .addImm(SPIRV::ImageOperand::Sample)
1733 .addUse(Call->Arguments[2]);
1734 } else {
1735 MIRBuilder.buildInstr(SPIRV::OpImageRead)
1736 .addDef(Call->ReturnRegister)
1737 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1738 .addUse(Image)
1739 .addUse(Call->Arguments[1]); // Coordinate.
1740 }
1741 return true;
1742}
1743
1745 MachineIRBuilder &MIRBuilder,
1746 SPIRVGlobalRegistry *GR) {
1747 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1748 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
1749 MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
1750 MIRBuilder.buildInstr(SPIRV::OpImageWrite)
1751 .addUse(Call->Arguments[0]) // Image.
1752 .addUse(Call->Arguments[1]) // Coordinate.
1753 .addUse(Call->Arguments[2]); // Texel.
1754 return true;
1755}
1756
1757static bool generateSampleImageInst(const StringRef DemangledCall,
1758 const SPIRV::IncomingCall *Call,
1759 MachineIRBuilder &MIRBuilder,
1760 SPIRVGlobalRegistry *GR) {
1761 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1762 if (Call->Builtin->Name.contains_insensitive(
1763 "__translate_sampler_initializer")) {
1764 // Build sampler literal.
1765 uint64_t Bitmask = getIConstVal(Call->Arguments[0], MRI);
1766 Register Sampler = GR->buildConstantSampler(
1767 Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask),
1769 getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType);
1770 return Sampler.isValid();
1771 } else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {
1772 // Create OpSampledImage.
1773 Register Image = Call->Arguments[0];
1774 SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1775 SPIRVType *SampledImageType =
1776 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1777 Register SampledImage =
1778 Call->ReturnRegister.isValid()
1779 ? Call->ReturnRegister
1780 : MRI->createVirtualRegister(&SPIRV::IDRegClass);
1781 MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1782 .addDef(SampledImage)
1783 .addUse(GR->getSPIRVTypeID(SampledImageType))
1784 .addUse(Image)
1785 .addUse(Call->Arguments[1]); // Sampler.
1786 return true;
1787 } else if (Call->Builtin->Name.contains_insensitive(
1788 "__spirv_ImageSampleExplicitLod")) {
1789 // Sample an image using an explicit level of detail.
1790 std::string ReturnType = DemangledCall.str();
1791 if (DemangledCall.contains("_R")) {
1792 ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);
1793 ReturnType = ReturnType.substr(0, ReturnType.find('('));
1794 }
1795 SPIRVType *Type =
1796 Call->ReturnType
1797 ? Call->ReturnType
1798 : GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder);
1799 if (!Type) {
1800 std::string DiagMsg =
1801 "Unable to recognize SPIRV type name: " + ReturnType;
1802 report_fatal_error(DiagMsg.c_str());
1803 }
1804 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1805 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
1806 MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass);
1807
1808 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1809 .addDef(Call->ReturnRegister)
1811 .addUse(Call->Arguments[0]) // Image.
1812 .addUse(Call->Arguments[1]) // Coordinate.
1813 .addImm(SPIRV::ImageOperand::Lod)
1814 .addUse(Call->Arguments[3]);
1815 return true;
1816 }
1817 return false;
1818}
1819
1821 MachineIRBuilder &MIRBuilder) {
1822 MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],
1823 Call->Arguments[1], Call->Arguments[2]);
1824 return true;
1825}
1826
1828 MachineIRBuilder &MIRBuilder,
1829 SPIRVGlobalRegistry *GR) {
1830 // Lookup the instruction opcode in the TableGen records.
1831 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1832 unsigned Opcode =
1833 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1834 const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1835
1836 switch (Opcode) {
1837 case SPIRV::OpSpecConstant: {
1838 // Build the SpecID decoration.
1839 unsigned SpecId =
1840 static_cast<unsigned>(getIConstVal(Call->Arguments[0], MRI));
1841 buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId,
1842 {SpecId});
1843 // Determine the constant MI.
1844 Register ConstRegister = Call->Arguments[1];
1845 const MachineInstr *Const = getDefInstrMaybeConstant(ConstRegister, MRI);
1846 assert(Const &&
1847 (Const->getOpcode() == TargetOpcode::G_CONSTANT ||
1848 Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&
1849 "Argument should be either an int or floating-point constant");
1850 // Determine the opcode and built the OpSpec MI.
1851 const MachineOperand &ConstOperand = Const->getOperand(1);
1852 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
1853 assert(ConstOperand.isCImm() && "Int constant operand is expected");
1854 Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
1855 ? SPIRV::OpSpecConstantTrue
1856 : SPIRV::OpSpecConstantFalse;
1857 }
1858 auto MIB = MIRBuilder.buildInstr(Opcode)
1859 .addDef(Call->ReturnRegister)
1860 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1861
1862 if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
1863 if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
1864 addNumImm(ConstOperand.getCImm()->getValue(), MIB);
1865 else
1866 addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
1867 }
1868 return true;
1869 }
1870 case SPIRV::OpSpecConstantComposite: {
1871 auto MIB = MIRBuilder.buildInstr(Opcode)
1872 .addDef(Call->ReturnRegister)
1873 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1874 for (unsigned i = 0; i < Call->Arguments.size(); i++)
1875 MIB.addUse(Call->Arguments[i]);
1876 return true;
1877 }
1878 default:
1879 return false;
1880 }
1881}
1882
1883static bool buildNDRange(const SPIRV::IncomingCall *Call,
1884 MachineIRBuilder &MIRBuilder,
1885 SPIRVGlobalRegistry *GR) {
1886 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1887 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
1888 SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1889 assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&
1890 PtrType->getOperand(2).isReg());
1891 Register TypeReg = PtrType->getOperand(2).getReg();
1893 MachineFunction &MF = MIRBuilder.getMF();
1894 Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1895 GR->assignSPIRVTypeToVReg(StructType, TmpReg, MF);
1896 // Skip the first arg, it's the destination pointer. OpBuildNDRange takes
1897 // three other arguments, so pass zero constant on absence.
1898 unsigned NumArgs = Call->Arguments.size();
1899 assert(NumArgs >= 2);
1900 Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];
1901 MRI->setRegClass(GlobalWorkSize, &SPIRV::IDRegClass);
1902 Register LocalWorkSize =
1903 NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];
1904 if (LocalWorkSize.isValid())
1905 MRI->setRegClass(LocalWorkSize, &SPIRV::IDRegClass);
1906 Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1];
1907 if (GlobalWorkOffset.isValid())
1908 MRI->setRegClass(GlobalWorkOffset, &SPIRV::IDRegClass);
1909 if (NumArgs < 4) {
1910 Register Const;
1911 SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);
1912 if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {
1913 MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize);
1914 assert(DefInstr && isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&
1915 DefInstr->getOperand(3).isReg());
1916 Register GWSPtr = DefInstr->getOperand(3).getReg();
1917 if (!MRI->getRegClassOrNull(GWSPtr))
1918 MRI->setRegClass(GWSPtr, &SPIRV::IDRegClass);
1919 // TODO: Maybe simplify generation of the type of the fields.
1920 unsigned Size = Call->Builtin->Name == "ndrange_3D" ? 3 : 2;
1921 unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32;
1923 Type *FieldTy = ArrayType::get(BaseTy, Size);
1924 SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder);
1925 GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1926 GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize, MF);
1927 MIRBuilder.buildInstr(SPIRV::OpLoad)
1928 .addDef(GlobalWorkSize)
1929 .addUse(GR->getSPIRVTypeID(SpvFieldTy))
1930 .addUse(GWSPtr);
1931 Const = GR->getOrCreateConsIntArray(0, MIRBuilder, SpvFieldTy);
1932 } else {
1933 Const = GR->buildConstantInt(0, MIRBuilder, SpvTy);
1934 }
1935 if (!LocalWorkSize.isValid())
1936 LocalWorkSize = Const;
1937 if (!GlobalWorkOffset.isValid())
1938 GlobalWorkOffset = Const;
1939 }
1940 assert(LocalWorkSize.isValid() && GlobalWorkOffset.isValid());
1941 MIRBuilder.buildInstr(SPIRV::OpBuildNDRange)
1942 .addDef(TmpReg)
1943 .addUse(TypeReg)
1944 .addUse(GlobalWorkSize)
1945 .addUse(LocalWorkSize)
1946 .addUse(GlobalWorkOffset);
1947 return MIRBuilder.buildInstr(SPIRV::OpStore)
1948 .addUse(Call->Arguments[0])
1949 .addUse(TmpReg);
1950}
1951
1952// TODO: maybe move to the global register.
1953static SPIRVType *
1955 SPIRVGlobalRegistry *GR) {
1956 LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
1957 Type *OpaqueType = StructType::getTypeByName(Context, "spirv.DeviceEvent");
1958 if (!OpaqueType)
1959 OpaqueType = StructType::getTypeByName(Context, "opencl.clk_event_t");
1960 if (!OpaqueType)
1961 OpaqueType = StructType::create(Context, "spirv.DeviceEvent");
1962 unsigned SC0 = storageClassToAddressSpace(SPIRV::StorageClass::Function);
1963 unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1964 Type *PtrType = PointerType::get(PointerType::get(OpaqueType, SC0), SC1);
1965 return GR->getOrCreateSPIRVType(PtrType, MIRBuilder);
1966}
1967
1969 MachineIRBuilder &MIRBuilder,
1970 SPIRVGlobalRegistry *GR) {
1971 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1972 const DataLayout &DL = MIRBuilder.getDataLayout();
1973 bool IsSpirvOp = Call->isSpirvOp();
1974 bool HasEvents = Call->Builtin->Name.contains("events") || IsSpirvOp;
1975 const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1976
1977 // Make vararg instructions before OpEnqueueKernel.
1978 // Local sizes arguments: Sizes of block invoke arguments. Clang generates
1979 // local size operands as an array, so we need to unpack them.
1980 SmallVector<Register, 16> LocalSizes;
1981 if (Call->Builtin->Name.contains("_varargs") || IsSpirvOp) {
1982 const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
1983 Register GepReg = Call->Arguments[LocalSizeArrayIdx];
1984 MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
1985 assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
1986 GepMI->getOperand(3).isReg());
1987 Register ArrayReg = GepMI->getOperand(3).getReg();
1988 MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
1989 const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
1990 assert(LocalSizeTy && "Local size type is expected");
1991 const uint64_t LocalSizeNum =
1992 cast<ArrayType>(LocalSizeTy)->getNumElements();
1993 unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1994 const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
1995 const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(
1996 Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
1997 for (unsigned I = 0; I < LocalSizeNum; ++I) {
1998 Register Reg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1999 MRI->setType(Reg, LLType);
2000 GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
2001 auto GEPInst = MIRBuilder.buildIntrinsic(
2002 Intrinsic::spv_gep, ArrayRef<Register>{Reg}, true, false);
2003 GEPInst
2004 .addImm(GepMI->getOperand(2).getImm()) // In bound.
2005 .addUse(ArrayMI->getOperand(0).getReg()) // Alloca.
2006 .addUse(buildConstantIntReg(0, MIRBuilder, GR)) // Indices.
2007 .addUse(buildConstantIntReg(I, MIRBuilder, GR));
2008 LocalSizes.push_back(Reg);
2009 }
2010 }
2011
2012 // SPIRV OpEnqueueKernel instruction has 10+ arguments.
2013 auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
2014 .addDef(Call->ReturnRegister)
2015 .addUse(GR->getSPIRVTypeID(Int32Ty));
2016
2017 // Copy all arguments before block invoke function pointer.
2018 const unsigned BlockFIdx = HasEvents ? 6 : 3;
2019 for (unsigned i = 0; i < BlockFIdx; i++)
2020 MIB.addUse(Call->Arguments[i]);
2021
2022 // If there are no event arguments in the original call, add dummy ones.
2023 if (!HasEvents) {
2024 MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Dummy num events.
2025 Register NullPtr = GR->getOrCreateConstNullPtr(
2026 MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
2027 MIB.addUse(NullPtr); // Dummy wait events.
2028 MIB.addUse(NullPtr); // Dummy ret event.
2029 }
2030
2031 MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
2032 assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
2033 // Invoke: Pointer to invoke function.
2034 MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
2035
2036 Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
2037 // Param: Pointer to block literal.
2038 MIB.addUse(BlockLiteralReg);
2039
2040 Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
2041 // TODO: these numbers should be obtained from block literal structure.
2042 // Param Size: Size of block literal structure.
2043 MIB.addUse(buildConstantIntReg(DL.getTypeStoreSize(PType), MIRBuilder, GR));
2044 // Param Aligment: Aligment of block literal structure.
2045 MIB.addUse(
2046 buildConstantIntReg(DL.getPrefTypeAlign(PType).value(), MIRBuilder, GR));
2047
2048 for (unsigned i = 0; i < LocalSizes.size(); i++)
2049 MIB.addUse(LocalSizes[i]);
2050 return true;
2051}
2052
2054 MachineIRBuilder &MIRBuilder,
2055 SPIRVGlobalRegistry *GR) {
2056 // Lookup the instruction opcode in the TableGen records.
2057 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2058 unsigned Opcode =
2059 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2060
2061 switch (Opcode) {
2062 case SPIRV::OpRetainEvent:
2063 case SPIRV::OpReleaseEvent:
2064 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
2065 return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);
2066 case SPIRV::OpCreateUserEvent:
2067 case SPIRV::OpGetDefaultQueue:
2068 return MIRBuilder.buildInstr(Opcode)
2069 .addDef(Call->ReturnRegister)
2070 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
2071 case SPIRV::OpIsValidEvent:
2072 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
2073 return MIRBuilder.buildInstr(Opcode)
2074 .addDef(Call->ReturnRegister)
2075 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2076 .addUse(Call->Arguments[0]);
2077 case SPIRV::OpSetUserEventStatus:
2078 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
2079 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
2080 return MIRBuilder.buildInstr(Opcode)
2081 .addUse(Call->Arguments[0])
2082 .addUse(Call->Arguments[1]);
2083 case SPIRV::OpCaptureEventProfilingInfo:
2084 MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
2085 MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
2086 MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
2087 return MIRBuilder.buildInstr(Opcode)
2088 .addUse(Call->Arguments[0])
2089 .addUse(Call->Arguments[1])
2090 .addUse(Call->Arguments[2]);
2091 case SPIRV::OpBuildNDRange:
2092 return buildNDRange(Call, MIRBuilder, GR);
2093 case SPIRV::OpEnqueueKernel:
2094 return buildEnqueueKernel(Call, MIRBuilder, GR);
2095 default:
2096 return false;
2097 }
2098}
2099
2101 MachineIRBuilder &MIRBuilder,
2102 SPIRVGlobalRegistry *GR) {
2103 // Lookup the instruction opcode in the TableGen records.
2104 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2105 unsigned Opcode =
2106 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2107
2108 bool IsSet = Opcode == SPIRV::OpGroupAsyncCopy;
2109 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
2110 if (Call->isSpirvOp())
2111 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2112 IsSet ? TypeReg : Register(0));
2113
2114 auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GR);
2115
2116 switch (Opcode) {
2117 case SPIRV::OpGroupAsyncCopy: {
2118 SPIRVType *NewType =
2119 Call->ReturnType->getOpcode() == SPIRV::OpTypeEvent
2120 ? nullptr
2121 : GR->getOrCreateSPIRVTypeByName("spirv.Event", MIRBuilder);
2122 Register TypeReg = GR->getSPIRVTypeID(NewType ? NewType : Call->ReturnType);
2123 unsigned NumArgs = Call->Arguments.size();
2124 Register EventReg = Call->Arguments[NumArgs - 1];
2125 bool Res = MIRBuilder.buildInstr(Opcode)
2126 .addDef(Call->ReturnRegister)
2127 .addUse(TypeReg)
2128 .addUse(Scope)
2129 .addUse(Call->Arguments[0])
2130 .addUse(Call->Arguments[1])
2131 .addUse(Call->Arguments[2])
2132 .addUse(Call->Arguments.size() > 4
2133 ? Call->Arguments[3]
2134 : buildConstantIntReg(1, MIRBuilder, GR))
2135 .addUse(EventReg);
2136 if (NewType != nullptr)
2137 insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder,
2138 MIRBuilder.getMF().getRegInfo());
2139 return Res;
2140 }
2141 case SPIRV::OpGroupWaitEvents:
2142 return MIRBuilder.buildInstr(Opcode)
2143 .addUse(Scope)
2144 .addUse(Call->Arguments[0])
2145 .addUse(Call->Arguments[1]);
2146 default:
2147 return false;
2148 }
2149}
2150
2151static bool generateConvertInst(const StringRef DemangledCall,
2152 const SPIRV::IncomingCall *Call,
2153 MachineIRBuilder &MIRBuilder,
2154 SPIRVGlobalRegistry *GR) {
2155 // Lookup the conversion builtin in the TableGen records.
2156 const SPIRV::ConvertBuiltin *Builtin =
2157 SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);
2158
2159 if (Builtin->IsSaturated)
2160 buildOpDecorate(Call->ReturnRegister, MIRBuilder,
2161 SPIRV::Decoration::SaturatedConversion, {});
2162 if (Builtin->IsRounded)
2163 buildOpDecorate(Call->ReturnRegister, MIRBuilder,
2164 SPIRV::Decoration::FPRoundingMode,
2165 {(unsigned)Builtin->RoundingMode});
2166
2167 std::string NeedExtMsg; // no errors if empty
2168 bool IsRightComponentsNumber = true; // check if input/output accepts vectors
2169 unsigned Opcode = SPIRV::OpNop;
2170 if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
2171 // Int -> ...
2172 if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
2173 // Int -> Int
2174 if (Builtin->IsSaturated)
2175 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
2176 : SPIRV::OpSatConvertSToU;
2177 else
2178 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert
2179 : SPIRV::OpSConvert;
2180 } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
2181 SPIRV::OpTypeFloat)) {
2182 // Int -> Float
2183 if (Builtin->IsBfloat16) {
2184 const auto *ST = static_cast<const SPIRVSubtarget *>(
2185 &MIRBuilder.getMF().getSubtarget());
2186 if (!ST->canUseExtension(
2187 SPIRV::Extension::SPV_INTEL_bfloat16_conversion))
2188 NeedExtMsg = "SPV_INTEL_bfloat16_conversion";
2189 IsRightComponentsNumber =
2190 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) ==
2191 GR->getScalarOrVectorComponentCount(Call->ReturnRegister);
2192 Opcode = SPIRV::OpConvertBF16ToFINTEL;
2193 } else {
2194 bool IsSourceSigned =
2195 DemangledCall[DemangledCall.find_first_of('(') + 1] != 'u';
2196 Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
2197 }
2198 }
2199 } else if (GR->isScalarOrVectorOfType(Call->Arguments[0],
2200 SPIRV::OpTypeFloat)) {
2201 // Float -> ...
2202 if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
2203 // Float -> Int
2204 if (Builtin->IsBfloat16) {
2205 const auto *ST = static_cast<const SPIRVSubtarget *>(
2206 &MIRBuilder.getMF().getSubtarget());
2207 if (!ST->canUseExtension(
2208 SPIRV::Extension::SPV_INTEL_bfloat16_conversion))
2209 NeedExtMsg = "SPV_INTEL_bfloat16_conversion";
2210 IsRightComponentsNumber =
2211 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) ==
2212 GR->getScalarOrVectorComponentCount(Call->ReturnRegister);
2213 Opcode = SPIRV::OpConvertFToBF16INTEL;
2214 } else {
2215 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
2216 : SPIRV::OpConvertFToU;
2217 }
2218 } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
2219 SPIRV::OpTypeFloat)) {
2220 // Float -> Float
2221 Opcode = SPIRV::OpFConvert;
2222 }
2223 }
2224
2225 if (!NeedExtMsg.empty()) {
2226 std::string DiagMsg = std::string(Builtin->Name) +
2227 ": the builtin requires the following SPIR-V "
2228 "extension: " +
2229 NeedExtMsg;
2230 report_fatal_error(DiagMsg.c_str(), false);
2231 }
2232 if (!IsRightComponentsNumber) {
2233 std::string DiagMsg =
2234 std::string(Builtin->Name) +
2235 ": result and argument must have the same number of components";
2236 report_fatal_error(DiagMsg.c_str(), false);
2237 }
2238 assert(Opcode != SPIRV::OpNop &&
2239 "Conversion between the types not implemented!");
2240
2241 MIRBuilder.buildInstr(Opcode)
2242 .addDef(Call->ReturnRegister)
2243 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2244 .addUse(Call->Arguments[0]);
2245 return true;
2246}
2247
2249 MachineIRBuilder &MIRBuilder,
2250 SPIRVGlobalRegistry *GR) {
2251 // Lookup the vector load/store builtin in the TableGen records.
2252 const SPIRV::VectorLoadStoreBuiltin *Builtin =
2253 SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
2254 Call->Builtin->Set);
2255 // Build extended instruction.
2256 auto MIB =
2257 MIRBuilder.buildInstr(SPIRV::OpExtInst)
2258 .addDef(Call->ReturnRegister)
2259 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2260 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
2261 .addImm(Builtin->Number);
2262 for (auto Argument : Call->Arguments)
2263 MIB.addUse(Argument);
2264 if (Builtin->Name.contains("load") && Builtin->ElementCount > 1)
2265 MIB.addImm(Builtin->ElementCount);
2266
2267 // Rounding mode should be passed as a last argument in the MI for builtins
2268 // like "vstorea_halfn_r".
2269 if (Builtin->IsRounded)
2270 MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));
2271 return true;
2272}
2273
2275 MachineIRBuilder &MIRBuilder,
2276 SPIRVGlobalRegistry *GR) {
2277 // Lookup the instruction opcode in the TableGen records.
2278 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2279 unsigned Opcode =
2280 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2281 bool IsLoad = Opcode == SPIRV::OpLoad;
2282 // Build the instruction.
2283 auto MIB = MIRBuilder.buildInstr(Opcode);
2284 if (IsLoad) {
2285 MIB.addDef(Call->ReturnRegister);
2286 MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType));
2287 }
2288 // Add a pointer to the value to load/store.
2289 MIB.addUse(Call->Arguments[0]);
2290 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2291 MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
2292 // Add a value to store.
2293 if (!IsLoad) {
2294 MIB.addUse(Call->Arguments[1]);
2295 MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
2296 }
2297 // Add optional memory attributes and an alignment.
2298 unsigned NumArgs = Call->Arguments.size();
2299 if ((IsLoad && NumArgs >= 2) || NumArgs >= 3) {
2300 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 1 : 2], MRI));
2301 MRI->setRegClass(Call->Arguments[IsLoad ? 1 : 2], &SPIRV::IDRegClass);
2302 }
2303 if ((IsLoad && NumArgs >= 3) || NumArgs >= 4) {
2304 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 2 : 3], MRI));
2305 MRI->setRegClass(Call->Arguments[IsLoad ? 2 : 3], &SPIRV::IDRegClass);
2306 }
2307 return true;
2308}
2309
2310/// Lowers a builtin funtion call using the provided \p DemangledCall skeleton
2311/// and external instruction \p Set.
2312namespace SPIRV {
2313std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
2314 SPIRV::InstructionSet::InstructionSet Set,
2315 MachineIRBuilder &MIRBuilder,
2316 const Register OrigRet, const Type *OrigRetTy,
2317 const SmallVectorImpl<Register> &Args,
2318 SPIRVGlobalRegistry *GR) {
2319 LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
2320
2321 // SPIR-V type and return register.
2322 Register ReturnRegister = OrigRet;
2323 SPIRVType *ReturnType = nullptr;
2324 if (OrigRetTy && !OrigRetTy->isVoidTy()) {
2325 ReturnType = GR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder);
2326 if (!MIRBuilder.getMRI()->getRegClassOrNull(ReturnRegister))
2327 MIRBuilder.getMRI()->setRegClass(ReturnRegister, &SPIRV::IDRegClass);
2328 } else if (OrigRetTy && OrigRetTy->isVoidTy()) {
2329 ReturnRegister = MIRBuilder.getMRI()->createVirtualRegister(&IDRegClass);
2330 MIRBuilder.getMRI()->setType(ReturnRegister, LLT::scalar(32));
2331 ReturnType = GR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder);
2332 }
2333
2334 // Lookup the builtin in the TableGen records.
2335 std::unique_ptr<const IncomingCall> Call =
2336 lookupBuiltin(DemangledCall, Set, ReturnRegister, ReturnType, Args);
2337
2338 if (!Call) {
2339 LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
2340 return std::nullopt;
2341 }
2342
2343 // TODO: check if the provided args meet the builtin requirments.
2344 assert(Args.size() >= Call->Builtin->MinNumArgs &&
2345 "Too few arguments to generate the builtin");
2346 if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)
2347 LLVM_DEBUG(dbgs() << "More arguments provided than required!\n");
2348
2349 // Match the builtin with implementation based on the grouping.
2350 switch (Call->Builtin->Group) {
2351 case SPIRV::Extended:
2352 return generateExtInst(Call.get(), MIRBuilder, GR);
2353 case SPIRV::Relational:
2354 return generateRelationalInst(Call.get(), MIRBuilder, GR);
2355 case SPIRV::Group:
2356 return generateGroupInst(Call.get(), MIRBuilder, GR);
2357 case SPIRV::Variable:
2358 return generateBuiltinVar(Call.get(), MIRBuilder, GR);
2359 case SPIRV::Atomic:
2360 return generateAtomicInst(Call.get(), MIRBuilder, GR);
2361 case SPIRV::AtomicFloating:
2362 return generateAtomicFloatingInst(Call.get(), MIRBuilder, GR);
2363 case SPIRV::Barrier:
2364 return generateBarrierInst(Call.get(), MIRBuilder, GR);
2365 case SPIRV::CastToPtr:
2366 return generateCastToPtrInst(Call.get(), MIRBuilder);
2367 case SPIRV::Dot:
2368 return generateDotOrFMulInst(Call.get(), MIRBuilder, GR);
2369 case SPIRV::Wave:
2370 return generateWaveInst(Call.get(), MIRBuilder, GR);
2371 case SPIRV::GetQuery:
2372 return generateGetQueryInst(Call.get(), MIRBuilder, GR);
2373 case SPIRV::ImageSizeQuery:
2374 return generateImageSizeQueryInst(Call.get(), MIRBuilder, GR);
2375 case SPIRV::ImageMiscQuery:
2376 return generateImageMiscQueryInst(Call.get(), MIRBuilder, GR);
2377 case SPIRV::ReadImage:
2378 return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
2379 case SPIRV::WriteImage:
2380 return generateWriteImageInst(Call.get(), MIRBuilder, GR);
2381 case SPIRV::SampleImage:
2382 return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
2383 case SPIRV::Select:
2384 return generateSelectInst(Call.get(), MIRBuilder);
2385 case SPIRV::SpecConstant:
2386 return generateSpecConstantInst(Call.get(), MIRBuilder, GR);
2387 case SPIRV::Enqueue:
2388 return generateEnqueueInst(Call.get(), MIRBuilder, GR);
2389 case SPIRV::AsyncCopy:
2390 return generateAsyncCopy(Call.get(), MIRBuilder, GR);
2391 case SPIRV::Convert:
2392 return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);
2393 case SPIRV::VectorLoadStore:
2394 return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);
2395 case SPIRV::LoadStore:
2396 return generateLoadStoreInst(Call.get(), MIRBuilder, GR);
2397 case SPIRV::IntelSubgroups:
2398 return generateIntelSubgroupsInst(Call.get(), MIRBuilder, GR);
2399 case SPIRV::GroupUniform:
2400 return generateGroupUniformInst(Call.get(), MIRBuilder, GR);
2401 case SPIRV::KernelClock:
2402 return generateKernelClockInst(Call.get(), MIRBuilder, GR);
2403 }
2404 return false;
2405}
2406
2408 unsigned ArgIdx, LLVMContext &Ctx) {
2409 SmallVector<StringRef, 10> BuiltinArgsTypeStrs;
2410 StringRef BuiltinArgs =
2411 DemangledCall.slice(DemangledCall.find('(') + 1, DemangledCall.find(')'));
2412 BuiltinArgs.split(BuiltinArgsTypeStrs, ',', -1, false);
2413 if (ArgIdx >= BuiltinArgsTypeStrs.size())
2414 return nullptr;
2415 StringRef TypeStr = BuiltinArgsTypeStrs[ArgIdx].trim();
2416
2417 // Parse strings representing OpenCL builtin types.
2418 if (hasBuiltinTypePrefix(TypeStr)) {
2419 // OpenCL builtin types in demangled call strings have the following format:
2420 // e.g. ocl_image2d_ro
2421 bool IsOCLBuiltinType = TypeStr.consume_front("ocl_");
2422 assert(IsOCLBuiltinType && "Invalid OpenCL builtin prefix");
2423
2424 // Check if this is pointer to a builtin type and not just pointer
2425 // representing a builtin type. In case it is a pointer to builtin type,
2426 // this will require additional handling in the method calling
2427 // parseBuiltinCallArgumentBaseType(...) as this function only retrieves the
2428 // base types.
2429 if (TypeStr.ends_with("*"))
2430 TypeStr = TypeStr.slice(0, TypeStr.find_first_of(" *"));
2431
2432 return parseBuiltinTypeNameToTargetExtType("opencl." + TypeStr.str() + "_t",
2433 Ctx);
2434 }
2435
2436 // Parse type name in either "typeN" or "type vector[N]" format, where
2437 // N is the number of elements of the vector.
2438 Type *BaseType;
2439 unsigned VecElts = 0;
2440
2441 BaseType = parseBasicTypeName(TypeStr, Ctx);
2442 if (!BaseType)
2443 // Unable to recognize SPIRV type name.
2444 return nullptr;
2445
2446 if (BaseType->isVoidTy())
2448
2449 // Handle "typeN*" or "type vector[N]*".
2450 TypeStr.consume_back("*");
2451
2452 if (TypeStr.consume_front(" vector["))
2453 TypeStr = TypeStr.substr(0, TypeStr.find(']'));
2454
2455 TypeStr.getAsInteger(10, VecElts);
2456 if (VecElts > 0)
2457 BaseType = VectorType::get(BaseType, VecElts, false);
2458
2459 return BaseType;
2460}
2461
2465};
2466
2467#define GET_BuiltinTypes_DECL
2468#define GET_BuiltinTypes_IMPL
2469
2473};
2474
2475#define GET_OpenCLTypes_DECL
2476#define GET_OpenCLTypes_IMPL
2477
2478#include "SPIRVGenTables.inc"
2479} // namespace SPIRV
2480
2481//===----------------------------------------------------------------------===//
2482// Misc functions for parsing builtin types.
2483//===----------------------------------------------------------------------===//
2484
2486 if (Name.starts_with("void"))
2487 return Type::getVoidTy(Context);
2488 else if (Name.starts_with("int") || Name.starts_with("uint"))
2489 return Type::getInt32Ty(Context);
2490 else if (Name.starts_with("float"))
2491 return Type::getFloatTy(Context);
2492 else if (Name.starts_with("half"))
2493 return Type::getHalfTy(Context);
2494 report_fatal_error("Unable to recognize type!");
2495}
2496
2497//===----------------------------------------------------------------------===//
2498// Implementation functions for builtin types.
2499//===----------------------------------------------------------------------===//
2500
2502 const SPIRV::BuiltinType *TypeRecord,
2503 MachineIRBuilder &MIRBuilder,
2504 SPIRVGlobalRegistry *GR) {
2505 unsigned Opcode = TypeRecord->Opcode;
2506 // Create or get an existing type from GlobalRegistry.
2507 return GR->getOrCreateOpTypeByOpcode(ExtensionType, MIRBuilder, Opcode);
2508}
2509
2511 SPIRVGlobalRegistry *GR) {
2512 // Create or get an existing type from GlobalRegistry.
2513 return GR->getOrCreateOpTypeSampler(MIRBuilder);
2514}
2515
2516static SPIRVType *getPipeType(const TargetExtType *ExtensionType,
2517 MachineIRBuilder &MIRBuilder,
2518 SPIRVGlobalRegistry *GR) {
2519 assert(ExtensionType->getNumIntParameters() == 1 &&
2520 "Invalid number of parameters for SPIR-V pipe builtin!");
2521 // Create or get an existing type from GlobalRegistry.
2522 return GR->getOrCreateOpTypePipe(MIRBuilder,
2523 SPIRV::AccessQualifier::AccessQualifier(
2524 ExtensionType->getIntParameter(0)));
2525}
2526
2527static SPIRVType *
2528getImageType(const TargetExtType *ExtensionType,
2529 const SPIRV::AccessQualifier::AccessQualifier Qualifier,
2530 MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
2531 assert(ExtensionType->getNumTypeParameters() == 1 &&
2532 "SPIR-V image builtin type must have sampled type parameter!");
2533 const SPIRVType *SampledType =
2534 GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder);
2535 assert(ExtensionType->getNumIntParameters() == 7 &&
2536 "Invalid number of parameters for SPIR-V image builtin!");
2537 // Create or get an existing type from GlobalRegistry.
2538 return GR->getOrCreateOpTypeImage(
2539 MIRBuilder, SampledType,
2540 SPIRV::Dim::Dim(ExtensionType->getIntParameter(0)),
2541 ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2),
2542 ExtensionType->getIntParameter(3), ExtensionType->getIntParameter(4),
2543 SPIRV::ImageFormat::ImageFormat(ExtensionType->getIntParameter(5)),
2544 Qualifier == SPIRV::AccessQualifier::WriteOnly
2545 ? SPIRV::AccessQualifier::WriteOnly
2546 : SPIRV::AccessQualifier::AccessQualifier(
2547 ExtensionType->getIntParameter(6)));
2548}
2549
2551 MachineIRBuilder &MIRBuilder,
2552 SPIRVGlobalRegistry *GR) {
2553 SPIRVType *OpaqueImageType = getImageType(
2554 OpaqueType, SPIRV::AccessQualifier::ReadOnly, MIRBuilder, GR);
2555 // Create or get an existing type from GlobalRegistry.
2556 return GR->getOrCreateOpTypeSampledImage(OpaqueImageType, MIRBuilder);
2557}
2558
2559namespace SPIRV {
2561 LLVMContext &Context) {
2562 StringRef NameWithParameters = TypeName;
2563
2564 // Pointers-to-opaque-structs representing OpenCL types are first translated
2565 // to equivalent SPIR-V types. OpenCL builtin type names should have the
2566 // following format: e.g. %opencl.event_t
2567 if (NameWithParameters.starts_with("opencl.")) {
2568 const SPIRV::OpenCLType *OCLTypeRecord =
2569 SPIRV::lookupOpenCLType(NameWithParameters);
2570 if (!OCLTypeRecord)
2571 report_fatal_error("Missing TableGen record for OpenCL type: " +
2572 NameWithParameters);
2573 NameWithParameters = OCLTypeRecord->SpirvTypeLiteral;
2574 // Continue with the SPIR-V builtin type...
2575 }
2576
2577 // Names of the opaque structs representing a SPIR-V builtins without
2578 // parameters should have the following format: e.g. %spirv.Event
2579 assert(NameWithParameters.starts_with("spirv.") &&
2580 "Unknown builtin opaque type!");
2581
2582 // Parameterized SPIR-V builtins names follow this format:
2583 // e.g. %spirv.Image._void_1_0_0_0_0_0_0, %spirv.Pipe._0
2584 if (!NameWithParameters.contains('_'))
2585 return TargetExtType::get(Context, NameWithParameters);
2586
2587 SmallVector<StringRef> Parameters;
2588 unsigned BaseNameLength = NameWithParameters.find('_') - 1;
2589 SplitString(NameWithParameters.substr(BaseNameLength + 1), Parameters, "_");
2590
2591 SmallVector<Type *, 1> TypeParameters;
2592 bool HasTypeParameter = !isDigit(Parameters[0][0]);
2593 if (HasTypeParameter)
2594 TypeParameters.push_back(parseTypeString(Parameters[0], Context));
2595 SmallVector<unsigned> IntParameters;
2596 for (unsigned i = HasTypeParameter ? 1 : 0; i < Parameters.size(); i++) {
2597 unsigned IntParameter = 0;
2598 bool ValidLiteral = !Parameters[i].getAsInteger(10, IntParameter);
2599 (void)ValidLiteral;
2600 assert(ValidLiteral &&
2601 "Invalid format of SPIR-V builtin parameter literal!");
2602 IntParameters.push_back(IntParameter);
2603 }
2604 return TargetExtType::get(Context,
2605 NameWithParameters.substr(0, BaseNameLength),
2606 TypeParameters, IntParameters);
2607}
2608
2610 SPIRV::AccessQualifier::AccessQualifier AccessQual,
2611 MachineIRBuilder &MIRBuilder,
2612 SPIRVGlobalRegistry *GR) {
2613 // In LLVM IR, SPIR-V and OpenCL builtin types are represented as either
2614 // target(...) target extension types or pointers-to-opaque-structs. The
2615 // approach relying on structs is deprecated and works only in the non-opaque
2616 // pointer mode (-opaque-pointers=0).
2617 // In order to maintain compatibility with LLVM IR generated by older versions
2618 // of Clang and LLVM/SPIR-V Translator, the pointers-to-opaque-structs are
2619 // "translated" to target extension types. This translation is temporary and
2620 // will be removed in the future release of LLVM.
2621 const TargetExtType *BuiltinType = dyn_cast<TargetExtType>(OpaqueType);
2622 if (!BuiltinType)
2624 OpaqueType->getStructName().str(), MIRBuilder.getContext());
2625
2626 unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();
2627
2628 const StringRef Name = BuiltinType->getName();
2629 LLVM_DEBUG(dbgs() << "Lowering builtin type: " << Name << "\n");
2630
2631 // Lookup the demangled builtin type in the TableGen records.
2632 const SPIRV::BuiltinType *TypeRecord = SPIRV::lookupBuiltinType(Name);
2633 if (!TypeRecord)
2634 report_fatal_error("Missing TableGen record for builtin type: " + Name);
2635
2636 // "Lower" the BuiltinType into TargetType. The following get<...>Type methods
2637 // use the implementation details from TableGen records or TargetExtType
2638 // parameters to either create a new OpType<...> machine instruction or get an
2639 // existing equivalent SPIRVType from GlobalRegistry.
2640 SPIRVType *TargetType;
2641 switch (TypeRecord->Opcode) {
2642 case SPIRV::OpTypeImage:
2643 TargetType = getImageType(BuiltinType, AccessQual, MIRBuilder, GR);
2644 break;
2645 case SPIRV::OpTypePipe:
2646 TargetType = getPipeType(BuiltinType, MIRBuilder, GR);
2647 break;
2648 case SPIRV::OpTypeDeviceEvent:
2649 TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
2650 break;
2651 case SPIRV::OpTypeSampler:
2652 TargetType = getSamplerType(MIRBuilder, GR);
2653 break;
2654 case SPIRV::OpTypeSampledImage:
2655 TargetType = getSampledImageType(BuiltinType, MIRBuilder, GR);
2656 break;
2657 default:
2658 TargetType =
2659 getNonParameterizedType(BuiltinType, TypeRecord, MIRBuilder, GR);
2660 break;
2661 }
2662
2663 // Emit OpName instruction if a new OpType<...> instruction was added
2664 // (equivalent type was not found in GlobalRegistry).
2665 if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
2666 buildOpName(GR->getSPIRVTypeID(TargetType), Name, MIRBuilder);
2667
2668 return TargetType;
2669}
2670} // namespace SPIRV
2671} // namespace llvm
unsigned const MachineRegisterInfo * MRI
MachineInstrBuilder MachineInstrBuilder & DefMI
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
AMDGPU Lower Kernel Arguments
return RetTy
#define LLVM_DEBUG(X)
Definition: Debug.h:101
std::string Name
uint64_t Size
IRTranslator LLVM IR MI
#define I(x, y, z)
Definition: MD5.cpp:58
unsigned Reg
static bool isDigit(const char C)
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
This file contains some functions that are useful when dealing with strings.
static bool contains(SmallPtrSetImpl< ConstantExpr * > &Cache, ConstantExpr *Expr, Constant *C)
Definition: Value.cpp:469
APInt bitcastToAPInt() const
Definition: APFloat.h:1246
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
Definition: APFloat.h:975
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
Definition: APInt.h:213
uint64_t getZExtValue() const
Get zero extended value.
Definition: APInt.h:1499
This class represents an incoming formal argument to a Function.
Definition: Argument.h:31
ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...
Definition: ArrayRef.h:41
static ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Definition: Type.cpp:647
@ ICMP_ULT
unsigned less than
Definition: InstrTypes.h:1018
@ ICMP_EQ
equal
Definition: InstrTypes.h:1014
const APFloat & getValueAPF() const
Definition: Constants.h:312
const APInt & getValue() const
Return the constant as an APInt value reference.
Definition: Constants.h:146
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:110
Tagged union holding either a T or a Error.
Definition: Error.h:474
Class to represent fixed width SIMD vectors.
Definition: DerivedTypes.h:539
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Definition: Function.cpp:356
static IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
Definition: Type.cpp:278
static constexpr LLT vector(ElementCount EC, unsigned ScalarSizeInBits)
Get a low-level vector of some number of elements and element width.
Definition: LowLevelType.h:64
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
Definition: LowLevelType.h:42
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
Definition: LowLevelType.h:57
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
Definition: LowLevelType.h:100
This is an important class for using LLVM in a threaded context.
Definition: LLVMContext.h:67
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
LLVMContext & getContext() const
MachineInstrBuilder buildSelect(const DstOp &Res, const SrcOp &Tst, const SrcOp &Op0, const SrcOp &Op1, std::optional< unsigned > Flags=std::nullopt)
Build and insert a Res = G_SELECT Tst, Op0, Op1.
MachineInstrBuilder buildICmp(CmpInst::Predicate Pred, const DstOp &Res, const SrcOp &Op0, const SrcOp &Op1)
Build and insert a Res = G_ICMP Pred, Op0, Op1.
MachineInstrBuilder buildIntrinsic(Intrinsic::ID ID, ArrayRef< Register > Res, bool HasSideEffects, bool isConvergent)
Build and insert a G_INTRINSIC instruction.
MachineInstrBuilder buildLoad(const DstOp &Res, const SrcOp &Addr, MachineMemOperand &MMO)
Build and insert Res = G_LOAD Addr, MMO.
MachineInstrBuilder buildZExtOrTrunc(const DstOp &Res, const SrcOp &Op)
Build and insert Res = G_ZEXT Op, Res = G_TRUNC Op, or Res = COPY Op depending on the differing sizes...
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
MachineInstrBuilder buildCopy(const DstOp &Res, const SrcOp &Op)
Build and insert Res = COPY Op.
const DataLayout & getDataLayout() const
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & addUse(Register RegNo, unsigned Flags=0, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addDef(Register RegNo, unsigned Flags=0, unsigned SubReg=0) const
Add a virtual register definition operand.
Representation of each machine instruction.
Definition: MachineInstr.h:69
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
Definition: MachineInstr.h:564
const MachineOperand & getOperand(unsigned i) const
Definition: MachineInstr.h:574
MachineOperand class - Representation of each machine instruction operand.
const GlobalValue * getGlobal() const
const ConstantInt * getCImm() const
bool isCImm() const
isCImm - Test if this is a MO_CImmediate operand.
int64_t getImm() const
bool isReg() const
isReg - Tests if this is a MO_Register operand.
const MDNode * getMetadata() const
Register getReg() const
getReg - Returns the register number.
const ConstantFP * getFPImm() const
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
static PointerType * get(Type *ElementType, unsigned AddressSpace)
This constructs a pointer to an object of the specified type in a numbered address space.
Wrapper class representing virtual and physical registers.
Definition: Register.h:19
constexpr bool isValid() const
Definition: Register.h:116
SPIRVType * getOrCreateOpTypePipe(MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AccQual)
SPIRVType * getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
SPIRVType * getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder)
Register getOrCreateConsIntVector(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType, bool EmitIR=true)
const Type * getTypeForSPIRVType(const SPIRVType *Ty) const
Register buildConstantSampler(Register Res, unsigned AddrMode, unsigned Param, unsigned FilerMode, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)
unsigned getScalarOrVectorComponentCount(Register VReg) const
SPIRVType * getOrCreateOpTypeImage(MachineIRBuilder &MIRBuilder, SPIRVType *SampledType, SPIRV::Dim::Dim Dim, uint32_t Depth, uint32_t Arrayed, uint32_t Multisampled, uint32_t Sampled, SPIRV::ImageFormat::ImageFormat ImageFormat, SPIRV::AccessQualifier::AccessQualifier AccQual)
SPIRVType * getOrCreateOpTypeByOpcode(const Type *Ty, MachineIRBuilder &MIRBuilder, unsigned Opcode)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType=nullptr)
Register getSPIRVTypeID(const SPIRVType *SpirvType) const
SPIRVType * getOrCreateSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite, bool EmitIR=true)
void assignSPIRVTypeToVReg(SPIRVType *Type, Register VReg, MachineFunction &MF)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register buildGlobalVariable(Register Reg, SPIRVType *BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, bool HasLinkageTy, SPIRV::LinkageType::LinkageType LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVType * getOrCreateOpTypeSampledImage(SPIRVType *ImageType, MachineIRBuilder &MIRBuilder)
SPIRVType * getOrCreateSPIRVTypeByName(StringRef TypeStr, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC=SPIRV::StorageClass::Function, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite)
SPIRVType * assignTypeToVReg(const Type *Type, Register VReg, MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite, bool EmitIR=true)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
SPIRVType * getOrCreateOpTypeDeviceEvent(MachineIRBuilder &MIRBuilder)
SPIRVType * getOrCreateSPIRVPointerType(SPIRVType *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SClass=SPIRV::StorageClass::Function)
SPIRVType * getOrCreateSPIRVVectorType(SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder)
SPIRVType * getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
Register getOrCreateConsIntArray(uint64_t Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
SPIRVType * getOrCreateOpTypeSampler(MachineIRBuilder &MIRBuilder)
Register buildConstantInt(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType=nullptr, bool EmitIR=true)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)
unsigned getScalarOrVectorBitWidth(const SPIRVType *Type) const
size_t size() const
Definition: SmallVector.h:91
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
Definition: SmallVector.h:586
void push_back(const T &Elt)
Definition: SmallVector.h:426
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1209
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:50
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
Definition: StringRef.h:693
bool consume_back(StringRef Suffix)
Returns true if this StringRef has the given suffix and removes that suffix.
Definition: StringRef.h:648
bool getAsInteger(unsigned Radix, T &Result) const
Parse the current string as an integer of the specified radix.
Definition: StringRef.h:463
std::string str() const
str - Get the contents as an std::string.
Definition: StringRef.h:223
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
Definition: StringRef.h:564
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
Definition: StringRef.h:258
bool contains_insensitive(StringRef Other) const
Return true if the given string is a substring of *this, and false otherwise.
Definition: StringRef.h:429
StringRef slice(size_t Start, size_t End) const
Return a reference to the substring from [Start, End).
Definition: StringRef.h:677
bool contains(StringRef Other) const
Return true if the given string is a substring of *this, and false otherwise.
Definition: StringRef.h:417
bool consume_front(StringRef Prefix)
Returns true if this StringRef has the given prefix and removes that prefix.
Definition: StringRef.h:628
size_t find_first_of(char C, size_t From=0) const
Find the first character in the string that is C, or npos if not found.
Definition: StringRef.h:370
size_t find(char C, size_t From=0) const
Search for the first character C in the string.
Definition: StringRef.h:290
bool ends_with(StringRef Suffix) const
Check if this string ends with the given Suffix.
Definition: StringRef.h:270
A switch()-like statement whose cases are string literals.
Definition: StringSwitch.h:44
StringSwitch & EndsWith(StringLiteral S, T Value)
Definition: StringSwitch.h:76
Class to represent struct types.
Definition: DerivedTypes.h:216
static StructType * getTypeByName(LLVMContext &C, StringRef Name)
Return the type with the specified name, or null if there is none by that name.
Definition: Type.cpp:632
static StructType * create(LLVMContext &Context, StringRef Name)
This creates an identified struct.
Definition: Type.cpp:513
Class to represent target extensions types, which are generally unintrospectable from target-independ...
Definition: DerivedTypes.h:720
unsigned getNumIntParameters() const
Definition: DerivedTypes.h:765
static TargetExtType * get(LLVMContext &Context, StringRef Name, ArrayRef< Type * > Types=std::nullopt, ArrayRef< unsigned > Ints=std::nullopt)
Return a target extension type having the specified name and optional type and integer parameters.
Definition: Type.cpp:796
Type * getTypeParameter(unsigned i) const
Definition: DerivedTypes.h:755
unsigned getNumTypeParameters() const
Definition: DerivedTypes.h:756
unsigned getIntParameter(unsigned i) const
Definition: DerivedTypes.h:764
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
static Type * getHalfTy(LLVMContext &C)
StringRef getStructName() const
static Type * getVoidTy(LLVMContext &C)
static IntegerType * getInt8Ty(LLVMContext &C)
static IntegerType * getInt32Ty(LLVMContext &C)
static Type * getFloatTy(LLVMContext &C)
bool isVoidTy() const
Return true if this is 'void'.
Definition: Type.h:140
LLVM Value Representation.
Definition: Value.h:74
Value(Type *Ty, unsigned scid)
Definition: Value.cpp:53
static VectorType * get(Type *ElementType, ElementCount EC)
This static method is the primary way to construct an VectorType.
Definition: Type.cpp:676
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
Definition: ilist_node.h:316
LLVMTypeRef LLVMVectorType(LLVMTypeRef ElementType, unsigned ElementCount)
Create a vector type that contains a defined type and has a specific number of elements.
Definition: Core.cpp:884
Type * parseBuiltinCallArgumentBaseType(const StringRef DemangledCall, unsigned ArgIdx, LLVMContext &Ctx)
Parses the provided ArgIdx argument base type in the DemangledCall skeleton.
TargetExtType * parseBuiltinTypeNameToTargetExtType(std::string TypeName, LLVMContext &Context)
Translates a string representing a SPIR-V or OpenCL builtin type to a TargetExtType that can be furth...
std::optional< bool > lowerBuiltin(const StringRef DemangledCall, SPIRV::InstructionSet::InstructionSet Set, MachineIRBuilder &MIRBuilder, const Register OrigRet, const Type *OrigRetTy, const SmallVectorImpl< Register > &Args, SPIRVGlobalRegistry *GR)
SPIRVType * lowerBuiltinType(const Type *OpaqueType, SPIRV::AccessQualifier::AccessQualifier AccessQual, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
StorageClass
Definition: XCOFF.h:170
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:18
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
Definition: SPIRVUtils.cpp:100
unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
Definition: SPIRVUtils.cpp:166
static bool generateGetQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static Register buildConstantIntReg(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, unsigned BitWidth=32)
static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building atomic flag instructions (e.g.
static Register buildBuiltinVariableLoad(MachineIRBuilder &MIRBuilder, SPIRVType *VariableType, SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, LLT LLType, Register Reg=Register(0), bool isConst=true, bool hasLinkageTy=true)
Helper function for building a load instruction for loading a builtin global variable of BuiltinValue...
static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRV::SamplerFilterMode::SamplerFilterMode getSamplerFilterModeFromBitmask(unsigned Bitmask)
static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic store instruction.
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
Definition: SPIRVUtils.cpp:80
static const Type * getBlockStructType(Register ParamReg, MachineRegisterInfo *MRI)
static bool generateGroupInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim)
Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy, SPIRVGlobalRegistry *GR, MachineIRBuilder &MIB, MachineRegisterInfo &MRI)
Helper external function for inserting ASSIGN_TYPE instuction between Reg and its definition,...
static Register buildScopeReg(Register CLScopeRegister, SPIRV::Scope::Scope Scope, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI)
static std::tuple< Register, SPIRVType * > buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType, SPIRVGlobalRegistry *GR)
Helper function building either a resulting scalar or vector bool register depending on the expected ...
static unsigned getNumSizeComponents(SPIRVType *imgType)
Helper function for obtaining the number of size components.
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
Definition: SPIRVUtils.cpp:269
static SPIRVType * getSampledImageType(const TargetExtType *OpaqueType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
Definition: SPIRVUtils.cpp:218
static bool generateDotOrFMulInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateSampleImageInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateBarrierInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateKernelClockInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateGroupUniformInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateWaveInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRVType * getImageType(const TargetExtType *ExtensionType, const SPIRV::AccessQualifier::AccessQualifier Qualifier, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
Definition: SPIRVUtils.cpp:117
static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building barriers, i.e., memory/control ordering operations.
static bool generateAsyncCopy(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope)
static SPIRVType * getSamplerType(MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition: Debug.cpp:163
static Register buildLoadInst(SPIRVType *BaseType, Register PtrRegister, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, LLT LowLevelType, Register DestinationReg=Register(0))
Helper function for building a load instruction loading into the DestinationReg.
static bool generateEnqueueInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
void report_fatal_error(Error Err, bool gen_crash_diag=true)
Report a serious error, calling any installed error handler.
Definition: Error.cpp:159
static bool buildSelectInst(MachineIRBuilder &MIRBuilder, Register ReturnRegister, Register SourceRegister, const SPIRVType *ReturnType, SPIRVGlobalRegistry *GR)
Helper function for building either a vector or scalar select instruction depending on the expected R...
static const Type * getMachineInstrType(MachineInstr *MI)
static SPIRV::SamplerAddressingMode::SamplerAddressingMode getSamplerAddressingModeFromBitmask(unsigned Bitmask)
static bool generateAtomicInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateConvertInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static Register buildMemSemanticsReg(Register SemanticsRegister, Register PtrRegister, unsigned &Semantics, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static unsigned getConstFromIntrinsic(Register Reg, MachineRegisterInfo *MRI)
static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateSelectInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)
static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic load instruction.
static bool generateIntelSubgroupsInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRVType * getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Type * parseBasicTypeName(StringRef &TypeName, LLVMContext &Ctx)
Definition: SPIRVUtils.cpp:400
static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, uint64_t DefaultValue)
static bool buildOpFromWrapper(MachineIRBuilder &MIRBuilder, unsigned Opcode, const SPIRV::IncomingCall *Call, Register TypeReg=Register(0))
static std::unique_ptr< const SPIRV::IncomingCall > lookupBuiltin(StringRef DemangledCall, SPIRV::InstructionSet::InstructionSet Set, Register ReturnRegister, const SPIRVType *ReturnType, const SmallVectorImpl< Register > &Arguments)
Looks up the demangled builtin call in the SPIRVBuiltins.td records using the provided DemangledCall ...
static bool buildAtomicFloatingRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic floating-type instruction.
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
Definition: SPIRVUtils.cpp:254
constexpr unsigned BitWidth
Definition: BitmaskEnum.h:191
const MachineInstr SPIRVType
static bool generateReadImageInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
bool hasBuiltinTypePrefix(StringRef Name)
Definition: SPIRVUtils.cpp:372
static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Type * getMDOperandAsType(const MDNode *N, unsigned I)
Definition: SPIRVUtils.cpp:281
static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic load instruction.
static SPIRV::MemorySemantics::MemorySemantics getSPIRVMemSemantics(std::memory_order MemOrder)
static bool generateRelationalInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)
Helper function for translating atomic init to OpStore.
static bool generateWriteImageInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRVType * getPipeType(const TargetExtType *ExtensionType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static Type * parseTypeString(const StringRef Name, LLVMContext &Context)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
Definition: SPIRVUtils.cpp:275
static bool generateCastToPtrInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)
static bool generateAtomicFloatingInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateExtInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool buildNDRange(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRVType * getNonParameterizedType(const TargetExtType *ExtensionType, const SPIRV::BuiltinType *TypeRecord, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static MachineInstr * getBlockStructInstr(Register ParamReg, MachineRegisterInfo *MRI)
static unsigned getSamplerParamFromBitmask(unsigned Bitmask)
static bool buildAtomicCompareExchangeInst(const SPIRV::IncomingCall *Call, const SPIRV::DemangledBuiltin *Builtin, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic compare-exchange instruction.
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
static bool generateBuiltinVar(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static const fltSemantics & IEEEsingle() LLVM_READNONE
Definition: APFloat.cpp:265
This struct is a compact representation of a valid (non-zero power of two) alignment.
Definition: Alignment.h:39
This class contains a discriminated union of information about pointers in memory operands,...
FPRoundingMode::FPRoundingMode RoundingMode
InstructionSet::InstructionSet Set
InstructionSet::InstructionSet Set
InstructionSet::InstructionSet Set
BuiltIn::BuiltIn Value
InstructionSet::InstructionSet Set
const SmallVectorImpl< Register > & Arguments
const std::string BuiltinName
const SPIRVType * ReturnType
const Register ReturnRegister
const DemangledBuiltin * Builtin
IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin, const Register ReturnRegister, const SPIRVType *ReturnType, const SmallVectorImpl< Register > &Arguments)
InstructionSet::InstructionSet Set
InstructionSet::InstructionSet Set
FPRoundingMode::FPRoundingMode RoundingMode