File: SPIRVInstructionSelector.cpp

package info (click to toggle)
llvm-toolchain-21 1%3A21.1.0-1
links: PTS, VCS
area: main
in suites: sid
size: 2,235,796 kB
sloc: cpp: 7,617,614; ansic: 1,433,901; asm: 1,058,726; python: 252,096; f90: 94,671; objc: 70,753; lisp: 42,813; pascal: 18,401; sh: 10,032; ml: 5,111; perl: 4,720; awk: 3,523; makefile: 3,401; javascript: 2,272; xml: 892; fortran: 770
file content (4174 lines) | stat: -rw-r--r-- 177,184 bytes
parent folder | download | duplicates (2)
//===- SPIRVInstructionSelector.cpp ------------------------------*- C++ -*-==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
// This file implements the targeting of the InstructionSelector class for
// SPIRV.
// TODO: This should be generated by TableGen.
//
//===----------------------------------------------------------------------===//

#include "MCTargetDesc/SPIRVBaseInfo.h"
#include "MCTargetDesc/SPIRVMCTargetDesc.h"
#include "SPIRV.h"
#include "SPIRVGlobalRegistry.h"
#include "SPIRVInstrInfo.h"
#include "SPIRVRegisterInfo.h"
#include "SPIRVTargetMachine.h"
#include "SPIRVUtils.h"
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/CodeGen/GlobalISel/GIMatchTableExecutorImpl.h"
#include "llvm/CodeGen/GlobalISel/GenericMachineInstrs.h"
#include "llvm/CodeGen/GlobalISel/InstructionSelector.h"
#include "llvm/CodeGen/MachineInstrBuilder.h"
#include "llvm/CodeGen/MachineRegisterInfo.h"
#include "llvm/CodeGen/Register.h"
#include "llvm/CodeGen/TargetOpcodes.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/ErrorHandling.h"

#define DEBUG_TYPE "spirv-isel"

using namespace llvm;
namespace CL = SPIRV::OpenCLExtInst;
namespace GL = SPIRV::GLSLExtInst;

using ExtInstList =
    std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;

namespace {

llvm::SPIRV::SelectionControl::SelectionControl
getSelectionOperandForImm(int Imm) {
  if (Imm == 2)
    return SPIRV::SelectionControl::Flatten;
  if (Imm == 1)
    return SPIRV::SelectionControl::DontFlatten;
  if (Imm == 0)
    return SPIRV::SelectionControl::None;
  llvm_unreachable("Invalid immediate");
}

#define GET_GLOBALISEL_PREDICATE_BITSET
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_PREDICATE_BITSET

class SPIRVInstructionSelector : public InstructionSelector {
  const SPIRVSubtarget &STI;
  const SPIRVInstrInfo &TII;
  const SPIRVRegisterInfo &TRI;
  const RegisterBankInfo &RBI;
  SPIRVGlobalRegistry &GR;
  MachineRegisterInfo *MRI;
  MachineFunction *HasVRegsReset = nullptr;

  /// We need to keep track of the number we give to anonymous global values to
  /// generate the same name every time when this is needed.
  mutable DenseMap<const GlobalValue *, unsigned> UnnamedGlobalIDs;
  SmallPtrSet<MachineInstr *, 8> DeadMIs;

public:
  SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
                           const SPIRVSubtarget &ST,
                           const RegisterBankInfo &RBI);
  void setupMF(MachineFunction &MF, GISelValueTracking *VT,
               CodeGenCoverage *CoverageInfo, ProfileSummaryInfo *PSI,
               BlockFrequencyInfo *BFI) override;
  // Common selection code. Instruction-specific selection occurs in spvSelect.
  bool select(MachineInstr &I) override;
  static const char *getName() { return DEBUG_TYPE; }

#define GET_GLOBALISEL_PREDICATES_DECL
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_PREDICATES_DECL

#define GET_GLOBALISEL_TEMPORARIES_DECL
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_TEMPORARIES_DECL

private:
  void resetVRegsType(MachineFunction &MF);

  // tblgen-erated 'select' implementation, used as the initial selector for
  // the patterns that don't require complex C++.
  bool selectImpl(MachineInstr &I, CodeGenCoverage &CoverageInfo) const;

  // All instruction-specific selection that didn't happen in "select()".
  // Is basically a large Switch/Case delegating to all other select method.
  bool spvSelect(Register ResVReg, const SPIRVType *ResType,
                 MachineInstr &I) const;

  bool selectFirstBitHigh(Register ResVReg, const SPIRVType *ResType,
                          MachineInstr &I, bool IsSigned) const;

  bool selectFirstBitLow(Register ResVReg, const SPIRVType *ResType,
                         MachineInstr &I) const;

  bool selectFirstBitSet16(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I, unsigned ExtendOpcode,
                           unsigned BitSetOpcode) const;

  bool selectFirstBitSet32(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I, Register SrcReg,
                           unsigned BitSetOpcode) const;

  bool selectFirstBitSet64(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I, Register SrcReg,
                           unsigned BitSetOpcode, bool SwapPrimarySide) const;

  bool selectFirstBitSet64Overflow(Register ResVReg, const SPIRVType *ResType,
                                   MachineInstr &I, Register SrcReg,
                                   unsigned BitSetOpcode,
                                   bool SwapPrimarySide) const;

  bool selectGlobalValue(Register ResVReg, MachineInstr &I,
                         const MachineInstr *Init = nullptr) const;

  bool selectOpWithSrcs(Register ResVReg, const SPIRVType *ResType,
                        MachineInstr &I, std::vector<Register> SrcRegs,
                        unsigned Opcode) const;

  bool selectUnOp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
                  unsigned Opcode) const;

  bool selectBitcast(Register ResVReg, const SPIRVType *ResType,
                     MachineInstr &I) const;

  bool selectLoad(Register ResVReg, const SPIRVType *ResType,
                  MachineInstr &I) const;
  bool selectStore(MachineInstr &I) const;

  bool selectStackSave(Register ResVReg, const SPIRVType *ResType,
                       MachineInstr &I) const;
  bool selectStackRestore(MachineInstr &I) const;

  bool selectMemOperation(Register ResVReg, MachineInstr &I) const;

  bool selectAtomicRMW(Register ResVReg, const SPIRVType *ResType,
                       MachineInstr &I, unsigned NewOpcode,
                       unsigned NegateOpcode = 0) const;

  bool selectAtomicCmpXchg(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I) const;

  bool selectFence(MachineInstr &I) const;

  bool selectAddrSpaceCast(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I) const;

  bool selectAnyOrAll(Register ResVReg, const SPIRVType *ResType,
                      MachineInstr &I, unsigned OpType) const;

  bool selectAll(Register ResVReg, const SPIRVType *ResType,
                 MachineInstr &I) const;

  bool selectAny(Register ResVReg, const SPIRVType *ResType,
                 MachineInstr &I) const;

  bool selectBitreverse(Register ResVReg, const SPIRVType *ResType,
                        MachineInstr &I) const;

  bool selectBuildVector(Register ResVReg, const SPIRVType *ResType,
                         MachineInstr &I) const;
  bool selectSplatVector(Register ResVReg, const SPIRVType *ResType,
                         MachineInstr &I) const;

  bool selectCmp(Register ResVReg, const SPIRVType *ResType,
                 unsigned comparisonOpcode, MachineInstr &I) const;
  bool selectDiscard(Register ResVReg, const SPIRVType *ResType,
                     MachineInstr &I) const;

  bool selectICmp(Register ResVReg, const SPIRVType *ResType,
                  MachineInstr &I) const;
  bool selectFCmp(Register ResVReg, const SPIRVType *ResType,
                  MachineInstr &I) const;

  bool selectSign(Register ResVReg, const SPIRVType *ResType,
                  MachineInstr &I) const;

  bool selectFloatDot(Register ResVReg, const SPIRVType *ResType,
                      MachineInstr &I) const;

  bool selectOverflowArith(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I, unsigned Opcode) const;

  bool selectIntegerDot(Register ResVReg, const SPIRVType *ResType,
                        MachineInstr &I, bool Signed) const;

  bool selectIntegerDotExpansion(Register ResVReg, const SPIRVType *ResType,
                                 MachineInstr &I) const;

  template <bool Signed>
  bool selectDot4AddPacked(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I) const;
  template <bool Signed>
  bool selectDot4AddPackedExpansion(Register ResVReg, const SPIRVType *ResType,
                                    MachineInstr &I) const;

  bool selectWaveReduceMax(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I, bool IsUnsigned) const;

  bool selectWaveReduceSum(Register ResVReg, const SPIRVType *ResType,
                           MachineInstr &I) const;

  bool selectConst(Register ResVReg, const SPIRVType *ResType,
                   MachineInstr &I) const;

  bool selectSelect(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
                    bool IsSigned) const;
  bool selectIToF(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
                  bool IsSigned, unsigned Opcode) const;
  bool selectExt(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
                 bool IsSigned) const;

  bool selectTrunc(Register ResVReg, const SPIRVType *ResType,
                   MachineInstr &I) const;

  bool selectSUCmp(Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
                   bool IsSigned) const;

  bool selectIntToBool(Register IntReg, Register ResVReg, MachineInstr &I,
                       const SPIRVType *intTy, const SPIRVType *boolTy) const;

  bool selectOpUndef(Register ResVReg, const SPIRVType *ResType,
                     MachineInstr &I) const;
  bool selectFreeze(Register ResVReg, const SPIRVType *ResType,
                    MachineInstr &I) const;
  bool selectIntrinsic(Register ResVReg, const SPIRVType *ResType,
                       MachineInstr &I) const;
  bool selectExtractVal(Register ResVReg, const SPIRVType *ResType,
                        MachineInstr &I) const;
  bool selectInsertVal(Register ResVReg, const SPIRVType *ResType,
                       MachineInstr &I) const;
  bool selectExtractElt(Register ResVReg, const SPIRVType *ResType,
                        MachineInstr &I) const;
  bool selectInsertElt(Register ResVReg, const SPIRVType *ResType,
                       MachineInstr &I) const;
  bool selectGEP(Register ResVReg, const SPIRVType *ResType,
                 MachineInstr &I) const;

  bool selectFrameIndex(Register ResVReg, const SPIRVType *ResType,
                        MachineInstr &I) const;
  bool selectAllocaArray(Register ResVReg, const SPIRVType *ResType,
                         MachineInstr &I) const;

  bool selectBranch(MachineInstr &I) const;
  bool selectBranchCond(MachineInstr &I) const;

  bool selectPhi(Register ResVReg, const SPIRVType *ResType,
                 MachineInstr &I) const;

  bool selectExtInst(Register ResVReg, const SPIRVType *RestType,
                     MachineInstr &I, GL::GLSLExtInst GLInst) const;
  bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
                     MachineInstr &I, CL::OpenCLExtInst CLInst) const;
  bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
                     MachineInstr &I, CL::OpenCLExtInst CLInst,
                     GL::GLSLExtInst GLInst) const;
  bool selectExtInst(Register ResVReg, const SPIRVType *ResType,
                     MachineInstr &I, const ExtInstList &ExtInsts) const;

  bool selectLog10(Register ResVReg, const SPIRVType *ResType,
                   MachineInstr &I) const;

  bool selectSaturate(Register ResVReg, const SPIRVType *ResType,
                      MachineInstr &I) const;

  bool selectWaveOpInst(Register ResVReg, const SPIRVType *ResType,
                        MachineInstr &I, unsigned Opcode) const;

  bool selectWaveActiveCountBits(Register ResVReg, const SPIRVType *ResType,
                                 MachineInstr &I) const;

  bool selectUnmergeValues(MachineInstr &I) const;

  bool selectHandleFromBinding(Register &ResVReg, const SPIRVType *ResType,
                               MachineInstr &I) const;

  bool selectReadImageIntrinsic(Register &ResVReg, const SPIRVType *ResType,
                                MachineInstr &I) const;
  bool selectImageWriteIntrinsic(MachineInstr &I) const;
  bool selectResourceGetPointer(Register &ResVReg, const SPIRVType *ResType,
                                MachineInstr &I) const;

  // Utilities
  std::pair<Register, bool>
  buildI32Constant(uint32_t Val, MachineInstr &I,
                   const SPIRVType *ResType = nullptr) const;

  Register buildZerosVal(const SPIRVType *ResType, MachineInstr &I) const;
  Register buildZerosValF(const SPIRVType *ResType, MachineInstr &I) const;
  Register buildOnesVal(bool AllOnes, const SPIRVType *ResType,
                        MachineInstr &I) const;
  Register buildOnesValF(const SPIRVType *ResType, MachineInstr &I) const;

  bool wrapIntoSpecConstantOp(MachineInstr &I,
                              SmallVector<Register> &CompositeArgs) const;

  Register getUcharPtrTypeReg(MachineInstr &I,
                              SPIRV::StorageClass::StorageClass SC) const;
  MachineInstrBuilder buildSpecConstantOp(MachineInstr &I, Register Dest,
                                          Register Src, Register DestType,
                                          uint32_t Opcode) const;
  MachineInstrBuilder buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
                                           SPIRVType *SrcPtrTy) const;
  Register buildPointerToResource(const SPIRVType *ResType,
                                  SPIRV::StorageClass::StorageClass SC,
                                  uint32_t Set, uint32_t Binding,
                                  uint32_t ArraySize, Register IndexReg,
                                  bool IsNonUniform, StringRef Name,
                                  MachineIRBuilder MIRBuilder) const;
  SPIRVType *widenTypeToVec4(const SPIRVType *Type, MachineInstr &I) const;
  bool extractSubvector(Register &ResVReg, const SPIRVType *ResType,
                        Register &ReadReg, MachineInstr &InsertionPoint) const;
  bool generateImageRead(Register &ResVReg, const SPIRVType *ResType,
                         Register ImageReg, Register IdxReg, DebugLoc Loc,
                         MachineInstr &Pos) const;
  bool BuildCOPY(Register DestReg, Register SrcReg, MachineInstr &I) const;
  bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
                              Register ResVReg, const SPIRVType *ResType,
                              MachineInstr &I) const;
  bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
                          Register ResVReg, const SPIRVType *ResType,
                          MachineInstr &I) const;
  bool loadHandleBeforePosition(Register &HandleReg, const SPIRVType *ResType,
                                GIntrinsic &HandleDef, MachineInstr &Pos) const;
};

bool sampledTypeIsSignedInteger(const llvm::Type *HandleType) {
  const TargetExtType *TET = cast<TargetExtType>(HandleType);
  if (TET->getTargetExtName() == "spirv.Image") {
    return false;
  }
  assert(TET->getTargetExtName() == "spirv.SignedImage");
  return TET->getTypeParameter(0)->isIntegerTy();
}
} // end anonymous namespace

#define GET_GLOBALISEL_IMPL
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_IMPL

SPIRVInstructionSelector::SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
                                                   const SPIRVSubtarget &ST,
                                                   const RegisterBankInfo &RBI)
    : InstructionSelector(), STI(ST), TII(*ST.getInstrInfo()),
      TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
#define GET_GLOBALISEL_PREDICATES_INIT
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_TEMPORARIES_INIT
{
}

void SPIRVInstructionSelector::setupMF(MachineFunction &MF,
                                       GISelValueTracking *VT,
                                       CodeGenCoverage *CoverageInfo,
                                       ProfileSummaryInfo *PSI,
                                       BlockFrequencyInfo *BFI) {
  MRI = &MF.getRegInfo();
  GR.setCurrentFunc(MF);
  InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
}

// Ensure that register classes correspond to pattern matching rules.
void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
  if (HasVRegsReset == &MF)
    return;
  HasVRegsReset = &MF;

  MachineRegisterInfo &MRI = MF.getRegInfo();
  for (unsigned I = 0, E = MRI.getNumVirtRegs(); I != E; ++I) {
    Register Reg = Register::index2VirtReg(I);
    LLT RegType = MRI.getType(Reg);
    if (RegType.isScalar())
      MRI.setType(Reg, LLT::scalar(64));
    else if (RegType.isPointer())
      MRI.setType(Reg, LLT::pointer(0, 64));
    else if (RegType.isVector())
      MRI.setType(Reg, LLT::fixed_vector(2, LLT::scalar(64)));
  }
  for (const auto &MBB : MF) {
    for (const auto &MI : MBB) {
      if (isPreISelGenericOpcode(MI.getOpcode()))
        GR.erase(&MI);
      if (MI.getOpcode() != SPIRV::ASSIGN_TYPE)
        continue;

      Register DstReg = MI.getOperand(0).getReg();
      LLT DstType = MRI.getType(DstReg);
      Register SrcReg = MI.getOperand(1).getReg();
      LLT SrcType = MRI.getType(SrcReg);
      if (DstType != SrcType)
        MRI.setType(DstReg, MRI.getType(SrcReg));

      const TargetRegisterClass *DstRC = MRI.getRegClassOrNull(DstReg);
      const TargetRegisterClass *SrcRC = MRI.getRegClassOrNull(SrcReg);
      if (DstRC != SrcRC && SrcRC)
        MRI.setRegClass(DstReg, SrcRC);
    }
  }
}

// Return true if the type represents a constant register
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef,
                       SmallPtrSet<SPIRVType *, 4> &Visited) {
  OpDef = passCopy(OpDef, MRI);

  if (Visited.contains(OpDef))
    return true;
  Visited.insert(OpDef);

  unsigned Opcode = OpDef->getOpcode();
  switch (Opcode) {
  case TargetOpcode::G_CONSTANT:
  case TargetOpcode::G_FCONSTANT:
    return true;
  case TargetOpcode::G_INTRINSIC:
  case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
  case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
    return cast<GIntrinsic>(*OpDef).getIntrinsicID() ==
           Intrinsic::spv_const_composite;
  case TargetOpcode::G_BUILD_VECTOR:
  case TargetOpcode::G_SPLAT_VECTOR: {
    for (unsigned i = OpDef->getNumExplicitDefs(); i < OpDef->getNumOperands();
         i++) {
      MachineInstr *OpNestedDef =
          OpDef->getOperand(i).isReg()
              ? MRI->getVRegDef(OpDef->getOperand(i).getReg())
              : nullptr;
      if (OpNestedDef && !isConstReg(MRI, OpNestedDef, Visited))
        return false;
    }
    return true;
  case SPIRV::OpConstantTrue:
  case SPIRV::OpConstantFalse:
  case SPIRV::OpConstantI:
  case SPIRV::OpConstantF:
  case SPIRV::OpConstantComposite:
  case SPIRV::OpConstantCompositeContinuedINTEL:
  case SPIRV::OpConstantSampler:
  case SPIRV::OpConstantNull:
  case SPIRV::OpUndef:
  case SPIRV::OpConstantFunctionPointerINTEL:
    return true;
  }
  }
  return false;
}

// Return true if the virtual register represents a constant
static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
  SmallPtrSet<SPIRVType *, 4> Visited;
  if (MachineInstr *OpDef = MRI->getVRegDef(OpReg))
    return isConstReg(MRI, OpDef, Visited);
  return false;
}

bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
  for (const auto &MO : MI.all_defs()) {
    Register Reg = MO.getReg();
    if (Reg.isPhysical() || !MRI.use_nodbg_empty(Reg))
      return false;
  }
  if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||
      MI.isLifetimeMarker())
    return false;
  if (MI.isPHI())
    return true;
  if (MI.mayStore() || MI.isCall() ||
      (MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||
      MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo())
    return false;
  return true;
}

