contrib/llvm-project/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp

   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"
  17 #include "llvm/Analysis/ValueTracking.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
  29 struct DemangledBuiltin {
  30   StringRef Name;
  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;
  42   const DemangledBuiltin *Builtin;
  43
  44   const Register ReturnRegister;
  45   const SPIRVType *ReturnType;
  46   const SmallVectorImpl<Register> &Arguments;
  47
  48   IncomingCall(const std::string BuiltinName, const DemangledBuiltin *Builtin,
  49                const Register ReturnRegister, const SPIRVType *ReturnType,
  50                const SmallVectorImpl<Register> &Arguments)
  51       : BuiltinName(BuiltinName), Builtin(Builtin),
  52         ReturnRegister(ReturnRegister), ReturnType(ReturnType),
  53         Arguments(Arguments) {}
  54 };
  55
  56 struct NativeBuiltin {
  57   StringRef Name;
  58   InstructionSet::InstructionSet Set;
  59   uint32_t Opcode;
  60 };
  61
  62 #define GET_NativeBuiltins_DECL
  63 #define GET_NativeBuiltins_IMPL
  64
  65 struct GroupBuiltin {
  66   StringRef Name;
  67   uint32_t Opcode;
  68   uint32_t GroupOperation;
  69   bool IsElect;
  70   bool IsAllOrAny;
  71   bool IsAllEqual;
  72   bool IsBallot;
  73   bool IsInverseBallot;
  74   bool IsBallotBitExtract;
  75   bool IsBallotFindBit;
  76   bool IsLogical;
  77   bool NoGroupOperation;
  78   bool HasBoolArg;
  79 };
  80
  81 #define GET_GroupBuiltins_DECL
  82 #define GET_GroupBuiltins_IMPL
  83
  84 struct GetBuiltin {
  85   StringRef Name;
  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
  94 struct ImageQueryBuiltin {
  95   StringRef Name;
  96   InstructionSet::InstructionSet Set;
  97   uint32_t Component;
  98 };
  99
 100 #define GET_ImageQueryBuiltins_DECL
 101 #define GET_ImageQueryBuiltins_IMPL
 102
 103 struct ConvertBuiltin {
 104   StringRef Name;
 105   InstructionSet::InstructionSet Set;
 106   bool IsDestinationSigned;
 107   bool IsSaturated;
 108   bool IsRounded;
 109   FPRoundingMode::FPRoundingMode RoundingMode;
 110 };
 111
 112 struct VectorLoadStoreBuiltin {
 113   StringRef Name;
 114   InstructionSet::InstructionSet Set;
 115   uint32_t Number;
 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,
 154               const SmallVectorImpl<Register> &Arguments) {
 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);
 285     const FixedVectorType *LLVMVectorType =
 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 =
 293       MIRBuilder.getMRI()->createGenericVirtualRegister(Type);
 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);
 308     uint64_t AllOnes = APInt::getAllOnesValue(Bits).getZExtValue();
 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.
 321 static Register buildLoadInst(SPIRVType *BaseType, Register PtrRegister,
 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).
 333   MachinePointerInfo PtrInfo = MachinePointerInfo();
 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.
 340 static Register buildBuiltinVariableLoad(MachineIRBuilder &MIRBuilder,
 341                                          SPIRVType *VariableType,
 342                                          SPIRVGlobalRegistry *GR,
 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
 412 static Register buildConstantIntReg(uint64_t Val, MachineIRBuilder &MIRBuilder,
 413                                     SPIRVGlobalRegistry *GR,
 414                                     unsigned BitWidth = 32) {
 415   SPIRVType *IntType = GR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
 416   return GR->buildConstantInt(Val, MIRBuilder, IntType);
 417 }
 418
 419 static Register buildScopeReg(Register CLScopeRegister,
 420                               MachineIRBuilder &MIRBuilder,
 421                               SPIRVGlobalRegistry *GR,
 422                               const MachineRegisterInfo *MRI) {
 423   auto CLScope =
 424       static_cast<SPIRV::CLMemoryScope>(getIConstVal(CLScopeRegister, MRI));
 425   SPIRV::Scope::Scope Scope = getSPIRVScope(CLScope);
 426
 427   if (CLScope == static_cast<unsigned>(Scope))
 428     return CLScopeRegister;
 429
 430   return buildConstantIntReg(Scope, MIRBuilder, GR);
 431 }
 432
 433 static Register buildMemSemanticsReg(Register SemanticsRegister,
 434                                      Register PtrRegister,
 435                                      const MachineRegisterInfo *MRI,
 436                                      SPIRVGlobalRegistry *GR) {
 437   std::memory_order Order =
 438       static_cast<std::memory_order>(getIConstVal(SemanticsRegister, MRI));
 439   unsigned Semantics =
 440       getSPIRVMemSemantics(Order) |
 441       getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
 442
 443   if (Order == Semantics)
 444     return SemanticsRegister;
 445
 446   return Register();
 447 }
 448
 449 /// Helper function for translating atomic init to OpStore.
 450 static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call,
 451                                 MachineIRBuilder &MIRBuilder) {
 452   assert(Call->Arguments.size() == 2 &&
 453          "Need 2 arguments for atomic init translation");
 454
 455   MIRBuilder.buildInstr(SPIRV::OpStore)
 456       .addUse(Call->Arguments[0])
 457       .addUse(Call->Arguments[1]);
 458   return true;
 459 }
 460
 461 /// Helper function for building an atomic load instruction.
 462 static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call,
 463                                 MachineIRBuilder &MIRBuilder,
 464                                 SPIRVGlobalRegistry *GR) {
 465   Register PtrRegister = Call->Arguments[0];
 466   // TODO: if true insert call to __translate_ocl_memory_sccope before
 467   // OpAtomicLoad and the function implementation. We can use Translator's
 468   // output for transcoding/atomic_explicit_arguments.cl as an example.
 469   Register ScopeRegister;
 470   if (Call->Arguments.size() > 1)
 471     ScopeRegister = Call->Arguments[1];
 472   else
 473     ScopeRegister = buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
 474
 475   Register MemSemanticsReg;
 476   if (Call->Arguments.size() > 2) {
 477     // TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
 478     MemSemanticsReg = Call->Arguments[2];
 479   } else {
 480     int Semantics =
 481         SPIRV::MemorySemantics::SequentiallyConsistent |
 482         getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
 483     MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
 484   }
 485
 486   MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
 487       .addDef(Call->ReturnRegister)
 488       .addUse(GR->getSPIRVTypeID(Call->ReturnType))
 489       .addUse(PtrRegister)
 490       .addUse(ScopeRegister)
 491       .addUse(MemSemanticsReg);
 492   return true;
 493 }
 494
 495 /// Helper function for building an atomic store instruction.
 496 static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call,
 497                                  MachineIRBuilder &MIRBuilder,
 498                                  SPIRVGlobalRegistry *GR) {
 499   Register ScopeRegister =
 500       buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
 501   Register PtrRegister = Call->Arguments[0];
 502   int Semantics =
 503       SPIRV::MemorySemantics::SequentiallyConsistent |
 504       getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
 505   Register MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
 506
 507   MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
 508       .addUse(PtrRegister)
 509       .addUse(ScopeRegister)
 510       .addUse(MemSemanticsReg)
 511       .addUse(Call->Arguments[1]);
 512   return true;
 513 }
 514
 515 /// Helper function for building an atomic compare-exchange instruction.
 516 static bool buildAtomicCompareExchangeInst(const SPIRV::IncomingCall *Call,
 517                                            MachineIRBuilder &MIRBuilder,
 518                                            SPIRVGlobalRegistry *GR) {
 519   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
 520   unsigned Opcode =
 521       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
 522   bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");
 523   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 524
 525   Register ObjectPtr = Call->Arguments[0];   // Pointer (volatile A *object.)
 526   Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
 527   Register Desired = Call->Arguments[2];     // Value (C Desired).
 528   SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);
 529   LLT DesiredLLT = MRI->getType(Desired);
 530
 531   assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
 532          SPIRV::OpTypePointer);
 533   unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
 534   assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
 535                    : ExpectedType == SPIRV::OpTypePointer);
 536   assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
 537
 538   SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);
 539   assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
 540   auto StorageClass = static_cast<SPIRV::StorageClass::StorageClass>(
 541       SpvObjectPtrTy->getOperand(1).getImm());
 542   auto MemSemStorage = getMemSemanticsForStorageClass(StorageClass);
 543
 544   Register MemSemEqualReg;
 545   Register MemSemUnequalReg;
 546   uint64_t MemSemEqual =
 547       IsCmpxchg
 548           ? SPIRV::MemorySemantics::None
 549           : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
 550   uint64_t MemSemUnequal =
 551       IsCmpxchg
 552           ? SPIRV::MemorySemantics::None
 553           : SPIRV::MemorySemantics::SequentiallyConsistent | MemSemStorage;
 554   if (Call->Arguments.size() >= 4) {
 555     assert(Call->Arguments.size() >= 5 &&
 556            "Need 5+ args for explicit atomic cmpxchg");
 557     auto MemOrdEq =
 558         static_cast<std::memory_order>(getIConstVal(Call->Arguments[3], MRI));
 559     auto MemOrdNeq =
 560         static_cast<std::memory_order>(getIConstVal(Call->Arguments[4], MRI));
 561     MemSemEqual = getSPIRVMemSemantics(MemOrdEq) | MemSemStorage;
 562     MemSemUnequal = getSPIRVMemSemantics(MemOrdNeq) | MemSemStorage;
 563     if (MemOrdEq == MemSemEqual)
 564       MemSemEqualReg = Call->Arguments[3];
 565     if (MemOrdNeq == MemSemEqual)
 566       MemSemUnequalReg = Call->Arguments[4];
 567   }
 568   if (!MemSemEqualReg.isValid())
 569     MemSemEqualReg = buildConstantIntReg(MemSemEqual, MIRBuilder, GR);
 570   if (!MemSemUnequalReg.isValid())
 571     MemSemUnequalReg = buildConstantIntReg(MemSemUnequal, MIRBuilder, GR);
 572
 573   Register ScopeReg;
 574   auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
 575   if (Call->Arguments.size() >= 6) {
 576     assert(Call->Arguments.size() == 6 &&
 577            "Extra args for explicit atomic cmpxchg");
 578     auto ClScope = static_cast<SPIRV::CLMemoryScope>(
 579         getIConstVal(Call->Arguments[5], MRI));
 580     Scope = getSPIRVScope(ClScope);
 581     if (ClScope == static_cast<unsigned>(Scope))
 582       ScopeReg = Call->Arguments[5];
 583   }
 584   if (!ScopeReg.isValid())
 585     ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
 586
 587   Register Expected = IsCmpxchg
 588                           ? ExpectedArg
 589                           : buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,
 590                                           GR, LLT::scalar(32));
 591   MRI->setType(Expected, DesiredLLT);
 592   Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT)
 593                             : Call->ReturnRegister;
 594   GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
 595
 596   SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
 597   MIRBuilder.buildInstr(Opcode)
 598       .addDef(Tmp)
 599       .addUse(GR->getSPIRVTypeID(IntTy))
 600       .addUse(ObjectPtr)
 601       .addUse(ScopeReg)
 602       .addUse(MemSemEqualReg)
 603       .addUse(MemSemUnequalReg)
 604       .addUse(Desired)
 605       .addUse(Expected);
 606   if (!IsCmpxchg) {
 607     MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);
 608     MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected);
 609   }
 610   return true;
 611 }
 612
 613 /// Helper function for building an atomic load instruction.
 614 static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
 615                                MachineIRBuilder &MIRBuilder,
 616                                SPIRVGlobalRegistry *GR) {
 617   const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 618   SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
 619   Register ScopeRegister;
 620
 621   if (Call->Arguments.size() >= 4) {
 622     assert(Call->Arguments.size() == 4 &&
 623            "Too many args for explicit atomic RMW");
 624     ScopeRegister = buildScopeReg(Call->Arguments[3], MIRBuilder, GR, MRI);
 625   }
 626
 627   if (!ScopeRegister.isValid())
 628     ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
 629
 630   Register PtrRegister = Call->Arguments[0];
 631   unsigned Semantics = SPIRV::MemorySemantics::None;
 632   Register MemSemanticsReg;
 633
 634   if (Call->Arguments.size() >= 3)
 635     MemSemanticsReg =
 636         buildMemSemanticsReg(Call->Arguments[2], PtrRegister, MRI, GR);
 637
 638   if (!MemSemanticsReg.isValid())
 639     MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
 640
 641   MIRBuilder.buildInstr(Opcode)
 642       .addDef(Call->ReturnRegister)
 643       .addUse(GR->getSPIRVTypeID(Call->ReturnType))
 644       .addUse(PtrRegister)
 645       .addUse(ScopeRegister)
 646       .addUse(MemSemanticsReg)
 647       .addUse(Call->Arguments[1]);
 648   return true;
 649 }
 650
 651 /// Helper function for building atomic flag instructions (e.g.
 652 /// OpAtomicFlagTestAndSet).
 653 static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call,
 654                                 unsigned Opcode, MachineIRBuilder &MIRBuilder,
 655                                 SPIRVGlobalRegistry *GR) {
 656   const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 657
 658   Register PtrRegister = Call->Arguments[0];
 659   unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;
 660   Register MemSemanticsReg;
 661
 662   if (Call->Arguments.size() >= 2)
 663     MemSemanticsReg =
 664         buildMemSemanticsReg(Call->Arguments[1], PtrRegister, MRI, GR);
 665
 666   if (!MemSemanticsReg.isValid())
 667     MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
 668
 669   assert((Opcode != SPIRV::OpAtomicFlagClear ||
 670           (Semantics != SPIRV::MemorySemantics::Acquire &&
 671            Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&
 672          "Invalid memory order argument!");
 673
 674   SPIRV::Scope::Scope Scope = SPIRV::Scope::Device;
 675   Register ScopeRegister;
 676
 677   if (Call->Arguments.size() >= 3)
 678     ScopeRegister = buildScopeReg(Call->Arguments[2], MIRBuilder, GR, MRI);
 679
 680   if (!ScopeRegister.isValid())
 681     ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
 682
 683   auto MIB = MIRBuilder.buildInstr(Opcode);
 684   if (Opcode == SPIRV::OpAtomicFlagTestAndSet)
 685     MIB.addDef(Call->ReturnRegister)
 686         .addUse(GR->getSPIRVTypeID(Call->ReturnType));
 687
 688   MIB.addUse(PtrRegister).addUse(ScopeRegister).addUse(MemSemanticsReg);
 689   return true;
 690 }
 691
 692 /// Helper function for building barriers, i.e., memory/control ordering
 693 /// operations.
 694 static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
 695                              MachineIRBuilder &MIRBuilder,
 696                              SPIRVGlobalRegistry *GR) {
 697   const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 698   unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI);
 699   unsigned MemSemantics = SPIRV::MemorySemantics::None;
 700
 701   if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
 702     MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
 703
 704   if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
 705     MemSemantics |= SPIRV::MemorySemantics::CrossWorkgroupMemory;
 706
 707   if (MemFlags & SPIRV::CLK_IMAGE_MEM_FENCE)
 708     MemSemantics |= SPIRV::MemorySemantics::ImageMemory;
 709
 710   if (Opcode == SPIRV::OpMemoryBarrier) {
 711     std::memory_order MemOrder =
 712         static_cast<std::memory_order>(getIConstVal(Call->Arguments[1], MRI));
 713     MemSemantics = getSPIRVMemSemantics(MemOrder) | MemSemantics;
 714   } else {
 715     MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
 716   }
 717
 718   Register MemSemanticsReg;
 719   if (MemFlags == MemSemantics)
 720     MemSemanticsReg = Call->Arguments[0];
 721   else
 722     MemSemanticsReg = buildConstantIntReg(MemSemantics, MIRBuilder, GR);
 723
 724   Register ScopeReg;
 725   SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
 726   SPIRV::Scope::Scope MemScope = Scope;
 727   if (Call->Arguments.size() >= 2) {
 728     assert(
 729         ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
 730          (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
 731         "Extra args for explicitly scoped barrier");
 732     Register ScopeArg = (Opcode == SPIRV::OpMemoryBarrier) ? Call->Arguments[2]
 733                                                            : Call->Arguments[1];
 734     SPIRV::CLMemoryScope CLScope =
 735         static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg, MRI));
 736     MemScope = getSPIRVScope(CLScope);
 737     if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
 738         (Opcode == SPIRV::OpMemoryBarrier))
 739       Scope = MemScope;
 740
 741     if (CLScope == static_cast<unsigned>(Scope))
 742       ScopeReg = Call->Arguments[1];
 743   }
 744
 745   if (!ScopeReg.isValid())
 746     ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
 747
 748   auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);
 749   if (Opcode != SPIRV::OpMemoryBarrier)
 750     MIB.addUse(buildConstantIntReg(MemScope, MIRBuilder, GR));
 751   MIB.addUse(MemSemanticsReg);
 752   return true;
 753 }
 754
 755 static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
 756   switch (dim) {
 757   case SPIRV::Dim::DIM_1D:
 758   case SPIRV::Dim::DIM_Buffer:
 759     return 1;
 760   case SPIRV::Dim::DIM_2D:
 761   case SPIRV::Dim::DIM_Cube:
 762   case SPIRV::Dim::DIM_Rect:
 763     return 2;
 764   case SPIRV::Dim::DIM_3D:
 765     return 3;
 766   default:
 767     llvm_unreachable("Cannot get num components for given Dim");
 768   }
 769 }
 770
 771 /// Helper function for obtaining the number of size components.
 772 static unsigned getNumSizeComponents(SPIRVType *imgType) {
 773   assert(imgType->getOpcode() == SPIRV::OpTypeImage);
 774   auto dim = static_cast<SPIRV::Dim::Dim>(imgType->getOperand(2).getImm());
 775   unsigned numComps = getNumComponentsForDim(dim);
 776   bool arrayed = imgType->getOperand(4).getImm() == 1;
 777   return arrayed ? numComps + 1 : numComps;
 778 }
 779
 780 //===----------------------------------------------------------------------===//
 781 // Implementation functions for each builtin group
 782 //===----------------------------------------------------------------------===//
 783
 784 static bool generateExtInst(const SPIRV::IncomingCall *Call,
 785                             MachineIRBuilder &MIRBuilder,
 786                             SPIRVGlobalRegistry *GR) {
 787   // Lookup the extended instruction number in the TableGen records.
 788   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
 789   uint32_t Number =
 790       SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;
 791
 792   // Build extended instruction.
 793   auto MIB =
 794       MIRBuilder.buildInstr(SPIRV::OpExtInst)
 795           .addDef(Call->ReturnRegister)
 796           .addUse(GR->getSPIRVTypeID(Call->ReturnType))
 797           .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
 798           .addImm(Number);
 799
 800   for (auto Argument : Call->Arguments)
 801     MIB.addUse(Argument);
 802   return true;
 803 }
 804
 805 static bool generateRelationalInst(const SPIRV::IncomingCall *Call,
 806                                    MachineIRBuilder &MIRBuilder,
 807                                    SPIRVGlobalRegistry *GR) {
 808   // Lookup the instruction opcode in the TableGen records.
 809   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
 810   unsigned Opcode =
 811       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
 812
 813   Register CompareRegister;
 814   SPIRVType *RelationType;
 815   std::tie(CompareRegister, RelationType) =
 816       buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
 817
 818   // Build relational instruction.
 819   auto MIB = MIRBuilder.buildInstr(Opcode)
 820                  .addDef(CompareRegister)
 821                  .addUse(GR->getSPIRVTypeID(RelationType));
 822
 823   for (auto Argument : Call->Arguments)
 824     MIB.addUse(Argument);
 825
 826   // Build select instruction.
 827   return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,
 828                          Call->ReturnType, GR);
 829 }
 830
 831 static bool generateGroupInst(const SPIRV::IncomingCall *Call,
 832                               MachineIRBuilder &MIRBuilder,
 833                               SPIRVGlobalRegistry *GR) {
 834   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
 835   const SPIRV::GroupBuiltin *GroupBuiltin =
 836       SPIRV::lookupGroupBuiltin(Builtin->Name);
 837   const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 838   Register Arg0;
 839   if (GroupBuiltin->HasBoolArg) {
 840     Register ConstRegister = Call->Arguments[0];
 841     auto ArgInstruction = getDefInstrMaybeConstant(ConstRegister, MRI);
 842     // TODO: support non-constant bool values.
 843     assert(ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT &&
 844            "Only constant bool value args are supported");
 845     if (GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() !=
 846         SPIRV::OpTypeBool)
 847       Arg0 = GR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder,
 848                                   GR->getOrCreateSPIRVBoolType(MIRBuilder));
 849   }
 850
 851   Register GroupResultRegister = Call->ReturnRegister;
 852   SPIRVType *GroupResultType = Call->ReturnType;
 853
 854   // TODO: maybe we need to check whether the result type is already boolean
 855   // and in this case do not insert select instruction.
 856   const bool HasBoolReturnTy =
 857       GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
 858       GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
 859       GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
 860
 861   if (HasBoolReturnTy)
 862     std::tie(GroupResultRegister, GroupResultType) =
 863         buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
 864
 865   auto Scope = Builtin->Name.startswith("sub_group") ? SPIRV::Scope::Subgroup
 866                                                      : SPIRV::Scope::Workgroup;
 867   Register ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
 868
 869   // Build work/sub group instruction.
 870   auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
 871                  .addDef(GroupResultRegister)
 872                  .addUse(GR->getSPIRVTypeID(GroupResultType))
 873                  .addUse(ScopeRegister);
 874
 875   if (!GroupBuiltin->NoGroupOperation)
 876     MIB.addImm(GroupBuiltin->GroupOperation);
 877   if (Call->Arguments.size() > 0) {
 878     MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);
 879     for (unsigned i = 1; i < Call->Arguments.size(); i++)
 880       MIB.addUse(Call->Arguments[i]);
 881   }
 882
 883   // Build select instruction.
 884   if (HasBoolReturnTy)
 885     buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
 886                     Call->ReturnType, GR);
 887   return true;
 888 }
 889
 890 // These queries ask for a single size_t result for a given dimension index, e.g
 891 // size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
 892 // these values are all vec3 types, so we need to extract the correct index or
 893 // return defaultVal (0 or 1 depending on the query). We also handle extending
 894 // or tuncating in case size_t does not match the expected result type's
 895 // bitwidth.
 896 //
 897 // For a constant index >= 3 we generate:
 898 //  %res = OpConstant %SizeT 0
 899 //
 900 // For other indices we generate:
 901 //  %g = OpVariable %ptr_V3_SizeT Input
 902 //  OpDecorate %g BuiltIn XXX
 903 //  OpDecorate %g LinkageAttributes "__spirv_BuiltInXXX"
 904 //  OpDecorate %g Constant
 905 //  %loadedVec = OpLoad %V3_SizeT %g
 906 //
 907 //  Then, if the index is constant < 3, we generate:
 908 //    %res = OpCompositeExtract %SizeT %loadedVec idx
 909 //  If the index is dynamic, we generate:
 910 //    %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
 911 //    %cmp = OpULessThan %bool %idx %const_3
 912 //    %res = OpSelect %SizeT %cmp %tmp %const_0
 913 //
 914 //  If the bitwidth of %res does not match the expected return type, we add an
 915 //  extend or truncate.
 916 static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call,
 917                               MachineIRBuilder &MIRBuilder,
 918                               SPIRVGlobalRegistry *GR,
 919                               SPIRV::BuiltIn::BuiltIn BuiltinValue,
 920                               uint64_t DefaultValue) {
 921   Register IndexRegister = Call->Arguments[0];
 922   const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();
 923   const unsigned PointerSize = GR->getPointerSize();
 924   const SPIRVType *PointerSizeType =
 925       GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
 926   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 927   auto IndexInstruction = getDefInstrMaybeConstant(IndexRegister, MRI);
 928
 929   // Set up the final register to do truncation or extension on at the end.
 930   Register ToTruncate = Call->ReturnRegister;
 931
 932   // If the index is constant, we can statically determine if it is in range.
 933   bool IsConstantIndex =
 934       IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
 935
 936   // If it's out of range (max dimension is 3), we can just return the constant
 937   // default value (0 or 1 depending on which query function).
 938   if (IsConstantIndex && getIConstVal(IndexRegister, MRI) >= 3) {
 939     Register defaultReg = Call->ReturnRegister;
 940     if (PointerSize != ResultWidth) {
 941       defaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
 942       GR->assignSPIRVTypeToVReg(PointerSizeType, defaultReg,
 943                                 MIRBuilder.getMF());
 944       ToTruncate = defaultReg;
 945     }
 946     auto NewRegister =
 947         GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
 948     MIRBuilder.buildCopy(defaultReg, NewRegister);
 949   } else { // If it could be in range, we need to load from the given builtin.
 950     auto Vec3Ty =
 951         GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);
 952     Register LoadedVector =
 953         buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,
 954                                  LLT::fixed_vector(3, PointerSize));
 955     // Set up the vreg to extract the result to (possibly a new temporary one).
 956     Register Extracted = Call->ReturnRegister;
 957     if (!IsConstantIndex || PointerSize != ResultWidth) {
 958       Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
 959       GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());
 960     }
 961     // Use Intrinsic::spv_extractelt so dynamic vs static extraction is
 962     // handled later: extr = spv_extractelt LoadedVector, IndexRegister.
 963     MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
 964         Intrinsic::spv_extractelt, ArrayRef<Register>{Extracted}, true);
 965     ExtractInst.addUse(LoadedVector).addUse(IndexRegister);
 966
 967     // If the index is dynamic, need check if it's < 3, and then use a select.
 968     if (!IsConstantIndex) {
 969       insertAssignInstr(Extracted, nullptr, PointerSizeType, GR, MIRBuilder,
 970                         *MRI);
 971
 972       auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);
 973       auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
 974
 975       Register CompareRegister =
 976           MRI->createGenericVirtualRegister(LLT::scalar(1));
 977       GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
 978
 979       // Use G_ICMP to check if idxVReg < 3.
 980       MIRBuilder.buildICmp(CmpInst::ICMP_ULT, CompareRegister, IndexRegister,
 981                            GR->buildConstantInt(3, MIRBuilder, IndexType));
 982
 983       // Get constant for the default value (0 or 1 depending on which
 984       // function).
 985       Register DefaultRegister =
 986           GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
 987
 988       // Get a register for the selection result (possibly a new temporary one).
 989       Register SelectionResult = Call->ReturnRegister;
 990       if (PointerSize != ResultWidth) {
 991         SelectionResult =
 992             MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
 993         GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
 994                                   MIRBuilder.getMF());
 995       }
 996       // Create the final G_SELECT to return the extracted value or the default.
 997       MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,
 998                              DefaultRegister);
 999       ToTruncate = SelectionResult;
