LLVM 22.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 <regex>
22#include <string>
23#include <tuple>
24
25#define DEBUG_TYPE "spirv-builtins"
26
27namespace llvm {
28namespace SPIRV {
29#define GET_BuiltinGroup_DECL
30#include "SPIRVGenTables.inc"
31
34 InstructionSet::InstructionSet Set;
35 BuiltinGroup Group;
38};
39
40#define GET_DemangledBuiltins_DECL
41#define GET_DemangledBuiltins_IMPL
42
60
63 InstructionSet::InstructionSet Set;
65};
66
67#define GET_NativeBuiltins_DECL
68#define GET_NativeBuiltins_IMPL
69
85
86#define GET_GroupBuiltins_DECL
87#define GET_GroupBuiltins_IMPL
88
96
97#define GET_IntelSubgroupsBuiltins_DECL
98#define GET_IntelSubgroupsBuiltins_IMPL
99
104
105#define GET_AtomicFloatingBuiltins_DECL
106#define GET_AtomicFloatingBuiltins_IMPL
112
113#define GET_GroupUniformBuiltins_DECL
114#define GET_GroupUniformBuiltins_IMPL
115
118 InstructionSet::InstructionSet Set;
119 BuiltIn::BuiltIn Value;
120};
121
122using namespace BuiltIn;
123#define GET_GetBuiltins_DECL
124#define GET_GetBuiltins_IMPL
125
128 InstructionSet::InstructionSet Set;
130};
131
132#define GET_ImageQueryBuiltins_DECL
133#define GET_ImageQueryBuiltins_IMPL
134
140
141#define GET_IntegerDotProductBuiltins_DECL
142#define GET_IntegerDotProductBuiltins_IMPL
143
146 InstructionSet::InstructionSet Set;
151 bool IsTF32;
152 FPRoundingMode::FPRoundingMode RoundingMode;
153};
154
157 InstructionSet::InstructionSet Set;
161 FPRoundingMode::FPRoundingMode RoundingMode;
162};
163
164using namespace FPRoundingMode;
165#define GET_ConvertBuiltins_DECL
166#define GET_ConvertBuiltins_IMPL
167
168using namespace InstructionSet;
169#define GET_VectorLoadStoreBuiltins_DECL
170#define GET_VectorLoadStoreBuiltins_IMPL
171
172#define GET_CLMemoryScope_DECL
173#define GET_CLSamplerAddressingMode_DECL
174#define GET_CLMemoryFenceFlags_DECL
175#define GET_ExtendedBuiltins_DECL
176#include "SPIRVGenTables.inc"
177} // namespace SPIRV
178
179//===----------------------------------------------------------------------===//
180// Misc functions for looking up builtins and veryfying requirements using
181// TableGen records
182//===----------------------------------------------------------------------===//
183
184namespace SPIRV {
185/// Parses the name part of the demangled builtin call.
186std::string lookupBuiltinNameHelper(StringRef DemangledCall,
187 FPDecorationId *DecorationId) {
188 StringRef PassPrefix = "(anonymous namespace)::";
189 std::string BuiltinName;
190 // Itanium Demangler result may have "(anonymous namespace)::" prefix
191 if (DemangledCall.starts_with(PassPrefix))
192 BuiltinName = DemangledCall.substr(PassPrefix.size());
193 else
194 BuiltinName = DemangledCall;
195 // Extract the builtin function name and types of arguments from the call
196 // skeleton.
197 BuiltinName = BuiltinName.substr(0, BuiltinName.find('('));
198
199 // Account for possible "__spirv_ocl_" prefix in SPIR-V friendly LLVM IR
200 if (BuiltinName.rfind("__spirv_ocl_", 0) == 0)
201 BuiltinName = BuiltinName.substr(12);
202
203 // Check if the extracted name contains type information between angle
204 // brackets. If so, the builtin is an instantiated template - needs to have
205 // the information after angle brackets and return type removed.
206 std::size_t Pos1 = BuiltinName.rfind('<');
207 if (Pos1 != std::string::npos && BuiltinName.back() == '>') {
208 std::size_t Pos2 = BuiltinName.rfind(' ', Pos1);
209 if (Pos2 == std::string::npos)
210 Pos2 = 0;
211 else
212 ++Pos2;
213 BuiltinName = BuiltinName.substr(Pos2, Pos1 - Pos2);
214 BuiltinName = BuiltinName.substr(BuiltinName.find_last_of(' ') + 1);
215 }
216
217 // Check if the extracted name begins with:
218 // - "__spirv_ImageSampleExplicitLod"
219 // - "__spirv_ImageRead"
220 // - "__spirv_ImageWrite"
221 // - "__spirv_ImageQuerySizeLod"
222 // - "__spirv_UDotKHR"
223 // - "__spirv_SDotKHR"
224 // - "__spirv_SUDotKHR"
225 // - "__spirv_SDotAccSatKHR"
226 // - "__spirv_UDotAccSatKHR"
227 // - "__spirv_SUDotAccSatKHR"
228 // - "__spirv_ReadClockKHR"
229 // - "__spirv_SubgroupBlockReadINTEL"
230 // - "__spirv_SubgroupImageBlockReadINTEL"
231 // - "__spirv_SubgroupImageMediaBlockReadINTEL"
232 // - "__spirv_SubgroupImageMediaBlockWriteINTEL"
233 // - "__spirv_Convert"
234 // - "__spirv_Round"
235 // - "__spirv_UConvert"
236 // - "__spirv_SConvert"
237 // - "__spirv_FConvert"
238 // - "__spirv_SatConvert"
239 // and maybe contains return type information at the end "_R<type>".
240 // If so, extract the plain builtin name without the type information.
241 static const std::regex SpvWithR(
242 "(__spirv_(ImageSampleExplicitLod|ImageRead|ImageWrite|ImageQuerySizeLod|"
243 "UDotKHR|"
244 "SDotKHR|SUDotKHR|SDotAccSatKHR|UDotAccSatKHR|SUDotAccSatKHR|"
245 "ReadClockKHR|SubgroupBlockReadINTEL|SubgroupImageBlockReadINTEL|"
246 "SubgroupImageMediaBlockReadINTEL|SubgroupImageMediaBlockWriteINTEL|"
247 "Convert|Round|"
248 "UConvert|SConvert|FConvert|SatConvert)[^_]*)(_R[^_]*_?(\\w+)?.*)?");
249 std::smatch Match;
250 if (std::regex_match(BuiltinName, Match, SpvWithR) && Match.size() > 1) {
251 std::ssub_match SubMatch;
252 if (DecorationId && Match.size() > 3) {
253 SubMatch = Match[4];
254 *DecorationId = demangledPostfixToDecorationId(SubMatch.str());
255 }
256 SubMatch = Match[1];
257 BuiltinName = SubMatch.str();
258 }
259
260 return BuiltinName;
261}
262} // namespace SPIRV
263
264/// Looks up the demangled builtin call in the SPIRVBuiltins.td records using
265/// the provided \p DemangledCall and specified \p Set.
266///
267/// The lookup follows the following algorithm, returning the first successful
268/// match:
269/// 1. Search with the plain demangled name (expecting a 1:1 match).
270/// 2. Search with the prefix before or suffix after the demangled name
271/// signyfying the type of the first argument.
272///
273/// \returns Wrapper around the demangled call and found builtin definition.
274static std::unique_ptr<const SPIRV::IncomingCall>
276 SPIRV::InstructionSet::InstructionSet Set,
277 Register ReturnRegister, const SPIRVType *ReturnType,
279 std::string BuiltinName = SPIRV::lookupBuiltinNameHelper(DemangledCall);
280
281 SmallVector<StringRef, 10> BuiltinArgumentTypes;
282 StringRef BuiltinArgs =
283 DemangledCall.slice(DemangledCall.find('(') + 1, DemangledCall.find(')'));
284 BuiltinArgs.split(BuiltinArgumentTypes, ',', -1, false);
285
286 // Look up the builtin in the defined set. Start with the plain demangled
287 // name, expecting a 1:1 match in the defined builtin set.
288 const SPIRV::DemangledBuiltin *Builtin;
289 if ((Builtin = SPIRV::lookupBuiltin(BuiltinName, Set)))
290 return std::make_unique<SPIRV::IncomingCall>(
291 BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
292
293 // If the initial look up was unsuccessful and the demangled call takes at
294 // least 1 argument, add a prefix or suffix signifying the type of the first
295 // argument and repeat the search.
296 if (BuiltinArgumentTypes.size() >= 1) {
297 char FirstArgumentType = BuiltinArgumentTypes[0][0];
298 // Prefix to be added to the builtin's name for lookup.
299 // For example, OpenCL "abs" taking an unsigned value has a prefix "u_".
300 std::string Prefix;
301
302 switch (FirstArgumentType) {
303 // Unsigned:
304 case 'u':
305 if (Set == SPIRV::InstructionSet::OpenCL_std)
306 Prefix = "u_";
307 else if (Set == SPIRV::InstructionSet::GLSL_std_450)
308 Prefix = "u";
309 break;
310 // Signed:
311 case 'c':
312 case 's':
313 case 'i':
314 case 'l':
315 if (Set == SPIRV::InstructionSet::OpenCL_std)
316 Prefix = "s_";
317 else if (Set == SPIRV::InstructionSet::GLSL_std_450)
318 Prefix = "s";
319 break;
320 // Floating-point:
321 case 'f':
322 case 'd':
323 case 'h':
324 if (Set == SPIRV::InstructionSet::OpenCL_std ||
325 Set == SPIRV::InstructionSet::GLSL_std_450)
326 Prefix = "f";
327 break;
328 }
329
330 // If argument-type name prefix was added, look up the builtin again.
331 if (!Prefix.empty() &&
332 (Builtin = SPIRV::lookupBuiltin(Prefix + BuiltinName, Set)))
333 return std::make_unique<SPIRV::IncomingCall>(
334 BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
335
336 // If lookup with a prefix failed, find a suffix to be added to the
337 // builtin's name for lookup. For example, OpenCL "group_reduce_max" taking
338 // an unsigned value has a suffix "u".
339 std::string Suffix;
340
341 switch (FirstArgumentType) {
342 // Unsigned:
343 case 'u':
344 Suffix = "u";
345 break;
346 // Signed:
347 case 'c':
348 case 's':
349 case 'i':
350 case 'l':
351 Suffix = "s";
352 break;
353 // Floating-point:
354 case 'f':
355 case 'd':
356 case 'h':
357 Suffix = "f";
358 break;
359 }
360
361 // If argument-type name suffix was added, look up the builtin again.
362 if (!Suffix.empty() &&
363 (Builtin = SPIRV::lookupBuiltin(BuiltinName + Suffix, Set)))
364 return std::make_unique<SPIRV::IncomingCall>(
365 BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
366 }
367
368 // No builtin with such name was found in the set.
369 return nullptr;
370}
371
374 // We expect the following sequence of instructions:
375 // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
376 // or = G_GLOBAL_VALUE @block_literal_global
377 // %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0
378 // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
379 MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
380 assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
381 MI->getOperand(1).isReg());
382 Register BitcastReg = MI->getOperand(1).getReg();
383 MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg);
384 assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&
385 BitcastMI->getOperand(2).isReg());
386 Register ValueReg = BitcastMI->getOperand(2).getReg();
387 MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
388 return ValueMI;
389}
390
391// Return an integer constant corresponding to the given register and
392// defined in spv_track_constant.
393// TODO: maybe unify with prelegalizer pass.
395 MachineInstr *DefMI = MRI->getUniqueVRegDef(Reg);
396 assert(DefMI->getOpcode() == TargetOpcode::G_CONSTANT &&
397 DefMI->getOperand(1).isCImm());
398 return DefMI->getOperand(1).getCImm()->getValue().getZExtValue();
399}
400
401// Return type of the instruction result from spv_assign_type intrinsic.
402// TODO: maybe unify with prelegalizer pass.
404 MachineInstr *NextMI = MI->getNextNode();
405 if (!NextMI)
406 return nullptr;
407 if (isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_name))
408 if ((NextMI = NextMI->getNextNode()) == nullptr)
409 return nullptr;
410 Register ValueReg = MI->getOperand(0).getReg();
411 if ((!isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_type) &&
412 !isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_ptr_type)) ||
413 NextMI->getOperand(1).getReg() != ValueReg)
414 return nullptr;
415 Type *Ty = getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0);
416 assert(Ty && "Type is expected");
417 return Ty;
418}
419
420static const Type *getBlockStructType(Register ParamReg,
422 // In principle, this information should be passed to us from Clang via
423 // an elementtype attribute. However, said attribute requires that
424 // the function call be an intrinsic, which is not. Instead, we rely on being
425 // able to trace this to the declaration of a variable: OpenCL C specification
426 // section 6.12.5 should guarantee that we can do this.
428 if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
429 return MI->getOperand(1).getGlobal()->getType();
430 assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
431 "Blocks in OpenCL C must be traceable to allocation site");
432 return getMachineInstrType(MI);
433}
434
435//===----------------------------------------------------------------------===//
436// Helper functions for building misc instructions
437//===----------------------------------------------------------------------===//
438
439/// Helper function building either a resulting scalar or vector bool register
440/// depending on the expected \p ResultType.
441///
442/// \returns Tuple of the resulting register and its type.
443static std::tuple<Register, SPIRVType *>
444buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType,
446 LLT Type;
447 SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder, true);
448
449 if (ResultType->getOpcode() == SPIRV::OpTypeVector) {
450 unsigned VectorElements = ResultType->getOperand(2).getImm();
451 BoolType = GR->getOrCreateSPIRVVectorType(BoolType, VectorElements,
452 MIRBuilder, true);
455 Type = LLT::vector(LLVMVectorType->getElementCount(), 1);
456 } else {
457 Type = LLT::scalar(1);
458 }
459
460 Register ResultRegister =
462 MIRBuilder.getMRI()->setRegClass(ResultRegister, GR->getRegClass(ResultType));
463 GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());
464 return std::make_tuple(ResultRegister, BoolType);
465}
466
467/// Helper function for building either a vector or scalar select instruction
468/// depending on the expected \p ResultType.
469static bool buildSelectInst(MachineIRBuilder &MIRBuilder,
470 Register ReturnRegister, Register SourceRegister,
471 const SPIRVType *ReturnType,
473 Register TrueConst, FalseConst;
474
475 if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {
476 unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType);
478 TrueConst =
479 GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType, true);
480 FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType, true);
481 } else {
482 TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType, true);
483 FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType, true);
484 }
485
486 return MIRBuilder.buildSelect(ReturnRegister, SourceRegister, TrueConst,
487 FalseConst);
488}
489
490/// Helper function for building a load instruction loading into the
491/// \p DestinationReg.
493 MachineIRBuilder &MIRBuilder,
494 SPIRVGlobalRegistry *GR, LLT LowLevelType,
495 Register DestinationReg = Register(0)) {
496 if (!DestinationReg.isValid())
497 DestinationReg = createVirtualRegister(BaseType, GR, MIRBuilder);
498 // TODO: consider using correct address space and alignment (p0 is canonical
499 // type for selection though).
501 MIRBuilder.buildLoad(DestinationReg, PtrRegister, PtrInfo, Align());
502 return DestinationReg;
503}
504
505/// Helper function for building a load instruction for loading a builtin global
506/// variable of \p BuiltinValue value.
508 MachineIRBuilder &MIRBuilder, SPIRVType *VariableType,
509 SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, LLT LLType,
510 Register Reg = Register(0), bool isConst = true,
511 const std::optional<SPIRV::LinkageType::LinkageType> &LinkageTy = {
512 SPIRV::LinkageType::Import}) {
513 Register NewRegister =
514 MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::pIDRegClass);
515 MIRBuilder.getMRI()->setType(
516 NewRegister,
517 LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Function),
518 GR->getPointerSize()));
520 VariableType, MIRBuilder, SPIRV::StorageClass::Input);
521 GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
522
523 // Set up the global OpVariable with the necessary builtin decorations.
524 Register Variable = GR->buildGlobalVariable(
525 NewRegister, PtrType, getLinkStringForBuiltIn(BuiltinValue), nullptr,
526 SPIRV::StorageClass::Input, nullptr, /* isConst= */ isConst, LinkageTy,
527 MIRBuilder, false);
528
529 // Load the value from the global variable.
530 Register LoadedRegister =
531 buildLoadInst(VariableType, Variable, MIRBuilder, GR, LLType, Reg);
532 MIRBuilder.getMRI()->setType(LoadedRegister, LLType);
533 return LoadedRegister;
534}
535
536/// Helper external function for inserting ASSIGN_TYPE instuction between \p Reg
537/// and its definition, set the new register as a destination of the definition,
538/// assign SPIRVType to both registers. If SpirvTy is provided, use it as
539/// SPIRVType in ASSIGN_TYPE, otherwise create it from \p Ty. Defined in
540/// SPIRVPreLegalizer.cpp.
541extern void insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy,
544
545// TODO: Move to TableGen.
546static SPIRV::MemorySemantics::MemorySemantics
547getSPIRVMemSemantics(std::memory_order MemOrder) {
548 switch (MemOrder) {
549 case std::memory_order_relaxed:
550 return SPIRV::MemorySemantics::None;
551 case std::memory_order_acquire:
552 return SPIRV::MemorySemantics::Acquire;
553 case std::memory_order_release:
554 return SPIRV::MemorySemantics::Release;
555 case std::memory_order_acq_rel:
556 return SPIRV::MemorySemantics::AcquireRelease;
557 case std::memory_order_seq_cst:
558 return SPIRV::MemorySemantics::SequentiallyConsistent;
559 default:
560 report_fatal_error("Unknown CL memory scope");
561 }
562}
563
564static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope) {
565 switch (ClScope) {
566 case SPIRV::CLMemoryScope::memory_scope_work_item:
567 return SPIRV::Scope::Invocation;
568 case SPIRV::CLMemoryScope::memory_scope_work_group:
569 return SPIRV::Scope::Workgroup;
570 case SPIRV::CLMemoryScope::memory_scope_device:
571 return SPIRV::Scope::Device;
572 case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:
573 return SPIRV::Scope::CrossDevice;
574 case SPIRV::CLMemoryScope::memory_scope_sub_group:
575 return SPIRV::Scope::Subgroup;
576 }
577 report_fatal_error("Unknown CL memory scope");
578}
579
581 MachineIRBuilder &MIRBuilder,
583 return GR->buildConstantInt(
584 Val, MIRBuilder, GR->getOrCreateSPIRVIntegerType(32, MIRBuilder), true);
585}
586
587static Register buildScopeReg(Register CLScopeRegister,
588 SPIRV::Scope::Scope Scope,
589 MachineIRBuilder &MIRBuilder,
592 if (CLScopeRegister.isValid()) {
593 auto CLScope =
594 static_cast<SPIRV::CLMemoryScope>(getIConstVal(CLScopeRegister, MRI));
595 Scope = getSPIRVScope(CLScope);
596
597 if (CLScope == static_cast<unsigned>(Scope)) {
598 MRI->setRegClass(CLScopeRegister, &SPIRV::iIDRegClass);
599 return CLScopeRegister;
600 }
601 }
602 return buildConstantIntReg32(Scope, MIRBuilder, GR);
603}
604
607 if (MRI->getRegClassOrNull(Reg))
608 return;
609 SPIRVType *SpvType = GR->getSPIRVTypeForVReg(Reg);
610 MRI->setRegClass(Reg,
611 SpvType ? GR->getRegClass(SpvType) : &SPIRV::iIDRegClass);
612}
613
614static Register buildMemSemanticsReg(Register SemanticsRegister,
615 Register PtrRegister, unsigned &Semantics,
616 MachineIRBuilder &MIRBuilder,
618 if (SemanticsRegister.isValid()) {
619 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
620 std::memory_order Order =
621 static_cast<std::memory_order>(getIConstVal(SemanticsRegister, MRI));
622 Semantics =
623 getSPIRVMemSemantics(Order) |
625 if (static_cast<unsigned>(Order) == Semantics) {
626 MRI->setRegClass(SemanticsRegister, &SPIRV::iIDRegClass);
627 return SemanticsRegister;
628 }
629 }
630 return buildConstantIntReg32(Semantics, MIRBuilder, GR);
631}
632
633static bool buildOpFromWrapper(MachineIRBuilder &MIRBuilder, unsigned Opcode,
635 Register TypeReg,
636 ArrayRef<uint32_t> ImmArgs = {}) {
637 auto MIB = MIRBuilder.buildInstr(Opcode);
638 if (TypeReg.isValid())
639 MIB.addDef(Call->ReturnRegister).addUse(TypeReg);
640 unsigned Sz = Call->Arguments.size() - ImmArgs.size();
641 for (unsigned i = 0; i < Sz; ++i)
642 MIB.addUse(Call->Arguments[i]);
643 for (uint32_t ImmArg : ImmArgs)
644 MIB.addImm(ImmArg);
645 return true;
646}
647
648/// Helper function for translating atomic init to OpStore.
650 MachineIRBuilder &MIRBuilder) {
651 if (Call->isSpirvOp())
652 return buildOpFromWrapper(MIRBuilder, SPIRV::OpStore, Call, Register(0));
653
654 assert(Call->Arguments.size() == 2 &&
655 "Need 2 arguments for atomic init translation");
656 MIRBuilder.buildInstr(SPIRV::OpStore)
657 .addUse(Call->Arguments[0])
658 .addUse(Call->Arguments[1]);
659 return true;
660}
661
662/// Helper function for building an atomic load instruction.
664 MachineIRBuilder &MIRBuilder,
666 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
667 if (Call->isSpirvOp())
668 return buildOpFromWrapper(MIRBuilder, SPIRV::OpAtomicLoad, Call, TypeReg);
669
670 Register PtrRegister = Call->Arguments[0];
671 // TODO: if true insert call to __translate_ocl_memory_sccope before
672 // OpAtomicLoad and the function implementation. We can use Translator's
673 // output for transcoding/atomic_explicit_arguments.cl as an example.
674 Register ScopeRegister =
675 Call->Arguments.size() > 1
676 ? Call->Arguments[1]
677 : buildConstantIntReg32(SPIRV::Scope::Device, MIRBuilder, GR);
678 Register MemSemanticsReg;
679 if (Call->Arguments.size() > 2) {
680 // TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
681 MemSemanticsReg = Call->Arguments[2];
682 } else {
683 int Semantics =
684 SPIRV::MemorySemantics::SequentiallyConsistent |
686 MemSemanticsReg = buildConstantIntReg32(Semantics, MIRBuilder, GR);
687 }
688
689 MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
690 .addDef(Call->ReturnRegister)
691 .addUse(TypeReg)
692 .addUse(PtrRegister)
693 .addUse(ScopeRegister)
694 .addUse(MemSemanticsReg);
695 return true;
696}
697
698/// Helper function for building an atomic store instruction.
700 MachineIRBuilder &MIRBuilder,
702 if (Call->isSpirvOp())
703 return buildOpFromWrapper(MIRBuilder, SPIRV::OpAtomicStore, Call,
704 Register(0));
705
706 Register ScopeRegister =
707 buildConstantIntReg32(SPIRV::Scope::Device, MIRBuilder, GR);
708 Register PtrRegister = Call->Arguments[0];
709 int Semantics =
710 SPIRV::MemorySemantics::SequentiallyConsistent |
712 Register MemSemanticsReg = buildConstantIntReg32(Semantics, MIRBuilder, GR);
713 MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
714 .addUse(PtrRegister)
715 .addUse(ScopeRegister)
716 .addUse(MemSemanticsReg)
717 .addUse(Call->Arguments[1]);
718 return true;
719}
720
721/// Helper function for building an atomic compare-exchange instruction.
723 const SPIRV::IncomingCall *Call, const SPIRV::DemangledBuiltin *Builtin,
724 unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
725 if (Call->isSpirvOp())
726 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
727 GR->getSPIRVTypeID(Call->ReturnType));
728
729 bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");
730 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
731
732 Register ObjectPtr = Call->Arguments[0]; // Pointer (volatile A *object.)
733 Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
734 Register Desired = Call->Arguments[2]; // Value (C Desired).
735 SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);
736 LLT DesiredLLT = MRI->getType(Desired);
737
738 assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
739 SPIRV::OpTypePointer);
740 unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
741 (void)ExpectedType;
742 assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
743 : ExpectedType == SPIRV::OpTypePointer);
744 assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
745
746 SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);
747 assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
748 auto StorageClass = static_cast<SPIRV::StorageClass::StorageClass>(
749 SpvObjectPtrTy->getOperand(1).getImm());
750 auto MemSemStorage = getMemSemanticsForStorageClass(StorageClass);
751
752 Register MemSemEqualReg;
753 Register MemSemUnequalReg;
754 uint64_t MemSemEqual =
755 IsCmpxchg
756 ? SPIRV::MemorySemantics::None
757 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
758 uint64_t MemSemUnequal =
759 IsCmpxchg
760 ? SPIRV::MemorySemantics::None
761 : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
762 if (Call->Arguments.size() >= 4) {
763 assert(Call->Arguments.size() >= 5 &&
764 "Need 5+ args for explicit atomic cmpxchg");
765 auto MemOrdEq =
766 static_cast<std::memory_order>(getIConstVal(Call->Arguments[3], MRI));
767 auto MemOrdNeq =
768 static_cast<std::memory_order>(getIConstVal(Call->Arguments[4], MRI));
769 MemSemEqual = getSPIRVMemSemantics(MemOrdEq) | MemSemStorage;
770 MemSemUnequal = getSPIRVMemSemantics(MemOrdNeq) | MemSemStorage;
771 if (static_cast<unsigned>(MemOrdEq) == MemSemEqual)
772 MemSemEqualReg = Call->Arguments[3];
773 if (static_cast<unsigned>(MemOrdNeq) == MemSemEqual)
774 MemSemUnequalReg = Call->Arguments[4];
775 }
776 if (!MemSemEqualReg.isValid())
777 MemSemEqualReg = buildConstantIntReg32(MemSemEqual, MIRBuilder, GR);
778 if (!MemSemUnequalReg.isValid())
779 MemSemUnequalReg = buildConstantIntReg32(MemSemUnequal, MIRBuilder, GR);
780
781 Register ScopeReg;
782 auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
783 if (Call->Arguments.size() >= 6) {
784 assert(Call->Arguments.size() == 6 &&
785 "Extra args for explicit atomic cmpxchg");
786 auto ClScope = static_cast<SPIRV::CLMemoryScope>(
787 getIConstVal(Call->Arguments[5], MRI));
788 Scope = getSPIRVScope(ClScope);
789 if (ClScope == static_cast<unsigned>(Scope))
790 ScopeReg = Call->Arguments[5];
791 }
792 if (!ScopeReg.isValid())
793 ScopeReg = buildConstantIntReg32(Scope, MIRBuilder, GR);
794
795 Register Expected = IsCmpxchg
796 ? ExpectedArg
797 : buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,
798 GR, LLT::scalar(64));
799 MRI->setType(Expected, DesiredLLT);
800 Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT)
801 : Call->ReturnRegister;
802 if (!MRI->getRegClassOrNull(Tmp))
803 MRI->setRegClass(Tmp, GR->getRegClass(SpvDesiredTy));
804 GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
805
806 MIRBuilder.buildInstr(Opcode)
807 .addDef(Tmp)
808 .addUse(GR->getSPIRVTypeID(SpvDesiredTy))
809 .addUse(ObjectPtr)
810 .addUse(ScopeReg)
811 .addUse(MemSemEqualReg)
812 .addUse(MemSemUnequalReg)
813 .addUse(Desired)
815 if (!IsCmpxchg) {
816 MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);
817 MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected);
818 }
819 return true;
820}
821
822/// Helper function for building atomic instructions.
823static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
824 MachineIRBuilder &MIRBuilder,
826 if (Call->isSpirvOp())
827 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
828 GR->getSPIRVTypeID(Call->ReturnType));
829
830 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
831 Register ScopeRegister =
832 Call->Arguments.size() >= 4 ? Call->Arguments[3] : Register();
833
834 assert(Call->Arguments.size() <= 4 &&
835 "Too many args for explicit atomic RMW");
836 ScopeRegister = buildScopeReg(ScopeRegister, SPIRV::Scope::Workgroup,
837 MIRBuilder, GR, MRI);
838
839 Register PtrRegister = Call->Arguments[0];
840 unsigned Semantics = SPIRV::MemorySemantics::None;
841 Register MemSemanticsReg =
842 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
843 MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
844 Semantics, MIRBuilder, GR);
845 Register ValueReg = Call->Arguments[1];
846 Register ValueTypeReg = GR->getSPIRVTypeID(Call->ReturnType);
847 // support cl_ext_float_atomics
848 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeFloat) {
849 if (Opcode == SPIRV::OpAtomicIAdd) {
850 Opcode = SPIRV::OpAtomicFAddEXT;
851 } else if (Opcode == SPIRV::OpAtomicISub) {
852 // Translate OpAtomicISub applied to a floating type argument to
853 // OpAtomicFAddEXT with the negative value operand
854 Opcode = SPIRV::OpAtomicFAddEXT;
855 Register NegValueReg =
856 MRI->createGenericVirtualRegister(MRI->getType(ValueReg));
857 MRI->setRegClass(NegValueReg, GR->getRegClass(Call->ReturnType));
858 GR->assignSPIRVTypeToVReg(Call->ReturnType, NegValueReg,
859 MIRBuilder.getMF());
860 MIRBuilder.buildInstr(TargetOpcode::G_FNEG)
861 .addDef(NegValueReg)
862 .addUse(ValueReg);
863 insertAssignInstr(NegValueReg, nullptr, Call->ReturnType, GR, MIRBuilder,
864 MIRBuilder.getMF().getRegInfo());
865 ValueReg = NegValueReg;
866 }
867 }
868 MIRBuilder.buildInstr(Opcode)
869 .addDef(Call->ReturnRegister)
870 .addUse(ValueTypeReg)
871 .addUse(PtrRegister)
872 .addUse(ScopeRegister)
873 .addUse(MemSemanticsReg)
874 .addUse(ValueReg);
875 return true;
876}
877
878/// Helper function for building an atomic floating-type instruction.
880 unsigned Opcode,
881 MachineIRBuilder &MIRBuilder,
883 assert(Call->Arguments.size() == 4 &&
884 "Wrong number of atomic floating-type builtin");
885 Register PtrReg = Call->Arguments[0];
886 Register ScopeReg = Call->Arguments[1];
887 Register MemSemanticsReg = Call->Arguments[2];
888 Register ValueReg = Call->Arguments[3];
889 MIRBuilder.buildInstr(Opcode)
890 .addDef(Call->ReturnRegister)
891 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
892 .addUse(PtrReg)
893 .addUse(ScopeReg)
894 .addUse(MemSemanticsReg)
895 .addUse(ValueReg);
896 return true;
897}
898
899/// Helper function for building atomic flag instructions (e.g.
900/// OpAtomicFlagTestAndSet).
902 unsigned Opcode, MachineIRBuilder &MIRBuilder,
904 bool IsSet = Opcode == SPIRV::OpAtomicFlagTestAndSet;
905 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
906 if (Call->isSpirvOp())
907 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
908 IsSet ? TypeReg : Register(0));
909
910 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
911 Register PtrRegister = Call->Arguments[0];
912 unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;
913 Register MemSemanticsReg =
914 Call->Arguments.size() >= 2 ? Call->Arguments[1] : Register();
915 MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
916 Semantics, MIRBuilder, GR);
917
918 assert((Opcode != SPIRV::OpAtomicFlagClear ||
919 (Semantics != SPIRV::MemorySemantics::Acquire &&
920 Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&
921 "Invalid memory order argument!");
922
923 Register ScopeRegister =
924 Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
925 ScopeRegister =
926 buildScopeReg(ScopeRegister, SPIRV::Scope::Device, MIRBuilder, GR, MRI);
927
928 auto MIB = MIRBuilder.buildInstr(Opcode);
929 if (IsSet)
930 MIB.addDef(Call->ReturnRegister).addUse(TypeReg);
931
932 MIB.addUse(PtrRegister).addUse(ScopeRegister).addUse(MemSemanticsReg);
933 return true;
934}
935
936/// Helper function for building barriers, i.e., memory/control ordering
937/// operations.
938static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
939 MachineIRBuilder &MIRBuilder,
941 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
942 const auto *ST =
943 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
944 if ((Opcode == SPIRV::OpControlBarrierArriveINTEL ||
945 Opcode == SPIRV::OpControlBarrierWaitINTEL) &&
946 !ST->canUseExtension(SPIRV::Extension::SPV_INTEL_split_barrier)) {
947 std::string DiagMsg = std::string(Builtin->Name) +
948 ": the builtin requires the following SPIR-V "
949 "extension: SPV_INTEL_split_barrier";
950 report_fatal_error(DiagMsg.c_str(), false);
951 }
952
953 if (Call->isSpirvOp())
954 return buildOpFromWrapper(MIRBuilder, Opcode, Call, Register(0));
955
956 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
957 unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI);
958 unsigned MemSemantics = SPIRV::MemorySemantics::None;
959
960 if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
961 MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
962
963 if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
964 MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;
965
966 if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)
967 MemSemantics |= SPIRV::MemorySemantics::ImageMemory;
968
969 if (Opcode == SPIRV::OpMemoryBarrier)
970 MemSemantics = getSPIRVMemSemantics(static_cast<std::memory_order>(
971 getIConstVal(Call->Arguments[1], MRI))) |
972 MemSemantics;
973 else if (Opcode == SPIRV::OpControlBarrierArriveINTEL)
974 MemSemantics |= SPIRV::MemorySemantics::Release;
975 else if (Opcode == SPIRV::OpControlBarrierWaitINTEL)
976 MemSemantics |= SPIRV::MemorySemantics::Acquire;
977 else
978 MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
979
980 Register MemSemanticsReg =
981 MemFlags == MemSemantics
982 ? Call->Arguments[0]
983 : buildConstantIntReg32(MemSemantics, MIRBuilder, GR);
984 Register ScopeReg;
985 SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
986 SPIRV::Scope::Scope MemScope = Scope;
987 if (Call->Arguments.size() >= 2) {
988 assert(
989 ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
990 (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
991 "Extra args for explicitly scoped barrier");
992 Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]
993 : Call->Arguments[1];
994 SPIRV::CLMemoryScope CLScope =
995 static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg, MRI));
996 MemScope = getSPIRVScope(CLScope);
997 if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
998 (Opcode == SPIRV::OpMemoryBarrier))
999 Scope = MemScope;
1000 if (CLScope == static_cast<unsigned>(Scope))
1001 ScopeReg = Call->Arguments[1];
1002 }
1003
1004 if (!ScopeReg.isValid())
1005 ScopeReg = buildConstantIntReg32(Scope, MIRBuilder, GR);
1006
1007 auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);
1008 if (Opcode != SPIRV::OpMemoryBarrier)
1009 MIB.addUse(buildConstantIntReg32(MemScope, MIRBuilder, GR));
1010 MIB.addUse(MemSemanticsReg);
1011 return true;
1012}
1013
1014/// Helper function for building extended bit operations.
1016 unsigned Opcode,
1017 MachineIRBuilder &MIRBuilder,
1018 SPIRVGlobalRegistry *GR) {
1019 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1020 const auto *ST =
1021 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
1022 if ((Opcode == SPIRV::OpBitFieldInsert ||
1023 Opcode == SPIRV::OpBitFieldSExtract ||
1024 Opcode == SPIRV::OpBitFieldUExtract || Opcode == SPIRV::OpBitReverse) &&
1025 !ST->canUseExtension(SPIRV::Extension::SPV_KHR_bit_instructions)) {
1026 std::string DiagMsg = std::string(Builtin->Name) +
1027 ": the builtin requires the following SPIR-V "
1028 "extension: SPV_KHR_bit_instructions";
1029 report_fatal_error(DiagMsg.c_str(), false);
1030 }
1031
1032 // Generate SPIRV instruction accordingly.
1033 if (Call->isSpirvOp())
1034 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1035 GR->getSPIRVTypeID(Call->ReturnType));
1036
1037 auto MIB = MIRBuilder.buildInstr(Opcode)
1038 .addDef(Call->ReturnRegister)
1039 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1040 for (unsigned i = 0; i < Call->Arguments.size(); ++i)
1041 MIB.addUse(Call->Arguments[i]);
1042
1043 return true;
1044}
1045
1046/// Helper function for building Intel's bindless image instructions.
1048 unsigned Opcode,
1049 MachineIRBuilder &MIRBuilder,
1050 SPIRVGlobalRegistry *GR) {
1051 // Generate SPIRV instruction accordingly.
1052 if (Call->isSpirvOp())
1053 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1054 GR->getSPIRVTypeID(Call->ReturnType));
1055
1056 MIRBuilder.buildInstr(Opcode)
1057 .addDef(Call->ReturnRegister)
1058 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1059 .addUse(Call->Arguments[0]);
1060
1061 return true;
1062}
1063
1064/// Helper function for building Intel's OpBitwiseFunctionINTEL instruction.
1066 const SPIRV::IncomingCall *Call, unsigned Opcode,
1067 MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
1068 // Generate SPIRV instruction accordingly.
1069 if (Call->isSpirvOp())
1070 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1071 GR->getSPIRVTypeID(Call->ReturnType));
1072
1073 auto MIB = MIRBuilder.buildInstr(Opcode)
1074 .addDef(Call->ReturnRegister)
1075 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1076 for (unsigned i = 0; i < Call->Arguments.size(); ++i)
1077 MIB.addUse(Call->Arguments[i]);
1078
1079 return true;
1080}
1081
1082/// Helper function for building Intel's 2d block io instructions.
1084 unsigned Opcode,
1085 MachineIRBuilder &MIRBuilder,
1086 SPIRVGlobalRegistry *GR) {
1087 // Generate SPIRV instruction accordingly.
1088 if (Call->isSpirvOp())
1089 return buildOpFromWrapper(MIRBuilder, Opcode, Call, Register(0));
1090
1091 auto MIB = MIRBuilder.buildInstr(Opcode)
1092 .addDef(Call->ReturnRegister)
1093 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1094 for (unsigned i = 0; i < Call->Arguments.size(); ++i)
1095 MIB.addUse(Call->Arguments[i]);
1096
1097 return true;
1098}
1099
1100static bool buildPipeInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
1101 unsigned Scope, MachineIRBuilder &MIRBuilder,
1102 SPIRVGlobalRegistry *GR) {
1103 switch (Opcode) {
1104 case SPIRV::OpCommitReadPipe:
1105 case SPIRV::OpCommitWritePipe:
1106 return buildOpFromWrapper(MIRBuilder, Opcode, Call, Register(0));
1107 case SPIRV::OpGroupCommitReadPipe:
1108 case SPIRV::OpGroupCommitWritePipe:
1109 case SPIRV::OpGroupReserveReadPipePackets:
1110 case SPIRV::OpGroupReserveWritePipePackets: {
1111 Register ScopeConstReg =
1112 MIRBuilder.buildConstant(LLT::scalar(32), Scope).getReg(0);
1113 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1114 MRI->setRegClass(ScopeConstReg, &SPIRV::iIDRegClass);
1116 MIB = MIRBuilder.buildInstr(Opcode);
1117 // Add Return register and type.
1118 if (Opcode == SPIRV::OpGroupReserveReadPipePackets ||
1119 Opcode == SPIRV::OpGroupReserveWritePipePackets)
1120 MIB.addDef(Call->ReturnRegister)
1121 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1122
1123 MIB.addUse(ScopeConstReg);
1124 for (unsigned int i = 0; i < Call->Arguments.size(); ++i)
1125 MIB.addUse(Call->Arguments[i]);
1126
1127 return true;
1128 }
1129 default:
1130 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1131 GR->getSPIRVTypeID(Call->ReturnType));
1132 }
1133}
1134
1135static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
1136 switch (dim) {
1137 case SPIRV::Dim::DIM_1D:
1138 case SPIRV::Dim::DIM_Buffer:
1139 return 1;
1140 case SPIRV::Dim::DIM_2D:
1141 case SPIRV::Dim::DIM_Cube:
1142 case SPIRV::Dim::DIM_Rect:
1143 return 2;
1144 case SPIRV::Dim::DIM_3D:
1145 return 3;
1146 default:
1147 report_fatal_error("Cannot get num components for given Dim");
1148 }
1149}
1150
1151/// Helper function for obtaining the number of size components.
1152static unsigned getNumSizeComponents(SPIRVType *imgType) {
1153 assert(imgType->getOpcode() == SPIRV::OpTypeImage);
1154 auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm());
1155 unsigned numComps = getNumComponentsForDim(dim);
1156 bool arrayed = imgType->getOperand(4).getImm() == 1;
1157 return arrayed ? numComps + 1 : numComps;
1158}
1159
1160//===----------------------------------------------------------------------===//
1161// Implementation functions for each builtin group
1162//===----------------------------------------------------------------------===//
1163
1165 MachineIRBuilder &MIRBuilder,
1166 SPIRVGlobalRegistry *GR, const CallBase &CB) {
1167 // Lookup the extended instruction number in the TableGen records.
1168 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1170 SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;
1171 // fmin_common and fmax_common are now deprecated, and we should use fmin and
1172 // fmax with NotInf and NotNaN flags instead. Keep original number to add
1173 // later the NoNans and NoInfs flags.
1174 uint32_t OrigNumber = Number;
1175 const SPIRVSubtarget &ST =
1176 cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
1177 if (ST.canUseExtension(SPIRV::Extension::SPV_KHR_float_controls2) &&
1178 (Number == SPIRV::OpenCLExtInst::fmin_common ||
1179 Number == SPIRV::OpenCLExtInst::fmax_common)) {
1180 Number = (Number == SPIRV::OpenCLExtInst::fmin_common)
1181 ? SPIRV::OpenCLExtInst::fmin
1182 : SPIRV::OpenCLExtInst::fmax;
1183 }
1184
1185 // Build extended instruction.
1186 auto MIB =
1187 MIRBuilder.buildInstr(SPIRV::OpExtInst)
1188 .addDef(Call->ReturnRegister)
1189 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1190 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
1191 .addImm(Number);
1192
1193 for (auto Argument : Call->Arguments)
1194 MIB.addUse(Argument);
1195 MIB.getInstr()->copyIRFlags(CB);
1196 if (OrigNumber == SPIRV::OpenCLExtInst::fmin_common ||
1197 OrigNumber == SPIRV::OpenCLExtInst::fmax_common) {
1198 // Add NoNans and NoInfs flags to fmin/fmax instruction.
1199 MIB.getInstr()->setFlag(MachineInstr::MIFlag::FmNoNans);
1200 MIB.getInstr()->setFlag(MachineInstr::MIFlag::FmNoInfs);
1201 }
1202 return true;
1203}
1204
1206 MachineIRBuilder &MIRBuilder,
1207 SPIRVGlobalRegistry *GR) {
1208 // Lookup the instruction opcode in the TableGen records.
1209 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1210 unsigned Opcode =
1211 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1212
1213 Register CompareRegister;
1214 SPIRVType *RelationType;
1215 std::tie(CompareRegister, RelationType) =
1216 buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
1217
1218 // Build relational instruction.
1219 auto MIB = MIRBuilder.buildInstr(Opcode)
1220 .addDef(CompareRegister)
1221 .addUse(GR->getSPIRVTypeID(RelationType));
1222
1223 for (auto Argument : Call->Arguments)
1224 MIB.addUse(Argument);
1225
1226 // Build select instruction.
1227 return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,
1228 Call->ReturnType, GR);
1229}
1230
1232 MachineIRBuilder &MIRBuilder,
1233 SPIRVGlobalRegistry *GR) {
1234 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1235 const SPIRV::GroupBuiltin *GroupBuiltin =
1236 SPIRV::lookupGroupBuiltin(Builtin->Name);
1237
1238 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1239 if (Call->isSpirvOp()) {
1240 if (GroupBuiltin->NoGroupOperation) {
1242 if (GroupBuiltin->Opcode ==
1243 SPIRV::OpSubgroupMatrixMultiplyAccumulateINTEL &&
1244 Call->Arguments.size() > 4)
1245 ImmArgs.push_back(getConstFromIntrinsic(Call->Arguments[4], MRI));
1246 return buildOpFromWrapper(MIRBuilder, GroupBuiltin->Opcode, Call,
1247 GR->getSPIRVTypeID(Call->ReturnType), ImmArgs);
1248 }
1249
1250 // Group Operation is a literal
1251 Register GroupOpReg = Call->Arguments[1];
1252 const MachineInstr *MI = getDefInstrMaybeConstant(GroupOpReg, MRI);
1253 if (!MI || MI->getOpcode() != TargetOpcode::G_CONSTANT)
1255 "Group Operation parameter must be an integer constant");
1256 uint64_t GrpOp = MI->getOperand(1).getCImm()->getValue().getZExtValue();
1257 Register ScopeReg = Call->Arguments[0];
1258 auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
1259 .addDef(Call->ReturnRegister)
1260 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1261 .addUse(ScopeReg)
1262 .addImm(GrpOp);
1263 for (unsigned i = 2; i < Call->Arguments.size(); ++i)
1264 MIB.addUse(Call->Arguments[i]);
1265 return true;
1266 }
1267
1268 Register Arg0;
1269 if (GroupBuiltin->HasBoolArg) {
1270 SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder, true);
1271 Register BoolReg = Call->Arguments[0];
1272 SPIRVType *BoolRegType = GR->getSPIRVTypeForVReg(BoolReg);
1273 if (!BoolRegType)
1274 report_fatal_error("Can't find a register's type definition");
1275 MachineInstr *ArgInstruction = getDefInstrMaybeConstant(BoolReg, MRI);
1276 if (ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT) {
1277 if (BoolRegType->getOpcode() != SPIRV::OpTypeBool)
1278 Arg0 = GR->buildConstantInt(getIConstVal(BoolReg, MRI), MIRBuilder,
1279 BoolType, true);
1280 } else {
1281 if (BoolRegType->getOpcode() == SPIRV::OpTypeInt) {
1282 Arg0 = MRI->createGenericVirtualRegister(LLT::scalar(1));
1283 MRI->setRegClass(Arg0, &SPIRV::iIDRegClass);
1284 GR->assignSPIRVTypeToVReg(BoolType, Arg0, MIRBuilder.getMF());
1285 MIRBuilder.buildICmp(
1286 CmpInst::ICMP_NE, Arg0, BoolReg,
1287 GR->buildConstantInt(0, MIRBuilder, BoolRegType, true));
1288 insertAssignInstr(Arg0, nullptr, BoolType, GR, MIRBuilder,
1289 MIRBuilder.getMF().getRegInfo());
1290 } else if (BoolRegType->getOpcode() != SPIRV::OpTypeBool) {
1291 report_fatal_error("Expect a boolean argument");
1292 }
1293 // if BoolReg is a boolean register, we don't need to do anything
1294 }
1295 }
1296
1297 Register GroupResultRegister = Call->ReturnRegister;
1298 SPIRVType *GroupResultType = Call->ReturnType;
1299
1300 // TODO: maybe we need to check whether the result type is already boolean
1301 // and in this case do not insert select instruction.
1302 const bool HasBoolReturnTy =
1303 GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
1304 GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
1305 GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
1306
1307 if (HasBoolReturnTy)
1308 std::tie(GroupResultRegister, GroupResultType) =
1309 buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
1310
1311 auto Scope = Builtin->Name.starts_with("sub_group") ? SPIRV::Scope::Subgroup
1312 : SPIRV::Scope::Workgroup;
1313 Register ScopeRegister = buildConstantIntReg32(Scope, MIRBuilder, GR);
1314
1315 Register VecReg;
1316 if (GroupBuiltin->Opcode == SPIRV::OpGroupBroadcast &&
1317 Call->Arguments.size() > 2) {
1318 // For OpGroupBroadcast "LocalId must be an integer datatype. It must be a
1319 // scalar, a vector with 2 components, or a vector with 3 components.",
1320 // meaning that we must create a vector from the function arguments if
1321 // it's a work_group_broadcast(val, local_id_x, local_id_y) or
1322 // work_group_broadcast(val, local_id_x, local_id_y, local_id_z) call.
1323 Register ElemReg = Call->Arguments[1];
1324 SPIRVType *ElemType = GR->getSPIRVTypeForVReg(ElemReg);
1325 if (!ElemType || ElemType->getOpcode() != SPIRV::OpTypeInt)
1326 report_fatal_error("Expect an integer <LocalId> argument");
1327 unsigned VecLen = Call->Arguments.size() - 1;
1328 VecReg = MRI->createGenericVirtualRegister(
1329 LLT::fixed_vector(VecLen, MRI->getType(ElemReg)));
1330 MRI->setRegClass(VecReg, &SPIRV::vIDRegClass);
1331 SPIRVType *VecType =
1332 GR->getOrCreateSPIRVVectorType(ElemType, VecLen, MIRBuilder, true);
1333 GR->assignSPIRVTypeToVReg(VecType, VecReg, MIRBuilder.getMF());
1334 auto MIB =
1335 MIRBuilder.buildInstr(TargetOpcode::G_BUILD_VECTOR).addDef(VecReg);
1336 for (unsigned i = 1; i < Call->Arguments.size(); i++) {
1337 MIB.addUse(Call->Arguments[i]);
1338 setRegClassIfNull(Call->Arguments[i], MRI, GR);
1339 }
1340 insertAssignInstr(VecReg, nullptr, VecType, GR, MIRBuilder,
1341 MIRBuilder.getMF().getRegInfo());
1342 }
1343
1344 // Build work/sub group instruction.
1345 auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
1346 .addDef(GroupResultRegister)
1347 .addUse(GR->getSPIRVTypeID(GroupResultType))
1348 .addUse(ScopeRegister);
1349
1350 if (!GroupBuiltin->NoGroupOperation)
1351 MIB.addImm(GroupBuiltin->GroupOperation);
1352 if (Call->Arguments.size() > 0) {
1353 MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);
1354 setRegClassIfNull(Call->Arguments[0], MRI, GR);
1355 if (VecReg.isValid())
1356 MIB.addUse(VecReg);
1357 else
1358 for (unsigned i = 1; i < Call->Arguments.size(); i++)
1359 MIB.addUse(Call->Arguments[i]);
1360 }
1361
1362 // Build select instruction.
1363 if (HasBoolReturnTy)
1364 buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
1365 Call->ReturnType, GR);
1366 return true;
1367}
1368
1370 MachineIRBuilder &MIRBuilder,
1371 SPIRVGlobalRegistry *GR) {
1372 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1373 MachineFunction &MF = MIRBuilder.getMF();
1374 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1375 const SPIRV::IntelSubgroupsBuiltin *IntelSubgroups =
1376 SPIRV::lookupIntelSubgroupsBuiltin(Builtin->Name);
1377
1378 if (IntelSubgroups->IsMedia &&
1379 !ST->canUseExtension(SPIRV::Extension::SPV_INTEL_media_block_io)) {
1380 std::string DiagMsg = std::string(Builtin->Name) +
1381 ": the builtin requires the following SPIR-V "
1382 "extension: SPV_INTEL_media_block_io";
1383 report_fatal_error(DiagMsg.c_str(), false);
1384 } else if (!IntelSubgroups->IsMedia &&
1385 !ST->canUseExtension(SPIRV::Extension::SPV_INTEL_subgroups)) {
1386 std::string DiagMsg = std::string(Builtin->Name) +
1387 ": the builtin requires the following SPIR-V "
1388 "extension: SPV_INTEL_subgroups";
1389 report_fatal_error(DiagMsg.c_str(), false);
1390 }
1391
1392 uint32_t OpCode = IntelSubgroups->Opcode;
1393 if (Call->isSpirvOp()) {
1394 bool IsSet = OpCode != SPIRV::OpSubgroupBlockWriteINTEL &&
1395 OpCode != SPIRV::OpSubgroupImageBlockWriteINTEL &&
1396 OpCode != SPIRV::OpSubgroupImageMediaBlockWriteINTEL;
1397 return buildOpFromWrapper(MIRBuilder, OpCode, Call,
1398 IsSet ? GR->getSPIRVTypeID(Call->ReturnType)
1399 : Register(0));
1400 }
1401
1402 if (IntelSubgroups->IsBlock) {
1403 // Minimal number or arguments set in TableGen records is 1
1404 if (SPIRVType *Arg0Type = GR->getSPIRVTypeForVReg(Call->Arguments[0])) {
1405 if (Arg0Type->getOpcode() == SPIRV::OpTypeImage) {
1406 // TODO: add required validation from the specification:
1407 // "'Image' must be an object whose type is OpTypeImage with a 'Sampled'
1408 // operand of 0 or 2. If the 'Sampled' operand is 2, then some
1409 // dimensions require a capability."
1410 switch (OpCode) {
1411 case SPIRV::OpSubgroupBlockReadINTEL:
1412 OpCode = SPIRV::OpSubgroupImageBlockReadINTEL;
1413 break;
1414 case SPIRV::OpSubgroupBlockWriteINTEL:
1415 OpCode = SPIRV::OpSubgroupImageBlockWriteINTEL;
1416 break;
1417 }
1418 }
1419 }
1420 }
1421
1422 // TODO: opaque pointers types should be eventually resolved in such a way
1423 // that validation of block read is enabled with respect to the following
1424 // specification requirement:
1425 // "'Result Type' may be a scalar or vector type, and its component type must
1426 // be equal to the type pointed to by 'Ptr'."
1427 // For example, function parameter type should not be default i8 pointer, but
1428 // depend on the result type of the instruction where it is used as a pointer
1429 // argument of OpSubgroupBlockReadINTEL
1430
1431 // Build Intel subgroups instruction
1433 IntelSubgroups->IsWrite
1434 ? MIRBuilder.buildInstr(OpCode)
1435 : MIRBuilder.buildInstr(OpCode)
1436 .addDef(Call->ReturnRegister)
1437 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1438 for (size_t i = 0; i < Call->Arguments.size(); ++i)
1439 MIB.addUse(Call->Arguments[i]);
1440 return true;
1441}
1442
1444 MachineIRBuilder &MIRBuilder,
1445 SPIRVGlobalRegistry *GR) {
1446 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1447 MachineFunction &MF = MIRBuilder.getMF();
1448 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1449 if (!ST->canUseExtension(
1450 SPIRV::Extension::SPV_KHR_uniform_group_instructions)) {
1451 std::string DiagMsg = std::string(Builtin->Name) +
1452 ": the builtin requires the following SPIR-V "
1453 "extension: SPV_KHR_uniform_group_instructions";
1454 report_fatal_error(DiagMsg.c_str(), false);
1455 }
1456 const SPIRV::GroupUniformBuiltin *GroupUniform =
1457 SPIRV::lookupGroupUniformBuiltin(Builtin->Name);
1458 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1459
1460 Register GroupResultReg = Call->ReturnRegister;
1461 Register ScopeReg = Call->Arguments[0];
1462 Register ValueReg = Call->Arguments[2];
1463
1464 // Group Operation
1465 Register ConstGroupOpReg = Call->Arguments[1];
1466 const MachineInstr *Const = getDefInstrMaybeConstant(ConstGroupOpReg, MRI);
1467 if (!Const || Const->getOpcode() != TargetOpcode::G_CONSTANT)
1469 "expect a constant group operation for a uniform group instruction",
1470 false);
1471 const MachineOperand &ConstOperand = Const->getOperand(1);
1472 if (!ConstOperand.isCImm())
1473 report_fatal_error("uniform group instructions: group operation must be an "
1474 "integer constant",
1475 false);
1476
1477 auto MIB = MIRBuilder.buildInstr(GroupUniform->Opcode)
1478 .addDef(GroupResultReg)
1479 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1480 .addUse(ScopeReg);
1481 addNumImm(ConstOperand.getCImm()->getValue(), MIB);
1482 MIB.addUse(ValueReg);
1483
1484 return true;
1485}
1486
1488 MachineIRBuilder &MIRBuilder,
1489 SPIRVGlobalRegistry *GR) {
1490 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1491 MachineFunction &MF = MIRBuilder.getMF();
1492 const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
1493 if (!ST->canUseExtension(SPIRV::Extension::SPV_KHR_shader_clock)) {
1494 std::string DiagMsg = std::string(Builtin->Name) +
1495 ": the builtin requires the following SPIR-V "
1496 "extension: SPV_KHR_shader_clock";
1497 report_fatal_error(DiagMsg.c_str(), false);
1498 }
1499
1500 Register ResultReg = Call->ReturnRegister;
1501
1502 // Deduce the `Scope` operand from the builtin function name.
1503 SPIRV::Scope::Scope ScopeArg =
1505 .EndsWith("device", SPIRV::Scope::Scope::Device)
1506 .EndsWith("work_group", SPIRV::Scope::Scope::Workgroup)
1507 .EndsWith("sub_group", SPIRV::Scope::Scope::Subgroup);
1508 Register ScopeReg = buildConstantIntReg32(ScopeArg, MIRBuilder, GR);
1509
1510 MIRBuilder.buildInstr(SPIRV::OpReadClockKHR)
1511 .addDef(ResultReg)
1512 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1513 .addUse(ScopeReg);
1514
1515 return true;
1516}
1517
1518// These queries ask for a single size_t result for a given dimension index,
1519// e.g. size_t get_global_id(uint dimindex). In SPIR-V, the builtins
1520// corresponding to these values are all vec3 types, so we need to extract the
1521// correct index or return DefaultValue (0 or 1 depending on the query). We also
1522// handle extending or truncating in case size_t does not match the expected
1523// result type's bitwidth.
1524//
1525// For a constant index >= 3 we generate:
1526// %res = OpConstant %SizeT DefaultValue
1527//
1528// For other indices we generate:
1529// %g = OpVariable %ptr_V3_SizeT Input
1530// OpDecorate %g BuiltIn XXX
1531// OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"
1532// OpDecorate %g Constant
1533// %loadedVec = OpLoad %V3_SizeT %g
1534//
1535// Then, if the index is constant < 3, we generate:
1536// %res = OpCompositeExtract %SizeT %loadedVec idx
1537// If the index is dynamic, we generate:
1538// %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
1539// %cmp = OpULessThan %bool %idx %const_3
1540// %res = OpSelect %SizeT %cmp %tmp %const_<DefaultValue>
1541//
1542// If the bitwidth of %res does not match the expected return type, we add an
1543// extend or truncate.
1545 MachineIRBuilder &MIRBuilder,
1547 SPIRV::BuiltIn::BuiltIn BuiltinValue,
1548 uint64_t DefaultValue) {
1549 Register IndexRegister = Call->Arguments[0];
1550 const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();
1551 const unsigned PointerSize = GR->getPointerSize();
1552 const SPIRVType *PointerSizeType =
1553 GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
1554 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1555 auto IndexInstruction = getDefInstrMaybeConstant(IndexRegister, MRI);
1556
1557 // Set up the final register to do truncation or extension on at the end.
1558 Register ToTruncate = Call->ReturnRegister;
1559
1560 // If the index is constant, we can statically determine if it is in range.
1561 bool IsConstantIndex =
1562 IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
1563
1564 // If it's out of range (max dimension is 3), we can just return the constant
1565 // default value (0 or 1 depending on which query function).
1566 if (IsConstantIndex && getIConstVal(IndexRegister, MRI) >= 3) {
1567 Register DefaultReg = Call->ReturnRegister;
1568 if (PointerSize != ResultWidth) {
1569 DefaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
1570 MRI->setRegClass(DefaultReg, &SPIRV::iIDRegClass);
1571 GR->assignSPIRVTypeToVReg(PointerSizeType, DefaultReg,
1572 MIRBuilder.getMF());
1573 ToTruncate = DefaultReg;
1574 }
1575 auto NewRegister =
1576 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType, true);
1577 MIRBuilder.buildCopy(DefaultReg, NewRegister);
1578 } else { // If it could be in range, we need to load from the given builtin.
1579 auto Vec3Ty =
1580 GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder, true);
1581 Register LoadedVector =
1582 buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,
1583 LLT::fixed_vector(3, PointerSize));
1584 // Set up the vreg to extract the result to (possibly a new temporary one).
1585 Register Extracted = Call->ReturnRegister;
1586 if (!IsConstantIndex || PointerSize != ResultWidth) {
1587 Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
1588 MRI->setRegClass(Extracted, &SPIRV::iIDRegClass);
1589 GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());
1590 }
1591 // Use Intrinsic::spv_extractelt so dynamic vs static extraction is
1592 // handled later: extr = spv_extractelt LoadedVector, IndexRegister.
1593 MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
1594 Intrinsic::spv_extractelt, ArrayRef<Register>{Extracted}, true, false);
1595 ExtractInst.addUse(LoadedVector).addUse(IndexRegister);
1596
1597 // If the index is dynamic, need check if it's < 3, and then use a select.
1598 if (!IsConstantIndex) {
1599 insertAssignInstr(Extracted, nullptr, PointerSizeType, GR, MIRBuilder,
1600 *MRI);
1601
1602 auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);
1603 auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder, true);
1604
1605 Register CompareRegister =
1606 MRI->createGenericVirtualRegister(LLT::scalar(1));
1607 MRI->setRegClass(CompareRegister, &SPIRV::iIDRegClass);
1608 GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
1609
1610 // Use G_ICMP to check if idxVReg < 3.
1611 MIRBuilder.buildICmp(
1612 CmpInst::ICMP_ULT, CompareRegister, IndexRegister,
1613 GR->buildConstantInt(3, MIRBuilder, IndexType, true));
1614
1615 // Get constant for the default value (0 or 1 depending on which
1616 // function).
1617 Register DefaultRegister =
1618 GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType, true);
1619
1620 // Get a register for the selection result (possibly a new temporary one).
1621 Register SelectionResult = Call->ReturnRegister;
1622 if (PointerSize != ResultWidth) {
1623 SelectionResult =
1624 MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
1625 MRI->setRegClass(SelectionResult, &SPIRV::iIDRegClass);
1626 GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
1627 MIRBuilder.getMF());
1628 }
1629 // Create the final G_SELECT to return the extracted value or the default.
1630 MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,
1631 DefaultRegister);
1632 ToTruncate = SelectionResult;
1633 } else {
1634 ToTruncate = Extracted;
1635 }
1636 }
1637 // Alter the result's bitwidth if it does not match the SizeT value extracted.
1638 if (PointerSize != ResultWidth)
1639 MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);
1640 return true;
1641}
1642
1644 MachineIRBuilder &MIRBuilder,
1645 SPIRVGlobalRegistry *GR) {
1646 // Lookup the builtin variable record.
1647 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1648 SPIRV::BuiltIn::BuiltIn Value =
1649 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
1650
1651 if (Value == SPIRV::BuiltIn::GlobalInvocationId)
1652 return genWorkgroupQuery(Call, MIRBuilder, GR, Value, 0);
1653
1654 // Build a load instruction for the builtin variable.
1655 unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);
1656 LLT LLType;
1657 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
1658 LLType =
1659 LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth);
1660 else
1661 LLType = LLT::scalar(BitWidth);
1662
1663 return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value,
1664 LLType, Call->ReturnRegister);
1665}
1666
1668 MachineIRBuilder &MIRBuilder,
1669 SPIRVGlobalRegistry *GR) {
1670 // Lookup the instruction opcode in the TableGen records.
1671 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1672 unsigned Opcode =
1673 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1674
1675 switch (Opcode) {
1676 case SPIRV::OpStore:
1677 return buildAtomicInitInst(Call, MIRBuilder);
1678 case SPIRV::OpAtomicLoad:
1679 return buildAtomicLoadInst(Call, MIRBuilder, GR);
1680 case SPIRV::OpAtomicStore:
1681 return buildAtomicStoreInst(Call, MIRBuilder, GR);
1682 case SPIRV::OpAtomicCompareExchange:
1683 case SPIRV::OpAtomicCompareExchangeWeak:
1684 return buildAtomicCompareExchangeInst(Call, Builtin, Opcode, MIRBuilder,
1685 GR);
1686 case SPIRV::OpAtomicIAdd:
1687 case SPIRV::OpAtomicISub:
1688 case SPIRV::OpAtomicOr:
1689 case SPIRV::OpAtomicXor:
1690 case SPIRV::OpAtomicAnd:
1691 case SPIRV::OpAtomicExchange:
1692 return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
1693 case SPIRV::OpMemoryBarrier:
1694 return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
1695 case SPIRV::OpAtomicFlagTestAndSet:
1696 case SPIRV::OpAtomicFlagClear:
1697 return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);
1698 default:
1699 if (Call->isSpirvOp())
1700 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
1701 GR->getSPIRVTypeID(Call->ReturnType));
1702 return false;
1703 }
1704}
1705
1707 MachineIRBuilder &MIRBuilder,
1708 SPIRVGlobalRegistry *GR) {
1709 // Lookup the instruction opcode in the TableGen records.
1710 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1711 unsigned Opcode = SPIRV::lookupAtomicFloatingBuiltin(Builtin->Name)->Opcode;
1712
1713 switch (Opcode) {
1714 case SPIRV::OpAtomicFAddEXT:
1715 case SPIRV::OpAtomicFMinEXT:
1716 case SPIRV::OpAtomicFMaxEXT:
1717 return buildAtomicFloatingRMWInst(Call, Opcode, MIRBuilder, GR);
1718 default:
1719 return false;
1720 }
1721}
1722
1724 MachineIRBuilder &MIRBuilder,
1725 SPIRVGlobalRegistry *GR) {
1726 // Lookup the instruction opcode in the TableGen records.
1727 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1728 unsigned Opcode =
1729 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1730
1731 return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
1732}
1733
1735 MachineIRBuilder &MIRBuilder,
1736 SPIRVGlobalRegistry *GR) {
1737 // Lookup the instruction opcode in the TableGen records.
1738 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1739 unsigned Opcode =
1740 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1741
1742 if (Opcode == SPIRV::OpGenericCastToPtrExplicit) {
1743 SPIRV::StorageClass::StorageClass ResSC =
1744 GR->getPointerStorageClass(Call->ReturnRegister);
1745 if (!isGenericCastablePtr(ResSC))
1746 return false;
1747
1748 MIRBuilder.buildInstr(Opcode)
1749 .addDef(Call->ReturnRegister)
1750 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1751 .addUse(Call->Arguments[0])
1752 .addImm(ResSC);
1753 } else {
1754 MIRBuilder.buildInstr(TargetOpcode::G_ADDRSPACE_CAST)
1755 .addDef(Call->ReturnRegister)
1756 .addUse(Call->Arguments[0]);
1757 }
1758 return true;
1759}
1760
1761static bool generateDotOrFMulInst(const StringRef DemangledCall,
1763 MachineIRBuilder &MIRBuilder,
1764 SPIRVGlobalRegistry *GR) {
1765 if (Call->isSpirvOp())
1766 return buildOpFromWrapper(MIRBuilder, SPIRV::OpDot, Call,
1767 GR->getSPIRVTypeID(Call->ReturnType));
1768
1769 bool IsVec = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() ==
1770 SPIRV::OpTypeVector;
1771 // Use OpDot only in case of vector args and OpFMul in case of scalar args.
1772 uint32_t OC = IsVec ? SPIRV::OpDot : SPIRV::OpFMulS;
1773 bool IsSwapReq = false;
1774
1775 const auto *ST =
1776 static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
1777 if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt) &&
1778 (ST->canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
1779 ST->isAtLeastSPIRVVer(VersionTuple(1, 6)))) {
1780 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1781 const SPIRV::IntegerDotProductBuiltin *IntDot =
1782 SPIRV::lookupIntegerDotProductBuiltin(Builtin->Name);
1783 if (IntDot) {
1784 OC = IntDot->Opcode;
1785 IsSwapReq = IntDot->IsSwapReq;
1786 } else if (IsVec) {
1787 // Handling "dot" and "dot_acc_sat" builtins which use vectors of
1788 // integers.
1789 LLVMContext &Ctx = MIRBuilder.getContext();
1791 SPIRV::parseBuiltinTypeStr(TypeStrs, DemangledCall, Ctx);
1792 bool IsFirstSigned = TypeStrs[0].trim()[0] != 'u';
1793 bool IsSecondSigned = TypeStrs[1].trim()[0] != 'u';
1794
1795 if (Call->BuiltinName == "dot") {
1796 if (IsFirstSigned && IsSecondSigned)
1797 OC = SPIRV::OpSDot;
1798 else if (!IsFirstSigned && !IsSecondSigned)
1799 OC = SPIRV::OpUDot;
1800 else {
1801 OC = SPIRV::OpSUDot;
1802 if (!IsFirstSigned)
1803 IsSwapReq = true;
1804 }
1805 } else if (Call->BuiltinName == "dot_acc_sat") {
1806 if (IsFirstSigned && IsSecondSigned)
1807 OC = SPIRV::OpSDotAccSat;
1808 else if (!IsFirstSigned && !IsSecondSigned)
1809 OC = SPIRV::OpUDotAccSat;
1810 else {
1811 OC = SPIRV::OpSUDotAccSat;
1812 if (!IsFirstSigned)
1813 IsSwapReq = true;
1814 }
1815 }
1816 }
1817 }
1818
1819 MachineInstrBuilder MIB = MIRBuilder.buildInstr(OC)
1820 .addDef(Call->ReturnRegister)
1821 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1822
1823 if (IsSwapReq) {
1824 MIB.addUse(Call->Arguments[1]);
1825 MIB.addUse(Call->Arguments[0]);
1826 // needed for dot_acc_sat* builtins
1827 for (size_t i = 2; i < Call->Arguments.size(); ++i)
1828 MIB.addUse(Call->Arguments[i]);
1829 } else {
1830 for (size_t i = 0; i < Call->Arguments.size(); ++i)
1831 MIB.addUse(Call->Arguments[i]);
1832 }
1833
1834 // Add Packed Vector Format for Integer dot product builtins if arguments are
1835 // scalar
1836 if (!IsVec && OC != SPIRV::OpFMulS)
1837 MIB.addImm(SPIRV::PackedVectorFormat4x8Bit);
1838
1839 return true;
1840}
1841
1843 MachineIRBuilder &MIRBuilder,
1844 SPIRVGlobalRegistry *GR) {
1845 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1846 SPIRV::BuiltIn::BuiltIn Value =
1847 SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
1848
1849 // For now, we only support a single Wave intrinsic with a single return type.
1850 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt);
1851 LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(Call->ReturnType));
1852
1854 MIRBuilder, Call->ReturnType, GR, Value, LLType, Call->ReturnRegister,
1855 /* isConst= */ false, /* LinkageType= */ std::nullopt);
1856}
1857
1858// We expect a builtin
1859// Name(ptr sret([RetType]) %result, Type %operand1, Type %operand1)
1860// where %result is a pointer to where the result of the builtin execution
1861// is to be stored, and generate the following instructions:
1862// Res = Opcode RetType Operand1 Operand1
1863// OpStore RetVariable Res
1865 MachineIRBuilder &MIRBuilder,
1866 SPIRVGlobalRegistry *GR) {
1867 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1868 unsigned Opcode =
1869 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1870
1871 Register SRetReg = Call->Arguments[0];
1872 SPIRVType *PtrRetType = GR->getSPIRVTypeForVReg(SRetReg);
1873 SPIRVType *RetType = GR->getPointeeType(PtrRetType);
1874 if (!RetType)
1875 report_fatal_error("The first parameter must be a pointer");
1876 if (RetType->getOpcode() != SPIRV::OpTypeStruct)
1877 report_fatal_error("Expected struct type result for the arithmetic with "
1878 "overflow builtins");
1879
1880 SPIRVType *OpType1 = GR->getSPIRVTypeForVReg(Call->Arguments[1]);
1881 SPIRVType *OpType2 = GR->getSPIRVTypeForVReg(Call->Arguments[2]);
1882 if (!OpType1 || !OpType2 || OpType1 != OpType2)
1883 report_fatal_error("Operands must have the same type");
1884 if (OpType1->getOpcode() == SPIRV::OpTypeVector)
1885 switch (Opcode) {
1886 case SPIRV::OpIAddCarryS:
1887 Opcode = SPIRV::OpIAddCarryV;
1888 break;
1889 case SPIRV::OpISubBorrowS:
1890 Opcode = SPIRV::OpISubBorrowV;
1891 break;
1892 }
1893
1894 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1895 Register ResReg = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
1896 if (const TargetRegisterClass *DstRC =
1897 MRI->getRegClassOrNull(Call->Arguments[1])) {
1898 MRI->setRegClass(ResReg, DstRC);
1899 MRI->setType(ResReg, MRI->getType(Call->Arguments[1]));
1900 } else {
1901 MRI->setType(ResReg, LLT::scalar(64));
1902 }
1903 GR->assignSPIRVTypeToVReg(RetType, ResReg, MIRBuilder.getMF());
1904 MIRBuilder.buildInstr(Opcode)
1905 .addDef(ResReg)
1906 .addUse(GR->getSPIRVTypeID(RetType))
1907 .addUse(Call->Arguments[1])
1908 .addUse(Call->Arguments[2]);
1909 MIRBuilder.buildInstr(SPIRV::OpStore).addUse(SRetReg).addUse(ResReg);
1910 return true;
1911}
1912
1914 MachineIRBuilder &MIRBuilder,
1915 SPIRVGlobalRegistry *GR) {
1916 // Lookup the builtin record.
1917 SPIRV::BuiltIn::BuiltIn Value =
1918 SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;
1919 const bool IsDefaultOne = (Value == SPIRV::BuiltIn::GlobalSize ||
1920 Value == SPIRV::BuiltIn::NumWorkgroups ||
1921 Value == SPIRV::BuiltIn::WorkgroupSize ||
1922 Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
1923 return genWorkgroupQuery(Call, MIRBuilder, GR, Value, IsDefaultOne ? 1 : 0);
1924}
1925
1927 MachineIRBuilder &MIRBuilder,
1928 SPIRVGlobalRegistry *GR) {
1929 // Lookup the image size query component number in the TableGen records.
1930 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1931 uint32_t Component =
1932 SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;
1933 // Query result may either be a vector or a scalar. If return type is not a
1934 // vector, expect only a single size component. Otherwise get the number of
1935 // expected components.
1936 unsigned NumExpectedRetComponents =
1937 Call->ReturnType->getOpcode() == SPIRV::OpTypeVector
1938 ? Call->ReturnType->getOperand(2).getImm()
1939 : 1;
1940 // Get the actual number of query result/size components.
1941 SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1942 unsigned NumActualRetComponents = getNumSizeComponents(ImgType);
1943 Register QueryResult = Call->ReturnRegister;
1944 SPIRVType *QueryResultType = Call->ReturnType;
1945 if (NumExpectedRetComponents != NumActualRetComponents) {
1946 unsigned Bitwidth = Call->ReturnType->getOpcode() == SPIRV::OpTypeInt
1947 ? Call->ReturnType->getOperand(1).getImm()
1948 : 32;
1949 QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
1950 LLT::fixed_vector(NumActualRetComponents, Bitwidth));
1951 MIRBuilder.getMRI()->setRegClass(QueryResult, &SPIRV::vIDRegClass);
1952 SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(Bitwidth, MIRBuilder);
1953 QueryResultType = GR->getOrCreateSPIRVVectorType(
1954 IntTy, NumActualRetComponents, MIRBuilder, true);
1955 GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());
1956 }
1957 bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;
1958 unsigned Opcode =
1959 IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
1960 auto MIB = MIRBuilder.buildInstr(Opcode)
1961 .addDef(QueryResult)
1962 .addUse(GR->getSPIRVTypeID(QueryResultType))
1963 .addUse(Call->Arguments[0]);
1964 if (!IsDimBuf)
1965 MIB.addUse(buildConstantIntReg32(0, MIRBuilder, GR)); // Lod id.
1966 if (NumExpectedRetComponents == NumActualRetComponents)
1967 return true;
1968 if (NumExpectedRetComponents == 1) {
1969 // Only 1 component is expected, build OpCompositeExtract instruction.
1970 unsigned ExtractedComposite =
1971 Component == 3 ? NumActualRetComponents - 1 : Component;
1972 assert(ExtractedComposite < NumActualRetComponents &&
1973 "Invalid composite index!");
1974 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
1975 SPIRVType *NewType = nullptr;
1976 if (QueryResultType->getOpcode() == SPIRV::OpTypeVector) {
1977 Register NewTypeReg = QueryResultType->getOperand(1).getReg();
1978 if (TypeReg != NewTypeReg &&
1979 (NewType = GR->getSPIRVTypeForVReg(NewTypeReg)) != nullptr)
1980 TypeReg = NewTypeReg;
1981 }
1982 MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1983 .addDef(Call->ReturnRegister)
1984 .addUse(TypeReg)
1985 .addUse(QueryResult)
1986 .addImm(ExtractedComposite);
1987 if (NewType != nullptr)
1988 insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder,
1989 MIRBuilder.getMF().getRegInfo());
1990 } else {
1991 // More than 1 component is expected, fill a new vector.
1992 auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)
1993 .addDef(Call->ReturnRegister)
1994 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1995 .addUse(QueryResult)
1996 .addUse(QueryResult);
1997 for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
1998 MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);
1999 }
2000 return true;
2001}
2002
2004 MachineIRBuilder &MIRBuilder,
2005 SPIRVGlobalRegistry *GR) {
2006 assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
2007 "Image samples query result must be of int type!");
2008
2009 // Lookup the instruction opcode in the TableGen records.
2010 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2011 unsigned Opcode =
2012 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2013
2014 Register Image = Call->Arguments[0];
2015 SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
2016 GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
2017 (void)ImageDimensionality;
2018
2019 switch (Opcode) {
2020 case SPIRV::OpImageQuerySamples:
2021 assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
2022 "Image must be of 2D dimensionality");
2023 break;
2024 case SPIRV::OpImageQueryLevels:
2025 assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
2026 ImageDimensionality == SPIRV::Dim::DIM_2D ||
2027 ImageDimensionality == SPIRV::Dim::DIM_3D ||
2028 ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
2029 "Image must be of 1D/2D/3D/Cube dimensionality");
2030 break;
2031 }
2032
2033 MIRBuilder.buildInstr(Opcode)
2034 .addDef(Call->ReturnRegister)
2035 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2036 .addUse(Image);
2037 return true;
2038}
2039
2040// TODO: Move to TableGen.
2041static SPIRV::SamplerAddressingMode::SamplerAddressingMode
2043 switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
2044 case SPIRV::CLK_ADDRESS_CLAMP:
2045 return SPIRV::SamplerAddressingMode::Clamp;
2046 case SPIRV::CLK_ADDRESS_CLAMP_TO_EDGE:
2047 return SPIRV::SamplerAddressingMode::ClampToEdge;
2048 case SPIRV::CLK_ADDRESS_REPEAT:
2049 return SPIRV::SamplerAddressingMode::Repeat;
2050 case SPIRV::CLK_ADDRESS_MIRRORED_REPEAT:
2051 return SPIRV::SamplerAddressingMode::RepeatMirrored;
2052 case SPIRV::CLK_ADDRESS_NONE:
2053 return SPIRV::SamplerAddressingMode::None;
2054 default:
2055 report_fatal_error("Unknown CL address mode");
2056 }
2057}
2058
2059static unsigned getSamplerParamFromBitmask(unsigned Bitmask) {
2060 return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;
2061}
2062
2063static SPIRV::SamplerFilterMode::SamplerFilterMode
2065 if (Bitmask & SPIRV::CLK_FILTER_LINEAR)
2066 return SPIRV::SamplerFilterMode::Linear;
2067 if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
2068 return SPIRV::SamplerFilterMode::Nearest;
2069 return SPIRV::SamplerFilterMode::Nearest;
2070}
2071
2072static bool generateReadImageInst(const StringRef DemangledCall,
2074 MachineIRBuilder &MIRBuilder,
2075 SPIRVGlobalRegistry *GR) {
2076 if (Call->isSpirvOp())
2077 return buildOpFromWrapper(MIRBuilder, SPIRV::OpImageRead, Call,
2078 GR->getSPIRVTypeID(Call->ReturnType));
2079 Register Image = Call->Arguments[0];
2080 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2081 bool HasOclSampler = DemangledCall.contains_insensitive("ocl_sampler");
2082 bool HasMsaa = DemangledCall.contains_insensitive("msaa");
2083 if (HasOclSampler) {
2084 Register Sampler = Call->Arguments[1];
2085
2086 if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
2087 getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) {
2088 uint64_t SamplerMask = getIConstVal(Sampler, MRI);
2089 Sampler = GR->buildConstantSampler(
2091 getSamplerParamFromBitmask(SamplerMask),
2092 getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder);
2093 }
2094 SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
2095 SPIRVType *SampledImageType =
2096 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
2097 Register SampledImage = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
2098
2099 MIRBuilder.buildInstr(SPIRV::OpSampledImage)
2100 .addDef(SampledImage)
2101 .addUse(GR->getSPIRVTypeID(SampledImageType))
2102 .addUse(Image)
2103 .addUse(Sampler);
2104
2106 MIRBuilder);
2107
2108 if (Call->ReturnType->getOpcode() != SPIRV::OpTypeVector) {
2109 SPIRVType *TempType =
2110 GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder, true);
2111 Register TempRegister =
2112 MRI->createGenericVirtualRegister(GR->getRegType(TempType));
2113 MRI->setRegClass(TempRegister, GR->getRegClass(TempType));
2114 GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
2115 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
2116 .addDef(TempRegister)
2117 .addUse(GR->getSPIRVTypeID(TempType))
2118 .addUse(SampledImage)
2119 .addUse(Call->Arguments[2]) // Coordinate.
2120 .addImm(SPIRV::ImageOperand::Lod)
2121 .addUse(Lod);
2122 MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
2123 .addDef(Call->ReturnRegister)
2124 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2125 .addUse(TempRegister)
2126 .addImm(0);
2127 } else {
2128 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
2129 .addDef(Call->ReturnRegister)
2130 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2131 .addUse(SampledImage)
2132 .addUse(Call->Arguments[2]) // Coordinate.
2133 .addImm(SPIRV::ImageOperand::Lod)
2134 .addUse(Lod);
2135 }
2136 } else if (HasMsaa) {
2137 MIRBuilder.buildInstr(SPIRV::OpImageRead)
2138 .addDef(Call->ReturnRegister)
2139 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2140 .addUse(Image)
2141 .addUse(Call->Arguments[1]) // Coordinate.
2142 .addImm(SPIRV::ImageOperand::Sample)
2143 .addUse(Call->Arguments[2]);
2144 } else {
2145 MIRBuilder.buildInstr(SPIRV::OpImageRead)
2146 .addDef(Call->ReturnRegister)
2147 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2148 .addUse(Image)
2149 .addUse(Call->Arguments[1]); // Coordinate.
2150 }
2151 return true;
2152}
2153
2155 MachineIRBuilder &MIRBuilder,
2156 SPIRVGlobalRegistry *GR) {
2157 if (Call->isSpirvOp())
2158 return buildOpFromWrapper(MIRBuilder, SPIRV::OpImageWrite, Call,
2159 Register(0));
2160 MIRBuilder.buildInstr(SPIRV::OpImageWrite)
2161 .addUse(Call->Arguments[0]) // Image.
2162 .addUse(Call->Arguments[1]) // Coordinate.
2163 .addUse(Call->Arguments[2]); // Texel.
2164 return true;
2165}
2166
2167static bool generateSampleImageInst(const StringRef DemangledCall,
2169 MachineIRBuilder &MIRBuilder,
2170 SPIRVGlobalRegistry *GR) {
2171 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2172 if (Call->Builtin->Name.contains_insensitive(
2173 "__translate_sampler_initializer")) {
2174 // Build sampler literal.
2175 uint64_t Bitmask = getIConstVal(Call->Arguments[0], MRI);
2176 Register Sampler = GR->buildConstantSampler(
2177 Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask),
2179 getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder);
2180 return Sampler.isValid();
2181 } else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {
2182 // Create OpSampledImage.
2183 Register Image = Call->Arguments[0];
2184 SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
2185 SPIRVType *SampledImageType =
2186 GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
2187 Register SampledImage =
2188 Call->ReturnRegister.isValid()
2189 ? Call->ReturnRegister
2190 : MRI->createVirtualRegister(&SPIRV::iIDRegClass);
2191 MIRBuilder.buildInstr(SPIRV::OpSampledImage)
2192 .addDef(SampledImage)
2193 .addUse(GR->getSPIRVTypeID(SampledImageType))
2194 .addUse(Image)
2195 .addUse(Call->Arguments[1]); // Sampler.
2196 return true;
2197 } else if (Call->Builtin->Name.contains_insensitive(
2198 "__spirv_ImageSampleExplicitLod")) {
2199 // Sample an image using an explicit level of detail.
2200 std::string ReturnType = DemangledCall.str();
2201 if (DemangledCall.contains("_R")) {
2202 ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);
2203 ReturnType = ReturnType.substr(0, ReturnType.find('('));
2204 }
2205 SPIRVType *Type =
2206 Call->ReturnType
2207 ? Call->ReturnType
2208 : GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder, true);
2209 if (!Type) {
2210 std::string DiagMsg =
2211 "Unable to recognize SPIRV type name: " + ReturnType;
2212 report_fatal_error(DiagMsg.c_str());
2213 }
2214 MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
2215 .addDef(Call->ReturnRegister)
2217 .addUse(Call->Arguments[0]) // Image.
2218 .addUse(Call->Arguments[1]) // Coordinate.
2219 .addImm(SPIRV::ImageOperand::Lod)
2220 .addUse(Call->Arguments[3]);
2221 return true;
2222 }
2223 return false;
2224}
2225
2227 MachineIRBuilder &MIRBuilder) {
2228 MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],
2229 Call->Arguments[1], Call->Arguments[2]);
2230 return true;
2231}
2232
2234 MachineIRBuilder &MIRBuilder,
2235 SPIRVGlobalRegistry *GR) {
2236 createContinuedInstructions(MIRBuilder, SPIRV::OpCompositeConstruct, 3,
2237 SPIRV::OpCompositeConstructContinuedINTEL,
2238 Call->Arguments, Call->ReturnRegister,
2239 GR->getSPIRVTypeID(Call->ReturnType));
2240 return true;
2241}
2242
2244 MachineIRBuilder &MIRBuilder,
2245 SPIRVGlobalRegistry *GR) {
2246 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2247 unsigned Opcode =
2248 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2249 bool IsSet = Opcode != SPIRV::OpCooperativeMatrixStoreKHR &&
2250 Opcode != SPIRV::OpCooperativeMatrixStoreCheckedINTEL &&
2251 Opcode != SPIRV::OpCooperativeMatrixPrefetchINTEL;
2252 unsigned ArgSz = Call->Arguments.size();
2253 unsigned LiteralIdx = 0;
2254 switch (Opcode) {
2255 // Memory operand is optional and is literal.
2256 case SPIRV::OpCooperativeMatrixLoadKHR:
2257 LiteralIdx = ArgSz > 3 ? 3 : 0;
2258 break;
2259 case SPIRV::OpCooperativeMatrixStoreKHR:
2260 LiteralIdx = ArgSz > 4 ? 4 : 0;
2261 break;
2262 case SPIRV::OpCooperativeMatrixLoadCheckedINTEL:
2263 LiteralIdx = ArgSz > 7 ? 7 : 0;
2264 break;
2265 case SPIRV::OpCooperativeMatrixStoreCheckedINTEL:
2266 LiteralIdx = ArgSz > 8 ? 8 : 0;
2267 break;
2268 // Cooperative Matrix Operands operand is optional and is literal.
2269 case SPIRV::OpCooperativeMatrixMulAddKHR:
2270 LiteralIdx = ArgSz > 3 ? 3 : 0;
2271 break;
2272 };
2273
2275 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2276 if (Opcode == SPIRV::OpCooperativeMatrixPrefetchINTEL) {
2277 const uint32_t CacheLevel = getConstFromIntrinsic(Call->Arguments[3], MRI);
2278 auto MIB = MIRBuilder.buildInstr(SPIRV::OpCooperativeMatrixPrefetchINTEL)
2279 .addUse(Call->Arguments[0]) // pointer
2280 .addUse(Call->Arguments[1]) // rows
2281 .addUse(Call->Arguments[2]) // columns
2282 .addImm(CacheLevel) // cache level
2283 .addUse(Call->Arguments[4]); // memory layout
2284 if (ArgSz > 5)
2285 MIB.addUse(Call->Arguments[5]); // stride
2286 if (ArgSz > 6) {
2287 const uint32_t MemOp = getConstFromIntrinsic(Call->Arguments[6], MRI);
2288 MIB.addImm(MemOp); // memory operand
2289 }
2290 return true;
2291 }
2292 if (LiteralIdx > 0)
2293 ImmArgs.push_back(getConstFromIntrinsic(Call->Arguments[LiteralIdx], MRI));
2294 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
2295 if (Opcode == SPIRV::OpCooperativeMatrixLengthKHR) {
2296 SPIRVType *CoopMatrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
2297 if (!CoopMatrType)
2298 report_fatal_error("Can't find a register's type definition");
2299 MIRBuilder.buildInstr(Opcode)
2300 .addDef(Call->ReturnRegister)
2301 .addUse(TypeReg)
2302 .addUse(CoopMatrType->getOperand(0).getReg());
2303 return true;
2304 }
2305 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2306 IsSet ? TypeReg : Register(0), ImmArgs);
2307}
2308
2310 MachineIRBuilder &MIRBuilder,
2311 SPIRVGlobalRegistry *GR) {
2312 // Lookup the instruction opcode in the TableGen records.
2313 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2314 unsigned Opcode =
2315 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2316 const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2317
2318 switch (Opcode) {
2319 case SPIRV::OpSpecConstant: {
2320 // Build the SpecID decoration.
2321 unsigned SpecId =
2322 static_cast<unsigned>(getIConstVal(Call->Arguments[0], MRI));
2323 buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId,
2324 {SpecId});
2325 // Determine the constant MI.
2326 Register ConstRegister = Call->Arguments[1];
2327 const MachineInstr *Const = getDefInstrMaybeConstant(ConstRegister, MRI);
2328 assert(Const &&
2329 (Const->getOpcode() == TargetOpcode::G_CONSTANT ||
2330 Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&
2331 "Argument should be either an int or floating-point constant");
2332 // Determine the opcode and built the OpSpec MI.
2333 const MachineOperand &ConstOperand = Const->getOperand(1);
2334 if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
2335 assert(ConstOperand.isCImm() && "Int constant operand is expected");
2336 Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
2337 ? SPIRV::OpSpecConstantTrue
2338 : SPIRV::OpSpecConstantFalse;
2339 }
2340 auto MIB = MIRBuilder.buildInstr(Opcode)
2341 .addDef(Call->ReturnRegister)
2342 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
2343
2344 if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
2345 if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
2346 addNumImm(ConstOperand.getCImm()->getValue(), MIB);
2347 else
2348 addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
2349 }
2350 return true;
2351 }
2352 case SPIRV::OpSpecConstantComposite: {
2353 createContinuedInstructions(MIRBuilder, Opcode, 3,
2354 SPIRV::OpSpecConstantCompositeContinuedINTEL,
2355 Call->Arguments, Call->ReturnRegister,
2356 GR->getSPIRVTypeID(Call->ReturnType));
2357 return true;
2358 }
2359 default:
2360 return false;
2361 }
2362}
2363
2365 MachineIRBuilder &MIRBuilder,
2366 SPIRVGlobalRegistry *GR) {
2367 // Lookup the instruction opcode in the TableGen records.
2368 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2369 unsigned Opcode =
2370 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2371
2372 return buildExtendedBitOpsInst(Call, Opcode, MIRBuilder, GR);
2373}
2374
2376 MachineIRBuilder &MIRBuilder,
2377 SPIRVGlobalRegistry *GR) {
2378 // Lookup the instruction opcode in the TableGen records.
2379 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2380 unsigned Opcode =
2381 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2382
2383 return buildBindlessImageINTELInst(Call, Opcode, MIRBuilder, GR);
2384}
2385
2386static bool
2388 MachineIRBuilder &MIRBuilder,
2389 SPIRVGlobalRegistry *GR) {
2390 // Lookup the instruction opcode in the TableGen records.
2391 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2392 unsigned Opcode =
2393 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2394
2395 return buildTernaryBitwiseFunctionINTELInst(Call, Opcode, MIRBuilder, GR);
2396}
2397
2399 MachineIRBuilder &MIRBuilder,
2400 SPIRVGlobalRegistry *GR) {
2401 // Lookup the instruction opcode in the TableGen records.
2402 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2403 unsigned Opcode =
2404 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2405
2406 return build2DBlockIOINTELInst(Call, Opcode, MIRBuilder, GR);
2407}
2408
2410 MachineIRBuilder &MIRBuilder,
2411 SPIRVGlobalRegistry *GR) {
2412 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2413 unsigned Opcode =
2414 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2415
2416 unsigned Scope = SPIRV::Scope::Workgroup;
2417 if (Builtin->Name.contains("sub_group"))
2418 Scope = SPIRV::Scope::Subgroup;
2419
2420 return buildPipeInst(Call, Opcode, Scope, MIRBuilder, GR);
2421}
2422
2424 MachineIRBuilder &MIRBuilder,
2425 SPIRVGlobalRegistry *GR) {
2426 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2427 unsigned Opcode =
2428 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2429
2430 bool IsSet = Opcode != SPIRV::OpPredicatedStoreINTEL;
2431 unsigned ArgSz = Call->Arguments.size();
2433 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2434 // Memory operand is optional and is literal.
2435 if (ArgSz > 3)
2436 ImmArgs.push_back(
2437 getConstFromIntrinsic(Call->Arguments[/*Literal index*/ 3], MRI));
2438
2439 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
2440 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2441 IsSet ? TypeReg : Register(0), ImmArgs);
2442}
2443
2445 MachineIRBuilder &MIRBuilder,
2446 SPIRVGlobalRegistry *GR) {
2447 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2448 SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
2449 assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&
2450 PtrType->getOperand(2).isReg());
2451 Register TypeReg = PtrType->getOperand(2).getReg();
2453 MachineFunction &MF = MIRBuilder.getMF();
2454 Register TmpReg = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
2455 GR->assignSPIRVTypeToVReg(StructType, TmpReg, MF);
2456 // Skip the first arg, it's the destination pointer. OpBuildNDRange takes
2457 // three other arguments, so pass zero constant on absence.
2458 unsigned NumArgs = Call->Arguments.size();
2459 assert(NumArgs >= 2);
2460 Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];
2461 Register LocalWorkSize =
2462 NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];
2463 Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1];
2464 if (NumArgs < 4) {
2465 Register Const;
2466 SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);
2467 if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {
2468 MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize);
2469 assert(DefInstr && isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&
2470 DefInstr->getOperand(3).isReg());
2471 Register GWSPtr = DefInstr->getOperand(3).getReg();
2472 // TODO: Maybe simplify generation of the type of the fields.
2473 unsigned Size = Call->Builtin->Name == "ndrange_3D" ? 3 : 2;
2474 unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32;
2476 Type *FieldTy = ArrayType::get(BaseTy, Size);
2477 SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(
2478 FieldTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, true);
2479 GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
2480 GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize, MF);
2481 MIRBuilder.buildInstr(SPIRV::OpLoad)
2482 .addDef(GlobalWorkSize)
2483 .addUse(GR->getSPIRVTypeID(SpvFieldTy))
2484 .addUse(GWSPtr);
2485 const SPIRVSubtarget &ST =
2486 cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
2487 Const = GR->getOrCreateConstIntArray(0, Size, *MIRBuilder.getInsertPt(),
2488 SpvFieldTy, *ST.getInstrInfo());
2489 } else {
2490 Const = GR->buildConstantInt(0, MIRBuilder, SpvTy, true);
2491 }
2492 if (!LocalWorkSize.isValid())
2493 LocalWorkSize = Const;
2494 if (!GlobalWorkOffset.isValid())
2495 GlobalWorkOffset = Const;
2496 }
2497 assert(LocalWorkSize.isValid() && GlobalWorkOffset.isValid());
2498 MIRBuilder.buildInstr(SPIRV::OpBuildNDRange)
2499 .addDef(TmpReg)
2500 .addUse(TypeReg)
2501 .addUse(GlobalWorkSize)
2502 .addUse(LocalWorkSize)
2503 .addUse(GlobalWorkOffset);
2504 return MIRBuilder.buildInstr(SPIRV::OpStore)
2505 .addUse(Call->Arguments[0])
2506 .addUse(TmpReg);
2507}
2508
2509// TODO: maybe move to the global register.
2510static SPIRVType *
2512 SPIRVGlobalRegistry *GR) {
2513 LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
2514 unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
2515 Type *PtrType = PointerType::get(Context, SC1);
2516 return GR->getOrCreateSPIRVType(PtrType, MIRBuilder,
2517 SPIRV::AccessQualifier::ReadWrite, true);
2518}
2519
2521 MachineIRBuilder &MIRBuilder,
2522 SPIRVGlobalRegistry *GR) {
2523 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2524 const DataLayout &DL = MIRBuilder.getDataLayout();
2525 bool IsSpirvOp = Call->isSpirvOp();
2526 bool HasEvents = Call->Builtin->Name.contains("events") || IsSpirvOp;
2527 const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
2528
2529 // Make vararg instructions before OpEnqueueKernel.
2530 // Local sizes arguments: Sizes of block invoke arguments. Clang generates
2531 // local size operands as an array, so we need to unpack them.
2532 SmallVector<Register, 16> LocalSizes;
2533 if (Call->Builtin->Name.contains("_varargs") || IsSpirvOp) {
2534 const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
2535 Register GepReg = Call->Arguments[LocalSizeArrayIdx];
2536 MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
2537 assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
2538 GepMI->getOperand(3).isReg());
2539 Register ArrayReg = GepMI->getOperand(3).getReg();
2540 MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
2541 const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
2542 assert(LocalSizeTy && "Local size type is expected");
2543 const uint64_t LocalSizeNum =
2544 cast<ArrayType>(LocalSizeTy)->getNumElements();
2545 unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
2546 const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
2547 const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(
2548 Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
2549 for (unsigned I = 0; I < LocalSizeNum; ++I) {
2550 Register Reg = MRI->createVirtualRegister(&SPIRV::pIDRegClass);
2551 MRI->setType(Reg, LLType);
2552 GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
2553 auto GEPInst = MIRBuilder.buildIntrinsic(
2554 Intrinsic::spv_gep, ArrayRef<Register>{Reg}, true, false);
2555 GEPInst
2556 .addImm(GepMI->getOperand(2).getImm()) // In bound.
2557 .addUse(ArrayMI->getOperand(0).getReg()) // Alloca.
2558 .addUse(buildConstantIntReg32(0, MIRBuilder, GR)) // Indices.
2559 .addUse(buildConstantIntReg32(I, MIRBuilder, GR));
2560 LocalSizes.push_back(Reg);
2561 }
2562 }
2563
2564 // SPIRV OpEnqueueKernel instruction has 10+ arguments.
2565 auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
2566 .addDef(Call->ReturnRegister)
2568
2569 // Copy all arguments before block invoke function pointer.
2570 const unsigned BlockFIdx = HasEvents ? 6 : 3;
2571 for (unsigned i = 0; i < BlockFIdx; i++)
2572 MIB.addUse(Call->Arguments[i]);
2573
2574 // If there are no event arguments in the original call, add dummy ones.
2575 if (!HasEvents) {
2576 MIB.addUse(buildConstantIntReg32(0, MIRBuilder, GR)); // Dummy num events.
2577 Register NullPtr = GR->getOrCreateConstNullPtr(
2578 MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
2579 MIB.addUse(NullPtr); // Dummy wait events.
2580 MIB.addUse(NullPtr); // Dummy ret event.
2581 }
2582
2583 MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
2584 assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
2585 // Invoke: Pointer to invoke function.
2586 MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
2587
2588 Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
2589 // Param: Pointer to block literal.
2590 MIB.addUse(BlockLiteralReg);
2591
2592 Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
2593 // TODO: these numbers should be obtained from block literal structure.
2594 // Param Size: Size of block literal structure.
2595 MIB.addUse(buildConstantIntReg32(DL.getTypeStoreSize(PType), MIRBuilder, GR));
2596 // Param Aligment: Aligment of block literal structure.
2597 MIB.addUse(buildConstantIntReg32(DL.getPrefTypeAlign(PType).value(),
2598 MIRBuilder, GR));
2599
2600 for (unsigned i = 0; i < LocalSizes.size(); i++)
2601 MIB.addUse(LocalSizes[i]);
2602 return true;
2603}
2604
2606 MachineIRBuilder &MIRBuilder,
2607 SPIRVGlobalRegistry *GR) {
2608 // Lookup the instruction opcode in the TableGen records.
2609 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2610 unsigned Opcode =
2611 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2612
2613 switch (Opcode) {
2614 case SPIRV::OpRetainEvent:
2615 case SPIRV::OpReleaseEvent:
2616 return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);
2617 case SPIRV::OpCreateUserEvent:
2618 case SPIRV::OpGetDefaultQueue:
2619 return MIRBuilder.buildInstr(Opcode)
2620 .addDef(Call->ReturnRegister)
2621 .addUse(GR->getSPIRVTypeID(Call->ReturnType));
2622 case SPIRV::OpIsValidEvent:
2623 return MIRBuilder.buildInstr(Opcode)
2624 .addDef(Call->ReturnRegister)
2625 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2626 .addUse(Call->Arguments[0]);
2627 case SPIRV::OpSetUserEventStatus:
2628 return MIRBuilder.buildInstr(Opcode)
2629 .addUse(Call->Arguments[0])
2630 .addUse(Call->Arguments[1]);
2631 case SPIRV::OpCaptureEventProfilingInfo:
2632 return MIRBuilder.buildInstr(Opcode)
2633 .addUse(Call->Arguments[0])
2634 .addUse(Call->Arguments[1])
2635 .addUse(Call->Arguments[2]);
2636 case SPIRV::OpBuildNDRange:
2637 return buildNDRange(Call, MIRBuilder, GR);
2638 case SPIRV::OpEnqueueKernel:
2639 return buildEnqueueKernel(Call, MIRBuilder, GR);
2640 default:
2641 return false;
2642 }
2643}
2644
2646 MachineIRBuilder &MIRBuilder,
2647 SPIRVGlobalRegistry *GR) {
2648 // Lookup the instruction opcode in the TableGen records.
2649 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2650 unsigned Opcode =
2651 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2652
2653 bool IsSet = Opcode == SPIRV::OpGroupAsyncCopy;
2654 Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType);
2655 if (Call->isSpirvOp())
2656 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2657 IsSet ? TypeReg : Register(0));
2658
2659 auto Scope = buildConstantIntReg32(SPIRV::Scope::Workgroup, MIRBuilder, GR);
2660
2661 switch (Opcode) {
2662 case SPIRV::OpGroupAsyncCopy: {
2663 SPIRVType *NewType =
2664 Call->ReturnType->getOpcode() == SPIRV::OpTypeEvent
2665 ? nullptr
2666 : GR->getOrCreateSPIRVTypeByName("spirv.Event", MIRBuilder, true);
2667 Register TypeReg = GR->getSPIRVTypeID(NewType ? NewType : Call->ReturnType);
2668 unsigned NumArgs = Call->Arguments.size();
2669 Register EventReg = Call->Arguments[NumArgs - 1];
2670 bool Res = MIRBuilder.buildInstr(Opcode)
2671 .addDef(Call->ReturnRegister)
2672 .addUse(TypeReg)
2673 .addUse(Scope)
2674 .addUse(Call->Arguments[0])
2675 .addUse(Call->Arguments[1])
2676 .addUse(Call->Arguments[2])
2677 .addUse(Call->Arguments.size() > 4
2678 ? Call->Arguments[3]
2679 : buildConstantIntReg32(1, MIRBuilder, GR))
2680 .addUse(EventReg);
2681 if (NewType != nullptr)
2682 insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder,
2683 MIRBuilder.getMF().getRegInfo());
2684 return Res;
2685 }
2686 case SPIRV::OpGroupWaitEvents:
2687 return MIRBuilder.buildInstr(Opcode)
2688 .addUse(Scope)
2689 .addUse(Call->Arguments[0])
2690 .addUse(Call->Arguments[1]);
2691 default:
2692 return false;
2693 }
2694}
2695
2696static bool generateConvertInst(const StringRef DemangledCall,
2698 MachineIRBuilder &MIRBuilder,
2699 SPIRVGlobalRegistry *GR) {
2700 // Lookup the conversion builtin in the TableGen records.
2701 const SPIRV::ConvertBuiltin *Builtin =
2702 SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);
2703
2704 if (!Builtin && Call->isSpirvOp()) {
2705 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2706 unsigned Opcode =
2707 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2708 return buildOpFromWrapper(MIRBuilder, Opcode, Call,
2709 GR->getSPIRVTypeID(Call->ReturnType));
2710 }
2711
2712 assert(Builtin && "Conversion builtin not found.");
2713 if (Builtin->IsSaturated)
2714 buildOpDecorate(Call->ReturnRegister, MIRBuilder,
2715 SPIRV::Decoration::SaturatedConversion, {});
2716 if (Builtin->IsRounded)
2717 buildOpDecorate(Call->ReturnRegister, MIRBuilder,
2718 SPIRV::Decoration::FPRoundingMode,
2719 {(unsigned)Builtin->RoundingMode});
2720
2721 std::string NeedExtMsg; // no errors if empty
2722 bool IsRightComponentsNumber = true; // check if input/output accepts vectors
2723 unsigned Opcode = SPIRV::OpNop;
2724 if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
2725 // Int -> ...
2726 if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
2727 // Int -> Int
2728 if (Builtin->IsSaturated)
2729 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
2730 : SPIRV::OpSatConvertSToU;
2731 else
2732 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert
2733 : SPIRV::OpSConvert;
2734 } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
2735 SPIRV::OpTypeFloat)) {
2736 // Int -> Float
2737 if (Builtin->IsBfloat16) {
2738 const auto *ST = static_cast<const SPIRVSubtarget *>(
2739 &MIRBuilder.getMF().getSubtarget());
2740 if (!ST->canUseExtension(
2741 SPIRV::Extension::SPV_INTEL_bfloat16_conversion))
2742 NeedExtMsg = "SPV_INTEL_bfloat16_conversion";
2743 IsRightComponentsNumber =
2744 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) ==
2745 GR->getScalarOrVectorComponentCount(Call->ReturnRegister);
2746 Opcode = SPIRV::OpConvertBF16ToFINTEL;
2747 } else {
2748 bool IsSourceSigned =
2749 DemangledCall[DemangledCall.find_first_of('(') + 1] != 'u';
2750 Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
2751 }
2752 }
2753 } else if (GR->isScalarOrVectorOfType(Call->Arguments[0],
2754 SPIRV::OpTypeFloat)) {
2755 // Float -> ...
2756 if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
2757 // Float -> Int
2758 if (Builtin->IsBfloat16) {
2759 const auto *ST = static_cast<const SPIRVSubtarget *>(
2760 &MIRBuilder.getMF().getSubtarget());
2761 if (!ST->canUseExtension(
2762 SPIRV::Extension::SPV_INTEL_bfloat16_conversion))
2763 NeedExtMsg = "SPV_INTEL_bfloat16_conversion";
2764 IsRightComponentsNumber =
2765 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) ==
2766 GR->getScalarOrVectorComponentCount(Call->ReturnRegister);
2767 Opcode = SPIRV::OpConvertFToBF16INTEL;
2768 } else {
2769 Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
2770 : SPIRV::OpConvertFToU;
2771 }
2772 } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
2773 SPIRV::OpTypeFloat)) {
2774 if (Builtin->IsTF32) {
2775 const auto *ST = static_cast<const SPIRVSubtarget *>(
2776 &MIRBuilder.getMF().getSubtarget());
2777 if (!ST->canUseExtension(
2778 SPIRV::Extension::SPV_INTEL_tensor_float32_conversion))
2779 NeedExtMsg = "SPV_INTEL_tensor_float32_conversion";
2780 IsRightComponentsNumber =
2781 GR->getScalarOrVectorComponentCount(Call->Arguments[0]) ==
2782 GR->getScalarOrVectorComponentCount(Call->ReturnRegister);
2783 Opcode = SPIRV::OpRoundFToTF32INTEL;
2784 } else {
2785 // Float -> Float
2786 Opcode = SPIRV::OpFConvert;
2787 }
2788 }
2789 }
2790
2791 if (!NeedExtMsg.empty()) {
2792 std::string DiagMsg = std::string(Builtin->Name) +
2793 ": the builtin requires the following SPIR-V "
2794 "extension: " +
2795 NeedExtMsg;
2796 report_fatal_error(DiagMsg.c_str(), false);
2797 }
2798 if (!IsRightComponentsNumber) {
2799 std::string DiagMsg =
2800 std::string(Builtin->Name) +
2801 ": result and argument must have the same number of components";
2802 report_fatal_error(DiagMsg.c_str(), false);
2803 }
2804 assert(Opcode != SPIRV::OpNop &&
2805 "Conversion between the types not implemented!");
2806
2807 MIRBuilder.buildInstr(Opcode)
2808 .addDef(Call->ReturnRegister)
2809 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2810 .addUse(Call->Arguments[0]);
2811 return true;
2812}
2813
2815 MachineIRBuilder &MIRBuilder,
2816 SPIRVGlobalRegistry *GR) {
2817 // Lookup the vector load/store builtin in the TableGen records.
2818 const SPIRV::VectorLoadStoreBuiltin *Builtin =
2819 SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
2820 Call->Builtin->Set);
2821 // Build extended instruction.
2822 auto MIB =
2823 MIRBuilder.buildInstr(SPIRV::OpExtInst)
2824 .addDef(Call->ReturnRegister)
2825 .addUse(GR->getSPIRVTypeID(Call->ReturnType))
2826 .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
2827 .addImm(Builtin->Number);
2828 for (auto Argument : Call->Arguments)
2829 MIB.addUse(Argument);
2830 if (Builtin->Name.contains("load") && Builtin->ElementCount > 1)
2831 MIB.addImm(Builtin->ElementCount);
2832
2833 // Rounding mode should be passed as a last argument in the MI for builtins
2834 // like "vstorea_halfn_r".
2835 if (Builtin->IsRounded)
2836 MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));
2837 return true;
2838}
2839
2841 MachineIRBuilder &MIRBuilder,
2842 SPIRVGlobalRegistry *GR) {
2843 // Lookup the instruction opcode in the TableGen records.
2844 const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
2845 unsigned Opcode =
2846 SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
2847 bool IsLoad = Opcode == SPIRV::OpLoad;
2848 // Build the instruction.
2849 auto MIB = MIRBuilder.buildInstr(Opcode);
2850 if (IsLoad) {
2851 MIB.addDef(Call->ReturnRegister);
2852 MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType));
2853 }
2854 // Add a pointer to the value to load/store.
2855 MIB.addUse(Call->Arguments[0]);
2856 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
2857 // Add a value to store.
2858 if (!IsLoad)
2859 MIB.addUse(Call->Arguments[1]);
2860 // Add optional memory attributes and an alignment.
2861 unsigned NumArgs = Call->Arguments.size();
2862 if ((IsLoad && NumArgs >= 2) || NumArgs >= 3)
2863 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 1 : 2], MRI));
2864 if ((IsLoad && NumArgs >= 3) || NumArgs >= 4)
2865 MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 2 : 3], MRI));
2866 return true;
2867}
2868
2869namespace SPIRV {
2870// Try to find a builtin function attributes by a demangled function name and
2871// return a tuple <builtin group, op code, ext instruction number>, or a special
2872// tuple value <-1, 0, 0> if the builtin function is not found.
2873// Not all builtin functions are supported, only those with a ready-to-use op
2874// code or instruction number defined in TableGen.
2875// TODO: consider a major rework of mapping demangled calls into a builtin
2876// functions to unify search and decrease number of individual cases.
2877std::tuple<int, unsigned, unsigned>
2878mapBuiltinToOpcode(const StringRef DemangledCall,
2879 SPIRV::InstructionSet::InstructionSet Set) {
2880 Register Reg;
2882 std::unique_ptr<const IncomingCall> Call =
2883 lookupBuiltin(DemangledCall, Set, Reg, nullptr, Args);
2884 if (!Call)
2885 return std::make_tuple(-1, 0, 0);
2886
2887 switch (Call->Builtin->Group) {
2888 case SPIRV::Relational:
2889 case SPIRV::Atomic:
2890 case SPIRV::Barrier:
2891 case SPIRV::CastToPtr:
2892 case SPIRV::ImageMiscQuery:
2893 case SPIRV::SpecConstant:
2894 case SPIRV::Enqueue:
2895 case SPIRV::AsyncCopy:
2896 case SPIRV::LoadStore:
2897 case SPIRV::CoopMatr:
2898 if (const auto *R =
2899 SPIRV::lookupNativeBuiltin(Call->Builtin->Name, Call->Builtin->Set))
2900 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);
2901 break;
2902 case SPIRV::Extended:
2903 if (const auto *R = SPIRV::lookupExtendedBuiltin(Call->Builtin->Name,
2904 Call->Builtin->Set))
2905 return std::make_tuple(Call->Builtin->Group, 0, R->Number);
2906 break;
2907 case SPIRV::VectorLoadStore:
2908 if (const auto *R = SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
2909 Call->Builtin->Set))
2910 return std::make_tuple(SPIRV::Extended, 0, R->Number);
2911 break;
2912 case SPIRV::Group:
2913 if (const auto *R = SPIRV::lookupGroupBuiltin(Call->Builtin->Name))
2914 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);
2915 break;
2916 case SPIRV::AtomicFloating:
2917 if (const auto *R = SPIRV::lookupAtomicFloatingBuiltin(Call->Builtin->Name))
2918 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);
2919 break;
2920 case SPIRV::IntelSubgroups:
2921 if (const auto *R = SPIRV::lookupIntelSubgroupsBuiltin(Call->Builtin->Name))
2922 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);
2923 break;
2924 case SPIRV::GroupUniform:
2925 if (const auto *R = SPIRV::lookupGroupUniformBuiltin(Call->Builtin->Name))
2926 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);
2927 break;
2928 case SPIRV::IntegerDot:
2929 if (const auto *R =
2930 SPIRV::lookupIntegerDotProductBuiltin(Call->Builtin->Name))
2931 return std::make_tuple(Call->Builtin->Group, R->Opcode, 0);
2932 break;
2933 case SPIRV::WriteImage:
2934 return std::make_tuple(Call->Builtin->Group, SPIRV::OpImageWrite, 0);
2935 case SPIRV::Select:
2936 return std::make_tuple(Call->Builtin->Group, TargetOpcode::G_SELECT, 0);
2937 case SPIRV::Construct:
2938 return std::make_tuple(Call->Builtin->Group, SPIRV::OpCompositeConstruct,
2939 0);
2940 case SPIRV::KernelClock:
2941 return std::make_tuple(Call->Builtin->Group, SPIRV::OpReadClockKHR, 0);
2942 default:
2943 return std::make_tuple(-1, 0, 0);
2944 }
2945 return std::make_tuple(-1, 0, 0);
2946}
2947
2948std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
2949 SPIRV::InstructionSet::InstructionSet Set,
2950 MachineIRBuilder &MIRBuilder,
2951 const Register OrigRet, const Type *OrigRetTy,
2952 const SmallVectorImpl<Register> &Args,
2953 SPIRVGlobalRegistry *GR, const CallBase &CB) {
2954 LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
2955
2956 // Lookup the builtin in the TableGen records.
2957 SPIRVType *SpvType = GR->getSPIRVTypeForVReg(OrigRet);
2958 assert(SpvType && "Inconsistent return register: expected valid type info");
2959 std::unique_ptr<const IncomingCall> Call =
2960 lookupBuiltin(DemangledCall, Set, OrigRet, SpvType, Args);
2961
2962 if (!Call) {
2963 LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
2964 return std::nullopt;
2965 }
2966
2967 // TODO: check if the provided args meet the builtin requirments.
2968 assert(Args.size() >= Call->Builtin->MinNumArgs &&
2969 "Too few arguments to generate the builtin");
2970 if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)
2971 LLVM_DEBUG(dbgs() << "More arguments provided than required!\n");
2972
2973 // Match the builtin with implementation based on the grouping.
2974 switch (Call->Builtin->Group) {
2975 case SPIRV::Extended:
2976 return generateExtInst(Call.get(), MIRBuilder, GR, CB);
2977 case SPIRV::Relational:
2978 return generateRelationalInst(Call.get(), MIRBuilder, GR);
2979 case SPIRV::Group:
2980 return generateGroupInst(Call.get(), MIRBuilder, GR);
2981 case SPIRV::Variable:
2982 return generateBuiltinVar(Call.get(), MIRBuilder, GR);
2983 case SPIRV::Atomic:
2984 return generateAtomicInst(Call.get(), MIRBuilder, GR);
2985 case SPIRV::AtomicFloating:
2986 return generateAtomicFloatingInst(Call.get(), MIRBuilder, GR);
2987 case SPIRV::Barrier:
2988 return generateBarrierInst(Call.get(), MIRBuilder, GR);
2989 case SPIRV::CastToPtr:
2990 return generateCastToPtrInst(Call.get(), MIRBuilder, GR);
2991 case SPIRV::Dot:
2992 case SPIRV::IntegerDot:
2993 return generateDotOrFMulInst(DemangledCall, Call.get(), MIRBuilder, GR);
2994 case SPIRV::Wave:
2995 return generateWaveInst(Call.get(), MIRBuilder, GR);
2996 case SPIRV::ICarryBorrow:
2997 return generateICarryBorrowInst(Call.get(), MIRBuilder, GR);
2998 case SPIRV::GetQuery:
2999 return generateGetQueryInst(Call.get(), MIRBuilder, GR);
3000 case SPIRV::ImageSizeQuery:
3001 return generateImageSizeQueryInst(Call.get(), MIRBuilder, GR);
3002 case SPIRV::ImageMiscQuery:
3003 return generateImageMiscQueryInst(Call.get(), MIRBuilder, GR);
3004 case SPIRV::ReadImage:
3005 return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
3006 case SPIRV::WriteImage:
3007 return generateWriteImageInst(Call.get(), MIRBuilder, GR);
3008 case SPIRV::SampleImage:
3009 return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
3010 case SPIRV::Select:
3011 return generateSelectInst(Call.get(), MIRBuilder);
3012 case SPIRV::Construct:
3013 return generateConstructInst(Call.get(), MIRBuilder, GR);
3014 case SPIRV::SpecConstant:
3015 return generateSpecConstantInst(Call.get(), MIRBuilder, GR);
3016 case SPIRV::Enqueue:
3017 return generateEnqueueInst(Call.get(), MIRBuilder, GR);
3018 case SPIRV::AsyncCopy:
3019 return generateAsyncCopy(Call.get(), MIRBuilder, GR);
3020 case SPIRV::Convert:
3021 return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);
3022 case SPIRV::VectorLoadStore:
3023 return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);
3024 case SPIRV::LoadStore:
3025 return generateLoadStoreInst(Call.get(), MIRBuilder, GR);
3026 case SPIRV::IntelSubgroups:
3027 return generateIntelSubgroupsInst(Call.get(), MIRBuilder, GR);
3028 case SPIRV::GroupUniform:
3029 return generateGroupUniformInst(Call.get(), MIRBuilder, GR);
3030 case SPIRV::KernelClock:
3031 return generateKernelClockInst(Call.get(), MIRBuilder, GR);
3032 case SPIRV::CoopMatr:
3033 return generateCoopMatrInst(Call.get(), MIRBuilder, GR);
3034 case SPIRV::ExtendedBitOps:
3035 return generateExtendedBitOpsInst(Call.get(), MIRBuilder, GR);
3036 case SPIRV::BindlessINTEL:
3037 return generateBindlessImageINTELInst(Call.get(), MIRBuilder, GR);
3038 case SPIRV::TernaryBitwiseINTEL:
3039 return generateTernaryBitwiseFunctionINTELInst(Call.get(), MIRBuilder, GR);
3040 case SPIRV::Block2DLoadStore:
3041 return generate2DBlockIOINTELInst(Call.get(), MIRBuilder, GR);
3042 case SPIRV::Pipe:
3043 return generatePipeInst(Call.get(), MIRBuilder, GR);
3044 case SPIRV::PredicatedLoadStore:
3045 return generatePredicatedLoadStoreInst(Call.get(), MIRBuilder, GR);
3046 }
3047 return false;
3048}
3049
3051 // Parse strings representing OpenCL builtin types.
3052 if (hasBuiltinTypePrefix(TypeStr)) {
3053 // OpenCL builtin types in demangled call strings have the following format:
3054 // e.g. ocl_image2d_ro
3055 [[maybe_unused]] bool IsOCLBuiltinType = TypeStr.consume_front("ocl_");
3056 assert(IsOCLBuiltinType && "Invalid OpenCL builtin prefix");
3057
3058 // Check if this is pointer to a builtin type and not just pointer
3059 // representing a builtin type. In case it is a pointer to builtin type,
3060 // this will require additional handling in the method calling
3061 // parseBuiltinCallArgumentBaseType(...) as this function only retrieves the
3062 // base types.
3063 if (TypeStr.ends_with("*"))
3064 TypeStr = TypeStr.slice(0, TypeStr.find_first_of(" *"));
3065
3066 return parseBuiltinTypeNameToTargetExtType("opencl." + TypeStr.str() + "_t",
3067 Ctx);
3068 }
3069
3070 // Parse type name in either "typeN" or "type vector[N]" format, where
3071 // N is the number of elements of the vector.
3072 Type *BaseType;
3073 unsigned VecElts = 0;
3074
3075 BaseType = parseBasicTypeName(TypeStr, Ctx);
3076 if (!BaseType)
3077 // Unable to recognize SPIRV type name.
3078 return nullptr;
3079
3080 // Handle "typeN*" or "type vector[N]*".
3081 TypeStr.consume_back("*");
3082
3083 if (TypeStr.consume_front(" vector["))
3084 TypeStr = TypeStr.substr(0, TypeStr.find(']'));
3085
3086 TypeStr.getAsInteger(10, VecElts);
3087 if (VecElts > 0)
3089 BaseType->isVoidTy() ? Type::getInt8Ty(Ctx) : BaseType, VecElts, false);
3090
3091 return BaseType;
3092}
3093
3095 const StringRef DemangledCall, LLVMContext &Ctx) {
3096 auto Pos1 = DemangledCall.find('(');
3097 if (Pos1 == StringRef::npos)
3098 return false;
3099 auto Pos2 = DemangledCall.find(')');
3100 if (Pos2 == StringRef::npos || Pos1 > Pos2)
3101 return false;
3102 DemangledCall.slice(Pos1 + 1, Pos2)
3103 .split(BuiltinArgsTypeStrs, ',', -1, false);
3104 return true;
3105}
3106
3108 unsigned ArgIdx, LLVMContext &Ctx) {
3109 SmallVector<StringRef, 10> BuiltinArgsTypeStrs;
3110 parseBuiltinTypeStr(BuiltinArgsTypeStrs, DemangledCall, Ctx);
3111 if (ArgIdx >= BuiltinArgsTypeStrs.size())
3112 return nullptr;
3113 StringRef TypeStr = BuiltinArgsTypeStrs[ArgIdx].trim();
3114 return parseBuiltinCallArgumentType(TypeStr, Ctx);
3115}
3116
3121
3122#define GET_BuiltinTypes_DECL
3123#define GET_BuiltinTypes_IMPL
3124
3129
3130#define GET_OpenCLTypes_DECL
3131#define GET_OpenCLTypes_IMPL
3132
3133#include "SPIRVGenTables.inc"
3134} // namespace SPIRV
3135
3136//===----------------------------------------------------------------------===//
3137// Misc functions for parsing builtin types.
3138//===----------------------------------------------------------------------===//
3139
3140static Type *parseTypeString(const StringRef Name, LLVMContext &Context) {
3141 if (Name.starts_with("void"))
3142 return Type::getVoidTy(Context);
3143 else if (Name.starts_with("int") || Name.starts_with("uint"))
3144 return Type::getInt32Ty(Context);
3145 else if (Name.starts_with("float"))
3146 return Type::getFloatTy(Context);
3147 else if (Name.starts_with("half"))
3148 return Type::getHalfTy(Context);
3149 report_fatal_error("Unable to recognize type!");
3150}
3151
3152//===----------------------------------------------------------------------===//
3153// Implementation functions for builtin types.
3154//===----------------------------------------------------------------------===//
3155
3157 const SPIRV::BuiltinType *TypeRecord,
3158 MachineIRBuilder &MIRBuilder,
3159 SPIRVGlobalRegistry *GR) {
3160 unsigned Opcode = TypeRecord->Opcode;
3161 // Create or get an existing type from GlobalRegistry.
3162 return GR->getOrCreateOpTypeByOpcode(ExtensionType, MIRBuilder, Opcode);
3163}
3164
3166 SPIRVGlobalRegistry *GR) {
3167 // Create or get an existing type from GlobalRegistry.
3168 return GR->getOrCreateOpTypeSampler(MIRBuilder);
3169}
3170
3171static SPIRVType *getPipeType(const TargetExtType *ExtensionType,
3172 MachineIRBuilder &MIRBuilder,
3173 SPIRVGlobalRegistry *GR) {
3174 assert(ExtensionType->getNumIntParameters() == 1 &&
3175 "Invalid number of parameters for SPIR-V pipe builtin!");
3176 // Create or get an existing type from GlobalRegistry.
3177 return GR->getOrCreateOpTypePipe(MIRBuilder,
3178 SPIRV::AccessQualifier::AccessQualifier(
3179 ExtensionType->getIntParameter(0)));
3180}
3181
3182static SPIRVType *getCoopMatrType(const TargetExtType *ExtensionType,
3183 MachineIRBuilder &MIRBuilder,
3184 SPIRVGlobalRegistry *GR) {
3185 assert(ExtensionType->getNumIntParameters() == 4 &&
3186 "Invalid number of parameters for SPIR-V coop matrices builtin!");
3187 assert(ExtensionType->getNumTypeParameters() == 1 &&
3188 "SPIR-V coop matrices builtin type must have a type parameter!");
3189 const SPIRVType *ElemType =
3190 GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder,
3191 SPIRV::AccessQualifier::ReadWrite, true);
3192 // Create or get an existing type from GlobalRegistry.
3193 return GR->getOrCreateOpTypeCoopMatr(
3194 MIRBuilder, ExtensionType, ElemType, ExtensionType->getIntParameter(0),
3195 ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2),
3196 ExtensionType->getIntParameter(3), true);
3197}
3198
3200 MachineIRBuilder &MIRBuilder,
3201 SPIRVGlobalRegistry *GR) {
3202 SPIRVType *OpaqueImageType = GR->getImageType(
3203 OpaqueType, SPIRV::AccessQualifier::ReadOnly, MIRBuilder);
3204 // Create or get an existing type from GlobalRegistry.
3205 return GR->getOrCreateOpTypeSampledImage(OpaqueImageType, MIRBuilder);
3206}
3207
3208static SPIRVType *getInlineSpirvType(const TargetExtType *ExtensionType,
3209 MachineIRBuilder &MIRBuilder,
3210 SPIRVGlobalRegistry *GR) {
3211 assert(ExtensionType->getNumIntParameters() == 3 &&
3212 "Inline SPIR-V type builtin takes an opcode, size, and alignment "
3213 "parameter");
3214 auto Opcode = ExtensionType->getIntParameter(0);
3215
3216 SmallVector<MCOperand> Operands;
3217 for (Type *Param : ExtensionType->type_params()) {
3218 if (const TargetExtType *ParamEType = dyn_cast<TargetExtType>(Param)) {
3219 if (ParamEType->getName() == "spirv.IntegralConstant") {
3220 assert(ParamEType->getNumTypeParameters() == 1 &&
3221 "Inline SPIR-V integral constant builtin must have a type "
3222 "parameter");
3223 assert(ParamEType->getNumIntParameters() == 1 &&
3224 "Inline SPIR-V integral constant builtin must have a "
3225 "value parameter");
3226
3227 auto OperandValue = ParamEType->getIntParameter(0);
3228 auto *OperandType = ParamEType->getTypeParameter(0);
3229
3230 const SPIRVType *OperandSPIRVType = GR->getOrCreateSPIRVType(
3231 OperandType, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, true);
3232
3234 OperandValue, MIRBuilder, OperandSPIRVType, true)));
3235 continue;
3236 } else if (ParamEType->getName() == "spirv.Literal") {
3237 assert(ParamEType->getNumTypeParameters() == 0 &&
3238 "Inline SPIR-V literal builtin does not take type "
3239 "parameters");
3240 assert(ParamEType->getNumIntParameters() == 1 &&
3241 "Inline SPIR-V literal builtin must have an integer "
3242 "parameter");
3243
3244 auto OperandValue = ParamEType->getIntParameter(0);
3245
3246 Operands.push_back(MCOperand::createImm(OperandValue));
3247 continue;
3248 }
3249 }
3250 const SPIRVType *TypeOperand = GR->getOrCreateSPIRVType(
3251 Param, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, true);
3252 Operands.push_back(MCOperand::createReg(GR->getSPIRVTypeID(TypeOperand)));
3253 }
3254
3255 return GR->getOrCreateUnknownType(ExtensionType, MIRBuilder, Opcode,
3256 Operands);
3257}
3258
3259static SPIRVType *getVulkanBufferType(const TargetExtType *ExtensionType,
3260 MachineIRBuilder &MIRBuilder,
3261 SPIRVGlobalRegistry *GR) {
3262 assert(ExtensionType->getNumTypeParameters() == 1 &&
3263 "Vulkan buffers have exactly one type for the type of the buffer.");
3264 assert(ExtensionType->getNumIntParameters() == 2 &&
3265 "Vulkan buffer have 2 integer parameters: storage class and is "
3266 "writable.");
3267
3268 auto *T = ExtensionType->getTypeParameter(0);
3269 auto SC = static_cast<SPIRV::StorageClass::StorageClass>(
3270 ExtensionType->getIntParameter(0));
3271 bool IsWritable = ExtensionType->getIntParameter(1);
3272 return GR->getOrCreateVulkanBufferType(MIRBuilder, T, SC, IsWritable);
3273}
3274
3275static SPIRVType *getLayoutType(const TargetExtType *ExtensionType,
3276 MachineIRBuilder &MIRBuilder,
3277 SPIRVGlobalRegistry *GR) {
3278 return GR->getOrCreateLayoutType(MIRBuilder, ExtensionType);
3279}
3280
3281namespace SPIRV {
3283 LLVMContext &Context) {
3284 StringRef NameWithParameters = TypeName;
3285
3286 // Pointers-to-opaque-structs representing OpenCL types are first translated
3287 // to equivalent SPIR-V types. OpenCL builtin type names should have the
3288 // following format: e.g. %opencl.event_t
3289 if (NameWithParameters.starts_with("opencl.")) {
3290 const SPIRV::OpenCLType *OCLTypeRecord =
3291 SPIRV::lookupOpenCLType(NameWithParameters);
3292 if (!OCLTypeRecord)
3293 report_fatal_error("Missing TableGen record for OpenCL type: " +
3294 NameWithParameters);
3295 NameWithParameters = OCLTypeRecord->SpirvTypeLiteral;
3296 // Continue with the SPIR-V builtin type...
3297 }
3298
3299 // Names of the opaque structs representing a SPIR-V builtins without
3300 // parameters should have the following format: e.g. %spirv.Event
3301 assert(NameWithParameters.starts_with("spirv.") &&
3302 "Unknown builtin opaque type!");
3303
3304 // Parameterized SPIR-V builtins names follow this format:
3305 // e.g. %spirv.Image._void_1_0_0_0_0_0_0, %spirv.Pipe._0
3306 if (!NameWithParameters.contains('_'))
3307 return TargetExtType::get(Context, NameWithParameters);
3308
3309 SmallVector<StringRef> Parameters;
3310 unsigned BaseNameLength = NameWithParameters.find('_') - 1;
3311 SplitString(NameWithParameters.substr(BaseNameLength + 1), Parameters, "_");
3312
3313 SmallVector<Type *, 1> TypeParameters;
3314 bool HasTypeParameter = !isDigit(Parameters[0][0]);
3315 if (HasTypeParameter)
3316 TypeParameters.push_back(parseTypeString(Parameters[0], Context));
3317 SmallVector<unsigned> IntParameters;
3318 for (unsigned i = HasTypeParameter ? 1 : 0; i < Parameters.size(); i++) {
3319 unsigned IntParameter = 0;
3320 bool ValidLiteral = !Parameters[i].getAsInteger(10, IntParameter);
3321 (void)ValidLiteral;
3322 assert(ValidLiteral &&
3323 "Invalid format of SPIR-V builtin parameter literal!");
3324 IntParameters.push_back(IntParameter);
3325 }
3326 return TargetExtType::get(Context,
3327 NameWithParameters.substr(0, BaseNameLength),
3328 TypeParameters, IntParameters);
3329}
3330
3332 SPIRV::AccessQualifier::AccessQualifier AccessQual,
3333 MachineIRBuilder &MIRBuilder,
3334 SPIRVGlobalRegistry *GR) {
3335 // In LLVM IR, SPIR-V and OpenCL builtin types are represented as either
3336 // target(...) target extension types or pointers-to-opaque-structs. The
3337 // approach relying on structs is deprecated and works only in the non-opaque
3338 // pointer mode (-opaque-pointers=0).
3339 // In order to maintain compatibility with LLVM IR generated by older versions
3340 // of Clang and LLVM/SPIR-V Translator, the pointers-to-opaque-structs are
3341 // "translated" to target extension types. This translation is temporary and
3342 // will be removed in the future release of LLVM.
3344 if (!BuiltinType)
3346 OpaqueType->getStructName().str(), MIRBuilder.getContext());
3347
3348 unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();
3349
3350 const StringRef Name = BuiltinType->getName();
3351 LLVM_DEBUG(dbgs() << "Lowering builtin type: " << Name << "\n");
3352
3353 SPIRVType *TargetType;
3354 if (Name == "spirv.Type") {
3355 TargetType = getInlineSpirvType(BuiltinType, MIRBuilder, GR);
3356 } else if (Name == "spirv.VulkanBuffer") {
3357 TargetType = getVulkanBufferType(BuiltinType, MIRBuilder, GR);
3358 } else if (Name == "spirv.Layout") {
3359 TargetType = getLayoutType(BuiltinType, MIRBuilder, GR);
3360 } else {
3361 // Lookup the demangled builtin type in the TableGen records.
3362 const SPIRV::BuiltinType *TypeRecord = SPIRV::lookupBuiltinType(Name);
3363 if (!TypeRecord)
3364 report_fatal_error("Missing TableGen record for builtin type: " + Name);
3365
3366 // "Lower" the BuiltinType into TargetType. The following get<...>Type
3367 // methods use the implementation details from TableGen records or
3368 // TargetExtType parameters to either create a new OpType<...> machine
3369 // instruction or get an existing equivalent SPIRVType from
3370 // GlobalRegistry.
3371
3372 switch (TypeRecord->Opcode) {
3373 case SPIRV::OpTypeImage:
3374 TargetType = GR->getImageType(BuiltinType, AccessQual, MIRBuilder);
3375 break;
3376 case SPIRV::OpTypePipe:
3377 TargetType = getPipeType(BuiltinType, MIRBuilder, GR);
3378 break;
3379 case SPIRV::OpTypeDeviceEvent:
3380 TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
3381 break;
3382 case SPIRV::OpTypeSampler:
3383 TargetType = getSamplerType(MIRBuilder, GR);
3384 break;
3385 case SPIRV::OpTypeSampledImage:
3386 TargetType = getSampledImageType(BuiltinType, MIRBuilder, GR);
3387 break;
3388 case SPIRV::OpTypeCooperativeMatrixKHR:
3389 TargetType = getCoopMatrType(BuiltinType, MIRBuilder, GR);
3390 break;
3391 default:
3392 TargetType =
3393 getNonParameterizedType(BuiltinType, TypeRecord, MIRBuilder, GR);
3394 break;
3395 }
3396 }
3397
3398 // Emit OpName instruction if a new OpType<...> instruction was added
3399 // (equivalent type was not found in GlobalRegistry).
3400 if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
3401 buildOpName(GR->getSPIRVTypeID(TargetType), Name, MIRBuilder);
3402
3403 return TargetType;
3404}
3405} // namespace SPIRV
3406} // namespace llvm
unsigned const MachineRegisterInfo * MRI
MachineInstrBuilder MachineInstrBuilder & DefMI
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU Lower Kernel Arguments
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
IRTranslator LLVM IR MI
#define I(x, y, z)
Definition MD5.cpp:58
Register Reg
Promote Memory to Register
Definition Mem2Reg.cpp:110
#define T
spirv structurize SPIRV
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.
#define LLVM_DEBUG(...)
Definition Debug.h:114
static const fltSemantics & IEEEsingle()
Definition APFloat.h:296
APInt bitcastToAPInt() const
Definition APFloat.h:1335
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
Definition APFloat.h:1061
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
Definition APInt.h:234
uint64_t getZExtValue() const
Get zero extended value.
Definition APInt.h:1540
This class represents an incoming formal argument to a Function.
Definition Argument.h:32
ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...
Definition ArrayRef.h:41
static LLVM_ABI ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Base class for all callable instructions (InvokeInst and CallInst) Holds everything related to callin...
@ ICMP_ULT
unsigned less than
Definition InstrTypes.h:701
@ ICMP_NE
not equal
Definition InstrTypes.h:698
const APFloat & getValueAPF() const
Definition Constants.h:320
const APInt & getValue() const
Return the constant as an APInt value reference.
Definition Constants.h:154
A parsed version of the target data layout string in and methods for querying it.
Definition DataLayout.h:63
Tagged union holding either a T or a Error.
Definition Error.h:485
Class to represent fixed width SIMD vectors.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Definition Function.cpp:359
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
Definition Type.cpp:319
static constexpr LLT vector(ElementCount EC, unsigned ScalarSizeInBits)
Get a low-level vector of some number of elements and element width.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
This is an important class for using LLVM in a threaded context.
Definition LLVMContext.h:68
static MCOperand createReg(MCRegister Reg)
Definition MCInst.h:138
static MCOperand createImm(int64_t Val)
Definition MCInst.h:145
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, std::optional< unsigned > Flags=std::nullopt)
Build and insert a Res = G_ICMP Pred, Op0, Op1.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
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
virtual MachineInstrBuilder buildConstant(const DstOp &Res, const ConstantInt &Val)
Build and insert Res = G_CONSTANT Val.
Register getReg(unsigned Idx) const
Get the register for the operand index.
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.
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
const MachineInstrBuilder & addDef(Register RegNo, unsigned Flags=0, unsigned SubReg=0) const
Add a virtual register definition operand.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
LLVM_ABI void copyIRFlags(const Instruction &I)
Copy all flags to MachineInst MIFlags.
const MachineOperand & getOperand(unsigned i) const
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,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
static LLVM_ABI 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:107
SPIRVType * getOrCreateOpTypePipe(MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AccQual)
SPIRVType * getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
SPIRVType * getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
void assignSPIRVTypeToVReg(SPIRVType *Type, Register VReg, const MachineFunction &MF)
Register buildGlobalVariable(Register Reg, SPIRVType *BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
const Type * getTypeForSPIRVType(const SPIRVType *Ty) const
SPIRVType * getOrCreateUnknownType(const Type *Ty, MachineIRBuilder &MIRBuilder, unsigned Opcode, const ArrayRef< MCOperand > Operands)
unsigned getScalarOrVectorComponentCount(Register VReg) const
SPIRVType * getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
SPIRVType * getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
SPIRVType * getOrCreateOpTypeByOpcode(const Type *Ty, MachineIRBuilder &MIRBuilder, unsigned Opcode)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType=nullptr)
SPIRVType * getPointeeType(SPIRVType *PtrType)
Register getSPIRVTypeID(const SPIRVType *SpirvType) const
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
SPIRVType * getOrCreateOpTypeSampledImage(SPIRVType *ImageType, MachineIRBuilder &MIRBuilder)
SPIRVType * getOrCreateVulkanBufferType(MachineIRBuilder &MIRBuilder, Type *ElemType, SPIRV::StorageClass::StorageClass SC, bool IsWritable, bool EmitIr=false)
SPIRVType * getOrCreateSPIRVTypeByName(StringRef TypeStr, MachineIRBuilder &MIRBuilder, bool EmitIR, SPIRV::StorageClass::StorageClass SC=SPIRV::StorageClass::Function, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite)
SPIRVType * getOrCreateLayoutType(MachineIRBuilder &MIRBuilder, const TargetExtType *T, bool EmitIr=false)
Register getOrCreateConsIntVector(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType, bool EmitIR)
const TargetRegisterClass * getRegClass(SPIRVType *SpvType) const
SPIRVType * getOrCreateSPIRVVectorType(SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVType * getOrCreateOpTypeCoopMatr(MachineIRBuilder &MIRBuilder, const TargetExtType *ExtensionType, const SPIRVType *ElemType, uint32_t Scope, uint32_t Rows, uint32_t Columns, uint32_t Use, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
SPIRVType * getOrCreateOpTypeDeviceEvent(MachineIRBuilder &MIRBuilder)
SPIRVType * getImageType(const TargetExtType *ExtensionType, const SPIRV::AccessQualifier::AccessQualifier Qualifier, MachineIRBuilder &MIRBuilder)
SPIRVType * getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
Register buildConstantInt(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType, bool EmitIR, bool ZeroAsNull=true)
LLT getRegType(SPIRVType *SpvType) const
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
SPIRVType * getOrCreateOpTypeSampler(MachineIRBuilder &MIRBuilder)
Register buildConstantSampler(Register Res, unsigned AddrMode, unsigned Param, unsigned FilerMode, MachineIRBuilder &MIRBuilder)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)
unsigned getScalarOrVectorBitWidth(const SPIRVType *Type) const
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
StringRef - Represent a constant reference to a string, i.e.
Definition StringRef.h:55
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
Definition StringRef.h:702
bool consume_back(StringRef Suffix)
Returns true if this StringRef has the given suffix and removes that suffix.
Definition StringRef.h:657
bool getAsInteger(unsigned Radix, T &Result) const
Parse the current string as an integer of the specified radix.
Definition StringRef.h:472
std::string str() const
str - Get the contents as an std::string.
Definition StringRef.h:225
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
Definition StringRef.h:573
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
Definition StringRef.h:261
bool contains_insensitive(StringRef Other) const
Return true if the given string is a substring of *this, and false otherwise.
Definition StringRef.h:438
StringRef slice(size_t Start, size_t End) const
Return a reference to the substring from [Start, End).
Definition StringRef.h:686
constexpr size_t size() const
size - Get the string size.
Definition StringRef.h:146
bool contains(StringRef Other) const
Return true if the given string is a substring of *this, and false otherwise.
Definition StringRef.h:426
bool consume_front(StringRef Prefix)
Returns true if this StringRef has the given prefix and removes that prefix.
Definition StringRef.h:637
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:376
size_t rfind(char C, size_t From=npos) const
Search for the last character C in the string.
Definition StringRef.h:345
size_t find(char C, size_t From=0) const
Search for the first character C in the string.
Definition StringRef.h:293
bool ends_with(StringRef Suffix) const
Check if this string ends with the given Suffix.
Definition StringRef.h:273
static constexpr size_t npos
Definition StringRef.h:57
A switch()-like statement whose cases are string literals.
StringSwitch & EndsWith(StringLiteral S, T Value)
Class to represent struct types.
Class to represent target extensions types, which are generally unintrospectable from target-independ...
ArrayRef< Type * > type_params() const
Return the type parameters for this particular target extension type.
unsigned getNumIntParameters() const
static LLVM_ABI TargetExtType * get(LLVMContext &Context, StringRef Name, ArrayRef< Type * > Types={}, ArrayRef< unsigned > Ints={})
Return a target extension type having the specified name and optional type and integer parameters.
Definition Type.cpp:908
Type * getTypeParameter(unsigned i) const
unsigned getNumTypeParameters() const
unsigned getIntParameter(unsigned i) const
The instances of the Type class are immutable: once they are created, they are never changed.
Definition Type.h:45
static LLVM_ABI IntegerType * getInt32Ty(LLVMContext &C)
Definition Type.cpp:297
LLVM_ABI StringRef getStructName() const
static LLVM_ABI Type * getVoidTy(LLVMContext &C)
Definition Type.cpp:281
static LLVM_ABI IntegerType * getInt8Ty(LLVMContext &C)
Definition Type.cpp:295
static LLVM_ABI Type * getFloatTy(LLVMContext &C)
Definition Type.cpp:285
static LLVM_ABI Type * getHalfTy(LLVMContext &C)
Definition Type.cpp:283
LLVM Value Representation.
Definition Value.h:75
LLVM_ABI Value(Type *Ty, unsigned scid)
Definition Value.cpp:53
static LLVM_ABI VectorType * get(Type *ElementType, ElementCount EC)
This static method is the primary way to construct an VectorType.
Represents a version number in the form major[.minor[.subminor[.build]]].
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
Definition ilist_node.h:348
CallInst * Call
LLVM_C_ABI 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:883
std::string lookupBuiltinNameHelper(StringRef DemangledCall, FPDecorationId *DecorationId)
Parses the name part of the demangled builtin call.
Type * parseBuiltinCallArgumentType(StringRef TypeStr, LLVMContext &Ctx)
bool parseBuiltinTypeStr(SmallVector< StringRef, 10 > &BuiltinArgsTypeStrs, const StringRef DemangledCall, LLVMContext &Ctx)
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, const CallBase &CB)
std::tuple< int, unsigned, unsigned > mapBuiltinToOpcode(const StringRef DemangledCall, SPIRV::InstructionSet::InstructionSet Set)
Helper function for finding a builtin function attributes by a demangled function name.
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...
SPIRVType * lowerBuiltinType(const Type *OpaqueType, SPIRV::AccessQualifier::AccessQualifier AccessQual, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
This is an optimization pass for GlobalISel generic memory operations.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
static bool build2DBlockIOINTELInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building Intel's 2d block io instructions.
static SPIRVType * getVulkanBufferType(const TargetExtType *ExtensionType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateExtInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, const CallBase &CB)
void 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 bool generateBindlessImageINTELInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateGetQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRVType * getInlineSpirvType(const TargetExtType *ExtensionType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateConstructInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building atomic flag instructions (e.g.
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.
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
Definition Casting.h:643
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
Definition InstrProf.h:296
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
static bool buildExtendedBitOpsInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building extended bit operations.
static const Type * getBlockStructType(Register ParamReg, MachineRegisterInfo *MRI)
static bool generateGroupInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
FPDecorationId demangledPostfixToDecorationId(const std::string &S)
Definition SPIRVUtils.h:529
static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim)
static bool generateICarryBorrowInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static Register buildScopeReg(Register CLScopeRegister, SPIRV::Scope::Scope Scope, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI)
FPDecorationId
Definition SPIRVUtils.h:527
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)
static Register buildConstantIntReg32(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static SPIRVType * getSampledImageType(const TargetExtType *OpaqueType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
Definition SPIRVUtils.h:239
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)
LLVM_ABI void SplitString(StringRef Source, SmallVectorImpl< StringRef > &OutFragments, StringRef Delimiters=" \t\n\v\f\r")
SplitString - Split up the specified string according to the specified delimiters,...
static SPIRVType * getCoopMatrType(const TargetExtType *ExtensionType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateKernelClockInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static void setRegClassIfNull(Register Reg, MachineRegisterInfo *MRI, SPIRVGlobalRegistry *GR)
static bool generateGroupUniformInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateWaveInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
Register createVirtualRegister(SPIRVType *SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
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)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition Debug.cpp:207
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)
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
Definition Error.cpp:167
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)
bool isDigit(char C)
Checks if character C is one of the 10 decimal digits.
static SPIRV::SamplerAddressingMode::SamplerAddressingMode getSamplerAddressingModeFromBitmask(unsigned Bitmask)
static bool generateAtomicInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
const MachineInstr SPIRVType
static SPIRVType * getLayoutType(const TargetExtType *ExtensionType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateDotOrFMulInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static Register buildBuiltinVariableLoad(MachineIRBuilder &MIRBuilder, SPIRVType *VariableType, SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, LLT LLType, Register Reg=Register(0), bool isConst=true, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageTy={ SPIRV::LinkageType::Import})
Helper function for building a load instruction for loading a builtin global variable of BuiltinValue...
static bool generateConvertInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateTernaryBitwiseFunctionINTELInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateCastToPtrInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
Definition SPIRVUtils.h:224
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 generateExtendedBitOpsInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool buildPipeInst(const SPIRV::IncomingCall *Call, unsigned Opcode, unsigned Scope, 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)
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 generateCoopMatrInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool buildBindlessImageINTELInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building Intel's bindless image instructions.
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)
constexpr unsigned BitWidth
OutputIt move(R &&Range, OutputIt Out)
Provide wrappers to std::move which take ranges instead of having to pass begin/end explicitly.
Definition STLExtras.h:1867
static bool generate2DBlockIOINTELInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateReadImageInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
Definition Casting.h:559
bool hasBuiltinTypePrefix(StringRef Name)
static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Type * getMDOperandAsType(const MDNode *N, unsigned I)
static bool generatePipeInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool buildTernaryBitwiseFunctionINTELInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building Intel's OpBitwiseFunctionINTEL instruction.
static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building atomic instructions.
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)
static bool generatePredicatedLoadStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
static bool generateAtomicFloatingInst(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 bool buildOpFromWrapper(MachineIRBuilder &MIRBuilder, unsigned Opcode, const SPIRV::IncomingCall *Call, Register TypeReg, ArrayRef< uint32_t > ImmArgs={})
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)
Implement std::hash so that hash_code can be used in STL containers.
Definition BitVector.h:867
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
InstructionSet::InstructionSet Set
const SmallVectorImpl< Register > & Arguments
const std::string BuiltinName
const SPIRVType * ReturnType
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