bool SPIRVInstructionSelector::select(MachineInstr &I) {
  resetVRegsType(*I.getParent()->getParent());

  assert(I.getParent() && "Instruction should be in a basic block!");
  assert(I.getParent()->getParent() && "Instruction should be in a function!");

  Register Opcode = I.getOpcode();
  // If it's not a GMIR instruction, we've selected it already.
  if (!isPreISelGenericOpcode(Opcode)) {
    if (Opcode == SPIRV::ASSIGN_TYPE) { // These pseudos aren't needed any more.
      Register DstReg = I.getOperand(0).getReg();
      Register SrcReg = I.getOperand(1).getReg();
      auto *Def = MRI->getVRegDef(SrcReg);
      if (isTypeFoldingSupported(Def->getOpcode()) &&
          Def->getOpcode() != TargetOpcode::G_CONSTANT &&
          Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
        bool Res = selectImpl(I, *CoverageInfo);
        LLVM_DEBUG({
          if (!Res && Def->getOpcode() != TargetOpcode::G_CONSTANT) {
            dbgs() << "Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
            I.print(dbgs());
          }
        });
        assert(Res || Def->getOpcode() == TargetOpcode::G_CONSTANT);
        if (Res) {
          if (!isTriviallyDead(*Def, *MRI) && isDead(*Def, *MRI))
            DeadMIs.insert(Def);
          return Res;
        }
      }
      MRI->setRegClass(SrcReg, MRI->getRegClass(DstReg));
      MRI->replaceRegWith(SrcReg, DstReg);
      GR.invalidateMachineInstr(&I);
      I.removeFromParent();
      return true;
    } else if (I.getNumDefs() == 1) {
      // Make all vregs 64 bits (for SPIR-V IDs).
      MRI->setType(I.getOperand(0).getReg(), LLT::scalar(64));
    }
    return constrainSelectedInstRegOperands(I, TII, TRI, RBI);
  }

  if (DeadMIs.contains(&I)) {
    // if the instruction has been already made dead by folding it away
    // erase it
    LLVM_DEBUG(dbgs() << "Instruction is folded and dead.\n");
    salvageDebugInfo(*MRI, I);
    GR.invalidateMachineInstr(&I);
    I.eraseFromParent();
    return true;
  }

  if (I.getNumOperands() != I.getNumExplicitOperands()) {
    LLVM_DEBUG(errs() << "Generic instr has unexpected implicit operands\n");
    return false;
  }

  // Common code for getting return reg+type, and removing selected instr
  // from parent occurs here. Instr-specific selection happens in spvSelect().
  bool HasDefs = I.getNumDefs() > 0;
  Register ResVReg = HasDefs ? I.getOperand(0).getReg() : Register(0);
  SPIRVType *ResType = HasDefs ? GR.getSPIRVTypeForVReg(ResVReg) : nullptr;
  assert(!HasDefs || ResType || I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
         I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
  if (spvSelect(ResVReg, ResType, I)) {
    if (HasDefs) // Make all vregs 64 bits (for SPIR-V IDs).
      for (unsigned i = 0; i < I.getNumDefs(); ++i)
        MRI->setType(I.getOperand(i).getReg(), LLT::scalar(64));
    GR.invalidateMachineInstr(&I);
    I.removeFromParent();
    return true;
  }
  return false;
}

static bool mayApplyGenericSelection(unsigned Opcode) {
  switch (Opcode) {
  case TargetOpcode::G_CONSTANT:
  case TargetOpcode::G_FCONSTANT:
    return false;
  case TargetOpcode::G_SADDO:
  case TargetOpcode::G_SSUBO:
    return true;
  }
  return isTypeFoldingSupported(Opcode);
}

bool SPIRVInstructionSelector::BuildCOPY(Register DestReg, Register SrcReg,
                                         MachineInstr &I) const {
  const TargetRegisterClass *DstRC = MRI->getRegClassOrNull(DestReg);
  const TargetRegisterClass *SrcRC = MRI->getRegClassOrNull(SrcReg);
  if (DstRC != SrcRC && SrcRC)
    MRI->setRegClass(DestReg, SrcRC);
  return BuildMI(*I.getParent(), I, I.getDebugLoc(),
                 TII.get(TargetOpcode::COPY))
      .addDef(DestReg)
      .addUse(SrcReg)
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
                                         const SPIRVType *ResType,
                                         MachineInstr &I) const {
  const unsigned Opcode = I.getOpcode();
  if (mayApplyGenericSelection(Opcode))
    return selectImpl(I, *CoverageInfo);
  switch (Opcode) {
  case TargetOpcode::G_CONSTANT:
  case TargetOpcode::G_FCONSTANT:
    return selectConst(ResVReg, ResType, I);
  case TargetOpcode::G_GLOBAL_VALUE:
    return selectGlobalValue(ResVReg, I);
  case TargetOpcode::G_IMPLICIT_DEF:
    return selectOpUndef(ResVReg, ResType, I);
  case TargetOpcode::G_FREEZE:
    return selectFreeze(ResVReg, ResType, I);

  case TargetOpcode::G_INTRINSIC:
  case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
  case TargetOpcode::G_INTRINSIC_CONVERGENT:
  case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
    return selectIntrinsic(ResVReg, ResType, I);
  case TargetOpcode::G_BITREVERSE:
    return selectBitreverse(ResVReg, ResType, I);

  case TargetOpcode::G_BUILD_VECTOR:
    return selectBuildVector(ResVReg, ResType, I);
  case TargetOpcode::G_SPLAT_VECTOR:
    return selectSplatVector(ResVReg, ResType, I);

  case TargetOpcode::G_SHUFFLE_VECTOR: {
    MachineBasicBlock &BB = *I.getParent();
    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))
                   .addDef(ResVReg)
                   .addUse(GR.getSPIRVTypeID(ResType))
                   .addUse(I.getOperand(1).getReg())
                   .addUse(I.getOperand(2).getReg());
    for (auto V : I.getOperand(3).getShuffleMask())
      MIB.addImm(V);
    return MIB.constrainAllUses(TII, TRI, RBI);
  }
  case TargetOpcode::G_MEMMOVE:
  case TargetOpcode::G_MEMCPY:
  case TargetOpcode::G_MEMSET:
    return selectMemOperation(ResVReg, I);

  case TargetOpcode::G_ICMP:
    return selectICmp(ResVReg, ResType, I);
  case TargetOpcode::G_FCMP:
    return selectFCmp(ResVReg, ResType, I);

  case TargetOpcode::G_FRAME_INDEX:
    return selectFrameIndex(ResVReg, ResType, I);

  case TargetOpcode::G_LOAD:
    return selectLoad(ResVReg, ResType, I);
  case TargetOpcode::G_STORE:
    return selectStore(I);

  case TargetOpcode::G_BR:
    return selectBranch(I);
  case TargetOpcode::G_BRCOND:
    return selectBranchCond(I);

  case TargetOpcode::G_PHI:
    return selectPhi(ResVReg, ResType, I);

  case TargetOpcode::G_FPTOSI:
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToS);
  case TargetOpcode::G_FPTOUI:
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertFToU);

  case TargetOpcode::G_SITOFP:
    return selectIToF(ResVReg, ResType, I, true, SPIRV::OpConvertSToF);
  case TargetOpcode::G_UITOFP:
    return selectIToF(ResVReg, ResType, I, false, SPIRV::OpConvertUToF);

  case TargetOpcode::G_CTPOP:
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitCount);
  case TargetOpcode::G_SMIN:
    return selectExtInst(ResVReg, ResType, I, CL::s_min, GL::SMin);
  case TargetOpcode::G_UMIN:
    return selectExtInst(ResVReg, ResType, I, CL::u_min, GL::UMin);

  case TargetOpcode::G_SMAX:
    return selectExtInst(ResVReg, ResType, I, CL::s_max, GL::SMax);
  case TargetOpcode::G_UMAX:
    return selectExtInst(ResVReg, ResType, I, CL::u_max, GL::UMax);

  case TargetOpcode::G_SCMP:
    return selectSUCmp(ResVReg, ResType, I, true);
  case TargetOpcode::G_UCMP:
    return selectSUCmp(ResVReg, ResType, I, false);

  case TargetOpcode::G_STRICT_FMA:
  case TargetOpcode::G_FMA:
    return selectExtInst(ResVReg, ResType, I, CL::fma, GL::Fma);

  case TargetOpcode::G_STRICT_FLDEXP:
    return selectExtInst(ResVReg, ResType, I, CL::ldexp);

  case TargetOpcode::G_FPOW:
    return selectExtInst(ResVReg, ResType, I, CL::pow, GL::Pow);
  case TargetOpcode::G_FPOWI:
    return selectExtInst(ResVReg, ResType, I, CL::pown);

  case TargetOpcode::G_FEXP:
    return selectExtInst(ResVReg, ResType, I, CL::exp, GL::Exp);
  case TargetOpcode::G_FEXP2:
    return selectExtInst(ResVReg, ResType, I, CL::exp2, GL::Exp2);

  case TargetOpcode::G_FLOG:
    return selectExtInst(ResVReg, ResType, I, CL::log, GL::Log);
  case TargetOpcode::G_FLOG2:
    return selectExtInst(ResVReg, ResType, I, CL::log2, GL::Log2);
  case TargetOpcode::G_FLOG10:
    return selectLog10(ResVReg, ResType, I);

  case TargetOpcode::G_FABS:
    return selectExtInst(ResVReg, ResType, I, CL::fabs, GL::FAbs);
  case TargetOpcode::G_ABS:
    return selectExtInst(ResVReg, ResType, I, CL::s_abs, GL::SAbs);

  case TargetOpcode::G_FMINNUM:
  case TargetOpcode::G_FMINIMUM:
    return selectExtInst(ResVReg, ResType, I, CL::fmin, GL::NMin);
  case TargetOpcode::G_FMAXNUM:
  case TargetOpcode::G_FMAXIMUM:
    return selectExtInst(ResVReg, ResType, I, CL::fmax, GL::NMax);

  case TargetOpcode::G_FCOPYSIGN:
    return selectExtInst(ResVReg, ResType, I, CL::copysign);

  case TargetOpcode::G_FCEIL:
    return selectExtInst(ResVReg, ResType, I, CL::ceil, GL::Ceil);
  case TargetOpcode::G_FFLOOR:
    return selectExtInst(ResVReg, ResType, I, CL::floor, GL::Floor);

  case TargetOpcode::G_FCOS:
    return selectExtInst(ResVReg, ResType, I, CL::cos, GL::Cos);
  case TargetOpcode::G_FSIN:
    return selectExtInst(ResVReg, ResType, I, CL::sin, GL::Sin);
  case TargetOpcode::G_FTAN:
    return selectExtInst(ResVReg, ResType, I, CL::tan, GL::Tan);
  case TargetOpcode::G_FACOS:
    return selectExtInst(ResVReg, ResType, I, CL::acos, GL::Acos);
  case TargetOpcode::G_FASIN:
    return selectExtInst(ResVReg, ResType, I, CL::asin, GL::Asin);
  case TargetOpcode::G_FATAN:
    return selectExtInst(ResVReg, ResType, I, CL::atan, GL::Atan);
  case TargetOpcode::G_FATAN2:
    return selectExtInst(ResVReg, ResType, I, CL::atan2, GL::Atan2);
  case TargetOpcode::G_FCOSH:
    return selectExtInst(ResVReg, ResType, I, CL::cosh, GL::Cosh);
  case TargetOpcode::G_FSINH:
    return selectExtInst(ResVReg, ResType, I, CL::sinh, GL::Sinh);
  case TargetOpcode::G_FTANH:
    return selectExtInst(ResVReg, ResType, I, CL::tanh, GL::Tanh);

  case TargetOpcode::G_STRICT_FSQRT:
  case TargetOpcode::G_FSQRT:
    return selectExtInst(ResVReg, ResType, I, CL::sqrt, GL::Sqrt);

  case TargetOpcode::G_CTTZ:
  case TargetOpcode::G_CTTZ_ZERO_UNDEF:
    return selectExtInst(ResVReg, ResType, I, CL::ctz);
  case TargetOpcode::G_CTLZ:
  case TargetOpcode::G_CTLZ_ZERO_UNDEF:
    return selectExtInst(ResVReg, ResType, I, CL::clz);

  case TargetOpcode::G_INTRINSIC_ROUND:
    return selectExtInst(ResVReg, ResType, I, CL::round, GL::Round);
  case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
    return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);
  case TargetOpcode::G_INTRINSIC_TRUNC:
    return selectExtInst(ResVReg, ResType, I, CL::trunc, GL::Trunc);
  case TargetOpcode::G_FRINT:
  case TargetOpcode::G_FNEARBYINT:
    return selectExtInst(ResVReg, ResType, I, CL::rint, GL::RoundEven);

  case TargetOpcode::G_SMULH:
    return selectExtInst(ResVReg, ResType, I, CL::s_mul_hi);
  case TargetOpcode::G_UMULH:
    return selectExtInst(ResVReg, ResType, I, CL::u_mul_hi);

  case TargetOpcode::G_SADDSAT:
    return selectExtInst(ResVReg, ResType, I, CL::s_add_sat);
  case TargetOpcode::G_UADDSAT:
    return selectExtInst(ResVReg, ResType, I, CL::u_add_sat);
  case TargetOpcode::G_SSUBSAT:
    return selectExtInst(ResVReg, ResType, I, CL::s_sub_sat);
  case TargetOpcode::G_USUBSAT:
    return selectExtInst(ResVReg, ResType, I, CL::u_sub_sat);

  case TargetOpcode::G_UADDO:
    return selectOverflowArith(ResVReg, ResType, I,
                               ResType->getOpcode() == SPIRV::OpTypeVector
                                   ? SPIRV::OpIAddCarryV
                                   : SPIRV::OpIAddCarryS);
  case TargetOpcode::G_USUBO:
    return selectOverflowArith(ResVReg, ResType, I,
                               ResType->getOpcode() == SPIRV::OpTypeVector
                                   ? SPIRV::OpISubBorrowV
                                   : SPIRV::OpISubBorrowS);
  case TargetOpcode::G_UMULO:
    return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpUMulExtended);
  case TargetOpcode::G_SMULO:
    return selectOverflowArith(ResVReg, ResType, I, SPIRV::OpSMulExtended);

  case TargetOpcode::G_SEXT:
    return selectExt(ResVReg, ResType, I, true);
  case TargetOpcode::G_ANYEXT:
  case TargetOpcode::G_ZEXT:
    return selectExt(ResVReg, ResType, I, false);
  case TargetOpcode::G_TRUNC:
    return selectTrunc(ResVReg, ResType, I);
  case TargetOpcode::G_FPTRUNC:
  case TargetOpcode::G_FPEXT:
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpFConvert);

  case TargetOpcode::G_PTRTOINT:
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertPtrToU);
  case TargetOpcode::G_INTTOPTR:
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpConvertUToPtr);
  case TargetOpcode::G_BITCAST:
    return selectBitcast(ResVReg, ResType, I);
  case TargetOpcode::G_ADDRSPACE_CAST:
    return selectAddrSpaceCast(ResVReg, ResType, I);
  case TargetOpcode::G_PTR_ADD: {
    // Currently, we get G_PTR_ADD only applied to global variables.
    assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
    Register GV = I.getOperand(1).getReg();
    MachineRegisterInfo::def_instr_iterator II = MRI->def_instr_begin(GV);
    (void)II;
    assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
            (*II).getOpcode() == TargetOpcode::COPY ||
            (*II).getOpcode() == SPIRV::OpVariable) &&
           getImm(I.getOperand(2), MRI));
    // It may be the initialization of a global variable.
    bool IsGVInit = false;
    for (MachineRegisterInfo::use_instr_iterator
             UseIt = MRI->use_instr_begin(I.getOperand(0).getReg()),
             UseEnd = MRI->use_instr_end();
         UseIt != UseEnd; UseIt = std::next(UseIt)) {
      if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
          (*UseIt).getOpcode() == SPIRV::OpVariable) {
        IsGVInit = true;
        break;
      }
    }
    MachineBasicBlock &BB = *I.getParent();
    if (!IsGVInit) {
      SPIRVType *GVType = GR.getSPIRVTypeForVReg(GV);
      SPIRVType *GVPointeeType = GR.getPointeeType(GVType);
      SPIRVType *ResPointeeType = GR.getPointeeType(ResType);
      if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
        // Build a new virtual register that is associated with the required
        // data type.
        Register NewVReg = MRI->createGenericVirtualRegister(MRI->getType(GV));
        MRI->setRegClass(NewVReg, MRI->getRegClass(GV));
        //  Having a correctly typed base we are ready to build the actually
        //  required GEP. It may not be a constant though, because all Operands
        //  of OpSpecConstantOp is to originate from other const instructions,
        //  and only the AccessChain named opcodes accept a global OpVariable
        //  instruction. We can't use an AccessChain opcode because of the type
        //  mismatch between result and base types.
        if (!GR.isBitcastCompatible(ResType, GVType))
          report_fatal_error(
              "incompatible result and operand types in a bitcast");
        Register ResTypeReg = GR.getSPIRVTypeID(ResType);
        MachineInstrBuilder MIB =
            BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitcast))
                .addDef(NewVReg)
                .addUse(ResTypeReg)
                .addUse(GV);
        return MIB.constrainAllUses(TII, TRI, RBI) &&
               BuildMI(BB, I, I.getDebugLoc(),
                       TII.get(STI.isLogicalSPIRV()
                                   ? SPIRV::OpInBoundsAccessChain
                                   : SPIRV::OpInBoundsPtrAccessChain))
                   .addDef(ResVReg)
                   .addUse(ResTypeReg)
                   .addUse(NewVReg)
                   .addUse(I.getOperand(2).getReg())
                   .constrainAllUses(TII, TRI, RBI);
      } else {
        return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
            .addDef(ResVReg)
            .addUse(GR.getSPIRVTypeID(ResType))
            .addImm(
                static_cast<uint32_t>(SPIRV::Opcode::InBoundsPtrAccessChain))
            .addUse(GV)
            .addUse(I.getOperand(2).getReg())
            .constrainAllUses(TII, TRI, RBI);
      }
    }
    // It's possible to translate G_PTR_ADD to OpSpecConstantOp: either to
    // initialize a global variable with a constant expression (e.g., the test
    // case opencl/basic/progvar_prog_scope_init.ll), or for another use case
    Register Idx = buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
                   .addDef(ResVReg)
                   .addUse(GR.getSPIRVTypeID(ResType))
                   .addImm(static_cast<uint32_t>(
                       SPIRV::Opcode::InBoundsPtrAccessChain))
                   .addUse(GV)
                   .addUse(Idx)
                   .addUse(I.getOperand(2).getReg());
    return MIB.constrainAllUses(TII, TRI, RBI);
  }

  case TargetOpcode::G_ATOMICRMW_OR:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicOr);
  case TargetOpcode::G_ATOMICRMW_ADD:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicIAdd);
  case TargetOpcode::G_ATOMICRMW_AND:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicAnd);
  case TargetOpcode::G_ATOMICRMW_MAX:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMax);
  case TargetOpcode::G_ATOMICRMW_MIN:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicSMin);
  case TargetOpcode::G_ATOMICRMW_SUB:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicISub);
  case TargetOpcode::G_ATOMICRMW_XOR:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicXor);
  case TargetOpcode::G_ATOMICRMW_UMAX:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMax);
  case TargetOpcode::G_ATOMICRMW_UMIN:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicUMin);
  case TargetOpcode::G_ATOMICRMW_XCHG:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicExchange);
  case TargetOpcode::G_ATOMIC_CMPXCHG:
    return selectAtomicCmpXchg(ResVReg, ResType, I);

  case TargetOpcode::G_ATOMICRMW_FADD:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT);
  case TargetOpcode::G_ATOMICRMW_FSUB:
    // Translate G_ATOMICRMW_FSUB to OpAtomicFAddEXT with negative value operand
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFAddEXT,
                           SPIRV::OpFNegate);
  case TargetOpcode::G_ATOMICRMW_FMIN:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMinEXT);
  case TargetOpcode::G_ATOMICRMW_FMAX:
    return selectAtomicRMW(ResVReg, ResType, I, SPIRV::OpAtomicFMaxEXT);

  case TargetOpcode::G_FENCE:
    return selectFence(I);

  case TargetOpcode::G_STACKSAVE:
    return selectStackSave(ResVReg, ResType, I);
  case TargetOpcode::G_STACKRESTORE:
    return selectStackRestore(I);

  case TargetOpcode::G_UNMERGE_VALUES:
    return selectUnmergeValues(I);

  // Discard gen opcodes for intrinsics which we do not expect to actually
  // represent code after lowering or intrinsics which are not implemented but
  // should not crash when found in a customer's LLVM IR input.
  case TargetOpcode::G_TRAP:
  case TargetOpcode::G_DEBUGTRAP:
  case TargetOpcode::G_UBSANTRAP:
  case TargetOpcode::DBG_LABEL:
    return true;

  default:
    return false;
  }
}

bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I,
                                             GL::GLSLExtInst GLInst) const {
  if (!STI.canUseExtInstSet(
          SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
    std::string DiagMsg;
    raw_string_ostream OS(DiagMsg);
    I.print(OS, true, false, false, false);
    DiagMsg += " is only supported with the GLSL extended instruction set.\n";
    report_fatal_error(DiagMsg.c_str(), false);
  }
  return selectExtInst(ResVReg, ResType, I,
                       {{SPIRV::InstructionSet::GLSL_std_450, GLInst}});
}

bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I,
                                             CL::OpenCLExtInst CLInst) const {
  return selectExtInst(ResVReg, ResType, I,
                       {{SPIRV::InstructionSet::OpenCL_std, CLInst}});
}

bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I,
                                             CL::OpenCLExtInst CLInst,
                                             GL::GLSLExtInst GLInst) const {
  ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
                          {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
  return selectExtInst(ResVReg, ResType, I, ExtInsts);
}

bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I,
                                             const ExtInstList &Insts) const {

  for (const auto &Ex : Insts) {
    SPIRV::InstructionSet::InstructionSet Set = Ex.first;
    uint32_t Opcode = Ex.second;
    if (STI.canUseExtInstSet(Set)) {
      MachineBasicBlock &BB = *I.getParent();
      auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
                     .addDef(ResVReg)
                     .addUse(GR.getSPIRVTypeID(ResType))
                     .addImm(static_cast<uint32_t>(Set))
                     .addImm(Opcode);
      const unsigned NumOps = I.getNumOperands();
      unsigned Index = 1;
      if (Index < NumOps &&
          I.getOperand(Index).getType() ==
              MachineOperand::MachineOperandType::MO_IntrinsicID)
        Index = 2;
      for (; Index < NumOps; ++Index)
        MIB.add(I.getOperand(Index));
      return MIB.constrainAllUses(TII, TRI, RBI);
    }
  }
  return false;
}

bool SPIRVInstructionSelector::selectOpWithSrcs(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I,
                                                std::vector<Register> Srcs,
                                                unsigned Opcode) const {
  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType));
  for (Register SReg : Srcs) {
    MIB.addUse(SReg);
  }
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I,
                                          unsigned Opcode) const {
  if (STI.isPhysicalSPIRV() && I.getOperand(1).isReg()) {
    Register SrcReg = I.getOperand(1).getReg();
    bool IsGV = false;
    for (MachineRegisterInfo::def_instr_iterator DefIt =
             MRI->def_instr_begin(SrcReg);
         DefIt != MRI->def_instr_end(); DefIt = std::next(DefIt)) {
      if ((*DefIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
          (*DefIt).getOpcode() == SPIRV::OpVariable) {
        IsGV = true;
        break;
      }
    }
    if (IsGV) {
      uint32_t SpecOpcode = 0;
      switch (Opcode) {
      case SPIRV::OpConvertPtrToU:
        SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertPtrToU);
        break;
      case SPIRV::OpConvertUToPtr:
        SpecOpcode = static_cast<uint32_t>(SPIRV::Opcode::ConvertUToPtr);
        break;
      }
      if (SpecOpcode)
        return BuildMI(*I.getParent(), I, I.getDebugLoc(),
                       TII.get(SPIRV::OpSpecConstantOp))
            .addDef(ResVReg)
            .addUse(GR.getSPIRVTypeID(ResType))
            .addImm(SpecOpcode)
            .addUse(SrcReg)
            .constrainAllUses(TII, TRI, RBI);
    }
  }
  return selectOpWithSrcs(ResVReg, ResType, I, {I.getOperand(1).getReg()},
                          Opcode);
}

bool SPIRVInstructionSelector::selectBitcast(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I) const {
  Register OpReg = I.getOperand(1).getReg();
  SPIRVType *OpType = OpReg.isValid() ? GR.getSPIRVTypeForVReg(OpReg) : nullptr;
  if (!GR.isBitcastCompatible(ResType, OpType))
    report_fatal_error("incompatible result and operand types in a bitcast");
  return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitcast);
}

static void addMemoryOperands(MachineMemOperand *MemOp,
                              MachineInstrBuilder &MIB,
                              MachineIRBuilder &MIRBuilder,
                              SPIRVGlobalRegistry &GR) {
  uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
  if (MemOp->isVolatile())
    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
  if (MemOp->isNonTemporal())
    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
  if (MemOp->getAlign().value())
    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);

  [[maybe_unused]] MachineInstr *AliasList = nullptr;
  [[maybe_unused]] MachineInstr *NoAliasList = nullptr;
  const SPIRVSubtarget *ST =
      static_cast<const SPIRVSubtarget *>(&MIRBuilder.getMF().getSubtarget());
  if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
    if (auto *MD = MemOp->getAAInfo().Scope) {
      AliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
      if (AliasList)
        SpvMemOp |=
            static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
    }
    if (auto *MD = MemOp->getAAInfo().NoAlias) {
      NoAliasList = GR.getOrAddMemAliasingINTELInst(MIRBuilder, MD);
      if (NoAliasList)
        SpvMemOp |=
            static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
    }
  }

  if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
    MIB.addImm(SpvMemOp);
    if (SpvMemOp & static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
      MIB.addImm(MemOp->getAlign().value());
    if (AliasList)
      MIB.addUse(AliasList->getOperand(0).getReg());
    if (NoAliasList)
      MIB.addUse(NoAliasList->getOperand(0).getReg());
  }
}

static void addMemoryOperands(uint64_t Flags, MachineInstrBuilder &MIB) {
  uint32_t SpvMemOp = static_cast<uint32_t>(SPIRV::MemoryOperand::None);
  if (Flags & MachineMemOperand::Flags::MOVolatile)
    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
  if (Flags & MachineMemOperand::Flags::MONonTemporal)
    SpvMemOp |= static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);

  if (SpvMemOp != static_cast<uint32_t>(SPIRV::MemoryOperand::None))
    MIB.addImm(SpvMemOp);
}

bool SPIRVInstructionSelector::selectLoad(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I) const {
  unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
  Register Ptr = I.getOperand(1 + OpOffset).getReg();

  auto *PtrDef = getVRegDef(*MRI, Ptr);
  auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
  if (IntPtrDef &&
      IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
    Register HandleReg = IntPtrDef->getOperand(2).getReg();
    SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
    if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
      Register NewHandleReg =
          MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
      auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
      if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
        return false;
      }

      Register IdxReg = IntPtrDef->getOperand(3).getReg();
      return generateImageRead(ResVReg, ResType, NewHandleReg, IdxReg,
                               I.getDebugLoc(), I);
    }
  }

  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 .addUse(Ptr);
  if (!I.getNumMemOperands()) {
    assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
           I.getOpcode() ==
               TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
    addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
  } else {
    MachineIRBuilder MIRBuilder(I);
    addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
  }
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectStore(MachineInstr &I) const {
  unsigned OpOffset = isa<GIntrinsic>(I) ? 1 : 0;
  Register StoreVal = I.getOperand(0 + OpOffset).getReg();
  Register Ptr = I.getOperand(1 + OpOffset).getReg();

  auto *PtrDef = getVRegDef(*MRI, Ptr);
  auto *IntPtrDef = dyn_cast<GIntrinsic>(PtrDef);
  if (IntPtrDef &&
      IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
    Register HandleReg = IntPtrDef->getOperand(2).getReg();
    Register NewHandleReg =
        MRI->createVirtualRegister(MRI->getRegClass(HandleReg));
    auto *HandleDef = cast<GIntrinsic>(getVRegDef(*MRI, HandleReg));
    SPIRVType *HandleType = GR.getSPIRVTypeForVReg(HandleReg);
    if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef, I)) {
      return false;
    }

    Register IdxReg = IntPtrDef->getOperand(3).getReg();
    if (HandleType->getOpcode() == SPIRV::OpTypeImage) {
      auto BMI = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                         TII.get(SPIRV::OpImageWrite))
                     .addUse(NewHandleReg)
                     .addUse(IdxReg)
                     .addUse(StoreVal);

      const llvm::Type *LLVMHandleType = GR.getTypeForSPIRVType(HandleType);
      if (sampledTypeIsSignedInteger(LLVMHandleType))
        BMI.addImm(0x1000); // SignExtend

      return BMI.constrainAllUses(TII, TRI, RBI);
    }
  }

  MachineBasicBlock &BB = *I.getParent();
  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpStore))
                 .addUse(Ptr)
                 .addUse(StoreVal);
  if (!I.getNumMemOperands()) {
    assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
           I.getOpcode() ==
               TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
    addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
  } else {
    MachineIRBuilder MIRBuilder(I);
    addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
  }
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,
                                               const SPIRVType *ResType,
                                               MachineInstr &I) const {
  if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
    report_fatal_error(
        "llvm.stacksave intrinsic: this instruction requires the following "
        "SPIR-V extension: SPV_INTEL_variable_length_array",
        false);
  MachineBasicBlock &BB = *I.getParent();
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSaveMemoryINTEL))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {
  if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
    report_fatal_error(
        "llvm.stackrestore intrinsic: this instruction requires the following "
        "SPIR-V extension: SPV_INTEL_variable_length_array",
        false);
  if (!I.getOperand(0).isReg())
    return false;
  MachineBasicBlock &BB = *I.getParent();
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpRestoreMemoryINTEL))
      .addUse(I.getOperand(0).getReg())
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,
                                                  MachineInstr &I) const {
  MachineBasicBlock &BB = *I.getParent();
  Register SrcReg = I.getOperand(1).getReg();
  bool Result = true;
  if (I.getOpcode() == TargetOpcode::G_MEMSET) {
    MachineIRBuilder MIRBuilder(I);
    assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
    unsigned Val = getIConstVal(I.getOperand(1).getReg(), MRI);
    unsigned Num = getIConstVal(I.getOperand(2).getReg(), MRI);
    Type *ValTy = Type::getInt8Ty(I.getMF()->getFunction().getContext());
    Type *ArrTy = ArrayType::get(ValTy, Num);
    SPIRVType *VarTy = GR.getOrCreateSPIRVPointerType(
        ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);

    SPIRVType *SpvArrTy = GR.getOrCreateSPIRVType(
        ArrTy, MIRBuilder, SPIRV::AccessQualifier::None, false);
    Register Const = GR.getOrCreateConstIntArray(Val, Num, I, SpvArrTy, TII);
    // TODO: check if we have such GV, add init, use buildGlobalVariable.
    Function &CurFunction = GR.CurMF->getFunction();
    Type *LLVMArrTy =
        ArrayType::get(IntegerType::get(CurFunction.getContext(), 8), Num);
    // Module takes ownership of the global var.
    GlobalVariable *GV = new GlobalVariable(*CurFunction.getParent(), LLVMArrTy,
                                            true, GlobalValue::InternalLinkage,
                                            Constant::getNullValue(LLVMArrTy));
    Register VarReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
    auto MIBVar =
        BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
            .addDef(VarReg)
            .addUse(GR.getSPIRVTypeID(VarTy))
            .addImm(SPIRV::StorageClass::UniformConstant)
            .addUse(Const);
    Result &= MIBVar.constrainAllUses(TII, TRI, RBI);

    GR.add(GV, MIBVar);
    GR.addGlobalObject(GV, GR.CurMF, VarReg);

    buildOpDecorate(VarReg, I, TII, SPIRV::Decoration::Constant, {});
    SPIRVType *SourceTy = GR.getOrCreateSPIRVPointerType(
        ValTy, I, SPIRV::StorageClass::UniformConstant);
    SrcReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
    selectOpWithSrcs(SrcReg, SourceTy, I, {VarReg}, SPIRV::OpBitcast);
  }
  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemorySized))
                 .addUse(I.getOperand(0).getReg())
                 .addUse(SrcReg)
                 .addUse(I.getOperand(2).getReg());
  if (I.getNumMemOperands()) {
    MachineIRBuilder MIRBuilder(I);
    addMemoryOperands(*I.memoperands_begin(), MIB, MIRBuilder, GR);
  }
  Result &= MIB.constrainAllUses(TII, TRI, RBI);
  if (ResVReg.isValid() && ResVReg != MIB->getOperand(0).getReg())
    Result &= BuildCOPY(ResVReg, MIB->getOperand(0).getReg(), I);
  return Result;
}

bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,
                                               const SPIRVType *ResType,
                                               MachineInstr &I,
                                               unsigned NewOpcode,
                                               unsigned NegateOpcode) const {
  bool Result = true;
  assert(I.hasOneMemOperand());
  const MachineMemOperand *MemOp = *I.memoperands_begin();
  uint32_t Scope = static_cast<uint32_t>(getMemScope(
      GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
  auto ScopeConstant = buildI32Constant(Scope, I);
  Register ScopeReg = ScopeConstant.first;
  Result &= ScopeConstant.second;

  Register Ptr = I.getOperand(1).getReg();
  // TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll
  // auto ScSem =
  // getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));
  AtomicOrdering AO = MemOp->getSuccessOrdering();
  uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
  auto MemSemConstant = buildI32Constant(MemSem /*| ScSem*/, I);
  Register MemSemReg = MemSemConstant.first;
  Result &= MemSemConstant.second;

  Register ValueReg = I.getOperand(2).getReg();
  if (NegateOpcode != 0) {
    // Translation with negative value operand is requested
    Register TmpReg = createVirtualRegister(ResType, &GR, MRI, MRI->getMF());
    Result &= selectOpWithSrcs(TmpReg, ResType, I, {ValueReg}, NegateOpcode);
    ValueReg = TmpReg;
  }

  return Result &&
         BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(NewOpcode))
             .addDef(ResVReg)
             .addUse(GR.getSPIRVTypeID(ResType))
             .addUse(Ptr)
             .addUse(ScopeReg)
             .addUse(MemSemReg)
             .addUse(ValueReg)
             .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &I) const {
  unsigned ArgI = I.getNumOperands() - 1;
  Register SrcReg =
      I.getOperand(ArgI).isReg() ? I.getOperand(ArgI).getReg() : Register(0);
  SPIRVType *DefType =
      SrcReg.isValid() ? GR.getSPIRVTypeForVReg(SrcReg) : nullptr;
  if (!DefType || DefType->getOpcode() != SPIRV::OpTypeVector)
    report_fatal_error(
        "cannot select G_UNMERGE_VALUES with a non-vector argument");

  SPIRVType *ScalarType =
      GR.getSPIRVTypeForVReg(DefType->getOperand(1).getReg());
  MachineBasicBlock &BB = *I.getParent();
  bool Res = false;
  for (unsigned i = 0; i < I.getNumDefs(); ++i) {
    Register ResVReg = I.getOperand(i).getReg();
    SPIRVType *ResType = GR.getSPIRVTypeForVReg(ResVReg);
    if (!ResType) {
      // There was no "assign type" actions, let's fix this now
      ResType = ScalarType;
      MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
      MRI->setType(ResVReg, LLT::scalar(GR.getScalarOrVectorBitWidth(ResType)));
      GR.assignSPIRVTypeToVReg(ResType, ResVReg, *GR.CurMF);
    }
    auto MIB =
        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
            .addDef(ResVReg)
            .addUse(GR.getSPIRVTypeID(ResType))
            .addUse(SrcReg)
            .addImm(static_cast<int64_t>(i));
    Res |= MIB.constrainAllUses(TII, TRI, RBI);
  }
  return Res;
}

bool SPIRVInstructionSelector::selectFence(MachineInstr &I) const {
  AtomicOrdering AO = AtomicOrdering(I.getOperand(0).getImm());
  uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
  auto MemSemConstant = buildI32Constant(MemSem, I);
  Register MemSemReg = MemSemConstant.first;
  bool Result = MemSemConstant.second;
  SyncScope::ID Ord = SyncScope::ID(I.getOperand(1).getImm());
  uint32_t Scope = static_cast<uint32_t>(
      getMemScope(GR.CurMF->getFunction().getContext(), Ord));
  auto ScopeConstant = buildI32Constant(Scope, I);
  Register ScopeReg = ScopeConstant.first;
  Result &= ScopeConstant.second;
  MachineBasicBlock &BB = *I.getParent();
  return Result &&
         BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpMemoryBarrier))
             .addUse(ScopeReg)
             .addUse(MemSemReg)
             .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectOverflowArith(Register ResVReg,
                                                   const SPIRVType *ResType,
                                                   MachineInstr &I,
                                                   unsigned Opcode) const {
  Type *ResTy = nullptr;
  StringRef ResName;
  if (!GR.findValueAttrs(&I, ResTy, ResName))
    report_fatal_error(
        "Not enough info to select the arithmetic with overflow instruction");
  if (!ResTy || !ResTy->isStructTy())
    report_fatal_error("Expect struct type result for the arithmetic "
                       "with overflow instruction");
  // "Result Type must be from OpTypeStruct. The struct must have two members,
  // and the two members must be the same type."
  Type *ResElemTy = cast<StructType>(ResTy)->getElementType(0);
  ResTy = StructType::get(ResElemTy, ResElemTy);
  // Build SPIR-V types and constant(s) if needed.
  MachineIRBuilder MIRBuilder(I);
  SPIRVType *StructType = GR.getOrCreateSPIRVType(
      ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);
  assert(I.getNumDefs() > 1 && "Not enought operands");
  SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
  unsigned N = GR.getScalarOrVectorComponentCount(ResType);
  if (N > 1)
    BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
  Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
  Register ZeroReg = buildZerosVal(ResType, I);
  // A new virtual register to store the result struct.
  Register StructVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
  MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);
  // Build the result name if needed.
  if (ResName.size() > 0)
    buildOpName(StructVReg, ResName, MIRBuilder);
  // Build the arithmetic with overflow instruction.
  MachineBasicBlock &BB = *I.getParent();
  auto MIB =
      BuildMI(BB, MIRBuilder.getInsertPt(), I.getDebugLoc(), TII.get(Opcode))
          .addDef(StructVReg)
          .addUse(GR.getSPIRVTypeID(StructType));
  for (unsigned i = I.getNumDefs(); i < I.getNumOperands(); ++i)
    MIB.addUse(I.getOperand(i).getReg());
  bool Result = MIB.constrainAllUses(TII, TRI, RBI);
  // Build instructions to extract fields of the instruction's result.
  // A new virtual register to store the higher part of the result struct.
  Register HigherVReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
  MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);
  for (unsigned i = 0; i < I.getNumDefs(); ++i) {
    auto MIB =
        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
            .addDef(i == 1 ? HigherVReg : I.getOperand(i).getReg())
            .addUse(GR.getSPIRVTypeID(ResType))
            .addUse(StructVReg)
            .addImm(i);
    Result &= MIB.constrainAllUses(TII, TRI, RBI);
  }
  // Build boolean value from the higher part.
  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
                       .addDef(I.getOperand(1).getReg())
                       .addUse(BoolTypeReg)
                       .addUse(HigherVReg)
                       .addUse(ZeroReg)
                       .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectAtomicCmpXchg(Register ResVReg,
                                                   const SPIRVType *ResType,
                                                   MachineInstr &I) const {
  bool Result = true;
  Register ScopeReg;
  Register MemSemEqReg;
  Register MemSemNeqReg;
  Register Ptr = I.getOperand(2).getReg();
  if (!isa<GIntrinsic>(I)) {
    assert(I.hasOneMemOperand());
    const MachineMemOperand *MemOp = *I.memoperands_begin();
    unsigned Scope = static_cast<uint32_t>(getMemScope(
        GR.CurMF->getFunction().getContext(), MemOp->getSyncScopeID()));
    auto ScopeConstant = buildI32Constant(Scope, I);
    ScopeReg = ScopeConstant.first;
    Result &= ScopeConstant.second;

    unsigned ScSem = static_cast<uint32_t>(
        getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr)));
    AtomicOrdering AO = MemOp->getSuccessOrdering();
    unsigned MemSemEq = static_cast<uint32_t>(getMemSemantics(AO)) | ScSem;
    auto MemSemEqConstant = buildI32Constant(MemSemEq, I);
    MemSemEqReg = MemSemEqConstant.first;
    Result &= MemSemEqConstant.second;
    AtomicOrdering FO = MemOp->getFailureOrdering();
    unsigned MemSemNeq = static_cast<uint32_t>(getMemSemantics(FO)) | ScSem;
    if (MemSemEq == MemSemNeq)
      MemSemNeqReg = MemSemEqReg;
    else {
      auto MemSemNeqConstant = buildI32Constant(MemSemEq, I);
      MemSemNeqReg = MemSemNeqConstant.first;
      Result &= MemSemNeqConstant.second;
    }
  } else {
    ScopeReg = I.getOperand(5).getReg();
    MemSemEqReg = I.getOperand(6).getReg();
    MemSemNeqReg = I.getOperand(7).getReg();
  }

  Register Cmp = I.getOperand(3).getReg();
  Register Val = I.getOperand(4).getReg();
  SPIRVType *SpvValTy = GR.getSPIRVTypeForVReg(Val);
  Register ACmpRes = createVirtualRegister(SpvValTy, &GR, MRI, *I.getMF());
  const DebugLoc &DL = I.getDebugLoc();
  Result &=
      BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpAtomicCompareExchange))
          .addDef(ACmpRes)
          .addUse(GR.getSPIRVTypeID(SpvValTy))
          .addUse(Ptr)
          .addUse(ScopeReg)
          .addUse(MemSemEqReg)
          .addUse(MemSemNeqReg)
          .addUse(Val)
          .addUse(Cmp)
          .constrainAllUses(TII, TRI, RBI);
  SPIRVType *BoolTy = GR.getOrCreateSPIRVBoolType(I, TII);
  Register CmpSuccReg = createVirtualRegister(BoolTy, &GR, MRI, *I.getMF());
  Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpIEqual))
                .addDef(CmpSuccReg)
                .addUse(GR.getSPIRVTypeID(BoolTy))
                .addUse(ACmpRes)
                .addUse(Cmp)
                .constrainAllUses(TII, TRI, RBI);
  Register TmpReg = createVirtualRegister(ResType, &GR, MRI, *I.getMF());
  Result &= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
                .addDef(TmpReg)
                .addUse(GR.getSPIRVTypeID(ResType))
                .addUse(ACmpRes)
                .addUse(GR.getOrCreateUndef(I, ResType, TII))
                .addImm(0)
                .constrainAllUses(TII, TRI, RBI);
  return Result &&
         BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
             .addDef(ResVReg)
             .addUse(GR.getSPIRVTypeID(ResType))
             .addUse(CmpSuccReg)
             .addUse(TmpReg)
             .addImm(1)
             .constrainAllUses(TII, TRI, RBI);
}

