brintos / llvm-project-archived public Read only

1//===- SPIRVInstructionSelector.cpp ------------------------------*- 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-exception6//7//===----------------------------------------------------------------------===//8//9// This file implements the targeting of the InstructionSelector class for10// SPIRV.11// TODO: This should be generated by TableGen.12//13//===----------------------------------------------------------------------===//14 15#include "MCTargetDesc/SPIRVBaseInfo.h"16#include "MCTargetDesc/SPIRVMCTargetDesc.h"17#include "SPIRV.h"18#include "SPIRVGlobalRegistry.h"19#include "SPIRVInstrInfo.h"20#include "SPIRVRegisterInfo.h"21#include "SPIRVTargetMachine.h"22#include "SPIRVUtils.h"23#include "llvm/ADT/APFloat.h"24#include "llvm/ADT/StringExtras.h"25#include "llvm/CodeGen/GlobalISel/GIMatchTableExecutorImpl.h"26#include "llvm/CodeGen/GlobalISel/GenericMachineInstrs.h"27#include "llvm/CodeGen/GlobalISel/InstructionSelector.h"28#include "llvm/CodeGen/MachineInstrBuilder.h"29#include "llvm/CodeGen/MachineRegisterInfo.h"30#include "llvm/CodeGen/Register.h"31#include "llvm/CodeGen/TargetOpcodes.h"32#include "llvm/IR/IntrinsicsSPIRV.h"33#include "llvm/Support/Debug.h"34#include "llvm/Support/ErrorHandling.h"35 36#define DEBUG_TYPE "spirv-isel"37 38using namespace llvm;39namespace CL = SPIRV::OpenCLExtInst;40namespace GL = SPIRV::GLSLExtInst;41 42using ExtInstList =43    std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;44 45namespace {46 47llvm::SPIRV::SelectionControl::SelectionControl48getSelectionOperandForImm(int Imm) {49  if (Imm == 2)50    return SPIRV::SelectionControl::Flatten;51  if (Imm == 1)52    return SPIRV::SelectionControl::DontFlatten;53  if (Imm == 0)54    return SPIRV::SelectionControl::None;55  llvm_unreachable("Invalid immediate");56}57 58#define GET_GLOBALISEL_PREDICATE_BITSET59#include "SPIRVGenGlobalISel.inc"60#undef GET_GLOBALISEL_PREDICATE_BITSET61 62class SPIRVInstructionSelector : public InstructionSelector {63  const SPIRVSubtarget &STI;64  const SPIRVInstrInfo &TII;65  const SPIRVRegisterInfo &TRI;66  const RegisterBankInfo &RBI;67  SPIRVGlobalRegistry &GR;68  MachineRegisterInfo *MRI;69  MachineFunction *HasVRegsReset = nullptr;70 71  /// We need to keep track of the number we give to anonymous global values to72  /// generate the same name every time when this is needed.73  mutable DenseMap<const GlobalValue *, unsigned> UnnamedGlobalIDs;74  SmallPtrSet<MachineInstr *, 8> DeadMIs;75 76public:77  SPIRVInstructionSelector(const SPIRVTargetMachine &TM,78                           const SPIRVSubtarget &ST,79                           const RegisterBankInfo &RBI);80  void setupMF(MachineFunction &MF, GISelValueTracking *VT,81               CodeGenCoverage *CoverageInfo, ProfileSummaryInfo *PSI,82               BlockFrequencyInfo *BFI) override;83  // Common selection code. Instruction-specific selection occurs in spvSelect.84  bool select(MachineInstr &I) override;85  static const char *getName() { return DEBUG_TYPE; }86 87#define GET_GLOBALISEL_PREDICATES_DECL88#include "SPIRVGenGlobalISel.inc"89#undef GET_GLOBALISEL_PREDICATES_DECL90 91#define GET_GLOBALISEL_TEMPORARIES_DECL92#include "SPIRVGenGlobalISel.inc"93#undef GET_GLOBALISEL_TEMPORARIES_DECL94 95private:96  void resetVRegsType(MachineFunction &MF);97  void removeDeadInstruction(MachineInstr &MI) const;98  void removeOpNamesForDeadMI(MachineInstr &MI) const;99 100  // tblgen-erated 'select' implementation, used as the initial selector for101  // the patterns that don't require complex C++.102  bool selectImpl(MachineInstr &I, CodeGenCoverage &CoverageInfo) const;103 104  // All instruction-specific selection that didn't happen in "select()".105  // Is basically a large Switch/Case delegating to all other select method.106  bool spvSelect(Register ResVReg, const SPIRVType *ResType,107                 MachineInstr &I) const;108 109  bool selectFirstBitHigh(Register ResVReg, const SPIRVType *ResType,110                          MachineInstr &I, bool IsSigned) const;111 112  bool selectFirstBitLow(Register ResVReg, const SPIRVType *ResType,113                         MachineInstr &I) const;114 115  bool selectFirstBitSet16(Register ResVReg, const SPIRVType *ResType,116                           MachineInstr &I, unsigned ExtendOpcode,117                           unsigned BitSetOpcode) const;118 119  bool selectFirstBitSet32(Register ResVReg, const SPIRVType *ResType,120                           MachineInstr &I, Register SrcReg,121                           unsigned BitSetOpcode) const;122 123  bool selectFirstBitSet64(Register ResVReg, const SPIRVType *ResType,124                           MachineInstr &I, Register SrcReg,125                           unsigned BitSetOpcode, bool SwapPrimarySide) const;126 127  bool selectFirstBitSet64Overflow(Register ResVReg, const SPIRVType *ResType,128                                   MachineInstr &I, Register SrcReg,129                                   unsigned BitSetOpcode,130                                   bool SwapPrimarySide) const;131 132  bool selectGlobalValue(Register ResVReg, MachineInstr &I,133                         const MachineInstr *Init = nullptr) const;134 135  bool selectOpWithSrcs(Register ResVReg, const SPIRVType *ResType,136                        MachineInstr &I, std::vector<Register> SrcRegs,137                        unsigned Opcode) const;138 139  bool selectUnOp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,140                  unsigned Opcode) const;141 142  bool selectBitcast(Register ResVReg, const SPIRVType *ResType,143                     MachineInstr &I) const;144 145  bool selectLoad(Register ResVReg, const SPIRVType *ResType,146                  MachineInstr &I) const;147  bool selectStore(MachineInstr &I) const;148 149  bool selectStackSave(Register ResVReg, const SPIRVType *ResType,150                       MachineInstr &I) const;151  bool selectStackRestore(MachineInstr &I) const;152 153  bool selectMemOperation(Register ResVReg, MachineInstr &I) const;154  Register getOrCreateMemSetGlobal(MachineInstr &I) const;155  bool selectCopyMemory(MachineInstr &I, Register SrcReg) const;156  bool selectCopyMemorySized(MachineInstr &I, Register SrcReg) const;157 158  bool selectAtomicRMW(Register ResVReg, const SPIRVType *ResType,159                       MachineInstr &I, unsigned NewOpcode,160                       unsigned NegateOpcode = 0) const;161 162  bool selectAtomicCmpXchg(Register ResVReg, const SPIRVType *ResType,163                           MachineInstr &I) const;164 165  bool selectFence(MachineInstr &I) const;166 167  bool selectAddrSpaceCast(Register ResVReg, const SPIRVType *ResType,168                           MachineInstr &I) const;169 170  bool selectAnyOrAll(Register ResVReg, const SPIRVType *ResType,171                      MachineInstr &I, unsigned OpType) const;172 173  bool selectAll(Register ResVReg, const SPIRVType *ResType,174                 MachineInstr &I) const;175 176  bool selectAny(Register ResVReg, const SPIRVType *ResType,177                 MachineInstr &I) const;178 179  bool selectBitreverse(Register ResVReg, const SPIRVType *ResType,180                        MachineInstr &I) const;181 182  bool selectBuildVector(Register ResVReg, const SPIRVType *ResType,183                         MachineInstr &I) const;184  bool selectSplatVector(Register ResVReg, const SPIRVType *ResType,185                         MachineInstr &I) const;186 187  bool selectCmp(Register ResVReg, const SPIRVType *ResType,188                 unsigned comparisonOpcode, MachineInstr &I) const;189  bool selectDiscard(Register ResVReg, const SPIRVType *ResType,190                     MachineInstr &I) const;191 192  bool selectICmp(Register ResVReg, const SPIRVType *ResType,193                  MachineInstr &I) const;194  bool selectFCmp(Register ResVReg, const SPIRVType *ResType,195                  MachineInstr &I) const;196 197  bool selectSign(Register ResVReg, const SPIRVType *ResType,198                  MachineInstr &I) const;199 200  bool selectFloatDot(Register ResVReg, const SPIRVType *ResType,201                      MachineInstr &I) const;202 203  bool selectOverflowArith(Register ResVReg, const SPIRVType *ResType,204                           MachineInstr &I, unsigned Opcode) const;205  bool selectDebugTrap(Register ResVReg, const SPIRVType *ResType,206                       MachineInstr &I) const;207 208  bool selectIntegerDot(Register ResVReg, const SPIRVType *ResType,209                        MachineInstr &I, bool Signed) const;210 211  bool selectIntegerDotExpansion(Register ResVReg, const SPIRVType *ResType,212                                 MachineInstr &I) const;213 214  bool selectOpIsInf(Register ResVReg, const SPIRVType *ResType,215                     MachineInstr &I) const;216 217  bool selectOpIsNan(Register ResVReg, const SPIRVType *ResType,218                     MachineInstr &I) const;219 220  template <bool Signed>221  bool selectDot4AddPacked(Register ResVReg, const SPIRVType *ResType,222                           MachineInstr &I) const;223  template <bool Signed>224  bool selectDot4AddPackedExpansion(Register ResVReg, const SPIRVType *ResType,225                                    MachineInstr &I) const;226 227  bool selectWaveReduceMax(Register ResVReg, const SPIRVType *ResType,228                           MachineInstr &I, bool IsUnsigned) const;229 230  bool selectWaveReduceMin(Register ResVReg, const SPIRVType *ResType,231                           MachineInstr &I, bool IsUnsigned) const;232 233  bool selectWaveReduceSum(Register ResVReg, const SPIRVType *ResType,234                           MachineInstr &I) const;235 236  bool selectConst(Register ResVReg, const SPIRVType *ResType,237                   MachineInstr &I) const;238 239  bool selectSelect(Register ResVReg, const SPIRVType *ResType,240                    MachineInstr &I) const;241  bool selectSelectDefaultArgs(Register ResVReg, const SPIRVType *ResType,242                               MachineInstr &I, bool IsSigned) const;243  bool selectIToF(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,244                  bool IsSigned, unsigned Opcode) const;245  bool selectExt(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,246                 bool IsSigned) const;247 248  bool selectTrunc(Register ResVReg, const SPIRVType *ResType,249                   MachineInstr &I) const;250 251  bool selectSUCmp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,252                   bool IsSigned) const;253 254  bool selectIntToBool(Register IntReg, Register ResVReg, MachineInstr &I,255                       const SPIRVType *intTy, const SPIRVType *boolTy) const;256 257  bool selectOpUndef(Register ResVReg, const SPIRVType *ResType,258                     MachineInstr &I) const;259  bool selectFreeze(Register ResVReg, const SPIRVType *ResType,260                    MachineInstr &I) const;261  bool selectIntrinsic(Register ResVReg, const SPIRVType *ResType,262                       MachineInstr &I) const;263  bool selectExtractVal(Register ResVReg, const SPIRVType *ResType,264                        MachineInstr &I) const;265  bool selectInsertVal(Register ResVReg, const SPIRVType *ResType,266                       MachineInstr &I) const;267  bool selectExtractElt(Register ResVReg, const SPIRVType *ResType,268                        MachineInstr &I) const;269  bool selectInsertElt(Register ResVReg, const SPIRVType *ResType,270                       MachineInstr &I) const;271  bool selectGEP(Register ResVReg, const SPIRVType *ResType,272                 MachineInstr &I) const;273 274  bool selectFrameIndex(Register ResVReg, const SPIRVType *ResType,275                        MachineInstr &I) const;276  bool selectAllocaArray(Register ResVReg, const SPIRVType *ResType,277                         MachineInstr &I) const;278 279  bool selectBranch(MachineInstr &I) const;280  bool selectBranchCond(MachineInstr &I) const;281 282  bool selectPhi(Register ResVReg, const SPIRVType *ResType,283                 MachineInstr &I) const;284 285  bool selectExtInst(Register ResVReg, const SPIRVType *RestType,286                     MachineInstr &I, GL::GLSLExtInst GLInst) const;287  bool selectExtInst(Register ResVReg, const SPIRVType *ResType,288                     MachineInstr &I, CL::OpenCLExtInst CLInst) const;289  bool selectExtInst(Register ResVReg, const SPIRVType *ResType,290                     MachineInstr &I, CL::OpenCLExtInst CLInst,291                     GL::GLSLExtInst GLInst) const;292  bool selectExtInst(Register ResVReg, const SPIRVType *ResType,293                     MachineInstr &I, const ExtInstList &ExtInsts) const;294  bool selectExtInstForLRound(Register ResVReg, const SPIRVType *ResType,295                              MachineInstr &I, CL::OpenCLExtInst CLInst,296                              GL::GLSLExtInst GLInst) const;297  bool selectExtInstForLRound(Register ResVReg, const SPIRVType *ResType,298                              MachineInstr &I,299                              const ExtInstList &ExtInsts) const;300 301  bool selectLog10(Register ResVReg, const SPIRVType *ResType,302                   MachineInstr &I) const;303 304  bool selectSaturate(Register ResVReg, const SPIRVType *ResType,305                      MachineInstr &I) const;306 307  bool selectWaveOpInst(Register ResVReg, const SPIRVType *ResType,308                        MachineInstr &I, unsigned Opcode) const;309 310  bool selectWaveActiveCountBits(Register ResVReg, const SPIRVType *ResType,311                                 MachineInstr &I) const;312 313  bool selectUnmergeValues(MachineInstr &I) const;314 315  bool selectHandleFromBinding(Register &ResVReg, const SPIRVType *ResType,316                               MachineInstr &I) const;317 318  bool selectCounterHandleFromBinding(Register &ResVReg,319                                      const SPIRVType *ResType,320                                      MachineInstr &I) const;321 322  bool selectReadImageIntrinsic(Register &ResVReg, const SPIRVType *ResType,323                                MachineInstr &I) const;324  bool selectImageWriteIntrinsic(MachineInstr &I) const;325  bool selectResourceGetPointer(Register &ResVReg, const SPIRVType *ResType,326                                MachineInstr &I) const;327  bool selectResourceNonUniformIndex(Register &ResVReg,328                                     const SPIRVType *ResType,329                                     MachineInstr &I) const;330  bool selectModf(Register ResVReg, const SPIRVType *ResType,331                  MachineInstr &I) const;332  bool selectUpdateCounter(Register &ResVReg, const SPIRVType *ResType,333                           MachineInstr &I) const;334  bool selectFrexp(Register ResVReg, const SPIRVType *ResType,335                   MachineInstr &I) const;336  bool selectDerivativeInst(Register ResVReg, const SPIRVType *ResType,337                            MachineInstr &I, const unsigned DPdOpCode) const;338  // Utilities339  std::pair<Register, bool>340  buildI32Constant(uint32_t Val, MachineInstr &I,341                   const SPIRVType *ResType = nullptr) const;342 343  Register buildZerosVal(const SPIRVType *ResType, MachineInstr &I) const;344  Register buildZerosValF(const SPIRVType *ResType, MachineInstr &I) const;345  Register buildOnesVal(bool AllOnes, const SPIRVType *ResType,346                        MachineInstr &I) const;347  Register buildOnesValF(const SPIRVType *ResType, MachineInstr &I) const;348 349  bool wrapIntoSpecConstantOp(MachineInstr &I,350                              SmallVector<Register> &CompositeArgs) const;351 352  Register getUcharPtrTypeReg(MachineInstr &I,353                              SPIRV::StorageClass::StorageClass SC) const;354  MachineInstrBuilder buildSpecConstantOp(MachineInstr &I, Register Dest,355                                          Register Src, Register DestType,356                                          uint32_t Opcode) const;357  MachineInstrBuilder buildConstGenericPtr(MachineInstr &I, Register SrcPtr,358                                           SPIRVType *SrcPtrTy) const;359  Register buildPointerToResource(const SPIRVType *ResType,360                                  SPIRV::StorageClass::StorageClass SC,361                                  uint32_t Set, uint32_t Binding,362                                  uint32_t ArraySize, Register IndexReg,363                                  StringRef Name,364                                  MachineIRBuilder MIRBuilder) const;365  SPIRVType *widenTypeToVec4(const SPIRVType *Type, MachineInstr &I) const;366  bool extractSubvector(Register &ResVReg, const SPIRVType *ResType,367                        Register &ReadReg, MachineInstr &InsertionPoint) const;368  bool generateImageReadOrFetch(Register &ResVReg, const SPIRVType *ResType,369                                Register ImageReg, Register IdxReg,370                                DebugLoc Loc, MachineInstr &Pos) const;371  bool BuildCOPY(Register DestReg, Register SrcReg, MachineInstr &I) const;372  bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,373                              Register ResVReg, const SPIRVType *ResType,374                              MachineInstr &I) const;375  bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,376                          Register ResVReg, const SPIRVType *ResType,377                          MachineInstr &I) const;378  bool loadHandleBeforePosition(Register &HandleReg, const SPIRVType *ResType,379                                GIntrinsic &HandleDef, MachineInstr &Pos) const;380  void decorateUsesAsNonUniform(Register &NonUniformReg) const;381  void errorIfInstrOutsideShader(MachineInstr &I) const;382};383 384bool sampledTypeIsSignedInteger(const llvm::Type *HandleType) {385  const TargetExtType *TET = cast<TargetExtType>(HandleType);386  if (TET->getTargetExtName() == "spirv.Image") {387    return false;388  }389  assert(TET->getTargetExtName() == "spirv.SignedImage");390  return TET->getTypeParameter(0)->isIntegerTy();391}392} // end anonymous namespace393 394#define GET_GLOBALISEL_IMPL395#include "SPIRVGenGlobalISel.inc"396#undef GET_GLOBALISEL_IMPL397 398SPIRVInstructionSelector::SPIRVInstructionSelector(const SPIRVTargetMachine &TM,399                                                   const SPIRVSubtarget &ST,400                                                   const RegisterBankInfo &RBI)401    : InstructionSelector(), STI(ST), TII(*ST.getInstrInfo()),402      TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),403      MRI(nullptr),404#define GET_GLOBALISEL_PREDICATES_INIT405#include "SPIRVGenGlobalISel.inc"406#undef GET_GLOBALISEL_PREDICATES_INIT407#define GET_GLOBALISEL_TEMPORARIES_INIT408#include "SPIRVGenGlobalISel.inc"409#undef GET_GLOBALISEL_TEMPORARIES_INIT410{411}412 413void SPIRVInstructionSelector::setupMF(MachineFunction &MF,414                                       GISelValueTracking *VT,415                                       CodeGenCoverage *CoverageInfo,416                                       ProfileSummaryInfo *PSI,417                                       BlockFrequencyInfo *BFI) {418  MRI = &MF.getRegInfo();419  GR.setCurrentFunc(MF);420  InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);421}422 423// Ensure that register classes correspond to pattern matching rules.424void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {425  if (HasVRegsReset == &MF)426    return;427  HasVRegsReset = &MF;428 429  MachineRegisterInfo &MRI = MF.getRegInfo();430  for (unsigned I = 0, E = MRI.getNumVirtRegs(); I != E; ++I) {431    Register Reg = Register::index2VirtReg(I);432    LLT RegType = MRI.getType(Reg);433    if (RegType.isScalar())434      MRI.setType(Reg, LLT::scalar(64));435    else if (RegType.isPointer())436      MRI.setType(Reg, LLT::pointer(0, 64));437    else if (RegType.isVector())438      MRI.setType(Reg, LLT::fixed_vector(2, LLT::scalar(64)));439  }440  for (const auto &MBB : MF) {441    for (const auto &MI : MBB) {442      if (isPreISelGenericOpcode(MI.getOpcode()))443        GR.erase(&MI);444      if (MI.getOpcode() != SPIRV::ASSIGN_TYPE)445        continue;446 447      Register DstReg = MI.getOperand(0).getReg();448      LLT DstType = MRI.getType(DstReg);449      Register SrcReg = MI.getOperand(1).getReg();450      LLT SrcType = MRI.getType(SrcReg);451      if (DstType != SrcType)452        MRI.setType(DstReg, MRI.getType(SrcReg));453 454      const TargetRegisterClass *DstRC = MRI.getRegClassOrNull(DstReg);455      const TargetRegisterClass *SrcRC = MRI.getRegClassOrNull(SrcReg);456      if (DstRC != SrcRC && SrcRC)457        MRI.setRegClass(DstReg, SrcRC);458    }459  }460}461 462// Return true if the type represents a constant register463static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef,464                       SmallPtrSet<SPIRVType *, 4> &Visited) {465  OpDef = passCopy(OpDef, MRI);466 467  if (Visited.contains(OpDef))468    return true;469  Visited.insert(OpDef);470 471  unsigned Opcode = OpDef->getOpcode();472  switch (Opcode) {473  case TargetOpcode::G_CONSTANT:474  case TargetOpcode::G_FCONSTANT:475  case TargetOpcode::G_IMPLICIT_DEF:476    return true;477  case TargetOpcode::G_INTRINSIC:478  case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:479  case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:480    return cast<GIntrinsic>(*OpDef).getIntrinsicID() ==481           Intrinsic::spv_const_composite;482  case TargetOpcode::G_BUILD_VECTOR:483  case TargetOpcode::G_SPLAT_VECTOR: {484    for (unsigned i = OpDef->getNumExplicitDefs(); i < OpDef->getNumOperands();485         i++) {486      MachineInstr *OpNestedDef =487          OpDef->getOperand(i).isReg()488              ? MRI->getVRegDef(OpDef->getOperand(i).getReg())489              : nullptr;490      if (OpNestedDef && !isConstReg(MRI, OpNestedDef, Visited))491        return false;492    }493    return true;494  case SPIRV::OpConstantTrue:495  case SPIRV::OpConstantFalse:496  case SPIRV::OpConstantI:497  case SPIRV::OpConstantF:498  case SPIRV::OpConstantComposite:499  case SPIRV::OpConstantCompositeContinuedINTEL:500  case SPIRV::OpConstantSampler:501  case SPIRV::OpConstantNull:502  case SPIRV::OpUndef:503  case SPIRV::OpConstantFunctionPointerINTEL:504    return true;505  }506  }507  return false;508}509 510// Return true if the virtual register represents a constant511static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {512  SmallPtrSet<SPIRVType *, 4> Visited;513  if (MachineInstr *OpDef = MRI->getVRegDef(OpReg))514    return isConstReg(MRI, OpDef, Visited);515  return false;516}517 518// TODO(168736): We should make this either a flag in tabelgen519// or reduce our dependence on the global registry, so we can remove this520// function. It can easily be missed when new intrinsics are added.521 522// Most SPIR-V instrinsics are considered to have side-effects in their tablegen523// definition because they are referenced in the global registry. This is a list524// of intrinsics that have no side effects other than their references in the525// global registry.526static bool intrinsicHasSideEffects(Intrinsic::ID ID) {527  switch (ID) {528  // This is not an exhaustive list and may need to be updated.529  case Intrinsic::spv_all:530  case Intrinsic::spv_alloca:531  case Intrinsic::spv_any:532  case Intrinsic::spv_bitcast:533  case Intrinsic::spv_const_composite:534  case Intrinsic::spv_cross:535  case Intrinsic::spv_degrees:536  case Intrinsic::spv_distance:537  case Intrinsic::spv_extractelt:538  case Intrinsic::spv_extractv:539  case Intrinsic::spv_faceforward:540  case Intrinsic::spv_fdot:541  case Intrinsic::spv_firstbitlow:542  case Intrinsic::spv_firstbitshigh:543  case Intrinsic::spv_firstbituhigh:544  case Intrinsic::spv_frac:545  case Intrinsic::spv_gep:546  case Intrinsic::spv_global_offset:547  case Intrinsic::spv_global_size:548  case Intrinsic::spv_group_id:549  case Intrinsic::spv_insertelt:550  case Intrinsic::spv_insertv:551  case Intrinsic::spv_isinf:552  case Intrinsic::spv_isnan:553  case Intrinsic::spv_lerp:554  case Intrinsic::spv_length:555  case Intrinsic::spv_normalize:556  case Intrinsic::spv_num_subgroups:557  case Intrinsic::spv_num_workgroups:558  case Intrinsic::spv_ptrcast:559  case Intrinsic::spv_radians:560  case Intrinsic::spv_reflect:561  case Intrinsic::spv_refract:562  case Intrinsic::spv_resource_getpointer:563  case Intrinsic::spv_resource_handlefrombinding:564  case Intrinsic::spv_resource_handlefromimplicitbinding:565  case Intrinsic::spv_resource_nonuniformindex:566  case Intrinsic::spv_rsqrt:567  case Intrinsic::spv_saturate:568  case Intrinsic::spv_sdot:569  case Intrinsic::spv_sign:570  case Intrinsic::spv_smoothstep:571  case Intrinsic::spv_step:572  case Intrinsic::spv_subgroup_id:573  case Intrinsic::spv_subgroup_local_invocation_id:574  case Intrinsic::spv_subgroup_max_size:575  case Intrinsic::spv_subgroup_size:576  case Intrinsic::spv_thread_id:577  case Intrinsic::spv_thread_id_in_group:578  case Intrinsic::spv_udot:579  case Intrinsic::spv_undef:580  case Intrinsic::spv_value_md:581  case Intrinsic::spv_workgroup_size:582    return false;583  default:584    return true;585  }586}587 588// TODO(168736): We should make this either a flag in tabelgen589// or reduce our dependence on the global registry, so we can remove this590// function. It can easily be missed when new intrinsics are added.591static bool isOpcodeWithNoSideEffects(unsigned Opcode) {592  switch (Opcode) {593  case SPIRV::OpTypeVoid:594  case SPIRV::OpTypeBool:595  case SPIRV::OpTypeInt:596  case SPIRV::OpTypeFloat:597  case SPIRV::OpTypeVector:598  case SPIRV::OpTypeMatrix:599  case SPIRV::OpTypeImage:600  case SPIRV::OpTypeSampler:601  case SPIRV::OpTypeSampledImage:602  case SPIRV::OpTypeArray:603  case SPIRV::OpTypeRuntimeArray:604  case SPIRV::OpTypeStruct:605  case SPIRV::OpTypeOpaque:606  case SPIRV::OpTypePointer:607  case SPIRV::OpTypeFunction:608  case SPIRV::OpTypeEvent:609  case SPIRV::OpTypeDeviceEvent:610  case SPIRV::OpTypeReserveId:611  case SPIRV::OpTypeQueue:612  case SPIRV::OpTypePipe:613  case SPIRV::OpTypeForwardPointer:614  case SPIRV::OpTypePipeStorage:615  case SPIRV::OpTypeNamedBarrier:616  case SPIRV::OpTypeAccelerationStructureNV:617  case SPIRV::OpTypeCooperativeMatrixNV:618  case SPIRV::OpTypeCooperativeMatrixKHR:619    return true;620  default:621    return false;622  }623}624 625bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {626  // If there are no definitions, then assume there is some other627  // side-effect that makes this instruction live.628  if (MI.getNumDefs() == 0)629    return false;630 631  for (const auto &MO : MI.all_defs()) {632    Register Reg = MO.getReg();633    if (Reg.isPhysical()) {634      LLVM_DEBUG(dbgs() << "Not dead: def of physical register " << Reg);635      return false;636    }637    for (const auto &UseMI : MRI.use_nodbg_instructions(Reg)) {638      if (UseMI.getOpcode() != SPIRV::OpName) {639        LLVM_DEBUG(dbgs() << "Not dead: def " << MO << " has use in " << UseMI);640        return false;641      }642    }643  }644 645  if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||646      MI.isLifetimeMarker()) {647    LLVM_DEBUG(648        dbgs()649        << "Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");650    return false;651  }652  if (MI.isPHI()) {653    LLVM_DEBUG(dbgs() << "Dead: Phi instruction with no uses.\n");654    return true;655  }656 657  // It is possible that the only side effect is that the instruction is658  // referenced in the global registry. If that is the only side effect, the659  // intrinsic is dead.660  if (MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||661      MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {662    const auto &Intr = cast<GIntrinsic>(MI);663    if (!intrinsicHasSideEffects(Intr.getIntrinsicID())) {664      LLVM_DEBUG(dbgs() << "Dead: Intrinsic with no real side effects.\n");665      return true;666    }667  }668 669  if (MI.mayStore() || MI.isCall() ||670      (MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||671      MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo()) {672    LLVM_DEBUG(dbgs() << "Not dead: instruction has side effects.\n");673    return false;674  }675 676  if (isPreISelGenericOpcode(MI.getOpcode())) {677    // TODO: Is there a generic way to check if the opcode has side effects?678    LLVM_DEBUG(dbgs() << "Dead: Generic opcode with no uses.\n");679    return true;680  }681 682  if (isOpcodeWithNoSideEffects(MI.getOpcode())) {683    LLVM_DEBUG(dbgs() << "Dead: known opcode with no side effects\n");684    return true;685  }686 687  return false;688}689 690void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &MI) const {691  // Delete the OpName that uses the result if there is one.692  for (const auto &MO : MI.all_defs()) {693    Register Reg = MO.getReg();694    if (Reg.isPhysical())695      continue;696    SmallVector<MachineInstr *, 4> UselessOpNames;697    for (MachineInstr &UseMI : MRI->use_nodbg_instructions(Reg)) {698      assert(UseMI.getOpcode() == SPIRV::OpName &&699             "There is still a use of the dead function.");700      UselessOpNames.push_back(&UseMI);701    }702    for (MachineInstr *OpNameMI : UselessOpNames) {703      GR.invalidateMachineInstr(OpNameMI);704      OpNameMI->eraseFromParent();705    }706  }707}708 709void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &MI) const {710  salvageDebugInfo(*MRI, MI);711  GR.invalidateMachineInstr(&MI);712  removeOpNamesForDeadMI(MI);713  MI.eraseFromParent();714}715 716bool SPIRVInstructionSelector::select(MachineInstr &I) {717  resetVRegsType(*I.getParent()->getParent());718 719  assert(I.getParent() && "Instruction should be in a basic block!");720  assert(I.getParent()->getParent() && "Instruction should be in a function!");721 722  LLVM_DEBUG(dbgs() << "Checking if instruction is dead: " << I;);723  if (isDead(I, *MRI)) {724    LLVM_DEBUG(dbgs() << "Instruction is dead.\n");725    removeDeadInstruction(I);726    return true;727  }728 729  Register Opcode = I.getOpcode();730  // If it's not a GMIR instruction, we've selected it already.731  if (!isPreISelGenericOpcode(Opcode)) {732    if (Opcode == SPIRV::ASSIGN_TYPE) { // These pseudos aren't needed any more.733      Register DstReg = I.getOperand(0).getReg();734      Register SrcReg = I.getOperand(1).getReg();735      auto *Def = MRI->getVRegDef(SrcReg);736      if (isTypeFoldingSupported(Def->getOpcode()) &&737          Def->getOpcode() != TargetOpcode::G_CONSTANT &&738          Def->getOpcode() != TargetOpcode::G_FCONSTANT) {739        bool Res = false;740        if (Def->getOpcode() == TargetOpcode::G_SELECT) {741          Register SelectDstReg = Def->getOperand(0).getReg();742          Res = selectSelect(SelectDstReg, GR.getSPIRVTypeForVReg(SelectDstReg),743                             *Def);744          GR.invalidateMachineInstr(Def);745          Def->removeFromParent();746          MRI->replaceRegWith(DstReg, SelectDstReg);747          GR.invalidateMachineInstr(&I);748          I.removeFromParent();749        } else750          Res = selectImpl(I, *CoverageInfo);751        LLVM_DEBUG({752          if (!Res && Def->getOpcode() != TargetOpcode::G_CONSTANT) {753            dbgs() << "Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";754            I.print(dbgs());755          }756        });757        assert(Res || Def->getOpcode() == TargetOpcode::G_CONSTANT);758        if (Res) {759          if (!isTriviallyDead(*Def, *MRI) && isDead(*Def, *MRI))760            DeadMIs.insert(Def);761          return Res;762        }763      }764      MRI->setRegClass(SrcReg, MRI->getRegClass(DstReg));765      MRI->replaceRegWith(SrcReg, DstReg);766      GR.invalidateMachineInstr(&I);767      I.removeFromParent();768      return true;769    } else if (I.getNumDefs() == 1) {770      // Make all vregs 64 bits (for SPIR-V IDs).771      MRI->setType(I.getOperand(0).getReg(), LLT::scalar(64));772    }773    return constrainSelectedInstRegOperands(I, TII, TRI, RBI);774  }775 776  if (DeadMIs.contains(&I)) {777    // if the instruction has been already made dead by folding it away778    // erase it779    LLVM_DEBUG(dbgs() << "Instruction is folded and dead.\n");780    removeDeadInstruction(I);781    return true;782  }783 784  if (I.getNumOperands() != I.getNumExplicitOperands()) {785    LLVM_DEBUG(errs() << "Generic instr has unexpected implicit operands\n");786    return false;787  }788 789  // Common code for getting return reg+type, and removing selected instr790  // from parent occurs here. Instr-specific selection happens in spvSelect().791  bool HasDefs = I.getNumDefs() > 0;792  Register ResVReg = HasDefs ? I.getOperand(0).getReg() : Register(0);793  SPIRVType *ResType = HasDefs ? GR.getSPIRVTypeForVReg(ResVReg) : nullptr;794  assert(!HasDefs || ResType || I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||795         I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);796  if (spvSelect(ResVReg, ResType, I)) {797    if (HasDefs) // Make all vregs 64 bits (for SPIR-V IDs).798      for (unsigned i = 0; i < I.getNumDefs(); ++i)799        MRI->setType(I.getOperand(i).getReg(), LLT::scalar(64));800    GR.invalidateMachineInstr(&I);801    I.removeFromParent();802    return true;803  }804  return false;805}806 807static bool mayApplyGenericSelection(unsigned Opcode) {808  switch (Opcode) {809  case TargetOpcode::G_CONSTANT:810  case TargetOpcode::G_FCONSTANT:811    return false;812  case TargetOpcode::G_SADDO:813  case TargetOpcode::G_SSUBO:814    return true;815  }816  return isTypeFoldingSupported(Opcode);817}818 819bool SPIRVInstructionSelector::BuildCOPY(Register DestReg, Register SrcReg,820                                         MachineInstr &I) const {821  const TargetRegisterClass *DstRC = MRI->getRegClassOrNull(DestReg);822  const TargetRegisterClass *SrcRC = MRI->getRegClassOrNull(SrcReg);823  if (DstRC != SrcRC && SrcRC)824    MRI->setRegClass(DestReg, SrcRC);825  return BuildMI(*I.getParent(), I, I.getDebugLoc(),826                 TII.get(TargetOpcode::COPY))827      .addDef(DestReg)828      .addUse(SrcReg)829      .constrainAllUses(TII, TRI, RBI);830}831 832bool SPIRVInstructionSelector::spvSelect(Register ResVReg,833                                         const SPIRVType *ResType,834                                         MachineInstr &I) const {835  const unsigned Opcode = I.getOpcode();836  if (mayApplyGenericSelection(Opcode))837    return selectImpl(I, *CoverageInfo);838  switch (Opcode) {839  case TargetOpcode::G_CONSTANT:840  case TargetOpcode::G_FCONSTANT:841    return selectConst(ResVReg, ResType, I);842  case TargetOpcode::G_GLOBAL_VALUE:843    return selectGlobalValue(ResVReg, I);844  case TargetOpcode::G_IMPLICIT_DEF:845    return selectOpUndef(ResVReg, ResType, I);846  case TargetOpcode::G_FREEZE:847    return selectFreeze(ResVReg, ResType, I);848 849  case TargetOpcode::G_INTRINSIC:850  case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:851  case TargetOpcode::G_INTRINSIC_CONVERGENT:852  case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:853    return selectIntrinsic(ResVReg, ResType, I);854  case TargetOpcode::G_BITREVERSE:855    return selectBitreverse(ResVReg, ResType, I);856 857  case TargetOpcode::G_BUILD_VECTOR:858    return selectBuildVector(ResVReg, ResType, I);859  case TargetOpcode::G_SPLAT_VECTOR:860    return selectSplatVector(ResVReg, ResType, I);861 862  case TargetOpcode::G_SHUFFLE_VECTOR: {863    MachineBasicBlock &BB = *I.getParent();864    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))865                   .addDef(ResVReg)866                   .addUse(GR.getSPIRVTypeID(ResType))867                   .addUse(I.getOperand(1).getReg())868                   .addUse(I.getOperand(2).getReg());869    for (auto V : I.getOperand(3).getShuffleMask())870      MIB.addImm(V);871    return MIB.constrainAllUses(TII, TRI, RBI);872  }873  case TargetOpcode::G_MEMMOVE:874  case TargetOpcode::G_MEMCPY:875  case TargetOpcode::G_MEMSET:876    return selectMemOperation(ResVReg, I);877 878  case TargetOpcode::G_ICMP:879    return selectICmp(ResVReg, ResType, I);880  case TargetOpcode::G_FCMP:881    return selectFCmp(ResVReg, ResType, I);882 883  case TargetOpcode::G_FRAME_INDEX:884    return selectFrameIndex(ResVReg, ResType, I);885 886  case TargetOpcode::G_LOAD:887    return selectLoad(ResVReg, ResType, I);888  case TargetOpcode::G_STORE:889    return selectStore(I);890 891  case TargetOpcode::G_BR:892    return selectBranch(I);893  case TargetOpcode::G_BRCOND:894    return selectBranchCond(I);895 896  case TargetOpcode::G_PHI:897    return selectPhi(ResVReg, ResType, I);898 899  case TargetOpcode::G_FPTOSI:900    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);901  case TargetOpcode::G_FPTOUI:902    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);903 904  case TargetOpcode::G_FPTOSI_SAT:905    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);906  case TargetOpcode::G_FPTOUI_SAT:907    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);908 909  case TargetOpcode::G_SITOFP:910    return selectIToF(ResVReg, ResType, I, true, SPIRV::OpConvertSToF);911  case TargetOpcode::G_UITOFP:912    return selectIToF(ResVReg, ResType, I, false, SPIRV::OpConvertUToF);913 914  case TargetOpcode::G_CTPOP:915    return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitCount);916  case TargetOpcode::G_SMIN:917    return selectExtInst(ResVReg, ResType, I, CL::s_min, GL::SMin);918  case TargetOpcode::G_UMIN:919    return selectExtInst(ResVReg, ResType, I, CL::u_min, GL::UMin);920 921  case TargetOpcode::G_SMAX:922    return selectExtInst(ResVReg, ResType, I, CL::s_max, GL::SMax);923  case TargetOpcode::G_UMAX:924    return selectExtInst(ResVReg, ResType, I, CL::u_max, GL::UMax);925 926  case TargetOpcode::G_SCMP:927    return selectSUCmp(ResVReg, ResType, I, true);928  case TargetOpcode::G_UCMP:929    return selectSUCmp(ResVReg, ResType, I, false);930  case TargetOpcode::G_LROUND:931  case TargetOpcode::G_LLROUND: {932    Register regForLround =933        MRI->createVirtualRegister(MRI->getRegClass(ResVReg), "lround");934    MRI->setRegClass(regForLround, &SPIRV::iIDRegClass);935    GR.assignSPIRVTypeToVReg(GR.getSPIRVTypeForVReg(I.getOperand(1).getReg()),936                             regForLround, *(I.getParent()->getParent()));937    selectExtInstForLRound(regForLround, GR.getSPIRVTypeForVReg(regForLround),938                           I, CL::round, GL::Round);939    MachineBasicBlock &BB = *I.getParent();940    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConvertFToS))941                   .addDef(ResVReg)942                   .addUse(GR.getSPIRVTypeID(ResType))943                   .addUse(regForLround);944    return MIB.constrainAllUses(TII, TRI, RBI);945  }946  case TargetOpcode::G_STRICT_FMA:947  case TargetOpcode::G_FMA:948    return selectExtInst(ResVReg, ResType, I, CL::fma, GL::Fma);949 950  case TargetOpcode::G_STRICT_FLDEXP:951    return selectExtInst(ResVReg, ResType, I, CL::ldexp);952 953  case TargetOpcode::G_FPOW:954    return selectExtInst(ResVReg, ResType, I, CL::pow, GL::Pow);955  case TargetOpcode::G_FPOWI:956    return selectExtInst(ResVReg, ResType, I, CL::pown);957 958  case TargetOpcode::G_FEXP:959    return selectExtInst(ResVReg, ResType, I, CL::exp, GL::Exp);960  case TargetOpcode::G_FEXP2:961    return selectExtInst(ResVReg, ResType, I, CL::exp2, GL::Exp2);962  case TargetOpcode::G_FMODF:963    return selectModf(ResVReg, ResType, I);964 965  case TargetOpcode::G_FLOG:966    return selectExtInst(ResVReg, ResType, I, CL::log, GL::Log);967  case TargetOpcode::G_FLOG2:968    return selectExtInst(ResVReg, ResType, I, CL::log2, GL::Log2);969  case TargetOpcode::G_FLOG10:970    return selectLog10(ResVReg, ResType, I);971 972  case TargetOpcode::G_FABS:973    return selectExtInst(ResVReg, ResType, I, CL::fabs, GL::FAbs);974  case TargetOpcode::G_ABS:975    return selectExtInst(ResVReg, ResType, I, CL::s_abs, GL::SAbs);976 977  case TargetOpcode::G_FMINNUM:978  case TargetOpcode::G_FMINIMUM:979    return selectExtInst(ResVReg, ResType, I, CL::fmin, GL::NMin);980  case TargetOpcode::G_FMAXNUM:981  case TargetOpcode::G_FMAXIMUM:982    return selectExtInst(ResVReg, ResType, I, CL::fmax, GL::NMax);983 984  case TargetOpcode::G_FCOPYSIGN:985    return selectExtInst(ResVReg, ResType, I, CL::copysign);986 987  case TargetOpcode::G_FCEIL:988    return selectExtInst(ResVReg, ResType, I, CL::ceil, GL::Ceil);989  case TargetOpcode::G_FFLOOR:990    return selectExtInst(ResVReg, ResType, I, CL::floor, GL::Floor);991 992  case TargetOpcode::G_FCOS:993    return selectExtInst(ResVReg, ResType, I, CL::cos, GL::Cos);994  case TargetOpcode::G_FSIN:995    return selectExtInst(ResVReg, ResType, I, CL::sin, GL::Sin);996  case TargetOpcode::G_FTAN:997    return selectExtInst(ResVReg, ResType, I, CL::tan, GL::Tan);998  case TargetOpcode::G_FACOS:999    return selectExtInst(ResVReg, ResType, I, CL::acos, GL::Acos);1000  case TargetOpcode::G_FASIN:1001    return selectExtInst(ResVReg, ResType, I, CL::asin, GL::Asin);1002  case TargetOpcode::G_FATAN:1003    return selectExtInst(ResVReg, ResType, I, CL::atan, GL::Atan);1004  case TargetOpcode::G_FATAN2:1005    return selectExtInst(ResVReg, ResType, I, CL::atan2, GL::Atan2);1006  case TargetOpcode::G_FCOSH:1007    return selectExtInst(ResVReg, ResType, I, CL::cosh, GL::Cosh);1008  case TargetOpcode::G_FSINH:1009    return selectExtInst(ResVReg, ResType, I, CL::sinh, GL::Sinh);1010  case TargetOpcode::G_FTANH:1011    return selectExtInst(ResVReg, ResType, I, CL::tanh, GL::Tanh);1012 1013  case TargetOpcode::G_STRICT_FSQRT:1014  case TargetOpcode::G_FSQRT:1015    return selectExtInst(ResVReg, ResType, I, CL::sqrt, GL::Sqrt);1016 1017  case TargetOpcode::G_CTTZ:1018  case TargetOpcode::G_CTTZ_ZERO_UNDEF:1019    return selectExtInst(ResVReg, ResType, I, CL::ctz);1020  case TargetOpcode::G_CTLZ:1021  case TargetOpcode::G_CTLZ_ZERO_UNDEF:1022    return selectExtInst(ResVReg, ResType, I, CL::clz);1023 1024  case TargetOpcode::G_INTRINSIC_ROUND:1025    return selectExtInst(ResVReg, ResType, I, CL::round, GL::Round);1026  case TargetOpcode::G_INTRINSIC_ROUNDEVEN:1027    return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);1028  case TargetOpcode::G_INTRINSIC_TRUNC:1029    return selectExtInst(ResVReg, ResType, I, CL::trunc, GL::Trunc);1030  case TargetOpcode::G_FRINT:1031  case TargetOpcode::G_FNEARBYINT:1032    return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);1033 1034  case TargetOpcode::G_SMULH:1035    return selectExtInst(ResVReg, ResType, I, CL::s_mul_hi);1036  case TargetOpcode::G_UMULH:1037    return selectExtInst(ResVReg, ResType, I, CL::u_mul_hi);1038 1039  case TargetOpcode::G_SADDSAT:1040    return selectExtInst(ResVReg, ResType, I, CL::s_add_sat);1041  case TargetOpcode::G_UADDSAT:1042    return selectExtInst(ResVReg, ResType, I, CL::u_add_sat);1043  case TargetOpcode::G_SSUBSAT:1044    return selectExtInst(ResVReg, ResType, I, CL::s_sub_sat);1045  case TargetOpcode::G_USUBSAT:1046    return selectExtInst(ResVReg, ResType, I, CL::u_sub_sat);1047 1048  case TargetOpcode::G_FFREXP:1049    return selectFrexp(ResVReg, ResType, I);1050 1051  case TargetOpcode::G_UADDO:1052    return selectOverflowArith(ResVReg, ResType, I,1053                               ResType->getOpcode() == SPIRV::OpTypeVector1054                                   ? SPIRV::OpIAddCarryV1055                                   : SPIRV::OpIAddCarryS);1056  case TargetOpcode::G_USUBO:1057    return selectOverflowArith(ResVReg, ResType, I,1058                               ResType->getOpcode() == SPIRV::OpTypeVector1059                                   ? SPIRV::OpISubBorrowV1060                                   : SPIRV::OpISubBorrowS);1061  case TargetOpcode::G_UMULO:1062    return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpUMulExtended);1063  case TargetOpcode::G_SMULO:1064    return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpSMulExtended);1065 1066  case TargetOpcode::G_SEXT:1067    return selectExt(ResVReg, ResType, I, true);1068  case TargetOpcode::G_ANYEXT:1069  case TargetOpcode::G_ZEXT:1070    return selectExt(ResVReg, ResType, I, false);1071  case TargetOpcode::G_TRUNC:1072    return selectTrunc(ResVReg, ResType, I);1073  case TargetOpcode::G_FPTRUNC:1074  case TargetOpcode::G_FPEXT:1075    return selectUnOp(ResVReg, ResType, I, SPIRV::OpFConvert);1076 1077  case TargetOpcode::G_PTRTOINT:1078    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertPtrToU);1079  case TargetOpcode::G_INTTOPTR:1080    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertUToPtr);1081  case TargetOpcode::G_BITCAST:1082    return selectBitcast(ResVReg, ResType, I);1083  case TargetOpcode::G_ADDRSPACE_CAST:1084    return selectAddrSpaceCast(ResVReg, ResType, I);1085  case TargetOpcode::G_PTR_ADD: {1086    // Currently, we get G_PTR_ADD only applied to global variables.1087    assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());1088    Register GV = I.getOperand(1).getReg();1089    MachineRegisterInfo::def_instr_iterator II = MRI->def_instr_begin(GV);1090    (void)II;1091    assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||1092            (*II).getOpcode() == TargetOpcode::COPY ||1093            (*II).getOpcode() == SPIRV::OpVariable) &&1094           getImm(I.getOperand(2), MRI));1095    // It may be the initialization of a global variable.1096    bool IsGVInit = false;1097    for (MachineRegisterInfo::use_instr_iterator1098             UseIt = MRI->use_instr_begin(I.getOperand(0).getReg()),1099             UseEnd = MRI->use_instr_end();1100         UseIt != UseEnd; UseIt = std::next(UseIt)) {1101      if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||1102          (*UseIt).getOpcode() == SPIRV::OpVariable) {1103        IsGVInit = true;1104        break;1105      }1106    }1107    MachineBasicBlock &BB = *I.getParent();1108    if (!IsGVInit) {1109      SPIRVType *GVType = GR.getSPIRVTypeForVReg(GV);1110      SPIRVType *GVPointeeType = GR.getPointeeType(GVType);1111      SPIRVType *ResPointeeType = GR.getPointeeType(ResType);1112      if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {1113        // Build a new virtual register that is associated with the required1114        // data type.1115        Register NewVReg = MRI->createGenericVirtualRegister(MRI->getType(GV));1116        MRI->setRegClass(NewVReg, MRI->getRegClass(GV));1117        //  Having a correctly typed base we are ready to build the actually1118        //  required GEP. It may not be a constant though, because all Operands1119        //  of OpSpecConstantOp is to originate from other const instructions,1120        //  and only the AccessChain named opcodes accept a global OpVariable1121        //  instruction. We can't use an AccessChain opcode because of the type1122        //  mismatch between result and base types.1123        if (!GR.isBitcastCompatible(ResType, GVType))1124          report_fatal_error(1125              "incompatible result and operand types in a bitcast");1126        Register ResTypeReg = GR.getSPIRVTypeID(ResType);1127        MachineInstrBuilder MIB =1128            BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitcast))1129                .addDef(NewVReg)1130                .addUse(ResTypeReg)1131                .addUse(GV);1132        return MIB.constrainAllUses(TII, TRI, RBI) &&1133               BuildMI(BB, I, I.getDebugLoc(),1134                       TII.get(STI.isLogicalSPIRV()1135                                   ? SPIRV::OpInBoundsAccessChain1136                                   : SPIRV::OpInBoundsPtrAccessChain))1137                   .addDef(ResVReg)1138                   .addUse(ResTypeReg)1139                   .addUse(NewVReg)1140                   .addUse(I.getOperand(2).getReg())1141                   .constrainAllUses(TII, TRI, RBI);1142      } else {1143        return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))1144            .addDef(ResVReg)1145            .addUse(GR.getSPIRVTypeID(ResType))1146            .addImm(1147                static_cast<uint32_t>(SPIRV::Opcode::InBoundsPtrAccessChain))1148            .addUse(GV)1149            .addUse(I.getOperand(2).getReg())1150            .constrainAllUses(TII, TRI, RBI);1151      }1152    }1153    // It's possible to translate G_PTR_ADD to OpSpecConstantOp: either to1154    // initialize a global variable with a constant expression (e.g., the test1155    // case opencl/basic/progvar_prog_scope_init.ll), or for another use case1156    Register Idx = buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);1157    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))1158                   .addDef(ResVReg)1159                   .addUse(GR.getSPIRVTypeID(ResType))1160                   .addImm(static_cast<uint32_t>(1161                       SPIRV::Opcode::InBoundsPtrAccessChain))1162                   .addUse(GV)1163                   .addUse(Idx)1164                   .addUse(I.getOperand(2).getReg());1165    return MIB.constrainAllUses(TII, TRI, RBI);1166  }1167 1168  case TargetOpcode::G_ATOMICRMW_OR:1169    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicOr);1170  case TargetOpcode::G_ATOMICRMW_ADD:1171    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicIAdd);1172  case TargetOpcode::G_ATOMICRMW_AND:1173    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicAnd);1174  case TargetOpcode::G_ATOMICRMW_MAX:1175    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMax);1176  case TargetOpcode::G_ATOMICRMW_MIN:1177    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMin);1178  case TargetOpcode::G_ATOMICRMW_SUB:1179    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicISub);1180  case TargetOpcode::G_ATOMICRMW_XOR:1181    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicXor);1182  case TargetOpcode::G_ATOMICRMW_UMAX:1183    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMax);1184  case TargetOpcode::G_ATOMICRMW_UMIN:1185    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMin);1186  case TargetOpcode::G_ATOMICRMW_XCHG:1187    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicExchange);1188  case TargetOpcode::G_ATOMIC_CMPXCHG:1189    return selectAtomicCmpXchg(ResVReg, ResType, I);1190 1191  case TargetOpcode::G_ATOMICRMW_FADD:1192    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT);1193  case TargetOpcode::G_ATOMICRMW_FSUB:1194    // Translate G_ATOMICRMW_FSUB to OpAtomicFAddEXT with negative value operand1195    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT,1196                           SPIRV::OpFNegate);1197  case TargetOpcode::G_ATOMICRMW_FMIN:1198    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMinEXT);1199  case TargetOpcode::G_ATOMICRMW_FMAX:1200    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMaxEXT);1201 1202  case TargetOpcode::G_FENCE:1203    return selectFence(I);1204 1205  case TargetOpcode::G_STACKSAVE:1206    return selectStackSave(ResVReg, ResType, I);1207  case TargetOpcode::G_STACKRESTORE:1208    return selectStackRestore(I);1209 1210  case TargetOpcode::G_UNMERGE_VALUES:1211    return selectUnmergeValues(I);1212 1213  // Discard gen opcodes for intrinsics which we do not expect to actually1214  // represent code after lowering or intrinsics which are not implemented but1215  // should not crash when found in a customer's LLVM IR input.1216  case TargetOpcode::G_TRAP:1217  case TargetOpcode::G_UBSANTRAP:1218  case TargetOpcode::DBG_LABEL:1219    return true;1220  case TargetOpcode::G_DEBUGTRAP:1221    return selectDebugTrap(ResVReg, ResType, I);1222 1223  default:1224    return false;1225  }1226}1227 1228bool SPIRVInstructionSelector::selectDebugTrap(Register ResVReg,1229                                               const SPIRVType *ResType,1230                                               MachineInstr &I) const {1231  unsigned Opcode = SPIRV::OpNop;1232  MachineBasicBlock &BB = *I.getParent();1233  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))1234      .constrainAllUses(TII, TRI, RBI);1235}1236 1237bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,1238                                             const SPIRVType *ResType,1239                                             MachineInstr &I,1240                                             GL::GLSLExtInst GLInst) const {1241  if (!STI.canUseExtInstSet(1242          SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {1243    std::string DiagMsg;1244    raw_string_ostream OS(DiagMsg);1245    I.print(OS, true, false, false, false);1246    DiagMsg += " is only supported with the GLSL extended instruction set.\n";1247    report_fatal_error(DiagMsg.c_str(), false);1248  }1249  return selectExtInst(ResVReg, ResType, I,1250                       {{SPIRV::InstructionSet::GLSL_std_450, GLInst}});1251}1252 1253bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,1254                                             const SPIRVType *ResType,1255                                             MachineInstr &I,1256                                             CL::OpenCLExtInst CLInst) const {1257  return selectExtInst(ResVReg, ResType, I,1258                       {{SPIRV::InstructionSet::OpenCL_std, CLInst}});1259}1260 1261bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,1262                                             const SPIRVType *ResType,1263                                             MachineInstr &I,1264                                             CL::OpenCLExtInst CLInst,1265                                             GL::GLSLExtInst GLInst) const {1266  ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},1267                          {SPIRV::InstructionSet::GLSL_std_450, GLInst}};1268  return selectExtInst(ResVReg, ResType, I, ExtInsts);1269}1270 1271bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,1272                                             const SPIRVType *ResType,1273                                             MachineInstr &I,1274                                             const ExtInstList &Insts) const {1275 1276  for (const auto &Ex : Insts) {1277    SPIRV::InstructionSet::InstructionSet Set = Ex.first;1278    uint32_t Opcode = Ex.second;1279    if (STI.canUseExtInstSet(Set)) {1280      MachineBasicBlock &BB = *I.getParent();1281      auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))1282                     .addDef(ResVReg)1283                     .addUse(GR.getSPIRVTypeID(ResType))1284                     .addImm(static_cast<uint32_t>(Set))1285                     .addImm(Opcode)1286                     .setMIFlags(I.getFlags());1287      const unsigned NumOps = I.getNumOperands();1288      unsigned Index = 1;1289      if (Index < NumOps &&1290          I.getOperand(Index).getType() ==1291              MachineOperand::MachineOperandType::MO_IntrinsicID)1292        Index = 2;1293      for (; Index < NumOps; ++Index)1294        MIB.add(I.getOperand(Index));1295      return MIB.constrainAllUses(TII, TRI, RBI);1296    }1297  }1298  return false;1299}1300bool SPIRVInstructionSelector::selectExtInstForLRound(1301    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,1302    CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst) const {1303  ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},1304                          {SPIRV::InstructionSet::GLSL_std_450, GLInst}};1305  return selectExtInstForLRound(ResVReg, ResType, I, ExtInsts);1306}1307 1308bool SPIRVInstructionSelector::selectExtInstForLRound(1309    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,1310    const ExtInstList &Insts) const {1311  for (const auto &Ex : Insts) {1312    SPIRV::InstructionSet::InstructionSet Set = Ex.first;1313    uint32_t Opcode = Ex.second;1314    if (STI.canUseExtInstSet(Set)) {1315      MachineBasicBlock &BB = *I.getParent();1316      auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))1317                     .addDef(ResVReg)1318                     .addUse(GR.getSPIRVTypeID(ResType))1319                     .addImm(static_cast<uint32_t>(Set))1320                     .addImm(Opcode);1321      const unsigned NumOps = I.getNumOperands();1322      unsigned Index = 1;1323      if (Index < NumOps &&1324          I.getOperand(Index).getType() ==1325              MachineOperand::MachineOperandType::MO_IntrinsicID)1326        Index = 2;1327      for (; Index < NumOps; ++Index)1328        MIB.add(I.getOperand(Index));1329      MIB.constrainAllUses(TII, TRI, RBI);1330      return true;1331    }1332  }1333  return false;1334}1335 1336bool SPIRVInstructionSelector::selectFrexp(Register ResVReg,1337                                           const SPIRVType *ResType,1338                                           MachineInstr &I) const {1339  ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},1340                          {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};1341  for (const auto &Ex : ExtInsts) {1342    SPIRV::InstructionSet::InstructionSet Set = Ex.first;1343    uint32_t Opcode = Ex.second;1344    if (!STI.canUseExtInstSet(Set))1345      continue;1346 1347    MachineIRBuilder MIRBuilder(I);1348    SPIRVType *PointeeTy = GR.getSPIRVTypeForVReg(I.getOperand(1).getReg());1349    const SPIRVType *PointerType = GR.getOrCreateSPIRVPointerType(1350        PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);1351    Register PointerVReg =1352        createVirtualRegister(PointerType, &GR, MRI, MRI->getMF());1353 1354    auto It = getOpVariableMBBIt(I);1355    auto MIB = BuildMI(*It->getParent(), It, It->getDebugLoc(),1356                       TII.get(SPIRV::OpVariable))1357                   .addDef(PointerVReg)1358                   .addUse(GR.getSPIRVTypeID(PointerType))1359                   .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))1360                   .constrainAllUses(TII, TRI, RBI);1361 1362    MIB = MIB &1363          BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))1364              .addDef(ResVReg)1365              .addUse(GR.getSPIRVTypeID(ResType))1366              .addImm(static_cast<uint32_t>(Ex.first))1367              .addImm(Opcode)1368              .add(I.getOperand(2))1369              .addUse(PointerVReg)1370              .constrainAllUses(TII, TRI, RBI);1371 1372    MIB = MIB &1373          BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))1374              .addDef(I.getOperand(1).getReg())1375              .addUse(GR.getSPIRVTypeID(PointeeTy))1376              .addUse(PointerVReg)1377              .constrainAllUses(TII, TRI, RBI);1378    return MIB;1379  }1380  return false;1381}1382 1383bool SPIRVInstructionSelector::selectOpWithSrcs(Register ResVReg,1384                                                const SPIRVType *ResType,1385                                                MachineInstr &I,1386                                                std::vector<Register> Srcs,1387                                                unsigned Opcode) const {1388  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))1389                 .addDef(ResVReg)1390                 .addUse(GR.getSPIRVTypeID(ResType));1391  for (Register SReg : Srcs) {1392    MIB.addUse(SReg);1393  }1394  return MIB.constrainAllUses(TII, TRI, RBI);1395}1396 1397bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,1398                                          const SPIRVType *ResType,1399                                          MachineInstr &I,1400                                          unsigned Opcode) const {1401  if (STI.isPhysicalSPIRV() && I.getOperand(1).isReg()) {1402    Register SrcReg = I.getOperand(1).getReg();1403    bool IsGV = false;1404    for (MachineRegisterInfo::def_instr_iterator DefIt =1405             MRI->def_instr_begin(SrcReg);1406         DefIt != MRI->def_instr_end(); DefIt = std::next(DefIt)) {1407      unsigned DefOpCode = DefIt->getOpcode();1408      if (DefOpCode == SPIRV::ASSIGN_TYPE) {1409        // We need special handling to look through the type assignment and see1410        // if this is a constant or a global1411        if (auto *VRD = getVRegDef(*MRI, DefIt->getOperand(1).getReg()))1412          DefOpCode = VRD->getOpcode();1413      }1414      if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||1415          DefOpCode == TargetOpcode::G_CONSTANT ||1416          DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {1417        IsGV = true;1418        break;1419      }1420    }1421    if (IsGV) {1422      uint32_t SpecOpcode = 0;1423      switch (Opcode) {1424      case SPIRV::OpConvertPtrToU:1425        SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertPtrToU);1426        break;1427      case SPIRV::OpConvertUToPtr:1428        SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertUToPtr);1429        break;1430      }1431      if (SpecOpcode)1432        return BuildMI(*I.getParent(), I, I.getDebugLoc(),1433                       TII.get(SPIRV::OpSpecConstantOp))1434            .addDef(ResVReg)1435            .addUse(GR.getSPIRVTypeID(ResType))1436            .addImm(SpecOpcode)1437            .addUse(SrcReg)1438            .constrainAllUses(TII, TRI, RBI);1439    }1440  }1441  return selectOpWithSrcs(ResVReg, ResType, I, {I.getOperand(1).getReg()},1442                          Opcode);1443}1444 1445bool SPIRVInstructionSelector::selectBitcast(Register ResVReg,1446                                             const SPIRVType *ResType,1447                                             MachineInstr &I) const {1448  Register OpReg = I.getOperand(1).getReg();1449  SPIRVType *OpType = OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;1450  if (!GR.isBitcastCompatible(ResType, OpType))1451    report_fatal_error("incompatible result and operand types in a bitcast");1452  return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitcast);1453}1454 1455static void addMemoryOperands(MachineMemOperand *MemOp,1456                              MachineInstrBuilder &MIB,1457                              MachineIRBuilder &MIRBuilder,1458                              SPIRVGlobalRegistry &GR) {1459  uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);1460  if (MemOp->isVolatile())1461    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);1462  if (MemOp->isNonTemporal())1463    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);1464  if (MemOp->getAlign().value())1465    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);1466 1467  [[maybe_unused]] MachineInstr *AliasList = nullptr;1468  [[maybe_unused]] MachineInstr *NoAliasList = nullptr;1469  const SPIRVSubtarget *ST =1470      static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());1471  if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {1472    if (auto *MD = MemOp->getAAInfo().Scope) {1473      AliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);1474      if (AliasList)1475        SpvMemOp |=1476            static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);1477    }1478    if (auto *MD = MemOp->getAAInfo().NoAlias) {1479      NoAliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);1480      if (NoAliasList)1481        SpvMemOp |=1482            static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);1483    }1484  }1485 1486  if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {1487    MIB.addImm(SpvMemOp);1488    if (SpvMemOp & static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))1489      MIB.addImm(MemOp->getAlign().value());1490    if (AliasList)1491      MIB.addUse(AliasList->getOperand(0).getReg());1492    if (NoAliasList)1493      MIB.addUse(NoAliasList->getOperand(0).getReg());1494  }1495}1496 1497static void addMemoryOperands(uint64_t Flags, MachineInstrBuilder &MIB) {1498  uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);1499  if (Flags & MachineMemOperand::Flags::MOVolatile)1500    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);1501  if (Flags & MachineMemOperand::Flags::MONonTemporal)1502    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);1503 1504  if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None))1505    MIB.addImm(SpvMemOp);1506}1507 1508bool SPIRVInstructionSelector::selectLoad(Register ResVReg,1509                                          const SPIRVType *ResType,1510                                          MachineInstr &I) const {1511  unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;1512  Register Ptr = I.getOperand(1 + OpOffset).getReg();1513 1514  auto *PtrDef = getVRegDef(*MRI, Ptr);1515  auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);1516  if (IntPtrDef &&1517      IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {1518    Register HandleReg = IntPtrDef->getOperand(2).getReg();1519    SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);1520    if (HandleType->getOpcode() == SPIRV::OpTypeImage) {1521      Register NewHandleReg =1522          MRI->createVirtualRegister(MRI->getRegClass(HandleReg));1523      auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));1524      if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {1525        return false;1526      }1527 1528      Register IdxReg = IntPtrDef->getOperand(3).getReg();1529      return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,1530                                      I.getDebugLoc(), I);1531    }1532  }1533 1534  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))1535                 .addDef(ResVReg)1536                 .addUse(GR.getSPIRVTypeID(ResType))1537                 .addUse(Ptr);1538  if (!I.getNumMemOperands()) {1539    assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||1540           I.getOpcode() ==1541               TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);1542    addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);1543  } else {1544    MachineIRBuilder MIRBuilder(I);1545    addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);1546  }1547  return MIB.constrainAllUses(TII, TRI, RBI);1548}1549 1550bool SPIRVInstructionSelector::selectStore(MachineInstr &I) const {1551  unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;1552  Register StoreVal = I.getOperand(0 + OpOffset).getReg();1553  Register Ptr = I.getOperand(1 + OpOffset).getReg();1554 1555  auto *PtrDef = getVRegDef(*MRI, Ptr);1556  auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);1557  if (IntPtrDef &&1558      IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {1559    Register HandleReg = IntPtrDef->getOperand(2).getReg();1560    Register NewHandleReg =1561        MRI->createVirtualRegister(MRI->getRegClass(HandleReg));1562    auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));1563    SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);1564    if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {1565      return false;1566    }1567 1568    Register IdxReg = IntPtrDef->getOperand(3).getReg();1569    if (HandleType->getOpcode() == SPIRV::OpTypeImage) {1570      auto BMI = BuildMI(*I.getParent(), I, I.getDebugLoc(),1571                         TII.get(SPIRV::OpImageWrite))1572                     .addUse(NewHandleReg)1573                     .addUse(IdxReg)1574                     .addUse(StoreVal);1575 1576      const llvm::Type *LLVMHandleType = GR.getTypeForSPIRVType(HandleType);1577      if (sampledTypeIsSignedInteger(LLVMHandleType))1578        BMI.addImm(0x1000); // SignExtend1579 1580      return BMI.constrainAllUses(TII, TRI, RBI);1581    }1582  }1583 1584  MachineBasicBlock &BB = *I.getParent();1585  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpStore))1586                 .addUse(Ptr)1587                 .addUse(StoreVal);1588  if (!I.getNumMemOperands()) {1589    assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||1590           I.getOpcode() ==1591               TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);1592    addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);1593  } else {1594    MachineIRBuilder MIRBuilder(I);1595    addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);1596  }1597  return MIB.constrainAllUses(TII, TRI, RBI);1598}1599 1600bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,1601                                               const SPIRVType *ResType,1602                                               MachineInstr &I) const {1603  if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))1604    report_fatal_error(1605        "llvm.stacksave intrinsic: this instruction requires the following "1606        "SPIR-V extension: SPV_INTEL_variable_length_array",1607        false);1608  MachineBasicBlock &BB = *I.getParent();1609  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSaveMemoryINTEL))1610      .addDef(ResVReg)1611      .addUse(GR.getSPIRVTypeID(ResType))1612      .constrainAllUses(TII, TRI, RBI);1613}1614 1615bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {1616  if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))1617    report_fatal_error(1618        "llvm.stackrestore intrinsic: this instruction requires the following "1619        "SPIR-V extension: SPV_INTEL_variable_length_array",1620        false);1621  if (!I.getOperand(0).isReg())1622    return false;1623  MachineBasicBlock &BB = *I.getParent();1624  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpRestoreMemoryINTEL))1625      .addUse(I.getOperand(0).getReg())1626      .constrainAllUses(TII, TRI, RBI);1627}1628 1629Register1630SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &I) const {1631  MachineIRBuilder MIRBuilder(I);1632  assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());1633 1634  // TODO: check if we have such GV, add init, use buildGlobalVariable.1635  unsigned Num = getIConstVal(I.getOperand(2).getReg(), MRI);1636  Function &CurFunction = GR.CurMF->getFunction();1637  Type *LLVMArrTy =1638      ArrayType::get(IntegerType::get(CurFunction.getContext(), 8), Num);1639  GlobalVariable *GV = new GlobalVariable(*CurFunction.getParent(), LLVMArrTy,1640                                          true, GlobalValue::InternalLinkage,1641                                          Constant::getNullValue(LLVMArrTy));1642 1643  Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());1644  Type *ArrTy = ArrayType::get(ValTy, Num);1645  SPIRVType *VarTy = GR.getOrCreateSPIRVPointerType(1646      ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);1647 1648  SPIRVType *SpvArrTy = GR.getOrCreateSPIRVType(1649      ArrTy, MIRBuilder, SPIRV::AccessQualifier::None, false);1650 1651  unsigned Val = getIConstVal(I.getOperand(1).getReg(), MRI);1652  Register Const = GR.getOrCreateConstIntArray(Val, Num, I, SpvArrTy, TII);1653 1654  Register VarReg = MRI->createGenericVirtualRegister(LLT::scalar(64));1655  auto MIBVar =1656      BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpVariable))1657          .addDef(VarReg)1658          .addUse(GR.getSPIRVTypeID(VarTy))1659          .addImm(SPIRV::StorageClass::UniformConstant)1660          .addUse(Const);1661  if (!MIBVar.constrainAllUses(TII, TRI, RBI))1662    return Register();1663 1664  GR.add(GV, MIBVar);1665  GR.addGlobalObject(GV, GR.CurMF, VarReg);1666 1667  buildOpDecorate(VarReg, I, TII, SPIRV::Decoration::Constant, {});1668  return VarReg;1669}1670 1671bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &I,1672                                                Register SrcReg) const {1673  MachineBasicBlock &BB = *I.getParent();1674  Register DstReg = I.getOperand(0).getReg();1675  SPIRVType *DstTy = GR.getSPIRVTypeForVReg(DstReg);1676  SPIRVType *SrcTy = GR.getSPIRVTypeForVReg(SrcReg);1677  if (GR.getPointeeType(DstTy) != GR.getPointeeType(SrcTy))1678    report_fatal_error("OpCopyMemory requires operands to have the same type");1679  uint64_t CopySize = getIConstVal(I.getOperand(2).getReg(), MRI);1680  SPIRVType *PointeeTy = GR.getPointeeType(DstTy);1681  const Type *LLVMPointeeTy = GR.getTypeForSPIRVType(PointeeTy);1682  if (!LLVMPointeeTy)1683    report_fatal_error(1684        "Unable to determine pointee type size for OpCopyMemory");1685  const DataLayout &DL = I.getMF()->getFunction().getDataLayout();1686  if (CopySize != DL.getTypeStoreSize(const_cast<Type *>(LLVMPointeeTy)))1687    report_fatal_error(1688        "OpCopyMemory requires the size to match the pointee type size");1689  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemory))1690                 .addUse(DstReg)1691                 .addUse(SrcReg);1692  if (I.getNumMemOperands()) {1693    MachineIRBuilder MIRBuilder(I);1694    addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);1695  }1696  return MIB.constrainAllUses(TII, TRI, RBI);1697}1698 1699bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &I,1700                                                     Register SrcReg) const {1701  MachineBasicBlock &BB = *I.getParent();1702  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemorySized))1703                 .addUse(I.getOperand(0).getReg())1704                 .addUse(SrcReg)1705                 .addUse(I.getOperand(2).getReg());1706  if (I.getNumMemOperands()) {1707    MachineIRBuilder MIRBuilder(I);1708    addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);1709  }1710  return MIB.constrainAllUses(TII, TRI, RBI);1711}1712 1713bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,1714                                                  MachineInstr &I) const {1715  Register SrcReg = I.getOperand(1).getReg();1716  bool Result = true;1717  if (I.getOpcode() == TargetOpcode::G_MEMSET) {1718    Register VarReg = getOrCreateMemSetGlobal(I);1719    if (!VarReg.isValid())1720      return false;1721    Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());1722    SPIRVType *SourceTy = GR.getOrCreateSPIRVPointerType(1723        ValTy, I, SPIRV::StorageClass::UniformConstant);1724    SrcReg = MRI->createGenericVirtualRegister(LLT::scalar(64));1725    Result &= selectOpWithSrcs(SrcReg, SourceTy, I, {VarReg}, SPIRV::OpBitcast);1726  }1727  if (STI.isLogicalSPIRV()) {1728    Result &= selectCopyMemory(I, SrcReg);1729  } else {1730    Result &= selectCopyMemorySized(I, SrcReg);1731  }1732  if (ResVReg.isValid() && ResVReg != I.getOperand(0).getReg())1733    Result &= BuildCOPY(ResVReg, I.getOperand(0).getReg(), I);1734  return Result;1735}1736 1737bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,1738                                               const SPIRVType *ResType,1739                                               MachineInstr &I,1740                                               unsigned NewOpcode,1741                                               unsigned NegateOpcode) const {1742  bool Result = true;1743  assert(I.hasOneMemOperand());1744  const MachineMemOperand *MemOp = *I.memoperands_begin();1745  uint32_t Scope = static_cast<uint32_t>(getMemScope(1746      GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));1747  auto ScopeConstant = buildI32Constant(Scope, I);1748  Register ScopeReg = ScopeConstant.first;1749  Result &= ScopeConstant.second;1750 1751  Register Ptr = I.getOperand(1).getReg();1752  // TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll1753  // auto ScSem =1754  // getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));1755  AtomicOrdering AO = MemOp->getSuccessOrdering();1756  uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));1757  auto MemSemConstant = buildI32Constant(MemSem /*| ScSem*/, I);1758  Register MemSemReg = MemSemConstant.first;1759  Result &= MemSemConstant.second;1760 1761  Register ValueReg = I.getOperand(2).getReg();1762  if (NegateOpcode != 0) {1763    // Translation with negative value operand is requested1764    Register TmpReg = createVirtualRegister(ResType, &GR, MRI, MRI->getMF());1765    Result &= selectOpWithSrcs(TmpReg, ResType, I, {ValueReg}, NegateOpcode);1766    ValueReg = TmpReg;1767  }1768 1769  return Result &&1770         BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(NewOpcode))1771             .addDef(ResVReg)1772             .addUse(GR.getSPIRVTypeID(ResType))1773             .addUse(Ptr)1774             .addUse(ScopeReg)1775             .addUse(MemSemReg)1776             .addUse(ValueReg)1777             .constrainAllUses(TII, TRI, RBI);1778}1779 1780bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &I) const {1781  unsigned ArgI = I.getNumOperands() - 1;1782  Register SrcReg =1783      I.getOperand(ArgI).isReg() ? I.getOperand(ArgI).getReg() : Register(0);1784  SPIRVType *SrcType =1785      SrcReg.isValid() ? GR.getSPIRVTypeForVReg(SrcReg) : nullptr;1786  if (!SrcType || SrcType->getOpcode() != SPIRV::OpTypeVector)1787    report_fatal_error(1788        "cannot select G_UNMERGE_VALUES with a non-vector argument");1789 1790  SPIRVType *ScalarType =1791      GR.getSPIRVTypeForVReg(SrcType->getOperand(1).getReg());1792  MachineBasicBlock &BB = *I.getParent();1793  bool Res = false;1794  unsigned CurrentIndex = 0;1795  for (unsigned i = 0; i < I.getNumDefs(); ++i) {1796    Register ResVReg = I.getOperand(i).getReg();1797    SPIRVType *ResType = GR.getSPIRVTypeForVReg(ResVReg);1798    if (!ResType) {1799      LLT ResLLT = MRI->getType(ResVReg);1800      assert(ResLLT.isValid());1801      if (ResLLT.isVector()) {1802        ResType = GR.getOrCreateSPIRVVectorType(1803            ScalarType, ResLLT.getNumElements(), I, TII);1804      } else {1805        ResType = ScalarType;1806      }1807      MRI->setRegClass(ResVReg, GR.getRegClass(ResType));1808      GR.assignSPIRVTypeToVReg(ResType, ResVReg, *GR.CurMF);1809    }1810 1811    if (ResType->getOpcode() == SPIRV::OpTypeVector) {1812      Register UndefReg = GR.getOrCreateUndef(I, SrcType, TII);1813      auto MIB =1814          BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))1815              .addDef(ResVReg)1816              .addUse(GR.getSPIRVTypeID(ResType))1817              .addUse(SrcReg)1818              .addUse(UndefReg);1819      unsigned NumElements = GR.getScalarOrVectorComponentCount(ResType);1820      for (unsigned j = 0; j < NumElements; ++j) {1821        MIB.addImm(CurrentIndex + j);1822      }1823      CurrentIndex += NumElements;1824      Res |= MIB.constrainAllUses(TII, TRI, RBI);1825    } else {1826      auto MIB =1827          BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))1828              .addDef(ResVReg)1829              .addUse(GR.getSPIRVTypeID(ResType))1830              .addUse(SrcReg)1831              .addImm(CurrentIndex);1832      CurrentIndex++;1833      Res |= MIB.constrainAllUses(TII, TRI, RBI);1834    }1835  }1836  return Res;1837}1838 1839bool SPIRVInstructionSelector::selectFence(MachineInstr &I) const {1840  AtomicOrdering AO = AtomicOrdering(I.getOperand(0).getImm());1841  uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));1842  auto MemSemConstant = buildI32Constant(MemSem, I);1843  Register MemSemReg = MemSemConstant.first;1844  bool Result = MemSemConstant.second;1845  SyncScope::ID Ord = SyncScope::ID(I.getOperand(1).getImm());1846  uint32_t Scope = static_cast<uint32_t>(1847      getMemScope(GR.CurMF->getFunction().getContext(), Ord));1848  auto ScopeConstant = buildI32Constant(Scope, I);1849  Register ScopeReg = ScopeConstant.first;1850  Result &= ScopeConstant.second;1851  MachineBasicBlock &BB = *I.getParent();1852  return Result &&1853         BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpMemoryBarrier))1854             .addUse(ScopeReg)1855             .addUse(MemSemReg)1856             .constrainAllUses(TII, TRI, RBI);1857}1858 1859bool SPIRVInstructionSelector::selectOverflowArith(Register ResVReg,1860                                                   const SPIRVType *ResType,1861                                                   MachineInstr &I,1862                                                   unsigned Opcode) const {1863  Type *ResTy = nullptr;1864  StringRef ResName;1865  if (!GR.findValueAttrs(&I, ResTy, ResName))1866    report_fatal_error(1867        "Not enough info to select the arithmetic with overflow instruction");1868  if (!ResTy || !ResTy->isStructTy())1869    report_fatal_error("Expect struct type result for the arithmetic "1870                       "with overflow instruction");1871  // "Result Type must be from OpTypeStruct. The struct must have two members,1872  // and the two members must be the same type."1873  Type *ResElemTy = cast<StructType>(ResTy)->getElementType(0);1874  ResTy = StructType::get(ResElemTy, ResElemTy);1875  // Build SPIR-V types and constant(s) if needed.1876  MachineIRBuilder MIRBuilder(I);1877  SPIRVType *StructType = GR.getOrCreateSPIRVType(1878      ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);1879  assert(I.getNumDefs() > 1 && "Not enought operands");1880  SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);1881  unsigned N = GR.getScalarOrVectorComponentCount(ResType);1882  if (N > 1)1883    BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);1884  Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);1885  Register ZeroReg = buildZerosVal(ResType, I);1886  // A new virtual register to store the result struct.1887  Register StructVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));1888  MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);1889  // Build the result name if needed.1890  if (ResName.size() > 0)1891    buildOpName(StructVReg, ResName, MIRBuilder);1892  // Build the arithmetic with overflow instruction.1893  MachineBasicBlock &BB = *I.getParent();1894  auto MIB =1895      BuildMI(BB, MIRBuilder.getInsertPt(), I.getDebugLoc(), TII.get(Opcode))1896          .addDef(StructVReg)1897          .addUse(GR.getSPIRVTypeID(StructType));1898  for (unsigned i = I.getNumDefs(); i < I.getNumOperands(); ++i)1899    MIB.addUse(I.getOperand(i).getReg());1900  bool Result = MIB.constrainAllUses(TII, TRI, RBI);1901  // Build instructions to extract fields of the instruction's result.1902  // A new virtual register to store the higher part of the result struct.1903  Register HigherVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));1904  MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);1905  for (unsigned i = 0; i < I.getNumDefs(); ++i) {1906    auto MIB =1907        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))1908            .addDef(i == 1 ? HigherVReg : I.getOperand(i).getReg())1909            .addUse(GR.getSPIRVTypeID(ResType))1910            .addUse(StructVReg)1911            .addImm(i);1912    Result &= MIB.constrainAllUses(TII, TRI, RBI);1913  }1914  // Build boolean value from the higher part.1915  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))1916                       .addDef(I.getOperand(1).getReg())1917                       .addUse(BoolTypeReg)1918                       .addUse(HigherVReg)1919                       .addUse(ZeroReg)1920                       .constrainAllUses(TII, TRI, RBI);1921}1922 1923bool SPIRVInstructionSelector::selectAtomicCmpXchg(Register ResVReg,1924                                                   const SPIRVType *ResType,1925                                                   MachineInstr &I) const {1926  bool Result = true;1927  Register ScopeReg;1928  Register MemSemEqReg;1929  Register MemSemNeqReg;1930  Register Ptr = I.getOperand(2).getReg();1931  if (!isa<GIntrinsic>(I)) {1932    assert(I.hasOneMemOperand());1933    const MachineMemOperand *MemOp = *I.memoperands_begin();1934    unsigned Scope = static_cast<uint32_t>(getMemScope(1935        GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));1936    auto ScopeConstant = buildI32Constant(Scope, I);1937    ScopeReg = ScopeConstant.first;1938    Result &= ScopeConstant.second;1939 1940    unsigned ScSem = static_cast<uint32_t>(1941        getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr)));1942    AtomicOrdering AO = MemOp->getSuccessOrdering();1943    unsigned MemSemEq = static_cast<uint32_t>(getMemSemantics(AO)) | ScSem;1944    auto MemSemEqConstant = buildI32Constant(MemSemEq, I);1945    MemSemEqReg = MemSemEqConstant.first;1946    Result &= MemSemEqConstant.second;1947    AtomicOrdering FO = MemOp->getFailureOrdering();1948    unsigned MemSemNeq = static_cast<uint32_t>(getMemSemantics(FO)) | ScSem;1949    if (MemSemEq == MemSemNeq)1950      MemSemNeqReg = MemSemEqReg;1951    else {1952      auto MemSemNeqConstant = buildI32Constant(MemSemEq, I);1953      MemSemNeqReg = MemSemNeqConstant.first;1954      Result &= MemSemNeqConstant.second;1955    }1956  } else {1957    ScopeReg = I.getOperand(5).getReg();1958    MemSemEqReg = I.getOperand(6).getReg();1959    MemSemNeqReg = I.getOperand(7).getReg();1960  }1961 1962  Register Cmp = I.getOperand(3).getReg();1963  Register Val = I.getOperand(4).getReg();1964  SPIRVType *SpvValTy = GR.getSPIRVTypeForVReg(Val);1965  Register ACmpRes = createVirtualRegister(SpvValTy, &GR, MRI, *I.getMF());1966  const DebugLoc &DL = I.getDebugLoc();1967  Result &=1968      BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpAtomicCompareExchange))1969          .addDef(ACmpRes)1970          .addUse(GR.getSPIRVTypeID(SpvValTy))1971          .addUse(Ptr)1972          .addUse(ScopeReg)1973          .addUse(MemSemEqReg)1974          .addUse(MemSemNeqReg)1975          .addUse(Val)1976          .addUse(Cmp)1977          .constrainAllUses(TII, TRI, RBI);1978  SPIRVType *BoolTy = GR.getOrCreateSPIRVBoolType(I, TII);1979  Register CmpSuccReg = createVirtualRegister(BoolTy, &GR, MRI, *I.getMF());1980  Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpIEqual))1981                .addDef(CmpSuccReg)1982                .addUse(GR.getSPIRVTypeID(BoolTy))1983                .addUse(ACmpRes)1984                .addUse(Cmp)1985                .constrainAllUses(TII, TRI, RBI);1986  Register TmpReg = createVirtualRegister(ResType, &GR, MRI, *I.getMF());1987  Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))1988                .addDef(TmpReg)1989                .addUse(GR.getSPIRVTypeID(ResType))1990                .addUse(ACmpRes)1991                .addUse(GR.getOrCreateUndef(I, ResType, TII))1992                .addImm(0)1993                .constrainAllUses(TII, TRI, RBI);1994  return Result &&1995         BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))1996             .addDef(ResVReg)1997             .addUse(GR.getSPIRVTypeID(ResType))1998             .addUse(CmpSuccReg)1999             .addUse(TmpReg)2000             .addImm(1)2001             .constrainAllUses(TII, TRI, RBI);2002}2003 2004static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC) {2005  switch (SC) {2006  case SPIRV::StorageClass::DeviceOnlyINTEL:2007  case SPIRV::StorageClass::HostOnlyINTEL:2008    return true;2009  default:2010    return false;2011  }2012}2013 2014// Returns true ResVReg is referred only from global vars and OpName's.2015static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg) {2016  bool IsGRef = false;2017  bool IsAllowedRefs =2018      llvm::all_of(MRI->use_instructions(ResVReg), [&IsGRef](auto const &It) {2019        unsigned Opcode = It.getOpcode();2020        if (Opcode == SPIRV::OpConstantComposite ||2021            Opcode == SPIRV::OpVariable ||2022            isSpvIntrinsic(It, Intrinsic::spv_init_global))2023          return IsGRef = true;2024        return Opcode == SPIRV::OpName;2025      });2026  return IsAllowedRefs && IsGRef;2027}2028 2029Register SPIRVInstructionSelector::getUcharPtrTypeReg(2030    MachineInstr &I, SPIRV::StorageClass::StorageClass SC) const {2031  return GR.getSPIRVTypeID(GR.getOrCreateSPIRVPointerType(2032      Type::getInt8Ty(I.getMF()->getFunction().getContext()), I, SC));2033}2034 2035MachineInstrBuilder2036SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &I, Register Dest,2037                                              Register Src, Register DestType,2038                                              uint32_t Opcode) const {2039  return BuildMI(*I.getParent(), I, I.getDebugLoc(),2040                 TII.get(SPIRV::OpSpecConstantOp))2041      .addDef(Dest)2042      .addUse(DestType)2043      .addImm(Opcode)2044      .addUse(Src);2045}2046 2047MachineInstrBuilder2048SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &I, Register SrcPtr,2049                                               SPIRVType *SrcPtrTy) const {2050  SPIRVType *GenericPtrTy =2051      GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);2052  Register Tmp = MRI->createVirtualRegister(&SPIRV::pIDRegClass);2053  MRI->setType(Tmp, LLT::pointer(storageClassToAddressSpace(2054                                     SPIRV::StorageClass::Generic),2055                                 GR.getPointerSize()));2056  MachineFunction *MF = I.getParent()->getParent();2057  GR.assignSPIRVTypeToVReg(GenericPtrTy, Tmp, *MF);2058  MachineInstrBuilder MIB = buildSpecConstantOp(2059      I, Tmp, SrcPtr, GR.getSPIRVTypeID(GenericPtrTy),2060      static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric));2061  GR.add(MIB.getInstr(), MIB);2062  return MIB;2063}2064 2065// In SPIR-V address space casting can only happen to and from the Generic2066// storage class. We can also only cast Workgroup, CrossWorkgroup, or Function2067// pointers to and from Generic pointers. As such, we can convert e.g. from2068// Workgroup to Function by going via a Generic pointer as an intermediary. All2069// other combinations can only be done by a bitcast, and are probably not safe.2070bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,2071                                                   const SPIRVType *ResType,2072                                                   MachineInstr &I) const {2073  MachineBasicBlock &BB = *I.getParent();2074  const DebugLoc &DL = I.getDebugLoc();2075 2076  Register SrcPtr = I.getOperand(1).getReg();2077  SPIRVType *SrcPtrTy = GR.getSPIRVTypeForVReg(SrcPtr);2078 2079  // don't generate a cast for a null that may be represented by OpTypeInt2080  if (SrcPtrTy->getOpcode() != SPIRV::OpTypePointer ||2081      ResType->getOpcode() != SPIRV::OpTypePointer)2082    return BuildCOPY(ResVReg, SrcPtr, I);2083 2084  SPIRV::StorageClass::StorageClass SrcSC = GR.getPointerStorageClass(SrcPtrTy);2085  SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResType);2086 2087  if (isASCastInGVar(MRI, ResVReg)) {2088    // AddrSpaceCast uses within OpVariable and OpConstantComposite instructions2089    // are expressed by OpSpecConstantOp with an Opcode.2090    // TODO: maybe insert a check whether the Kernel capability was declared and2091    // so PtrCastToGeneric/GenericCastToPtr are available.2092    unsigned SpecOpcode =2093        DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC)2094            ? static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric)2095            : (SrcSC == SPIRV::StorageClass::Generic &&2096                       isGenericCastablePtr(DstSC)2097                   ? static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr)2098                   : 0);2099    // TODO: OpConstantComposite expects i8*, so we are forced to forget a2100    // correct value of ResType and use general i8* instead. Maybe this should2101    // be addressed in the emit-intrinsic step to infer a correct2102    // OpConstantComposite type.2103    if (SpecOpcode) {2104      return buildSpecConstantOp(I, ResVReg, SrcPtr,2105                                 getUcharPtrTypeReg(I, DstSC), SpecOpcode)2106          .constrainAllUses(TII, TRI, RBI);2107    } else if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {2108      MachineInstrBuilder MIB = buildConstGenericPtr(I, SrcPtr, SrcPtrTy);2109      return MIB.constrainAllUses(TII, TRI, RBI) &&2110             buildSpecConstantOp(2111                 I, ResVReg, MIB->getOperand(0).getReg(),2112                 getUcharPtrTypeReg(I, DstSC),2113                 static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr))2114                 .constrainAllUses(TII, TRI, RBI);2115    }2116  }2117 2118  // don't generate a cast between identical storage classes2119  if (SrcSC == DstSC)2120    return BuildCOPY(ResVReg, SrcPtr, I);2121 2122  if ((SrcSC == SPIRV::StorageClass::Function &&2123       DstSC == SPIRV::StorageClass::Private) ||2124      (DstSC == SPIRV::StorageClass::Function &&2125       SrcSC == SPIRV::StorageClass::Private))2126    return BuildCOPY(ResVReg, SrcPtr, I);2127 2128  // Casting from an eligible pointer to Generic.2129  if (DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC))2130    return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);2131  // Casting from Generic to an eligible pointer.2132  if (SrcSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(DstSC))2133    return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);2134  // Casting between 2 eligible pointers using Generic as an intermediary.2135  if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {2136    SPIRVType *GenericPtrTy =2137        GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);2138    Register Tmp = createVirtualRegister(GenericPtrTy, &GR, MRI, MRI->getMF());2139    bool Result = BuildMI(BB, I, DL, TII.get(SPIRV::OpPtrCastToGeneric))2140                      .addDef(Tmp)2141                      .addUse(GR.getSPIRVTypeID(GenericPtrTy))2142                      .addUse(SrcPtr)2143                      .constrainAllUses(TII, TRI, RBI);2144    return Result && BuildMI(BB, I, DL, TII.get(SPIRV::OpGenericCastToPtr))2145                         .addDef(ResVReg)2146                         .addUse(GR.getSPIRVTypeID(ResType))2147                         .addUse(Tmp)2148                         .constrainAllUses(TII, TRI, RBI);2149  }2150 2151  // Check if instructions from the SPV_INTEL_usm_storage_classes extension may2152  // be applied2153  if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::CrossWorkgroup)2154    return selectUnOp(ResVReg, ResType, I,2155                      SPIRV::OpPtrCastToCrossWorkgroupINTEL);2156  if (SrcSC == SPIRV::StorageClass::CrossWorkgroup && isUSMStorageClass(DstSC))2157    return selectUnOp(ResVReg, ResType, I,2158                      SPIRV::OpCrossWorkgroupCastToPtrINTEL);2159  if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::Generic)2160    return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);2161  if (SrcSC == SPIRV::StorageClass::Generic && isUSMStorageClass(DstSC))2162    return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);2163 2164  // Bitcast for pointers requires that the address spaces must match2165  return false;2166}2167 2168static unsigned getFCmpOpcode(unsigned PredNum) {2169  auto Pred = static_cast<CmpInst::Predicate>(PredNum);2170  switch (Pred) {2171  case CmpInst::FCMP_OEQ:2172    return SPIRV::OpFOrdEqual;2173  case CmpInst::FCMP_OGE:2174    return SPIRV::OpFOrdGreaterThanEqual;2175  case CmpInst::FCMP_OGT:2176    return SPIRV::OpFOrdGreaterThan;2177  case CmpInst::FCMP_OLE:2178    return SPIRV::OpFOrdLessThanEqual;2179  case CmpInst::FCMP_OLT:2180    return SPIRV::OpFOrdLessThan;2181  case CmpInst::FCMP_ONE:2182    return SPIRV::OpFOrdNotEqual;2183  case CmpInst::FCMP_ORD:2184    return SPIRV::OpOrdered;2185  case CmpInst::FCMP_UEQ:2186    return SPIRV::OpFUnordEqual;2187  case CmpInst::FCMP_UGE:2188    return SPIRV::OpFUnordGreaterThanEqual;2189  case CmpInst::FCMP_UGT:2190    return SPIRV::OpFUnordGreaterThan;2191  case CmpInst::FCMP_ULE:2192    return SPIRV::OpFUnordLessThanEqual;2193  case CmpInst::FCMP_ULT:2194    return SPIRV::OpFUnordLessThan;2195  case CmpInst::FCMP_UNE:2196    return SPIRV::OpFUnordNotEqual;2197  case CmpInst::FCMP_UNO:2198    return SPIRV::OpUnordered;2199  default:2200    llvm_unreachable("Unknown predicate type for FCmp");2201  }2202}2203 2204static unsigned getICmpOpcode(unsigned PredNum) {2205  auto Pred = static_cast<CmpInst::Predicate>(PredNum);2206  switch (Pred) {2207  case CmpInst::ICMP_EQ:2208    return SPIRV::OpIEqual;2209  case CmpInst::ICMP_NE:2210    return SPIRV::OpINotEqual;2211  case CmpInst::ICMP_SGE:2212    return SPIRV::OpSGreaterThanEqual;2213  case CmpInst::ICMP_SGT:2214    return SPIRV::OpSGreaterThan;2215  case CmpInst::ICMP_SLE:2216    return SPIRV::OpSLessThanEqual;2217  case CmpInst::ICMP_SLT:2218    return SPIRV::OpSLessThan;2219  case CmpInst::ICMP_UGE:2220    return SPIRV::OpUGreaterThanEqual;2221  case CmpInst::ICMP_UGT:2222    return SPIRV::OpUGreaterThan;2223  case CmpInst::ICMP_ULE:2224    return SPIRV::OpULessThanEqual;2225  case CmpInst::ICMP_ULT:2226    return SPIRV::OpULessThan;2227  default:2228    llvm_unreachable("Unknown predicate type for ICmp");2229  }2230}2231 2232static unsigned getPtrCmpOpcode(unsigned Pred) {2233  switch (static_cast<CmpInst::Predicate>(Pred)) {2234  case CmpInst::ICMP_EQ:2235    return SPIRV::OpPtrEqual;2236  case CmpInst::ICMP_NE:2237    return SPIRV::OpPtrNotEqual;2238  default:2239    llvm_unreachable("Unknown predicate type for pointer comparison");2240  }2241}2242 2243// Return the logical operation, or abort if none exists.2244static unsigned getBoolCmpOpcode(unsigned PredNum) {2245  auto Pred = static_cast<CmpInst::Predicate>(PredNum);2246  switch (Pred) {2247  case CmpInst::ICMP_EQ:2248    return SPIRV::OpLogicalEqual;2249  case CmpInst::ICMP_NE:2250    return SPIRV::OpLogicalNotEqual;2251  default:2252    llvm_unreachable("Unknown predicate type for Bool comparison");2253  }2254}2255 2256static APFloat getZeroFP(const Type *LLVMFloatTy) {2257  if (!LLVMFloatTy)2258    return APFloat::getZero(APFloat::IEEEsingle());2259  switch (LLVMFloatTy->getScalarType()->getTypeID()) {2260  case Type::HalfTyID:2261    return APFloat::getZero(APFloat::IEEEhalf());2262  default:2263  case Type::FloatTyID:2264    return APFloat::getZero(APFloat::IEEEsingle());2265  case Type::DoubleTyID:2266    return APFloat::getZero(APFloat::IEEEdouble());2267  }2268}2269 2270static APFloat getOneFP(const Type *LLVMFloatTy) {2271  if (!LLVMFloatTy)2272    return APFloat::getOne(APFloat::IEEEsingle());2273  switch (LLVMFloatTy->getScalarType()->getTypeID()) {2274  case Type::HalfTyID:2275    return APFloat::getOne(APFloat::IEEEhalf());2276  default:2277  case Type::FloatTyID:2278    return APFloat::getOne(APFloat::IEEEsingle());2279  case Type::DoubleTyID:2280    return APFloat::getOne(APFloat::IEEEdouble());2281  }2282}2283 2284bool SPIRVInstructionSelector::selectAnyOrAll(Register ResVReg,2285                                              const SPIRVType *ResType,2286                                              MachineInstr &I,2287                                              unsigned OpAnyOrAll) const {2288  assert(I.getNumOperands() == 3);2289  assert(I.getOperand(2).isReg());2290  MachineBasicBlock &BB = *I.getParent();2291  Register InputRegister = I.getOperand(2).getReg();2292  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);2293 2294  if (!InputType)2295    report_fatal_error("Input Type could not be determined.");2296 2297  bool IsBoolTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeBool);2298  bool IsVectorTy = InputType->getOpcode() == SPIRV::OpTypeVector;2299  if (IsBoolTy && !IsVectorTy) {2300    assert(ResVReg == I.getOperand(0).getReg());2301    return BuildCOPY(ResVReg, InputRegister, I);2302  }2303 2304  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);2305  unsigned SpirvNotEqualId =2306      IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;2307  SPIRVType *SpvBoolScalarTy = GR.getOrCreateSPIRVBoolType(I, TII);2308  SPIRVType *SpvBoolTy = SpvBoolScalarTy;2309  Register NotEqualReg = ResVReg;2310 2311  if (IsVectorTy) {2312    NotEqualReg =2313        IsBoolTy ? InputRegister2314                 : createVirtualRegister(SpvBoolTy, &GR, MRI, MRI->getMF());2315    const unsigned NumElts = InputType->getOperand(2).getImm();2316    SpvBoolTy = GR.getOrCreateSPIRVVectorType(SpvBoolTy, NumElts, I, TII);2317  }2318 2319  bool Result = true;2320  if (!IsBoolTy) {2321    Register ConstZeroReg =2322        IsFloatTy ? buildZerosValF(InputType, I) : buildZerosVal(InputType, I);2323 2324    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SpirvNotEqualId))2325                  .addDef(NotEqualReg)2326                  .addUse(GR.getSPIRVTypeID(SpvBoolTy))2327                  .addUse(InputRegister)2328                  .addUse(ConstZeroReg)2329                  .constrainAllUses(TII, TRI, RBI);2330  }2331 2332  if (!IsVectorTy)2333    return Result;2334 2335  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(OpAnyOrAll))2336                       .addDef(ResVReg)2337                       .addUse(GR.getSPIRVTypeID(SpvBoolScalarTy))2338                       .addUse(NotEqualReg)2339                       .constrainAllUses(TII, TRI, RBI);2340}2341 2342bool SPIRVInstructionSelector::selectAll(Register ResVReg,2343                                         const SPIRVType *ResType,2344                                         MachineInstr &I) const {2345  return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAll);2346}2347 2348bool SPIRVInstructionSelector::selectAny(Register ResVReg,2349                                         const SPIRVType *ResType,2350                                         MachineInstr &I) const {2351  return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAny);2352}2353 2354// Select the OpDot instruction for the given float dot2355bool SPIRVInstructionSelector::selectFloatDot(Register ResVReg,2356                                              const SPIRVType *ResType,2357                                              MachineInstr &I) const {2358  assert(I.getNumOperands() == 4);2359  assert(I.getOperand(2).isReg());2360  assert(I.getOperand(3).isReg());2361 2362  [[maybe_unused]] SPIRVType *VecType =2363      GR.getSPIRVTypeForVReg(I.getOperand(2).getReg());2364 2365  assert(VecType->getOpcode() == SPIRV::OpTypeVector &&2366         GR.getScalarOrVectorComponentCount(VecType) > 1 &&2367         "dot product requires a vector of at least 2 components");2368 2369  [[maybe_unused]] SPIRVType *EltType =2370      GR.getSPIRVTypeForVReg(VecType->getOperand(1).getReg());2371 2372  assert(EltType->getOpcode() == SPIRV::OpTypeFloat);2373 2374  MachineBasicBlock &BB = *I.getParent();2375  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpDot))2376      .addDef(ResVReg)2377      .addUse(GR.getSPIRVTypeID(ResType))2378      .addUse(I.getOperand(2).getReg())2379      .addUse(I.getOperand(3).getReg())2380      .constrainAllUses(TII, TRI, RBI);2381}2382 2383bool SPIRVInstructionSelector::selectIntegerDot(Register ResVReg,2384                                                const SPIRVType *ResType,2385                                                MachineInstr &I,2386                                                bool Signed) const {2387  assert(I.getNumOperands() == 4);2388  assert(I.getOperand(2).isReg());2389  assert(I.getOperand(3).isReg());2390  MachineBasicBlock &BB = *I.getParent();2391 2392  auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;2393  return BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))2394      .addDef(ResVReg)2395      .addUse(GR.getSPIRVTypeID(ResType))2396      .addUse(I.getOperand(2).getReg())2397      .addUse(I.getOperand(3).getReg())2398      .constrainAllUses(TII, TRI, RBI);2399}2400 2401// Since pre-1.6 SPIRV has no integer dot implementation,2402// expand by piecewise multiplying and adding the results2403bool SPIRVInstructionSelector::selectIntegerDotExpansion(2404    Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {2405  assert(I.getNumOperands() == 4);2406  assert(I.getOperand(2).isReg());2407  assert(I.getOperand(3).isReg());2408  MachineBasicBlock &BB = *I.getParent();2409 2410  // Multiply the vectors, then sum the results2411  Register Vec0 = I.getOperand(2).getReg();2412  Register Vec1 = I.getOperand(3).getReg();2413  Register TmpVec = MRI->createVirtualRegister(GR.getRegClass(ResType));2414  SPIRVType *VecType = GR.getSPIRVTypeForVReg(Vec0);2415 2416  bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulV))2417                    .addDef(TmpVec)2418                    .addUse(GR.getSPIRVTypeID(VecType))2419                    .addUse(Vec0)2420                    .addUse(Vec1)2421                    .constrainAllUses(TII, TRI, RBI);2422 2423  assert(VecType->getOpcode() == SPIRV::OpTypeVector &&2424         GR.getScalarOrVectorComponentCount(VecType) > 1 &&2425         "dot product requires a vector of at least 2 components");2426 2427  Register Res = MRI->createVirtualRegister(GR.getRegClass(ResType));2428  Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))2429                .addDef(Res)2430                .addUse(GR.getSPIRVTypeID(ResType))2431                .addUse(TmpVec)2432                .addImm(0)2433                .constrainAllUses(TII, TRI, RBI);2434 2435  for (unsigned i = 1; i < GR.getScalarOrVectorComponentCount(VecType); i++) {2436    Register Elt = MRI->createVirtualRegister(GR.getRegClass(ResType));2437 2438    Result &=2439        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))2440            .addDef(Elt)2441            .addUse(GR.getSPIRVTypeID(ResType))2442            .addUse(TmpVec)2443            .addImm(i)2444            .constrainAllUses(TII, TRI, RBI);2445 2446    Register Sum = i < GR.getScalarOrVectorComponentCount(VecType) - 12447                       ? MRI->createVirtualRegister(GR.getRegClass(ResType))2448                       : ResVReg;2449 2450    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))2451                  .addDef(Sum)2452                  .addUse(GR.getSPIRVTypeID(ResType))2453                  .addUse(Res)2454                  .addUse(Elt)2455                  .constrainAllUses(TII, TRI, RBI);2456    Res = Sum;2457  }2458 2459  return Result;2460}2461 2462bool SPIRVInstructionSelector::selectOpIsInf(Register ResVReg,2463                                             const SPIRVType *ResType,2464                                             MachineInstr &I) const {2465  MachineBasicBlock &BB = *I.getParent();2466  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIsInf))2467      .addDef(ResVReg)2468      .addUse(GR.getSPIRVTypeID(ResType))2469      .addUse(I.getOperand(2).getReg())2470      .constrainAllUses(TII, TRI, RBI);2471}2472 2473bool SPIRVInstructionSelector::selectOpIsNan(Register ResVReg,2474                                             const SPIRVType *ResType,2475                                             MachineInstr &I) const {2476  MachineBasicBlock &BB = *I.getParent();2477  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIsNan))2478      .addDef(ResVReg)2479      .addUse(GR.getSPIRVTypeID(ResType))2480      .addUse(I.getOperand(2).getReg())2481      .constrainAllUses(TII, TRI, RBI);2482}2483 2484template <bool Signed>2485bool SPIRVInstructionSelector::selectDot4AddPacked(Register ResVReg,2486                                                   const SPIRVType *ResType,2487                                                   MachineInstr &I) const {2488  assert(I.getNumOperands() == 5);2489  assert(I.getOperand(2).isReg());2490  assert(I.getOperand(3).isReg());2491  assert(I.getOperand(4).isReg());2492  MachineBasicBlock &BB = *I.getParent();2493 2494  Register Acc = I.getOperand(2).getReg();2495  Register X = I.getOperand(3).getReg();2496  Register Y = I.getOperand(4).getReg();2497 2498  auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;2499  Register Dot = MRI->createVirtualRegister(GR.getRegClass(ResType));2500  bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))2501                    .addDef(Dot)2502                    .addUse(GR.getSPIRVTypeID(ResType))2503                    .addUse(X)2504                    .addUse(Y)2505                    .constrainAllUses(TII, TRI, RBI);2506 2507  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))2508                       .addDef(ResVReg)2509                       .addUse(GR.getSPIRVTypeID(ResType))2510                       .addUse(Dot)2511                       .addUse(Acc)2512                       .constrainAllUses(TII, TRI, RBI);2513}2514 2515// Since pre-1.6 SPIRV has no DotProductInput4x8BitPacked implementation,2516// extract the elements of the packed inputs, multiply them and add the result2517// to the accumulator.2518template <bool Signed>2519bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(2520    Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {2521  assert(I.getNumOperands() == 5);2522  assert(I.getOperand(2).isReg());2523  assert(I.getOperand(3).isReg());2524  assert(I.getOperand(4).isReg());2525  MachineBasicBlock &BB = *I.getParent();2526 2527  bool Result = true;2528 2529  Register Acc = I.getOperand(2).getReg();2530  Register X = I.getOperand(3).getReg();2531  Register Y = I.getOperand(4).getReg();2532 2533  SPIRVType *EltType = GR.getOrCreateSPIRVIntegerType(8, I, TII);2534  auto ExtractOp =2535      Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;2536 2537  bool ZeroAsNull = !STI.isShader();2538  // Extract the i8 element, multiply and add it to the accumulator2539  for (unsigned i = 0; i < 4; i++) {2540    // A[i]2541    Register AElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);2542    Result &=2543        BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))2544            .addDef(AElt)2545            .addUse(GR.getSPIRVTypeID(ResType))2546            .addUse(X)2547            .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))2548            .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))2549            .constrainAllUses(TII, TRI, RBI);2550 2551    // B[i]2552    Register BElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);2553    Result &=2554        BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))2555            .addDef(BElt)2556            .addUse(GR.getSPIRVTypeID(ResType))2557            .addUse(Y)2558            .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))2559            .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))2560            .constrainAllUses(TII, TRI, RBI);2561 2562    // A[i] * B[i]2563    Register Mul = MRI->createVirtualRegister(&SPIRV::IDRegClass);2564    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulS))2565                  .addDef(Mul)2566                  .addUse(GR.getSPIRVTypeID(ResType))2567                  .addUse(AElt)2568                  .addUse(BElt)2569                  .constrainAllUses(TII, TRI, RBI);2570 2571    // Discard 24 highest-bits so that stored i32 register is i8 equivalent2572    Register MaskMul = MRI->createVirtualRegister(&SPIRV::IDRegClass);2573    Result &=2574        BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))2575            .addDef(MaskMul)2576            .addUse(GR.getSPIRVTypeID(ResType))2577            .addUse(Mul)2578            .addUse(GR.getOrCreateConstInt(0, I, EltType, TII, ZeroAsNull))2579            .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))2580            .constrainAllUses(TII, TRI, RBI);2581 2582    // Acc = Acc + A[i] * B[i]2583    Register Sum =2584        i < 3 ? MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;2585    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))2586                  .addDef(Sum)2587                  .addUse(GR.getSPIRVTypeID(ResType))2588                  .addUse(Acc)2589                  .addUse(MaskMul)2590                  .constrainAllUses(TII, TRI, RBI);2591 2592    Acc = Sum;2593  }2594 2595  return Result;2596}2597 2598/// Transform saturate(x) to clamp(x, 0.0f, 1.0f) as SPIRV2599/// does not have a saturate builtin.2600bool SPIRVInstructionSelector::selectSaturate(Register ResVReg,2601                                              const SPIRVType *ResType,2602                                              MachineInstr &I) const {2603  assert(I.getNumOperands() == 3);2604  assert(I.getOperand(2).isReg());2605  MachineBasicBlock &BB = *I.getParent();2606  Register VZero = buildZerosValF(ResType, I);2607  Register VOne = buildOnesValF(ResType, I);2608 2609  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))2610      .addDef(ResVReg)2611      .addUse(GR.getSPIRVTypeID(ResType))2612      .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))2613      .addImm(GL::FClamp)2614      .addUse(I.getOperand(2).getReg())2615      .addUse(VZero)2616      .addUse(VOne)2617      .constrainAllUses(TII, TRI, RBI);2618}2619 2620bool SPIRVInstructionSelector::selectSign(Register ResVReg,2621                                          const SPIRVType *ResType,2622                                          MachineInstr &I) const {2623  assert(I.getNumOperands() == 3);2624  assert(I.getOperand(2).isReg());2625  MachineBasicBlock &BB = *I.getParent();2626  Register InputRegister = I.getOperand(2).getReg();2627  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);2628  auto &DL = I.getDebugLoc();2629 2630  if (!InputType)2631    report_fatal_error("Input Type could not be determined.");2632 2633  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);2634 2635  unsigned SignBitWidth = GR.getScalarOrVectorBitWidth(InputType);2636  unsigned ResBitWidth = GR.getScalarOrVectorBitWidth(ResType);2637 2638  bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;2639 2640  auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;2641  Register SignReg = NeedsConversion2642                         ? MRI->createVirtualRegister(&SPIRV::IDRegClass)2643                         : ResVReg;2644 2645  bool Result =2646      BuildMI(BB, I, DL, TII.get(SPIRV::OpExtInst))2647          .addDef(SignReg)2648          .addUse(GR.getSPIRVTypeID(InputType))2649          .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))2650          .addImm(SignOpcode)2651          .addUse(InputRegister)2652          .constrainAllUses(TII, TRI, RBI);2653 2654  if (NeedsConversion) {2655    auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;2656    Result &= BuildMI(*I.getParent(), I, DL, TII.get(ConvertOpcode))2657                  .addDef(ResVReg)2658                  .addUse(GR.getSPIRVTypeID(ResType))2659                  .addUse(SignReg)2660                  .constrainAllUses(TII, TRI, RBI);2661  }2662 2663  return Result;2664}2665 2666bool SPIRVInstructionSelector::selectWaveOpInst(Register ResVReg,2667                                                const SPIRVType *ResType,2668                                                MachineInstr &I,2669                                                unsigned Opcode) const {2670  MachineBasicBlock &BB = *I.getParent();2671  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);2672 2673  auto BMI = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))2674                 .addDef(ResVReg)2675                 .addUse(GR.getSPIRVTypeID(ResType))2676                 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I,2677                                                IntTy, TII, !STI.isShader()));2678 2679  for (unsigned J = 2; J < I.getNumOperands(); J++) {2680    BMI.addUse(I.getOperand(J).getReg());2681  }2682 2683  return BMI.constrainAllUses(TII, TRI, RBI);2684}2685 2686bool SPIRVInstructionSelector::selectWaveActiveCountBits(2687    Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {2688 2689  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);2690  SPIRVType *BallotType = GR.getOrCreateSPIRVVectorType(IntTy, 4, I, TII);2691  Register BallotReg = MRI->createVirtualRegister(GR.getRegClass(BallotType));2692  bool Result = selectWaveOpInst(BallotReg, BallotType, I,2693                                 SPIRV::OpGroupNonUniformBallot);2694 2695  MachineBasicBlock &BB = *I.getParent();2696  Result &= BuildMI(BB, I, I.getDebugLoc(),2697                    TII.get(SPIRV::OpGroupNonUniformBallotBitCount))2698                .addDef(ResVReg)2699                .addUse(GR.getSPIRVTypeID(ResType))2700                .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy,2701                                               TII, !STI.isShader()))2702                .addImm(SPIRV::GroupOperation::Reduce)2703                .addUse(BallotReg)2704                .constrainAllUses(TII, TRI, RBI);2705 2706  return Result;2707}2708 2709bool SPIRVInstructionSelector::selectWaveReduceMax(Register ResVReg,2710                                                   const SPIRVType *ResType,2711                                                   MachineInstr &I,2712                                                   bool IsUnsigned) const {2713  assert(I.getNumOperands() == 3);2714  assert(I.getOperand(2).isReg());2715  MachineBasicBlock &BB = *I.getParent();2716  Register InputRegister = I.getOperand(2).getReg();2717  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);2718 2719  if (!InputType)2720    report_fatal_error("Input Type could not be determined.");2721 2722  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);2723  // Retreive the operation to use based on input type2724  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);2725  auto IntegerOpcodeType =2726      IsUnsigned ? SPIRV::OpGroupNonUniformUMax : SPIRV::OpGroupNonUniformSMax;2727  auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntegerOpcodeType;2728  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))2729      .addDef(ResVReg)2730      .addUse(GR.getSPIRVTypeID(ResType))2731      .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,2732                                     !STI.isShader()))2733      .addImm(SPIRV::GroupOperation::Reduce)2734      .addUse(I.getOperand(2).getReg())2735      .constrainAllUses(TII, TRI, RBI);2736}2737 2738bool SPIRVInstructionSelector::selectWaveReduceMin(Register ResVReg,2739                                                   const SPIRVType *ResType,2740                                                   MachineInstr &I,2741                                                   bool IsUnsigned) const {2742  assert(I.getNumOperands() == 3);2743  assert(I.getOperand(2).isReg());2744  MachineBasicBlock &BB = *I.getParent();2745  Register InputRegister = I.getOperand(2).getReg();2746  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);2747 2748  if (!InputType)2749    report_fatal_error("Input Type could not be determined.");2750 2751  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);2752  // Retreive the operation to use based on input type2753  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);2754  auto IntegerOpcodeType =2755      IsUnsigned ? SPIRV::OpGroupNonUniformUMin : SPIRV::OpGroupNonUniformSMin;2756  auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntegerOpcodeType;2757  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))2758      .addDef(ResVReg)2759      .addUse(GR.getSPIRVTypeID(ResType))2760      .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,2761                                     !STI.isShader()))2762      .addImm(SPIRV::GroupOperation::Reduce)2763      .addUse(I.getOperand(2).getReg())2764      .constrainAllUses(TII, TRI, RBI);2765}2766 2767bool SPIRVInstructionSelector::selectWaveReduceSum(Register ResVReg,2768                                                   const SPIRVType *ResType,2769                                                   MachineInstr &I) const {2770  assert(I.getNumOperands() == 3);2771  assert(I.getOperand(2).isReg());2772  MachineBasicBlock &BB = *I.getParent();2773  Register InputRegister = I.getOperand(2).getReg();2774  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);2775 2776  if (!InputType)2777    report_fatal_error("Input Type could not be determined.");2778 2779  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);2780  // Retreive the operation to use based on input type2781  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);2782  auto Opcode =2783      IsFloatTy ? SPIRV::OpGroupNonUniformFAdd : SPIRV::OpGroupNonUniformIAdd;2784  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))2785      .addDef(ResVReg)2786      .addUse(GR.getSPIRVTypeID(ResType))2787      .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,2788                                     !STI.isShader()))2789      .addImm(SPIRV::GroupOperation::Reduce)2790      .addUse(I.getOperand(2).getReg());2791}2792 2793bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,2794                                                const SPIRVType *ResType,2795                                                MachineInstr &I) const {2796  MachineBasicBlock &BB = *I.getParent();2797  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitReverse))2798      .addDef(ResVReg)2799      .addUse(GR.getSPIRVTypeID(ResType))2800      .addUse(I.getOperand(1).getReg())2801      .constrainAllUses(TII, TRI, RBI);2802}2803 2804bool SPIRVInstructionSelector::selectFreeze(Register ResVReg,2805                                            const SPIRVType *ResType,2806                                            MachineInstr &I) const {2807  // There is no way to implement `freeze` correctly without support on SPIR-V2808  // standard side, but we may at least address a simple (static) case when2809  // undef/poison value presence is obvious. The main benefit of even2810  // incomplete `freeze` support is preventing of translation from crashing due2811  // to lack of support on legalization and instruction selection steps.2812  if (!I.getOperand(0).isReg() || !I.getOperand(1).isReg())2813    return false;2814  Register OpReg = I.getOperand(1).getReg();2815  if (MachineInstr *Def = MRI->getVRegDef(OpReg)) {2816    if (Def->getOpcode() == TargetOpcode::COPY)2817      Def = MRI->getVRegDef(Def->getOperand(1).getReg());2818    Register Reg;2819    switch (Def->getOpcode()) {2820    case SPIRV::ASSIGN_TYPE:2821      if (MachineInstr *AssignToDef =2822              MRI->getVRegDef(Def->getOperand(1).getReg())) {2823        if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)2824          Reg = Def->getOperand(2).getReg();2825      }2826      break;2827    case SPIRV::OpUndef:2828      Reg = Def->getOperand(1).getReg();2829      break;2830    }2831    unsigned DestOpCode;2832    if (Reg.isValid()) {2833      DestOpCode = SPIRV::OpConstantNull;2834    } else {2835      DestOpCode = TargetOpcode::COPY;2836      Reg = OpReg;2837    }2838    return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DestOpCode))2839        .addDef(I.getOperand(0).getReg())2840        .addUse(Reg)2841        .constrainAllUses(TII, TRI, RBI);2842  }2843  return false;2844}2845 2846bool SPIRVInstructionSelector::selectBuildVector(Register ResVReg,2847                                                 const SPIRVType *ResType,2848                                                 MachineInstr &I) const {2849  unsigned N = 0;2850  if (ResType->getOpcode() == SPIRV::OpTypeVector)2851    N = GR.getScalarOrVectorComponentCount(ResType);2852  else if (ResType->getOpcode() == SPIRV::OpTypeArray)2853    N = getArrayComponentCount(MRI, ResType);2854  else2855    report_fatal_error("Cannot select G_BUILD_VECTOR with a non-vector result");2856  if (I.getNumExplicitOperands() - I.getNumExplicitDefs() != N)2857    report_fatal_error("G_BUILD_VECTOR and the result type are inconsistent");2858 2859  // check if we may construct a constant vector2860  bool IsConst = true;2861  for (unsigned i = I.getNumExplicitDefs();2862       i < I.getNumExplicitOperands() && IsConst; ++i)2863    if (!isConstReg(MRI, I.getOperand(i).getReg()))2864      IsConst = false;2865 2866  if (!IsConst && N < 2)2867    report_fatal_error(2868        "There must be at least two constituent operands in a vector");2869 2870  MRI->setRegClass(ResVReg, GR.getRegClass(ResType));2871  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),2872                     TII.get(IsConst ? SPIRV::OpConstantComposite2873                                     : SPIRV::OpCompositeConstruct))2874                 .addDef(ResVReg)2875                 .addUse(GR.getSPIRVTypeID(ResType));2876  for (unsigned i = I.getNumExplicitDefs(); i < I.getNumExplicitOperands(); ++i)2877    MIB.addUse(I.getOperand(i).getReg());2878  return MIB.constrainAllUses(TII, TRI, RBI);2879}2880 2881bool SPIRVInstructionSelector::selectSplatVector(Register ResVReg,2882                                                 const SPIRVType *ResType,2883                                                 MachineInstr &I) const {2884  unsigned N = 0;2885  if (ResType->getOpcode() == SPIRV::OpTypeVector)2886    N = GR.getScalarOrVectorComponentCount(ResType);2887  else if (ResType->getOpcode() == SPIRV::OpTypeArray)2888    N = getArrayComponentCount(MRI, ResType);2889  else2890    report_fatal_error("Cannot select G_SPLAT_VECTOR with a non-vector result");2891 2892  unsigned OpIdx = I.getNumExplicitDefs();2893  if (!I.getOperand(OpIdx).isReg())2894    report_fatal_error("Unexpected argument in G_SPLAT_VECTOR");2895 2896  // check if we may construct a constant vector2897  Register OpReg = I.getOperand(OpIdx).getReg();2898  bool IsConst = isConstReg(MRI, OpReg);2899 2900  if (!IsConst && N < 2)2901    report_fatal_error(2902        "There must be at least two constituent operands in a vector");2903 2904  MRI->setRegClass(ResVReg, GR.getRegClass(ResType));2905  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),2906                     TII.get(IsConst ? SPIRV::OpConstantComposite2907                                     : SPIRV::OpCompositeConstruct))2908                 .addDef(ResVReg)2909                 .addUse(GR.getSPIRVTypeID(ResType));2910  for (unsigned i = 0; i < N; ++i)2911    MIB.addUse(OpReg);2912  return MIB.constrainAllUses(TII, TRI, RBI);2913}2914 2915bool SPIRVInstructionSelector::selectDiscard(Register ResVReg,2916                                             const SPIRVType *ResType,2917                                             MachineInstr &I) const {2918 2919  unsigned Opcode;2920 2921  if (STI.canUseExtension(2922          SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||2923      STI.isAtLeastSPIRVVer(llvm::VersionTuple(1, 6))) {2924    Opcode = SPIRV::OpDemoteToHelperInvocation;2925  } else {2926    Opcode = SPIRV::OpKill;2927    // OpKill must be the last operation of any basic block.2928    if (MachineInstr *NextI = I.getNextNode()) {2929      GR.invalidateMachineInstr(NextI);2930      NextI->removeFromParent();2931    }2932  }2933 2934  MachineBasicBlock &BB = *I.getParent();2935  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))2936      .constrainAllUses(TII, TRI, RBI);2937}2938 2939bool SPIRVInstructionSelector::selectCmp(Register ResVReg,2940                                         const SPIRVType *ResType,2941                                         unsigned CmpOpc,2942                                         MachineInstr &I) const {2943  Register Cmp0 = I.getOperand(2).getReg();2944  Register Cmp1 = I.getOperand(3).getReg();2945  assert(GR.getSPIRVTypeForVReg(Cmp0)->getOpcode() ==2946             GR.getSPIRVTypeForVReg(Cmp1)->getOpcode() &&2947         "CMP operands should have the same type");2948  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(CmpOpc))2949      .addDef(ResVReg)2950      .addUse(GR.getSPIRVTypeID(ResType))2951      .addUse(Cmp0)2952      .addUse(Cmp1)2953      .setMIFlags(I.getFlags())2954      .constrainAllUses(TII, TRI, RBI);2955}2956 2957bool SPIRVInstructionSelector::selectICmp(Register ResVReg,2958                                          const SPIRVType *ResType,2959                                          MachineInstr &I) const {2960  auto Pred = I.getOperand(1).getPredicate();2961  unsigned CmpOpc;2962 2963  Register CmpOperand = I.getOperand(2).getReg();2964  if (GR.isScalarOfType(CmpOperand, SPIRV::OpTypePointer))2965    CmpOpc = getPtrCmpOpcode(Pred);2966  else if (GR.isScalarOrVectorOfType(CmpOperand, SPIRV::OpTypeBool))2967    CmpOpc = getBoolCmpOpcode(Pred);2968  else2969    CmpOpc = getICmpOpcode(Pred);2970  return selectCmp(ResVReg, ResType, CmpOpc, I);2971}2972 2973std::pair<Register, bool>2974SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &I,2975                                           const SPIRVType *ResType) const {2976  Type *LLVMTy = IntegerType::get(GR.CurMF->getFunction().getContext(), 32);2977  const SPIRVType *SpvI32Ty =2978      ResType ? ResType : GR.getOrCreateSPIRVIntegerType(32, I, TII);2979  // Find a constant in DT or build a new one.2980  auto ConstInt = ConstantInt::get(LLVMTy, Val);2981  Register NewReg = GR.find(ConstInt, GR.CurMF);2982  bool Result = true;2983  if (!NewReg.isValid()) {2984    NewReg = MRI->createGenericVirtualRegister(LLT::scalar(64));2985    MachineBasicBlock &BB = *I.getParent();2986    MachineInstr *MI =2987        Val == 02988            ? BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))2989                  .addDef(NewReg)2990                  .addUse(GR.getSPIRVTypeID(SpvI32Ty))2991            : BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantI))2992                  .addDef(NewReg)2993                  .addUse(GR.getSPIRVTypeID(SpvI32Ty))2994                  .addImm(APInt(32, Val).getZExtValue());2995    Result &= constrainSelectedInstRegOperands(*MI, TII, TRI, RBI);2996    GR.add(ConstInt, MI);2997  }2998  return {NewReg, Result};2999}3000 3001bool SPIRVInstructionSelector::selectFCmp(Register ResVReg,3002                                          const SPIRVType *ResType,3003                                          MachineInstr &I) const {3004  unsigned CmpOp = getFCmpOpcode(I.getOperand(1).getPredicate());3005  return selectCmp(ResVReg, ResType, CmpOp, I);3006}3007 3008Register SPIRVInstructionSelector::buildZerosVal(const SPIRVType *ResType,3009                                                 MachineInstr &I) const {3010  // OpenCL uses nulls for Zero. In HLSL we don't use null constants.3011  bool ZeroAsNull = !STI.isShader();3012  if (ResType->getOpcode() == SPIRV::OpTypeVector)3013    return GR.getOrCreateConstVector(0UL, I, ResType, TII, ZeroAsNull);3014  return GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);3015}3016 3017Register SPIRVInstructionSelector::buildZerosValF(const SPIRVType *ResType,3018                                                  MachineInstr &I) const {3019  // OpenCL uses nulls for Zero. In HLSL we don't use null constants.3020  bool ZeroAsNull = !STI.isShader();3021  APFloat VZero = getZeroFP(GR.getTypeForSPIRVType(ResType));3022  if (ResType->getOpcode() == SPIRV::OpTypeVector)3023    return GR.getOrCreateConstVector(VZero, I, ResType, TII, ZeroAsNull);3024  return GR.getOrCreateConstFP(VZero, I, ResType, TII, ZeroAsNull);3025}3026 3027Register SPIRVInstructionSelector::buildOnesValF(const SPIRVType *ResType,3028                                                 MachineInstr &I) const {3029  // OpenCL uses nulls for Zero. In HLSL we don't use null constants.3030  bool ZeroAsNull = !STI.isShader();3031  APFloat VOne = getOneFP(GR.getTypeForSPIRVType(ResType));3032  if (ResType->getOpcode() == SPIRV::OpTypeVector)3033    return GR.getOrCreateConstVector(VOne, I, ResType, TII, ZeroAsNull);3034  return GR.getOrCreateConstFP(VOne, I, ResType, TII, ZeroAsNull);3035}3036 3037Register SPIRVInstructionSelector::buildOnesVal(bool AllOnes,3038                                                const SPIRVType *ResType,3039                                                MachineInstr &I) const {3040  unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);3041  APInt One =3042      AllOnes ? APInt::getAllOnes(BitWidth) : APInt::getOneBitSet(BitWidth, 0);3043  if (ResType->getOpcode() == SPIRV::OpTypeVector)3044    return GR.getOrCreateConstVector(One.getZExtValue(), I, ResType, TII);3045  return GR.getOrCreateConstInt(One.getZExtValue(), I, ResType, TII);3046}3047 3048bool SPIRVInstructionSelector::selectSelect(Register ResVReg,3049                                            const SPIRVType *ResType,3050                                            MachineInstr &I) const {3051  Register SelectFirstArg = I.getOperand(2).getReg();3052  Register SelectSecondArg = I.getOperand(3).getReg();3053  assert(ResType == GR.getSPIRVTypeForVReg(SelectFirstArg) &&3054         ResType == GR.getSPIRVTypeForVReg(SelectSecondArg));3055 3056  bool IsFloatTy =3057      GR.isScalarOrVectorOfType(SelectFirstArg, SPIRV::OpTypeFloat);3058  bool IsPtrTy =3059      GR.isScalarOrVectorOfType(SelectFirstArg, SPIRV::OpTypePointer);3060  bool IsVectorTy = GR.getSPIRVTypeForVReg(SelectFirstArg)->getOpcode() ==3061                    SPIRV::OpTypeVector;3062 3063  bool IsScalarBool =3064      GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);3065  unsigned Opcode;3066  if (IsVectorTy) {3067    if (IsFloatTy) {3068      Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;3069    } else if (IsPtrTy) {3070      Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;3071    } else {3072      Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;3073    }3074  } else {3075    if (IsFloatTy) {3076      Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;3077    } else if (IsPtrTy) {3078      Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;3079    } else {3080      Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;3081    }3082  }3083  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))3084      .addDef(ResVReg)3085      .addUse(GR.getSPIRVTypeID(ResType))3086      .addUse(I.getOperand(1).getReg())3087      .addUse(SelectFirstArg)3088      .addUse(SelectSecondArg)3089      .constrainAllUses(TII, TRI, RBI);3090}3091 3092bool SPIRVInstructionSelector::selectSelectDefaultArgs(Register ResVReg,3093                                                       const SPIRVType *ResType,3094                                                       MachineInstr &I,3095                                                       bool IsSigned) const {3096  // To extend a bool, we need to use OpSelect between constants.3097  Register ZeroReg = buildZerosVal(ResType, I);3098  Register OneReg = buildOnesVal(IsSigned, ResType, I);3099  bool IsScalarBool =3100      GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);3101  unsigned Opcode =3102      IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;3103  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))3104      .addDef(ResVReg)3105      .addUse(GR.getSPIRVTypeID(ResType))3106      .addUse(I.getOperand(1).getReg())3107      .addUse(OneReg)3108      .addUse(ZeroReg)3109      .constrainAllUses(TII, TRI, RBI);3110}3111 3112bool SPIRVInstructionSelector::selectIToF(Register ResVReg,3113                                          const SPIRVType *ResType,3114                                          MachineInstr &I, bool IsSigned,3115                                          unsigned Opcode) const {3116  Register SrcReg = I.getOperand(1).getReg();3117  // We can convert bool value directly to float type without OpConvert*ToF,3118  // however the translator generates OpSelect+OpConvert*ToF, so we do the same.3119  if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {3120    unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);3121    SPIRVType *TmpType = GR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);3122    if (ResType->getOpcode() == SPIRV::OpTypeVector) {3123      const unsigned NumElts = ResType->getOperand(2).getImm();3124      TmpType = GR.getOrCreateSPIRVVectorType(TmpType, NumElts, I, TII);3125    }3126    SrcReg = createVirtualRegister(TmpType, &GR, MRI, MRI->getMF());3127    selectSelectDefaultArgs(SrcReg, TmpType, I, false);3128  }3129  return selectOpWithSrcs(ResVReg, ResType, I, {SrcReg}, Opcode);3130}3131 3132bool SPIRVInstructionSelector::selectExt(Register ResVReg,3133                                         const SPIRVType *ResType,3134                                         MachineInstr &I, bool IsSigned) const {3135  Register SrcReg = I.getOperand(1).getReg();3136  if (GR.isScalarOrVectorOfType(SrcReg, SPIRV::OpTypeBool))3137    return selectSelectDefaultArgs(ResVReg, ResType, I, IsSigned);3138 3139  SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);3140  if (SrcType == ResType)3141    return BuildCOPY(ResVReg, SrcReg, I);3142 3143  unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;3144  return selectUnOp(ResVReg, ResType, I, Opcode);3145}3146 3147bool SPIRVInstructionSelector::selectSUCmp(Register ResVReg,3148                                           const SPIRVType *ResType,3149                                           MachineInstr &I,3150                                           bool IsSigned) const {3151  MachineIRBuilder MIRBuilder(I);3152  MachineRegisterInfo *MRI = MIRBuilder.getMRI();3153  MachineBasicBlock &BB = *I.getParent();3154  // Ensure we have bool.3155  SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);3156  unsigned N = GR.getScalarOrVectorComponentCount(ResType);3157  if (N > 1)3158    BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);3159  Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);3160  // Build less-than-equal and less-than.3161  // TODO: replace with one-liner createVirtualRegister() from3162  // llvm/lib/Target/SPIRV/SPIRVUtils.cpp when PR #116609 is merged.3163  Register IsLessEqReg = MRI->createVirtualRegister(GR.getRegClass(ResType));3164  MRI->setType(IsLessEqReg, LLT::scalar(64));3165  GR.assignSPIRVTypeToVReg(ResType, IsLessEqReg, MIRBuilder.getMF());3166  bool Result = BuildMI(BB, I, I.getDebugLoc(),3167                        TII.get(IsSigned ? SPIRV::OpSLessThanEqual3168                                         : SPIRV::OpULessThanEqual))3169                    .addDef(IsLessEqReg)3170                    .addUse(BoolTypeReg)3171                    .addUse(I.getOperand(1).getReg())3172                    .addUse(I.getOperand(2).getReg())3173                    .constrainAllUses(TII, TRI, RBI);3174  Register IsLessReg = MRI->createVirtualRegister(GR.getRegClass(ResType));3175  MRI->setType(IsLessReg, LLT::scalar(64));3176  GR.assignSPIRVTypeToVReg(ResType, IsLessReg, MIRBuilder.getMF());3177  Result &= BuildMI(BB, I, I.getDebugLoc(),3178                    TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))3179                .addDef(IsLessReg)3180                .addUse(BoolTypeReg)3181                .addUse(I.getOperand(1).getReg())3182                .addUse(I.getOperand(2).getReg())3183                .constrainAllUses(TII, TRI, RBI);3184  // Build selects.3185  Register ResTypeReg = GR.getSPIRVTypeID(ResType);3186  Register NegOneOrZeroReg =3187      MRI->createVirtualRegister(GR.getRegClass(ResType));3188  MRI->setType(NegOneOrZeroReg, LLT::scalar(64));3189  GR.assignSPIRVTypeToVReg(ResType, NegOneOrZeroReg, MIRBuilder.getMF());3190  unsigned SelectOpcode =3191      N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;3192  Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))3193                .addDef(NegOneOrZeroReg)3194                .addUse(ResTypeReg)3195                .addUse(IsLessReg)3196                .addUse(buildOnesVal(true, ResType, I)) // -13197                .addUse(buildZerosVal(ResType, I))3198                .constrainAllUses(TII, TRI, RBI);3199  return Result & BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))3200                      .addDef(ResVReg)3201                      .addUse(ResTypeReg)3202                      .addUse(IsLessEqReg)3203                      .addUse(NegOneOrZeroReg) // -1 or 03204                      .addUse(buildOnesVal(false, ResType, I))3205                      .constrainAllUses(TII, TRI, RBI);3206}3207 3208bool SPIRVInstructionSelector::selectIntToBool(Register IntReg,3209                                               Register ResVReg,3210                                               MachineInstr &I,3211                                               const SPIRVType *IntTy,3212                                               const SPIRVType *BoolTy) const {3213  // To truncate to a bool, we use OpBitwiseAnd 1 and OpINotEqual to zero.3214  Register BitIntReg = createVirtualRegister(IntTy, &GR, MRI, MRI->getMF());3215  bool IsVectorTy = IntTy->getOpcode() == SPIRV::OpTypeVector;3216  unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;3217  Register Zero = buildZerosVal(IntTy, I);3218  Register One = buildOnesVal(false, IntTy, I);3219  MachineBasicBlock &BB = *I.getParent();3220  bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))3221                    .addDef(BitIntReg)3222                    .addUse(GR.getSPIRVTypeID(IntTy))3223                    .addUse(IntReg)3224                    .addUse(One)3225                    .constrainAllUses(TII, TRI, RBI);3226  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))3227                       .addDef(ResVReg)3228                       .addUse(GR.getSPIRVTypeID(BoolTy))3229                       .addUse(BitIntReg)3230                       .addUse(Zero)3231                       .constrainAllUses(TII, TRI, RBI);3232}3233 3234bool SPIRVInstructionSelector::selectTrunc(Register ResVReg,3235                                           const SPIRVType *ResType,3236                                           MachineInstr &I) const {3237  Register IntReg = I.getOperand(1).getReg();3238  const SPIRVType *ArgType = GR.getSPIRVTypeForVReg(IntReg);3239  if (GR.isScalarOrVectorOfType(ResVReg, SPIRV::OpTypeBool))3240    return selectIntToBool(IntReg, ResVReg, I, ArgType, ResType);3241  if (ArgType == ResType)3242    return BuildCOPY(ResVReg, IntReg, I);3243  bool IsSigned = GR.isScalarOrVectorSigned(ResType);3244  unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;3245  return selectUnOp(ResVReg, ResType, I, Opcode);3246}3247 3248bool SPIRVInstructionSelector::selectConst(Register ResVReg,3249                                           const SPIRVType *ResType,3250                                           MachineInstr &I) const {3251  unsigned Opcode = I.getOpcode();3252  unsigned TpOpcode = ResType->getOpcode();3253  Register Reg;3254  if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {3255    assert(Opcode == TargetOpcode::G_CONSTANT &&3256           I.getOperand(1).getCImm()->isZero());3257    MachineBasicBlock &DepMBB = I.getMF()->front();3258    MachineIRBuilder MIRBuilder(DepMBB, DepMBB.getFirstNonPHI());3259    Reg = GR.getOrCreateConstNullPtr(MIRBuilder, ResType);3260  } else if (Opcode == TargetOpcode::G_FCONSTANT) {3261    Reg = GR.getOrCreateConstFP(I.getOperand(1).getFPImm()->getValue(), I,3262                                ResType, TII, !STI.isShader());3263  } else {3264    Reg = GR.getOrCreateConstInt(I.getOperand(1).getCImm()->getZExtValue(), I,3265                                 ResType, TII, !STI.isShader());3266  }3267  return Reg == ResVReg ? true : BuildCOPY(ResVReg, Reg, I);3268}3269 3270bool SPIRVInstructionSelector::selectOpUndef(Register ResVReg,3271                                             const SPIRVType *ResType,3272                                             MachineInstr &I) const {3273  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))3274      .addDef(ResVReg)3275      .addUse(GR.getSPIRVTypeID(ResType))3276      .constrainAllUses(TII, TRI, RBI);3277}3278 3279bool SPIRVInstructionSelector::selectInsertVal(Register ResVReg,3280                                               const SPIRVType *ResType,3281                                               MachineInstr &I) const {3282  MachineBasicBlock &BB = *I.getParent();3283  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeInsert))3284                 .addDef(ResVReg)3285                 .addUse(GR.getSPIRVTypeID(ResType))3286                 // object to insert3287                 .addUse(I.getOperand(3).getReg())3288                 // composite to insert into3289                 .addUse(I.getOperand(2).getReg());3290  for (unsigned i = 4; i < I.getNumOperands(); i++)3291    MIB.addImm(foldImm(I.getOperand(i), MRI));3292  return MIB.constrainAllUses(TII, TRI, RBI);3293}3294 3295bool SPIRVInstructionSelector::selectExtractVal(Register ResVReg,3296                                                const SPIRVType *ResType,3297                                                MachineInstr &I) const {3298  MachineBasicBlock &BB = *I.getParent();3299  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))3300                 .addDef(ResVReg)3301                 .addUse(GR.getSPIRVTypeID(ResType))3302                 .addUse(I.getOperand(2).getReg());3303  for (unsigned i = 3; i < I.getNumOperands(); i++)3304    MIB.addImm(foldImm(I.getOperand(i), MRI));3305  return MIB.constrainAllUses(TII, TRI, RBI);3306}3307 3308bool SPIRVInstructionSelector::selectInsertElt(Register ResVReg,3309                                               const SPIRVType *ResType,3310                                               MachineInstr &I) const {3311  if (getImm(I.getOperand(4), MRI))3312    return selectInsertVal(ResVReg, ResType, I);3313  MachineBasicBlock &BB = *I.getParent();3314  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorInsertDynamic))3315      .addDef(ResVReg)3316      .addUse(GR.getSPIRVTypeID(ResType))3317      .addUse(I.getOperand(2).getReg())3318      .addUse(I.getOperand(3).getReg())3319      .addUse(I.getOperand(4).getReg())3320      .constrainAllUses(TII, TRI, RBI);3321}3322 3323bool SPIRVInstructionSelector::selectExtractElt(Register ResVReg,3324                                                const SPIRVType *ResType,3325                                                MachineInstr &I) const {3326  if (getImm(I.getOperand(3), MRI))3327    return selectExtractVal(ResVReg, ResType, I);3328  MachineBasicBlock &BB = *I.getParent();3329  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorExtractDynamic))3330      .addDef(ResVReg)3331      .addUse(GR.getSPIRVTypeID(ResType))3332      .addUse(I.getOperand(2).getReg())3333      .addUse(I.getOperand(3).getReg())3334      .constrainAllUses(TII, TRI, RBI);3335}3336 3337bool SPIRVInstructionSelector::selectGEP(Register ResVReg,3338                                         const SPIRVType *ResType,3339                                         MachineInstr &I) const {3340  const bool IsGEPInBounds = I.getOperand(2).getImm();3341 3342  // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator only3343  // relies on PtrAccessChain, so we'll try not to deviate. For Vulkan however,3344  // we have to use Op[InBounds]AccessChain.3345  const unsigned Opcode = STI.isLogicalSPIRV()3346                              ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain3347                                               : SPIRV::OpAccessChain)3348                              : (IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain3349                                               : SPIRV::OpPtrAccessChain);3350 3351  auto Res = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))3352                 .addDef(ResVReg)3353                 .addUse(GR.getSPIRVTypeID(ResType))3354                 // Object to get a pointer to.3355                 .addUse(I.getOperand(3).getReg());3356  assert(Opcode == SPIRV::OpPtrAccessChain ||3357         Opcode == SPIRV::OpInBoundsPtrAccessChain ||3358         (getImm(I.getOperand(4), MRI) && foldImm(I.getOperand(4), MRI) == 0) &&3359             "Cannot translate GEP to OpAccessChain. First index must be 0.");3360 3361  // Adding indices.3362  const unsigned StartingIndex =3363      (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)3364          ? 53365          : 4;3366  for (unsigned i = StartingIndex; i < I.getNumExplicitOperands(); ++i)3367    Res.addUse(I.getOperand(i).getReg());3368  return Res.constrainAllUses(TII, TRI, RBI);3369}3370 3371// Maybe wrap a value into OpSpecConstantOp3372bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(3373    MachineInstr &I, SmallVector<Register> &CompositeArgs) const {3374  bool Result = true;3375  unsigned Lim = I.getNumExplicitOperands();3376  for (unsigned i = I.getNumExplicitDefs() + 1; i < Lim; ++i) {3377    Register OpReg = I.getOperand(i).getReg();3378    MachineInstr *OpDefine = MRI->getVRegDef(OpReg);3379    SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);3380    SmallPtrSet<SPIRVType *, 4> Visited;3381    if (!OpDefine || !OpType || isConstReg(MRI, OpDefine, Visited) ||3382        OpDefine->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||3383        OpDefine->getOpcode() == TargetOpcode::G_INTTOPTR ||3384        GR.isAggregateType(OpType)) {3385      // The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed3386      // by selectAddrSpaceCast(), and G_INTTOPTR is processed by selectUnOp()3387      CompositeArgs.push_back(OpReg);3388      continue;3389    }3390    MachineFunction *MF = I.getMF();3391    Register WrapReg = GR.find(OpDefine, MF);3392    if (WrapReg.isValid()) {3393      CompositeArgs.push_back(WrapReg);3394      continue;3395    }3396    // Create a new register for the wrapper3397    WrapReg = MRI->createVirtualRegister(GR.getRegClass(OpType));3398    CompositeArgs.push_back(WrapReg);3399    // Decorate the wrapper register and generate a new instruction3400    MRI->setType(WrapReg, LLT::pointer(0, 64));3401    GR.assignSPIRVTypeToVReg(OpType, WrapReg, *MF);3402    auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),3403                       TII.get(SPIRV::OpSpecConstantOp))3404                   .addDef(WrapReg)3405                   .addUse(GR.getSPIRVTypeID(OpType))3406                   .addImm(static_cast<uint32_t>(SPIRV::Opcode::Bitcast))3407                   .addUse(OpReg);3408    GR.add(OpDefine, MIB);3409    Result = MIB.constrainAllUses(TII, TRI, RBI);3410    if (!Result)3411      break;3412  }3413  return Result;3414}3415 3416bool SPIRVInstructionSelector::selectDerivativeInst(3417    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,3418    const unsigned DPdOpCode) const {3419  // TODO: This should check specifically for Fragment Execution Model, but STI3420  // doesn't provide that information yet. See #1675623421  errorIfInstrOutsideShader(I);3422 3423  // If the arg/result types are half then we need to wrap the instr in3424  // conversions to float3425  // This case occurs because a half arg/result is legal in HLSL but not spirv.3426  Register SrcReg = I.getOperand(2).getReg();3427  SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);3428  unsigned BitWidth = std::min(GR.getScalarOrVectorBitWidth(SrcType),3429                               GR.getScalarOrVectorBitWidth(ResType));3430  if (BitWidth == 32)3431    return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DPdOpCode))3432        .addDef(ResVReg)3433        .addUse(GR.getSPIRVTypeID(ResType))3434        .addUse(I.getOperand(2).getReg());3435 3436  MachineIRBuilder MIRBuilder(I);3437  unsigned componentCount = GR.getScalarOrVectorComponentCount(SrcType);3438  SPIRVType *F32ConvertTy = GR.getOrCreateSPIRVFloatType(32, I, TII);3439  if (componentCount != 1)3440    F32ConvertTy = GR.getOrCreateSPIRVVectorType(F32ConvertTy, componentCount,3441                                                 MIRBuilder, false);3442 3443  const TargetRegisterClass *RegClass = GR.getRegClass(SrcType);3444  Register ConvertToVReg = MRI->createVirtualRegister(RegClass);3445  Register DpdOpVReg = MRI->createVirtualRegister(RegClass);3446 3447  bool Result =3448      BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpFConvert))3449          .addDef(ConvertToVReg)3450          .addUse(GR.getSPIRVTypeID(F32ConvertTy))3451          .addUse(SrcReg)3452          .constrainAllUses(TII, TRI, RBI);3453  Result &= BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DPdOpCode))3454                .addDef(DpdOpVReg)3455                .addUse(GR.getSPIRVTypeID(F32ConvertTy))3456                .addUse(ConvertToVReg)3457                .constrainAllUses(TII, TRI, RBI);3458  Result &=3459      BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpFConvert))3460          .addDef(ResVReg)3461          .addUse(GR.getSPIRVTypeID(ResType))3462          .addUse(DpdOpVReg)3463          .constrainAllUses(TII, TRI, RBI);3464  return Result;3465}3466 3467bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,3468                                               const SPIRVType *ResType,3469                                               MachineInstr &I) const {3470  MachineBasicBlock &BB = *I.getParent();3471  Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();3472  switch (IID) {3473  case Intrinsic::spv_load:3474    return selectLoad(ResVReg, ResType, I);3475  case Intrinsic::spv_store:3476    return selectStore(I);3477  case Intrinsic::spv_extractv:3478    return selectExtractVal(ResVReg, ResType, I);3479  case Intrinsic::spv_insertv:3480    return selectInsertVal(ResVReg, ResType, I);3481  case Intrinsic::spv_extractelt:3482    return selectExtractElt(ResVReg, ResType, I);3483  case Intrinsic::spv_insertelt:3484    return selectInsertElt(ResVReg, ResType, I);3485  case Intrinsic::spv_gep:3486    return selectGEP(ResVReg, ResType, I);3487  case Intrinsic::spv_bitcast: {3488    Register OpReg = I.getOperand(2).getReg();3489    SPIRVType *OpType =3490        OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;3491    if (!GR.isBitcastCompatible(ResType, OpType))3492      report_fatal_error("incompatible result and operand types in a bitcast");3493    return selectOpWithSrcs(ResVReg, ResType, I, {OpReg}, SPIRV::OpBitcast);3494  }3495  case Intrinsic::spv_unref_global:3496  case Intrinsic::spv_init_global: {3497    MachineInstr *MI = MRI->getVRegDef(I.getOperand(1).getReg());3498    MachineInstr *Init = I.getNumExplicitOperands() > 23499                             ? MRI->getVRegDef(I.getOperand(2).getReg())3500                             : nullptr;3501    assert(MI);3502    Register GVarVReg = MI->getOperand(0).getReg();3503    bool Res = selectGlobalValue(GVarVReg, *MI, Init);3504    // We violate SSA form by inserting OpVariable and still having a gMIR3505    // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by erasing3506    // the duplicated definition.3507    if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {3508      GR.invalidateMachineInstr(MI);3509      MI->removeFromParent();3510    }3511    return Res;3512  }3513  case Intrinsic::spv_undef: {3514    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))3515                   .addDef(ResVReg)3516                   .addUse(GR.getSPIRVTypeID(ResType));3517    return MIB.constrainAllUses(TII, TRI, RBI);3518  }3519  case Intrinsic::spv_const_composite: {3520    // If no values are attached, the composite is null constant.3521    bool IsNull = I.getNumExplicitDefs() + 1 == I.getNumExplicitOperands();3522    SmallVector<Register> CompositeArgs;3523    MRI->setRegClass(ResVReg, GR.getRegClass(ResType));3524 3525    // skip type MD node we already used when generated assign.type for this3526    if (!IsNull) {3527      if (!wrapIntoSpecConstantOp(I, CompositeArgs))3528        return false;3529      MachineIRBuilder MIR(I);3530      SmallVector<MachineInstr *, 4> Instructions = createContinuedInstructions(3531          MIR, SPIRV::OpConstantComposite, 3,3532          SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,3533          GR.getSPIRVTypeID(ResType));3534      for (auto *Instr : Instructions) {3535        Instr->setDebugLoc(I.getDebugLoc());3536        if (!constrainSelectedInstRegOperands(*Instr, TII, TRI, RBI))3537          return false;3538      }3539      return true;3540    } else {3541      auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))3542                     .addDef(ResVReg)3543                     .addUse(GR.getSPIRVTypeID(ResType));3544      return MIB.constrainAllUses(TII, TRI, RBI);3545    }3546  }3547  case Intrinsic::spv_assign_name: {3548    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpName));3549    MIB.addUse(I.getOperand(I.getNumExplicitDefs() + 1).getReg());3550    for (unsigned i = I.getNumExplicitDefs() + 2;3551         i < I.getNumExplicitOperands(); ++i) {3552      MIB.addImm(I.getOperand(i).getImm());3553    }3554    return MIB.constrainAllUses(TII, TRI, RBI);3555  }3556  case Intrinsic::spv_switch: {3557    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSwitch));3558    for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {3559      if (I.getOperand(i).isReg())3560        MIB.addReg(I.getOperand(i).getReg());3561      else if (I.getOperand(i).isCImm())3562        addNumImm(I.getOperand(i).getCImm()->getValue(), MIB);3563      else if (I.getOperand(i).isMBB())3564        MIB.addMBB(I.getOperand(i).getMBB());3565      else3566        llvm_unreachable("Unexpected OpSwitch operand");3567    }3568    return MIB.constrainAllUses(TII, TRI, RBI);3569  }3570  case Intrinsic::spv_loop_merge: {3571    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoopMerge));3572    for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {3573      if (I.getOperand(i).isMBB())3574        MIB.addMBB(I.getOperand(i).getMBB());3575      else3576        MIB.addImm(foldImm(I.getOperand(i), MRI));3577    }3578    return MIB.constrainAllUses(TII, TRI, RBI);3579  }3580  case Intrinsic::spv_selection_merge: {3581    auto MIB =3582        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSelectionMerge));3583    assert(I.getOperand(1).isMBB() &&3584           "operand 1 to spv_selection_merge must be a basic block");3585    MIB.addMBB(I.getOperand(1).getMBB());3586    MIB.addImm(getSelectionOperandForImm(I.getOperand(2).getImm()));3587    return MIB.constrainAllUses(TII, TRI, RBI);3588  }3589  case Intrinsic::spv_cmpxchg:3590    return selectAtomicCmpXchg(ResVReg, ResType, I);3591  case Intrinsic::spv_unreachable:3592    return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUnreachable))3593        .constrainAllUses(TII, TRI, RBI);3594  case Intrinsic::spv_alloca:3595    return selectFrameIndex(ResVReg, ResType, I);3596  case Intrinsic::spv_alloca_array:3597    return selectAllocaArray(ResVReg, ResType, I);3598  case Intrinsic::spv_assume:3599    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))3600      return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAssumeTrueKHR))3601          .addUse(I.getOperand(1).getReg())3602          .constrainAllUses(TII, TRI, RBI);3603    break;3604  case Intrinsic::spv_expect:3605    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))3606      return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExpectKHR))3607          .addDef(ResVReg)3608          .addUse(GR.getSPIRVTypeID(ResType))3609          .addUse(I.getOperand(2).getReg())3610          .addUse(I.getOperand(3).getReg())3611          .constrainAllUses(TII, TRI, RBI);3612    break;3613  case Intrinsic::arithmetic_fence:3614    if (STI.canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence))3615      return BuildMI(BB, I, I.getDebugLoc(),3616                     TII.get(SPIRV::OpArithmeticFenceEXT))3617          .addDef(ResVReg)3618          .addUse(GR.getSPIRVTypeID(ResType))3619          .addUse(I.getOperand(2).getReg())3620          .constrainAllUses(TII, TRI, RBI);3621    else3622      return BuildCOPY(ResVReg, I.getOperand(2).getReg(), I);3623    break;3624  case Intrinsic::spv_thread_id:3625    // The HLSL SV_DispatchThreadID semantic is lowered to llvm.spv.thread.id3626    // intrinsic in LLVM IR for SPIR-V backend.3627    //3628    // In SPIR-V backend, llvm.spv.thread.id is now correctly translated to a3629    // `GlobalInvocationId` builtin variable3630    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,3631                                  ResType, I);3632  case Intrinsic::spv_thread_id_in_group:3633    // The HLSL SV_GroupThreadId semantic is lowered to3634    // llvm.spv.thread.id.in.group intrinsic in LLVM IR for SPIR-V backend.3635    //3636    // In SPIR-V backend, llvm.spv.thread.id.in.group is now correctly3637    // translated to a `LocalInvocationId` builtin variable3638    return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,3639                                  ResType, I);3640  case Intrinsic::spv_group_id:3641    // The HLSL SV_GroupId semantic is lowered to3642    // llvm.spv.group.id intrinsic in LLVM IR for SPIR-V backend.3643    //3644    // In SPIR-V backend, llvm.spv.group.id is now translated to a `WorkgroupId`3645    // builtin variable3646    return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,3647                                  I);3648  case Intrinsic::spv_flattened_thread_id_in_group:3649    // The HLSL SV_GroupIndex semantic is lowered to3650    // llvm.spv.flattened.thread.id.in.group() intrinsic in LLVM IR for SPIR-V3651    // backend.3652    //3653    // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated to3654    // a `LocalInvocationIndex` builtin variable3655    return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,3656                              ResType, I);3657  case Intrinsic::spv_workgroup_size:3658    return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,3659                                  ResType, I);3660  case Intrinsic::spv_global_size:3661    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,3662                                  I);3663  case Intrinsic::spv_global_offset:3664    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,3665                                  ResType, I);3666  case Intrinsic::spv_num_workgroups:3667    return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,3668                                  ResType, I);3669  case Intrinsic::spv_subgroup_size:3670    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,3671                              I);3672  case Intrinsic::spv_num_subgroups:3673    return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,3674                              I);3675  case Intrinsic::spv_subgroup_id:3676    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType, I);3677  case Intrinsic::spv_subgroup_local_invocation_id:3678    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,3679                              ResVReg, ResType, I);3680  case Intrinsic::spv_subgroup_max_size:3681    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,3682                              I);3683  case Intrinsic::spv_fdot:3684    return selectFloatDot(ResVReg, ResType, I);3685  case Intrinsic::spv_udot:3686  case Intrinsic::spv_sdot:3687    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||3688        STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))3689      return selectIntegerDot(ResVReg, ResType, I,3690                              /*Signed=*/IID == Intrinsic::spv_sdot);3691    return selectIntegerDotExpansion(ResVReg, ResType, I);3692  case Intrinsic::spv_dot4add_i8packed:3693    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||3694        STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))3695      return selectDot4AddPacked<true>(ResVReg, ResType, I);3696    return selectDot4AddPackedExpansion<true>(ResVReg, ResType, I);3697  case Intrinsic::spv_dot4add_u8packed:3698    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||3699        STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))3700      return selectDot4AddPacked<false>(ResVReg, ResType, I);3701    return selectDot4AddPackedExpansion<false>(ResVReg, ResType, I);3702  case Intrinsic::spv_all:3703    return selectAll(ResVReg, ResType, I);3704  case Intrinsic::spv_any:3705    return selectAny(ResVReg, ResType, I);3706  case Intrinsic::spv_cross:3707    return selectExtInst(ResVReg, ResType, I, CL::cross, GL::Cross);3708  case Intrinsic::spv_distance:3709    return selectExtInst(ResVReg, ResType, I, CL::distance, GL::Distance);3710  case Intrinsic::spv_lerp:3711    return selectExtInst(ResVReg, ResType, I, CL::mix, GL::FMix);3712  case Intrinsic::spv_length:3713    return selectExtInst(ResVReg, ResType, I, CL::length, GL::Length);3714  case Intrinsic::spv_degrees:3715    return selectExtInst(ResVReg, ResType, I, CL::degrees, GL::Degrees);3716  case Intrinsic::spv_faceforward:3717    return selectExtInst(ResVReg, ResType, I, GL::FaceForward);3718  case Intrinsic::spv_frac:3719    return selectExtInst(ResVReg, ResType, I, CL::fract, GL::Fract);3720  case Intrinsic::spv_isinf:3721    return selectOpIsInf(ResVReg, ResType, I);3722  case Intrinsic::spv_isnan:3723    return selectOpIsNan(ResVReg, ResType, I);3724  case Intrinsic::spv_normalize:3725    return selectExtInst(ResVReg, ResType, I, CL::normalize, GL::Normalize);3726  case Intrinsic::spv_refract:3727    return selectExtInst(ResVReg, ResType, I, GL::Refract);3728  case Intrinsic::spv_reflect:3729    return selectExtInst(ResVReg, ResType, I, GL::Reflect);3730  case Intrinsic::spv_rsqrt:3731    return selectExtInst(ResVReg, ResType, I, CL::rsqrt, GL::InverseSqrt);3732  case Intrinsic::spv_sign:3733    return selectSign(ResVReg, ResType, I);3734  case Intrinsic::spv_smoothstep:3735    return selectExtInst(ResVReg, ResType, I, CL::smoothstep, GL::SmoothStep);3736  case Intrinsic::spv_firstbituhigh: // There is no CL equivalent of FindUMsb3737    return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/false);3738  case Intrinsic::spv_firstbitshigh: // There is no CL equivalent of FindSMsb3739    return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/true);3740  case Intrinsic::spv_firstbitlow: // There is no CL equivlent of FindILsb3741    return selectFirstBitLow(ResVReg, ResType, I);3742  case Intrinsic::spv_group_memory_barrier_with_group_sync: {3743    bool Result = true;3744    auto MemSemConstant =3745        buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent, I);3746    Register MemSemReg = MemSemConstant.first;3747    Result &= MemSemConstant.second;3748    auto ScopeConstant = buildI32Constant(SPIRV::Scope::Workgroup, I);3749    Register ScopeReg = ScopeConstant.first;3750    Result &= ScopeConstant.second;3751    MachineBasicBlock &BB = *I.getParent();3752    return Result &&3753           BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpControlBarrier))3754               .addUse(ScopeReg)3755               .addUse(ScopeReg)3756               .addUse(MemSemReg)3757               .constrainAllUses(TII, TRI, RBI);3758  }3759  case Intrinsic::spv_generic_cast_to_ptr_explicit: {3760    Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 1).getReg();3761    SPIRV::StorageClass::StorageClass ResSC =3762        GR.getPointerStorageClass(ResType);3763    if (!isGenericCastablePtr(ResSC))3764      report_fatal_error("The target storage class is not castable from the "3765                         "Generic storage class");3766    return BuildMI(BB, I, I.getDebugLoc(),3767                   TII.get(SPIRV::OpGenericCastToPtrExplicit))3768        .addDef(ResVReg)3769        .addUse(GR.getSPIRVTypeID(ResType))3770        .addUse(PtrReg)3771        .addImm(ResSC)3772        .constrainAllUses(TII, TRI, RBI);3773  }3774  case Intrinsic::spv_lifetime_start:3775  case Intrinsic::spv_lifetime_end: {3776    unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart3777                                                       : SPIRV::OpLifetimeStop;3778    int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();3779    Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();3780    if (Size == -1)3781      Size = 0;3782    return BuildMI(BB, I, I.getDebugLoc(), TII.get(Op))3783        .addUse(PtrReg)3784        .addImm(Size)3785        .constrainAllUses(TII, TRI, RBI);3786  }3787  case Intrinsic::spv_saturate:3788    return selectSaturate(ResVReg, ResType, I);3789  case Intrinsic::spv_nclamp:3790    return selectExtInst(ResVReg, ResType, I, CL::fclamp, GL::NClamp);3791  case Intrinsic::spv_uclamp:3792    return selectExtInst(ResVReg, ResType, I, CL::u_clamp, GL::UClamp);3793  case Intrinsic::spv_sclamp:3794    return selectExtInst(ResVReg, ResType, I, CL::s_clamp, GL::SClamp);3795  case Intrinsic::spv_wave_active_countbits:3796    return selectWaveActiveCountBits(ResVReg, ResType, I);3797  case Intrinsic::spv_wave_all:3798    return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAll);3799  case Intrinsic::spv_wave_any:3800    return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAny);3801  case Intrinsic::spv_wave_is_first_lane:3802    return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformElect);3803  case Intrinsic::spv_wave_reduce_umax:3804    return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ true);3805  case Intrinsic::spv_wave_reduce_max:3806    return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ false);3807  case Intrinsic::spv_wave_reduce_umin:3808    return selectWaveReduceMin(ResVReg, ResType, I, /*IsUnsigned*/ true);3809  case Intrinsic::spv_wave_reduce_min:3810    return selectWaveReduceMin(ResVReg, ResType, I, /*IsUnsigned*/ false);3811  case Intrinsic::spv_wave_reduce_sum:3812    return selectWaveReduceSum(ResVReg, ResType, I);3813  case Intrinsic::spv_wave_readlane:3814    return selectWaveOpInst(ResVReg, ResType, I,3815                            SPIRV::OpGroupNonUniformShuffle);3816  case Intrinsic::spv_step:3817    return selectExtInst(ResVReg, ResType, I, CL::step, GL::Step);3818  case Intrinsic::spv_radians:3819    return selectExtInst(ResVReg, ResType, I, CL::radians, GL::Radians);3820  // Discard intrinsics which we do not expect to actually represent code after3821  // lowering or intrinsics which are not implemented but should not crash when3822  // found in a customer's LLVM IR input.3823  case Intrinsic::instrprof_increment:3824  case Intrinsic::instrprof_increment_step:3825  case Intrinsic::instrprof_value_profile:3826    break;3827  // Discard internal intrinsics.3828  case Intrinsic::spv_value_md:3829    break;3830  case Intrinsic::spv_resource_handlefrombinding: {3831    return selectHandleFromBinding(ResVReg, ResType, I);3832  }3833  case Intrinsic::spv_resource_counterhandlefrombinding:3834    return selectCounterHandleFromBinding(ResVReg, ResType, I);3835  case Intrinsic::spv_resource_updatecounter:3836    return selectUpdateCounter(ResVReg, ResType, I);3837  case Intrinsic::spv_resource_store_typedbuffer: {3838    return selectImageWriteIntrinsic(I);3839  }3840  case Intrinsic::spv_resource_load_typedbuffer: {3841    return selectReadImageIntrinsic(ResVReg, ResType, I);3842  }3843  case Intrinsic::spv_resource_getpointer: {3844    return selectResourceGetPointer(ResVReg, ResType, I);3845  }3846  case Intrinsic::spv_discard: {3847    return selectDiscard(ResVReg, ResType, I);3848  }3849  case Intrinsic::spv_resource_nonuniformindex: {3850    return selectResourceNonUniformIndex(ResVReg, ResType, I);3851  }3852  case Intrinsic::spv_unpackhalf2x16: {3853    return selectExtInst(ResVReg, ResType, I, GL::UnpackHalf2x16);3854  }3855  case Intrinsic::spv_ddx_coarse:3856    return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdxCoarse);3857  case Intrinsic::spv_ddy_coarse:3858    return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpDPdyCoarse);3859  case Intrinsic::spv_fwidth:3860    return selectDerivativeInst(ResVReg, ResType, I, SPIRV::OpFwidth);3861  default: {3862    std::string DiagMsg;3863    raw_string_ostream OS(DiagMsg);3864    I.print(OS);3865    DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;3866    report_fatal_error(DiagMsg.c_str(), false);3867  }3868  }3869  return true;3870}3871 3872bool SPIRVInstructionSelector::selectHandleFromBinding(Register &ResVReg,3873                                                       const SPIRVType *ResType,3874                                                       MachineInstr &I) const {3875  // The images need to be loaded in the same basic block as their use. We defer3876  // loading the image to the intrinsic that uses it.3877  if (ResType->getOpcode() == SPIRV::OpTypeImage)3878    return true;3879 3880  return loadHandleBeforePosition(ResVReg, GR.getSPIRVTypeForVReg(ResVReg),3881                                  *cast<GIntrinsic>(&I), I);3882}3883 3884bool SPIRVInstructionSelector::selectCounterHandleFromBinding(3885    Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {3886  auto &Intr = cast<GIntrinsic>(I);3887  assert(Intr.getIntrinsicID() ==3888         Intrinsic::spv_resource_counterhandlefrombinding);3889 3890  // Extract information from the intrinsic call.3891  Register MainHandleReg = Intr.getOperand(2).getReg();3892  auto *MainHandleDef = cast<GIntrinsic>(getVRegDef(*MRI, MainHandleReg));3893  assert(MainHandleDef->getIntrinsicID() ==3894         Intrinsic::spv_resource_handlefrombinding);3895 3896  uint32_t Set = getIConstVal(Intr.getOperand(4).getReg(), MRI);3897  uint32_t Binding = getIConstVal(Intr.getOperand(3).getReg(), MRI);3898  uint32_t ArraySize = getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);3899  Register IndexReg = MainHandleDef->getOperand(5).getReg();3900  std::string CounterName =3901      getStringValueFromReg(MainHandleDef->getOperand(6).getReg(), *MRI) +3902      ".counter";3903 3904  // Create the counter variable.3905  MachineIRBuilder MIRBuilder(I);3906  Register CounterVarReg = buildPointerToResource(3907      GR.getPointeeType(ResType), GR.getPointerStorageClass(ResType), Set,3908      Binding, ArraySize, IndexReg, CounterName, MIRBuilder);3909 3910  return BuildCOPY(ResVReg, CounterVarReg, I);3911}3912 3913bool SPIRVInstructionSelector::selectUpdateCounter(Register &ResVReg,3914                                                   const SPIRVType *ResType,3915                                                   MachineInstr &I) const {3916  auto &Intr = cast<GIntrinsic>(I);3917  assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);3918 3919  Register CounterHandleReg = Intr.getOperand(2).getReg();3920  Register IncrReg = Intr.getOperand(3).getReg();3921 3922  // The counter handle is a pointer to the counter variable (which is a struct3923  // containing an i32). We need to get a pointer to that i32 member to do the3924  // atomic operation.3925#ifndef NDEBUG3926  SPIRVType *CounterVarType = GR.getSPIRVTypeForVReg(CounterHandleReg);3927  SPIRVType *CounterVarPointeeType = GR.getPointeeType(CounterVarType);3928  assert(CounterVarPointeeType &&3929         CounterVarPointeeType->getOpcode() == SPIRV::OpTypeStruct &&3930         "Counter variable must be a struct");3931  assert(GR.getPointerStorageClass(CounterVarType) ==3932             SPIRV::StorageClass::StorageBuffer &&3933         "Counter variable must be in the storage buffer storage class");3934  assert(CounterVarPointeeType->getNumOperands() == 2 &&3935         "Counter variable must have exactly 1 member in the struct");3936  const SPIRVType *MemberType =3937      GR.getSPIRVTypeForVReg(CounterVarPointeeType->getOperand(1).getReg());3938  assert(MemberType->getOpcode() == SPIRV::OpTypeInt &&3939         "Counter variable struct must have a single i32 member");3940#endif3941 3942  // The struct has a single i32 member.3943  MachineIRBuilder MIRBuilder(I);3944  const Type *LLVMIntType =3945      Type::getInt32Ty(I.getMF()->getFunction().getContext());3946 3947  SPIRVType *IntPtrType = GR.getOrCreateSPIRVPointerType(3948      LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);3949 3950  auto Zero = buildI32Constant(0, I);3951  if (!Zero.second)3952    return false;3953 3954  Register PtrToCounter =3955      MRI->createVirtualRegister(GR.getRegClass(IntPtrType));3956  if (!BuildMI(*I.getParent(), I, I.getDebugLoc(),3957               TII.get(SPIRV::OpAccessChain))3958           .addDef(PtrToCounter)3959           .addUse(GR.getSPIRVTypeID(IntPtrType))3960           .addUse(CounterHandleReg)3961           .addUse(Zero.first)3962           .constrainAllUses(TII, TRI, RBI)) {3963    return false;3964  }3965 3966  // For UAV/SSBO counters, the scope is Device. The counter variable is not3967  // used as a flag. So the memory semantics can be None.3968  auto Scope = buildI32Constant(SPIRV::Scope::Device, I);3969  if (!Scope.second)3970    return false;3971  auto Semantics = buildI32Constant(SPIRV::MemorySemantics::None, I);3972  if (!Semantics.second)3973    return false;3974 3975  int64_t IncrVal = getIConstValSext(IncrReg, MRI);3976  auto Incr = buildI32Constant(static_cast<uint32_t>(IncrVal), I);3977  if (!Incr.second)3978    return false;3979 3980  Register AtomicRes = MRI->createVirtualRegister(GR.getRegClass(ResType));3981  if (!BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpAtomicIAdd))3982           .addDef(AtomicRes)3983           .addUse(GR.getSPIRVTypeID(ResType))3984           .addUse(PtrToCounter)3985           .addUse(Scope.first)3986           .addUse(Semantics.first)3987           .addUse(Incr.first)3988           .constrainAllUses(TII, TRI, RBI)) {3989    return false;3990  }3991  if (IncrVal >= 0) {3992    return BuildCOPY(ResVReg, AtomicRes, I);3993  }3994 3995  // In HLSL, IncrementCounter returns the value *before* the increment, while3996  // DecrementCounter returns the value *after* the decrement. Both are lowered3997  // to the same atomic intrinsic which returns the value *before* the3998  // operation. So for decrements (negative IncrVal), we must subtract the3999  // increment value from the result to get the post-decrement value.4000  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))4001      .addDef(ResVReg)4002      .addUse(GR.getSPIRVTypeID(ResType))4003      .addUse(AtomicRes)4004      .addUse(Incr.first)4005      .constrainAllUses(TII, TRI, RBI);4006}4007bool SPIRVInstructionSelector::selectReadImageIntrinsic(4008    Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {4009 4010  // If the load of the image is in a different basic block, then4011  // this will generate invalid code. A proper solution is to move4012  // the OpLoad from selectHandleFromBinding here. However, to do4013  // that we will need to change the return type of the intrinsic.4014  // We will do that when we can, but for now trying to move forward with other4015  // issues.4016  Register ImageReg = I.getOperand(2).getReg();4017  auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));4018  Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));4019  if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),4020                                *ImageDef, I)) {4021    return false;4022  }4023 4024  Register IdxReg = I.getOperand(3).getReg();4025  DebugLoc Loc = I.getDebugLoc();4026  MachineInstr &Pos = I;4027 4028  return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,4029                                  Pos);4030}4031 4032bool SPIRVInstructionSelector::generateImageReadOrFetch(4033    Register &ResVReg, const SPIRVType *ResType, Register ImageReg,4034    Register IdxReg, DebugLoc Loc, MachineInstr &Pos) const {4035  SPIRVType *ImageType = GR.getSPIRVTypeForVReg(ImageReg);4036  assert(ImageType && ImageType->getOpcode() == SPIRV::OpTypeImage &&4037         "ImageReg is not an image type.");4038 4039  bool IsSignedInteger =4040      sampledTypeIsSignedInteger(GR.getTypeForSPIRVType(ImageType));4041  // Check if the "sampled" operand of the image type is 1.4042  // https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpImageFetch4043  auto SampledOp = ImageType->getOperand(6);4044  bool IsFetch = (SampledOp.getImm() == 1);4045 4046  uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);4047  if (ResultSize == 4) {4048    auto BMI =4049        BuildMI(*Pos.getParent(), Pos, Loc,4050                TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))4051            .addDef(ResVReg)4052            .addUse(GR.getSPIRVTypeID(ResType))4053            .addUse(ImageReg)4054            .addUse(IdxReg);4055 4056    if (IsSignedInteger)4057      BMI.addImm(0x1000); // SignExtend4058    return BMI.constrainAllUses(TII, TRI, RBI);4059  }4060 4061  SPIRVType *ReadType = widenTypeToVec4(ResType, Pos);4062  Register ReadReg = MRI->createVirtualRegister(GR.getRegClass(ReadType));4063  auto BMI =4064      BuildMI(*Pos.getParent(), Pos, Loc,4065              TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))4066          .addDef(ReadReg)4067          .addUse(GR.getSPIRVTypeID(ReadType))4068          .addUse(ImageReg)4069          .addUse(IdxReg);4070  if (IsSignedInteger)4071    BMI.addImm(0x1000); // SignExtend4072  bool Succeed = BMI.constrainAllUses(TII, TRI, RBI);4073  if (!Succeed)4074    return false;4075 4076  if (ResultSize == 1) {4077    return BuildMI(*Pos.getParent(), Pos, Loc,4078                   TII.get(SPIRV::OpCompositeExtract))4079        .addDef(ResVReg)4080        .addUse(GR.getSPIRVTypeID(ResType))4081        .addUse(ReadReg)4082        .addImm(0)4083        .constrainAllUses(TII, TRI, RBI);4084  }4085  return extractSubvector(ResVReg, ResType, ReadReg, Pos);4086}4087 4088bool SPIRVInstructionSelector::selectResourceGetPointer(4089    Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {4090  Register ResourcePtr = I.getOperand(2).getReg();4091  SPIRVType *RegType = GR.getSPIRVTypeForVReg(ResourcePtr, I.getMF());4092  if (RegType->getOpcode() == SPIRV::OpTypeImage) {4093    // For texel buffers, the index into the image is part of the OpImageRead or4094    // OpImageWrite instructions. So we will do nothing in this case. This4095    // intrinsic will be combined with the load or store when selecting the load4096    // or store.4097    return true;4098  }4099 4100  assert(ResType->getOpcode() == SPIRV::OpTypePointer);4101  MachineIRBuilder MIRBuilder(I);4102 4103  Register IndexReg = I.getOperand(3).getReg();4104  Register ZeroReg =4105      buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);4106  return BuildMI(*I.getParent(), I, I.getDebugLoc(),4107                 TII.get(SPIRV::OpAccessChain))4108      .addDef(ResVReg)4109      .addUse(GR.getSPIRVTypeID(ResType))4110      .addUse(ResourcePtr)4111      .addUse(ZeroReg)4112      .addUse(IndexReg)4113      .constrainAllUses(TII, TRI, RBI);4114}4115 4116bool SPIRVInstructionSelector::selectResourceNonUniformIndex(4117    Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {4118  Register ObjReg = I.getOperand(2).getReg();4119  if (!BuildCOPY(ResVReg, ObjReg, I))4120    return false;4121 4122  buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::NonUniformEXT, {});4123  // Check for the registers that use the index marked as non-uniform4124  // and recursively mark them as non-uniform.4125  // Per the spec, it's necessary that the final argument used for4126  // load/store/sample/atomic must be decorated, so we need to propagate the4127  // decoration through access chains and copies.4128  // https://docs.vulkan.org/samples/latest/samples/extensions/descriptor_indexing/README.html#_when_to_use_non_uniform_indexing_qualifier4129  decorateUsesAsNonUniform(ResVReg);4130  return true;4131}4132 4133void SPIRVInstructionSelector::decorateUsesAsNonUniform(4134    Register &NonUniformReg) const {4135  llvm::SmallVector<Register> WorkList = {NonUniformReg};4136  while (WorkList.size() > 0) {4137    Register CurrentReg = WorkList.back();4138    WorkList.pop_back();4139 4140    bool IsDecorated = false;4141    for (MachineInstr &Use : MRI->use_instructions(CurrentReg)) {4142      if (Use.getOpcode() == SPIRV::OpDecorate &&4143          Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {4144        IsDecorated = true;4145        continue;4146      }4147      // Check if the instruction has the result register and add it to the4148      // worklist.4149      if (Use.getOperand(0).isReg() && Use.getOperand(0).isDef()) {4150        Register ResultReg = Use.getOperand(0).getReg();4151        if (ResultReg == CurrentReg)4152          continue;4153        WorkList.push_back(ResultReg);4154      }4155    }4156 4157    if (!IsDecorated) {4158      buildOpDecorate(CurrentReg, *MRI->getVRegDef(CurrentReg), TII,4159                      SPIRV::Decoration::NonUniformEXT, {});4160    }4161  }4162}4163 4164bool SPIRVInstructionSelector::extractSubvector(4165    Register &ResVReg, const SPIRVType *ResType, Register &ReadReg,4166    MachineInstr &InsertionPoint) const {4167  SPIRVType *InputType = GR.getResultType(ReadReg);4168  [[maybe_unused]] uint64_t InputSize =4169      GR.getScalarOrVectorComponentCount(InputType);4170  uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);4171  assert(InputSize > 1 && "The input must be a vector.");4172  assert(ResultSize > 1 && "The result must be a vector.");4173  assert(ResultSize < InputSize &&4174         "Cannot extract more element than there are in the input.");4175  SmallVector<Register> ComponentRegisters;4176  SPIRVType *ScalarType = GR.getScalarOrVectorComponentType(ResType);4177  const TargetRegisterClass *ScalarRegClass = GR.getRegClass(ScalarType);4178  for (uint64_t I = 0; I < ResultSize; I++) {4179    Register ComponentReg = MRI->createVirtualRegister(ScalarRegClass);4180    bool Succeed = BuildMI(*InsertionPoint.getParent(), InsertionPoint,4181                           InsertionPoint.getDebugLoc(),4182                           TII.get(SPIRV::OpCompositeExtract))4183                       .addDef(ComponentReg)4184                       .addUse(ScalarType->getOperand(0).getReg())4185                       .addUse(ReadReg)4186                       .addImm(I)4187                       .constrainAllUses(TII, TRI, RBI);4188    if (!Succeed)4189      return false;4190    ComponentRegisters.emplace_back(ComponentReg);4191  }4192 4193  MachineInstrBuilder MIB = BuildMI(*InsertionPoint.getParent(), InsertionPoint,4194                                    InsertionPoint.getDebugLoc(),4195                                    TII.get(SPIRV::OpCompositeConstruct))4196                                .addDef(ResVReg)4197                                .addUse(GR.getSPIRVTypeID(ResType));4198 4199  for (Register ComponentReg : ComponentRegisters)4200    MIB.addUse(ComponentReg);4201  return MIB.constrainAllUses(TII, TRI, RBI);4202}4203 4204bool SPIRVInstructionSelector::selectImageWriteIntrinsic(4205    MachineInstr &I) const {4206  // If the load of the image is in a different basic block, then4207  // this will generate invalid code. A proper solution is to move4208  // the OpLoad from selectHandleFromBinding here. However, to do4209  // that we will need to change the return type of the intrinsic.4210  // We will do that when we can, but for now trying to move forward with other4211  // issues.4212  Register ImageReg = I.getOperand(1).getReg();4213  auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));4214  Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));4215  if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),4216                                *ImageDef, I)) {4217    return false;4218  }4219 4220  Register CoordinateReg = I.getOperand(2).getReg();4221  Register DataReg = I.getOperand(3).getReg();4222  assert(GR.getResultType(DataReg)->getOpcode() == SPIRV::OpTypeVector);4223  assert(GR.getScalarOrVectorComponentCount(GR.getResultType(DataReg)) == 4);4224  return BuildMI(*I.getParent(), I, I.getDebugLoc(),4225                 TII.get(SPIRV::OpImageWrite))4226      .addUse(NewImageReg)4227      .addUse(CoordinateReg)4228      .addUse(DataReg)4229      .constrainAllUses(TII, TRI, RBI);4230}4231 4232Register SPIRVInstructionSelector::buildPointerToResource(4233    const SPIRVType *SpirvResType, SPIRV::StorageClass::StorageClass SC,4234    uint32_t Set, uint32_t Binding, uint32_t ArraySize, Register IndexReg,4235    StringRef Name, MachineIRBuilder MIRBuilder) const {4236  const Type *ResType = GR.getTypeForSPIRVType(SpirvResType);4237  if (ArraySize == 1) {4238    SPIRVType *PtrType =4239        GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);4240    assert(GR.getPointeeType(PtrType) == SpirvResType &&4241           "SpirvResType did not have an explicit layout.");4242    return GR.getOrCreateGlobalVariableWithBinding(PtrType, Set, Binding, Name,4243                                                   MIRBuilder);4244  }4245 4246  const Type *VarType = ArrayType::get(const_cast<Type *>(ResType), ArraySize);4247  SPIRVType *VarPointerType =4248      GR.getOrCreateSPIRVPointerType(VarType, MIRBuilder, SC);4249  Register VarReg = GR.getOrCreateGlobalVariableWithBinding(4250      VarPointerType, Set, Binding, Name, MIRBuilder);4251 4252  SPIRVType *ResPointerType =4253      GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);4254  Register AcReg = MRI->createVirtualRegister(GR.getRegClass(ResPointerType));4255 4256  MIRBuilder.buildInstr(SPIRV::OpAccessChain)4257      .addDef(AcReg)4258      .addUse(GR.getSPIRVTypeID(ResPointerType))4259      .addUse(VarReg)4260      .addUse(IndexReg);4261 4262  return AcReg;4263}4264 4265bool SPIRVInstructionSelector::selectFirstBitSet16(4266    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,4267    unsigned ExtendOpcode, unsigned BitSetOpcode) const {4268  Register ExtReg = MRI->createVirtualRegister(GR.getRegClass(ResType));4269  bool Result = selectOpWithSrcs(ExtReg, ResType, I, {I.getOperand(2).getReg()},4270                                 ExtendOpcode);4271 4272  return Result &&4273         selectFirstBitSet32(ResVReg, ResType, I, ExtReg, BitSetOpcode);4274}4275 4276bool SPIRVInstructionSelector::selectFirstBitSet32(4277    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,4278    Register SrcReg, unsigned BitSetOpcode) const {4279  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))4280      .addDef(ResVReg)4281      .addUse(GR.getSPIRVTypeID(ResType))4282      .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))4283      .addImm(BitSetOpcode)4284      .addUse(SrcReg)4285      .constrainAllUses(TII, TRI, RBI);4286}4287 4288bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(4289    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,4290    Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {4291 4292  // SPIR-V allow vectors of size 2,3,4 only. Calling with a larger vectors4293  // requires creating a param register and return register with an invalid4294  // vector size. If that is resolved, then this function can be used for4295  // vectors of any component size.4296  unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);4297  assert(ComponentCount < 5 && "Vec 5+ will generate invalid SPIR-V ops");4298 4299  MachineIRBuilder MIRBuilder(I);4300  SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);4301  SPIRVType *I64Type = GR.getOrCreateSPIRVIntegerType(64, MIRBuilder);4302  SPIRVType *I64x2Type =4303      GR.getOrCreateSPIRVVectorType(I64Type, 2, MIRBuilder, false);4304  SPIRVType *Vec2ResType =4305      GR.getOrCreateSPIRVVectorType(BaseType, 2, MIRBuilder, false);4306 4307  std::vector<Register> PartialRegs;4308 4309  // Loops 0, 2, 4, ... but stops one loop early when ComponentCount is odd4310  unsigned CurrentComponent = 0;4311  for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {4312    // This register holds the firstbitX result for each of the i64x2 vectors4313    // extracted from SrcReg4314    Register BitSetResult =4315        MRI->createVirtualRegister(GR.getRegClass(I64x2Type));4316 4317    auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),4318                       TII.get(SPIRV::OpVectorShuffle))4319                   .addDef(BitSetResult)4320                   .addUse(GR.getSPIRVTypeID(I64x2Type))4321                   .addUse(SrcReg)4322                   .addUse(SrcReg)4323                   .addImm(CurrentComponent)4324                   .addImm(CurrentComponent + 1);4325 4326    if (!MIB.constrainAllUses(TII, TRI, RBI))4327      return false;4328 4329    Register SubVecBitSetReg =4330        MRI->createVirtualRegister(GR.getRegClass(Vec2ResType));4331 4332    if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType, I, BitSetResult,4333                             BitSetOpcode, SwapPrimarySide))4334      return false;4335 4336    PartialRegs.push_back(SubVecBitSetReg);4337  }4338 4339  // On odd component counts we need to handle one more component4340  if (CurrentComponent != ComponentCount) {4341    bool ZeroAsNull = !STI.isShader();4342    Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I64Type));4343    Register ConstIntLastIdx = GR.getOrCreateConstInt(4344        ComponentCount - 1, I, BaseType, TII, ZeroAsNull);4345 4346    if (!selectOpWithSrcs(FinalElemReg, I64Type, I, {SrcReg, ConstIntLastIdx},4347                          SPIRV::OpVectorExtractDynamic))4348      return false;4349 4350    Register FinalElemBitSetReg =4351        MRI->createVirtualRegister(GR.getRegClass(BaseType));4352 4353    if (!selectFirstBitSet64(FinalElemBitSetReg, BaseType, I, FinalElemReg,4354                             BitSetOpcode, SwapPrimarySide))4355      return false;4356 4357    PartialRegs.push_back(FinalElemBitSetReg);4358  }4359 4360  // Join all the resulting registers back into the return type in order4361  // (ie i32x2, i32x2, i32x1 -> i32x5)4362  return selectOpWithSrcs(ResVReg, ResType, I, std::move(PartialRegs),4363                          SPIRV::OpCompositeConstruct);4364}4365 4366bool SPIRVInstructionSelector::selectFirstBitSet64(4367    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,4368    Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {4369  unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);4370  SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);4371  bool ZeroAsNull = !STI.isShader();4372  Register ConstIntZero =4373      GR.getOrCreateConstInt(0, I, BaseType, TII, ZeroAsNull);4374  Register ConstIntOne =4375      GR.getOrCreateConstInt(1, I, BaseType, TII, ZeroAsNull);4376 4377  // SPIRV doesn't support vectors with more than 4 components. Since the4378  // algoritm below converts i64 -> i32x2 and i64x4 -> i32x8 it can only4379  // operate on vectors with 2 or less components. When largers vectors are4380  // seen. Split them, recurse, then recombine them.4381  if (ComponentCount > 2) {4382    return selectFirstBitSet64Overflow(ResVReg, ResType, I, SrcReg,4383                                       BitSetOpcode, SwapPrimarySide);4384  }4385 4386  // 1. Split int64 into 2 pieces using a bitcast4387  MachineIRBuilder MIRBuilder(I);4388  SPIRVType *PostCastType = GR.getOrCreateSPIRVVectorType(4389      BaseType, 2 * ComponentCount, MIRBuilder, false);4390  Register BitcastReg =4391      MRI->createVirtualRegister(GR.getRegClass(PostCastType));4392 4393  if (!selectOpWithSrcs(BitcastReg, PostCastType, I, {SrcReg},4394                        SPIRV::OpBitcast))4395    return false;4396 4397  // 2. Find the first set bit from the primary side for all the pieces in #14398  Register FBSReg = MRI->createVirtualRegister(GR.getRegClass(PostCastType));4399  if (!selectFirstBitSet32(FBSReg, PostCastType, I, BitcastReg, BitSetOpcode))4400    return false;4401 4402  // 3. Split result vector into high bits and low bits4403  Register HighReg = MRI->createVirtualRegister(GR.getRegClass(ResType));4404  Register LowReg = MRI->createVirtualRegister(GR.getRegClass(ResType));4405 4406  bool IsScalarRes = ResType->getOpcode() != SPIRV::OpTypeVector;4407  if (IsScalarRes) {4408    // if scalar do a vector extract4409    if (!selectOpWithSrcs(HighReg, ResType, I, {FBSReg, ConstIntZero},4410                          SPIRV::OpVectorExtractDynamic))4411      return false;4412    if (!selectOpWithSrcs(LowReg, ResType, I, {FBSReg, ConstIntOne},4413                          SPIRV::OpVectorExtractDynamic))4414      return false;4415  } else {4416    // if vector do a shufflevector4417    auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),4418                       TII.get(SPIRV::OpVectorShuffle))4419                   .addDef(HighReg)4420                   .addUse(GR.getSPIRVTypeID(ResType))4421                   .addUse(FBSReg)4422                   // Per the spec, repeat the vector if only one vec is needed4423                   .addUse(FBSReg);4424 4425    // high bits are stored in even indexes. Extract them from FBSReg4426    for (unsigned J = 0; J < ComponentCount * 2; J += 2) {4427      MIB.addImm(J);4428    }4429 4430    if (!MIB.constrainAllUses(TII, TRI, RBI))4431      return false;4432 4433    MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),4434                  TII.get(SPIRV::OpVectorShuffle))4435              .addDef(LowReg)4436              .addUse(GR.getSPIRVTypeID(ResType))4437              .addUse(FBSReg)4438              // Per the spec, repeat the vector if only one vec is needed4439              .addUse(FBSReg);4440 4441    // low bits are stored in odd indexes. Extract them from FBSReg4442    for (unsigned J = 1; J < ComponentCount * 2; J += 2) {4443      MIB.addImm(J);4444    }4445    if (!MIB.constrainAllUses(TII, TRI, RBI))4446      return false;4447  }4448 4449  // 4. Check the result. When primary bits == -1 use secondary, otherwise use4450  // primary4451  SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);4452  Register NegOneReg;4453  Register Reg0;4454  Register Reg32;4455  unsigned SelectOp;4456  unsigned AddOp;4457 4458  if (IsScalarRes) {4459    NegOneReg =4460        GR.getOrCreateConstInt((unsigned)-1, I, ResType, TII, ZeroAsNull);4461    Reg0 = GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);4462    Reg32 = GR.getOrCreateConstInt(32, I, ResType, TII, ZeroAsNull);4463    SelectOp = SPIRV::OpSelectSISCond;4464    AddOp = SPIRV::OpIAddS;4465  } else {4466    BoolType = GR.getOrCreateSPIRVVectorType(BoolType, ComponentCount,4467                                             MIRBuilder, false);4468    NegOneReg =4469        GR.getOrCreateConstVector((unsigned)-1, I, ResType, TII, ZeroAsNull);4470    Reg0 = GR.getOrCreateConstVector(0, I, ResType, TII, ZeroAsNull);4471    Reg32 = GR.getOrCreateConstVector(32, I, ResType, TII, ZeroAsNull);4472    SelectOp = SPIRV::OpSelectVIVCond;4473    AddOp = SPIRV::OpIAddV;4474  }4475 4476  Register PrimaryReg = HighReg;4477  Register SecondaryReg = LowReg;4478  Register PrimaryShiftReg = Reg32;4479  Register SecondaryShiftReg = Reg0;4480 4481  // By default the emitted opcodes check for the set bit from the MSB side.4482  // Setting SwapPrimarySide checks the set bit from the LSB side4483  if (SwapPrimarySide) {4484    PrimaryReg = LowReg;4485    SecondaryReg = HighReg;4486    PrimaryShiftReg = Reg0;4487    SecondaryShiftReg = Reg32;4488  }4489 4490  // Check if the primary bits are == -14491  Register BReg = MRI->createVirtualRegister(GR.getRegClass(BoolType));4492  if (!selectOpWithSrcs(BReg, BoolType, I, {PrimaryReg, NegOneReg},4493                        SPIRV::OpIEqual))4494    return false;4495 4496  // Select secondary bits if true in BReg, otherwise primary bits4497  Register TmpReg = MRI->createVirtualRegister(GR.getRegClass(ResType));4498  if (!selectOpWithSrcs(TmpReg, ResType, I, {BReg, SecondaryReg, PrimaryReg},4499                        SelectOp))4500    return false;4501 4502  // 5. Add 32 when high bits are used, otherwise 0 for low bits4503  Register ValReg = MRI->createVirtualRegister(GR.getRegClass(ResType));4504  if (!selectOpWithSrcs(ValReg, ResType, I,4505                        {BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))4506    return false;4507 4508  return selectOpWithSrcs(ResVReg, ResType, I, {ValReg, TmpReg}, AddOp);4509}4510 4511bool SPIRVInstructionSelector::selectFirstBitHigh(Register ResVReg,4512                                                  const SPIRVType *ResType,4513                                                  MachineInstr &I,4514                                                  bool IsSigned) const {4515  // FindUMsb and FindSMsb intrinsics only support 32 bit integers4516  Register OpReg = I.getOperand(2).getReg();4517  SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);4518  // zero or sign extend4519  unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;4520  unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;4521 4522  switch (GR.getScalarOrVectorBitWidth(OpType)) {4523  case 16:4524    return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);4525  case 32:4526    return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);4527  case 64:4528    return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,4529                               /*SwapPrimarySide=*/false);4530  default:4531    report_fatal_error(4532        "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");4533  }4534}4535 4536bool SPIRVInstructionSelector::selectFirstBitLow(Register ResVReg,4537                                                 const SPIRVType *ResType,4538                                                 MachineInstr &I) const {4539  // FindILsb intrinsic only supports 32 bit integers4540  Register OpReg = I.getOperand(2).getReg();4541  SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);4542  // OpUConvert treats the operand bits as an unsigned i16 and zero extends it4543  // to an unsigned i32. As this leaves all the least significant bits unchanged4544  // so the first set bit from the LSB side doesn't change.4545  unsigned ExtendOpcode = SPIRV::OpUConvert;4546  unsigned BitSetOpcode = GL::FindILsb;4547 4548  switch (GR.getScalarOrVectorBitWidth(OpType)) {4549  case 16:4550    return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);4551  case 32:4552    return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);4553  case 64:4554    return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,4555                               /*SwapPrimarySide=*/true);4556  default:4557    report_fatal_error("spv_firstbitlow only supports 16,32,64 bits.");4558  }4559}4560 4561bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,4562                                                 const SPIRVType *ResType,4563                                                 MachineInstr &I) const {4564  // there was an allocation size parameter to the allocation instruction4565  // that is not 14566  MachineBasicBlock &BB = *I.getParent();4567  bool Res = BuildMI(BB, I, I.getDebugLoc(),4568                     TII.get(SPIRV::OpVariableLengthArrayINTEL))4569                 .addDef(ResVReg)4570                 .addUse(GR.getSPIRVTypeID(ResType))4571                 .addUse(I.getOperand(2).getReg())4572                 .constrainAllUses(TII, TRI, RBI);4573  if (!STI.isShader()) {4574    unsigned Alignment = I.getOperand(3).getImm();4575    buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::Alignment, {Alignment});4576  }4577  return Res;4578}4579 4580bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,4581                                                const SPIRVType *ResType,4582                                                MachineInstr &I) const {4583  // Change order of instructions if needed: all OpVariable instructions in a4584  // function must be the first instructions in the first block4585  auto It = getOpVariableMBBIt(I);4586  bool Res = BuildMI(*It->getParent(), It, It->getDebugLoc(),4587                     TII.get(SPIRV::OpVariable))4588                 .addDef(ResVReg)4589                 .addUse(GR.getSPIRVTypeID(ResType))4590                 .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))4591                 .constrainAllUses(TII, TRI, RBI);4592  if (!STI.isShader()) {4593    unsigned Alignment = I.getOperand(2).getImm();4594    buildOpDecorate(ResVReg, *It, TII, SPIRV::Decoration::Alignment,4595                    {Alignment});4596  }4597  return Res;4598}4599 4600bool SPIRVInstructionSelector::selectBranch(MachineInstr &I) const {4601  // InstructionSelector walks backwards through the instructions. We can use4602  // both a G_BR and a G_BRCOND to create an OpBranchConditional. We hit G_BR4603  // first, so can generate an OpBranchConditional here. If there is no4604  // G_BRCOND, we just use OpBranch for a regular unconditional branch.4605  const MachineInstr *PrevI = I.getPrevNode();4606  MachineBasicBlock &MBB = *I.getParent();4607  if (PrevI != nullptr && PrevI->getOpcode() == TargetOpcode::G_BRCOND) {4608    return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))4609        .addUse(PrevI->getOperand(0).getReg())4610        .addMBB(PrevI->getOperand(1).getMBB())4611        .addMBB(I.getOperand(0).getMBB())4612        .constrainAllUses(TII, TRI, RBI);4613  }4614  return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranch))4615      .addMBB(I.getOperand(0).getMBB())4616      .constrainAllUses(TII, TRI, RBI);4617}4618 4619bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &I) const {4620  // InstructionSelector walks backwards through the instructions. For an4621  // explicit conditional branch with no fallthrough, we use both a G_BR and a4622  // G_BRCOND to create an OpBranchConditional. We should hit G_BR first, and4623  // generate the OpBranchConditional in selectBranch above.4624  //4625  // If an OpBranchConditional has been generated, we simply return, as the work4626  // is alread done. If there is no OpBranchConditional, LLVM must be relying on4627  // implicit fallthrough to the next basic block, so we need to create an4628  // OpBranchConditional with an explicit "false" argument pointing to the next4629  // basic block that LLVM would fall through to.4630  const MachineInstr *NextI = I.getNextNode();4631  // Check if this has already been successfully selected.4632  if (NextI != nullptr && NextI->getOpcode() == SPIRV::OpBranchConditional)4633    return true;4634  // Must be relying on implicit block fallthrough, so generate an4635  // OpBranchConditional with the "next" basic block as the "false" target.4636  MachineBasicBlock &MBB = *I.getParent();4637  unsigned NextMBBNum = MBB.getNextNode()->getNumber();4638  MachineBasicBlock *NextMBB = I.getMF()->getBlockNumbered(NextMBBNum);4639  return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))4640      .addUse(I.getOperand(0).getReg())4641      .addMBB(I.getOperand(1).getMBB())4642      .addMBB(NextMBB)4643      .constrainAllUses(TII, TRI, RBI);4644}4645 4646bool SPIRVInstructionSelector::selectPhi(Register ResVReg,4647                                         const SPIRVType *ResType,4648                                         MachineInstr &I) const {4649  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpPhi))4650                 .addDef(ResVReg)4651                 .addUse(GR.getSPIRVTypeID(ResType));4652  const unsigned NumOps = I.getNumOperands();4653  for (unsigned i = 1; i < NumOps; i += 2) {4654    MIB.addUse(I.getOperand(i + 0).getReg());4655    MIB.addMBB(I.getOperand(i + 1).getMBB());4656  }4657  bool Res = MIB.constrainAllUses(TII, TRI, RBI);4658  MIB->setDesc(TII.get(TargetOpcode::PHI));4659  MIB->removeOperand(1);4660  return Res;4661}4662 4663bool SPIRVInstructionSelector::selectGlobalValue(4664    Register ResVReg, MachineInstr &I, const MachineInstr *Init) const {4665  // FIXME: don't use MachineIRBuilder here, replace it with BuildMI.4666  MachineIRBuilder MIRBuilder(I);4667  const GlobalValue *GV = I.getOperand(1).getGlobal();4668  Type *GVType = toTypedPointer(GR.getDeducedGlobalValueType(GV));4669 4670  std::string GlobalIdent;4671  if (!GV->hasName()) {4672    unsigned &ID = UnnamedGlobalIDs[GV];4673    if (ID == 0)4674      ID = UnnamedGlobalIDs.size();4675    GlobalIdent = "__unnamed_" + Twine(ID).str();4676  } else {4677    GlobalIdent = GV->getName();4678  }4679 4680  // Behaviour of functions as operands depends on availability of the4681  // corresponding extension (SPV_INTEL_function_pointers):4682  // - If there is an extension to operate with functions as operands:4683  // We create a proper constant operand and evaluate a correct type for a4684  // function pointer.4685  // - Without the required extension:4686  // We have functions as operands in tests with blocks of instruction e.g. in4687  // transcoding/global_block.ll. These operands are not used and should be4688  // substituted by zero constants. Their type is expected to be always4689  // OpTypePointer Function %uchar.4690  if (isa<Function>(GV)) {4691    const Constant *ConstVal = GV;4692    MachineBasicBlock &BB = *I.getParent();4693    Register NewReg = GR.find(ConstVal, GR.CurMF);4694    if (!NewReg.isValid()) {4695      Register NewReg = ResVReg;4696      const Function *GVFun =4697          STI.canUseExtension(SPIRV::Extension::SPV_INTEL_function_pointers)4698              ? dyn_cast<Function>(GV)4699              : nullptr;4700      SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(4701          GVType, I,4702          GVFun ? SPIRV::StorageClass::CodeSectionINTEL4703                : addressSpaceToStorageClass(GV->getAddressSpace(), STI));4704      if (GVFun) {4705        // References to a function via function pointers generate virtual4706        // registers without a definition. We will resolve it later, during4707        // module analysis stage.4708        Register ResTypeReg = GR.getSPIRVTypeID(ResType);4709        MachineRegisterInfo *MRI = MIRBuilder.getMRI();4710        Register FuncVReg =4711            MRI->createGenericVirtualRegister(GR.getRegType(ResType));4712        MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);4713        MachineInstrBuilder MIB1 =4714            BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))4715                .addDef(FuncVReg)4716                .addUse(ResTypeReg);4717        MachineInstrBuilder MIB2 =4718            BuildMI(BB, I, I.getDebugLoc(),4719                    TII.get(SPIRV::OpConstantFunctionPointerINTEL))4720                .addDef(NewReg)4721                .addUse(ResTypeReg)4722                .addUse(FuncVReg);4723        GR.add(ConstVal, MIB2);4724        // mapping the function pointer to the used Function4725        GR.recordFunctionPointer(&MIB2.getInstr()->getOperand(2), GVFun);4726        return MIB1.constrainAllUses(TII, TRI, RBI) &&4727               MIB2.constrainAllUses(TII, TRI, RBI);4728      }4729      MachineInstrBuilder MIB3 =4730          BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))4731              .addDef(NewReg)4732              .addUse(GR.getSPIRVTypeID(ResType));4733      GR.add(ConstVal, MIB3);4734      return MIB3.constrainAllUses(TII, TRI, RBI);4735    }4736    assert(NewReg != ResVReg);4737    return BuildCOPY(ResVReg, NewReg, I);4738  }4739  auto GlobalVar = cast<GlobalVariable>(GV);4740  assert(GlobalVar->getName() != "llvm.global.annotations");4741 4742  // Skip empty declaration for GVs with initializers till we get the decl with4743  // passed initializer.4744  if (hasInitializer(GlobalVar) && !Init)4745    return true;4746 4747  const std::optional<SPIRV::LinkageType::LinkageType> LnkType =4748      getSpirvLinkageTypeFor(STI, *GV);4749 4750  const unsigned AddrSpace = GV->getAddressSpace();4751  SPIRV::StorageClass::StorageClass StorageClass =4752      addressSpaceToStorageClass(AddrSpace, STI);4753  SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(GVType, I, StorageClass);4754  Register Reg = GR.buildGlobalVariable(4755      ResVReg, ResType, GlobalIdent, GV, StorageClass, Init,4756      GlobalVar->isConstant(), LnkType, MIRBuilder, true);4757  return Reg.isValid();4758}4759 4760bool SPIRVInstructionSelector::selectLog10(Register ResVReg,4761                                           const SPIRVType *ResType,4762                                           MachineInstr &I) const {4763  if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {4764    return selectExtInst(ResVReg, ResType, I, CL::log10);4765  }4766 4767  // There is no log10 instruction in the GLSL Extended Instruction set, so it4768  // is implemented as:4769  // log10(x) = log2(x) * (1 / log2(10))4770  //          = log2(x) * 0.301034771 4772  MachineIRBuilder MIRBuilder(I);4773  MachineBasicBlock &BB = *I.getParent();4774 4775  // Build log2(x).4776  Register VarReg = MRI->createVirtualRegister(GR.getRegClass(ResType));4777  bool Result =4778      BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))4779          .addDef(VarReg)4780          .addUse(GR.getSPIRVTypeID(ResType))4781          .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))4782          .addImm(GL::Log2)4783          .add(I.getOperand(1))4784          .constrainAllUses(TII, TRI, RBI);4785 4786  // Build 0.30103.4787  assert(ResType->getOpcode() == SPIRV::OpTypeVector ||4788         ResType->getOpcode() == SPIRV::OpTypeFloat);4789  // TODO: Add matrix implementation once supported by the HLSL frontend.4790  const SPIRVType *SpirvScalarType =4791      ResType->getOpcode() == SPIRV::OpTypeVector4792          ? GR.getSPIRVTypeForVReg(ResType->getOperand(1).getReg())4793          : ResType;4794  Register ScaleReg =4795      GR.buildConstantFP(APFloat(0.30103f), MIRBuilder, SpirvScalarType);4796 4797  // Multiply log2(x) by 0.30103 to get log10(x) result.4798  auto Opcode = ResType->getOpcode() == SPIRV::OpTypeVector4799                    ? SPIRV::OpVectorTimesScalar4800                    : SPIRV::OpFMulS;4801  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))4802                       .addDef(ResVReg)4803                       .addUse(GR.getSPIRVTypeID(ResType))4804                       .addUse(VarReg)4805                       .addUse(ScaleReg)4806                       .constrainAllUses(TII, TRI, RBI);4807}4808 4809bool SPIRVInstructionSelector::selectModf(Register ResVReg,4810                                          const SPIRVType *ResType,4811                                          MachineInstr &I) const {4812  // llvm.modf has a single arg --the number to be decomposed-- and returns a4813  // struct { restype, restype }, while OpenCLLIB::modf has two args --the4814  // number to be decomposed and a pointer--, returns the fractional part and4815  // the integral part is stored in the pointer argument. Therefore, we can't4816  // use directly the OpenCLLIB::modf intrinsic. However, we can do some4817  // scaffolding to make it work. The idea is to create an alloca instruction4818  // to get a ptr, pass this ptr to OpenCL::modf, and then load the value4819  // from this ptr to place it in the struct. llvm.modf returns the fractional4820  // part as the first element of the result, and the integral part as the4821  // second element of the result.4822 4823  // At this point, the return type is not a struct anymore, but rather two4824  // independent elements of SPIRVResType. We can get each independent element4825  // from I.getDefs() or I.getOperands().4826  if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {4827    MachineIRBuilder MIRBuilder(I);4828    // Get pointer type for alloca variable.4829    const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(4830        ResType, MIRBuilder, SPIRV::StorageClass::Function);4831    // Create new register for the pointer type of alloca variable.4832    Register PtrTyReg =4833        MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);4834    MIRBuilder.getMRI()->setType(4835        PtrTyReg,4836        LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Function),4837                     GR.getPointerSize()));4838 4839    // Assign SPIR-V type of the pointer type of the alloca variable to the4840    // new register.4841    GR.assignSPIRVTypeToVReg(PtrType, PtrTyReg, MIRBuilder.getMF());4842    MachineBasicBlock &EntryBB = I.getMF()->front();4843    MachineBasicBlock::iterator VarPos =4844        getFirstValidInstructionInsertPoint(EntryBB);4845    auto AllocaMIB =4846        BuildMI(EntryBB, VarPos, I.getDebugLoc(), TII.get(SPIRV::OpVariable))4847            .addDef(PtrTyReg)4848            .addUse(GR.getSPIRVTypeID(PtrType))4849            .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function));4850    Register Variable = AllocaMIB->getOperand(0).getReg();4851 4852    MachineBasicBlock &BB = *I.getParent();4853    // Create the OpenCLLIB::modf instruction.4854    auto MIB =4855        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))4856            .addDef(ResVReg)4857            .addUse(GR.getSPIRVTypeID(ResType))4858            .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))4859            .addImm(CL::modf)4860            .setMIFlags(I.getFlags())4861            .add(I.getOperand(I.getNumExplicitDefs())) // Floating point value.4862            .addUse(Variable); // Pointer to integral part.4863    // Assign the integral part stored in the ptr to the second element of the4864    // result.4865    Register IntegralPartReg = I.getOperand(1).getReg();4866    if (IntegralPartReg.isValid()) {4867      // Load the value from the pointer to integral part.4868      auto LoadMIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))4869                         .addDef(IntegralPartReg)4870                         .addUse(GR.getSPIRVTypeID(ResType))4871                         .addUse(Variable);4872      return LoadMIB.constrainAllUses(TII, TRI, RBI);4873    }4874 4875    return MIB.constrainAllUses(TII, TRI, RBI);4876  } else if (STI.canUseExtInstSet(SPIRV::InstructionSet::GLSL_std_450)) {4877    assert(false && "GLSL::Modf is deprecated.");4878    // FIXME: GL::Modf is deprecated, use Modfstruct instead.4879    return false;4880  }4881  return false;4882}4883 4884// Generate the instructions to load 3-element vector builtin input4885// IDs/Indices.4886// Like: GlobalInvocationId, LocalInvocationId, etc....4887 4888bool SPIRVInstructionSelector::loadVec3BuiltinInputID(4889    SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,4890    const SPIRVType *ResType, MachineInstr &I) const {4891  MachineIRBuilder MIRBuilder(I);4892  const SPIRVType *Vec3Ty =4893      GR.getOrCreateSPIRVVectorType(ResType, 3, MIRBuilder, false);4894  const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(4895      Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);4896 4897  // Create new register for the input ID builtin variable.4898  Register NewRegister =4899      MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);4900  MIRBuilder.getMRI()->setType(NewRegister, LLT::pointer(0, 64));4901  GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());4902 4903  // Build global variable with the necessary decorations for the input ID4904  // builtin variable.4905  Register Variable = GR.buildGlobalVariable(4906      NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,4907      SPIRV::StorageClass::Input, nullptr, true, std::nullopt, MIRBuilder,4908      false);4909 4910  // Create new register for loading value.4911  MachineRegisterInfo *MRI = MIRBuilder.getMRI();4912  Register LoadedRegister = MRI->createVirtualRegister(&SPIRV::iIDRegClass);4913  MIRBuilder.getMRI()->setType(LoadedRegister, LLT::pointer(0, 64));4914  GR.assignSPIRVTypeToVReg(Vec3Ty, LoadedRegister, MIRBuilder.getMF());4915 4916  // Load v3uint value from the global variable.4917  bool Result =4918      BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))4919          .addDef(LoadedRegister)4920          .addUse(GR.getSPIRVTypeID(Vec3Ty))4921          .addUse(Variable);4922 4923  // Get the input ID index. Expecting operand is a constant immediate value,4924  // wrapped in a type assignment.4925  assert(I.getOperand(2).isReg());4926  const uint32_t ThreadId = foldImm(I.getOperand(2), MRI);4927 4928  // Extract the input ID from the loaded vector value.4929  MachineBasicBlock &BB = *I.getParent();4930  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))4931                 .addDef(ResVReg)4932                 .addUse(GR.getSPIRVTypeID(ResType))4933                 .addUse(LoadedRegister)4934                 .addImm(ThreadId);4935  return Result && MIB.constrainAllUses(TII, TRI, RBI);4936}4937 4938// Generate the instructions to load 32-bit integer builtin input IDs/Indices.4939// Like LocalInvocationIndex4940bool SPIRVInstructionSelector::loadBuiltinInputID(4941    SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,4942    const SPIRVType *ResType, MachineInstr &I) const {4943  MachineIRBuilder MIRBuilder(I);4944  const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(4945      ResType, MIRBuilder, SPIRV::StorageClass::Input);4946 4947  // Create new register for the input ID builtin variable.4948  Register NewRegister =4949      MIRBuilder.getMRI()->createVirtualRegister(GR.getRegClass(PtrType));4950  MIRBuilder.getMRI()->setType(4951      NewRegister,4952      LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Input),4953                   GR.getPointerSize()));4954  GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());4955 4956  // Build global variable with the necessary decorations for the input ID4957  // builtin variable.4958  Register Variable = GR.buildGlobalVariable(4959      NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,4960      SPIRV::StorageClass::Input, nullptr, true, std::nullopt, MIRBuilder,4961      false);4962 4963  // Load uint value from the global variable.4964  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))4965                 .addDef(ResVReg)4966                 .addUse(GR.getSPIRVTypeID(ResType))4967                 .addUse(Variable);4968 4969  return MIB.constrainAllUses(TII, TRI, RBI);4970}4971 4972SPIRVType *SPIRVInstructionSelector::widenTypeToVec4(const SPIRVType *Type,4973                                                     MachineInstr &I) const {4974  MachineIRBuilder MIRBuilder(I);4975  if (Type->getOpcode() != SPIRV::OpTypeVector)4976    return GR.getOrCreateSPIRVVectorType(Type, 4, MIRBuilder, false);4977 4978  uint64_t VectorSize = Type->getOperand(2).getImm();4979  if (VectorSize == 4)4980    return Type;4981 4982  Register ScalarTypeReg = Type->getOperand(1).getReg();4983  const SPIRVType *ScalarType = GR.getSPIRVTypeForVReg(ScalarTypeReg);4984  return GR.getOrCreateSPIRVVectorType(ScalarType, 4, MIRBuilder, false);4985}4986 4987bool SPIRVInstructionSelector::loadHandleBeforePosition(4988    Register &HandleReg, const SPIRVType *ResType, GIntrinsic &HandleDef,4989    MachineInstr &Pos) const {4990 4991  assert(HandleDef.getIntrinsicID() ==4992         Intrinsic::spv_resource_handlefrombinding);4993  uint32_t Set = foldImm(HandleDef.getOperand(2), MRI);4994  uint32_t Binding = foldImm(HandleDef.getOperand(3), MRI);4995  uint32_t ArraySize = foldImm(HandleDef.getOperand(4), MRI);4996  Register IndexReg = HandleDef.getOperand(5).getReg();4997  std::string Name =4998      getStringValueFromReg(HandleDef.getOperand(6).getReg(), *MRI);4999 5000  bool IsStructuredBuffer = ResType->getOpcode() == SPIRV::OpTypePointer;5001  MachineIRBuilder MIRBuilder(HandleDef);5002  SPIRVType *VarType = ResType;5003  SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;5004 5005  if (IsStructuredBuffer) {5006    VarType = GR.getPointeeType(ResType);5007    SC = GR.getPointerStorageClass(ResType);5008  }5009 5010  Register VarReg = buildPointerToResource(VarType, SC, Set, Binding, ArraySize,5011                                           IndexReg, Name, MIRBuilder);5012 5013  // The handle for the buffer is the pointer to the resource. For an image, the5014  // handle is the image object. So images get an extra load.5015  uint32_t LoadOpcode =5016      IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;5017  GR.assignSPIRVTypeToVReg(ResType, HandleReg, *Pos.getMF());5018  return BuildMI(*Pos.getParent(), Pos, HandleDef.getDebugLoc(),5019                 TII.get(LoadOpcode))5020      .addDef(HandleReg)5021      .addUse(GR.getSPIRVTypeID(ResType))5022      .addUse(VarReg)5023      .constrainAllUses(TII, TRI, RBI);5024}5025 5026void SPIRVInstructionSelector::errorIfInstrOutsideShader(5027    MachineInstr &I) const {5028  if (!STI.isShader()) {5029    std::string DiagMsg;5030    raw_string_ostream OS(DiagMsg);5031    I.print(OS, true, false, false, false);5032    DiagMsg += " is only supported in shaders.\n";5033    report_fatal_error(DiagMsg.c_str(), false);5034  }5035}5036 5037namespace llvm {5038InstructionSelector *5039createSPIRVInstructionSelector(const SPIRVTargetMachine &TM,5040                               const SPIRVSubtarget &Subtarget,5041                               const RegisterBankInfo &RBI) {5042  return new SPIRVInstructionSelector(TM, Subtarget, RBI);5043}5044} // namespace llvm5045