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