static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC) {
  switch (SC) {
  case SPIRV::StorageClass::DeviceOnlyINTEL:
  case SPIRV::StorageClass::HostOnlyINTEL:
    return true;
  default:
    return false;
  }
}

// Returns true ResVReg is referred only from global vars and OpName's.
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg) {
  bool IsGRef = false;
  bool IsAllowedRefs =
      llvm::all_of(MRI->use_instructions(ResVReg), [&IsGRef](auto const &It) {
        unsigned Opcode = It.getOpcode();
        if (Opcode == SPIRV::OpConstantComposite ||
            Opcode == SPIRV::OpVariable ||
            isSpvIntrinsic(It, Intrinsic::spv_init_global))
          return IsGRef = true;
        return Opcode == SPIRV::OpName;
      });
  return IsAllowedRefs && IsGRef;
}

Register SPIRVInstructionSelector::getUcharPtrTypeReg(
    MachineInstr &I, SPIRV::StorageClass::StorageClass SC) const {
  return GR.getSPIRVTypeID(GR.getOrCreateSPIRVPointerType(
      Type::getInt8Ty(I.getMF()->getFunction().getContext()), I, SC));
}

MachineInstrBuilder
SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &I, Register Dest,
                                              Register Src, Register DestType,
                                              uint32_t Opcode) const {
  return BuildMI(*I.getParent(), I, I.getDebugLoc(),
                 TII.get(SPIRV::OpSpecConstantOp))
      .addDef(Dest)
      .addUse(DestType)
      .addImm(Opcode)
      .addUse(Src);
}

MachineInstrBuilder
SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
                                               SPIRVType *SrcPtrTy) const {
  SPIRVType *GenericPtrTy =
      GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
  Register Tmp = MRI->createVirtualRegister(&SPIRV::pIDRegClass);
  MRI->setType(Tmp, LLT::pointer(storageClassToAddressSpace(
                                     SPIRV::StorageClass::Generic),
                                 GR.getPointerSize()));
  MachineFunction *MF = I.getParent()->getParent();
  GR.assignSPIRVTypeToVReg(GenericPtrTy, Tmp, *MF);
  MachineInstrBuilder MIB = buildSpecConstantOp(
      I, Tmp, SrcPtr, GR.getSPIRVTypeID(GenericPtrTy),
      static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric));
  GR.add(MIB.getInstr(), MIB);
  return MIB;
}

// In SPIR-V address space casting can only happen to and from the Generic
// storage class. We can also only cast Workgroup, CrossWorkgroup, or Function
// pointers to and from Generic pointers. As such, we can convert e.g. from
// Workgroup to Function by going via a Generic pointer as an intermediary. All
// other combinations can only be done by a bitcast, and are probably not safe.
bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
                                                   const SPIRVType *ResType,
                                                   MachineInstr &I) const {
  MachineBasicBlock &BB = *I.getParent();
  const DebugLoc &DL = I.getDebugLoc();

  Register SrcPtr = I.getOperand(1).getReg();
  SPIRVType *SrcPtrTy = GR.getSPIRVTypeForVReg(SrcPtr);

  // don't generate a cast for a null that may be represented by OpTypeInt
  if (SrcPtrTy->getOpcode() != SPIRV::OpTypePointer ||
      ResType->getOpcode() != SPIRV::OpTypePointer)
    return BuildCOPY(ResVReg, SrcPtr, I);

  SPIRV::StorageClass::StorageClass SrcSC = GR.getPointerStorageClass(SrcPtrTy);
  SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResType);

  if (isASCastInGVar(MRI, ResVReg)) {
    // AddrSpaceCast uses within OpVariable and OpConstantComposite instructions
    // are expressed by OpSpecConstantOp with an Opcode.
    // TODO: maybe insert a check whether the Kernel capability was declared and
    // so PtrCastToGeneric/GenericCastToPtr are available.
    unsigned SpecOpcode =
        DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC)
            ? static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric)
            : (SrcSC == SPIRV::StorageClass::Generic &&
                       isGenericCastablePtr(DstSC)
                   ? static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr)
                   : 0);
    // TODO: OpConstantComposite expects i8*, so we are forced to forget a
    // correct value of ResType and use general i8* instead. Maybe this should
    // be addressed in the emit-intrinsic step to infer a correct
    // OpConstantComposite type.
    if (SpecOpcode) {
      return buildSpecConstantOp(I, ResVReg, SrcPtr,
                                 getUcharPtrTypeReg(I, DstSC), SpecOpcode)
          .constrainAllUses(TII, TRI, RBI);
    } else if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
      MachineInstrBuilder MIB = buildConstGenericPtr(I, SrcPtr, SrcPtrTy);
      return MIB.constrainAllUses(TII, TRI, RBI) &&
             buildSpecConstantOp(
                 I, ResVReg, MIB->getOperand(0).getReg(),
                 getUcharPtrTypeReg(I, DstSC),
                 static_cast<uint32_t>(SPIRV::Opcode::GenericCastToPtr))
                 .constrainAllUses(TII, TRI, RBI);
    }
  }

  // don't generate a cast between identical storage classes
  if (SrcSC == DstSC)
    return BuildCOPY(ResVReg, SrcPtr, I);

  if ((SrcSC == SPIRV::StorageClass::Function &&
       DstSC == SPIRV::StorageClass::Private) ||
      (DstSC == SPIRV::StorageClass::Function &&
       SrcSC == SPIRV::StorageClass::Private))
    return BuildCOPY(ResVReg, SrcPtr, I);

  // Casting from an eligible pointer to Generic.
  if (DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC))
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
  // Casting from Generic to an eligible pointer.
  if (SrcSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(DstSC))
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
  // Casting between 2 eligible pointers using Generic as an intermediary.
  if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
    SPIRVType *GenericPtrTy =
        GR.changePointerStorageClass(SrcPtrTy, SPIRV::StorageClass::Generic, I);
    Register Tmp = createVirtualRegister(GenericPtrTy, &GR, MRI, MRI->getMF());
    bool Result = BuildMI(BB, I, DL, TII.get(SPIRV::OpPtrCastToGeneric))
                      .addDef(Tmp)
                      .addUse(GR.getSPIRVTypeID(GenericPtrTy))
                      .addUse(SrcPtr)
                      .constrainAllUses(TII, TRI, RBI);
    return Result && BuildMI(BB, I, DL, TII.get(SPIRV::OpGenericCastToPtr))
                         .addDef(ResVReg)
                         .addUse(GR.getSPIRVTypeID(ResType))
                         .addUse(Tmp)
                         .constrainAllUses(TII, TRI, RBI);
  }

  // Check if instructions from the SPV_INTEL_usm_storage_classes extension may
  // be applied
  if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::CrossWorkgroup)
    return selectUnOp(ResVReg, ResType, I,
                      SPIRV::OpPtrCastToCrossWorkgroupINTEL);
  if (SrcSC == SPIRV::StorageClass::CrossWorkgroup && isUSMStorageClass(DstSC))
    return selectUnOp(ResVReg, ResType, I,
                      SPIRV::OpCrossWorkgroupCastToPtrINTEL);
  if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::Generic)
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
  if (SrcSC == SPIRV::StorageClass::Generic && isUSMStorageClass(DstSC))
    return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);

  // Bitcast for pointers requires that the address spaces must match
  return false;
}

static unsigned getFCmpOpcode(unsigned PredNum) {
  auto Pred = static_cast<CmpInst::Predicate>(PredNum);
  switch (Pred) {
  case CmpInst::FCMP_OEQ:
    return SPIRV::OpFOrdEqual;
  case CmpInst::FCMP_OGE:
    return SPIRV::OpFOrdGreaterThanEqual;
  case CmpInst::FCMP_OGT:
    return SPIRV::OpFOrdGreaterThan;
  case CmpInst::FCMP_OLE:
    return SPIRV::OpFOrdLessThanEqual;
  case CmpInst::FCMP_OLT:
    return SPIRV::OpFOrdLessThan;
  case CmpInst::FCMP_ONE:
    return SPIRV::OpFOrdNotEqual;
  case CmpInst::FCMP_ORD:
    return SPIRV::OpOrdered;
  case CmpInst::FCMP_UEQ:
    return SPIRV::OpFUnordEqual;
  case CmpInst::FCMP_UGE:
    return SPIRV::OpFUnordGreaterThanEqual;
  case CmpInst::FCMP_UGT:
    return SPIRV::OpFUnordGreaterThan;
  case CmpInst::FCMP_ULE:
    return SPIRV::OpFUnordLessThanEqual;
  case CmpInst::FCMP_ULT:
    return SPIRV::OpFUnordLessThan;
  case CmpInst::FCMP_UNE:
    return SPIRV::OpFUnordNotEqual;
  case CmpInst::FCMP_UNO:
    return SPIRV::OpUnordered;
  default:
    llvm_unreachable("Unknown predicate type for FCmp");
  }
}

static unsigned getICmpOpcode(unsigned PredNum) {
  auto Pred = static_cast<CmpInst::Predicate>(PredNum);
  switch (Pred) {
  case CmpInst::ICMP_EQ:
    return SPIRV::OpIEqual;
  case CmpInst::ICMP_NE:
    return SPIRV::OpINotEqual;
  case CmpInst::ICMP_SGE:
    return SPIRV::OpSGreaterThanEqual;
  case CmpInst::ICMP_SGT:
    return SPIRV::OpSGreaterThan;
  case CmpInst::ICMP_SLE:
    return SPIRV::OpSLessThanEqual;
  case CmpInst::ICMP_SLT:
    return SPIRV::OpSLessThan;
  case CmpInst::ICMP_UGE:
    return SPIRV::OpUGreaterThanEqual;
  case CmpInst::ICMP_UGT:
    return SPIRV::OpUGreaterThan;
  case CmpInst::ICMP_ULE:
    return SPIRV::OpULessThanEqual;
  case CmpInst::ICMP_ULT:
    return SPIRV::OpULessThan;
  default:
    llvm_unreachable("Unknown predicate type for ICmp");
  }
}

static unsigned getPtrCmpOpcode(unsigned Pred) {
  switch (static_cast<CmpInst::Predicate>(Pred)) {
  case CmpInst::ICMP_EQ:
    return SPIRV::OpPtrEqual;
  case CmpInst::ICMP_NE:
    return SPIRV::OpPtrNotEqual;
  default:
    llvm_unreachable("Unknown predicate type for pointer comparison");
  }
}

// Return the logical operation, or abort if none exists.
static unsigned getBoolCmpOpcode(unsigned PredNum) {
  auto Pred = static_cast<CmpInst::Predicate>(PredNum);
  switch (Pred) {
  case CmpInst::ICMP_EQ:
    return SPIRV::OpLogicalEqual;
  case CmpInst::ICMP_NE:
    return SPIRV::OpLogicalNotEqual;
  default:
    llvm_unreachable("Unknown predicate type for Bool comparison");
  }
}

static APFloat getZeroFP(const Type *LLVMFloatTy) {
  if (!LLVMFloatTy)
    return APFloat::getZero(APFloat::IEEEsingle());
  switch (LLVMFloatTy->getScalarType()->getTypeID()) {
  case Type::HalfTyID:
    return APFloat::getZero(APFloat::IEEEhalf());
  default:
  case Type::FloatTyID:
    return APFloat::getZero(APFloat::IEEEsingle());
  case Type::DoubleTyID:
    return APFloat::getZero(APFloat::IEEEdouble());
  }
}

static APFloat getOneFP(const Type *LLVMFloatTy) {
  if (!LLVMFloatTy)
    return APFloat::getOne(APFloat::IEEEsingle());
  switch (LLVMFloatTy->getScalarType()->getTypeID()) {
  case Type::HalfTyID:
    return APFloat::getOne(APFloat::IEEEhalf());
  default:
  case Type::FloatTyID:
    return APFloat::getOne(APFloat::IEEEsingle());
  case Type::DoubleTyID:
    return APFloat::getOne(APFloat::IEEEdouble());
  }
}

bool SPIRVInstructionSelector::selectAnyOrAll(Register ResVReg,
                                              const SPIRVType *ResType,
                                              MachineInstr &I,
                                              unsigned OpAnyOrAll) const {
  assert(I.getNumOperands() == 3);
  assert(I.getOperand(2).isReg());
  MachineBasicBlock &BB = *I.getParent();
  Register InputRegister = I.getOperand(2).getReg();
  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);

  if (!InputType)
    report_fatal_error("Input Type could not be determined.");

  bool IsBoolTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeBool);
  bool IsVectorTy = InputType->getOpcode() == SPIRV::OpTypeVector;
  if (IsBoolTy && !IsVectorTy) {
    assert(ResVReg == I.getOperand(0).getReg());
    return BuildCOPY(ResVReg, InputRegister, I);
  }

  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
  unsigned SpirvNotEqualId =
      IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
  SPIRVType *SpvBoolScalarTy = GR.getOrCreateSPIRVBoolType(I, TII);
  SPIRVType *SpvBoolTy = SpvBoolScalarTy;
  Register NotEqualReg = ResVReg;

  if (IsVectorTy) {
    NotEqualReg =
        IsBoolTy ? InputRegister
                 : createVirtualRegister(SpvBoolTy, &GR, MRI, MRI->getMF());
    const unsigned NumElts = InputType->getOperand(2).getImm();
    SpvBoolTy = GR.getOrCreateSPIRVVectorType(SpvBoolTy, NumElts, I, TII);
  }

  bool Result = true;
  if (!IsBoolTy) {
    Register ConstZeroReg =
        IsFloatTy ? buildZerosValF(InputType, I) : buildZerosVal(InputType, I);

    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SpirvNotEqualId))
                  .addDef(NotEqualReg)
                  .addUse(GR.getSPIRVTypeID(SpvBoolTy))
                  .addUse(InputRegister)
                  .addUse(ConstZeroReg)
                  .constrainAllUses(TII, TRI, RBI);
  }

  if (!IsVectorTy)
    return Result;

  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(OpAnyOrAll))
                       .addDef(ResVReg)
                       .addUse(GR.getSPIRVTypeID(SpvBoolScalarTy))
                       .addUse(NotEqualReg)
                       .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectAll(Register ResVReg,
                                         const SPIRVType *ResType,
                                         MachineInstr &I) const {
  return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAll);
}

bool SPIRVInstructionSelector::selectAny(Register ResVReg,
                                         const SPIRVType *ResType,
                                         MachineInstr &I) const {
  return selectAnyOrAll(ResVReg, ResType, I, SPIRV::OpAny);
}

// Select the OpDot instruction for the given float dot
bool SPIRVInstructionSelector::selectFloatDot(Register ResVReg,
                                              const SPIRVType *ResType,
                                              MachineInstr &I) const {
  assert(I.getNumOperands() == 4);
  assert(I.getOperand(2).isReg());
  assert(I.getOperand(3).isReg());

  [[maybe_unused]] SPIRVType *VecType =
      GR.getSPIRVTypeForVReg(I.getOperand(2).getReg());

  assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
         GR.getScalarOrVectorComponentCount(VecType) > 1 &&
         "dot product requires a vector of at least 2 components");

  [[maybe_unused]] SPIRVType *EltType =
      GR.getSPIRVTypeForVReg(VecType->getOperand(1).getReg());

  assert(EltType->getOpcode() == SPIRV::OpTypeFloat);

  MachineBasicBlock &BB = *I.getParent();
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpDot))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(I.getOperand(2).getReg())
      .addUse(I.getOperand(3).getReg())
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectIntegerDot(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I,
                                                bool Signed) const {
  assert(I.getNumOperands() == 4);
  assert(I.getOperand(2).isReg());
  assert(I.getOperand(3).isReg());
  MachineBasicBlock &BB = *I.getParent();

  auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(I.getOperand(2).getReg())
      .addUse(I.getOperand(3).getReg())
      .constrainAllUses(TII, TRI, RBI);
}

// Since pre-1.6 SPIRV has no integer dot implementation,
// expand by piecewise multiplying and adding the results
bool SPIRVInstructionSelector::selectIntegerDotExpansion(
    Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
  assert(I.getNumOperands() == 4);
  assert(I.getOperand(2).isReg());
  assert(I.getOperand(3).isReg());
  MachineBasicBlock &BB = *I.getParent();

  // Multiply the vectors, then sum the results
  Register Vec0 = I.getOperand(2).getReg();
  Register Vec1 = I.getOperand(3).getReg();
  Register TmpVec = MRI->createVirtualRegister(GR.getRegClass(ResType));
  SPIRVType *VecType = GR.getSPIRVTypeForVReg(Vec0);

  bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulV))
                    .addDef(TmpVec)
                    .addUse(GR.getSPIRVTypeID(VecType))
                    .addUse(Vec0)
                    .addUse(Vec1)
                    .constrainAllUses(TII, TRI, RBI);

  assert(VecType->getOpcode() == SPIRV::OpTypeVector &&
         GR.getScalarOrVectorComponentCount(VecType) > 1 &&
         "dot product requires a vector of at least 2 components");

  Register Res = MRI->createVirtualRegister(GR.getRegClass(ResType));
  Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
                .addDef(Res)
                .addUse(GR.getSPIRVTypeID(ResType))
                .addUse(TmpVec)
                .addImm(0)
                .constrainAllUses(TII, TRI, RBI);

  for (unsigned i = 1; i < GR.getScalarOrVectorComponentCount(VecType); i++) {
    Register Elt = MRI->createVirtualRegister(GR.getRegClass(ResType));

    Result &=
        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
            .addDef(Elt)
            .addUse(GR.getSPIRVTypeID(ResType))
            .addUse(TmpVec)
            .addImm(i)
            .constrainAllUses(TII, TRI, RBI);

    Register Sum = i < GR.getScalarOrVectorComponentCount(VecType) - 1
                       ? MRI->createVirtualRegister(GR.getRegClass(ResType))
                       : ResVReg;

    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
                  .addDef(Sum)
                  .addUse(GR.getSPIRVTypeID(ResType))
                  .addUse(Res)
                  .addUse(Elt)
                  .constrainAllUses(TII, TRI, RBI);
    Res = Sum;
  }

  return Result;
}

template <bool Signed>
bool SPIRVInstructionSelector::selectDot4AddPacked(Register ResVReg,
                                                   const SPIRVType *ResType,
                                                   MachineInstr &I) const {
  assert(I.getNumOperands() == 5);
  assert(I.getOperand(2).isReg());
  assert(I.getOperand(3).isReg());
  assert(I.getOperand(4).isReg());
  MachineBasicBlock &BB = *I.getParent();

  Register Acc = I.getOperand(2).getReg();
  Register X = I.getOperand(3).getReg();
  Register Y = I.getOperand(4).getReg();

  auto DotOp = Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
  Register Dot = MRI->createVirtualRegister(GR.getRegClass(ResType));
  bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(DotOp))
                    .addDef(Dot)
                    .addUse(GR.getSPIRVTypeID(ResType))
                    .addUse(X)
                    .addUse(Y)
                    .constrainAllUses(TII, TRI, RBI);

  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
                       .addDef(ResVReg)
                       .addUse(GR.getSPIRVTypeID(ResType))
                       .addUse(Dot)
                       .addUse(Acc)
                       .constrainAllUses(TII, TRI, RBI);
}

