LLVM  16.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 "SPIRVUtils.h"
18 #include "llvm/IR/IntrinsicsSPIRV.h"
19 #include <string>
20 #include <tuple>
21 
22 #define DEBUG_TYPE "spirv-builtins"
23 
24 namespace llvm {
25 namespace SPIRV {
26 #define GET_BuiltinGroup_DECL
27 #include "SPIRVGenTables.inc"
28 
31  InstructionSet::InstructionSet Set;
32  BuiltinGroup Group;
33  uint8_t MinNumArgs;
34  uint8_t MaxNumArgs;
35 };
36 
37 #define GET_DemangledBuiltins_DECL
38 #define GET_DemangledBuiltins_IMPL
39 
40 struct IncomingCall {
41  const std::string BuiltinName;
43 
47 
48  IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin,
54 };
55 
56 struct NativeBuiltin {
58  InstructionSet::InstructionSet Set;
60 };
61 
62 #define GET_NativeBuiltins_DECL
63 #define GET_NativeBuiltins_IMPL
64 
65 struct GroupBuiltin {
69  bool IsElect;
70  bool IsAllOrAny;
71  bool IsAllEqual;
72  bool IsBallot;
76  bool IsLogical;
78  bool HasBoolArg;
79 };
80 
81 #define GET_GroupBuiltins_DECL
82 #define GET_GroupBuiltins_IMPL
83 
84 struct GetBuiltin {
86  InstructionSet::InstructionSet Set;
87  BuiltIn::BuiltIn Value;
88 };
89 
90 using namespace BuiltIn;
91 #define GET_GetBuiltins_DECL
92 #define GET_GetBuiltins_IMPL
93 
96  InstructionSet::InstructionSet Set;
98 };
99 
100 #define GET_ImageQueryBuiltins_DECL
101 #define GET_ImageQueryBuiltins_IMPL
102 
105  InstructionSet::InstructionSet Set;
108  bool IsRounded;
109  FPRoundingMode::FPRoundingMode RoundingMode;
110 };
111 
114  InstructionSet::InstructionSet Set;
116  bool IsRounded;
117  FPRoundingMode::FPRoundingMode RoundingMode;
118 };
119 
120 using namespace FPRoundingMode;
121 #define GET_ConvertBuiltins_DECL
122 #define GET_ConvertBuiltins_IMPL
123 
124 using namespace InstructionSet;
125 #define GET_VectorLoadStoreBuiltins_DECL
126 #define GET_VectorLoadStoreBuiltins_IMPL
127 
128 #define GET_CLMemoryScope_DECL
129 #define GET_CLSamplerAddressingMode_DECL
130 #define GET_CLMemoryFenceFlags_DECL
131 #define GET_ExtendedBuiltins_DECL
132 #include "SPIRVGenTables.inc"
133 } // namespace SPIRV
134 
135 //===----------------------------------------------------------------------===//
136 // Misc functions for looking up builtins and veryfying requirements using
137 // TableGen records
138 //===----------------------------------------------------------------------===//
139 
140 /// Looks up the demangled builtin call in the SPIRVBuiltins.td records using
141 /// the provided \p DemangledCall and specified \p Set.
142 ///
143 /// The lookup follows the following algorithm, returning the first successful
144 /// match:
145 /// 1. Search with the plain demangled name (expecting a 1:1 match).
146 /// 2. Search with the prefix before or suffix after the demangled name
147 /// signyfying the type of the first argument.
148 ///
149 /// \returns Wrapper around the demangled call and found builtin definition.
150 static std::unique_ptr<const SPIRV::IncomingCall>
151 lookupBuiltin(StringRef DemangledCall,
152  SPIRV::InstructionSet::InstructionSet Set,
153  Register ReturnRegister, const SPIRVType *ReturnType,
155  // Extract the builtin function name and types of arguments from the call
156  // skeleton.
157  std::string BuiltinName =
158  DemangledCall.substr(0, DemangledCall.find('(')).str();
159 
160  // Check if the extracted name contains type information between angle
161  // brackets. If so, the builtin is an instantiated template - needs to have
162  // the information after angle brackets and return type removed.
163  if (BuiltinName.find('<') && BuiltinName.back() == '>') {
164  BuiltinName = BuiltinName.substr(0, BuiltinName.find('<'));
165  BuiltinName = BuiltinName.substr(BuiltinName.find_last_of(" ") + 1);
166  }
167 
168  // Check if the extracted name begins with "__spirv_ImageSampleExplicitLod"
169  // contains return type information at the end "_R<type>", if so extract the
170  // plain builtin name without the type information.
171  if (StringRef(BuiltinName).contains("__spirv_ImageSampleExplicitLod") &&
172  StringRef(BuiltinName).contains("_R")) {
173  BuiltinName = BuiltinName.substr(0, BuiltinName.find("_R"));
174  }
175 
176  SmallVector<StringRef, 10> BuiltinArgumentTypes;
177  StringRef BuiltinArgs =
178  DemangledCall.slice(DemangledCall.find('(') + 1, DemangledCall.find(')'));
179  BuiltinArgs.split(BuiltinArgumentTypes, ',', -1, false);
180 
181  // Look up the builtin in the defined set. Start with the plain demangled
182  // name, expecting a 1:1 match in the defined builtin set.
183  const SPIRV::DemangledBuiltin *Builtin;
184  if ((Builtin = SPIRV::lookupBuiltin(BuiltinName, Set)))
185  return std::make_unique<SPIRV::IncomingCall>(
186  BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
187 
188  // If the initial look up was unsuccessful and the demangled call takes at
189  // least 1 argument, add a prefix or suffix signifying the type of the first
190  // argument and repeat the search.
191  if (BuiltinArgumentTypes.size() >= 1) {
192  char FirstArgumentType = BuiltinArgumentTypes[0][0];
193  // Prefix to be added to the builtin's name for lookup.
194  // For example, OpenCL "abs" taking an unsigned value has a prefix "u_".
195  std::string Prefix;
196 
197  switch (FirstArgumentType) {
198  // Unsigned:
199  case 'u':
200  if (Set == SPIRV::InstructionSet::OpenCL_std)
201  Prefix = "u_";
202  else if (Set == SPIRV::InstructionSet::GLSL_std_450)
203  Prefix = "u";
204  break;
205  // Signed:
206  case 'c':
207  case 's':
208  case 'i':
209  case 'l':
210  if (Set == SPIRV::InstructionSet::OpenCL_std)
211  Prefix = "s_";
212  else if (Set == SPIRV::InstructionSet::GLSL_std_450)
213  Prefix = "s";
214  break;
215  // Floating-point:
216  case 'f':
217  case 'd':
218  case 'h':
219  if (Set == SPIRV::InstructionSet::OpenCL_std ||
220  Set == SPIRV::InstructionSet::GLSL_std_450)
221  Prefix = "f";
222  break;
223  }
224 
225  // If argument-type name prefix was added, look up the builtin again.
226  if (!Prefix.empty() &&
227  (Builtin = SPIRV::lookupBuiltin(Prefix + BuiltinName, Set)))
228  return std::make_unique<SPIRV::IncomingCall>(
229  BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
230 
231  // If lookup with a prefix failed, find a suffix to be added to the
232  // builtin's name for lookup. For example, OpenCL "group_reduce_max" taking
233  // an unsigned value has a suffix "u".
234  std::string Suffix;
235 
236  switch (FirstArgumentType) {
237  // Unsigned:
238  case 'u':
239  Suffix = "u";
240  break;
241  // Signed:
242  case 'c':
243  case 's':
244  case 'i':
245  case 'l':
246  Suffix = "s";
247  break;
248  // Floating-point:
249  case 'f':
250  case 'd':
251  case 'h':
252  Suffix = "f";
253  break;
254  }
255 
256  // If argument-type name suffix was added, look up the builtin again.
257  if (!Suffix.empty() &&
258  (Builtin = SPIRV::lookupBuiltin(BuiltinName + Suffix, Set)))
259  return std::make_unique<SPIRV::IncomingCall>(
260  BuiltinName, Builtin, ReturnRegister, ReturnType, Arguments);
261  }
262 
263  // No builtin with such name was found in the set.
264  return nullptr;
265 }
266 
267 //===----------------------------------------------------------------------===//
268 // Helper functions for building misc instructions
269 //===----------------------------------------------------------------------===//
270 
271 /// Helper function building either a resulting scalar or vector bool register
272 /// depending on the expected \p ResultType.
273 ///
274 /// \returns Tuple of the resulting register and its type.
275 static std::tuple<Register, SPIRVType *>
276 buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType,
277  SPIRVGlobalRegistry *GR) {
278  LLT Type;
279  SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
280 
281  if (ResultType->getOpcode() == SPIRV::OpTypeVector) {
282  unsigned VectorElements = ResultType->getOperand(2).getImm();
283  BoolType =
284  GR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder);
286  cast<FixedVectorType>(GR->getTypeForSPIRVType(BoolType));
287  Type = LLT::vector(LLVMVectorType->getElementCount(), 1);
288  } else {
289  Type = LLT::scalar(1);
290  }
291 
292  Register ResultRegister =
294  GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());
295  return std::make_tuple(ResultRegister, BoolType);
296 }
297 
298 /// Helper function for building either a vector or scalar select instruction
299 /// depending on the expected \p ResultType.
300 static bool buildSelectInst(MachineIRBuilder &MIRBuilder,
301  Register ReturnRegister, Register SourceRegister,
302  const SPIRVType *ReturnType,
303  SPIRVGlobalRegistry *GR) {
304  Register TrueConst, FalseConst;
305 
306  if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {
307  unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType);
309  TrueConst = GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType);
310  FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType);
311  } else {
312  TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType);
313  FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType);
314  }
315  return MIRBuilder.buildSelect(ReturnRegister, SourceRegister, TrueConst,
316  FalseConst);
317 }
318 
319 /// Helper function for building a load instruction loading into the
320 /// \p DestinationReg.
322  MachineIRBuilder &MIRBuilder,
323  SPIRVGlobalRegistry *GR, LLT LowLevelType,
324  Register DestinationReg = Register(0)) {
325  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
326  if (!DestinationReg.isValid()) {
327  DestinationReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
328  MRI->setType(DestinationReg, LLT::scalar(32));
329  GR->assignSPIRVTypeToVReg(BaseType, DestinationReg, MIRBuilder.getMF());
330  }
331  // TODO: consider using correct address space and alignment (p0 is canonical
332  // type for selection though).
334  MIRBuilder.buildLoad(DestinationReg, PtrRegister, PtrInfo, Align());
335  return DestinationReg;
336 }
337 
338 /// Helper function for building a load instruction for loading a builtin global
339 /// variable of \p BuiltinValue value.
341  SPIRVType *VariableType,
343  SPIRV::BuiltIn::BuiltIn BuiltinValue,
344  LLT LLType,
345  Register Reg = Register(0)) {
346  Register NewRegister =
347  MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
348  MIRBuilder.getMRI()->setType(NewRegister,
349  LLT::pointer(0, GR->getPointerSize()));
350  SPIRVType *PtrType = GR->getOrCreateSPIRVPointerType(
351  VariableType, MIRBuilder, SPIRV::StorageClass::Input);
352  GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
353 
354  // Set up the global OpVariable with the necessary builtin decorations.
355  Register Variable = GR->buildGlobalVariable(
356  NewRegister, PtrType, getLinkStringForBuiltIn(BuiltinValue), nullptr,
357  SPIRV::StorageClass::Input, nullptr, true, true,
358  SPIRV::LinkageType::Import, MIRBuilder, false);
359 
360  // Load the value from the global variable.
361  Register LoadedRegister =
362  buildLoadInst(VariableType, Variable, MIRBuilder, GR, LLType, Reg);
363  MIRBuilder.getMRI()->setType(LoadedRegister, LLType);
364  return LoadedRegister;
365 }
366 
367 /// Helper external function for inserting ASSIGN_TYPE instuction between \p Reg
368 /// and its definition, set the new register as a destination of the definition,
369 /// assign SPIRVType to both registers. If SpirvTy is provided, use it as
370 /// SPIRVType in ASSIGN_TYPE, otherwise create it from \p Ty. Defined in
371 /// SPIRVPreLegalizer.cpp.
372 extern Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy,
373  SPIRVGlobalRegistry *GR,
374  MachineIRBuilder &MIB,
375  MachineRegisterInfo &MRI);
376 
377 // TODO: Move to TableGen.
378 static SPIRV::MemorySemantics::MemorySemantics
379 getSPIRVMemSemantics(std::memory_order MemOrder) {
380  switch (MemOrder) {
381  case std::memory_order::memory_order_relaxed:
382  return SPIRV::MemorySemantics::None;
383  case std::memory_order::memory_order_acquire:
384  return SPIRV::MemorySemantics::Acquire;
385  case std::memory_order::memory_order_release:
386  return SPIRV::MemorySemantics::Release;
387  case std::memory_order::memory_order_acq_rel:
388  return SPIRV::MemorySemantics::AcquireRelease;
389  case std::memory_order::memory_order_seq_cst:
390  return SPIRV::MemorySemantics::SequentiallyConsistent;
391  default:
392  llvm_unreachable("Unknown CL memory scope");
393  }
394 }
395 
396 static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope) {
397  switch (ClScope) {
398  case SPIRV::CLMemoryScope::memory_scope_work_item:
399  return SPIRV::Scope::Invocation;
400  case SPIRV::CLMemoryScope::memory_scope_work_group:
401  return SPIRV::Scope::Workgroup;
402  case SPIRV::CLMemoryScope::memory_scope_device:
403  return SPIRV::Scope::Device;
404  case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:
405  return SPIRV::Scope::CrossDevice;
406  case SPIRV::CLMemoryScope::memory_scope_sub_group:
407  return SPIRV::Scope::Subgroup;
408  }
409  llvm_unreachable("Unknown CL memory scope");
410 }
411 
414  unsigned BitWidth = 32) {
415  SPIRVType *IntType = GR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
416  return GR->buildConstantInt(Val, MIRBuilder, IntType);
417 }
418 
419 /// Helper function for translating atomic init to OpStore.
420 static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call,
421  MachineIRBuilder &MIRBuilder) {
422  assert(Call->Arguments.size() == 2 &&
423  "Need 2 arguments for atomic init translation");
424 
425  MIRBuilder.buildInstr(SPIRV::OpStore)
426  .addUse(Call->Arguments[0])
427  .addUse(Call->Arguments[1]);
428  return true;
429 }
430 
431 /// Helper function for building an atomic load instruction.
432 static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call,
433  MachineIRBuilder &MIRBuilder,
434  SPIRVGlobalRegistry *GR) {
435  Register PtrRegister = Call->Arguments[0];
436  // TODO: if true insert call to __translate_ocl_memory_sccope before
437  // OpAtomicLoad and the function implementation. We can use Translator's
438  // output for transcoding/atomic_explicit_arguments.cl as an example.
439  Register ScopeRegister;
440  if (Call->Arguments.size() > 1)
441  ScopeRegister = Call->Arguments[1];
442  else
443  ScopeRegister = buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
444 
445  Register MemSemanticsReg;
446  if (Call->Arguments.size() > 2) {
447  // TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
448  MemSemanticsReg = Call->Arguments[2];
449  } else {
450  int Semantics =
451  SPIRV::MemorySemantics::SequentiallyConsistent |
453  MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
454  }
455 
456  MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
457  .addDef(Call->ReturnRegister)
458  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
459  .addUse(PtrRegister)
460  .addUse(ScopeRegister)
461  .addUse(MemSemanticsReg);
462  return true;
463 }
464 
465 /// Helper function for building an atomic store instruction.
467  MachineIRBuilder &MIRBuilder,
468  SPIRVGlobalRegistry *GR) {
469  Register ScopeRegister =
470  buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
471  Register PtrRegister = Call->Arguments[0];
472  int Semantics =
473  SPIRV::MemorySemantics::SequentiallyConsistent |
475  Register MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
476 
477  MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
478  .addUse(PtrRegister)
479  .addUse(ScopeRegister)
480  .addUse(MemSemanticsReg)
481  .addUse(Call->Arguments[1]);
482  return true;
483 }
484 
485 /// Helper function for building an atomic compare-exchange instruction.
487  MachineIRBuilder &MIRBuilder,
488  SPIRVGlobalRegistry *GR) {
489  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
490  unsigned Opcode =
491  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
492  bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");
493  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
494 
495  Register ObjectPtr = Call->Arguments[0]; // Pointer (volatile A *object.)
496  Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
497  Register Desired = Call->Arguments[2]; // Value (C Desired).
498  SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);
499  LLT DesiredLLT = MRI->getType(Desired);
500 
501  assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
502  SPIRV::OpTypePointer);
503  unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
504  assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
505  : ExpectedType == SPIRV::OpTypePointer);
506  assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
507 
508  SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);
509  assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
511  SpvObjectPtrTy->getOperand(1).getImm());
512  auto MemSemStorage = getMemSemanticsForStorageClass(StorageClass);
513 
514  Register MemSemEqualReg;
515  Register MemSemUnequalReg;
516  uint64_t MemSemEqual =
517  IsCmpxchg
518  ? SPIRV::MemorySemantics::None
519  : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
520  uint64_t MemSemUnequal =
521  IsCmpxchg
522  ? SPIRV::MemorySemantics::None
523  : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
524  if (Call->Arguments.size() >= 4) {
525  assert(Call->Arguments.size() >= 5 &&
526  "Need 5+ args for explicit atomic cmpxchg");
527  auto MemOrdEq =
528  static_cast<std::memory_order>(getIConstVal(Call->Arguments[3], MRI));
529  auto MemOrdNeq =
530  static_cast<std::memory_order>(getIConstVal(Call->Arguments[4], MRI));
531  MemSemEqual = getSPIRVMemSemantics(MemOrdEq) | MemSemStorage;
532  MemSemUnequal = getSPIRVMemSemantics(MemOrdNeq) | MemSemStorage;
533  if (MemOrdEq == MemSemEqual)
534  MemSemEqualReg = Call->Arguments[3];
535  if (MemOrdNeq == MemSemEqual)
536  MemSemUnequalReg = Call->Arguments[4];
537  }
538  if (!MemSemEqualReg.isValid())
539  MemSemEqualReg = buildConstantIntReg(MemSemEqual, MIRBuilder, GR);
540  if (!MemSemUnequalReg.isValid())
541  MemSemUnequalReg = buildConstantIntReg(MemSemUnequal, MIRBuilder, GR);
542 
543  Register ScopeReg;
544  auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
545  if (Call->Arguments.size() >= 6) {
546  assert(Call->Arguments.size() == 6 &&
547  "Extra args for explicit atomic cmpxchg");
548  auto ClScope = static_cast<SPIRV::CLMemoryScope>(
549  getIConstVal(Call->Arguments[5], MRI));
550  Scope = getSPIRVScope(ClScope);
551  if (ClScope == static_cast<unsigned>(Scope))
552  ScopeReg = Call->Arguments[5];
553  }
554  if (!ScopeReg.isValid())
555  ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
556 
557  Register Expected = IsCmpxchg
558  ? ExpectedArg
559  : buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,
560  GR, LLT::scalar(32));
561  MRI->setType(Expected, DesiredLLT);
562  Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT)
563  : Call->ReturnRegister;
564  GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
565 
566  SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
567  MIRBuilder.buildInstr(Opcode)
568  .addDef(Tmp)
569  .addUse(GR->getSPIRVTypeID(IntTy))
570  .addUse(ObjectPtr)
571  .addUse(ScopeReg)
572  .addUse(MemSemEqualReg)
573  .addUse(MemSemUnequalReg)
574  .addUse(Desired)
575  .addUse(Expected);
576  if (!IsCmpxchg) {
577  MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);
578  MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected);
579  }
580  return true;
581 }
582 
583 /// Helper function for building an atomic load instruction.
584 static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
585  MachineIRBuilder &MIRBuilder,
586  SPIRVGlobalRegistry *GR) {
587  const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
588  Register ScopeRegister;
589  SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
590  if (Call->Arguments.size() >= 4) {
591  assert(Call->Arguments.size() == 4 && "Extra args for explicit atomic RMW");
592  auto CLScope = static_cast<SPIRV::CLMemoryScope>(
593  getIConstVal(Call->Arguments[3], MRI));
594  Scope = getSPIRVScope(CLScope);
595  if (CLScope == static_cast<unsigned>(Scope))
596  ScopeRegister = Call->Arguments[3];
597  }
598  if (!ScopeRegister.isValid())
599  ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
600 
601  Register PtrRegister = Call->Arguments[0];
602  Register MemSemanticsReg;
603  unsigned Semantics = SPIRV::MemorySemantics::None;
604  if (Call->Arguments.size() >= 3) {
605  std::memory_order Order =
606  static_cast<std::memory_order>(getIConstVal(Call->Arguments[2], MRI));
607  Semantics =
608  getSPIRVMemSemantics(Order) |
610  if (Order == Semantics)
611  MemSemanticsReg = Call->Arguments[2];
612  }
613  if (!MemSemanticsReg.isValid())
614  MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
615 
616  MIRBuilder.buildInstr(Opcode)
617  .addDef(Call->ReturnRegister)
618  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
619  .addUse(PtrRegister)
620  .addUse(ScopeRegister)
621  .addUse(MemSemanticsReg)
622  .addUse(Call->Arguments[1]);
623  return true;
624 }
625 
626 /// Helper function for building barriers, i.e., memory/control ordering
627 /// operations.
628 static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
629  MachineIRBuilder &MIRBuilder,
630  SPIRVGlobalRegistry *GR) {
631  const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
632  unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI);
633  unsigned MemSemantics = SPIRV::MemorySemantics::None;
634 
635  if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
636  MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
637 
638  if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
639  MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;
640 
641  if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)
642  MemSemantics |= SPIRV::MemorySemantics::ImageMemory;
643 
644  if (Opcode == SPIRV::OpMemoryBarrier) {
645  std::memory_order MemOrder =
646  static_cast<std::memory_order>(getIConstVal(Call->Arguments[1], MRI));
647  MemSemantics = getSPIRVMemSemantics(MemOrder) | MemSemantics;
648  } else {
649  MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
650  }
651 
652  Register MemSemanticsReg;
653  if (MemFlags == MemSemantics)
654  MemSemanticsReg = Call->Arguments[0];
655  else
656  MemSemanticsReg = buildConstantIntReg(MemSemantics, MIRBuilder, GR);
657 
658  Register ScopeReg;
659  SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
660  SPIRV::Scope::Scope MemScope = Scope;
661  if (Call->Arguments.size() >= 2) {
662  assert(
663  ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
664  (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
665  "Extra args for explicitly scoped barrier");
666  Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]
667  : Call->Arguments[1];
668  SPIRV::CLMemoryScope CLScope =
669  static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg, MRI));
670  MemScope = getSPIRVScope(CLScope);
671  if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
672  (Opcode == SPIRV::OpMemoryBarrier))
673  Scope = MemScope;
674 
675  if (CLScope == static_cast<unsigned>(Scope))
676  ScopeReg = Call->Arguments[1];
677  }
678 
679  if (!ScopeReg.isValid())
680  ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
681 
682  auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);
683  if (Opcode != SPIRV::OpMemoryBarrier)
684  MIB.addUse(buildConstantIntReg(MemScope, MIRBuilder, GR));
685  MIB.addUse(MemSemanticsReg);
686  return true;
687 }
688 
689 static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
690  switch (dim) {
691  case SPIRV::Dim::DIM_1D:
692  case SPIRV::Dim::DIM_Buffer:
693  return 1;
694  case SPIRV::Dim::DIM_2D:
695  case SPIRV::Dim::DIM_Cube:
696  case SPIRV::Dim::DIM_Rect:
697  return 2;
698  case SPIRV::Dim::DIM_3D:
699  return 3;
700  default:
701  llvm_unreachable("Cannot get num components for given Dim");
702  }
703 }
704 
705 /// Helper function for obtaining the number of size components.
706 static unsigned getNumSizeComponents(SPIRVType *imgType) {
707  assert(imgType->getOpcode() == SPIRV::OpTypeImage);
708  auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm());
709  unsigned numComps = getNumComponentsForDim(dim);
710  bool arrayed = imgType->getOperand(4).getImm() == 1;
711  return arrayed ? numComps + 1 : numComps;
712 }
713 
714 //===----------------------------------------------------------------------===//
715 // Implementation functions for each builtin group
716 //===----------------------------------------------------------------------===//
717 
718 static bool generateExtInst(const SPIRV::IncomingCall *Call,
719  MachineIRBuilder &MIRBuilder,
720  SPIRVGlobalRegistry *GR) {
721  // Lookup the extended instruction number in the TableGen records.
722  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
723  uint32_t Number =
724  SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;
725 
726  // Build extended instruction.
727  auto MIB =
728  MIRBuilder.buildInstr(SPIRV::OpExtInst)
729  .addDef(Call->ReturnRegister)
730  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
731  .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
732  .addImm(Number);
733 
734  for (auto Argument : Call->Arguments)
735  MIB.addUse(Argument);
736  return true;
737 }
738 
740  MachineIRBuilder &MIRBuilder,
741  SPIRVGlobalRegistry *GR) {
742  // Lookup the instruction opcode in the TableGen records.
743  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
744  unsigned Opcode =
745  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
746 
747  Register CompareRegister;
748  SPIRVType *RelationType;
749  std::tie(CompareRegister, RelationType) =
750  buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
751 
752  // Build relational instruction.
753  auto MIB = MIRBuilder.buildInstr(Opcode)
754  .addDef(CompareRegister)
755  .addUse(GR->getSPIRVTypeID(RelationType));
756 
757  for (auto Argument : Call->Arguments)
758  MIB.addUse(Argument);
759 
760  // Build select instruction.
761  return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,
762  Call->ReturnType, GR);
763 }
764 
765 static bool generateGroupInst(const SPIRV::IncomingCall *Call,
766  MachineIRBuilder &MIRBuilder,
767  SPIRVGlobalRegistry *GR) {
768  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
769  const SPIRV::GroupBuiltin *GroupBuiltin =
770  SPIRV::lookupGroupBuiltin(Builtin->Name);
771  const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
772  Register Arg0;
773  if (GroupBuiltin->HasBoolArg) {
774  Register ConstRegister = Call->Arguments[0];
775  auto ArgInstruction = getDefInstrMaybeConstant(ConstRegister, MRI);
776  // TODO: support non-constant bool values.
777  assert(ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT &&
778  "Only constant bool value args are supported");
779  if (GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() !=
780  SPIRV::OpTypeBool)
781  Arg0 = GR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder,
782  GR->getOrCreateSPIRVBoolType(MIRBuilder));
783  }
784 
785  Register GroupResultRegister = Call->ReturnRegister;
786  SPIRVType *GroupResultType = Call->ReturnType;
787 
788  // TODO: maybe we need to check whether the result type is already boolean
789  // and in this case do not insert select instruction.
790  const bool HasBoolReturnTy =
791  GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
792  GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
793  GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
794 
795  if (HasBoolReturnTy)
796  std::tie(GroupResultRegister, GroupResultType) =
797  buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
798 
799  auto Scope = Builtin->Name.startswith("sub_group") ? SPIRV::Scope::Subgroup
800  : SPIRV::Scope::Workgroup;
801  Register ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
802 
803  // Build work/sub group instruction.
804  auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
805  .addDef(GroupResultRegister)
806  .addUse(GR->getSPIRVTypeID(GroupResultType))
807  .addUse(ScopeRegister);
808 
809  if (!GroupBuiltin->NoGroupOperation)
810  MIB.addImm(GroupBuiltin->GroupOperation);
811  if (Call->Arguments.size() > 0) {
812  MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);
813  for (unsigned i = 1; i < Call->Arguments.size(); i++)
814  MIB.addUse(Call->Arguments[i]);
815  }
816 
817  // Build select instruction.
818  if (HasBoolReturnTy)
819  buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
820  Call->ReturnType, GR);
821  return true;
822 }
823 
824 // These queries ask for a single size_t result for a given dimension index, e.g
825 // size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
826 // these values are all vec3 types, so we need to extract the correct index or
827 // return defaultVal (0 or 1 depending on the query). We also handle extending
828 // or tuncating in case size_t does not match the expected result type's
829 // bitwidth.
830 //
831 // For a constant index >= 3 we generate:
832 // %res = OpConstant %SizeT 0
833 //
834 // For other indices we generate:
835 // %g = OpVariable %ptr_V3_SizeT Input
836 // OpDecorate %g BuiltIn XXX
837 // OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"
838 // OpDecorate %g Constant
839 // %loadedVec = OpLoad %V3_SizeT %g
840 //
841 // Then, if the index is constant < 3, we generate:
842 // %res = OpCompositeExtract %SizeT %loadedVec idx
843 // If the index is dynamic, we generate:
844 // %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
845 // %cmp = OpULessThan %bool %idx %const_3
846 // %res = OpSelect %SizeT %cmp %tmp %const_0
847 //
848 // If the bitwidth of %res does not match the expected return type, we add an
849 // extend or truncate.
850 static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call,
851  MachineIRBuilder &MIRBuilder,
853  SPIRV::BuiltIn::BuiltIn BuiltinValue,
854  uint64_t DefaultValue) {
855  Register IndexRegister = Call->Arguments[0];
856  const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();
857  const unsigned PointerSize = GR->getPointerSize();
858  const SPIRVType *PointerSizeType =
859  GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
860  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
861  auto IndexInstruction = getDefInstrMaybeConstant(IndexRegister, MRI);
862 
863  // Set up the final register to do truncation or extension on at the end.
864  Register ToTruncate = Call->ReturnRegister;
865 
866  // If the index is constant, we can statically determine if it is in range.
867  bool IsConstantIndex =
868  IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
869 
870  // If it's out of range (max dimension is 3), we can just return the constant
871  // default value (0 or 1 depending on which query function).
872  if (IsConstantIndex && getIConstVal(IndexRegister, MRI) >= 3) {
873  Register defaultReg = Call->ReturnRegister;
874  if (PointerSize != ResultWidth) {
875  defaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
876  GR->assignSPIRVTypeToVReg(PointerSizeType, defaultReg,
877  MIRBuilder.getMF());
878  ToTruncate = defaultReg;
879  }
880  auto NewRegister =
881  GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
882  MIRBuilder.buildCopy(defaultReg, NewRegister);
883  } else { // If it could be in range, we need to load from the given builtin.
884  auto Vec3Ty =
885  GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);
886  Register LoadedVector =
887  buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,
889  // Set up the vreg to extract the result to (possibly a new temporary one).
890  Register Extracted = Call->ReturnRegister;
891  if (!IsConstantIndex || PointerSize != ResultWidth) {
892  Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
893  GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());
894  }
895  // Use Intrinsic::spv_extractelt so dynamic vs static extraction is
896  // handled later: extr = spv_extractelt LoadedVector, IndexRegister.
897  MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
898  Intrinsic::spv_extractelt, ArrayRef<Register>{Extracted}, true);
899  ExtractInst.addUse(LoadedVector).addUse(IndexRegister);
900 
901  // If the index is dynamic, need check if it's < 3, and then use a select.
902  if (!IsConstantIndex) {
903  insertAssignInstr(Extracted, nullptr, PointerSizeType, GR, MIRBuilder,
904  *MRI);
905 
906  auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);
907  auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
908 
909  Register CompareRegister =
910  MRI->createGenericVirtualRegister(LLT::scalar(1));
911  GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
912 
913  // Use G_ICMP to check if idxVReg < 3.
914  MIRBuilder.buildICmp(CmpInst::ICMP_ULT, CompareRegister, IndexRegister,
915  GR->buildConstantInt(3, MIRBuilder, IndexType));
916 
917  // Get constant for the default value (0 or 1 depending on which
918  // function).
919  Register DefaultRegister =
920  GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
921 
922  // Get a register for the selection result (possibly a new temporary one).
923  Register SelectionResult = Call->ReturnRegister;
924  if (PointerSize != ResultWidth) {
925  SelectionResult =
926  MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
927  GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
928  MIRBuilder.getMF());
929  }
930  // Create the final G_SELECT to return the extracted value or the default.
931  MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,
932  DefaultRegister);
933  ToTruncate = SelectionResult;
934  } else {
935  ToTruncate = Extracted;
936  }
937  }
938  // Alter the result's bitwidth if it does not match the SizeT value extracted.
939  if (PointerSize != ResultWidth)
940  MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);
941  return true;
942 }
943 
944 static bool generateBuiltinVar(const SPIRV::IncomingCall *Call,
945  MachineIRBuilder &MIRBuilder,
946  SPIRVGlobalRegistry *GR) {
947  // Lookup the builtin variable record.
948  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
949  SPIRV::BuiltIn::BuiltIn Value =
950  SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
951 
952  if (Value == SPIRV::BuiltIn::GlobalInvocationId)
953  return genWorkgroupQuery(Call, MIRBuilder, GR, Value, 0);
954 
955  // Build a load instruction for the builtin variable.
956  unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);
957  LLT LLType;
958  if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
959  LLType =
960  LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth);
961  else
962  LLType = LLT::scalar(BitWidth);
963 
964  return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value,
965  LLType, Call->ReturnRegister);
966 }
967 
968 static bool generateAtomicInst(const SPIRV::IncomingCall *Call,
969  MachineIRBuilder &MIRBuilder,
970  SPIRVGlobalRegistry *GR) {
971  // Lookup the instruction opcode in the TableGen records.
972  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
973  unsigned Opcode =
974  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
975 
976  switch (Opcode) {
977  case SPIRV::OpStore:
978  return buildAtomicInitInst(Call, MIRBuilder);
979  case SPIRV::OpAtomicLoad:
980  return buildAtomicLoadInst(Call, MIRBuilder, GR);
981  case SPIRV::OpAtomicStore:
982  return buildAtomicStoreInst(Call, MIRBuilder, GR);
983  case SPIRV::OpAtomicCompareExchange:
984  case SPIRV::OpAtomicCompareExchangeWeak:
985  return buildAtomicCompareExchangeInst(Call, MIRBuilder, GR);
986  case SPIRV::OpAtomicIAdd:
987  case SPIRV::OpAtomicISub:
988  case SPIRV::OpAtomicOr:
989  case SPIRV::OpAtomicXor:
990  case SPIRV::OpAtomicAnd:
991  case SPIRV::OpAtomicExchange:
992  return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
993  case SPIRV::OpMemoryBarrier:
994  return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
995  default:
996  return false;
997  }
998 }
999 
1001  MachineIRBuilder &MIRBuilder,
1002  SPIRVGlobalRegistry *GR) {
1003  // Lookup the instruction opcode in the TableGen records.
1004  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1005  unsigned Opcode =
1006  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1007 
1008  return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
1009 }
1010 
1012  MachineIRBuilder &MIRBuilder,
1013  SPIRVGlobalRegistry *GR) {
1014  unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();
1015  bool IsVec = Opcode == SPIRV::OpTypeVector;
1016  // Use OpDot only in case of vector args and OpFMul in case of scalar args.
1017  MIRBuilder.buildInstr(IsVec ? SPIRV::OpDot : SPIRV::OpFMulS)
1018  .addDef(Call->ReturnRegister)
1019  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1020  .addUse(Call->Arguments[0])
1021  .addUse(Call->Arguments[1]);
1022  return true;
1023 }
1024 
1026  MachineIRBuilder &MIRBuilder,
1027  SPIRVGlobalRegistry *GR) {
1028  // Lookup the builtin record.
1029  SPIRV::BuiltIn::BuiltIn Value =
1030  SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;
1031  uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||
1032  Value == SPIRV::BuiltIn::WorkgroupSize ||
1033  Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
1034  return genWorkgroupQuery(Call, MIRBuilder, GR, Value, IsDefault ? 1 : 0);
1035 }
1036 
1038  MachineIRBuilder &MIRBuilder,
1039  SPIRVGlobalRegistry *GR) {
1040  // Lookup the image size query component number in the TableGen records.
1041  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1043  SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;
1044  // Query result may either be a vector or a scalar. If return type is not a
1045  // vector, expect only a single size component. Otherwise get the number of
1046  // expected components.
1047  SPIRVType *RetTy = Call->ReturnType;
1048  unsigned NumExpectedRetComponents = RetTy->getOpcode() == SPIRV::OpTypeVector
1049  ? RetTy->getOperand(2).getImm()
1050  : 1;
1051  // Get the actual number of query result/size components.
1052  SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1053  unsigned NumActualRetComponents = getNumSizeComponents(ImgType);
1054  Register QueryResult = Call->ReturnRegister;
1055  SPIRVType *QueryResultType = Call->ReturnType;
1056  if (NumExpectedRetComponents != NumActualRetComponents) {
1057  QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
1058  LLT::fixed_vector(NumActualRetComponents, 32));
1059  SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1060  QueryResultType = GR->getOrCreateSPIRVVectorType(
1061  IntTy, NumActualRetComponents, MIRBuilder);
1062  GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());
1063  }
1064  bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;
1065  unsigned Opcode =
1066  IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
1067  auto MIB = MIRBuilder.buildInstr(Opcode)
1068  .addDef(QueryResult)
1069  .addUse(GR->getSPIRVTypeID(QueryResultType))
1070  .addUse(Call->Arguments[0]);
1071  if (!IsDimBuf)
1072  MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Lod id.
1073  if (NumExpectedRetComponents == NumActualRetComponents)
1074  return true;
1075  if (NumExpectedRetComponents == 1) {
1076  // Only 1 component is expected, build OpCompositeExtract instruction.
1077  unsigned ExtractedComposite =
1078  Component == 3 ? NumActualRetComponents - 1 : Component;
1079  assert(ExtractedComposite < NumActualRetComponents &&
1080  "Invalid composite index!");
1081  MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1082  .addDef(Call->ReturnRegister)
1083  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1084  .addUse(QueryResult)
1085  .addImm(ExtractedComposite);
1086  } else {
1087  // More than 1 component is expected, fill a new vector.
1088  auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)
1089  .addDef(Call->ReturnRegister)
1090  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1091  .addUse(QueryResult)
1092  .addUse(QueryResult);
1093  for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
1094  MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);
1095  }
1096  return true;
1097 }
1098 
1100  MachineIRBuilder &MIRBuilder,
1101  SPIRVGlobalRegistry *GR) {
1102  assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
1103  "Image samples query result must be of int type!");
1104 
1105  // Lookup the instruction opcode in the TableGen records.
1106  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1107  unsigned Opcode =
1108  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1109 
1110  Register Image = Call->Arguments[0];
1111  SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
1112  GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
1113 
1114  switch (Opcode) {
1115  case SPIRV::OpImageQuerySamples:
1116  assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
1117  "Image must be of 2D dimensionality");
1118  break;
1119  case SPIRV::OpImageQueryLevels:
1120  assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
1121  ImageDimensionality == SPIRV::Dim::DIM_2D ||
1122  ImageDimensionality == SPIRV::Dim::DIM_3D ||
1123  ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
1124  "Image must be of 1D/2D/3D/Cube dimensionality");
1125  break;
1126  }
1127 
1128  MIRBuilder.buildInstr(Opcode)
1129  .addDef(Call->ReturnRegister)
1130  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1131  .addUse(Image);
1132  return true;
1133 }
1134 
1135 // TODO: Move to TableGen.
1136 static SPIRV::SamplerAddressingMode::SamplerAddressingMode
1138  switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
1140  return SPIRV::SamplerAddressingMode::Clamp;
1142  return SPIRV::SamplerAddressingMode::ClampToEdge;
1144  return SPIRV::SamplerAddressingMode::Repeat;
1146  return SPIRV::SamplerAddressingMode::RepeatMirrored;
1148  return SPIRV::SamplerAddressingMode::None;
1149  default:
1150  llvm_unreachable("Unknown CL address mode");
1151  }
1152 }
1153 
1154 static unsigned getSamplerParamFromBitmask(unsigned Bitmask) {
1155  return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;
1156 }
1157 
1158 static SPIRV::SamplerFilterMode::SamplerFilterMode
1160  if (Bitmask & SPIRV::CLK_FILTER_LINEAR)
1161  return SPIRV::SamplerFilterMode::Linear;
1162  if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
1163  return SPIRV::SamplerFilterMode::Nearest;
1164  return SPIRV::SamplerFilterMode::Nearest;
1165 }
1166 
1167 static bool generateReadImageInst(const StringRef DemangledCall,
1168  const SPIRV::IncomingCall *Call,
1169  MachineIRBuilder &MIRBuilder,
1170  SPIRVGlobalRegistry *GR) {
1171  Register Image = Call->Arguments[0];
1172  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1173 
1174  if (DemangledCall.contains_insensitive("ocl_sampler")) {
1175  Register Sampler = Call->Arguments[1];
1176 
1177  if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
1178  getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) {
1179  uint64_t SamplerMask = getIConstVal(Sampler, MRI);
1180  Sampler = GR->buildConstantSampler(
1182  getSamplerParamFromBitmask(SamplerMask),
1183  getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder,
1184  GR->getSPIRVTypeForVReg(Sampler));
1185  }
1186  SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1187  SPIRVType *SampledImageType =
1188  GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1189  Register SampledImage = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1190 
1191  MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1192  .addDef(SampledImage)
1193  .addUse(GR->getSPIRVTypeID(SampledImageType))
1194  .addUse(Image)
1195  .addUse(Sampler);
1196 
1198  MIRBuilder);
1199  SPIRVType *TempType = Call->ReturnType;
1200  bool NeedsExtraction = false;
1201  if (TempType->getOpcode() != SPIRV::OpTypeVector) {
1202  TempType =
1203  GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);
1204  NeedsExtraction = true;
1205  }
1206  LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(TempType));
1207  Register TempRegister = MRI->createGenericVirtualRegister(LLType);
1208  GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
1209 
1210  MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1211  .addDef(NeedsExtraction ? TempRegister : Call->ReturnRegister)
1212  .addUse(GR->getSPIRVTypeID(TempType))
1213  .addUse(SampledImage)
1214  .addUse(Call->Arguments[2]) // Coordinate.
1215  .addImm(SPIRV::ImageOperand::Lod)
1216  .addUse(Lod);
1217 
1218  if (NeedsExtraction)
1219  MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1220  .addDef(Call->ReturnRegister)
1221  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1222  .addUse(TempRegister)
1223  .addImm(0);
1224  } else if (DemangledCall.contains_insensitive("msaa")) {
1225  MIRBuilder.buildInstr(SPIRV::OpImageRead)
1226  .addDef(Call->ReturnRegister)
1227  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1228  .addUse(Image)
1229  .addUse(Call->Arguments[1]) // Coordinate.
1230  .addImm(SPIRV::ImageOperand::Sample)
1231  .addUse(Call->Arguments[2]);
1232  } else {
1233  MIRBuilder.buildInstr(SPIRV::OpImageRead)
1234  .addDef(Call->ReturnRegister)
1235  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1236  .addUse(Image)
1237  .addUse(Call->Arguments[1]); // Coordinate.
1238  }
1239  return true;
1240 }
1241 
1243  MachineIRBuilder &MIRBuilder,
1244  SPIRVGlobalRegistry *GR) {
1245  MIRBuilder.buildInstr(SPIRV::OpImageWrite)
1246  .addUse(Call->Arguments[0]) // Image.
1247  .addUse(Call->Arguments[1]) // Coordinate.
1248  .addUse(Call->Arguments[2]); // Texel.
1249  return true;
1250 }
1251 
1252 static bool generateSampleImageInst(const StringRef DemangledCall,
1253  const SPIRV::IncomingCall *Call,
1254  MachineIRBuilder &MIRBuilder,
1255  SPIRVGlobalRegistry *GR) {
1256  if (Call->Builtin->Name.contains_insensitive(
1257  "__translate_sampler_initializer")) {
1258  // Build sampler literal.
1259  uint64_t Bitmask = getIConstVal(Call->Arguments[0], MIRBuilder.getMRI());
1260  Register Sampler = GR->buildConstantSampler(
1261  Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask),
1262  getSamplerParamFromBitmask(Bitmask),
1263  getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType);
1264  return Sampler.isValid();
1265  } else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {
1266  // Create OpSampledImage.
1267  Register Image = Call->Arguments[0];
1268  SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1269  SPIRVType *SampledImageType =
1270  GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1271  Register SampledImage =
1272  Call->ReturnRegister.isValid()
1273  ? Call->ReturnRegister
1274  : MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
1275  MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1276  .addDef(SampledImage)
1277  .addUse(GR->getSPIRVTypeID(SampledImageType))
1278  .addUse(Image)
1279  .addUse(Call->Arguments[1]); // Sampler.
1280  return true;
1281  } else if (Call->Builtin->Name.contains_insensitive(
1282  "__spirv_ImageSampleExplicitLod")) {
1283  // Sample an image using an explicit level of detail.
1284  std::string ReturnType = DemangledCall.str();
1285  if (DemangledCall.contains("_R")) {
1286  ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);
1287  ReturnType = ReturnType.substr(0, ReturnType.find('('));
1288  }
1290  MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1291  .addDef(Call->ReturnRegister)
1292  .addUse(GR->getSPIRVTypeID(Type))
1293  .addUse(Call->Arguments[0]) // Image.
1294  .addUse(Call->Arguments[1]) // Coordinate.
1295  .addImm(SPIRV::ImageOperand::Lod)
1296  .addUse(Call->Arguments[3]);
1297  return true;
1298  }
1299  return false;
1300 }
1301 
1302 static bool generateSelectInst(const SPIRV::IncomingCall *Call,
1303  MachineIRBuilder &MIRBuilder) {
1304  MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],
1305  Call->Arguments[1], Call->Arguments[2]);
1306  return true;
1307 }
1308 
1310  MachineIRBuilder &MIRBuilder,
1311  SPIRVGlobalRegistry *GR) {
1312  // Lookup the instruction opcode in the TableGen records.
1313  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1314  unsigned Opcode =
1315  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1316  const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1317 
1318  switch (Opcode) {
1319  case SPIRV::OpSpecConstant: {
1320  // Build the SpecID decoration.
1321  unsigned SpecId =
1322  static_cast<unsigned>(getIConstVal(Call->Arguments[0], MRI));
1323  buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId,
1324  {SpecId});
1325  // Determine the constant MI.
1326  Register ConstRegister = Call->Arguments[1];
1327  const MachineInstr *Const = getDefInstrMaybeConstant(ConstRegister, MRI);
1328  assert(Const &&
1329  (Const->getOpcode() == TargetOpcode::G_CONSTANT ||
1330  Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&
1331  "Argument should be either an int or floating-point constant");
1332  // Determine the opcode and built the OpSpec MI.
1333  const MachineOperand &ConstOperand = Const->getOperand(1);
1334  if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
1335  assert(ConstOperand.isCImm() && "Int constant operand is expected");
1336  Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
1337  ? SPIRV::OpSpecConstantTrue
1338  : SPIRV::OpSpecConstantFalse;
1339  }
1340  auto MIB = MIRBuilder.buildInstr(Opcode)
1341  .addDef(Call->ReturnRegister)
1342  .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1343 
1344  if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
1345  if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
1346  addNumImm(ConstOperand.getCImm()->getValue(), MIB);
1347  else
1348  addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
1349  }
1350  return true;
1351  }
1352  case SPIRV::OpSpecConstantComposite: {
1353  auto MIB = MIRBuilder.buildInstr(Opcode)
1354  .addDef(Call->ReturnRegister)
1355  .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1356  for (unsigned i = 0; i < Call->Arguments.size(); i++)
1357  MIB.addUse(Call->Arguments[i]);
1358  return true;
1359  }
1360  default:
1361  return false;
1362  }
1363 }
1364 
1367  // We expect the following sequence of instructions:
1368  // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
1369  // or = G_GLOBAL_VALUE @block_literal_global
1370  // %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0
1371  // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
1372  MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
1373  assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
1374  MI->getOperand(1).isReg());
1375  Register BitcastReg = MI->getOperand(1).getReg();
1376  MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg);
1377  assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&
1378  BitcastMI->getOperand(2).isReg());
1379  Register ValueReg = BitcastMI->getOperand(2).getReg();
1380  MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
1381  return ValueMI;
1382 }
1383 
1384 // Return type of the instruction result from spv_assign_type intrinsic.
1385 // TODO: maybe unify with prelegalizer pass.
1387  MachineInstr *NextMI = MI->getNextNode();
1388  if (isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_name))
1389  NextMI = NextMI->getNextNode();
1390  Register ValueReg = MI->getOperand(0).getReg();
1391  if (!isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_type) ||
1392  NextMI->getOperand(1).getReg() != ValueReg)
1393  return nullptr;
1394  Type *Ty = getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0);
1395  assert(Ty && "Type is expected");
1396  return getTypedPtrEltType(Ty);
1397 }
1398 
1399 static const Type *getBlockStructType(Register ParamReg,
1401  // In principle, this information should be passed to us from Clang via
1402  // an elementtype attribute. However, said attribute requires that
1403  // the function call be an intrinsic, which is not. Instead, we rely on being
1404  // able to trace this to the declaration of a variable: OpenCL C specification
1405  // section 6.12.5 should guarantee that we can do this.
1406  MachineInstr *MI = getBlockStructInstr(ParamReg, MRI);
1407  if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
1408  return getTypedPtrEltType(MI->getOperand(1).getGlobal()->getType());
1409  assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
1410  "Blocks in OpenCL C must be traceable to allocation site");
1411  return getMachineInstrType(MI);
1412 }
1413 
1414 // TODO: maybe move to the global register.
1415 static SPIRVType *
1417  SPIRVGlobalRegistry *GR) {
1418  LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
1419  Type *OpaqueType = StructType::getTypeByName(Context, "spirv.DeviceEvent");
1420  if (!OpaqueType)
1421  OpaqueType = StructType::getTypeByName(Context, "opencl.clk_event_t");
1422  if (!OpaqueType)
1423  OpaqueType = StructType::create(Context, "spirv.DeviceEvent");
1424  unsigned SC0 = storageClassToAddressSpace(SPIRV::StorageClass::Function);
1425  unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1426  Type *PtrType = PointerType::get(PointerType::get(OpaqueType, SC0), SC1);
1427  return GR->getOrCreateSPIRVType(PtrType, MIRBuilder);
1428 }
1429 
1430 static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
1431  MachineIRBuilder &MIRBuilder,
1432  SPIRVGlobalRegistry *GR) {
1433  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1434  const DataLayout &DL = MIRBuilder.getDataLayout();
1435  bool HasEvents = Call->Builtin->Name.find("events") != StringRef::npos;
1436  const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1437 
1438  // Make vararg instructions before OpEnqueueKernel.
1439  // Local sizes arguments: Sizes of block invoke arguments. Clang generates
1440  // local size operands as an array, so we need to unpack them.
1441  SmallVector<Register, 16> LocalSizes;
1442  if (Call->Builtin->Name.find("_varargs") != StringRef::npos) {
1443  const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
1444  Register GepReg = Call->Arguments[LocalSizeArrayIdx];
1445  MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
1446  assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
1447  GepMI->getOperand(3).isReg());
1448  Register ArrayReg = GepMI->getOperand(3).getReg();
1449  MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
1450  const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
1451  assert(LocalSizeTy && "Local size type is expected");
1452  const uint64_t LocalSizeNum =
1453  cast<ArrayType>(LocalSizeTy)->getNumElements();
1454  unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1455  const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
1456  const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(
1457  Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
1458  for (unsigned I = 0; I < LocalSizeNum; ++I) {
1459  Register Reg =
1460  MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
1461  MIRBuilder.getMRI()->setType(Reg, LLType);
1462  GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
1463  auto GEPInst = MIRBuilder.buildIntrinsic(Intrinsic::spv_gep,
1464  ArrayRef<Register>{Reg}, true);
1465  GEPInst
1466  .addImm(GepMI->getOperand(2).getImm()) // In bound.
1467  .addUse(ArrayMI->getOperand(0).getReg()) // Alloca.
1468  .addUse(buildConstantIntReg(0, MIRBuilder, GR)) // Indices.
1469  .addUse(buildConstantIntReg(I, MIRBuilder, GR));
1470  LocalSizes.push_back(Reg);
1471  }
1472  }
1473 
1474  // SPIRV OpEnqueueKernel instruction has 10+ arguments.
1475  auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
1476  .addDef(Call->ReturnRegister)
1477  .addUse(GR->getSPIRVTypeID(Int32Ty));
1478 
1479  // Copy all arguments before block invoke function pointer.
1480  const unsigned BlockFIdx = HasEvents ? 6 : 3;
1481  for (unsigned i = 0; i < BlockFIdx; i++)
1482  MIB.addUse(Call->Arguments[i]);
1483 
1484  // If there are no event arguments in the original call, add dummy ones.
1485  if (!HasEvents) {
1486  MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Dummy num events.
1487  Register NullPtr = GR->getOrCreateConstNullPtr(
1488  MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
1489  MIB.addUse(NullPtr); // Dummy wait events.
1490  MIB.addUse(NullPtr); // Dummy ret event.
1491  }
1492 
1493  MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
1494  assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
1495  // Invoke: Pointer to invoke function.
1496  MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
1497 
1498  Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
1499  // Param: Pointer to block literal.
1500  MIB.addUse(BlockLiteralReg);
1501 
1502  Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
1503  // TODO: these numbers should be obtained from block literal structure.
1504  // Param Size: Size of block literal structure.
1505  MIB.addUse(buildConstantIntReg(DL.getTypeStoreSize(PType), MIRBuilder, GR));
1506  // Param Aligment: Aligment of block literal structure.
1507  MIB.addUse(
1508  buildConstantIntReg(DL.getPrefTypeAlignment(PType), MIRBuilder, GR));
1509 
1510  for (unsigned i = 0; i < LocalSizes.size(); i++)
1511  MIB.addUse(LocalSizes[i]);
1512  return true;
1513 }
1514 
1516  MachineIRBuilder &MIRBuilder,
1517  SPIRVGlobalRegistry *GR) {
1518  // Lookup the instruction opcode in the TableGen records.
1519  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1520  unsigned Opcode =
1521  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1522 
1523  switch (Opcode) {
1524  case SPIRV::OpRetainEvent:
1525  case SPIRV::OpReleaseEvent:
1526  return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);
1527  case SPIRV::OpCreateUserEvent:
1528  case SPIRV::OpGetDefaultQueue:
1529  return MIRBuilder.buildInstr(Opcode)
1530  .addDef(Call->ReturnRegister)
1531  .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1532  case SPIRV::OpIsValidEvent:
1533  return MIRBuilder.buildInstr(Opcode)
1534  .addDef(Call->ReturnRegister)
1535  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1536  .addUse(Call->Arguments[0]);
1537  case SPIRV::OpSetUserEventStatus:
1538  return MIRBuilder.buildInstr(Opcode)
1539  .addUse(Call->Arguments[0])
1540  .addUse(Call->Arguments[1]);
1541  case SPIRV::OpCaptureEventProfilingInfo:
1542  return MIRBuilder.buildInstr(Opcode)
1543  .addUse(Call->Arguments[0])
1544  .addUse(Call->Arguments[1])
1545  .addUse(Call->Arguments[2]);
1546  case SPIRV::OpBuildNDRange: {
1547  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1548  SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1549  assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&
1550  PtrType->getOperand(2).isReg());
1551  Register TypeReg = PtrType->getOperand(2).getReg();
1552  SPIRVType *StructType = GR->getSPIRVTypeForVReg(TypeReg);
1553  Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1554  GR->assignSPIRVTypeToVReg(StructType, TmpReg, MIRBuilder.getMF());
1555  // Skip the first arg, it's the destination pointer. OpBuildNDRange takes
1556  // three other arguments, so pass zero constant on absence.
1557  unsigned NumArgs = Call->Arguments.size();
1558  assert(NumArgs >= 2);
1559  Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];
1560  Register LocalWorkSize =
1561  NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];
1562  Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1];
1563  if (NumArgs < 4) {
1564  Register Const;
1565  SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);
1566  if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {
1567  MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize);
1568  assert(DefInstr && isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&
1569  DefInstr->getOperand(3).isReg());
1570  Register GWSPtr = DefInstr->getOperand(3).getReg();
1571  // TODO: Maybe simplify generation of the type of the fields.
1572  unsigned Size = Call->Builtin->Name.equals("ndrange_3D") ? 3 : 2;
1573  unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32;
1574  Type *BaseTy = IntegerType::get(
1575  MIRBuilder.getMF().getFunction().getContext(), BitWidth);
1576  Type *FieldTy = ArrayType::get(BaseTy, Size);
1577  SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder);
1578  GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1579  GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize,
1580  MIRBuilder.getMF());
1581  MIRBuilder.buildInstr(SPIRV::OpLoad)
1582  .addDef(GlobalWorkSize)
1583  .addUse(GR->getSPIRVTypeID(SpvFieldTy))
1584  .addUse(GWSPtr);
1585  Const = GR->getOrCreateConsIntArray(0, MIRBuilder, SpvFieldTy);
1586  } else {
1587  Const = GR->buildConstantInt(0, MIRBuilder, SpvTy);
1588  }
1589  if (!LocalWorkSize.isValid())
1590  LocalWorkSize = Const;
1591  if (!GlobalWorkOffset.isValid())
1592  GlobalWorkOffset = Const;
1593  }
1594  MIRBuilder.buildInstr(Opcode)
1595  .addDef(TmpReg)
1596  .addUse(TypeReg)
1597  .addUse(GlobalWorkSize)
1598  .addUse(LocalWorkSize)
1599  .addUse(GlobalWorkOffset);
1600  return MIRBuilder.buildInstr(SPIRV::OpStore)
1601  .addUse(Call->Arguments[0])
1602  .addUse(TmpReg);
1603  }
1604  case SPIRV::OpEnqueueKernel:
1605  return buildEnqueueKernel(Call, MIRBuilder, GR);
1606  default:
1607  return false;
1608  }
1609 }
1610 
1611 static bool generateAsyncCopy(const SPIRV::IncomingCall *Call,
1612  MachineIRBuilder &MIRBuilder,
1613  SPIRVGlobalRegistry *GR) {
1614  // Lookup the instruction opcode in the TableGen records.
1615  const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1616  unsigned Opcode =
1617  SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1618  auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GR);
1619 
1620  switch (Opcode) {
1621  case SPIRV::OpGroupAsyncCopy:
1622  return MIRBuilder.buildInstr(Opcode)
1623  .addDef(Call->ReturnRegister)
1624  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1625  .addUse(Scope)
1626  .addUse(Call->Arguments[0])
1627  .addUse(Call->Arguments[1])
1628  .addUse(Call->Arguments[2])
1629  .addUse(buildConstantIntReg(1, MIRBuilder, GR))
1630  .addUse(Call->Arguments[3]);
1631  case SPIRV::OpGroupWaitEvents:
1632  return MIRBuilder.buildInstr(Opcode)
1633  .addUse(Scope)
1634  .addUse(Call->Arguments[0])
1635  .addUse(Call->Arguments[1]);
1636  default:
1637  return false;
1638  }
1639 }
1640 
1641 static bool generateConvertInst(const StringRef DemangledCall,
1642  const SPIRV::IncomingCall *Call,
1643  MachineIRBuilder &MIRBuilder,
1644  SPIRVGlobalRegistry *GR) {
1645  // Lookup the conversion builtin in the TableGen records.
1646  const SPIRV::ConvertBuiltin *Builtin =
1647  SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);
1648 
1649  if (Builtin->IsSaturated)
1650  buildOpDecorate(Call->ReturnRegister, MIRBuilder,
1651  SPIRV::Decoration::SaturatedConversion, {});
1652  if (Builtin->IsRounded)
1653  buildOpDecorate(Call->ReturnRegister, MIRBuilder,
1654  SPIRV::Decoration::FPRoundingMode, {Builtin->RoundingMode});
1655 
1656  unsigned Opcode = SPIRV::OpNop;
1657  if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
1658  // Int -> ...
1659  if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
1660  // Int -> Int
1661  if (Builtin->IsSaturated)
1662  Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
1663  : SPIRV::OpSatConvertSToU;
1664  else
1665  Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert
1666  : SPIRV::OpSConvert;
1667  } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
1668  SPIRV::OpTypeFloat)) {
1669  // Int -> Float
1670  bool IsSourceSigned =
1671  DemangledCall[DemangledCall.find_first_of('(') + 1] != 'u';
1672  Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
1673  }
1674  } else if (GR->isScalarOrVectorOfType(Call->Arguments[0],
1675  SPIRV::OpTypeFloat)) {
1676  // Float -> ...
1677  if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt))
1678  // Float -> Int
1679  Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
1680  : SPIRV::OpConvertFToU;
1681  else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
1682  SPIRV::OpTypeFloat))
1683  // Float -> Float
1684  Opcode = SPIRV::OpFConvert;
1685  }
1686 
1687  assert(Opcode != SPIRV::OpNop &&
1688  "Conversion between the types not implemented!");
1689 
1690  MIRBuilder.buildInstr(Opcode)
1691  .addDef(Call->ReturnRegister)
1692  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1693  .addUse(Call->Arguments[0]);
1694  return true;
1695 }
1696 
1698  MachineIRBuilder &MIRBuilder,
1699  SPIRVGlobalRegistry *GR) {
1700  // Lookup the vector load/store builtin in the TableGen records.
1701  const SPIRV::VectorLoadStoreBuiltin *Builtin =
1702  SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
1703  Call->Builtin->Set);
1704  // Build extended instruction.
1705  auto MIB =
1706  MIRBuilder.buildInstr(SPIRV::OpExtInst)
1707  .addDef(Call->ReturnRegister)
1708  .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1709  .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
1710  .addImm(Builtin->Number);
1711  for (auto Argument : Call->Arguments)
1712  MIB.addUse(Argument);
1713 
1714  // Rounding mode should be passed as a last argument in the MI for builtins
1715  // like "vstorea_halfn_r".
1716  if (Builtin->IsRounded)
1717  MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));
1718  return true;
1719 }
1720 
1721 /// Lowers a builtin funtion call using the provided \p DemangledCall skeleton
1722 /// and external instruction \p Set.
1723 namespace SPIRV {
1725  SPIRV::InstructionSet::InstructionSet Set,
1726  MachineIRBuilder &MIRBuilder,
1727  const Register OrigRet, const Type *OrigRetTy,
1729  SPIRVGlobalRegistry *GR) {
1730  LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
1731 
1732  // SPIR-V type and return register.
1733  Register ReturnRegister = OrigRet;
1734  SPIRVType *ReturnType = nullptr;
1735  if (OrigRetTy && !OrigRetTy->isVoidTy()) {
1736  ReturnType = GR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder);
1737  } else if (OrigRetTy && OrigRetTy->isVoidTy()) {
1738  ReturnRegister = MIRBuilder.getMRI()->createVirtualRegister(&IDRegClass);
1739  MIRBuilder.getMRI()->setType(ReturnRegister, LLT::scalar(32));
1740  ReturnType = GR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder);
1741  }
1742 
1743  // Lookup the builtin in the TableGen records.
1744  std::unique_ptr<const IncomingCall> Call =
1745  lookupBuiltin(DemangledCall, Set, ReturnRegister, ReturnType, Args);
1746 
1747  if (!Call) {
1748  LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
1749  return None;
1750  }
1751 
1752  // TODO: check if the provided args meet the builtin requirments.
1753  assert(Args.size() >= Call->Builtin->MinNumArgs &&
1754  "Too few arguments to generate the builtin");
1755  if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)
1756  LLVM_DEBUG(dbgs() << "More arguments provided than required!\n");
1757 
1758  // Match the builtin with implementation based on the grouping.
1759  switch (Call->Builtin->Group) {
1760  case SPIRV::Extended:
1761  return generateExtInst(Call.get(), MIRBuilder, GR);
1762  case SPIRV::Relational:
1763  return generateRelationalInst(Call.get(), MIRBuilder, GR);
1764  case SPIRV::Group:
1765  return generateGroupInst(Call.get(), MIRBuilder, GR);
1766  case SPIRV::Variable:
1767  return generateBuiltinVar(Call.get(), MIRBuilder, GR);
1768  case SPIRV::Atomic:
1769  return generateAtomicInst(Call.get(), MIRBuilder, GR);
1770  case SPIRV::Barrier:
1771  return generateBarrierInst(Call.get(), MIRBuilder, GR);
1772  case SPIRV::Dot:
1773  return generateDotOrFMulInst(Call.get(), MIRBuilder, GR);
1774  case SPIRV::GetQuery:
1775  return generateGetQueryInst(Call.get(), MIRBuilder, GR);
1776  case SPIRV::ImageSizeQuery:
1777  return generateImageSizeQueryInst(Call.get(), MIRBuilder, GR);
1778  case SPIRV::ImageMiscQuery:
1779  return generateImageMiscQueryInst(Call.get(), MIRBuilder, GR);
1780  case SPIRV::ReadImage:
1781  return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
1782  case SPIRV::WriteImage:
1783  return generateWriteImageInst(Call.get(), MIRBuilder, GR);
1784  case SPIRV::SampleImage:
1785  return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
1786  case SPIRV::Select:
1787  return generateSelectInst(Call.get(), MIRBuilder);
1788  case SPIRV::SpecConstant:
1789  return generateSpecConstantInst(Call.get(), MIRBuilder, GR);
1790  case SPIRV::Enqueue:
1791  return generateEnqueueInst(Call.get(), MIRBuilder, GR);
1792  case SPIRV::AsyncCopy:
1793  return generateAsyncCopy(Call.get(), MIRBuilder, GR);
1794  case SPIRV::Convert:
1795  return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);
1796  case SPIRV::VectorLoadStore:
1797  return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);
1798  }
1799  return false;
1800 }
1801 
1805 };
1806 
1807 #define GET_DemangledTypes_DECL
1808 #define GET_DemangledTypes_IMPL
1809 
1810 struct ImageType {
1814  Dim::Dim Dimensionality;
1815  bool Arrayed;
1816  bool Depth;
1818  bool Sampled;
1819  ImageFormat::ImageFormat Format;
1820 };
1821 
1822 struct PipeType {
1825 };
1826 
1827 using namespace AccessQualifier;
1828 using namespace Dim;
1829 using namespace ImageFormat;
1830 #define GET_ImageTypes_DECL
1831 #define GET_ImageTypes_IMPL
1832 #define GET_PipeTypes_DECL
1833 #define GET_PipeTypes_IMPL
1834 #include "SPIRVGenTables.inc"
1835 } // namespace SPIRV
1836 
1837 //===----------------------------------------------------------------------===//
1838 // Misc functions for parsing builtin types and looking up implementation
1839 // details in TableGenerated tables.
1840 //===----------------------------------------------------------------------===//
1841 
1843  if (Name.startswith("opencl."))
1844  return SPIRV::lookupBuiltinType(Name);
1845  if (!Name.startswith("spirv."))
1846  return nullptr;
1847  // Some SPIR-V builtin types have a complex list of parameters as part of
1848  // their name (e.g. spirv.Image._void_1_0_0_0_0_0_0). Those parameters often
1849  // are numeric literals which cannot be easily represented by TableGen
1850  // records and should be parsed instead.
1851  unsigned BaseTypeNameLength =
1852  Name.contains('_') ? Name.find('_') - 1 : Name.size();
1853  return SPIRV::lookupBuiltinType(Name.substr(0, BaseTypeNameLength).str());
1854 }
1855 
1856 static std::unique_ptr<const SPIRV::ImageType>
1858  if (Name.startswith("opencl.")) {
1859  // Lookup OpenCL builtin image type lowering details in TableGen records.
1860  const SPIRV::ImageType *Record = SPIRV::lookupImageType(Name);
1861  return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType(*Record));
1862  }
1863  if (!Name.startswith("spirv."))
1864  llvm_unreachable("Unknown builtin image type name/literal");
1865  // Parse the literals of SPIR-V image builtin parameters. The name should
1866  // have the following format:
1867  // spirv.Image._Type_Dim_Depth_Arrayed_MS_Sampled_ImageFormat_AccessQualifier
1868  // e.g. %spirv.Image._void_1_0_0_0_0_0_0
1869  StringRef TypeParametersString = Name.substr(strlen("spirv.Image."));
1870  SmallVector<StringRef> TypeParameters;
1871  SplitString(TypeParametersString, TypeParameters, "_");
1872  assert(TypeParameters.size() == 8 &&
1873  "Wrong number of literals in SPIR-V builtin image type");
1874 
1875  StringRef SampledType = TypeParameters[0];
1876  unsigned Dim, Depth, Arrayed, Multisampled, Sampled, Format, AccessQual;
1877  bool AreParameterLiteralsValid =
1878  !(TypeParameters[1].getAsInteger(10, Dim) ||
1879  TypeParameters[2].getAsInteger(10, Depth) ||
1880  TypeParameters[3].getAsInteger(10, Arrayed) ||
1881  TypeParameters[4].getAsInteger(10, Multisampled) ||
1882  TypeParameters[5].getAsInteger(10, Sampled) ||
1883  TypeParameters[6].getAsInteger(10, Format) ||
1884  TypeParameters[7].getAsInteger(10, AccessQual));
1885  assert(AreParameterLiteralsValid &&
1886  "Invalid format of SPIR-V image type parameter literals.");
1887 
1888  return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType{
1889  Name, SampledType, SPIRV::AccessQualifier::AccessQualifier(AccessQual),
1890  SPIRV::Dim::Dim(Dim), static_cast<bool>(Arrayed),
1891  static_cast<bool>(Depth), static_cast<bool>(Multisampled),
1892  static_cast<bool>(Sampled), SPIRV::ImageFormat::ImageFormat(Format)});
1893 }
1894 
1895 static std::unique_ptr<const SPIRV::PipeType>
1897  if (Name.startswith("opencl.")) {
1898  // Lookup OpenCL builtin pipe type lowering details in TableGen records.
1899  const SPIRV::PipeType *Record = SPIRV::lookupPipeType(Name);
1900  return std::unique_ptr<SPIRV::PipeType>(new SPIRV::PipeType(*Record));
1901  }
1902  if (!Name.startswith("spirv."))
1903  llvm_unreachable("Unknown builtin pipe type name/literal");
1904  // Parse the access qualifier literal in the name of the SPIR-V pipe type.
1905  // The name should have the following format:
1906  // spirv.Pipe._AccessQualifier
1907  // e.g. %spirv.Pipe._1
1908  if (Name.endswith("_0"))
1909  return std::unique_ptr<SPIRV::PipeType>(
1910  new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadOnly});
1911  if (Name.endswith("_1"))
1912  return std::unique_ptr<SPIRV::PipeType>(
1913  new SPIRV::PipeType{Name, SPIRV::AccessQualifier::WriteOnly});
1914  if (Name.endswith("_2"))
1915  return std::unique_ptr<SPIRV::PipeType>(
1916  new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadWrite});
1917  llvm_unreachable("Unknown pipe type access qualifier literal");
1918 }
1919 
1920 //===----------------------------------------------------------------------===//
1921 // Implementation functions for builtin types.
1922 //===----------------------------------------------------------------------===//
1923 
1924 static SPIRVType *getNonParametrizedType(const StructType *OpaqueType,
1925  const SPIRV::DemangledType *TypeRecord,
1926  MachineIRBuilder &MIRBuilder,
1927  SPIRVGlobalRegistry *GR) {
1928  unsigned Opcode = TypeRecord->Opcode;
1929  // Create or get an existing type from GlobalRegistry.
1930  return GR->getOrCreateOpTypeByOpcode(OpaqueType, MIRBuilder, Opcode);
1931 }
1932 
1934  SPIRVGlobalRegistry *GR) {
1935  // Create or get an existing type from GlobalRegistry.
1936  return GR->getOrCreateOpTypeSampler(MIRBuilder);
1937 }
1938 
1939 static SPIRVType *getPipeType(const StructType *OpaqueType,
1940  MachineIRBuilder &MIRBuilder,
1941  SPIRVGlobalRegistry *GR) {
1942  // Lookup pipe type lowering details in TableGen records or parse the
1943  // name/literal for details.
1944  std::unique_ptr<const SPIRV::PipeType> Record =
1945  lookupOrParseBuiltinPipeType(OpaqueType->getName());
1946  // Create or get an existing type from GlobalRegistry.
1947  return GR->getOrCreateOpTypePipe(MIRBuilder, Record.get()->Qualifier);
1948 }
1949 
1950 static SPIRVType *
1951 getImageType(const StructType *OpaqueType,
1953  MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
1954  // Lookup image type lowering details in TableGen records or parse the
1955  // name/literal for details.
1956  std::unique_ptr<const SPIRV::ImageType> Record =
1957  lookupOrParseBuiltinImageType(OpaqueType->getName());
1958 
1959  SPIRVType *SampledType =
1960  GR->getOrCreateSPIRVTypeByName(Record.get()->SampledType, MIRBuilder);
1961  return GR->getOrCreateOpTypeImage(
1962  MIRBuilder, SampledType, Record.get()->Dimensionality,
1963  Record.get()->Depth, Record.get()->Arrayed, Record.get()->Multisampled,
1964  Record.get()->Sampled, Record.get()->Format,
1965  AccessQual == SPIRV::AccessQualifier::WriteOnly
1966  ? SPIRV::AccessQualifier::WriteOnly
1967  : Record.get()->Qualifier);
1968 }
1969 
1970 static SPIRVType *getSampledImageType(const StructType *OpaqueType,
1971  MachineIRBuilder &MIRBuilder,
1972  SPIRVGlobalRegistry *GR) {
1973  StringRef TypeParametersString =
1974  OpaqueType->getName().substr(strlen("spirv.SampledImage."));
1975  LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
1976  Type *ImageOpaqueType = StructType::getTypeByName(
1977  Context, "spirv.Image." + TypeParametersString.str());
1978  SPIRVType *TargetImageType =
1979  GR->getOrCreateSPIRVType(ImageOpaqueType, MIRBuilder);
1980  return GR->getOrCreateOpTypeSampledImage(TargetImageType, MIRBuilder);
1981 }
1982 
1983 namespace SPIRV {
1986  MachineIRBuilder &MIRBuilder,
1987  SPIRVGlobalRegistry *GR) {
1988  assert(OpaqueType->hasName() &&
1989  "Structs representing builtin types must have a parsable name");
1990  unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();
1991 
1992  const StringRef Name = OpaqueType->getName();
1993  LLVM_DEBUG(dbgs() << "Lowering builtin type: " << Name << "\n");
1994 
1995  // Lookup the demangled builtin type in the TableGen records.
1996  const SPIRV::DemangledType *TypeRecord = findBuiltinType(Name);
1997  if (!TypeRecord)
1998  report_fatal_error("Missing TableGen record for builtin type: " + Name);
1999 
2000  // "Lower" the BuiltinType into TargetType. The following get<...>Type methods
2001  // use the implementation details from TableGen records to either create a new
2002  // OpType<...> machine instruction or get an existing equivalent SPIRVType
2003  // from GlobalRegistry.
2004  SPIRVType *TargetType;
2005  switch (TypeRecord->Opcode) {
2006  case SPIRV::OpTypeImage:
2007  TargetType = getImageType(OpaqueType, AccessQual, MIRBuilder, GR);
2008  break;
2009  case SPIRV::OpTypePipe:
2010  TargetType = getPipeType(OpaqueType, MIRBuilder, GR);
2011  break;
2012  case SPIRV::OpTypeDeviceEvent:
2013  TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
2014  break;
2015  case SPIRV::OpTypeSampler:
2016  TargetType = getSamplerType(MIRBuilder, GR);
2017  break;
2018  case SPIRV::OpTypeSampledImage:
2019  TargetType = getSampledImageType(OpaqueType, MIRBuilder, GR);
2020  break;
2021  default:
2022  TargetType = getNonParametrizedType(OpaqueType, TypeRecord, MIRBuilder, GR);
2023  break;
2024  }
2025 
2026  // Emit OpName instruction if a new OpType<...> instruction was added
2027  // (equivalent type was not found in GlobalRegistry).
2028  if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
2029  buildOpName(GR->getSPIRVTypeID(TargetType), OpaqueType->getName(),
2030  MIRBuilder);
2031 
2032  return TargetType;
2033 }
2034 } // namespace SPIRV
2035 } // namespace llvm
i
i
Definition: README.txt:29
llvm::SPIRV::ImageType::Format
ImageFormat::ImageFormat Format
Definition: SPIRVBuiltins.cpp:1819
Int32Ty
IntegerType * Int32Ty
Definition: NVVMIntrRange.cpp:67
llvm::Argument
This class represents an incoming formal argument to a Function.
Definition: Argument.h:28
llvm::SPIRV::lowerBuiltin
Optional< bool > lowerBuiltin(const StringRef DemangledCall, SPIRV::InstructionSet::InstructionSet Set, MachineIRBuilder &MIRBuilder, const Register OrigRet, const Type *OrigRetTy, const SmallVectorImpl< Register > &Args, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1724
MI
IRTranslator LLVM IR MI
Definition: IRTranslator.cpp:108
llvm::MachineInstrBuilder::addImm
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
Definition: MachineInstrBuilder.h:131
llvm::SPIRVGlobalRegistry::getOrCreateOpTypeSampledImage
SPIRVType * getOrCreateOpTypeSampledImage(SPIRVType *ImageType, MachineIRBuilder &MIRBuilder)
Definition: SPIRVGlobalRegistry.cpp:880
llvm
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:18
llvm::SPIRV::ImageType::Dimensionality
Dim::Dim Dimensionality
Definition: SPIRVBuiltins.cpp:1814
llvm::buildConstantIntReg
static Register buildConstantIntReg(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, unsigned BitWidth=32)
Definition: SPIRVBuiltins.cpp:412
llvm::CmpInst::ICMP_EQ
@ ICMP_EQ
equal
Definition: InstrTypes.h:741
llvm::SPIRVGlobalRegistry::isScalarOrVectorOfType
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
Definition: SPIRVGlobalRegistry.cpp:774
llvm::SPIRV::VectorLoadStoreBuiltin
Definition: SPIRVBuiltins.cpp:112
llvm::MachineIRBuilder::buildICmp
MachineInstrBuilder buildICmp(CmpInst::Predicate Pred, const DstOp &Res, const SrcOp &Op0, const SrcOp &Op1)
Build and insert a Res = G_ICMP Pred, Op0, Op1.
Definition: MachineIRBuilder.cpp:771
llvm::StructType::getName
StringRef getName() const
Return the name for this struct type if it has an identity.
Definition: Type.cpp:581
llvm::DataLayout
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:113
llvm::SPIRV::ImageQueryBuiltin::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:95
llvm::SPIRV::PipeType::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:1823
llvm::SPIRVGlobalRegistry::isScalarOfType
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Definition: SPIRVGlobalRegistry.cpp:767
llvm::MachineRegisterInfo::createVirtualRegister
Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
Definition: MachineRegisterInfo.cpp:156
llvm::APFloatBase::IEEEsingle
static const fltSemantics & IEEEsingle() LLVM_READNONE
Definition: APFloat.cpp:201
llvm::cl::Prefix
@ Prefix
Definition: CommandLine.h:161
llvm::getSPIRVScope
static SPIRV::Scope::Scope getSPIRVScope(SPIRV::CLMemoryScope ClScope)
Definition: SPIRVBuiltins.cpp:396
llvm::SPIRV::GetBuiltin::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:85
llvm::MachineOperand::getGlobal
const GlobalValue * getGlobal() const
Definition: MachineOperand.h:572
llvm::MachineRegisterInfo
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
Definition: MachineRegisterInfo.h:50
llvm::getSampledImageType
static SPIRVType * getSampledImageType(const StructType *OpaqueType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1970
llvm::SPIRVGlobalRegistry::getOrCreateSPIRVPointerType
SPIRVType * getOrCreateSPIRVPointerType(SPIRVType *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SClass=SPIRV::StorageClass::Function)
Definition: SPIRVGlobalRegistry.cpp:1045
CLK_NORMALIZED_COORDS_TRUE
@ CLK_NORMALIZED_COORDS_TRUE
Definition: cl_common_defines.h:87
llvm::PointerType::get
static PointerType * get(Type *ElementType, unsigned AddressSpace)
This constructs a pointer to an object of the specified type in a numbered address space.
Definition: Type.cpp:727
contains
return AArch64::GPR64RegClass contains(Reg)
llvm::ConstantInt::getValue
const APInt & getValue() const
Return the constant as an APInt value reference.
Definition: Constants.h:133
llvm::StringRef::npos
static constexpr size_t npos
Definition: StringRef.h:52
llvm::ilist_node_with_parent::getNextNode
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
Definition: ilist_node.h:289
llvm::SmallVector
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1199
llvm::MachineIRBuilder::getMRI
MachineRegisterInfo * getMRI()
Getter for MRI.
Definition: MachineIRBuilder.h:289
llvm::generateSpecConstantInst
static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1309
llvm::SPIRVGlobalRegistry::getScalarOrVectorBitWidth
unsigned getScalarOrVectorBitWidth(const SPIRVType *Type) const
Definition: SPIRVGlobalRegistry.cpp:789
llvm::SPIRV::ImageQueryBuiltin
Definition: SPIRVBuiltins.cpp:94
llvm::X86Disassembler::Reg
Reg
All possible values of the reg field in the ModR/M byte.
Definition: X86DisassemblerDecoder.h:462
llvm::generateGroupInst
static bool generateGroupInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:765
llvm::MachineIRBuilder::buildZExtOrTrunc
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...
Definition: MachineIRBuilder.cpp:532
ValueTracking.h
llvm::generateExtInst
static bool generateExtInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:718
llvm::SPIRV::VectorLoadStoreBuiltin::IsRounded
bool IsRounded
Definition: SPIRVBuiltins.cpp:116
llvm::Depth
@ Depth
Definition: SIMachineScheduler.h:36
llvm::Function::getContext
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Definition: Function.cpp:321
llvm::SPIRV::ImageType
Definition: SPIRVBuiltins.cpp:1810
llvm::StringRef::find_first_of
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:381
llvm::Type
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
llvm::SPIRVGlobalRegistry::assignSPIRVTypeToVReg
void assignSPIRVTypeToVReg(SPIRVType *Type, Register VReg, MachineFunction &MF)
Definition: SPIRVGlobalRegistry.cpp:56
llvm::tgtok::Bits
@ Bits
Definition: TGLexer.h:50
llvm::StructType::getTypeByName
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:623
llvm::AMDGPU::VOPD::Component
Component
Definition: AMDGPUBaseInfo.h:520
llvm::buildBarrierInst
static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building barriers, i.e., memory/control ordering operations.
Definition: SPIRVBuiltins.cpp:628
llvm::SPIRV::GroupBuiltin::GroupOperation
uint32_t GroupOperation
Definition: SPIRVBuiltins.cpp:68
llvm::SPIRV::DemangledBuiltin::MinNumArgs
uint8_t MinNumArgs
Definition: SPIRVBuiltins.cpp:33
llvm::MachineOperand::isCImm
bool isCImm() const
isCImm - Test if this is a MO_CImmediate operand.
Definition: MachineOperand.h:324
llvm::SPIRVGlobalRegistry::getOrCreateOpTypeSampler
SPIRVType * getOrCreateOpTypeSampler(MachineIRBuilder &MIRBuilder)
Definition: SPIRVGlobalRegistry.cpp:848
llvm::ConstantFP::getValueAPF
const APFloat & getValueAPF() const
Definition: Constants.h:298
llvm::SPIRV::IncomingCall::Arguments
const SmallVectorImpl< Register > & Arguments
Definition: SPIRVBuiltins.cpp:46
llvm::Optional< bool >
llvm::buildEnqueueKernel
static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1430
llvm::MachineRegisterInfo::getNumVirtRegs
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
Definition: MachineRegisterInfo.h:770
llvm::LLT::vector
static LLT vector(ElementCount EC, unsigned ScalarSizeInBits)
Get a low-level vector of some number of elements and element width.
Definition: LowLevelTypeImpl.h:56
llvm::generateVectorLoadStoreInst
static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1697
llvm::SPIRVGlobalRegistry::assignTypeToVReg
SPIRVType * assignTypeToVReg(const Type *Type, Register VReg, MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite, bool EmitIR=true)
Definition: SPIRVGlobalRegistry.cpp:46
llvm::SPIRV::DemangledType
Definition: SPIRVBuiltins.cpp:1802
llvm::Expected
Tagged union holding either a T or a Error.
Definition: APFloat.h:41
llvm::buildAtomicLoadInst
static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic load instruction.
Definition: SPIRVBuiltins.cpp:432
llvm::generateEnqueueInst
static bool generateEnqueueInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1515
SPIRVBuiltins.h
llvm::StructType::create
static StructType * create(LLVMContext &Context, StringRef Name)
This creates an identified struct.
Definition: Type.cpp:513
llvm::buildBuiltinVariableLoad
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...
Definition: SPIRVBuiltins.cpp:340
BaseType
llvm::getNumComponentsForDim
static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim)
Definition: SPIRVBuiltins.cpp:689
llvm::FixedVectorType
Class to represent fixed width SIMD vectors.
Definition: DerivedTypes.h:525
llvm::SPIRV::VectorLoadStoreBuiltin::RoundingMode
FPRoundingMode::FPRoundingMode RoundingMode
Definition: SPIRVBuiltins.cpp:117
llvm::SPIRV::DemangledBuiltin
Definition: SPIRVBuiltins.cpp:29
llvm::SPIRV::DemangledBuiltin::Set
InstructionSet::InstructionSet Set
Definition: SPIRVBuiltins.cpp:31
CLK_ADDRESS_CLAMP
@ CLK_ADDRESS_CLAMP
Definition: cl_common_defines.h:71
LLVM_DEBUG
#define LLVM_DEBUG(X)
Definition: Debug.h:101
llvm::SPIRVGlobalRegistry::getSPIRVTypeForVReg
SPIRVType * getSPIRVTypeForVReg(Register VReg) const
Definition: SPIRVGlobalRegistry.cpp:734
llvm::SPIRV::GroupBuiltin::Opcode
uint32_t Opcode
Definition: SPIRVBuiltins.cpp:67
llvm::generateSelectInst
static bool generateSelectInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)
Definition: SPIRVBuiltins.cpp:1302
llvm::generateBuiltinVar
static bool generateBuiltinVar(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:944
llvm::SPIRVGlobalRegistry::getTypeForSPIRVType
const Type * getTypeForSPIRVType(const SPIRVType *Ty) const
Definition: SPIRVGlobalRegistry.h:133
llvm::AMDGPU::HSAMD::AccessQualifier
AccessQualifier
Access qualifiers.
Definition: AMDGPUMetadata.h:58
llvm::findBuiltinType
static const SPIRV::DemangledType * findBuiltinType(StringRef Name)
Definition: SPIRVBuiltins.cpp:1842
llvm::SPIRV::GroupBuiltin::IsInverseBallot
bool IsInverseBallot
Definition: SPIRVBuiltins.cpp:73
Context
LLVMContext & Context
Definition: NVVMIntrRange.cpp:66
llvm::dbgs
raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition: Debug.cpp:163
llvm::getNonParametrizedType
static SPIRVType * getNonParametrizedType(const StructType *OpaqueType, const SPIRV::DemangledType *TypeRecord, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1924
llvm::LLT::fixed_vector
static LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
Definition: LowLevelTypeImpl.h:74
llvm::SPIRV::GroupBuiltin::IsLogical
bool IsLogical
Definition: SPIRVBuiltins.cpp:76
llvm::MachineInstrBuilder::addDef
const MachineInstrBuilder & addDef(Register RegNo, unsigned Flags=0, unsigned SubReg=0) const
Add a virtual register definition operand.
Definition: MachineInstrBuilder.h:116
llvm::SPIRV::GroupBuiltin
Definition: SPIRVBuiltins.cpp:65
llvm::SPIRVGlobalRegistry::getOrCreateOpTypeByOpcode
SPIRVType * getOrCreateOpTypeByOpcode(const Type *Ty, MachineIRBuilder &MIRBuilder, unsigned Opcode)
Definition: SPIRVGlobalRegistry.cpp:895
llvm::APFloat::getZero
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
Definition: APFloat.h:900
CLK_FILTER_LINEAR
@ CLK_FILTER_LINEAR
Definition: cl_common_defines.h:95
llvm::SPIRVGlobalRegistry::buildConstantFP
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType=nullptr)
Definition: SPIRVGlobalRegistry.cpp:223
llvm::generateAsyncCopy
static bool generateAsyncCopy(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1611
llvm::buildSelectInst
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...
Definition: SPIRVBuiltins.cpp:300
llvm::MachineIRBuilder::buildLoad
MachineInstrBuilder buildLoad(const DstOp &Res, const SrcOp &Addr, MachineMemOperand &MMO)
Build and insert Res = G_LOAD Addr, MMO.
Definition: MachineIRBuilder.h:904
llvm::generateBarrierInst
static bool generateBarrierInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1000
llvm::buildOpName
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
Definition: SPIRVUtils.cpp:100
llvm::PPCISD::SC
@ SC
CHAIN = SC CHAIN, Imm128 - System call.
Definition: PPCISelLowering.h:420
llvm::StringRef::startswith
bool startswith(StringRef Prefix) const
Definition: StringRef.h:260
llvm::MachineRegisterInfo::setType
void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
Definition: MachineRegisterInfo.cpp:180
llvm::SPIRV::PipeType
Definition: SPIRVBuiltins.cpp:1822
llvm::SPIRVType
const MachineInstr SPIRVType
Definition: SPIRVGlobalRegistry.h:25
llvm::MachineOperand::getImm
int64_t getImm() const
Definition: MachineOperand.h:546
CLK_LOCAL_MEM_FENCE
#define CLK_LOCAL_MEM_FENCE
Definition: cl_common_defines.h:119
llvm::getLinkStringForBuiltIn
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
Definition: SPIRVBaseInfo.cpp:149
llvm::MachineInstr::getOperand
const MachineOperand & getOperand(unsigned i) const
Definition: MachineInstr.h:526
LLVMVectorType
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:790
StorageClass
COFF::SymbolStorageClass StorageClass
Definition: COFFYAML.cpp:361
llvm::buildAtomicCompareExchangeInst
static bool buildAtomicCompareExchangeInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic compare-exchange instruction.
Definition: SPIRVBuiltins.cpp:486
llvm::generateImageMiscQueryInst
static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1099
llvm::StringRef::contains_insensitive
bool contains_insensitive(StringRef Other) const
Return true if the given string is a substring of *this, and false otherwise.
Definition: StringRef.h:440
llvm::getTypedPtrEltType
const Type * getTypedPtrEltType(const Type *Ty)
Definition: SPIRVUtils.cpp:345
llvm::generateSampleImageInst
static bool generateSampleImageInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1252
CLK_ADDRESS_MIRRORED_REPEAT
@ CLK_ADDRESS_MIRRORED_REPEAT
Definition: cl_common_defines.h:77
llvm::getNumSizeComponents
static unsigned getNumSizeComponents(SPIRVType *imgType)
Helper function for obtaining the number of size components.
Definition: SPIRVBuiltins.cpp:706
llvm::StringRef::substr
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
Definition: StringRef.h:564
llvm::SPIRV::ImageType::Arrayed
bool Arrayed
Definition: SPIRVBuiltins.cpp:1815
llvm::getSPIRVMemSemantics
static SPIRV::MemorySemantics::MemorySemantics getSPIRVMemSemantics(std::memory_order MemOrder)
Definition: SPIRVBuiltins.cpp:379
llvm::SPIRV::GroupBuiltin::IsAllEqual
bool IsAllEqual
Definition: SPIRVBuiltins.cpp:71
llvm::MachineIRBuilder::getDataLayout
const DataLayout & getDataLayout() const
Definition: MachineIRBuilder.h:281
llvm::MachineOperand
MachineOperand class - Representation of each machine instruction operand.
Definition: MachineOperand.h:48
llvm::SPIRVGlobalRegistry::getOrCreateSPIRVTypeByName
SPIRVType * getOrCreateSPIRVTypeByName(StringRef TypeStr, MachineIRBuilder &MIRBuilder)
Definition: SPIRVGlobalRegistry.cpp:916
llvm::MachineIRBuilder::getMF
MachineFunction & getMF()
Getter for the function we currently build.
Definition: MachineIRBuilder.h:271
SPIRVUtils.h
llvm::SPIRV::lowerBuiltinType
SPIRVType * lowerBuiltinType(const StructType *OpaqueType, SPIRV::AccessQualifier::AccessQualifier AccessQual, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1984
llvm::report_fatal_error
void report_fatal_error(Error Err, bool gen_crash_diag=true)
Report a serious error, calling any installed error handler.
Definition: Error.cpp:145
llvm::APInt::getZExtValue
uint64_t getZExtValue() const
Get zero extended value.
Definition: APInt.h:1486
llvm::AMDGPU::CPol::SC0
@ SC0
Definition: SIDefines.h:308
llvm::SPIRVGlobalRegistry::getOrCreateConstNullPtr
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)
Definition: SPIRVGlobalRegistry.cpp:394
llvm::SPIRV::ConvertBuiltin::IsDestinationSigned
bool IsDestinationSigned
Definition: SPIRVBuiltins.cpp:106
CLK_ADDRESS_CLAMP_TO_EDGE
@ CLK_ADDRESS_CLAMP_TO_EDGE
Definition: cl_common_defines.h:72
llvm::SPIRV::GroupBuiltin::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:66
Align
uint64_t Align
Definition: ELFObjHandler.cpp:82
llvm::APFloat::bitcastToAPInt
APInt bitcastToAPInt() const
Definition: APFloat.h:1145
llvm::lookupBuiltin
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 ...
Definition: SPIRVBuiltins.cpp:151
llvm::SPIRV::DemangledBuiltin::MaxNumArgs
uint8_t MaxNumArgs
Definition: SPIRVBuiltins.cpp:34
llvm::SPIRV::ImageType::Qualifier
AccessQualifier::AccessQualifier Qualifier
Definition: SPIRVBuiltins.cpp:1813
llvm::SPIRVGlobalRegistry::getPointerSize
unsigned getPointerSize() const
Definition: SPIRVGlobalRegistry.h:177
llvm::SPIRV::NativeBuiltin
Definition: SPIRVBuiltins.cpp:56
CLK_ADDRESS_NONE
@ CLK_ADDRESS_NONE
Definition: cl_common_defines.h:70
llvm::getMachineInstrType
static const Type * getMachineInstrType(MachineInstr *MI)
Definition: SPIRVBuiltins.cpp:1386
llvm::StringRef::slice
StringRef slice(size_t Start, size_t End) const
Return a reference to the substring from [Start, End).
Definition: StringRef.h:677
llvm::getMDOperandAsType
Type * getMDOperandAsType(const MDNode *N, unsigned I)
Definition: SPIRVUtils.cpp:239
llvm::MachineOperand::getMetadata
const MDNode * getMetadata() const
Definition: MachineOperand.h:655
llvm::LLT::pointer
static LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
Definition: LowLevelTypeImpl.h:49
llvm::Value::Value
Value(Type *Ty, unsigned scid)
Definition: Value.cpp:52
llvm::XCOFF::StorageClass
StorageClass
Definition: XCOFF.h:157
llvm::generateImageSizeQueryInst
static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1037
llvm::SPIRV::ImageQueryBuiltin::Component
uint32_t Component
Definition: SPIRVBuiltins.cpp:97
llvm::SPIRV::GroupBuiltin::NoGroupOperation
bool NoGroupOperation
Definition: SPIRVBuiltins.cpp:77
llvm::MachineIRBuilder
Helper class to build MachineInstr.
Definition: MachineIRBuilder.h:221
llvm::getIConstVal
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
Definition: SPIRVUtils.cpp:228
llvm::MachineOperand::isReg
bool isReg() const
isReg - Tests if this is a MO_Register operand.
Definition: MachineOperand.h:320
llvm::MachineOperand::getCImm
const ConstantInt * getCImm() const
Definition: MachineOperand.h:551
llvm::MachineInstr
Representation of each machine instruction.
Definition: MachineInstr.h:66
llvm::MachineInstrBuilder
Definition: MachineInstrBuilder.h:69
uint64_t
llvm::getSamplerParamFromBitmask
static unsigned getSamplerParamFromBitmask(unsigned Bitmask)
Definition: SPIRVBuiltins.cpp:1154
llvm::MCID::Barrier
@ Barrier
Definition: MCInstrDesc.h:156
llvm::SPIRV::GroupBuiltin::IsBallotFindBit
bool IsBallotFindBit
Definition: SPIRVBuiltins.cpp:75
llvm::generateAtomicInst
static bool generateAtomicInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:968
llvm::SPIRV::GetBuiltin::Set
InstructionSet::InstructionSet Set
Definition: SPIRVBuiltins.cpp:86
llvm::getOrCreateSPIRVDeviceEventPointer
static SPIRVType * getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1416
llvm::MachinePointerInfo
This class contains a discriminated union of information about pointers in memory operands,...
Definition: MachineMemOperand.h:39
llvm::LLVMContext
This is an important class for using LLVM in a threaded context.
Definition: LLVMContext.h:67
llvm::generateGetQueryInst
static bool generateGetQueryInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1025
llvm::SPIRV::DemangledBuiltin::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:30
I
#define I(x, y, z)
Definition: MD5.cpp:58
llvm::SPIRV::IncomingCall::IncomingCall
IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin, const Register ReturnRegister, const SPIRVType *ReturnType, const SmallVectorImpl< Register > &Arguments)
Definition: SPIRVBuiltins.cpp:48
llvm::buildBoolRegister
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 ...
Definition: SPIRVBuiltins.cpp:276
llvm::SPIRV::GroupBuiltin::IsBallotBitExtract
bool IsBallotBitExtract
Definition: SPIRVBuiltins.cpp:74
llvm::SPIRV::IncomingCall
Definition: SPIRVBuiltins.cpp:40
llvm::SPIRV::IncomingCall::Builtin
const DemangledBuiltin * Builtin
Definition: SPIRVBuiltins.cpp:42
llvm::SPIRVGlobalRegistry::getOrCreateConsIntVector
Register getOrCreateConsIntVector(uint64_t Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
Definition: SPIRVGlobalRegistry.cpp:290
llvm::generateConvertInst
static bool generateConvertInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1641
CLK_ADDRESS_REPEAT
@ CLK_ADDRESS_REPEAT
Definition: cl_common_defines.h:73
llvm::SPIRV::ConvertBuiltin::IsSaturated
bool IsSaturated
Definition: SPIRVBuiltins.cpp:107
llvm::getSamplerAddressingModeFromBitmask
static SPIRV::SamplerAddressingMode::SamplerAddressingMode getSamplerAddressingModeFromBitmask(unsigned Bitmask)
Definition: SPIRVBuiltins.cpp:1137
llvm::MachineOperand::getFPImm
const ConstantFP * getFPImm() const
Definition: MachineOperand.h:556
llvm::MachineIRBuilder::buildSelect
MachineInstrBuilder buildSelect(const DstOp &Res, const SrcOp &Tst, const SrcOp &Op0, const SrcOp &Op1, Optional< unsigned > Flags=None)
Build and insert a Res = G_SELECT Tst, Op0, Op1.
Definition: MachineIRBuilder.cpp:787
SPIRV.h
assert
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
llvm::getDefInstrMaybeConstant
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
Definition: SPIRVUtils.cpp:214
llvm::getBlockStructType
static const Type * getBlockStructType(Register ParamReg, MachineRegisterInfo *MRI)
Definition: SPIRVBuiltins.cpp:1399
llvm::getBlockStructInstr
static MachineInstr * getBlockStructInstr(Register ParamReg, MachineRegisterInfo *MRI)
Definition: SPIRVBuiltins.cpp:1365
llvm::SPIRV::GetBuiltin
Definition: SPIRVBuiltins.cpp:84
llvm::Type::isVoidTy
bool isVoidTy() const
Return true if this is 'void'.
Definition: Type.h:139
llvm::MachineRegisterInfo::createGenericVirtualRegister
Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
Definition: MachineRegisterInfo.cpp:186
llvm::SPIRVGlobalRegistry
Definition: SPIRVGlobalRegistry.h:27
llvm::ArrayType::get
static ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Definition: Type.cpp:638
llvm::Record
Definition: Record.h:1573
llvm::MachineInstrBuilder::addUse
const MachineInstrBuilder & addUse(Register RegNo, unsigned Flags=0, unsigned SubReg=0) const
Add a virtual register use operand.
Definition: MachineInstrBuilder.h:123
llvm::MachineOperand::getReg
Register getReg() const
getReg - Returns the register number.
Definition: MachineOperand.h:359
llvm::lookupOrParseBuiltinPipeType
static std::unique_ptr< const SPIRV::PipeType > lookupOrParseBuiltinPipeType(StringRef Name)
Definition: SPIRVBuiltins.cpp:1896
llvm::getSamplerFilterModeFromBitmask
static SPIRV::SamplerFilterMode::SamplerFilterMode getSamplerFilterModeFromBitmask(unsigned Bitmask)
Definition: SPIRVBuiltins.cpp:1159
llvm::SPIRV::GroupBuiltin::IsElect
bool IsElect
Definition: SPIRVBuiltins.cpp:69
llvm::SPIRV::ConvertBuiltin::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:104
llvm::ARM::WinEH::ReturnType
ReturnType
Definition: ARMWinEH.h:25
llvm::MachineIRBuilder::buildInstr
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
Definition: MachineIRBuilder.h:383
llvm::SPIRV::ImageType::SampledType
StringRef SampledType
Definition: SPIRVBuiltins.cpp:1812
llvm::generateReadImageInst
static bool generateReadImageInst(const StringRef DemangledCall, const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1167
llvm::ArrayRef
ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...
Definition: APInt.h:32
llvm::SPIRV::PipeType::Qualifier
AccessQualifier::AccessQualifier Qualifier
Definition: SPIRVBuiltins.cpp:1824
llvm::APInt::getAllOnesValue
static APInt getAllOnesValue(unsigned numBits)
NOTE: This is soft-deprecated. Please use getAllOnes() instead.
Definition: APInt.h:219
llvm::SPIRV::GroupBuiltin::IsBallot
bool IsBallot
Definition: SPIRVBuiltins.cpp:72
llvm::StructType
Class to represent struct types.
Definition: DerivedTypes.h:213
llvm::StringRef
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:50
llvm::MachineInstr::getOpcode
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
Definition: MachineInstr.h:516
llvm::CmpInst::ICMP_ULT
@ ICMP_ULT
unsigned less than
Definition: InstrTypes.h:745
llvm_unreachable
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
Definition: ErrorHandling.h:143
llvm::MachineIRBuilder::buildCopy
MachineInstrBuilder buildCopy(const DstOp &Res, const SrcOp &Op)
Build and insert Res = COPY Op.
Definition: MachineIRBuilder.cpp:288
Number
uint32_t Number
Definition: Profile.cpp:47
llvm::SPIRV::ConvertBuiltin::Set
InstructionSet::InstructionSet Set
Definition: SPIRVBuiltins.cpp:105
uint32_t
DL
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
Definition: AArch64SLSHardening.cpp:76
llvm::getPipeType
static SPIRVType * getPipeType(const StructType *OpaqueType, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1939
llvm::addNumImm
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
Definition: SPIRVUtils.cpp:78
llvm::SPIRV::VectorLoadStoreBuiltin::Number
uint32_t Number
Definition: SPIRVBuiltins.cpp:115
llvm::SPIRV::ImageType::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:1811
CLK_GLOBAL_MEM_FENCE
#define CLK_GLOBAL_MEM_FENCE
Definition: cl_common_defines.h:120
llvm::SPIRV::ConvertBuiltin::IsRounded
bool IsRounded
Definition: SPIRVBuiltins.cpp:108
llvm::SPIRVGlobalRegistry::getOrCreateSPIRVBoolType
SPIRVType * getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder)
Definition: SPIRVGlobalRegistry.cpp:984
llvm::SPIRVGlobalRegistry::getSPIRVTypeID
Register getSPIRVTypeID(const SPIRVType *SpirvType) const
Definition: SPIRVGlobalRegistry.cpp:632
MRI
unsigned const MachineRegisterInfo * MRI
Definition: AArch64AdvSIMDScalarPass.cpp:105
llvm::SPIRV::NativeBuiltin::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:57
llvm::Register
Wrapper class representing virtual and physical registers.
Definition: Register.h:19
llvm::buildAtomicRMWInst
static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic load instruction.
Definition: SPIRVBuiltins.cpp:584
llvm::MCID::Select
@ Select
Definition: MCInstrDesc.h:164
llvm::SPIRV::ImageType::Sampled
bool Sampled
Definition: SPIRVBuiltins.cpp:1818
llvm::generateWriteImageInst
static bool generateWriteImageInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1242
CLK_FILTER_NEAREST
@ CLK_FILTER_NEAREST
Definition: cl_common_defines.h:94
llvm::SPIRV::DemangledBuiltin::Group
BuiltinGroup Group
Definition: SPIRVBuiltins.cpp:32
llvm::SPIRV::DemangledType::Opcode
uint32_t Opcode
Definition: SPIRVBuiltins.cpp:1804
llvm::StructType::hasName
bool hasName() const
Return true if this is a named struct that has a non-empty name.
Definition: DerivedTypes.h:290
llvm::SPIRV::ConvertBuiltin
Definition: SPIRVBuiltins.cpp:103
llvm::insertAssignInstr
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,...
Definition: SPIRVPreLegalizer.cpp:195
llvm::SPIRV::ImageType::Multisampled
bool Multisampled
Definition: SPIRVBuiltins.cpp:1817
llvm::SPIRV::VectorLoadStoreBuiltin::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:113
llvm::MachineFunction::getFunction
Function & getFunction()
Return the LLVM function that this machine code represents.
Definition: MachineFunction.h:623
llvm::None
constexpr std::nullopt_t None
Definition: None.h:27
llvm::SPIRV::IncomingCall::ReturnType
const SPIRVType * ReturnType
Definition: SPIRVBuiltins.cpp:45
llvm::buildAtomicInitInst
static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder)
Helper function for translating atomic init to OpStore.
Definition: SPIRVBuiltins.cpp:420
llvm::SPIRVGlobalRegistry::getOrCreateSPIRVIntegerType
SPIRVType * getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
Definition: SPIRVGlobalRegistry.cpp:953
llvm::SPIRVGlobalRegistry::getOrCreateOpTypeDeviceEvent
SPIRVType * getOrCreateOpTypeDeviceEvent(MachineIRBuilder &MIRBuilder)
Definition: SPIRVGlobalRegistry.cpp:870
llvm::SPIRVGlobalRegistry::getOrCreateOpTypePipe
SPIRVType * getOrCreateOpTypePipe(MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AccQual)
Definition: SPIRVGlobalRegistry.cpp:857
llvm::BitWidth
constexpr unsigned BitWidth
Definition: BitmaskEnum.h:147
Arguments
AMDGPU Lower Kernel Arguments
Definition: AMDGPULowerKernelArguments.cpp:242
llvm::generateRelationalInst
static bool generateRelationalInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:739
llvm::SPIRV::IncomingCall::ReturnRegister
const Register ReturnRegister
Definition: SPIRVBuiltins.cpp:44
llvm::SPIRV::IncomingCall::BuiltinName
const std::string BuiltinName
Definition: SPIRVBuiltins.cpp:41
llvm::genWorkgroupQuery
static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR, SPIRV::BuiltIn::BuiltIn BuiltinValue, uint64_t DefaultValue)
Definition: SPIRVBuiltins.cpp:850
llvm::StringRef::contains
bool contains(StringRef Other) const
Return true if the given string is a substring of *this, and false otherwise.
Definition: StringRef.h:428
llvm::MachineIRBuilder::buildIntrinsic
MachineInstrBuilder buildIntrinsic(Intrinsic::ID ID, ArrayRef< Register > Res, bool HasSideEffects)
Build and insert either a G_INTRINSIC (if HasSideEffects is false) or G_INTRINSIC_W_SIDE_EFFECTS inst...
Definition: MachineIRBuilder.cpp:736
llvm::getMemSemanticsForStorageClass
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
Definition: SPIRVUtils.cpp:178
llvm::SPIRV::GroupBuiltin::HasBoolArg
bool HasBoolArg
Definition: SPIRVBuiltins.cpp:78
llvm::SPIRV::GetBuiltin::Value
BuiltIn::BuiltIn Value
Definition: SPIRVBuiltins.cpp:87
llvm::SPIRV::GroupBuiltin::IsAllOrAny
bool IsAllOrAny
Definition: SPIRVBuiltins.cpp:70
llvm::SPIRVGlobalRegistry::buildConstantInt
Register buildConstantInt(uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType=nullptr, bool EmitIR=true)
Definition: SPIRVGlobalRegistry.cpp:177
llvm::Register::isValid
bool isValid() const
Definition: Register.h:126
llvm::storageClassToAddressSpace
unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
Definition: SPIRVUtils.cpp:138
llvm::SPIRV::ConvertBuiltin::RoundingMode
FPRoundingMode::FPRoundingMode RoundingMode
Definition: SPIRVBuiltins.cpp:109
llvm::SPIRV::VectorLoadStoreBuiltin::Set
InstructionSet::InstructionSet Set
Definition: SPIRVBuiltins.cpp:114
llvm::SPIRV::ImageType::Depth
bool Depth
Definition: SPIRVBuiltins.cpp:1816
llvm::SmallVectorImpl
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
Definition: APFloat.h:42
llvm::SPIRV::ImageQueryBuiltin::Set
InstructionSet::InstructionSet Set
Definition: SPIRVBuiltins.cpp:96
llvm::buildOpDecorate
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
Definition: SPIRVUtils.cpp:117
llvm::SPIRVGlobalRegistry::getOrCreateConsIntArray
Register getOrCreateConsIntArray(uint64_t Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
Definition: SPIRVGlobalRegistry.cpp:306
llvm::SPIRV::NativeBuiltin::Opcode
uint32_t Opcode
Definition: SPIRVBuiltins.cpp:59
llvm::getSamplerType
static SPIRVType * getSamplerType(MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1933
llvm::StringRef::str
std::string str() const
str - Get the contents as an std::string.
Definition: StringRef.h:221
llvm::StringRef::find
size_t find(char C, size_t From=0) const
Search for the first character C in the string.
Definition: StringRef.h:294
llvm::IntegerType::get
static IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
Definition: Type.cpp:311
llvm::SPIRV::DemangledType::Name
StringRef Name
Definition: SPIRVBuiltins.cpp:1803
llvm::getImageType
static SPIRVType * getImageType(const StructType *OpaqueType, SPIRV::AccessQualifier::AccessQualifier AccessQual, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1951
llvm::AMDGPU::CPol::SC1
@ SC1
Definition: SIDefines.h:309
llvm::isSpvIntrinsic
bool isSpvIntrinsic(MachineInstr &MI, Intrinsic::ID IntrinsicID)
Definition: SPIRVUtils.cpp:234
llvm::SPIRV::NativeBuiltin::Set
InstructionSet::InstructionSet Set
Definition: SPIRVBuiltins.cpp:58
llvm::AMDGPU::HSAMD::Kernel::Key::Args
constexpr char Args[]
Key for Kernel::Metadata::mArgs.
Definition: AMDGPUMetadata.h:394
llvm::buildAtomicStoreInst
static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Helper function for building an atomic store instruction.
Definition: SPIRVBuiltins.cpp:466
llvm::LLT::scalar
static LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
Definition: LowLevelTypeImpl.h:42
llvm::SPIRVGlobalRegistry::getPointerStorageClass
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
Definition: SPIRVGlobalRegistry.cpp:815
llvm::SPIRVGlobalRegistry::getOrCreateSPIRVVectorType
SPIRVType * getOrCreateSPIRVVectorType(SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder)
Definition: SPIRVGlobalRegistry.cpp:1003
llvm::generateDotOrFMulInst
static bool generateDotOrFMulInst(const SPIRV::IncomingCall *Call, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR)
Definition: SPIRVBuiltins.cpp:1011
llvm::StringRef::split
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
Definition: StringRef.h:693
llvm::SPIRVGlobalRegistry::buildConstantSampler
Register buildConstantSampler(Register Res, unsigned AddrMode, unsigned Param, unsigned FilerMode, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)
Definition: SPIRVGlobalRegistry.cpp:413
llvm::Value
LLVM Value Representation.
Definition: Value.h:74
llvm::buildLoadInst
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.
Definition: SPIRVBuiltins.cpp:321
llvm::SPIRVGlobalRegistry::buildGlobalVariable
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)
Definition: SPIRVGlobalRegistry.cpp:436
llvm::SPIRVGlobalRegistry::getOrCreateOpTypeImage
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)
Definition: SPIRVGlobalRegistry.cpp:823
llvm::lookupOrParseBuiltinImageType
static std::unique_ptr< const SPIRV::ImageType > lookupOrParseBuiltinImageType(StringRef Name)
Definition: SPIRVBuiltins.cpp:1857
llvm::LLT
Definition: LowLevelTypeImpl.h:39
llvm::SPIRVGlobalRegistry::getOrCreateSPIRVType
SPIRVType * getOrCreateSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder, SPIRV::AccessQualifier::AccessQualifier AQ=SPIRV::AccessQualifier::ReadWrite, bool EmitIR=true)
Definition: SPIRVGlobalRegistry.cpp:744