1000     } else {
1001       ToTruncate = Extracted;
1002     }
1003   }
1004   // Alter the result's bitwidth if it does not match the SizeT value extracted.
1005   if (PointerSize != ResultWidth)
1006     MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);
1007   return true;
1008 }
1009
1010 static bool generateBuiltinVar(const SPIRV::IncomingCall *Call,
1011                                MachineIRBuilder &MIRBuilder,
1012                                SPIRVGlobalRegistry *GR) {
1013   // Lookup the builtin variable record.
1014   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1015   SPIRV::BuiltIn::BuiltIn Value =
1016       SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
1017
1018   if (Value == SPIRV::BuiltIn::GlobalInvocationId)
1019     return genWorkgroupQuery(Call, MIRBuilder, GR, Value, 0);
1020
1021   // Build a load instruction for the builtin variable.
1022   unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);
1023   LLT LLType;
1024   if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
1025     LLType =
1026         LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth);
1027   else
1028     LLType = LLT::scalar(BitWidth);
1029
1030   return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value,
1031                                   LLType, Call->ReturnRegister);
1032 }
1033
1034 static bool generateAtomicInst(const SPIRV::IncomingCall *Call,
1035                                MachineIRBuilder &MIRBuilder,
1036                                SPIRVGlobalRegistry *GR) {
1037   // Lookup the instruction opcode in the TableGen records.
1038   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1039   unsigned Opcode =
1040       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1041
1042   switch (Opcode) {
1043   case SPIRV::OpStore:
1044     return buildAtomicInitInst(Call, MIRBuilder);
1045   case SPIRV::OpAtomicLoad:
1046     return buildAtomicLoadInst(Call, MIRBuilder, GR);
1047   case SPIRV::OpAtomicStore:
1048     return buildAtomicStoreInst(Call, MIRBuilder, GR);
1049   case SPIRV::OpAtomicCompareExchange:
1050   case SPIRV::OpAtomicCompareExchangeWeak:
1051     return buildAtomicCompareExchangeInst(Call, MIRBuilder, GR);
1052   case SPIRV::OpAtomicIAdd:
1053   case SPIRV::OpAtomicISub:
1054   case SPIRV::OpAtomicOr:
1055   case SPIRV::OpAtomicXor:
1056   case SPIRV::OpAtomicAnd:
1057   case SPIRV::OpAtomicExchange:
1058     return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
1059   case SPIRV::OpMemoryBarrier:
1060     return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
1061   case SPIRV::OpAtomicFlagTestAndSet:
1062   case SPIRV::OpAtomicFlagClear:
1063     return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);
1064   default:
1065     return false;
1066   }
1067 }
1068
1069 static bool generateBarrierInst(const SPIRV::IncomingCall *Call,
1070                                 MachineIRBuilder &MIRBuilder,
1071                                 SPIRVGlobalRegistry *GR) {
1072   // Lookup the instruction opcode in the TableGen records.
1073   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1074   unsigned Opcode =
1075       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1076
1077   return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
1078 }
1079
1080 static bool generateDotOrFMulInst(const SPIRV::IncomingCall *Call,
1081                                   MachineIRBuilder &MIRBuilder,
1082                                   SPIRVGlobalRegistry *GR) {
1083   unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();
1084   bool IsVec = Opcode == SPIRV::OpTypeVector;
1085   // Use OpDot only in case of vector args and OpFMul in case of scalar args.
1086   MIRBuilder.buildInstr(IsVec ? SPIRV::OpDot : SPIRV::OpFMulS)
1087       .addDef(Call->ReturnRegister)
1088       .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1089       .addUse(Call->Arguments[0])
1090       .addUse(Call->Arguments[1]);
1091   return true;
1092 }
1093
1094 static bool generateGetQueryInst(const SPIRV::IncomingCall *Call,
1095                                  MachineIRBuilder &MIRBuilder,
1096                                  SPIRVGlobalRegistry *GR) {
1097   // Lookup the builtin record.
1098   SPIRV::BuiltIn::BuiltIn Value =
1099       SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;
1100   uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||
1101                         Value == SPIRV::BuiltIn::WorkgroupSize ||
1102                         Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
1103   return genWorkgroupQuery(Call, MIRBuilder, GR, Value, IsDefault ? 1 : 0);
1104 }
1105
1106 static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call,
1107                                        MachineIRBuilder &MIRBuilder,
1108                                        SPIRVGlobalRegistry *GR) {
1109   // Lookup the image size query component number in the TableGen records.
1110   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1111   uint32_t Component =
1112       SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;
1113   // Query result may either be a vector or a scalar. If return type is not a
1114   // vector, expect only a single size component. Otherwise get the number of
1115   // expected components.
1116   SPIRVType *RetTy = Call->ReturnType;
1117   unsigned NumExpectedRetComponents = RetTy->getOpcode() == SPIRV::OpTypeVector
1118                                           ? RetTy->getOperand(2).getImm()
1119                                           : 1;
1120   // Get the actual number of query result/size components.
1121   SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1122   unsigned NumActualRetComponents = getNumSizeComponents(ImgType);
1123   Register QueryResult = Call->ReturnRegister;
1124   SPIRVType *QueryResultType = Call->ReturnType;
1125   if (NumExpectedRetComponents != NumActualRetComponents) {
1126     QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
1127         LLT::fixed_vector(NumActualRetComponents, 32));
1128     SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1129     QueryResultType = GR->getOrCreateSPIRVVectorType(
1130         IntTy, NumActualRetComponents, MIRBuilder);
1131     GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());
1132   }
1133   bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;
1134   unsigned Opcode =
1135       IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
1136   auto MIB = MIRBuilder.buildInstr(Opcode)
1137                  .addDef(QueryResult)
1138                  .addUse(GR->getSPIRVTypeID(QueryResultType))
1139                  .addUse(Call->Arguments[0]);
1140   if (!IsDimBuf)
1141     MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Lod id.
1142   if (NumExpectedRetComponents == NumActualRetComponents)
1143     return true;
1144   if (NumExpectedRetComponents == 1) {
1145     // Only 1 component is expected, build OpCompositeExtract instruction.
1146     unsigned ExtractedComposite =
1147         Component == 3 ? NumActualRetComponents - 1 : Component;
1148     assert(ExtractedComposite < NumActualRetComponents &&
1149            "Invalid composite index!");
1150     MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1151         .addDef(Call->ReturnRegister)
1152         .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1153         .addUse(QueryResult)
1154         .addImm(ExtractedComposite);
1155   } else {
1156     // More than 1 component is expected, fill a new vector.
1157     auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)
1158                    .addDef(Call->ReturnRegister)
1159                    .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1160                    .addUse(QueryResult)
1161                    .addUse(QueryResult);
1162     for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
1163       MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);
1164   }
1165   return true;
1166 }
1167
1168 static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call,
1169                                        MachineIRBuilder &MIRBuilder,
1170                                        SPIRVGlobalRegistry *GR) {
1171   assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
1172          "Image samples query result must be of int type!");
1173
1174   // Lookup the instruction opcode in the TableGen records.
1175   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1176   unsigned Opcode =
1177       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1178
1179   Register Image = Call->Arguments[0];
1180   SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
1181       GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
1182
1183   switch (Opcode) {
1184   case SPIRV::OpImageQuerySamples:
1185     assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
1186            "Image must be of 2D dimensionality");
1187     break;
1188   case SPIRV::OpImageQueryLevels:
1189     assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
1190             ImageDimensionality == SPIRV::Dim::DIM_2D ||
1191             ImageDimensionality == SPIRV::Dim::DIM_3D ||
1192             ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
1193            "Image must be of 1D/2D/3D/Cube dimensionality");
1194     break;
1195   }
1196
1197   MIRBuilder.buildInstr(Opcode)
1198       .addDef(Call->ReturnRegister)
1199       .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1200       .addUse(Image);
1201   return true;
1202 }
1203
1204 // TODO: Move to TableGen.
1205 static SPIRV::SamplerAddressingMode::SamplerAddressingMode
1206 getSamplerAddressingModeFromBitmask(unsigned Bitmask) {
1207   switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
1208   case SPIRV::CLK_ADDRESS_CLAMP:
1209     return SPIRV::SamplerAddressingMode::Clamp;
1210   case SPIRV::CLK_ADDRESS_CLAMP_TO_EDGE:
1211     return SPIRV::SamplerAddressingMode::ClampToEdge;
1212   case SPIRV::CLK_ADDRESS_REPEAT:
1213     return SPIRV::SamplerAddressingMode::Repeat;
1214   case SPIRV::CLK_ADDRESS_MIRRORED_REPEAT:
1215     return SPIRV::SamplerAddressingMode::RepeatMirrored;
1216   case SPIRV::CLK_ADDRESS_NONE:
1217     return SPIRV::SamplerAddressingMode::None;
1218   default:
1219     llvm_unreachable("Unknown CL address mode");
1220   }
1221 }
1222
1223 static unsigned getSamplerParamFromBitmask(unsigned Bitmask) {
1224   return (Bitmask & SPIRV::CLK_NORMALIZED_COORDS_TRUE) ? 1 : 0;
1225 }
1226
1227 static SPIRV::SamplerFilterMode::SamplerFilterMode
1228 getSamplerFilterModeFromBitmask(unsigned Bitmask) {
1229   if (Bitmask & SPIRV::CLK_FILTER_LINEAR)
1230     return SPIRV::SamplerFilterMode::Linear;
1231   if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
1232     return SPIRV::SamplerFilterMode::Nearest;
1233   return SPIRV::SamplerFilterMode::Nearest;
1234 }
1235
1236 static bool generateReadImageInst(const StringRef DemangledCall,
1237                                   const SPIRV::IncomingCall *Call,
1238                                   MachineIRBuilder &MIRBuilder,
1239                                   SPIRVGlobalRegistry *GR) {
1240   Register Image = Call->Arguments[0];
1241   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1242
1243   if (DemangledCall.contains_insensitive("ocl_sampler")) {
1244     Register Sampler = Call->Arguments[1];
1245
1246     if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
1247         getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) {
1248       uint64_t SamplerMask = getIConstVal(Sampler, MRI);
1249       Sampler = GR->buildConstantSampler(
1250           Register(), getSamplerAddressingModeFromBitmask(SamplerMask),
1251           getSamplerParamFromBitmask(SamplerMask),
1252           getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder,
1253           GR->getSPIRVTypeForVReg(Sampler));
1254     }
1255     SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1256     SPIRVType *SampledImageType =
1257         GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1258     Register SampledImage = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1259
1260     MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1261         .addDef(SampledImage)
1262         .addUse(GR->getSPIRVTypeID(SampledImageType))
1263         .addUse(Image)
1264         .addUse(Sampler);
1265
1266     Register Lod = GR->buildConstantFP(APFloat::getZero(APFloat::IEEEsingle()),
1267                                        MIRBuilder);
1268     SPIRVType *TempType = Call->ReturnType;
1269     bool NeedsExtraction = false;
1270     if (TempType->getOpcode() != SPIRV::OpTypeVector) {
1271       TempType =
1272           GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);
1273       NeedsExtraction = true;
1274     }
1275     LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(TempType));
1276     Register TempRegister = MRI->createGenericVirtualRegister(LLType);
1277     GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
1278
1279     MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1280         .addDef(NeedsExtraction ? TempRegister : Call->ReturnRegister)
1281         .addUse(GR->getSPIRVTypeID(TempType))
1282         .addUse(SampledImage)
1283         .addUse(Call->Arguments[2]) // Coordinate.
1284         .addImm(SPIRV::ImageOperand::Lod)
1285         .addUse(Lod);
1286
1287     if (NeedsExtraction)
1288       MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
1289           .addDef(Call->ReturnRegister)
1290           .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1291           .addUse(TempRegister)
1292           .addImm(0);
1293   } else if (DemangledCall.contains_insensitive("msaa")) {
1294     MIRBuilder.buildInstr(SPIRV::OpImageRead)
1295         .addDef(Call->ReturnRegister)
1296         .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1297         .addUse(Image)
1298         .addUse(Call->Arguments[1]) // Coordinate.
1299         .addImm(SPIRV::ImageOperand::Sample)
1300         .addUse(Call->Arguments[2]);
1301   } else {
1302     MIRBuilder.buildInstr(SPIRV::OpImageRead)
1303         .addDef(Call->ReturnRegister)
1304         .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1305         .addUse(Image)
1306         .addUse(Call->Arguments[1]); // Coordinate.
1307   }
1308   return true;
1309 }
1310
1311 static bool generateWriteImageInst(const SPIRV::IncomingCall *Call,
1312                                    MachineIRBuilder &MIRBuilder,
1313                                    SPIRVGlobalRegistry *GR) {
1314   MIRBuilder.buildInstr(SPIRV::OpImageWrite)
1315       .addUse(Call->Arguments[0])  // Image.
1316       .addUse(Call->Arguments[1])  // Coordinate.
1317       .addUse(Call->Arguments[2]); // Texel.
1318   return true;
1319 }
1320
1321 static bool generateSampleImageInst(const StringRef DemangledCall,
1322                                     const SPIRV::IncomingCall *Call,
1323                                     MachineIRBuilder &MIRBuilder,
1324                                     SPIRVGlobalRegistry *GR) {
1325   if (Call->Builtin->Name.contains_insensitive(
1326           "__translate_sampler_initializer")) {
1327     // Build sampler literal.
1328     uint64_t Bitmask = getIConstVal(Call->Arguments[0], MIRBuilder.getMRI());
1329     Register Sampler = GR->buildConstantSampler(
1330         Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask),
1331         getSamplerParamFromBitmask(Bitmask),
1332         getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType);
1333     return Sampler.isValid();
1334   } else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {
1335     // Create OpSampledImage.
1336     Register Image = Call->Arguments[0];
1337     SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
1338     SPIRVType *SampledImageType =
1339         GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
1340     Register SampledImage =
1341         Call->ReturnRegister.isValid()
1342             ? Call->ReturnRegister
1343             : MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
1344     MIRBuilder.buildInstr(SPIRV::OpSampledImage)
1345         .addDef(SampledImage)
1346         .addUse(GR->getSPIRVTypeID(SampledImageType))
1347         .addUse(Image)
1348         .addUse(Call->Arguments[1]); // Sampler.
1349     return true;
1350   } else if (Call->Builtin->Name.contains_insensitive(
1351                  "__spirv_ImageSampleExplicitLod")) {
1352     // Sample an image using an explicit level of detail.
1353     std::string ReturnType = DemangledCall.str();
1354     if (DemangledCall.contains("_R")) {
1355       ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);
1356       ReturnType = ReturnType.substr(0, ReturnType.find('('));
1357     }
1358     SPIRVType *Type = GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder);
1359     MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
1360         .addDef(Call->ReturnRegister)
1361         .addUse(GR->getSPIRVTypeID(Type))
1362         .addUse(Call->Arguments[0]) // Image.
1363         .addUse(Call->Arguments[1]) // Coordinate.
1364         .addImm(SPIRV::ImageOperand::Lod)
1365         .addUse(Call->Arguments[3]);
1366     return true;
1367   }
1368   return false;
1369 }
1370
1371 static bool generateSelectInst(const SPIRV::IncomingCall *Call,
1372                                MachineIRBuilder &MIRBuilder) {
1373   MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],
1374                          Call->Arguments[1], Call->Arguments[2]);
1375   return true;
1376 }
1377
1378 static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call,
1379                                      MachineIRBuilder &MIRBuilder,
1380                                      SPIRVGlobalRegistry *GR) {
1381   // Lookup the instruction opcode in the TableGen records.
1382   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1383   unsigned Opcode =
1384       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1385   const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1386
1387   switch (Opcode) {
1388   case SPIRV::OpSpecConstant: {
1389     // Build the SpecID decoration.
1390     unsigned SpecId =
1391         static_cast<unsigned>(getIConstVal(Call->Arguments[0], MRI));
1392     buildOpDecorate(Call->ReturnRegister, MIRBuilder, SPIRV::Decoration::SpecId,
1393                     {SpecId});
1394     // Determine the constant MI.
1395     Register ConstRegister = Call->Arguments[1];
1396     const MachineInstr *Const = getDefInstrMaybeConstant(ConstRegister, MRI);
1397     assert(Const &&
1398            (Const->getOpcode() == TargetOpcode::G_CONSTANT ||
1399             Const->getOpcode() == TargetOpcode::G_FCONSTANT) &&
1400            "Argument should be either an int or floating-point constant");
1401     // Determine the opcode and built the OpSpec MI.
1402     const MachineOperand &ConstOperand = Const->getOperand(1);
1403     if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
1404       assert(ConstOperand.isCImm() && "Int constant operand is expected");
1405       Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
1406                    ? SPIRV::OpSpecConstantTrue
1407                    : SPIRV::OpSpecConstantFalse;
1408     }
1409     auto MIB = MIRBuilder.buildInstr(Opcode)
1410                    .addDef(Call->ReturnRegister)
1411                    .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1412
1413     if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
1414       if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
1415         addNumImm(ConstOperand.getCImm()->getValue(), MIB);
1416       else
1417         addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
1418     }
1419     return true;
1420   }
1421   case SPIRV::OpSpecConstantComposite: {
1422     auto MIB = MIRBuilder.buildInstr(Opcode)
1423                    .addDef(Call->ReturnRegister)
1424                    .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1425     for (unsigned i = 0; i < Call->Arguments.size(); i++)
1426       MIB.addUse(Call->Arguments[i]);
1427     return true;
1428   }
1429   default:
1430     return false;
1431   }
1432 }
1433
1434 static MachineInstr *getBlockStructInstr(Register ParamReg,
1435                                          MachineRegisterInfo *MRI) {
1436   // We expect the following sequence of instructions:
1437   //   %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
1438   //   or       = G_GLOBAL_VALUE @block_literal_global
1439   //   %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0
1440   //   %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
1441   MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
1442   assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
1443          MI->getOperand(1).isReg());
1444   Register BitcastReg = MI->getOperand(1).getReg();
1445   MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg);
1446   assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&
1447          BitcastMI->getOperand(2).isReg());
1448   Register ValueReg = BitcastMI->getOperand(2).getReg();
1449   MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
1450   return ValueMI;
1451 }
1452
1453 // Return an integer constant corresponding to the given register and
1454 // defined in spv_track_constant.
1455 // TODO: maybe unify with prelegalizer pass.
1456 static unsigned getConstFromIntrinsic(Register Reg, MachineRegisterInfo *MRI) {
1457   MachineInstr *DefMI = MRI->getUniqueVRegDef(Reg);
1458   assert(isSpvIntrinsic(*DefMI, Intrinsic::spv_track_constant) &&
1459          DefMI->getOperand(2).isReg());
1460   MachineInstr *DefMI2 = MRI->getUniqueVRegDef(DefMI->getOperand(2).getReg());
1461   assert(DefMI2->getOpcode() == TargetOpcode::G_CONSTANT &&
1462          DefMI2->getOperand(1).isCImm());
1463   return DefMI2->getOperand(1).getCImm()->getValue().getZExtValue();
1464 }
1465
1466 // Return type of the instruction result from spv_assign_type intrinsic.
1467 // TODO: maybe unify with prelegalizer pass.
1468 static const Type *getMachineInstrType(MachineInstr *MI) {
1469   MachineInstr *NextMI = MI->getNextNode();
1470   if (isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_name))
1471     NextMI = NextMI->getNextNode();
1472   Register ValueReg = MI->getOperand(0).getReg();
1473   if (!isSpvIntrinsic(*NextMI, Intrinsic::spv_assign_type) ||
1474       NextMI->getOperand(1).getReg() != ValueReg)
1475     return nullptr;
1476   Type *Ty = getMDOperandAsType(NextMI->getOperand(2).getMetadata(), 0);
1477   assert(Ty && "Type is expected");
1478   return getTypedPtrEltType(Ty);
1479 }
1480
1481 static const Type *getBlockStructType(Register ParamReg,
1482                                       MachineRegisterInfo *MRI) {
1483   // In principle, this information should be passed to us from Clang via
1484   // an elementtype attribute. However, said attribute requires that
1485   // the function call be an intrinsic, which is not. Instead, we rely on being
1486   // able to trace this to the declaration of a variable: OpenCL C specification
1487   // section 6.12.5 should guarantee that we can do this.
1488   MachineInstr *MI = getBlockStructInstr(ParamReg, MRI);
1489   if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
1490     return getTypedPtrEltType(MI->getOperand(1).getGlobal()->getType());
1491   assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
1492          "Blocks in OpenCL C must be traceable to allocation site");
1493   return getMachineInstrType(MI);
1494 }
1495
1496 // TODO: maybe move to the global register.
1497 static SPIRVType *
1498 getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder,
1499                                    SPIRVGlobalRegistry *GR) {
1500   LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
1501   Type *OpaqueType = StructType::getTypeByName(Context, "spirv.DeviceEvent");
1502   if (!OpaqueType)
1503     OpaqueType = StructType::getTypeByName(Context, "opencl.clk_event_t");
1504   if (!OpaqueType)
1505     OpaqueType = StructType::create(Context, "spirv.DeviceEvent");
1506   unsigned SC0 = storageClassToAddressSpace(SPIRV::StorageClass::Function);
1507   unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1508   Type *PtrType = PointerType::get(PointerType::get(OpaqueType, SC0), SC1);
1509   return GR->getOrCreateSPIRVType(PtrType, MIRBuilder);
1510 }
1511
1512 static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
1513                                MachineIRBuilder &MIRBuilder,
1514                                SPIRVGlobalRegistry *GR) {
1515   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1516   const DataLayout &DL = MIRBuilder.getDataLayout();
1517   bool HasEvents = Call->Builtin->Name.find("events") != StringRef::npos;
1518   const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
1519
1520   // Make vararg instructions before OpEnqueueKernel.
1521   // Local sizes arguments: Sizes of block invoke arguments. Clang generates
1522   // local size operands as an array, so we need to unpack them.
1523   SmallVector<Register, 16> LocalSizes;
1524   if (Call->Builtin->Name.find("_varargs") != StringRef::npos) {
1525     const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
1526     Register GepReg = Call->Arguments[LocalSizeArrayIdx];
1527     MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
1528     assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
1529            GepMI->getOperand(3).isReg());
1530     Register ArrayReg = GepMI->getOperand(3).getReg();
1531     MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
1532     const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
1533     assert(LocalSizeTy && "Local size type is expected");
1534     const uint64_t LocalSizeNum =
1535         cast<ArrayType>(LocalSizeTy)->getNumElements();
1536     unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
1537     const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
1538     const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(
1539         Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
1540     for (unsigned I = 0; I < LocalSizeNum; ++I) {
1541       Register Reg =
1542           MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
1543       MIRBuilder.getMRI()->setType(Reg, LLType);
1544       GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
1545       auto GEPInst = MIRBuilder.buildIntrinsic(Intrinsic::spv_gep,
1546                                                ArrayRef<Register>{Reg}, true);
1547       GEPInst
1548           .addImm(GepMI->getOperand(2).getImm())          // In bound.
1549           .addUse(ArrayMI->getOperand(0).getReg())        // Alloca.
1550           .addUse(buildConstantIntReg(0, MIRBuilder, GR)) // Indices.
1551           .addUse(buildConstantIntReg(I, MIRBuilder, GR));
1552       LocalSizes.push_back(Reg);
1553     }
1554   }
1555
1556   // SPIRV OpEnqueueKernel instruction has 10+ arguments.
1557   auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
1558                  .addDef(Call->ReturnRegister)
1559                  .addUse(GR->getSPIRVTypeID(Int32Ty));
1560
1561   // Copy all arguments before block invoke function pointer.
1562   const unsigned BlockFIdx = HasEvents ? 6 : 3;
1563   for (unsigned i = 0; i < BlockFIdx; i++)
1564     MIB.addUse(Call->Arguments[i]);
1565
1566   // If there are no event arguments in the original call, add dummy ones.
1567   if (!HasEvents) {
1568     MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Dummy num events.
1569     Register NullPtr = GR->getOrCreateConstNullPtr(
1570         MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
1571     MIB.addUse(NullPtr); // Dummy wait events.
1572     MIB.addUse(NullPtr); // Dummy ret event.
1573   }
1574
1575   MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
1576   assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
1577   // Invoke: Pointer to invoke function.
1578   MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
1579
1580   Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
1581   // Param: Pointer to block literal.
1582   MIB.addUse(BlockLiteralReg);
1583
1584   Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
1585   // TODO: these numbers should be obtained from block literal structure.
1586   // Param Size: Size of block literal structure.
1587   MIB.addUse(buildConstantIntReg(DL.getTypeStoreSize(PType), MIRBuilder, GR));
1588   // Param Aligment: Aligment of block literal structure.
1589   MIB.addUse(
1590       buildConstantIntReg(DL.getPrefTypeAlignment(PType), MIRBuilder, GR));
1591
1592   for (unsigned i = 0; i < LocalSizes.size(); i++)
1593     MIB.addUse(LocalSizes[i]);
1594   return true;
1595 }
1596
1597 static bool generateEnqueueInst(const SPIRV::IncomingCall *Call,
1598                                 MachineIRBuilder &MIRBuilder,
1599                                 SPIRVGlobalRegistry *GR) {
1600   // Lookup the instruction opcode in the TableGen records.
1601   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1602   unsigned Opcode =
1603       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1604
1605   switch (Opcode) {
1606   case SPIRV::OpRetainEvent:
1607   case SPIRV::OpReleaseEvent:
1608     return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);
1609   case SPIRV::OpCreateUserEvent:
1610   case SPIRV::OpGetDefaultQueue:
1611     return MIRBuilder.buildInstr(Opcode)
1612         .addDef(Call->ReturnRegister)
1613         .addUse(GR->getSPIRVTypeID(Call->ReturnType));
1614   case SPIRV::OpIsValidEvent:
1615     return MIRBuilder.buildInstr(Opcode)
1616         .addDef(Call->ReturnRegister)
1617         .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1618         .addUse(Call->Arguments[0]);
1619   case SPIRV::OpSetUserEventStatus:
1620     return MIRBuilder.buildInstr(Opcode)
1621         .addUse(Call->Arguments[0])
1622         .addUse(Call->Arguments[1]);
1623   case SPIRV::OpCaptureEventProfilingInfo:
1624     return MIRBuilder.buildInstr(Opcode)
1625         .addUse(Call->Arguments[0])
1626         .addUse(Call->Arguments[1])
1627         .addUse(Call->Arguments[2]);
1628   case SPIRV::OpBuildNDRange: {
1629     MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1630     SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
1631     assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&
1632            PtrType->getOperand(2).isReg());
1633     Register TypeReg = PtrType->getOperand(2).getReg();
1634     SPIRVType *StructType = GR->getSPIRVTypeForVReg(TypeReg);
1635     Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1636     GR->assignSPIRVTypeToVReg(StructType, TmpReg, MIRBuilder.getMF());
1637     // Skip the first arg, it's the destination pointer. OpBuildNDRange takes
1638     // three other arguments, so pass zero constant on absence.
1639     unsigned NumArgs = Call->Arguments.size();
1640     assert(NumArgs >= 2);
1641     Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];
1642     Register LocalWorkSize =
1643         NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];
1644     Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1];
1645     if (NumArgs < 4) {
1646       Register Const;
1647       SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);
1648       if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {
1649         MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize);
1650         assert(DefInstr && isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&
1651                DefInstr->getOperand(3).isReg());
1652         Register GWSPtr = DefInstr->getOperand(3).getReg();
1653         // TODO: Maybe simplify generation of the type of the fields.
1654         unsigned Size = Call->Builtin->Name.equals("ndrange_3D") ? 3 : 2;
1655         unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32;
1656         Type *BaseTy = IntegerType::get(
1657             MIRBuilder.getMF().getFunction().getContext(), BitWidth);
1658         Type *FieldTy = ArrayType::get(BaseTy, Size);
1659         SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder);
1660         GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::IDRegClass);
1661         GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize,
1662                                   MIRBuilder.getMF());
1663         MIRBuilder.buildInstr(SPIRV::OpLoad)
1664             .addDef(GlobalWorkSize)
1665             .addUse(GR->getSPIRVTypeID(SpvFieldTy))
1666             .addUse(GWSPtr);
1667         Const = GR->getOrCreateConsIntArray(0, MIRBuilder, SpvFieldTy);
1668       } else {
1669         Const = GR->buildConstantInt(0, MIRBuilder, SpvTy);
1670       }
1671       if (!LocalWorkSize.isValid())
1672         LocalWorkSize = Const;
1673       if (!GlobalWorkOffset.isValid())
1674         GlobalWorkOffset = Const;
1675     }
1676     MIRBuilder.buildInstr(Opcode)
1677         .addDef(TmpReg)
1678         .addUse(TypeReg)
1679         .addUse(GlobalWorkSize)
1680         .addUse(LocalWorkSize)
1681         .addUse(GlobalWorkOffset);
1682     return MIRBuilder.buildInstr(SPIRV::OpStore)
1683         .addUse(Call->Arguments[0])
1684         .addUse(TmpReg);
1685   }
1686   case SPIRV::OpEnqueueKernel:
1687     return buildEnqueueKernel(Call, MIRBuilder, GR);
1688   default:
1689     return false;
1690   }
1691 }
1692
1693 static bool generateAsyncCopy(const SPIRV::IncomingCall *Call,
1694                               MachineIRBuilder &MIRBuilder,
1695                               SPIRVGlobalRegistry *GR) {
1696   // Lookup the instruction opcode in the TableGen records.
1697   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1698   unsigned Opcode =
1699       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1700   auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GR);
1701
1702   switch (Opcode) {
1703   case SPIRV::OpGroupAsyncCopy:
1704     return MIRBuilder.buildInstr(Opcode)
1705         .addDef(Call->ReturnRegister)
1706         .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1707         .addUse(Scope)
1708         .addUse(Call->Arguments[0])
1709         .addUse(Call->Arguments[1])
1710         .addUse(Call->Arguments[2])
1711         .addUse(buildConstantIntReg(1, MIRBuilder, GR))
1712         .addUse(Call->Arguments[3]);
1713   case SPIRV::OpGroupWaitEvents:
1714     return MIRBuilder.buildInstr(Opcode)
1715         .addUse(Scope)
1716         .addUse(Call->Arguments[0])
1717         .addUse(Call->Arguments[1]);
1718   default:
1719     return false;
1720   }
1721 }
1722
1723 static bool generateConvertInst(const StringRef DemangledCall,
1724                                 const SPIRV::IncomingCall *Call,
1725                                 MachineIRBuilder &MIRBuilder,
1726                                 SPIRVGlobalRegistry *GR) {
1727   // Lookup the conversion builtin in the TableGen records.
1728   const SPIRV::ConvertBuiltin *Builtin =
1729       SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);
1730
1731   if (Builtin->IsSaturated)
1732     buildOpDecorate(Call->ReturnRegister, MIRBuilder,
1733                     SPIRV::Decoration::SaturatedConversion, {});
1734   if (Builtin->IsRounded)
1735     buildOpDecorate(Call->ReturnRegister, MIRBuilder,
1736                     SPIRV::Decoration::FPRoundingMode, {Builtin->RoundingMode});
1737
1738   unsigned Opcode = SPIRV::OpNop;
1739   if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
1740     // Int -> ...
1741     if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
1742       // Int -> Int
1743       if (Builtin->IsSaturated)
1744         Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
1745                                               : SPIRV::OpSatConvertSToU;
1746       else
1747         Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert
1748                                               : SPIRV::OpSConvert;
1749     } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
1750                                           SPIRV::OpTypeFloat)) {
1751       // Int -> Float
1752       bool IsSourceSigned =
1753           DemangledCall[DemangledCall.find_first_of('(') + 1] != 'u';
1754       Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
1755     }
1756   } else if (GR->isScalarOrVectorOfType(Call->Arguments[0],
1757                                         SPIRV::OpTypeFloat)) {
1758     // Float -> ...
1759     if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt))
1760       // Float -> Int
1761       Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
1762                                             : SPIRV::OpConvertFToU;
1763     else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
1764                                         SPIRV::OpTypeFloat))
1765       // Float -> Float
1766       Opcode = SPIRV::OpFConvert;
1767   }
1768
1769   assert(Opcode != SPIRV::OpNop &&
1770          "Conversion between the types not implemented!");
1771
1772   MIRBuilder.buildInstr(Opcode)
1773       .addDef(Call->ReturnRegister)
1774       .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1775       .addUse(Call->Arguments[0]);
1776   return true;
1777 }
1778
1779 static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call,
1780                                         MachineIRBuilder &MIRBuilder,
1781                                         SPIRVGlobalRegistry *GR) {
1782   // Lookup the vector load/store builtin in the TableGen records.
1783   const SPIRV::VectorLoadStoreBuiltin *Builtin =
1784       SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
1785                                           Call->Builtin->Set);
1786   // Build extended instruction.
1787   auto MIB =
1788       MIRBuilder.buildInstr(SPIRV::OpExtInst)
1789           .addDef(Call->ReturnRegister)
1790           .addUse(GR->getSPIRVTypeID(Call->ReturnType))
1791           .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
1792           .addImm(Builtin->Number);
1793   for (auto Argument : Call->Arguments)
1794     MIB.addUse(Argument);
1795
1796   // Rounding mode should be passed as a last argument in the MI for builtins
1797   // like "vstorea_halfn_r".
1798   if (Builtin->IsRounded)
1799     MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));
1800   return true;
1801 }
1802
1803 static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call,
1804                                   MachineIRBuilder &MIRBuilder,
1805                                   SPIRVGlobalRegistry *GR) {
1806   // Lookup the instruction opcode in the TableGen records.
1807   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
1808   unsigned Opcode =
1809       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
1810   bool IsLoad = Opcode == SPIRV::OpLoad;
1811   // Build the instruction.
1812   auto MIB = MIRBuilder.buildInstr(Opcode);
1813   if (IsLoad) {
1814     MIB.addDef(Call->ReturnRegister);
1815     MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType));
1816   }
1817   // Add a pointer to the value to load/store.
1818   MIB.addUse(Call->Arguments[0]);
1819   // Add a value to store.
1820   if (!IsLoad)
1821     MIB.addUse(Call->Arguments[1]);
1822   // Add optional memory attributes and an alignment.
1823   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
1824   unsigned NumArgs = Call->Arguments.size();
1825   if ((IsLoad && NumArgs >= 2) || NumArgs >= 3)
1826     MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 1 : 2], MRI));
1827   if ((IsLoad && NumArgs >= 3) || NumArgs >= 4)
1828     MIB.addImm(getConstFromIntrinsic(Call->Arguments[IsLoad ? 2 : 3], MRI));
1829   return true;
1830 }
1831
1832 /// Lowers a builtin funtion call using the provided \p DemangledCall skeleton
1833 /// and external instruction \p Set.
1834 namespace SPIRV {
1835 std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
1836                                  SPIRV::InstructionSet::InstructionSet Set,
1837                                  MachineIRBuilder &MIRBuilder,
1838                                  const Register OrigRet, const Type *OrigRetTy,
1839                                  const SmallVectorImpl<Register> &Args,
1840                                  SPIRVGlobalRegistry *GR) {
1841   LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
1842
1843   // SPIR-V type and return register.
1844   Register ReturnRegister = OrigRet;
1845   SPIRVType *ReturnType = nullptr;
1846   if (OrigRetTy && !OrigRetTy->isVoidTy()) {
1847     ReturnType = GR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder);
1848   } else if (OrigRetTy && OrigRetTy->isVoidTy()) {
1849     ReturnRegister = MIRBuilder.getMRI()->createVirtualRegister(&IDRegClass);
1850     MIRBuilder.getMRI()->setType(ReturnRegister, LLT::scalar(32));
1851     ReturnType = GR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder);
1852   }
1853
1854   // Lookup the builtin in the TableGen records.
1855   std::unique_ptr<const IncomingCall> Call =
1856       lookupBuiltin(DemangledCall, Set, ReturnRegister, ReturnType, Args);
1857
1858   if (!Call) {
1859     LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
1860     return std::nullopt;
1861   }
1862
1863   // TODO: check if the provided args meet the builtin requirments.
1864   assert(Args.size() >= Call->Builtin->MinNumArgs &&
1865          "Too few arguments to generate the builtin");
1866   if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)
1867     LLVM_DEBUG(dbgs() << "More arguments provided than required!\n");
1868
1869   // Match the builtin with implementation based on the grouping.
1870   switch (Call->Builtin->Group) {
1871   case SPIRV::Extended:
1872     return generateExtInst(Call.get(), MIRBuilder, GR);
1873   case SPIRV::Relational:
1874     return generateRelationalInst(Call.get(), MIRBuilder, GR);
1875   case SPIRV::Group:
1876     return generateGroupInst(Call.get(), MIRBuilder, GR);
1877   case SPIRV::Variable:
1878     return generateBuiltinVar(Call.get(), MIRBuilder, GR);
1879   case SPIRV::Atomic:
1880     return generateAtomicInst(Call.get(), MIRBuilder, GR);
1881   case SPIRV::Barrier:
1882     return generateBarrierInst(Call.get(), MIRBuilder, GR);
1883   case SPIRV::Dot:
1884     return generateDotOrFMulInst(Call.get(), MIRBuilder, GR);
1885   case SPIRV::GetQuery:
1886     return generateGetQueryInst(Call.get(), MIRBuilder, GR);
1887   case SPIRV::ImageSizeQuery:
1888     return generateImageSizeQueryInst(Call.get(), MIRBuilder, GR);
1889   case SPIRV::ImageMiscQuery:
1890     return generateImageMiscQueryInst(Call.get(), MIRBuilder, GR);
1891   case SPIRV::ReadImage:
1892     return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
1893   case SPIRV::WriteImage:
1894     return generateWriteImageInst(Call.get(), MIRBuilder, GR);
1895   case SPIRV::SampleImage:
1896     return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
1897   case SPIRV::Select:
1898     return generateSelectInst(Call.get(), MIRBuilder);
1899   case SPIRV::SpecConstant:
1900     return generateSpecConstantInst(Call.get(), MIRBuilder, GR);
1901   case SPIRV::Enqueue:
1902     return generateEnqueueInst(Call.get(), MIRBuilder, GR);
1903   case SPIRV::AsyncCopy:
1904     return generateAsyncCopy(Call.get(), MIRBuilder, GR);
1905   case SPIRV::Convert:
1906     return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);
1907   case SPIRV::VectorLoadStore:
1908     return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);
1909   case SPIRV::LoadStore:
1910     return generateLoadStoreInst(Call.get(), MIRBuilder, GR);
1911   }
1912   return false;
1913 }
1914
1915 struct DemangledType {
1916   StringRef Name;
1917   uint32_t Opcode;
1918 };
1919
1920 #define GET_DemangledTypes_DECL
1921 #define GET_DemangledTypes_IMPL
1922
1923 struct ImageType {
1924   StringRef Name;
1925   StringRef SampledType;
1926   AccessQualifier::AccessQualifier Qualifier;
1927   Dim::Dim Dimensionality;
1928   bool Arrayed;
1929   bool Depth;
1930   bool Multisampled;
1931   bool Sampled;
1932   ImageFormat::ImageFormat Format;
1933 };
1934
1935 struct PipeType {
1936   StringRef Name;
1937   AccessQualifier::AccessQualifier Qualifier;
1938 };
1939
1940 using namespace AccessQualifier;
1941 using namespace Dim;
1942 using namespace ImageFormat;
1943 #define GET_ImageTypes_DECL
1944 #define GET_ImageTypes_IMPL
1945 #define GET_PipeTypes_DECL
1946 #define GET_PipeTypes_IMPL
1947 #include "SPIRVGenTables.inc"
1948 } // namespace SPIRV
1949
1950 //===----------------------------------------------------------------------===//
1951 // Misc functions for parsing builtin types and looking up implementation
1952 // details in TableGenerated tables.
1953 //===----------------------------------------------------------------------===//
1954
1955 static const SPIRV::DemangledType *findBuiltinType(StringRef Name) {
1956   if (Name.startswith("opencl."))
1957     return SPIRV::lookupBuiltinType(Name);
1958   if (!Name.startswith("spirv."))
1959     return nullptr;
1960   // Some SPIR-V builtin types have a complex list of parameters as part of
1961   // their name (e.g. spirv.Image._void_1_0_0_0_0_0_0). Those parameters often
1962   // are numeric literals which cannot be easily represented by TableGen
1963   // records and should be parsed instead.
1964   unsigned BaseTypeNameLength =
1965       Name.contains('_') ? Name.find('_') - 1 : Name.size();
1966   return SPIRV::lookupBuiltinType(Name.substr(0, BaseTypeNameLength).str());
1967 }
1968
1969 static std::unique_ptr<const SPIRV::ImageType>
1970 lookupOrParseBuiltinImageType(StringRef Name) {
1971   if (Name.startswith("opencl.")) {
1972     // Lookup OpenCL builtin image type lowering details in TableGen records.
1973     const SPIRV::ImageType *Record = SPIRV::lookupImageType(Name);
1974     return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType(*Record));
1975   }
1976   if (!Name.startswith("spirv."))
1977     llvm_unreachable("Unknown builtin image type name/literal");
1978   // Parse the literals of SPIR-V image builtin parameters. The name should
1979   // have the following format:
1980   // spirv.Image._Type_Dim_Depth_Arrayed_MS_Sampled_ImageFormat_AccessQualifier
1981   // e.g. %spirv.Image._void_1_0_0_0_0_0_0
1982   StringRef TypeParametersString = Name.substr(strlen("spirv.Image."));
1983   SmallVector<StringRef> TypeParameters;
1984   SplitString(TypeParametersString, TypeParameters, "_");
1985   assert(TypeParameters.size() == 8 &&
1986          "Wrong number of literals in SPIR-V builtin image type");
1987
1988   StringRef SampledType = TypeParameters[0];
1989   unsigned Dim, Depth, Arrayed, Multisampled, Sampled, Format, AccessQual;
1990   bool AreParameterLiteralsValid =
1991       !(TypeParameters[1].getAsInteger(10, Dim) ||
1992         TypeParameters[2].getAsInteger(10, Depth) ||
1993         TypeParameters[3].getAsInteger(10, Arrayed) ||
1994         TypeParameters[4].getAsInteger(10, Multisampled) ||
1995         TypeParameters[5].getAsInteger(10, Sampled) ||
1996         TypeParameters[6].getAsInteger(10, Format) ||
1997         TypeParameters[7].getAsInteger(10, AccessQual));
1998   assert(AreParameterLiteralsValid &&
1999          "Invalid format of SPIR-V image type parameter literals.");
2000
2001   return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType{
2002       Name, SampledType, SPIRV::AccessQualifier::AccessQualifier(AccessQual),
2003       SPIRV::Dim::Dim(Dim), static_cast<bool>(Arrayed),
2004       static_cast<bool>(Depth), static_cast<bool>(Multisampled),
2005       static_cast<bool>(Sampled), SPIRV::ImageFormat::ImageFormat(Format)});
2006 }
2007
2008 static std::unique_ptr<const SPIRV::PipeType>
2009 lookupOrParseBuiltinPipeType(StringRef Name) {
2010   if (Name.startswith("opencl.")) {
2011     // Lookup OpenCL builtin pipe type lowering details in TableGen records.
2012     const SPIRV::PipeType *Record = SPIRV::lookupPipeType(Name);
2013     return std::unique_ptr<SPIRV::PipeType>(new SPIRV::PipeType(*Record));
2014   }
2015   if (!Name.startswith("spirv."))
2016     llvm_unreachable("Unknown builtin pipe type name/literal");
2017   // Parse the access qualifier literal in the name of the SPIR-V pipe type.
2018   // The name should have the following format:
2019   // spirv.Pipe._AccessQualifier
2020   // e.g. %spirv.Pipe._1
2021   if (Name.endswith("_0"))
2022     return std::unique_ptr<SPIRV::PipeType>(
2023         new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadOnly});
2024   if (Name.endswith("_1"))
2025     return std::unique_ptr<SPIRV::PipeType>(
2026         new SPIRV::PipeType{Name, SPIRV::AccessQualifier::WriteOnly});
2027   if (Name.endswith("_2"))
2028     return std::unique_ptr<SPIRV::PipeType>(
2029         new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadWrite});
2030   llvm_unreachable("Unknown pipe type access qualifier literal");
2031 }
2032
2033 //===----------------------------------------------------------------------===//
2034 // Implementation functions for builtin types.
2035 //===----------------------------------------------------------------------===//
2036
2037 static SPIRVType *getNonParametrizedType(const StructType *OpaqueType,
2038                                          const SPIRV::DemangledType *TypeRecord,
2039                                          MachineIRBuilder &MIRBuilder,
2040                                          SPIRVGlobalRegistry *GR) {
2041   unsigned Opcode = TypeRecord->Opcode;
2042   // Create or get an existing type from GlobalRegistry.
2043   return GR->getOrCreateOpTypeByOpcode(OpaqueType, MIRBuilder, Opcode);
2044 }
2045
2046 static SPIRVType *getSamplerType(MachineIRBuilder &MIRBuilder,
2047                                  SPIRVGlobalRegistry *GR) {
2048   // Create or get an existing type from GlobalRegistry.
2049   return GR->getOrCreateOpTypeSampler(MIRBuilder);
2050 }
2051
2052 static SPIRVType *getPipeType(const StructType *OpaqueType,
2053                               MachineIRBuilder &MIRBuilder,
2054                               SPIRVGlobalRegistry *GR) {
2055   // Lookup pipe type lowering details in TableGen records or parse the
2056   // name/literal for details.
2057   std::unique_ptr<const SPIRV::PipeType> Record =
2058       lookupOrParseBuiltinPipeType(OpaqueType->getName());
2059   // Create or get an existing type from GlobalRegistry.
2060   return GR->getOrCreateOpTypePipe(MIRBuilder, Record.get()->Qualifier);
2061 }
2062
2063 static SPIRVType *
2064 getImageType(const StructType *OpaqueType,
2065              SPIRV::AccessQualifier::AccessQualifier AccessQual,
2066              MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
2067   // Lookup image type lowering details in TableGen records or parse the
2068   // name/literal for details.
2069   std::unique_ptr<const SPIRV::ImageType> Record =
2070       lookupOrParseBuiltinImageType(OpaqueType->getName());
2071
2072   SPIRVType *SampledType =
2073       GR->getOrCreateSPIRVTypeByName(Record.get()->SampledType, MIRBuilder);
2074   return GR->getOrCreateOpTypeImage(
2075       MIRBuilder, SampledType, Record.get()->Dimensionality,
2076       Record.get()->Depth, Record.get()->Arrayed, Record.get()->Multisampled,
2077       Record.get()->Sampled, Record.get()->Format,
2078       AccessQual == SPIRV::AccessQualifier::WriteOnly
2079           ? SPIRV::AccessQualifier::WriteOnly
2080           : Record.get()->Qualifier);
2081 }
2082
2083 static SPIRVType *getSampledImageType(const StructType *OpaqueType,
2084                                       MachineIRBuilder &MIRBuilder,
2085                                       SPIRVGlobalRegistry *GR) {
2086   StringRef TypeParametersString =
2087       OpaqueType->getName().substr(strlen("spirv.SampledImage."));
2088   LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
2089   Type *ImageOpaqueType = StructType::getTypeByName(
2090       Context, "spirv.Image." + TypeParametersString.str());
2091   SPIRVType *TargetImageType =
2092       GR->getOrCreateSPIRVType(ImageOpaqueType, MIRBuilder);
2093   return GR->getOrCreateOpTypeSampledImage(TargetImageType, MIRBuilder);
2094 }
2095
2096 namespace SPIRV {
2097 SPIRVType *lowerBuiltinType(const StructType *OpaqueType,
2098                             SPIRV::AccessQualifier::AccessQualifier AccessQual,
2099                             MachineIRBuilder &MIRBuilder,
2100                             SPIRVGlobalRegistry *GR) {
2101   assert(OpaqueType->hasName() &&
2102          "Structs representing builtin types must have a parsable name");
2103   unsigned NumStartingVRegs = MIRBuilder.getMRI()->getNumVirtRegs();
2104
2105   const StringRef Name = OpaqueType->getName();
2106   LLVM_DEBUG(dbgs() << "Lowering builtin type: " << Name << "\n");
2107
2108   // Lookup the demangled builtin type in the TableGen records.
2109   const SPIRV::DemangledType *TypeRecord = findBuiltinType(Name);
2110   if (!TypeRecord)
2111     report_fatal_error("Missing TableGen record for builtin type: " + Name);
2112
2113   // "Lower" the BuiltinType into TargetType. The following get<...>Type methods
2114   // use the implementation details from TableGen records to either create a new
2115   // OpType<...> machine instruction or get an existing equivalent SPIRVType
2116   // from GlobalRegistry.
2117   SPIRVType *TargetType;
2118   switch (TypeRecord->Opcode) {
2119   case SPIRV::OpTypeImage:
2120     TargetType = getImageType(OpaqueType, AccessQual, MIRBuilder, GR);
2121     break;
2122   case SPIRV::OpTypePipe:
2123     TargetType = getPipeType(OpaqueType, MIRBuilder, GR);
2124     break;
2125   case SPIRV::OpTypeDeviceEvent:
2126     TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
2127     break;
2128   case SPIRV::OpTypeSampler:
2129     TargetType = getSamplerType(MIRBuilder, GR);
2130     break;
2131   case SPIRV::OpTypeSampledImage:
2132     TargetType = getSampledImageType(OpaqueType, MIRBuilder, GR);
2133     break;
2134   default:
2135     TargetType = getNonParametrizedType(OpaqueType, TypeRecord, MIRBuilder, GR);
2136     break;
2137   }
2138
2139   // Emit OpName instruction if a new OpType<...> instruction was added
2140   // (equivalent type was not found in GlobalRegistry).
2141   if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
2142     buildOpName(GR->getSPIRVTypeID(TargetType), OpaqueType->getName(),
2143                 MIRBuilder);
2144
2145   return TargetType;
2146 }
2147 } // namespace SPIRV
2148 } // namespace llvm