// Since pre-1.6 SPIRV has no DotProductInput4x8BitPacked implementation,
// extract the elements of the packed inputs, multiply them and add the result
// to the accumulator.
template <bool Signed>
bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
    Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
  assert(I.getNumOperands() == 5);
  assert(I.getOperand(2).isReg());
  assert(I.getOperand(3).isReg());
  assert(I.getOperand(4).isReg());
  MachineBasicBlock &BB = *I.getParent();

  bool Result = true;

  Register Acc = I.getOperand(2).getReg();
  Register X = I.getOperand(3).getReg();
  Register Y = I.getOperand(4).getReg();

  SPIRVType *EltType = GR.getOrCreateSPIRVIntegerType(8, I, TII);
  auto ExtractOp =
      Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;

  bool ZeroAsNull = !STI.isShader();
  // Extract the i8 element, multiply and add it to the accumulator
  for (unsigned i = 0; i < 4; i++) {
    // A[i]
    Register AElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
    Result &=
        BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
            .addDef(AElt)
            .addUse(GR.getSPIRVTypeID(ResType))
            .addUse(X)
            .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
            .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
            .constrainAllUses(TII, TRI, RBI);

    // B[i]
    Register BElt = MRI->createVirtualRegister(&SPIRV::IDRegClass);
    Result &=
        BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
            .addDef(BElt)
            .addUse(GR.getSPIRVTypeID(ResType))
            .addUse(Y)
            .addUse(GR.getOrCreateConstInt(i * 8, I, EltType, TII, ZeroAsNull))
            .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
            .constrainAllUses(TII, TRI, RBI);

    // A[i] * B[i]
    Register Mul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIMulS))
                  .addDef(Mul)
                  .addUse(GR.getSPIRVTypeID(ResType))
                  .addUse(AElt)
                  .addUse(BElt)
                  .constrainAllUses(TII, TRI, RBI);

    // Discard 24 highest-bits so that stored i32 register is i8 equivalent
    Register MaskMul = MRI->createVirtualRegister(&SPIRV::IDRegClass);
    Result &=
        BuildMI(BB, I, I.getDebugLoc(), TII.get(ExtractOp))
            .addDef(MaskMul)
            .addUse(GR.getSPIRVTypeID(ResType))
            .addUse(Mul)
            .addUse(GR.getOrCreateConstInt(0, I, EltType, TII, ZeroAsNull))
            .addUse(GR.getOrCreateConstInt(8, I, EltType, TII, ZeroAsNull))
            .constrainAllUses(TII, TRI, RBI);

    // Acc = Acc + A[i] * B[i]
    Register Sum =
        i < 3 ? MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;
    Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpIAddS))
                  .addDef(Sum)
                  .addUse(GR.getSPIRVTypeID(ResType))
                  .addUse(Acc)
                  .addUse(MaskMul)
                  .constrainAllUses(TII, TRI, RBI);

    Acc = Sum;
  }

  return Result;
}

/// Transform saturate(x) to clamp(x, 0.0f, 1.0f) as SPIRV
/// does not have a saturate builtin.
bool SPIRVInstructionSelector::selectSaturate(Register ResVReg,
                                              const SPIRVType *ResType,
                                              MachineInstr &I) const {
  assert(I.getNumOperands() == 3);
  assert(I.getOperand(2).isReg());
  MachineBasicBlock &BB = *I.getParent();
  Register VZero = buildZerosValF(ResType, I);
  Register VOne = buildOnesValF(ResType, I);

  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
      .addImm(GL::FClamp)
      .addUse(I.getOperand(2).getReg())
      .addUse(VZero)
      .addUse(VOne)
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectSign(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I) const {
  assert(I.getNumOperands() == 3);
  assert(I.getOperand(2).isReg());
  MachineBasicBlock &BB = *I.getParent();
  Register InputRegister = I.getOperand(2).getReg();
  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);
  auto &DL = I.getDebugLoc();

  if (!InputType)
    report_fatal_error("Input Type could not be determined.");

  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);

  unsigned SignBitWidth = GR.getScalarOrVectorBitWidth(InputType);
  unsigned ResBitWidth = GR.getScalarOrVectorBitWidth(ResType);

  bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;

  auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
  Register SignReg = NeedsConversion
                         ? MRI->createVirtualRegister(&SPIRV::IDRegClass)
                         : ResVReg;

  bool Result =
      BuildMI(BB, I, DL, TII.get(SPIRV::OpExtInst))
          .addDef(SignReg)
          .addUse(GR.getSPIRVTypeID(InputType))
          .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
          .addImm(SignOpcode)
          .addUse(InputRegister)
          .constrainAllUses(TII, TRI, RBI);

  if (NeedsConversion) {
    auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
    Result &= BuildMI(*I.getParent(), I, DL, TII.get(ConvertOpcode))
                  .addDef(ResVReg)
                  .addUse(GR.getSPIRVTypeID(ResType))
                  .addUse(SignReg)
                  .constrainAllUses(TII, TRI, RBI);
  }

  return Result;
}

bool SPIRVInstructionSelector::selectWaveOpInst(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I,
                                                unsigned Opcode) const {
  MachineBasicBlock &BB = *I.getParent();
  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);

  auto BMI = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I,
                                                IntTy, TII, !STI.isShader()));

  for (unsigned J = 2; J < I.getNumOperands(); J++) {
    BMI.addUse(I.getOperand(J).getReg());
  }

  return BMI.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectWaveActiveCountBits(
    Register ResVReg, const SPIRVType *ResType, MachineInstr &I) const {

  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
  SPIRVType *BallotType = GR.getOrCreateSPIRVVectorType(IntTy, 4, I, TII);
  Register BallotReg = MRI->createVirtualRegister(GR.getRegClass(BallotType));
  bool Result = selectWaveOpInst(BallotReg, BallotType, I,
                                 SPIRV::OpGroupNonUniformBallot);

  MachineBasicBlock &BB = *I.getParent();
  Result &= BuildMI(BB, I, I.getDebugLoc(),
                    TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
                .addDef(ResVReg)
                .addUse(GR.getSPIRVTypeID(ResType))
                .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy,
                                               TII, !STI.isShader()))
                .addImm(SPIRV::GroupOperation::Reduce)
                .addUse(BallotReg)
                .constrainAllUses(TII, TRI, RBI);

  return Result;
}

bool SPIRVInstructionSelector::selectWaveReduceMax(Register ResVReg,
                                                   const SPIRVType *ResType,
                                                   MachineInstr &I,
                                                   bool IsUnsigned) const {
  assert(I.getNumOperands() == 3);
  assert(I.getOperand(2).isReg());
  MachineBasicBlock &BB = *I.getParent();
  Register InputRegister = I.getOperand(2).getReg();
  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);

  if (!InputType)
    report_fatal_error("Input Type could not be determined.");

  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
  // Retreive the operation to use based on input type
  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
  auto IntegerOpcodeType =
      IsUnsigned ? SPIRV::OpGroupNonUniformUMax : SPIRV::OpGroupNonUniformSMax;
  auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntegerOpcodeType;
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
                                     !STI.isShader()))
      .addImm(SPIRV::GroupOperation::Reduce)
      .addUse(I.getOperand(2).getReg())
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectWaveReduceSum(Register ResVReg,
                                                   const SPIRVType *ResType,
                                                   MachineInstr &I) const {
  assert(I.getNumOperands() == 3);
  assert(I.getOperand(2).isReg());
  MachineBasicBlock &BB = *I.getParent();
  Register InputRegister = I.getOperand(2).getReg();
  SPIRVType *InputType = GR.getSPIRVTypeForVReg(InputRegister);

  if (!InputType)
    report_fatal_error("Input Type could not be determined.");

  SPIRVType *IntTy = GR.getOrCreateSPIRVIntegerType(32, I, TII);
  // Retreive the operation to use based on input type
  bool IsFloatTy = GR.isScalarOrVectorOfType(InputRegister, SPIRV::OpTypeFloat);
  auto Opcode =
      IsFloatTy ? SPIRV::OpGroupNonUniformFAdd : SPIRV::OpGroupNonUniformIAdd;
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(GR.getOrCreateConstInt(SPIRV::Scope::Subgroup, I, IntTy, TII,
                                     !STI.isShader()))
      .addImm(SPIRV::GroupOperation::Reduce)
      .addUse(I.getOperand(2).getReg());
}

bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
  MachineBasicBlock &BB = *I.getParent();
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitReverse))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(I.getOperand(1).getReg())
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectFreeze(Register ResVReg,
                                            const SPIRVType *ResType,
                                            MachineInstr &I) const {
  // There is no way to implement `freeze` correctly without support on SPIR-V
  // standard side, but we may at least address a simple (static) case when
  // undef/poison value presence is obvious. The main benefit of even
  // incomplete `freeze` support is preventing of translation from crashing due
  // to lack of support on legalization and instruction selection steps.
  if (!I.getOperand(0).isReg() || !I.getOperand(1).isReg())
    return false;
  Register OpReg = I.getOperand(1).getReg();
  if (MachineInstr *Def = MRI->getVRegDef(OpReg)) {
    if (Def->getOpcode() == TargetOpcode::COPY)
      Def = MRI->getVRegDef(Def->getOperand(1).getReg());
    Register Reg;
    switch (Def->getOpcode()) {
    case SPIRV::ASSIGN_TYPE:
      if (MachineInstr *AssignToDef =
              MRI->getVRegDef(Def->getOperand(1).getReg())) {
        if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
          Reg = Def->getOperand(2).getReg();
      }
      break;
    case SPIRV::OpUndef:
      Reg = Def->getOperand(1).getReg();
      break;
    }
    unsigned DestOpCode;
    if (Reg.isValid()) {
      DestOpCode = SPIRV::OpConstantNull;
    } else {
      DestOpCode = TargetOpcode::COPY;
      Reg = OpReg;
    }
    return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(DestOpCode))
        .addDef(I.getOperand(0).getReg())
        .addUse(Reg)
        .constrainAllUses(TII, TRI, RBI);
  }
  return false;
}

bool SPIRVInstructionSelector::selectBuildVector(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
  unsigned N = 0;
  if (ResType->getOpcode() == SPIRV::OpTypeVector)
    N = GR.getScalarOrVectorComponentCount(ResType);
  else if (ResType->getOpcode() == SPIRV::OpTypeArray)
    N = getArrayComponentCount(MRI, ResType);
  else
    report_fatal_error("Cannot select G_BUILD_VECTOR with a non-vector result");
  if (I.getNumExplicitOperands() - I.getNumExplicitDefs() != N)
    report_fatal_error("G_BUILD_VECTOR and the result type are inconsistent");

  // check if we may construct a constant vector
  bool IsConst = true;
  for (unsigned i = I.getNumExplicitDefs();
       i < I.getNumExplicitOperands() && IsConst; ++i)
    if (!isConstReg(MRI, I.getOperand(i).getReg()))
      IsConst = false;

  if (!IsConst && N < 2)
    report_fatal_error(
        "There must be at least two constituent operands in a vector");

  MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                     TII.get(IsConst ? SPIRV::OpConstantComposite
                                     : SPIRV::OpCompositeConstruct))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType));
  for (unsigned i = I.getNumExplicitDefs(); i < I.getNumExplicitOperands(); ++i)
    MIB.addUse(I.getOperand(i).getReg());
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectSplatVector(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
  unsigned N = 0;
  if (ResType->getOpcode() == SPIRV::OpTypeVector)
    N = GR.getScalarOrVectorComponentCount(ResType);
  else if (ResType->getOpcode() == SPIRV::OpTypeArray)
    N = getArrayComponentCount(MRI, ResType);
  else
    report_fatal_error("Cannot select G_SPLAT_VECTOR with a non-vector result");

  unsigned OpIdx = I.getNumExplicitDefs();
  if (!I.getOperand(OpIdx).isReg())
    report_fatal_error("Unexpected argument in G_SPLAT_VECTOR");

  // check if we may construct a constant vector
  Register OpReg = I.getOperand(OpIdx).getReg();
  bool IsConst = isConstReg(MRI, OpReg);

  if (!IsConst && N < 2)
    report_fatal_error(
        "There must be at least two constituent operands in a vector");

  MRI->setRegClass(ResVReg, GR.getRegClass(ResType));
  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                     TII.get(IsConst ? SPIRV::OpConstantComposite
                                     : SPIRV::OpCompositeConstruct))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType));
  for (unsigned i = 0; i < N; ++i)
    MIB.addUse(OpReg);
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectDiscard(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I) const {

  unsigned Opcode;

  if (STI.canUseExtension(
          SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
      STI.isAtLeastSPIRVVer(llvm::VersionTuple(1, 6))) {
    Opcode = SPIRV::OpDemoteToHelperInvocation;
  } else {
    Opcode = SPIRV::OpKill;
    // OpKill must be the last operation of any basic block.
    if (MachineInstr *NextI = I.getNextNode()) {
      GR.invalidateMachineInstr(NextI);
      NextI->removeFromParent();
    }
  }

  MachineBasicBlock &BB = *I.getParent();
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectCmp(Register ResVReg,
                                         const SPIRVType *ResType,
                                         unsigned CmpOpc,
                                         MachineInstr &I) const {
  Register Cmp0 = I.getOperand(2).getReg();
  Register Cmp1 = I.getOperand(3).getReg();
  assert(GR.getSPIRVTypeForVReg(Cmp0)->getOpcode() ==
             GR.getSPIRVTypeForVReg(Cmp1)->getOpcode() &&
         "CMP operands should have the same type");
  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(CmpOpc))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(Cmp0)
      .addUse(Cmp1)
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectICmp(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I) const {
  auto Pred = I.getOperand(1).getPredicate();
  unsigned CmpOpc;

  Register CmpOperand = I.getOperand(2).getReg();
  if (GR.isScalarOfType(CmpOperand, SPIRV::OpTypePointer))
    CmpOpc = getPtrCmpOpcode(Pred);
  else if (GR.isScalarOrVectorOfType(CmpOperand, SPIRV::OpTypeBool))
    CmpOpc = getBoolCmpOpcode(Pred);
  else
    CmpOpc = getICmpOpcode(Pred);
  return selectCmp(ResVReg, ResType, CmpOpc, I);
}

std::pair<Register, bool>
SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &I,
                                           const SPIRVType *ResType) const {
  Type *LLVMTy = IntegerType::get(GR.CurMF->getFunction().getContext(), 32);
  const SPIRVType *SpvI32Ty =
      ResType ? ResType : GR.getOrCreateSPIRVIntegerType(32, I, TII);
  // Find a constant in DT or build a new one.
  auto ConstInt = ConstantInt::get(LLVMTy, Val);
  Register NewReg = GR.find(ConstInt, GR.CurMF);
  bool Result = true;
  if (!NewReg.isValid()) {
    NewReg = MRI->createGenericVirtualRegister(LLT::scalar(64));
    MachineBasicBlock &BB = *I.getParent();
    MachineInstr *MI =
        Val == 0
            ? BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
                  .addDef(NewReg)
                  .addUse(GR.getSPIRVTypeID(SpvI32Ty))
            : BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantI))
                  .addDef(NewReg)
                  .addUse(GR.getSPIRVTypeID(SpvI32Ty))
                  .addImm(APInt(32, Val).getZExtValue());
    Result &= constrainSelectedInstRegOperands(*MI, TII, TRI, RBI);
    GR.add(ConstInt, MI);
  }
  return {NewReg, Result};
}

bool SPIRVInstructionSelector::selectFCmp(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I) const {
  unsigned CmpOp = getFCmpOpcode(I.getOperand(1).getPredicate());
  return selectCmp(ResVReg, ResType, CmpOp, I);
}

Register SPIRVInstructionSelector::buildZerosVal(const SPIRVType *ResType,
                                                 MachineInstr &I) const {
  // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
  bool ZeroAsNull = !STI.isShader();
  if (ResType->getOpcode() == SPIRV::OpTypeVector)
    return GR.getOrCreateConstVector(0UL, I, ResType, TII, ZeroAsNull);
  return GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
}

Register SPIRVInstructionSelector::buildZerosValF(const SPIRVType *ResType,
                                                  MachineInstr &I) const {
  // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
  bool ZeroAsNull = !STI.isShader();
  APFloat VZero = getZeroFP(GR.getTypeForSPIRVType(ResType));
  if (ResType->getOpcode() == SPIRV::OpTypeVector)
    return GR.getOrCreateConstVector(VZero, I, ResType, TII, ZeroAsNull);
  return GR.getOrCreateConstFP(VZero, I, ResType, TII, ZeroAsNull);
}

Register SPIRVInstructionSelector::buildOnesValF(const SPIRVType *ResType,
                                                 MachineInstr &I) const {
  // OpenCL uses nulls for Zero. In HLSL we don't use null constants.
  bool ZeroAsNull = !STI.isShader();
  APFloat VOne = getOneFP(GR.getTypeForSPIRVType(ResType));
  if (ResType->getOpcode() == SPIRV::OpTypeVector)
    return GR.getOrCreateConstVector(VOne, I, ResType, TII, ZeroAsNull);
  return GR.getOrCreateConstFP(VOne, I, ResType, TII, ZeroAsNull);
}

Register SPIRVInstructionSelector::buildOnesVal(bool AllOnes,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
  unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
  APInt One =
      AllOnes ? APInt::getAllOnes(BitWidth) : APInt::getOneBitSet(BitWidth, 0);
  if (ResType->getOpcode() == SPIRV::OpTypeVector)
    return GR.getOrCreateConstVector(One.getZExtValue(), I, ResType, TII);
  return GR.getOrCreateConstInt(One.getZExtValue(), I, ResType, TII);
}

bool SPIRVInstructionSelector::selectSelect(Register ResVReg,
                                            const SPIRVType *ResType,
                                            MachineInstr &I,
                                            bool IsSigned) const {
  // To extend a bool, we need to use OpSelect between constants.
  Register ZeroReg = buildZerosVal(ResType, I);
  Register OneReg = buildOnesVal(IsSigned, ResType, I);
  bool IsScalarBool =
      GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
  unsigned Opcode =
      IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(I.getOperand(1).getReg())
      .addUse(OneReg)
      .addUse(ZeroReg)
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectIToF(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I, bool IsSigned,
                                          unsigned Opcode) const {
  Register SrcReg = I.getOperand(1).getReg();
  // We can convert bool value directly to float type without OpConvert*ToF,
  // however the translator generates OpSelect+OpConvert*ToF, so we do the same.
  if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
    unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
    SPIRVType *TmpType = GR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);
    if (ResType->getOpcode() == SPIRV::OpTypeVector) {
      const unsigned NumElts = ResType->getOperand(2).getImm();
      TmpType = GR.getOrCreateSPIRVVectorType(TmpType, NumElts, I, TII);
    }
    SrcReg = createVirtualRegister(TmpType, &GR, MRI, MRI->getMF());
    selectSelect(SrcReg, TmpType, I, false);
  }
  return selectOpWithSrcs(ResVReg, ResType, I, {SrcReg}, Opcode);
}

bool SPIRVInstructionSelector::selectExt(Register ResVReg,
                                         const SPIRVType *ResType,
                                         MachineInstr &I, bool IsSigned) const {
  Register SrcReg = I.getOperand(1).getReg();
  if (GR.isScalarOrVectorOfType(SrcReg, SPIRV::OpTypeBool))
    return selectSelect(ResVReg, ResType, I, IsSigned);

  SPIRVType *SrcType = GR.getSPIRVTypeForVReg(SrcReg);
  if (SrcType == ResType)
    return BuildCOPY(ResVReg, SrcReg, I);

  unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
  return selectUnOp(ResVReg, ResType, I, Opcode);
}

bool SPIRVInstructionSelector::selectSUCmp(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I,
                                           bool IsSigned) const {
  MachineIRBuilder MIRBuilder(I);
  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
  MachineBasicBlock &BB = *I.getParent();
  // Ensure we have bool.
  SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
  unsigned N = GR.getScalarOrVectorComponentCount(ResType);
  if (N > 1)
    BoolType = GR.getOrCreateSPIRVVectorType(BoolType, N, I, TII);
  Register BoolTypeReg = GR.getSPIRVTypeID(BoolType);
  // Build less-than-equal and less-than.
  // TODO: replace with one-liner createVirtualRegister() from
  // llvm/lib/Target/SPIRV/SPIRVUtils.cpp when PR #116609 is merged.
  Register IsLessEqReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
  MRI->setType(IsLessEqReg, LLT::scalar(64));
  GR.assignSPIRVTypeToVReg(ResType, IsLessEqReg, MIRBuilder.getMF());
  bool Result = BuildMI(BB, I, I.getDebugLoc(),
                        TII.get(IsSigned ? SPIRV::OpSLessThanEqual
                                         : SPIRV::OpULessThanEqual))
                    .addDef(IsLessEqReg)
                    .addUse(BoolTypeReg)
                    .addUse(I.getOperand(1).getReg())
                    .addUse(I.getOperand(2).getReg())
                    .constrainAllUses(TII, TRI, RBI);
  Register IsLessReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
  MRI->setType(IsLessReg, LLT::scalar(64));
  GR.assignSPIRVTypeToVReg(ResType, IsLessReg, MIRBuilder.getMF());
  Result &= BuildMI(BB, I, I.getDebugLoc(),
                    TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
                .addDef(IsLessReg)
                .addUse(BoolTypeReg)
                .addUse(I.getOperand(1).getReg())
                .addUse(I.getOperand(2).getReg())
                .constrainAllUses(TII, TRI, RBI);
  // Build selects.
  Register ResTypeReg = GR.getSPIRVTypeID(ResType);
  Register NegOneOrZeroReg =
      MRI->createVirtualRegister(GR.getRegClass(ResType));
  MRI->setType(NegOneOrZeroReg, LLT::scalar(64));
  GR.assignSPIRVTypeToVReg(ResType, NegOneOrZeroReg, MIRBuilder.getMF());
  unsigned SelectOpcode =
      N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
  Result &= BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
                .addDef(NegOneOrZeroReg)
                .addUse(ResTypeReg)
                .addUse(IsLessReg)
                .addUse(buildOnesVal(true, ResType, I)) // -1
                .addUse(buildZerosVal(ResType, I))
                .constrainAllUses(TII, TRI, RBI);
  return Result & BuildMI(BB, I, I.getDebugLoc(), TII.get(SelectOpcode))
                      .addDef(ResVReg)
                      .addUse(ResTypeReg)
                      .addUse(IsLessEqReg)
                      .addUse(NegOneOrZeroReg) // -1 or 0
                      .addUse(buildOnesVal(false, ResType, I))
                      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectIntToBool(Register IntReg,
                                               Register ResVReg,
                                               MachineInstr &I,
                                               const SPIRVType *IntTy,
                                               const SPIRVType *BoolTy) const {
  // To truncate to a bool, we use OpBitwiseAnd 1 and OpINotEqual to zero.
  Register BitIntReg = createVirtualRegister(IntTy, &GR, MRI, MRI->getMF());
  bool IsVectorTy = IntTy->getOpcode() == SPIRV::OpTypeVector;
  unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
  Register Zero = buildZerosVal(IntTy, I);
  Register One = buildOnesVal(false, IntTy, I);
  MachineBasicBlock &BB = *I.getParent();
  bool Result = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
                    .addDef(BitIntReg)
                    .addUse(GR.getSPIRVTypeID(IntTy))
                    .addUse(IntReg)
                    .addUse(One)
                    .constrainAllUses(TII, TRI, RBI);
  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
                       .addDef(ResVReg)
                       .addUse(GR.getSPIRVTypeID(BoolTy))
                       .addUse(BitIntReg)
                       .addUse(Zero)
                       .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectTrunc(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I) const {
  Register IntReg = I.getOperand(1).getReg();
  const SPIRVType *ArgType = GR.getSPIRVTypeForVReg(IntReg);
  if (GR.isScalarOrVectorOfType(ResVReg, SPIRV::OpTypeBool))
    return selectIntToBool(IntReg, ResVReg, I, ArgType, ResType);
  if (ArgType == ResType)
    return BuildCOPY(ResVReg, IntReg, I);
  bool IsSigned = GR.isScalarOrVectorSigned(ResType);
  unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
  return selectUnOp(ResVReg, ResType, I, Opcode);
}

bool SPIRVInstructionSelector::selectConst(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I) const {
  unsigned Opcode = I.getOpcode();
  unsigned TpOpcode = ResType->getOpcode();
  Register Reg;
  if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
    assert(Opcode == TargetOpcode::G_CONSTANT &&
           I.getOperand(1).getCImm()->isZero());
    MachineBasicBlock &DepMBB = I.getMF()->front();
    MachineIRBuilder MIRBuilder(DepMBB, DepMBB.getFirstNonPHI());
    Reg = GR.getOrCreateConstNullPtr(MIRBuilder, ResType);
  } else if (Opcode == TargetOpcode::G_FCONSTANT) {
    Reg = GR.getOrCreateConstFP(I.getOperand(1).getFPImm()->getValue(), I,
                                ResType, TII, !STI.isShader());
  } else {
    Reg = GR.getOrCreateConstInt(I.getOperand(1).getCImm()->getZExtValue(), I,
                                 ResType, TII, !STI.isShader());
  }
  return Reg == ResVReg ? true : BuildCOPY(ResVReg, Reg, I);
}

bool SPIRVInstructionSelector::selectOpUndef(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I) const {
  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectInsertVal(Register ResVReg,
                                               const SPIRVType *ResType,
                                               MachineInstr &I) const {
  MachineBasicBlock &BB = *I.getParent();
  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeInsert))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 // object to insert
                 .addUse(I.getOperand(3).getReg())
                 // composite to insert into
                 .addUse(I.getOperand(2).getReg());
  for (unsigned i = 4; i < I.getNumOperands(); i++)
    MIB.addImm(foldImm(I.getOperand(i), MRI));
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectExtractVal(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
  MachineBasicBlock &BB = *I.getParent();
  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 .addUse(I.getOperand(2).getReg());
  for (unsigned i = 3; i < I.getNumOperands(); i++)
    MIB.addImm(foldImm(I.getOperand(i), MRI));
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectInsertElt(Register ResVReg,
                                               const SPIRVType *ResType,
                                               MachineInstr &I) const {
  if (getImm(I.getOperand(4), MRI))
    return selectInsertVal(ResVReg, ResType, I);
  MachineBasicBlock &BB = *I.getParent();
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorInsertDynamic))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(I.getOperand(2).getReg())
      .addUse(I.getOperand(3).getReg())
      .addUse(I.getOperand(4).getReg())
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectExtractElt(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
  if (getImm(I.getOperand(3), MRI))
    return selectExtractVal(ResVReg, ResType, I);
  MachineBasicBlock &BB = *I.getParent();
  return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorExtractDynamic))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(I.getOperand(2).getReg())
      .addUse(I.getOperand(3).getReg())
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectGEP(Register ResVReg,
                                         const SPIRVType *ResType,
                                         MachineInstr &I) const {
  const bool IsGEPInBounds = I.getOperand(2).getImm();

  // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator only
  // relies on PtrAccessChain, so we'll try not to deviate. For Vulkan however,
  // we have to use Op[InBounds]AccessChain.
  const unsigned Opcode = STI.isLogicalSPIRV()
                              ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
                                               : SPIRV::OpAccessChain)
                              : (IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
                                               : SPIRV::OpPtrAccessChain);

  auto Res = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 // Object to get a pointer to.
                 .addUse(I.getOperand(3).getReg());
  // Adding indices.
  const unsigned StartingIndex =
      (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
          ? 5
          : 4;
  for (unsigned i = StartingIndex; i < I.getNumExplicitOperands(); ++i)
    Res.addUse(I.getOperand(i).getReg());
  return Res.constrainAllUses(TII, TRI, RBI);
}

// Maybe wrap a value into OpSpecConstantOp
bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
    MachineInstr &I, SmallVector<Register> &CompositeArgs) const {
  bool Result = true;
  unsigned Lim = I.getNumExplicitOperands();
  for (unsigned i = I.getNumExplicitDefs() + 1; i < Lim; ++i) {
    Register OpReg = I.getOperand(i).getReg();
    MachineInstr *OpDefine = MRI->getVRegDef(OpReg);
    SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
    SmallPtrSet<SPIRVType *, 4> Visited;
    if (!OpDefine || !OpType || isConstReg(MRI, OpDefine, Visited) ||
        OpDefine->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
        GR.isAggregateType(OpType)) {
      // The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed
      // by selectAddrSpaceCast()
      CompositeArgs.push_back(OpReg);
      continue;
    }
    MachineFunction *MF = I.getMF();
    Register WrapReg = GR.find(OpDefine, MF);
    if (WrapReg.isValid()) {
      CompositeArgs.push_back(WrapReg);
      continue;
    }
    // Create a new register for the wrapper
    WrapReg = MRI->createVirtualRegister(GR.getRegClass(OpType));
    CompositeArgs.push_back(WrapReg);
    // Decorate the wrapper register and generate a new instruction
    MRI->setType(WrapReg, LLT::pointer(0, 64));
    GR.assignSPIRVTypeToVReg(OpType, WrapReg, *MF);
    auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                       TII.get(SPIRV::OpSpecConstantOp))
                   .addDef(WrapReg)
                   .addUse(GR.getSPIRVTypeID(OpType))
                   .addImm(static_cast<uint32_t>(SPIRV::Opcode::Bitcast))
                   .addUse(OpReg);
    GR.add(OpDefine, MIB);
    Result = MIB.constrainAllUses(TII, TRI, RBI);
    if (!Result)
      break;
  }
  return Result;
}

bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
                                               const SPIRVType *ResType,
                                               MachineInstr &I) const {
  MachineBasicBlock &BB = *I.getParent();
  Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();
  switch (IID) {
  case Intrinsic::spv_load:
    return selectLoad(ResVReg, ResType, I);
  case Intrinsic::spv_store:
    return selectStore(I);
  case Intrinsic::spv_extractv:
    return selectExtractVal(ResVReg, ResType, I);
  case Intrinsic::spv_insertv:
    return selectInsertVal(ResVReg, ResType, I);
  case Intrinsic::spv_extractelt:
    return selectExtractElt(ResVReg, ResType, I);
  case Intrinsic::spv_insertelt:
    return selectInsertElt(ResVReg, ResType, I);
  case Intrinsic::spv_gep:
    return selectGEP(ResVReg, ResType, I);
  case Intrinsic::spv_unref_global:
  case Intrinsic::spv_init_global: {
    MachineInstr *MI = MRI->getVRegDef(I.getOperand(1).getReg());
    MachineInstr *Init = I.getNumExplicitOperands() > 2
                             ? MRI->getVRegDef(I.getOperand(2).getReg())
                             : nullptr;
    assert(MI);
    Register GVarVReg = MI->getOperand(0).getReg();
    bool Res = selectGlobalValue(GVarVReg, *MI, Init);
    // We violate SSA form by inserting OpVariable and still having a gMIR
    // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by erasing
    // the duplicated definition.
    if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
      GR.invalidateMachineInstr(MI);
      MI->removeFromParent();
    }
    return Res;
  }
  case Intrinsic::spv_undef: {
    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
                   .addDef(ResVReg)
                   .addUse(GR.getSPIRVTypeID(ResType));
    return MIB.constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_const_composite: {
    // If no values are attached, the composite is null constant.
    bool IsNull = I.getNumExplicitDefs() + 1 == I.getNumExplicitOperands();
    SmallVector<Register> CompositeArgs;
    MRI->setRegClass(ResVReg, GR.getRegClass(ResType));

    // skip type MD node we already used when generated assign.type for this
    if (!IsNull) {
      if (!wrapIntoSpecConstantOp(I, CompositeArgs))
        return false;
      MachineIRBuilder MIR(I);
      SmallVector<MachineInstr *, 4> Instructions = createContinuedInstructions(
          MIR, SPIRV::OpConstantComposite, 3,
          SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
          GR.getSPIRVTypeID(ResType));
      for (auto *Instr : Instructions) {
        Instr->setDebugLoc(I.getDebugLoc());
        if (!constrainSelectedInstRegOperands(*Instr, TII, TRI, RBI))
          return false;
      }
      return true;
    } else {
      auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
                     .addDef(ResVReg)
                     .addUse(GR.getSPIRVTypeID(ResType));
      return MIB.constrainAllUses(TII, TRI, RBI);
    }
  }
  case Intrinsic::spv_assign_name: {
    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpName));
    MIB.addUse(I.getOperand(I.getNumExplicitDefs() + 1).getReg());
    for (unsigned i = I.getNumExplicitDefs() + 2;
         i < I.getNumExplicitOperands(); ++i) {
      MIB.addImm(I.getOperand(i).getImm());
    }
    return MIB.constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_switch: {
    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSwitch));
    for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
      if (I.getOperand(i).isReg())
        MIB.addReg(I.getOperand(i).getReg());
      else if (I.getOperand(i).isCImm())
        addNumImm(I.getOperand(i).getCImm()->getValue(), MIB);
      else if (I.getOperand(i).isMBB())
        MIB.addMBB(I.getOperand(i).getMBB());
      else
        llvm_unreachable("Unexpected OpSwitch operand");
    }
    return MIB.constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_loop_merge: {
    auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpLoopMerge));
    for (unsigned i = 1; i < I.getNumExplicitOperands(); ++i) {
      if (I.getOperand(i).isMBB())
        MIB.addMBB(I.getOperand(i).getMBB());
      else
        MIB.addImm(foldImm(I.getOperand(i), MRI));
    }
    return MIB.constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_selection_merge: {
    auto MIB =
        BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSelectionMerge));
    assert(I.getOperand(1).isMBB() &&
           "operand 1 to spv_selection_merge must be a basic block");
    MIB.addMBB(I.getOperand(1).getMBB());
    MIB.addImm(getSelectionOperandForImm(I.getOperand(2).getImm()));
    return MIB.constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_cmpxchg:
    return selectAtomicCmpXchg(ResVReg, ResType, I);
  case Intrinsic::spv_unreachable:
    return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUnreachable))
        .constrainAllUses(TII, TRI, RBI);
  case Intrinsic::spv_alloca:
    return selectFrameIndex(ResVReg, ResType, I);
  case Intrinsic::spv_alloca_array:
    return selectAllocaArray(ResVReg, ResType, I);
  case Intrinsic::spv_assume:
    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
      return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAssumeTrueKHR))
          .addUse(I.getOperand(1).getReg())
          .constrainAllUses(TII, TRI, RBI);
    break;
  case Intrinsic::spv_expect:
    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
      return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExpectKHR))
          .addDef(ResVReg)
          .addUse(GR.getSPIRVTypeID(ResType))
          .addUse(I.getOperand(2).getReg())
          .addUse(I.getOperand(3).getReg())
          .constrainAllUses(TII, TRI, RBI);
    break;
  case Intrinsic::arithmetic_fence:
    if (STI.canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence))
      return BuildMI(BB, I, I.getDebugLoc(),
                     TII.get(SPIRV::OpArithmeticFenceEXT))
          .addDef(ResVReg)
          .addUse(GR.getSPIRVTypeID(ResType))
          .addUse(I.getOperand(2).getReg())
          .constrainAllUses(TII, TRI, RBI);
    else
      return BuildCOPY(ResVReg, I.getOperand(2).getReg(), I);
    break;
  case Intrinsic::spv_thread_id:
    // The HLSL SV_DispatchThreadID semantic is lowered to llvm.spv.thread.id
    // intrinsic in LLVM IR for SPIR-V backend.
    //
    // In SPIR-V backend, llvm.spv.thread.id is now correctly translated to a
    // `GlobalInvocationId` builtin variable
    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
                                  ResType, I);
  case Intrinsic::spv_thread_id_in_group:
    // The HLSL SV_GroupThreadId semantic is lowered to
    // llvm.spv.thread.id.in.group intrinsic in LLVM IR for SPIR-V backend.
    //
    // In SPIR-V backend, llvm.spv.thread.id.in.group is now correctly
    // translated to a `LocalInvocationId` builtin variable
    return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
                                  ResType, I);
  case Intrinsic::spv_group_id:
    // The HLSL SV_GroupId semantic is lowered to
    // llvm.spv.group.id intrinsic in LLVM IR for SPIR-V backend.
    //
    // In SPIR-V backend, llvm.spv.group.id is now translated to a `WorkgroupId`
    // builtin variable
    return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
                                  I);
  case Intrinsic::spv_flattened_thread_id_in_group:
    // The HLSL SV_GroupIndex semantic is lowered to
    // llvm.spv.flattened.thread.id.in.group() intrinsic in LLVM IR for SPIR-V
    // backend.
    //
    // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated to
    // a `LocalInvocationIndex` builtin variable
    return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
                              ResType, I);
  case Intrinsic::spv_workgroup_size:
    return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
                                  ResType, I);
  case Intrinsic::spv_global_size:
    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
                                  I);
  case Intrinsic::spv_global_offset:
    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
                                  ResType, I);
  case Intrinsic::spv_num_workgroups:
    return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
                                  ResType, I);
  case Intrinsic::spv_subgroup_size:
    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
                              I);
  case Intrinsic::spv_num_subgroups:
    return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
                              I);
  case Intrinsic::spv_subgroup_id:
    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType, I);
  case Intrinsic::spv_subgroup_local_invocation_id:
    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
                              ResVReg, ResType, I);
  case Intrinsic::spv_subgroup_max_size:
    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
                              I);
  case Intrinsic::spv_fdot:
    return selectFloatDot(ResVReg, ResType, I);
  case Intrinsic::spv_udot:
  case Intrinsic::spv_sdot:
    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
        STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
      return selectIntegerDot(ResVReg, ResType, I,
                              /*Signed=*/IID == Intrinsic::spv_sdot);
    return selectIntegerDotExpansion(ResVReg, ResType, I);
  case Intrinsic::spv_dot4add_i8packed:
    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
        STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
      return selectDot4AddPacked<true>(ResVReg, ResType, I);
    return selectDot4AddPackedExpansion<true>(ResVReg, ResType, I);
  case Intrinsic::spv_dot4add_u8packed:
    if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
        STI.isAtLeastSPIRVVer(VersionTuple(1, 6)))
      return selectDot4AddPacked<false>(ResVReg, ResType, I);
    return selectDot4AddPackedExpansion<false>(ResVReg, ResType, I);
  case Intrinsic::spv_all:
    return selectAll(ResVReg, ResType, I);
  case Intrinsic::spv_any:
    return selectAny(ResVReg, ResType, I);
  case Intrinsic::spv_cross:
    return selectExtInst(ResVReg, ResType, I, CL::cross, GL::Cross);
  case Intrinsic::spv_distance:
    return selectExtInst(ResVReg, ResType, I, CL::distance, GL::Distance);
  case Intrinsic::spv_lerp:
    return selectExtInst(ResVReg, ResType, I, CL::mix, GL::FMix);
  case Intrinsic::spv_length:
    return selectExtInst(ResVReg, ResType, I, CL::length, GL::Length);
  case Intrinsic::spv_degrees:
    return selectExtInst(ResVReg, ResType, I, CL::degrees, GL::Degrees);
  case Intrinsic::spv_faceforward:
    return selectExtInst(ResVReg, ResType, I, GL::FaceForward);
  case Intrinsic::spv_frac:
    return selectExtInst(ResVReg, ResType, I, CL::fract, GL::Fract);
  case Intrinsic::spv_normalize:
    return selectExtInst(ResVReg, ResType, I, CL::normalize, GL::Normalize);
  case Intrinsic::spv_reflect:
    return selectExtInst(ResVReg, ResType, I, GL::Reflect);
  case Intrinsic::spv_rsqrt:
    return selectExtInst(ResVReg, ResType, I, CL::rsqrt, GL::InverseSqrt);
  case Intrinsic::spv_sign:
    return selectSign(ResVReg, ResType, I);
  case Intrinsic::spv_smoothstep:
    return selectExtInst(ResVReg, ResType, I, CL::smoothstep, GL::SmoothStep);
  case Intrinsic::spv_firstbituhigh: // There is no CL equivalent of FindUMsb
    return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/false);
  case Intrinsic::spv_firstbitshigh: // There is no CL equivalent of FindSMsb
    return selectFirstBitHigh(ResVReg, ResType, I, /*IsSigned=*/true);
  case Intrinsic::spv_firstbitlow: // There is no CL equivlent of FindILsb
    return selectFirstBitLow(ResVReg, ResType, I);
  case Intrinsic::spv_group_memory_barrier_with_group_sync: {
    bool Result = true;
    auto MemSemConstant =
        buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent, I);
    Register MemSemReg = MemSemConstant.first;
    Result &= MemSemConstant.second;
    auto ScopeConstant = buildI32Constant(SPIRV::Scope::Workgroup, I);
    Register ScopeReg = ScopeConstant.first;
    Result &= ScopeConstant.second;
    MachineBasicBlock &BB = *I.getParent();
    return Result &&
           BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpControlBarrier))
               .addUse(ScopeReg)
               .addUse(ScopeReg)
               .addUse(MemSemReg)
               .constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_generic_cast_to_ptr_explicit: {
    Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 1).getReg();
    SPIRV::StorageClass::StorageClass ResSC =
        GR.getPointerStorageClass(ResType);
    if (!isGenericCastablePtr(ResSC))
      report_fatal_error("The target storage class is not castable from the "
                         "Generic storage class");
    return BuildMI(BB, I, I.getDebugLoc(),
                   TII.get(SPIRV::OpGenericCastToPtrExplicit))
        .addDef(ResVReg)
        .addUse(GR.getSPIRVTypeID(ResType))
        .addUse(PtrReg)
        .addImm(ResSC)
        .constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_lifetime_start:
  case Intrinsic::spv_lifetime_end: {
    unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
                                                       : SPIRV::OpLifetimeStop;
    int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();
    Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();
    if (Size == -1)
      Size = 0;
    return BuildMI(BB, I, I.getDebugLoc(), TII.get(Op))
        .addUse(PtrReg)
        .addImm(Size)
        .constrainAllUses(TII, TRI, RBI);
  }
  case Intrinsic::spv_saturate:
    return selectSaturate(ResVReg, ResType, I);
  case Intrinsic::spv_nclamp:
    return selectExtInst(ResVReg, ResType, I, CL::fclamp, GL::NClamp);
  case Intrinsic::spv_uclamp:
    return selectExtInst(ResVReg, ResType, I, CL::u_clamp, GL::UClamp);
  case Intrinsic::spv_sclamp:
    return selectExtInst(ResVReg, ResType, I, CL::s_clamp, GL::SClamp);
  case Intrinsic::spv_wave_active_countbits:
    return selectWaveActiveCountBits(ResVReg, ResType, I);
  case Intrinsic::spv_wave_all:
    return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAll);
  case Intrinsic::spv_wave_any:
    return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformAny);
  case Intrinsic::spv_wave_is_first_lane:
    return selectWaveOpInst(ResVReg, ResType, I, SPIRV::OpGroupNonUniformElect);
  case Intrinsic::spv_wave_reduce_umax:
    return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ true);
  case Intrinsic::spv_wave_reduce_max:
    return selectWaveReduceMax(ResVReg, ResType, I, /*IsUnsigned*/ false);
  case Intrinsic::spv_wave_reduce_sum:
    return selectWaveReduceSum(ResVReg, ResType, I);
  case Intrinsic::spv_wave_readlane:
    return selectWaveOpInst(ResVReg, ResType, I,
                            SPIRV::OpGroupNonUniformShuffle);
  case Intrinsic::spv_step:
    return selectExtInst(ResVReg, ResType, I, CL::step, GL::Step);
  case Intrinsic::spv_radians:
    return selectExtInst(ResVReg, ResType, I, CL::radians, GL::Radians);
  // Discard intrinsics which we do not expect to actually represent code after
  // lowering or intrinsics which are not implemented but should not crash when
  // found in a customer's LLVM IR input.
  case Intrinsic::instrprof_increment:
  case Intrinsic::instrprof_increment_step:
  case Intrinsic::instrprof_value_profile:
    break;
  // Discard internal intrinsics.
  case Intrinsic::spv_value_md:
    break;
  case Intrinsic::spv_resource_handlefrombinding: {
    return selectHandleFromBinding(ResVReg, ResType, I);
  }
  case Intrinsic::spv_resource_store_typedbuffer: {
    return selectImageWriteIntrinsic(I);
  }
  case Intrinsic::spv_resource_load_typedbuffer: {
    return selectReadImageIntrinsic(ResVReg, ResType, I);
  }
  case Intrinsic::spv_resource_getpointer: {
    return selectResourceGetPointer(ResVReg, ResType, I);
  }
  case Intrinsic::spv_discard: {
    return selectDiscard(ResVReg, ResType, I);
  }
  default: {
    std::string DiagMsg;
    raw_string_ostream OS(DiagMsg);
    I.print(OS);
    DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;
    report_fatal_error(DiagMsg.c_str(), false);
  }
  }
  return true;
}

bool SPIRVInstructionSelector::selectHandleFromBinding(Register &ResVReg,
                                                       const SPIRVType *ResType,
                                                       MachineInstr &I) const {
  // The images need to be loaded in the same basic block as their use. We defer
  // loading the image to the intrinsic that uses it.
  if (ResType->getOpcode() == SPIRV::OpTypeImage)
    return true;

  return loadHandleBeforePosition(ResVReg, GR.getSPIRVTypeForVReg(ResVReg),
                                  *cast<GIntrinsic>(&I), I);
}

bool SPIRVInstructionSelector::selectReadImageIntrinsic(
    Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {

  // If the load of the image is in a different basic block, then
  // this will generate invalid code. A proper solution is to move
  // the OpLoad from selectHandleFromBinding here. However, to do
  // that we will need to change the return type of the intrinsic.
  // We will do that when we can, but for now trying to move forward with other
  // issues.
  Register ImageReg = I.getOperand(2).getReg();
  auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
  Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
  if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
                                *ImageDef, I)) {
    return false;
  }

  Register IdxReg = I.getOperand(3).getReg();
  DebugLoc Loc = I.getDebugLoc();
  MachineInstr &Pos = I;

  return generateImageRead(ResVReg, ResType, NewImageReg, IdxReg, Loc, Pos);
}

bool SPIRVInstructionSelector::generateImageRead(Register &ResVReg,
                                                 const SPIRVType *ResType,
                                                 Register ImageReg,
                                                 Register IdxReg, DebugLoc Loc,
                                                 MachineInstr &Pos) const {
  SPIRVType *ImageType = GR.getSPIRVTypeForVReg(ImageReg);
  assert(ImageType && ImageType->getOpcode() == SPIRV::OpTypeImage &&
         "ImageReg is not an image type.");
  bool IsSignedInteger =
      sampledTypeIsSignedInteger(GR.getTypeForSPIRVType(ImageType));

  uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
  if (ResultSize == 4) {
    auto BMI = BuildMI(*Pos.getParent(), Pos, Loc, TII.get(SPIRV::OpImageRead))
                   .addDef(ResVReg)
                   .addUse(GR.getSPIRVTypeID(ResType))
                   .addUse(ImageReg)
                   .addUse(IdxReg);

    if (IsSignedInteger)
      BMI.addImm(0x1000); // SignExtend
    return BMI.constrainAllUses(TII, TRI, RBI);
  }

  SPIRVType *ReadType = widenTypeToVec4(ResType, Pos);
  Register ReadReg = MRI->createVirtualRegister(GR.getRegClass(ReadType));
  auto BMI = BuildMI(*Pos.getParent(), Pos, Loc, TII.get(SPIRV::OpImageRead))
                 .addDef(ReadReg)
                 .addUse(GR.getSPIRVTypeID(ReadType))
                 .addUse(ImageReg)
                 .addUse(IdxReg);
  if (IsSignedInteger)
    BMI.addImm(0x1000); // SignExtend
  bool Succeed = BMI.constrainAllUses(TII, TRI, RBI);
  if (!Succeed)
    return false;

  if (ResultSize == 1) {
    return BuildMI(*Pos.getParent(), Pos, Loc,
                   TII.get(SPIRV::OpCompositeExtract))
        .addDef(ResVReg)
        .addUse(GR.getSPIRVTypeID(ResType))
        .addUse(ReadReg)
        .addImm(0)
        .constrainAllUses(TII, TRI, RBI);
  }
  return extractSubvector(ResVReg, ResType, ReadReg, Pos);
}

bool SPIRVInstructionSelector::selectResourceGetPointer(
    Register &ResVReg, const SPIRVType *ResType, MachineInstr &I) const {
  Register ResourcePtr = I.getOperand(2).getReg();
  SPIRVType *RegType = GR.getSPIRVTypeForVReg(ResourcePtr, I.getMF());
  if (RegType->getOpcode() == SPIRV::OpTypeImage) {
    // For texel buffers, the index into the image is part of the OpImageRead or
    // OpImageWrite instructions. So we will do nothing in this case. This
    // intrinsic will be combined with the load or store when selecting the load
    // or store.
    return true;
  }

  assert(ResType->getOpcode() == SPIRV::OpTypePointer);
  MachineIRBuilder MIRBuilder(I);

  Register IndexReg = I.getOperand(3).getReg();
  Register ZeroReg =
      buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
  return BuildMI(*I.getParent(), I, I.getDebugLoc(),
                 TII.get(SPIRV::OpAccessChain))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(ResourcePtr)
      .addUse(ZeroReg)
      .addUse(IndexReg)
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::extractSubvector(
    Register &ResVReg, const SPIRVType *ResType, Register &ReadReg,
    MachineInstr &InsertionPoint) const {
  SPIRVType *InputType = GR.getResultType(ReadReg);
  [[maybe_unused]] uint64_t InputSize =
      GR.getScalarOrVectorComponentCount(InputType);
  uint64_t ResultSize = GR.getScalarOrVectorComponentCount(ResType);
  assert(InputSize > 1 && "The input must be a vector.");
  assert(ResultSize > 1 && "The result must be a vector.");
  assert(ResultSize < InputSize &&
         "Cannot extract more element than there are in the input.");
  SmallVector<Register> ComponentRegisters;
  SPIRVType *ScalarType = GR.getScalarOrVectorComponentType(ResType);
  const TargetRegisterClass *ScalarRegClass = GR.getRegClass(ScalarType);
  for (uint64_t I = 0; I < ResultSize; I++) {
    Register ComponentReg = MRI->createVirtualRegister(ScalarRegClass);
    bool Succeed = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
                           InsertionPoint.getDebugLoc(),
                           TII.get(SPIRV::OpCompositeExtract))
                       .addDef(ComponentReg)
                       .addUse(ScalarType->getOperand(0).getReg())
                       .addUse(ReadReg)
                       .addImm(I)
                       .constrainAllUses(TII, TRI, RBI);
    if (!Succeed)
      return false;
    ComponentRegisters.emplace_back(ComponentReg);
  }

  MachineInstrBuilder MIB = BuildMI(*InsertionPoint.getParent(), InsertionPoint,
                                    InsertionPoint.getDebugLoc(),
                                    TII.get(SPIRV::OpCompositeConstruct))
                                .addDef(ResVReg)
                                .addUse(GR.getSPIRVTypeID(ResType));

  for (Register ComponentReg : ComponentRegisters)
    MIB.addUse(ComponentReg);
  return MIB.constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
    MachineInstr &I) const {
  // If the load of the image is in a different basic block, then
  // this will generate invalid code. A proper solution is to move
  // the OpLoad from selectHandleFromBinding here. However, to do
  // that we will need to change the return type of the intrinsic.
  // We will do that when we can, but for now trying to move forward with other
  // issues.
  Register ImageReg = I.getOperand(1).getReg();
  auto *ImageDef = cast<GIntrinsic>(getVRegDef(*MRI, ImageReg));
  Register NewImageReg = MRI->createVirtualRegister(MRI->getRegClass(ImageReg));
  if (!loadHandleBeforePosition(NewImageReg, GR.getSPIRVTypeForVReg(ImageReg),
                                *ImageDef, I)) {
    return false;
  }

  Register CoordinateReg = I.getOperand(2).getReg();
  Register DataReg = I.getOperand(3).getReg();
  assert(GR.getResultType(DataReg)->getOpcode() == SPIRV::OpTypeVector);
  assert(GR.getScalarOrVectorComponentCount(GR.getResultType(DataReg)) == 4);
  return BuildMI(*I.getParent(), I, I.getDebugLoc(),
                 TII.get(SPIRV::OpImageWrite))
      .addUse(NewImageReg)
      .addUse(CoordinateReg)
      .addUse(DataReg)
      .constrainAllUses(TII, TRI, RBI);
}

Register SPIRVInstructionSelector::buildPointerToResource(
    const SPIRVType *SpirvResType, SPIRV::StorageClass::StorageClass SC,
    uint32_t Set, uint32_t Binding, uint32_t ArraySize, Register IndexReg,
    bool IsNonUniform, StringRef Name, MachineIRBuilder MIRBuilder) const {
  const Type *ResType = GR.getTypeForSPIRVType(SpirvResType);
  if (ArraySize == 1) {
    SPIRVType *PtrType =
        GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);
    assert(GR.getPointeeType(PtrType) == SpirvResType &&
           "SpirvResType did not have an explicit layout.");
    return GR.getOrCreateGlobalVariableWithBinding(PtrType, Set, Binding, Name,
                                                   MIRBuilder);
  }

  const Type *VarType = ArrayType::get(const_cast<Type *>(ResType), ArraySize);
  SPIRVType *VarPointerType =
      GR.getOrCreateSPIRVPointerType(VarType, MIRBuilder, SC);
  Register VarReg = GR.getOrCreateGlobalVariableWithBinding(
      VarPointerType, Set, Binding, Name, MIRBuilder);

  SPIRVType *ResPointerType =
      GR.getOrCreateSPIRVPointerType(ResType, MIRBuilder, SC);

  Register AcReg = MRI->createVirtualRegister(GR.getRegClass(ResPointerType));
  if (IsNonUniform) {
    // It is unclear which value needs to be marked an non-uniform, so both
    // the index and the access changed are decorated as non-uniform.
    buildOpDecorate(IndexReg, MIRBuilder, SPIRV::Decoration::NonUniformEXT, {});
    buildOpDecorate(AcReg, MIRBuilder, SPIRV::Decoration::NonUniformEXT, {});
  }

  MIRBuilder.buildInstr(SPIRV::OpAccessChain)
      .addDef(AcReg)
      .addUse(GR.getSPIRVTypeID(ResPointerType))
      .addUse(VarReg)
      .addUse(IndexReg);

  return AcReg;
}

bool SPIRVInstructionSelector::selectFirstBitSet16(
    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
    unsigned ExtendOpcode, unsigned BitSetOpcode) const {
  Register ExtReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
  bool Result = selectOpWithSrcs(ExtReg, ResType, I, {I.getOperand(2).getReg()},
                                 ExtendOpcode);

  return Result &&
         selectFirstBitSet32(ResVReg, ResType, I, ExtReg, BitSetOpcode);
}

bool SPIRVInstructionSelector::selectFirstBitSet32(
    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
    Register SrcReg, unsigned BitSetOpcode) const {
  return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
      .addDef(ResVReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
      .addImm(BitSetOpcode)
      .addUse(SrcReg)
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
    Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {

  // SPIR-V allow vectors of size 2,3,4 only. Calling with a larger vectors
  // requires creating a param register and return register with an invalid
  // vector size. If that is resolved, then this function can be used for
  // vectors of any component size.
  unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
  assert(ComponentCount < 5 && "Vec 5+ will generate invalid SPIR-V ops");

  MachineIRBuilder MIRBuilder(I);
  SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);
  SPIRVType *I64Type = GR.getOrCreateSPIRVIntegerType(64, MIRBuilder);
  SPIRVType *I64x2Type =
      GR.getOrCreateSPIRVVectorType(I64Type, 2, MIRBuilder, false);
  SPIRVType *Vec2ResType =
      GR.getOrCreateSPIRVVectorType(BaseType, 2, MIRBuilder, false);

  std::vector<Register> PartialRegs;

  // Loops 0, 2, 4, ... but stops one loop early when ComponentCount is odd
  unsigned CurrentComponent = 0;
  for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
    // This register holds the firstbitX result for each of the i64x2 vectors
    // extracted from SrcReg
    Register BitSetResult =
        MRI->createVirtualRegister(GR.getRegClass(I64x2Type));

    auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                       TII.get(SPIRV::OpVectorShuffle))
                   .addDef(BitSetResult)
                   .addUse(GR.getSPIRVTypeID(I64x2Type))
                   .addUse(SrcReg)
                   .addUse(SrcReg)
                   .addImm(CurrentComponent)
                   .addImm(CurrentComponent + 1);

    if (!MIB.constrainAllUses(TII, TRI, RBI))
      return false;

    Register SubVecBitSetReg =
        MRI->createVirtualRegister(GR.getRegClass(Vec2ResType));

    if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType, I, BitSetResult,
                             BitSetOpcode, SwapPrimarySide))
      return false;

    PartialRegs.push_back(SubVecBitSetReg);
  }

  // On odd component counts we need to handle one more component
  if (CurrentComponent != ComponentCount) {
    bool ZeroAsNull = !STI.isShader();
    Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I64Type));
    Register ConstIntLastIdx = GR.getOrCreateConstInt(
        ComponentCount - 1, I, BaseType, TII, ZeroAsNull);

    if (!selectOpWithSrcs(FinalElemReg, I64Type, I, {SrcReg, ConstIntLastIdx},
                          SPIRV::OpVectorExtractDynamic))
      return false;

    Register FinalElemBitSetReg =
        MRI->createVirtualRegister(GR.getRegClass(BaseType));

    if (!selectFirstBitSet64(FinalElemBitSetReg, BaseType, I, FinalElemReg,
                             BitSetOpcode, SwapPrimarySide))
      return false;

    PartialRegs.push_back(FinalElemBitSetReg);
  }

  // Join all the resulting registers back into the return type in order
  // (ie i32x2, i32x2, i32x1 -> i32x5)
  return selectOpWithSrcs(ResVReg, ResType, I, PartialRegs,
                          SPIRV::OpCompositeConstruct);
}

bool SPIRVInstructionSelector::selectFirstBitSet64(
    Register ResVReg, const SPIRVType *ResType, MachineInstr &I,
    Register SrcReg, unsigned BitSetOpcode, bool SwapPrimarySide) const {
  unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
  SPIRVType *BaseType = GR.retrieveScalarOrVectorIntType(ResType);
  bool ZeroAsNull = !STI.isShader();
  Register ConstIntZero =
      GR.getOrCreateConstInt(0, I, BaseType, TII, ZeroAsNull);
  Register ConstIntOne =
      GR.getOrCreateConstInt(1, I, BaseType, TII, ZeroAsNull);

  // SPIRV doesn't support vectors with more than 4 components. Since the
  // algoritm below converts i64 -> i32x2 and i64x4 -> i32x8 it can only
  // operate on vectors with 2 or less components. When largers vectors are
  // seen. Split them, recurse, then recombine them.
  if (ComponentCount > 2) {
    return selectFirstBitSet64Overflow(ResVReg, ResType, I, SrcReg,
                                       BitSetOpcode, SwapPrimarySide);
  }

  // 1. Split int64 into 2 pieces using a bitcast
  MachineIRBuilder MIRBuilder(I);
  SPIRVType *PostCastType = GR.getOrCreateSPIRVVectorType(
      BaseType, 2 * ComponentCount, MIRBuilder, false);
  Register BitcastReg =
      MRI->createVirtualRegister(GR.getRegClass(PostCastType));

  if (!selectOpWithSrcs(BitcastReg, PostCastType, I, {SrcReg},
                        SPIRV::OpBitcast))
    return false;

  // 2. Find the first set bit from the primary side for all the pieces in #1
  Register FBSReg = MRI->createVirtualRegister(GR.getRegClass(PostCastType));
  if (!selectFirstBitSet32(FBSReg, PostCastType, I, BitcastReg, BitSetOpcode))
    return false;

  // 3. Split result vector into high bits and low bits
  Register HighReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
  Register LowReg = MRI->createVirtualRegister(GR.getRegClass(ResType));

  bool IsScalarRes = ResType->getOpcode() != SPIRV::OpTypeVector;
  if (IsScalarRes) {
    // if scalar do a vector extract
    if (!selectOpWithSrcs(HighReg, ResType, I, {FBSReg, ConstIntZero},
                          SPIRV::OpVectorExtractDynamic))
      return false;
    if (!selectOpWithSrcs(LowReg, ResType, I, {FBSReg, ConstIntOne},
                          SPIRV::OpVectorExtractDynamic))
      return false;
  } else {
    // if vector do a shufflevector
    auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                       TII.get(SPIRV::OpVectorShuffle))
                   .addDef(HighReg)
                   .addUse(GR.getSPIRVTypeID(ResType))
                   .addUse(FBSReg)
                   // Per the spec, repeat the vector if only one vec is needed
                   .addUse(FBSReg);

    // high bits are stored in even indexes. Extract them from FBSReg
    for (unsigned J = 0; J < ComponentCount * 2; J += 2) {
      MIB.addImm(J);
    }

    if (!MIB.constrainAllUses(TII, TRI, RBI))
      return false;

    MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                  TII.get(SPIRV::OpVectorShuffle))
              .addDef(LowReg)
              .addUse(GR.getSPIRVTypeID(ResType))
              .addUse(FBSReg)
              // Per the spec, repeat the vector if only one vec is needed
              .addUse(FBSReg);

    // low bits are stored in odd indexes. Extract them from FBSReg
    for (unsigned J = 1; J < ComponentCount * 2; J += 2) {
      MIB.addImm(J);
    }
    if (!MIB.constrainAllUses(TII, TRI, RBI))
      return false;
  }

  // 4. Check the result. When primary bits == -1 use secondary, otherwise use
  // primary
  SPIRVType *BoolType = GR.getOrCreateSPIRVBoolType(I, TII);
  Register NegOneReg;
  Register Reg0;
  Register Reg32;
  unsigned SelectOp;
  unsigned AddOp;

  if (IsScalarRes) {
    NegOneReg =
        GR.getOrCreateConstInt((unsigned)-1, I, ResType, TII, ZeroAsNull);
    Reg0 = GR.getOrCreateConstInt(0, I, ResType, TII, ZeroAsNull);
    Reg32 = GR.getOrCreateConstInt(32, I, ResType, TII, ZeroAsNull);
    SelectOp = SPIRV::OpSelectSISCond;
    AddOp = SPIRV::OpIAddS;
  } else {
    BoolType = GR.getOrCreateSPIRVVectorType(BoolType, ComponentCount,
                                             MIRBuilder, false);
    NegOneReg =
        GR.getOrCreateConstVector((unsigned)-1, I, ResType, TII, ZeroAsNull);
    Reg0 = GR.getOrCreateConstVector(0, I, ResType, TII, ZeroAsNull);
    Reg32 = GR.getOrCreateConstVector(32, I, ResType, TII, ZeroAsNull);
    SelectOp = SPIRV::OpSelectVIVCond;
    AddOp = SPIRV::OpIAddV;
  }

  Register PrimaryReg = HighReg;
  Register SecondaryReg = LowReg;
  Register PrimaryShiftReg = Reg32;
  Register SecondaryShiftReg = Reg0;

  // By default the emitted opcodes check for the set bit from the MSB side.
  // Setting SwapPrimarySide checks the set bit from the LSB side
  if (SwapPrimarySide) {
    PrimaryReg = LowReg;
    SecondaryReg = HighReg;
    PrimaryShiftReg = Reg0;
    SecondaryShiftReg = Reg32;
  }

  // Check if the primary bits are == -1
  Register BReg = MRI->createVirtualRegister(GR.getRegClass(BoolType));
  if (!selectOpWithSrcs(BReg, BoolType, I, {PrimaryReg, NegOneReg},
                        SPIRV::OpIEqual))
    return false;

  // Select secondary bits if true in BReg, otherwise primary bits
  Register TmpReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
  if (!selectOpWithSrcs(TmpReg, ResType, I, {BReg, SecondaryReg, PrimaryReg},
                        SelectOp))
    return false;

  // 5. Add 32 when high bits are used, otherwise 0 for low bits
  Register ValReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
  if (!selectOpWithSrcs(ValReg, ResType, I,
                        {BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
    return false;

  return selectOpWithSrcs(ResVReg, ResType, I, {ValReg, TmpReg}, AddOp);
}

bool SPIRVInstructionSelector::selectFirstBitHigh(Register ResVReg,
                                                  const SPIRVType *ResType,
                                                  MachineInstr &I,
                                                  bool IsSigned) const {
  // FindUMsb and FindSMsb intrinsics only support 32 bit integers
  Register OpReg = I.getOperand(2).getReg();
  SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
  // zero or sign extend
  unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
  unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;

  switch (GR.getScalarOrVectorBitWidth(OpType)) {
  case 16:
    return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
  case 32:
    return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
  case 64:
    return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
                               /*SwapPrimarySide=*/false);
  default:
    report_fatal_error(
        "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
  }
}

bool SPIRVInstructionSelector::selectFirstBitLow(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
  // FindILsb intrinsic only supports 32 bit integers
  Register OpReg = I.getOperand(2).getReg();
  SPIRVType *OpType = GR.getSPIRVTypeForVReg(OpReg);
  // OpUConvert treats the operand bits as an unsigned i16 and zero extends it
  // to an unsigned i32. As this leaves all the least significant bits unchanged
  // so the first set bit from the LSB side doesn't change.
  unsigned ExtendOpcode = SPIRV::OpUConvert;
  unsigned BitSetOpcode = GL::FindILsb;

  switch (GR.getScalarOrVectorBitWidth(OpType)) {
  case 16:
    return selectFirstBitSet16(ResVReg, ResType, I, ExtendOpcode, BitSetOpcode);
  case 32:
    return selectFirstBitSet32(ResVReg, ResType, I, OpReg, BitSetOpcode);
  case 64:
    return selectFirstBitSet64(ResVReg, ResType, I, OpReg, BitSetOpcode,
                               /*SwapPrimarySide=*/true);
  default:
    report_fatal_error("spv_firstbitlow only supports 16,32,64 bits.");
  }
}

bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
  // there was an allocation size parameter to the allocation instruction
  // that is not 1
  MachineBasicBlock &BB = *I.getParent();
  bool Res = BuildMI(BB, I, I.getDebugLoc(),
                     TII.get(SPIRV::OpVariableLengthArrayINTEL))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 .addUse(I.getOperand(2).getReg())
                 .constrainAllUses(TII, TRI, RBI);
  if (!STI.isShader()) {
    unsigned Alignment = I.getOperand(3).getImm();
    buildOpDecorate(ResVReg, I, TII, SPIRV::Decoration::Alignment, {Alignment});
  }
  return Res;
}

bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
  // Change order of instructions if needed: all OpVariable instructions in a
  // function must be the first instructions in the first block
  auto It = getOpVariableMBBIt(I);
  bool Res = BuildMI(*It->getParent(), It, It->getDebugLoc(),
                     TII.get(SPIRV::OpVariable))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))
                 .constrainAllUses(TII, TRI, RBI);
  if (!STI.isShader()) {
    unsigned Alignment = I.getOperand(2).getImm();
    buildOpDecorate(ResVReg, *It, TII, SPIRV::Decoration::Alignment,
                    {Alignment});
  }
  return Res;
}

bool SPIRVInstructionSelector::selectBranch(MachineInstr &I) const {
  // InstructionSelector walks backwards through the instructions. We can use
  // both a G_BR and a G_BRCOND to create an OpBranchConditional. We hit G_BR
  // first, so can generate an OpBranchConditional here. If there is no
  // G_BRCOND, we just use OpBranch for a regular unconditional branch.
  const MachineInstr *PrevI = I.getPrevNode();
  MachineBasicBlock &MBB = *I.getParent();
  if (PrevI != nullptr && PrevI->getOpcode() == TargetOpcode::G_BRCOND) {
    return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
        .addUse(PrevI->getOperand(0).getReg())
        .addMBB(PrevI->getOperand(1).getMBB())
        .addMBB(I.getOperand(0).getMBB())
        .constrainAllUses(TII, TRI, RBI);
  }
  return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranch))
      .addMBB(I.getOperand(0).getMBB())
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &I) const {
  // InstructionSelector walks backwards through the instructions. For an
  // explicit conditional branch with no fallthrough, we use both a G_BR and a
  // G_BRCOND to create an OpBranchConditional. We should hit G_BR first, and
  // generate the OpBranchConditional in selectBranch above.
  //
  // If an OpBranchConditional has been generated, we simply return, as the work
  // is alread done. If there is no OpBranchConditional, LLVM must be relying on
  // implicit fallthrough to the next basic block, so we need to create an
  // OpBranchConditional with an explicit "false" argument pointing to the next
  // basic block that LLVM would fall through to.
  const MachineInstr *NextI = I.getNextNode();
  // Check if this has already been successfully selected.
  if (NextI != nullptr && NextI->getOpcode() == SPIRV::OpBranchConditional)
    return true;
  // Must be relying on implicit block fallthrough, so generate an
  // OpBranchConditional with the "next" basic block as the "false" target.
  MachineBasicBlock &MBB = *I.getParent();
  unsigned NextMBBNum = MBB.getNextNode()->getNumber();
  MachineBasicBlock *NextMBB = I.getMF()->getBlockNumbered(NextMBBNum);
  return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
      .addUse(I.getOperand(0).getReg())
      .addMBB(I.getOperand(1).getMBB())
      .addMBB(NextMBB)
      .constrainAllUses(TII, TRI, RBI);
}

bool SPIRVInstructionSelector::selectPhi(Register ResVReg,
                                         const SPIRVType *ResType,
                                         MachineInstr &I) const {
  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpPhi))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType));
  const unsigned NumOps = I.getNumOperands();
  for (unsigned i = 1; i < NumOps; i += 2) {
    MIB.addUse(I.getOperand(i + 0).getReg());
    MIB.addMBB(I.getOperand(i + 1).getMBB());
  }
  bool Res = MIB.constrainAllUses(TII, TRI, RBI);
  MIB->setDesc(TII.get(TargetOpcode::PHI));
  MIB->removeOperand(1);
  return Res;
}

bool SPIRVInstructionSelector::selectGlobalValue(
    Register ResVReg, MachineInstr &I, const MachineInstr *Init) const {
  // FIXME: don't use MachineIRBuilder here, replace it with BuildMI.
  MachineIRBuilder MIRBuilder(I);
  const GlobalValue *GV = I.getOperand(1).getGlobal();
  Type *GVType = toTypedPointer(GR.getDeducedGlobalValueType(GV));

  std::string GlobalIdent;
  if (!GV->hasName()) {
    unsigned &ID = UnnamedGlobalIDs[GV];
    if (ID == 0)
      ID = UnnamedGlobalIDs.size();
    GlobalIdent = "__unnamed_" + Twine(ID).str();
  } else {
    GlobalIdent = GV->getName();
  }

  // Behaviour of functions as operands depends on availability of the
  // corresponding extension (SPV_INTEL_function_pointers):
  // - If there is an extension to operate with functions as operands:
  // We create a proper constant operand and evaluate a correct type for a
  // function pointer.
  // - Without the required extension:
  // We have functions as operands in tests with blocks of instruction e.g. in
  // transcoding/global_block.ll. These operands are not used and should be
  // substituted by zero constants. Their type is expected to be always
  // OpTypePointer Function %uchar.
  if (isa<Function>(GV)) {
    const Constant *ConstVal = GV;
    MachineBasicBlock &BB = *I.getParent();
    Register NewReg = GR.find(ConstVal, GR.CurMF);
    if (!NewReg.isValid()) {
      Register NewReg = ResVReg;
      const Function *GVFun =
          STI.canUseExtension(SPIRV::Extension::SPV_INTEL_function_pointers)
              ? dyn_cast<Function>(GV)
              : nullptr;
      SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(
          GVType, I,
          GVFun ? SPIRV::StorageClass::CodeSectionINTEL
                : addressSpaceToStorageClass(GV->getAddressSpace(), STI));
      if (GVFun) {
        // References to a function via function pointers generate virtual
        // registers without a definition. We will resolve it later, during
        // module analysis stage.
        Register ResTypeReg = GR.getSPIRVTypeID(ResType);
        MachineRegisterInfo *MRI = MIRBuilder.getMRI();
        Register FuncVReg =
            MRI->createGenericVirtualRegister(GR.getRegType(ResType));
        MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);
        MachineInstrBuilder MIB1 =
            BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
                .addDef(FuncVReg)
                .addUse(ResTypeReg);
        MachineInstrBuilder MIB2 =
            BuildMI(BB, I, I.getDebugLoc(),
                    TII.get(SPIRV::OpConstantFunctionPointerINTEL))
                .addDef(NewReg)
                .addUse(ResTypeReg)
                .addUse(FuncVReg);
        GR.add(ConstVal, MIB2);
        // mapping the function pointer to the used Function
        GR.recordFunctionPointer(&MIB2.getInstr()->getOperand(2), GVFun);
        return MIB1.constrainAllUses(TII, TRI, RBI) &&
               MIB2.constrainAllUses(TII, TRI, RBI);
      }
      MachineInstrBuilder MIB3 =
          BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
              .addDef(NewReg)
              .addUse(GR.getSPIRVTypeID(ResType));
      GR.add(ConstVal, MIB3);
      return MIB3.constrainAllUses(TII, TRI, RBI);
    }
    assert(NewReg != ResVReg);
    return BuildCOPY(ResVReg, NewReg, I);
  }
  auto GlobalVar = cast<GlobalVariable>(GV);
  assert(GlobalVar->getName() != "llvm.global.annotations");

  // Skip empty declaration for GVs with initializers till we get the decl with
  // passed initializer.
  if (hasInitializer(GlobalVar) && !Init)
    return true;

  bool HasLnkTy = !GV->hasInternalLinkage() && !GV->hasPrivateLinkage() &&
                  !GV->hasHiddenVisibility();
  SPIRV::LinkageType::LinkageType LnkType =
      GV->isDeclarationForLinker()
          ? SPIRV::LinkageType::Import
          : (GV->hasLinkOnceODRLinkage() &&
                     STI.canUseExtension(SPIRV::Extension::SPV_KHR_linkonce_odr)
                 ? SPIRV::LinkageType::LinkOnceODR
                 : SPIRV::LinkageType::Export);

  const unsigned AddrSpace = GV->getAddressSpace();
  SPIRV::StorageClass::StorageClass StorageClass =
      addressSpaceToStorageClass(AddrSpace, STI);
  SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(GVType, I, StorageClass);
  Register Reg = GR.buildGlobalVariable(
      ResVReg, ResType, GlobalIdent, GV, StorageClass, Init,
      GlobalVar->isConstant(), HasLnkTy, LnkType, MIRBuilder, true);
  return Reg.isValid();
}

bool SPIRVInstructionSelector::selectLog10(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I) const {
  if (STI.canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
    return selectExtInst(ResVReg, ResType, I, CL::log10);
  }

  // There is no log10 instruction in the GLSL Extended Instruction set, so it
  // is implemented as:
  // log10(x) = log2(x) * (1 / log2(10))
  //          = log2(x) * 0.30103

  MachineIRBuilder MIRBuilder(I);
  MachineBasicBlock &BB = *I.getParent();

  // Build log2(x).
  Register VarReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
  bool Result =
      BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
          .addDef(VarReg)
          .addUse(GR.getSPIRVTypeID(ResType))
          .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::GLSL_std_450))
          .addImm(GL::Log2)
          .add(I.getOperand(1))
          .constrainAllUses(TII, TRI, RBI);

  // Build 0.30103.
  assert(ResType->getOpcode() == SPIRV::OpTypeVector ||
         ResType->getOpcode() == SPIRV::OpTypeFloat);
  // TODO: Add matrix implementation once supported by the HLSL frontend.
  const SPIRVType *SpirvScalarType =
      ResType->getOpcode() == SPIRV::OpTypeVector
          ? GR.getSPIRVTypeForVReg(ResType->getOperand(1).getReg())
          : ResType;
  Register ScaleReg =
      GR.buildConstantFP(APFloat(0.30103f), MIRBuilder, SpirvScalarType);

  // Multiply log2(x) by 0.30103 to get log10(x) result.
  auto Opcode = ResType->getOpcode() == SPIRV::OpTypeVector
                    ? SPIRV::OpVectorTimesScalar
                    : SPIRV::OpFMulS;
  return Result && BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
                       .addDef(ResVReg)
                       .addUse(GR.getSPIRVTypeID(ResType))
                       .addUse(VarReg)
                       .addUse(ScaleReg)
                       .constrainAllUses(TII, TRI, RBI);
}

// Generate the instructions to load 3-element vector builtin input
// IDs/Indices.
// Like: GlobalInvocationId, LocalInvocationId, etc....

bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
    SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
    const SPIRVType *ResType, MachineInstr &I) const {
  MachineIRBuilder MIRBuilder(I);
  const SPIRVType *Vec3Ty =
      GR.getOrCreateSPIRVVectorType(ResType, 3, MIRBuilder, false);
  const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
      Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);

  // Create new register for the input ID builtin variable.
  Register NewRegister =
      MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);
  MIRBuilder.getMRI()->setType(NewRegister, LLT::pointer(0, 64));
  GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());

  // Build global variable with the necessary decorations for the input ID
  // builtin variable.
  Register Variable = GR.buildGlobalVariable(
      NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
      SPIRV::StorageClass::Input, nullptr, true, false,
      SPIRV::LinkageType::Import, MIRBuilder, false);

  // Create new register for loading value.
  MachineRegisterInfo *MRI = MIRBuilder.getMRI();
  Register LoadedRegister = MRI->createVirtualRegister(&SPIRV::iIDRegClass);
  MIRBuilder.getMRI()->setType(LoadedRegister, LLT::pointer(0, 64));
  GR.assignSPIRVTypeToVReg(Vec3Ty, LoadedRegister, MIRBuilder.getMF());

  // Load v3uint value from the global variable.
  bool Result =
      BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
          .addDef(LoadedRegister)
          .addUse(GR.getSPIRVTypeID(Vec3Ty))
          .addUse(Variable);

  // Get the input ID index. Expecting operand is a constant immediate value,
  // wrapped in a type assignment.
  assert(I.getOperand(2).isReg());
  const uint32_t ThreadId = foldImm(I.getOperand(2), MRI);

  // Extract the input ID from the loaded vector value.
  MachineBasicBlock &BB = *I.getParent();
  auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 .addUse(LoadedRegister)
                 .addImm(ThreadId);
  return Result && MIB.constrainAllUses(TII, TRI, RBI);
}

// Generate the instructions to load 32-bit integer builtin input IDs/Indices.
// Like LocalInvocationIndex
bool SPIRVInstructionSelector::loadBuiltinInputID(
    SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
    const SPIRVType *ResType, MachineInstr &I) const {
  MachineIRBuilder MIRBuilder(I);
  const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
      ResType, MIRBuilder, SPIRV::StorageClass::Input);

  // Create new register for the input ID builtin variable.
  Register NewRegister =
      MIRBuilder.getMRI()->createVirtualRegister(GR.getRegClass(PtrType));
  MIRBuilder.getMRI()->setType(
      NewRegister,
      LLT::pointer(storageClassToAddressSpace(SPIRV::StorageClass::Input),
                   GR.getPointerSize()));
  GR.assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());

  // Build global variable with the necessary decorations for the input ID
  // builtin variable.
  Register Variable = GR.buildGlobalVariable(
      NewRegister, PtrType, getLinkStringForBuiltIn(BuiltInValue), nullptr,
      SPIRV::StorageClass::Input, nullptr, true, false,
      SPIRV::LinkageType::Import, MIRBuilder, false);

  // Load uint value from the global variable.
  auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
                 .addDef(ResVReg)
                 .addUse(GR.getSPIRVTypeID(ResType))
                 .addUse(Variable);

  return MIB.constrainAllUses(TII, TRI, RBI);
}

SPIRVType *SPIRVInstructionSelector::widenTypeToVec4(const SPIRVType *Type,
                                                     MachineInstr &I) const {
  MachineIRBuilder MIRBuilder(I);
  if (Type->getOpcode() != SPIRV::OpTypeVector)
    return GR.getOrCreateSPIRVVectorType(Type, 4, MIRBuilder, false);

  uint64_t VectorSize = Type->getOperand(2).getImm();
  if (VectorSize == 4)
    return Type;

  Register ScalarTypeReg = Type->getOperand(1).getReg();
  const SPIRVType *ScalarType = GR.getSPIRVTypeForVReg(ScalarTypeReg);
  return GR.getOrCreateSPIRVVectorType(ScalarType, 4, MIRBuilder, false);
}

bool SPIRVInstructionSelector::loadHandleBeforePosition(
    Register &HandleReg, const SPIRVType *ResType, GIntrinsic &HandleDef,
    MachineInstr &Pos) const {

  assert(HandleDef.getIntrinsicID() ==
         Intrinsic::spv_resource_handlefrombinding);
  uint32_t Set = foldImm(HandleDef.getOperand(2), MRI);
  uint32_t Binding = foldImm(HandleDef.getOperand(3), MRI);
  uint32_t ArraySize = foldImm(HandleDef.getOperand(4), MRI);
  Register IndexReg = HandleDef.getOperand(5).getReg();
  bool IsNonUniform = ArraySize > 1 && foldImm(HandleDef.getOperand(6), MRI);
  std::string Name =
      getStringValueFromReg(HandleDef.getOperand(7).getReg(), *MRI);

  bool IsStructuredBuffer = ResType->getOpcode() == SPIRV::OpTypePointer;
  MachineIRBuilder MIRBuilder(HandleDef);
  SPIRVType *VarType = ResType;
  SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;

  if (IsStructuredBuffer) {
    VarType = GR.getPointeeType(ResType);
    SC = GR.getPointerStorageClass(ResType);
  }

  Register VarReg =
      buildPointerToResource(VarType, SC, Set, Binding, ArraySize, IndexReg,
                             IsNonUniform, Name, MIRBuilder);

  if (IsNonUniform)
    buildOpDecorate(HandleReg, HandleDef, TII, SPIRV::Decoration::NonUniformEXT,
                    {});

  // The handle for the buffer is the pointer to the resource. For an image, the
  // handle is the image object. So images get an extra load.
  uint32_t LoadOpcode =
      IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
  GR.assignSPIRVTypeToVReg(ResType, HandleReg, *Pos.getMF());
  return BuildMI(*Pos.getParent(), Pos, HandleDef.getDebugLoc(),
                 TII.get(LoadOpcode))
      .addDef(HandleReg)
      .addUse(GR.getSPIRVTypeID(ResType))
      .addUse(VarReg)
      .constrainAllUses(TII, TRI, RBI);
}

namespace llvm {
InstructionSelector *
createSPIRVInstructionSelector(const SPIRVTargetMachine &TM,
                               const SPIRVSubtarget &Subtarget,
                               const RegisterBankInfo &RBI) {
  return new SPIRVInstructionSelector(TM, Subtarget, RBI);
}
} // namespace llvm