OCLToSPIRV.cpp

//===- OCLToSPIRV.cpp - Transform OCL to SPIR-V builtins --------*- C++ -*-===//
//
//                     The LLVM/SPIRV Translator
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
// Copyright (c) 2014 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a
// copy of this software and associated documentation files (the "Software"),
// to deal with the Software without restriction, including without limitation
// the rights to use, copy, modify, merge, publish, distribute, sublicense,
// and/or sell copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following conditions:
//
// Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimers in the documentation
// and/or other materials provided with the distribution.
// Neither the names of Advanced Micro Devices, Inc., nor the names of its
// contributors may be used to endorse or promote products derived from this
// Software without specific prior written permission.
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH
// THE SOFTWARE.
//
//===----------------------------------------------------------------------===//
//
// This file implements preprocessing of OpenCL C built-in functions into SPIR-V
// friendly IR form for further translation into SPIR-V
//
//===----------------------------------------------------------------------===//

#include "OCLToSPIRV.h"
#include "OCLTypeToSPIRV.h"
#include "SPIRVInternal.h"
#include "libSPIRV/SPIRVDebug.h"

#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringSwitch.h"
#include "llvm/Analysis/ValueTracking.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instruction.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Operator.h"
#include "llvm/IR/PatternMatch.h"
#include "llvm/IR/TypedPointerType.h"
#include "llvm/Support/Debug.h"

#include <algorithm>
#include <optional>
#include <regex>
#include <set>

#define DEBUG_TYPE "ocl-to-spv"

using namespace llvm;
using namespace PatternMatch;
using namespace SPIRV;
using namespace OCLUtil;

namespace SPIRV {

static unsigned getAddressSpaceFromType(const Type *Ty) {
  assert(Ty && "Can't deduce pointer AS");
  if (auto *TypedPtr = dyn_cast<TypedPointerType>(Ty))
    return TypedPtr->getAddressSpace();
  if (auto *Ptr = dyn_cast<PointerType>(Ty))
    return Ptr->getAddressSpace();
  llvm_unreachable("Can't deduce pointer AS");
}

// Performs an address space inference analysis.
static unsigned getAddressSpaceFromValue(const Value *Ptr) {
  assert(Ptr && "Can't deduce pointer AS");

  SmallPtrSet<const Value *, 8> Visited;
  SmallVector<const Value *, 8> Worklist;
  Worklist.push_back(Ptr);
  unsigned AS = SPIRAS_Generic;

  while (!Worklist.empty()) {
    const Value *Current = Worklist.pop_back_val();
    if (!Visited.insert(Current).second)
      continue;

    unsigned DeducedAS = getAddressSpaceFromType(Current->getType());
    if (DeducedAS != SPIRAS_Generic)
      return DeducedAS;
    AS = DeducedAS;

    // Find origins of the pointer and add to the worklist.
    if (auto *Op = dyn_cast<Operator>(Current)) {
      switch (Op->getOpcode()) {
      case Instruction::AddrSpaceCast:
      case Instruction::BitCast:
      case Instruction::GetElementPtr:
        Worklist.push_back(Op->getOperand(0));
        break;
      case Instruction::Select:
        Worklist.push_back(Op->getOperand(1));
        Worklist.push_back(Op->getOperand(2));
        break;
      case Instruction::PHI: {
        auto *Phi = cast<PHINode>(Op);
        for (Value *Incoming : Phi->incoming_values())
          Worklist.push_back(Incoming);
        break;
      }
      default:
        break;
      }
    }
  }

  return AS;
}

// Sets memory semantic mask of an atomic depending on a pointer argument
// address space.
static unsigned
getAtomicPointerMemorySemanticsMemoryMask(const Value *Ptr,
                                          const Type *RecordedType) {
  assert((Ptr && RecordedType) &&
         "Can't evaluate atomic builtin's memory semantic");
  unsigned AddrSpace = getAddressSpaceFromType(RecordedType);
  if (AddrSpace == SPIRAS_Generic)
    AddrSpace = getAddressSpaceFromValue(Ptr);

  switch (AddrSpace) {
  case SPIRAS_Global:
  case SPIRAS_GlobalDevice:
  case SPIRAS_GlobalHost:
    return MemorySemanticsCrossWorkgroupMemoryMask;
  case SPIRAS_Local:
    return MemorySemanticsWorkgroupMemoryMask;
  case SPIRAS_Generic:
    return MemorySemanticsCrossWorkgroupMemoryMask |
           MemorySemanticsWorkgroupMemoryMask;
  default:
    return MemorySemanticsMaskNone;
  }
}

static size_t getOCLCpp11AtomicMaxNumOps(StringRef Name) {
  return StringSwitch<size_t>(Name)
      .Cases({"load", "flag_test_and_set", "flag_clear"}, 3)
      .Cases({"store", "exchange"}, 4)
      .StartsWith("compare_exchange", 6)
      .StartsWith("fetch", 4)
      .Default(0);
}

static Type *getBlockStructType(Value *Parameter) {
  // In principle, this information should be passed to us from Clang via
  // an elementtype attribute. However, said attribute requires that the
  // function call be an intrinsic, which it is not. Instead, we rely on being
  // able to trace this to the declaration of a variable: OpenCL C specification
  // section 6.12.5 should guarantee that we can do this.
  Value *UnderlyingObject = Parameter->stripPointerCasts();
  Type *ParamType = nullptr;
  if (auto *GV = dyn_cast<GlobalValue>(UnderlyingObject))
    ParamType = GV->getValueType();
  else if (auto *Alloca = dyn_cast<AllocaInst>(UnderlyingObject))
    ParamType = Alloca->getAllocatedType();
  else
    llvm_unreachable("Blocks in OpenCL C must be traceable to allocation site");
  return ParamType;
}

/// Return one of the SPIR-V 1.4 SignExtend or ZeroExtend image operands
/// for a demangled function name, or 0 if the function does not return an
/// integer type (e.g. read_imagef).
static unsigned getImageSignZeroExt(StringRef DemangledName) {
  bool IsSigned = !DemangledName.ends_with("ui") && DemangledName.back() == 'i';
  bool IsUnsigned = DemangledName.ends_with("ui");

  if (IsSigned)
    return ImageOperandsMask::ImageOperandsSignExtendMask;
  if (IsUnsigned)
    return ImageOperandsMask::ImageOperandsZeroExtendMask;
  return 0;
}

bool OCLToSPIRVLegacy::runOnModule(Module &M) {
  setOCLTypeToSPIRV(&getAnalysis<OCLTypeToSPIRVLegacy>());
  return runOCLToSPIRV(M);
}

void OCLToSPIRVLegacy::getAnalysisUsage(AnalysisUsage &AU) const {
  AU.addRequired<OCLTypeToSPIRVLegacy>();
}

llvm::PreservedAnalyses OCLToSPIRVPass::run(llvm::Module &M,
                                            llvm::ModuleAnalysisManager &MAM) {
  setOCLTypeToSPIRV(&MAM.getResult<OCLTypeToSPIRVPass>(M));
  return runOCLToSPIRV(M) ? llvm::PreservedAnalyses::none()
                          : llvm::PreservedAnalyses::all();
}

/// Get vector width from OpenCL vload* function name.
SPIRVWord OCLToSPIRVBase::getVecLoadWidth(const std::string &DemangledName) {
  SPIRVWord Width = 0;
  if (DemangledName == "vloada_half")
    Width = 1;
  else {
    unsigned Loc = 5;
    if (DemangledName.find("vload_half") == 0)
      Loc = 10;
    else if (DemangledName.find("vloada_half") == 0)
      Loc = 11;

    std::stringstream SS(DemangledName.substr(Loc));
    SS >> Width;
  }
  return Width;
}

/// Transform OpenCL vload/vstore function name.
void OCLToSPIRVBase::transVecLoadStoreName(std::string &DemangledName,
                                           const std::string &Stem,
                                           bool AlwaysN) {
  auto HalfStem = Stem + "_half";
  auto HalfStemR = HalfStem + "_r";
  if (!AlwaysN && DemangledName == HalfStem)
    return;
  if (!AlwaysN && DemangledName.find(HalfStemR) == 0) {
    DemangledName = HalfStemR;
    return;
  }
  if (DemangledName.find(HalfStem) == 0) {
    auto OldName = DemangledName;
    DemangledName = HalfStem + "n";
    if (OldName.find("_r") != std::string::npos)
      DemangledName += "_r";
    return;
  }
  if (DemangledName.find(Stem) == 0) {
    DemangledName = Stem + "n";
    return;
  }
}

char OCLToSPIRVLegacy::ID = 0;

bool OCLToSPIRVBase::runOCLToSPIRV(Module &Module) {
  initialize(Module);
  Ctx = &M->getContext();
  auto Src = getSPIRVSource(&Module);
  // This is a pre-processing pass, which transform LLVM IR module to a more
  // suitable form for the SPIR-V translation: it is specifically designed to
  // handle OpenCL C built-in functions and shouldn't be launched for other
  // source languages
  if (std::get<0>(Src) != spv::SourceLanguageOpenCL_C)
    return false;

  CLVer = std::get<1>(Src);

  LLVM_DEBUG(dbgs() << "Enter OCLToSPIRV:\n");

  visit(*M);

  for (Instruction *I : ValuesToDelete)
    I->eraseFromParent();

  eraseUselessFunctions(M); // remove unused functions declarations
  LLVM_DEBUG(dbgs() << "After OCLToSPIRV:\n" << *M);

  verifyRegularizationPass(*M, "OCLToSPIRV");

  return true;
}

// The order of handling OCL builtin functions is important.
// Workgroup functions need to be handled before pipe functions since
// there are functions fall into both categories.
void OCLToSPIRVBase::visitCallInst(CallInst &CI) {
  LLVM_DEBUG(dbgs() << "[visistCallInst] " << CI << '\n');
  auto *F = CI.getCalledFunction();
  if (!F)
    return;

  auto MangledName = F->getName();
  StringRef DemangledName;
  if (!oclIsBuiltin(MangledName, DemangledName))
    return;

  LLVM_DEBUG(dbgs() << "DemangledName: " << DemangledName << '\n');
  if (DemangledName.find(kOCLBuiltinName::NDRangePrefix) == 0) {
    visitCallNDRange(&CI, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::All) {
    visitCallAllAny(OpAll, &CI);
    return;
  }
  if (DemangledName == kOCLBuiltinName::Any) {
    visitCallAllAny(OpAny, &CI);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::AsyncWorkGroupCopy) == 0 ||
      DemangledName.find(kOCLBuiltinName::AsyncWorkGroupStridedCopy) == 0) {
    visitCallAsyncWorkGroupCopy(&CI, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 ||
      DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) {

    // Compute atomic builtins do not support floating types.
    if (CI.getType()->isFloatingPointTy() &&
        isComputeAtomicOCLBuiltin(DemangledName))
      return;

    auto *PCI = &CI;
    if (DemangledName == kOCLBuiltinName::AtomicInit) {
      visitCallAtomicInit(PCI);
      return;
    }
    if (DemangledName == kOCLBuiltinName::AtomicWorkItemFence) {
      visitCallAtomicWorkItemFence(PCI);
      return;
    }
    if (DemangledName == kOCLBuiltinName::AtomicCmpXchgWeak ||
        DemangledName == kOCLBuiltinName::AtomicCmpXchgStrong ||
        DemangledName == kOCLBuiltinName::AtomicCmpXchgWeakExplicit ||
        DemangledName == kOCLBuiltinName::AtomicCmpXchgStrongExplicit) {
      assert((CLVer == kOCLVer::CL20 || CLVer == kOCLVer::CL30) &&
             "Wrong version of OpenCL");
      PCI = visitCallAtomicCmpXchg(PCI);
    }
    visitCallAtomicLegacy(PCI, MangledName, DemangledName);
    visitCallAtomicCpp11(PCI, MangledName, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::ConvertPrefix) == 0) {
    visitCallConvert(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::GetImageWidth ||
      DemangledName == kOCLBuiltinName::GetImageHeight ||
      DemangledName == kOCLBuiltinName::GetImageDepth ||
      DemangledName == kOCLBuiltinName::GetImageDim ||
      DemangledName == kOCLBuiltinName::GetImageArraySize) {
    visitCallGetImageSize(&CI, DemangledName);
    return;
  }
  if ((DemangledName.find(kOCLBuiltinName::WorkGroupPrefix) == 0 &&
       DemangledName != kOCLBuiltinName::WorkGroupBarrier) ||
      DemangledName == kOCLBuiltinName::WaitGroupEvent ||
      (DemangledName.find(kOCLBuiltinName::SubGroupPrefix) == 0 &&
       DemangledName != kOCLBuiltinName::SubGroupBarrier)) {
    visitCallGroupBuiltin(&CI, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::MemFence ||
      DemangledName == kOCLBuiltinName::ReadMemFence ||
      DemangledName == kOCLBuiltinName::WriteMemFence) {
    visitCallMemFence(&CI, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0) {
    if (MangledName.find(kMangledName::Sampler) != StringRef::npos) {
      visitCallReadImageWithSampler(&CI, MangledName, DemangledName);
      return;
    }
    if (MangledName.find("msaa") != StringRef::npos) {
      visitCallReadImageMSAA(&CI, MangledName);
      return;
    }
  }
  if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0 ||
      DemangledName.find(kOCLBuiltinName::WriteImage) == 0) {
    visitCallReadWriteImage(&CI, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::ToGlobal ||
      DemangledName == kOCLBuiltinName::ToLocal ||
      DemangledName == kOCLBuiltinName::ToPrivate) {
    visitCallToAddr(&CI, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::VLoadPrefix) == 0 ||
      DemangledName.find(kOCLBuiltinName::VStorePrefix) == 0) {
    visitCallVecLoadStore(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::IsFinite ||
      DemangledName == kOCLBuiltinName::IsInf ||
      DemangledName == kOCLBuiltinName::IsNan ||
      DemangledName == kOCLBuiltinName::IsNormal ||
      DemangledName == kOCLBuiltinName::Signbit) {
    visitCallRelational(&CI, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::WorkGroupBarrier ||
      DemangledName == kOCLBuiltinName::Barrier ||
      DemangledName == kOCLBuiltinName::SubGroupBarrier) {
    visitCallBarrier(&CI);
    return;
  }
  if (DemangledName == kOCLBuiltinName::GetFence) {
    visitCallGetFence(&CI, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::Dot &&
      CI.getOperand(0)->getType()->isFloatingPointTy()) {
    visitCallDot(&CI);
    return;
  }
  if (DemangledName == kOCLBuiltinName::Dot ||
      DemangledName == kOCLBuiltinName::DotAccSat ||
      DemangledName.starts_with(kOCLBuiltinName::Dot4x8PackedPrefix) ||
      DemangledName.starts_with(kOCLBuiltinName::DotAccSat4x8PackedPrefix)) {
    if (CI.getOperand(0)->getType()->isVectorTy()) {
      auto *VT = (VectorType *)(CI.getOperand(0)->getType());
      if (!isa<llvm::IntegerType>(VT->getElementType())) {
        visitCallBuiltinSimple(&CI, MangledName, DemangledName);
        return;
      }
    }
    visitCallDot(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName.starts_with(kOCLBuiltinName::ClockReadPrefix)) {
    visitCallClockRead(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::FMin ||
      DemangledName == kOCLBuiltinName::FMax ||
      DemangledName == kOCLBuiltinName::Min ||
      DemangledName == kOCLBuiltinName::Max ||
      DemangledName == kOCLBuiltinName::Step ||
      DemangledName == kOCLBuiltinName::SmoothStep ||
      DemangledName == kOCLBuiltinName::Clamp ||
      DemangledName == kOCLBuiltinName::Mix) {
    visitCallScalToVec(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::GetImageChannelDataType) {
    visitCallGetImageChannel(&CI, DemangledName, OCLImageChannelDataTypeOffset);
    return;
  }
  if (DemangledName == kOCLBuiltinName::GetImageChannelOrder) {
    visitCallGetImageChannel(&CI, DemangledName, OCLImageChannelOrderOffset);
    return;
  }
  if (isEnqueueKernelBI(MangledName)) {
    visitCallEnqueueKernel(&CI, DemangledName);
    return;
  }
  if (isKernelQueryBI(MangledName)) {
    visitCallKernelQuery(&CI, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::SubgroupBlockReadINTELPrefix) == 0) {
    visitSubgroupBlockReadINTEL(&CI);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::SubgroupBlockWriteINTELPrefix) == 0) {
    visitSubgroupBlockWriteINTEL(&CI);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::SubgroupImageMediaBlockINTELPrefix) ==
      0) {
    visitSubgroupImageMediaBlockINTEL(&CI, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::SplitBarrierINTELPrefix) == 0) {
    visitCallSplitBarrierINTEL(&CI, DemangledName);
    return;
  }
  // Handle 'cl_intel_device_side_avc_motion_estimation' extension built-ins
  if (DemangledName.find(kOCLSubgroupsAVCIntel::Prefix) == 0 ||
      // Workaround for a bug in the extension specification
      DemangledName.find("intel_sub_group_ime_ref_window_size") == 0) {
    if (MangledName.find(kMangledName::Sampler) != StringRef::npos)
      visitSubgroupAVCBuiltinCallWithSampler(&CI, DemangledName);
    else
      visitSubgroupAVCBuiltinCall(&CI, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::LDEXP) == 0) {
    visitCallLdexp(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::ConvertBFloat16AsUShort ||
      DemangledName == kOCLBuiltinName::ConvertBFloat162AsUShort2 ||
      DemangledName == kOCLBuiltinName::ConvertBFloat163AsUShort3 ||
      DemangledName == kOCLBuiltinName::ConvertBFloat164AsUShort4 ||
      DemangledName == kOCLBuiltinName::ConvertBFloat168AsUShort8 ||
      DemangledName == kOCLBuiltinName::ConvertBFloat1616AsUShort16) {
    visitCallConvertBFloat16AsUshort(&CI, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::ConvertAsBFloat16Float ||
      DemangledName == kOCLBuiltinName::ConvertAsBFloat162Float2 ||
      DemangledName == kOCLBuiltinName::ConvertAsBFloat163Float3 ||
      DemangledName == kOCLBuiltinName::ConvertAsBFloat164Float4 ||
      DemangledName == kOCLBuiltinName::ConvertAsBFloat168Float8 ||
      DemangledName == kOCLBuiltinName::ConvertAsBFloat1616Float16) {
    visitCallConvertAsBFloat16Float(&CI, DemangledName);
    return;
  }
  visitCallBuiltinSimple(&CI, MangledName, DemangledName);
}

void OCLToSPIRVBase::visitCallNDRange(CallInst *CI, StringRef DemangledName) {
  assert(DemangledName.find(kOCLBuiltinName::NDRangePrefix) == 0);
  StringRef LenStr = DemangledName.substr(8, 1);
  auto Len = atoi(LenStr.data());
  assert(Len >= 1 && Len <= 3);
  // Translate ndrange_ND into differently named SPIR-V
  // decorated functions because they have array arugments
  // of different dimension which mangled the same way.
  std::string Postfix("_");
  Postfix += LenStr;
  Postfix += 'D';
  std::string FuncName = getSPIRVFuncName(OpBuildNDRange, Postfix);
  auto Mutator = mutateCallInst(CI, FuncName);

  // SPIR-V ndrange structure requires 3 members in the following order:
  //   global work offset
  //   global work size
  //   local work size
  // The arguments need to add missing members.
  for (size_t I = 1, E = CI->arg_size(); I != E; ++I)
    Mutator.mapArg(I, [=](Value *V) {
      return getScalarOrArray(V, Len, CI->getIterator());
    });
  switch (CI->arg_size()) {
  case 2: {
    // Has global work size.
    auto *T = Mutator.getArg(1)->getType();
    auto *C = getScalarOrArrayConstantInt(CI->getIterator(), T, Len, 0);
    Mutator.appendArg(C);
    Mutator.appendArg(C);
    break;
  }
  case 3: {
    // Has global and local work size.
    auto *T = Mutator.getArg(1)->getType();
    Mutator.appendArg(
        getScalarOrArrayConstantInt(CI->getIterator(), T, Len, 0));
    break;
  }
  case 4: {
    // Move offset arg to the end
    Mutator.moveArg(1, CI->arg_size() - 1);
    break;
  }
  default:
    assert(0 && "Invalid number of arguments");
  }
}

void OCLToSPIRVBase::visitCallAsyncWorkGroupCopy(CallInst *CI,
                                                 StringRef DemangledName) {
  assert(CI->getCalledFunction() && "Unexpected indirect call");
  auto Mutator = mutateCallInst(CI, OpGroupAsyncCopy);
  if (DemangledName == OCLUtil::kOCLBuiltinName::AsyncWorkGroupCopy)
    Mutator.insertArg(3, addSizet(1));
  Mutator.insertArg(0, addInt32(ScopeWorkgroup));
}

CallInst *OCLToSPIRVBase::visitCallAtomicCmpXchg(CallInst *CI) {
  CallInst *NewCI = nullptr;
  {
    auto Mutator = mutateCallInst(CI, kOCLBuiltinName::AtomicCmpXchgStrong);
    Value *Expected = Mutator.getArg(1);
    Type *MemTy = Mutator.getArg(2)->getType();
    if (MemTy->isFloatTy() || MemTy->isDoubleTy()) {
      MemTy =
          MemTy->isFloatTy() ? Type::getInt32Ty(*Ctx) : Type::getInt64Ty(*Ctx);
      Mutator.replaceArg(
          0,
          {Mutator.getArg(0),
           TypedPointerType::get(
               MemTy, Mutator.getArg(0)->getType()->getPointerAddressSpace())});
      Mutator.mapArg(2, [=](IRBuilder<> &Builder, Value *V) {
        return Builder.CreateBitCast(V, MemTy);
      });
    }
    assert(MemTy->isIntegerTy() &&
           "In SPIR-V 1.0 arguments of OpAtomicCompareExchange must be "
           "an integer type scalars");
    Mutator.mapArg(1, [=](IRBuilder<> &Builder, Value *V) {
      return Builder.CreateLoad(MemTy, V, "exp");
    });
    Mutator.changeReturnType(
        MemTy, [Expected, &NewCI](IRBuilder<> &Builder, CallInst *NCI) {
          NewCI = NCI;
          Builder.CreateStore(NCI, Expected);
          return Builder.CreateICmpEQ(NCI, NCI->getArgOperand(1));
        });
  }
  return NewCI;
}

void OCLToSPIRVBase::visitCallAtomicInit(CallInst *CI) {
  auto *ST = new StoreInst(CI->getArgOperand(1), CI->getArgOperand(0),
                           CI->getIterator());
  ST->takeName(CI);
  CI->dropAllReferences();
  CI->eraseFromParent();
}

void OCLToSPIRVBase::visitCallAllAny(spv::Op OC, CallInst *CI) {
  assert(CI->getCalledFunction() && "Unexpected indirect call");

  auto Args = getArguments(CI);
  assert(Args.size() == 1);

  auto *ArgTy = Args[0]->getType();
  auto *Zero = Constant::getNullValue(Args[0]->getType());

  auto *Cmp = CmpInst::Create(CmpInst::ICmp, CmpInst::ICMP_SLT, Args[0], Zero,
                              "cast", CI->getIterator());

  if (!isa<VectorType>(ArgTy)) {
    auto *Cast = CastInst::CreateZExtOrBitCast(
        Cmp, Type::getInt32Ty(*Ctx), "", Cmp->getNextNode()->getIterator());
    CI->replaceAllUsesWith(Cast);
    CI->eraseFromParent();
  } else {
    mutateCallInst(CI, OC).setArgs({Cmp}).changeReturnType(
        Type::getInt32Ty(*Ctx), [](IRBuilder<> &Builder, CallInst *CI) {
          return Builder.CreateZExtOrBitCast(CI, Builder.getInt32Ty());
        });
  }
}

void OCLToSPIRVBase::visitCallAtomicWorkItemFence(CallInst *CI) {
  transMemoryBarrier(CI, getAtomicWorkItemFenceLiterals(CI));
}

void OCLToSPIRVBase::visitCallMemFence(CallInst *CI, StringRef DemangledName) {
  OCLMemOrderKind MO = StringSwitch<OCLMemOrderKind>(DemangledName)
                           .Case(kOCLBuiltinName::ReadMemFence, OCLMO_acquire)
                           .Case(kOCLBuiltinName::WriteMemFence, OCLMO_release)
                           .Default(OCLMO_acq_rel); // kOCLBuiltinName::MemFence
  transMemoryBarrier(
      CI,
      std::make_tuple(cast<ConstantInt>(CI->getArgOperand(0))->getZExtValue(),
                      MO, OCLMS_work_group));
}

void OCLToSPIRVBase::transMemoryBarrier(CallInst *CI,
                                        AtomicWorkItemFenceLiterals Lit) {
  assert(CI->getCalledFunction() && "Unexpected indirect call");
  mutateCallInst(CI, OpMemoryBarrier)
      .setArgs({addInt32(map<Scope>(std::get<2>(Lit))),
                addInt32(mapOCLMemSemanticToSPIRV(std::get<0>(Lit),
                                                  std::get<1>(Lit)))});
}

void OCLToSPIRVBase::visitCallAtomicLegacy(CallInst *CI, StringRef MangledName,
                                           StringRef DemangledName) {
  StringRef Stem = DemangledName;
  if (Stem.starts_with("atom_"))
    Stem = Stem.drop_front(strlen("atom_"));
  else if (Stem.starts_with("atomic_"))
    Stem = Stem.drop_front(strlen("atomic_"));
  else
    return;

  std::string Sign;
  std::string Postfix;
  std::string Prefix;
  if (Stem == "add" || Stem == "sub" || Stem == "and" || Stem == "or" ||
      Stem == "xor" || Stem == "min" || Stem == "max") {
    if ((Stem == "min" || Stem == "max") &&
        isMangledTypeUnsigned(MangledName.back()))
      Sign = 'u';
    Prefix = "fetch_";
    Postfix = "_explicit";
  } else if (Stem == "xchg") {
    Stem = "exchange";
    Postfix = "_explicit";
  } else if (Stem == "cmpxchg") {
    Stem = "compare_exchange_strong";
    Postfix = "_explicit";
  } else if (Stem == "inc" || Stem == "dec") {
    // do nothing
  } else
    return;

  OCLBuiltinTransInfo Info;
  Info.UniqName = "atomic_" + Prefix + Sign + Stem.str() + Postfix;
  std::vector<int> PostOps;
  PostOps.push_back(OCLLegacyAtomicMemOrder);
  if (Stem.starts_with("compare_exchange"))
    PostOps.push_back(OCLLegacyAtomicMemOrder);
  PostOps.push_back(OCLLegacyAtomicMemScope);

  Info.PostProc = [=](BuiltinCallMutator &Mutator) {
    for (auto &I : PostOps) {
      Mutator.appendArg(addInt32(I));
    }
  };
  transAtomicBuiltin(CI, Info);
}

void OCLToSPIRVBase::visitCallAtomicCpp11(CallInst *CI, StringRef MangledName,
                                          StringRef DemangledName) {
  StringRef Stem = DemangledName;
  if (Stem.starts_with("atomic_"))
    Stem = Stem.drop_front(strlen("atomic_"));
  else
    return;

  std::string NewStem(Stem);
  std::vector<int> PostOps;
  if (Stem.starts_with("store") || Stem.starts_with("load") ||
      Stem.starts_with("exchange") || Stem.starts_with("compare_exchange") ||
      Stem.starts_with("fetch") || Stem.starts_with("flag")) {
    if ((Stem.starts_with("fetch_min") || Stem.starts_with("fetch_max")) &&
        containsUnsignedAtomicType(MangledName))
      NewStem.insert(NewStem.begin() + strlen("fetch_"), 'u');

    if (!Stem.ends_with("_explicit")) {
      NewStem = NewStem + "_explicit";
      PostOps.push_back(OCLMO_seq_cst);
      if (Stem.starts_with("compare_exchange"))
        PostOps.push_back(OCLMO_seq_cst);
      PostOps.push_back(OCLMS_device);
    } else {
      auto MaxOps =
          getOCLCpp11AtomicMaxNumOps(Stem.drop_back(strlen("_explicit")));
      if (CI->arg_size() < MaxOps)
        PostOps.push_back(OCLMS_device);
    }
  } else if (Stem == "work_item_fence") {
    // do nothing
  } else
    return;

  OCLBuiltinTransInfo Info;
  Info.UniqName = std::string("atomic_") + NewStem;
  Info.PostProc = [=](BuiltinCallMutator &Mutator) {
    for (auto &I : PostOps) {
      Mutator.appendArg(addInt32(I));
    }
  };

  transAtomicBuiltin(CI, Info);
}

void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
                                        OCLBuiltinTransInfo &Info) {
  llvm::Type *AtomicBuiltinsReturnType = CI->getType();
  auto SPIRVFunctionName =
      getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
  bool NeedsNegate = false;
  if (AtomicBuiltinsReturnType->isFloatingPointTy()) {
    // Translate FP-typed atomic builtins. Currently we only need to
    // translate atomic_fetch_[add, sub, max, min] and atomic_fetch_[add,
    // sub, max, min]_explicit to related float instructions.
    // Translate atomic_fetch_sub to OpAtomicFAddEXT with negative value
    // operand
    auto SPIRFunctionNameForFloatAtomics =
        llvm::StringSwitch<std::string>(SPIRVFunctionName)
            .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT")
            .Case("__spirv_AtomicISub", "__spirv_AtomicFAddEXT")
            .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT")
            .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT")
            .Default("others");
    if (SPIRVFunctionName == "__spirv_AtomicISub") {
      NeedsNegate = true;
    }
    if (SPIRFunctionNameForFloatAtomics != "others")
      SPIRVFunctionName = SPIRFunctionNameForFloatAtomics;
  }

  auto Mutator = mutateCallInst(CI, SPIRVFunctionName);
  Info.PostProc(Mutator);
  // Order of args in OCL20:
  // object, 0-2 other args, 1-2 order, scope
  const size_t NumOrder = getAtomicBuiltinNumMemoryOrderArgs(Info.UniqName);
  const size_t ArgsCount = Mutator.arg_size();
  const size_t ScopeIdx = ArgsCount - 1;
  const size_t OrderIdx = ScopeIdx - NumOrder;

  unsigned PtrMemSemantics = MemorySemanticsMaskNone;
  if (Mutator.arg_size() > 0)
    PtrMemSemantics = getAtomicPointerMemorySemanticsMemoryMask(
        Mutator.getArg(0), Mutator.getType(0));

  if (NeedsNegate) {
    Mutator.mapArg(1, [=](Value *V) {
      IRBuilder<> IRB(CI);
      return IRB.CreateFNeg(V);
    });
  }
  Mutator.mapArg(ScopeIdx, [=](Value *V) {
    return transOCLMemScopeIntoSPIRVScope(V, OCLMS_device, CI);
  });
  for (size_t I = 0; I < NumOrder; ++I) {
    Mutator.mapArg(
        OrderIdx + I, [=](IRBuilder<> &Builder, Value *V) -> Value * {
          Value *MemSem =
              transOCLMemOrderIntoSPIRVMemorySemantics(V, OCLMO_seq_cst, CI);
          if (PtrMemSemantics == MemorySemanticsMaskNone)
            return MemSem;

          auto *MemSemTy = cast<IntegerType>(MemSem->getType());
          auto *Mask = ConstantInt::get(MemSemTy, PtrMemSemantics);
          if (auto *Const = dyn_cast<ConstantInt>(MemSem))
            return static_cast<Value *>(ConstantInt::get(
                MemSemTy, Const->getZExtValue() | PtrMemSemantics));
          return Builder.CreateOr(MemSem, Mask);
        });
  }

  // Order of args in SPIR-V:
  // object, scope, 1-2 order, 0-2 other args
  for (size_t I = 0; I < NumOrder; ++I) {
    Mutator.moveArg(OrderIdx + I, I + 1);
  }
  Mutator.moveArg(ScopeIdx, 1);
  if (Info.UniqName.find("atomic_compare_exchange") == 0) {
    // For atomic_compare_exchange, the two "other args" are in the opposite
    // order from the SPIR-V order. Swap these two arguments.
    Mutator.moveArg(Mutator.arg_size() - 1, Mutator.arg_size() - 2);
  }
}

void OCLToSPIRVBase::visitCallBarrier(CallInst *CI) {
  auto Lit = getBarrierLiterals(CI);
  // Use sequential consistent memory order by default.
  // But if the flags argument is set to 0, we use
  // None(Relaxed) memory order.
  unsigned MemFenceFlag = std::get<0>(Lit);
  OCLMemOrderKind MemOrder = MemFenceFlag ? OCLMO_seq_cst : OCLMO_relaxed;
  mutateCallInst(CI, OpControlBarrier)
      .setArgs({// Execution scope
                addInt32(map<Scope>(std::get<2>(Lit))),
                // Memory scope
                addInt32(map<Scope>(std::get<1>(Lit))),
                // Memory semantics
                addInt32(mapOCLMemSemanticToSPIRV(MemFenceFlag, MemOrder))});
}

void OCLToSPIRVBase::visitCallConvert(CallInst *CI, StringRef MangledName,
                                      StringRef DemangledName) {
  // OpenCL Explicit Conversions (6.4.3) formed as below for scalars:
  // destType convert_destType<_sat><_roundingMode>(sourceType)
  // and for vector type:
  // destTypeN convert_destTypeN<_sat><_roundingMode>(sourceTypeN)
  // If the demangled name is not matching the suggested pattern and does not
  // meet allowed destination type restrictions - this is not an OpenCL builtin,
  // return from the function and translate such CallInst as a function call.
  if (eraseUselessConvert(CI, MangledName, DemangledName))
    return;
  Op OC = OpNop;
  auto *TargetTy = CI->getType();
  auto *SrcTy = CI->getArgOperand(0)->getType();
  if (auto *VecTy = dyn_cast<VectorType>(TargetTy))
    TargetTy = VecTy->getElementType();
  if (auto *VecTy = dyn_cast<VectorType>(SrcTy))
    SrcTy = VecTy->getElementType();
  auto IsTargetInt = isa<IntegerType>(TargetTy);

  // Validate conversion function name and vector size if present
  std::regex Expr(
      "convert_(float|double|half|u?char|u?short|u?int|u?long)(2|3|4|8|16)*"
      "(_sat)*(_rt[ezpn])*$");
  std::smatch DestTyMatch;
  std::string ConversionFunc(DemangledName.str());
  if (!std::regex_match(ConversionFunc, DestTyMatch, Expr))
    return;

  // The first sub_match is the whole string; the next
  // sub_matches are the parenthesized expressions.
  enum { TypeIdx = 1, VecSizeIdx = 2, SatIdx = 3, RoundingIdx = 4 };
  std::string DestTy = DestTyMatch[TypeIdx].str();
  std::string VecSize = DestTyMatch[VecSizeIdx].str();
  std::string Sat = DestTyMatch[SatIdx].str();
  std::string Rounding = DestTyMatch[RoundingIdx].str();

  bool TargetSigned = DestTy[0] != 'u';

  if (isa<IntegerType>(SrcTy)) {
    bool Signed = isLastFuncParamSigned(MangledName);
    if (IsTargetInt) {
      if (!Sat.empty() && TargetSigned != Signed) {
        OC = Signed ? OpSatConvertSToU : OpSatConvertUToS;
        Sat = "";
      } else
        OC = Signed ? OpSConvert : OpUConvert;
    } else
      OC = Signed ? OpConvertSToF : OpConvertUToF;
  } else {
    if (IsTargetInt) {
      OC = TargetSigned ? OpConvertFToS : OpConvertFToU;
    } else
      OC = OpFConvert;
  }

  assert(CI->getCalledFunction() && "Unexpected indirect call");
  mutateCallInst(
      CI, getSPIRVFuncName(OC, "_R" + DestTy + VecSize + Sat + Rounding));
}

void OCLToSPIRVBase::visitCallGroupBuiltin(CallInst *CI,
                                           StringRef OrigDemangledName) {
  auto *F = CI->getCalledFunction();
  std::vector<int> PreOps;
  std::string DemangledName{OrigDemangledName};

  if (DemangledName == kOCLBuiltinName::WorkGroupBarrier)
    return;
  if (DemangledName == kOCLBuiltinName::WaitGroupEvent) {
    PreOps.push_back(ScopeWorkgroup);
  } else if (DemangledName.find(kOCLBuiltinName::WorkGroupPrefix) == 0) {
    DemangledName.erase(0, strlen(kOCLBuiltinName::WorkPrefix));
    PreOps.push_back(ScopeWorkgroup);
  } else if (DemangledName.find(kOCLBuiltinName::SubGroupPrefix) == 0) {
    DemangledName.erase(0, strlen(kOCLBuiltinName::SubPrefix));
    PreOps.push_back(ScopeSubgroup);
  } else
    return;

  if (DemangledName != kOCLBuiltinName::WaitGroupEvent) {
    StringRef FuncName = DemangledName;
    FuncName = FuncName.drop_front(strlen(kSPIRVName::GroupPrefix));
    SPIRSPIRVGroupOperationMap::foreachConditional(
        [&](const std::string &S, SPIRVGroupOperationKind G) {
          if (!FuncName.starts_with(S))
            return true; // continue
          PreOps.push_back(G);
          StringRef Op =
              StringSwitch<StringRef>(FuncName)
                  .StartsWith("ballot", "group_ballot_bit_count_")
                  .StartsWith("non_uniform", kSPIRVName::GroupNonUniformPrefix)
                  .Default(kSPIRVName::GroupPrefix);
          // clustered functions are handled with non uniform group opcodes
          StringRef ClusteredOp =
              FuncName.contains("clustered_") ? "non_uniform_" : "";
          StringRef LogicalOp = FuncName.contains("logical_") ? "logical_" : "";
          StringRef GroupOp = StringSwitch<StringRef>(FuncName)
                                  .Case("ballot_bit_count", "add")
                                  .Case("ballot_inclusive_scan", "add")
                                  .Case("ballot_exclusive_scan", "add")
                                  .Default(FuncName.take_back(
                                      3));    // assumes op is three characters
          (void)(GroupOp.consume_front("_")); // when op is two characters
          assert(!GroupOp.empty() && "Invalid OpenCL group builtin function");
          char OpTyC = 0;
          auto *OpTy = F->getReturnType();
          if (OpTy->isFloatingPointTy())
            OpTyC = 'f';
          else if (OpTy->isIntegerTy()) {
            auto NeedSign = GroupOp == "max" || GroupOp == "min";
            if (!NeedSign)
              OpTyC = 'i';
            else {
              // clustered reduce args are (type, uint)
              // other operation args are (type)
              auto MangledName = F->getName();
              auto MangledTyC = ClusteredOp.empty()
                                    ? MangledName.back()
                                    : MangledName.take_back(2).front();
              if (isMangledTypeSigned(MangledTyC))
                OpTyC = 's';
              else
                OpTyC = 'u';
            }
          } else
            llvm_unreachable("Invalid OpenCL group builtin argument type");

          DemangledName = Op.str() + ClusteredOp.str() + LogicalOp.str() +
                          OpTyC + GroupOp.str();
          return false; // break out of loop
        });
  }

  const bool IsElect = DemangledName == "group_elect";
  const bool IsAllOrAny = (DemangledName.find("_all") != std::string::npos ||
                           DemangledName.find("_any") != std::string::npos);
  const bool IsAllEqual = DemangledName.find("_all_equal") != std::string::npos;
  const bool IsBallot = DemangledName == "group_ballot";
  const bool IsInverseBallot = DemangledName == "group_inverse_ballot";
  const bool IsBallotBitExtract = DemangledName == "group_ballot_bit_extract";
  const bool IsLogical = DemangledName.find("_logical") != std::string::npos;

  const bool HasBoolReturnType = IsElect || IsAllOrAny || IsAllEqual ||
                                 IsInverseBallot || IsBallotBitExtract ||
                                 IsLogical;
  const bool HasBoolArg = (IsAllOrAny && !IsAllEqual) || IsBallot || IsLogical;

  auto Consts = getInt32(M, PreOps);
  OCLBuiltinTransInfo Info;
  if (HasBoolReturnType)
    Info.RetTy = Type::getInt1Ty(*Ctx);
  Info.UniqName = DemangledName;
  Info.PostProc = [=](BuiltinCallMutator &Mutator) {
    if (HasBoolArg) {
      Mutator.mapArg(0, [&](Value *V) {
        IRBuilder<> IRB(CI);
        return IRB.CreateICmpNE(V, IRB.getInt32(0));
      });
    }
    size_t E = Mutator.arg_size();
    if (DemangledName == "group_broadcast" && E > 2) {
      assert(E == 3 || E == 4);
      std::vector<Value *> Ops = getArguments(CI);
      makeVector(CI, Ops, std::make_pair(Ops.begin() + 1, Ops.end()));
      while (Mutator.arg_size() > 1)
        Mutator.removeArg(1);
      Mutator.appendArg(Ops.back());
    }
    for (unsigned I = 0; I < Consts.size(); I++)
      Mutator.insertArg(I, Consts[I]);
  };
  transBuiltin(CI, Info);
}

void OCLToSPIRVBase::transBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
  Op OC = OpNop;
  unsigned ExtOp = ~0U;
  SPIRVBuiltinVariableKind BVKind = BuiltInMax;
  if (StringRef(Info.UniqName).starts_with(kSPIRVName::Prefix))
    return;
  if (OCLSPIRVBuiltinMap::find(Info.UniqName, &OC)) {
    if (OC == OpImageRead) {
      // There are several read_image* functions defined by OpenCL C spec, but
      // all of them use the same SPIR-V Instruction - some of them might only
      // differ by return type, so, we need to include return type into the
      // mangling scheme to get them differentiated.
      //
      // Example: int4 read_imagei(image2d_t, sampler_t, int2)
      //          uint4 read_imageui(image2d_t, sampler_t, int2)
      // Both functions above are represented by the same SPIR-V
      // instruction: argument types are the same, only return type is
      // different
      Info.UniqName = getSPIRVFuncName(OC, CI->getType());
    } else {
      Info.UniqName = getSPIRVFuncName(OC);
    }
  } else if ((ExtOp = getExtOp(Info.MangledName, Info.UniqName)) != ~0U)
    Info.UniqName = getSPIRVExtFuncName(SPIRVEIS_OpenCL, ExtOp);
  else if (SPIRSPIRVBuiltinVariableMap::find(Info.UniqName, &BVKind)) {
    // Map OCL work item builtins to SPV-IR work item builtins.
    // e.g. get_global_id() --> __spirv_BuiltinGlobalInvocationId()
    Info.UniqName = getSPIRVFuncName(BVKind);
  } else
    return;
  BuiltinCallMutator Mutator = mutateCallInst(CI, Info.UniqName + Info.Postfix);
  Info.PostProc(Mutator);
  if (Info.RetTy) {
    Type *OldRetTy = CI->getType();
    Mutator.changeReturnType(
        Info.RetTy, [OldRetTy, &Info](IRBuilder<> &Builder, CallInst *NewCI) {
          if (Info.RetTy->isIntegerTy() && OldRetTy->isIntegerTy()) {
            return Builder.CreateIntCast(NewCI, OldRetTy, false);
          }
          return Builder.CreateAddrSpaceCast(NewCI, OldRetTy);
        });
  }
}

void OCLToSPIRVBase::visitCallReadImageMSAA(CallInst *CI,
                                            StringRef MangledName) {
  assert(MangledName.find("msaa") != StringRef::npos);
  mutateCallInst(
      CI, getSPIRVFuncName(OpImageRead, std::string(kSPIRVPostfix::ExtDivider) +
                                            getPostfixForReturnType(CI)))
      .insertArg(2, getInt32(M, ImageOperandsSampleMask));
}

void OCLToSPIRVBase::visitCallReadImageWithSampler(CallInst *CI,
                                                   StringRef MangledName,
                                                   StringRef DemangledName) {
  assert(MangledName.find(kMangledName::Sampler) != StringRef::npos);
  assert(CI->getCalledFunction() && "Unexpected indirect call");
  Function *Func = CI->getCalledFunction();
  bool IsRetScalar = !CI->getType()->isVectorTy();
  Type *Ret = CI->getType();
  auto *ImageTy = OCLTypeToSPIRVPtr->getAdaptedArgumentType(Func, 0);
  if (!ImageTy)
    ImageTy = getCallValueType(CI, 0);

  auto Mutator = mutateCallInst(
      CI, getSPIRVFuncName(OpImageSampleExplicitLod,
                           std::string(kSPIRVPostfix::ExtDivider) +
                               getPostfixForReturnType(Ret)));
  Mutator.mapArg(0, [&](IRBuilder<> &Builder, Value *ImgArg, Type *ImgType) {
    auto *SampledImgTy = adjustImageType(ImageTy, kSPIRVTypeName::Image,
                                         kSPIRVTypeName::SampledImg);
    Value *SampledImgArgs[] = {CI->getArgOperand(0), CI->getArgOperand(1)};
    return addSPIRVCallPair(Builder, OpSampledImage, SampledImgTy,
                            SampledImgArgs, {ImgType, Mutator.getType(1)},
                            kSPIRVName::TempSampledImage);
  });
  Mutator.removeArg(1);
  unsigned ImgOpMask = getImageSignZeroExt(DemangledName);
  unsigned ImgOpMaskInsIndex = Mutator.arg_size();
  switch (Mutator.arg_size()) {
  case 2: // no lod
    ImgOpMask |= ImageOperandsMask::ImageOperandsLodMask;
    ImgOpMaskInsIndex = Mutator.arg_size();
    Mutator.appendArg(getFloat32(M, 0.f));
    break;
  case 3: // explicit lod
    ImgOpMask |= ImageOperandsMask::ImageOperandsLodMask;
    ImgOpMaskInsIndex = 2;
    break;
  case 4: // gradient
    ImgOpMask |= ImageOperandsMask::ImageOperandsGradMask;
    ImgOpMaskInsIndex = 2;
    break;
  default:
    assert(0 && "read_image* with unhandled number of args!");
  }
  Mutator.insertArg(ImgOpMaskInsIndex, getInt32(M, ImgOpMask));

  // SPIR-V instruction always returns 4-element vector
  if (IsRetScalar)
    Mutator.changeReturnType(FixedVectorType::get(Ret, 4),
                             [=](IRBuilder<> &Builder, CallInst *NewCI) {
                               return Builder.CreateExtractElement(
                                   NewCI, getSizet(M, 0));
                             });
}

void OCLToSPIRVBase::visitCallGetImageSize(CallInst *CI,
                                           StringRef DemangledName) {
  auto Desc = getImageDescriptor(getCallValueType(CI, 0));
  unsigned Dim = getImageDimension(Desc.Dim) + Desc.Arrayed;
  assert(Dim > 0 && "Invalid image dimension.");
  assert(CI->arg_size() == 1);
  Type *NewRet = CI->getType()->isIntegerTy(64) ? Type::getInt64Ty(*Ctx)
                                                : Type::getInt32Ty(*Ctx);
  if (Dim > 1)
    NewRet = FixedVectorType::get(NewRet, Dim);
  auto Mutator = mutateCallInst(CI, getSPIRVFuncName(Desc.Dim == DimBuffer
                                                         ? OpImageQuerySize
                                                         : OpImageQuerySizeLod,
                                                     CI->getType()));
  if (Desc.Dim != DimBuffer)
    Mutator.appendArg(getInt32(M, 0));
  Mutator.changeReturnType(
      NewRet, [&](IRBuilder<> &, CallInst *NCI) -> Value * {
        if (Dim == 1)
          return NCI;
        if (DemangledName == kOCLBuiltinName::GetImageDim) {
          if (Desc.Dim == Dim3D) {
            auto *ZeroVec = ConstantVector::getSplat(
                ElementCount::getFixed(3),
                Constant::getNullValue(
                    cast<VectorType>(NCI->getType())->getElementType()));
            Constant *Index[] = {getInt32(M, 0), getInt32(M, 1), getInt32(M, 2),
                                 getInt32(M, 3)};
            return new ShuffleVectorInst(NCI, ZeroVec,
                                         ConstantVector::get(Index), "",
                                         CI->getIterator());

          } else if (Desc.Dim == Dim2D && Desc.Arrayed) {
            Constant *Index[] = {getInt32(M, 0), getInt32(M, 1)};
            Constant *Mask = ConstantVector::get(Index);
            return new ShuffleVectorInst(NCI, PoisonValue::get(NCI->getType()),
                                         Mask, NCI->getName(),
                                         CI->getIterator());
          }
          return NCI;
        }
        unsigned I = StringSwitch<unsigned>(DemangledName)
                         .Case(kOCLBuiltinName::GetImageWidth, 0)
                         .Case(kOCLBuiltinName::GetImageHeight, 1)
                         .Case(kOCLBuiltinName::GetImageDepth, 2)
                         .Case(kOCLBuiltinName::GetImageArraySize, Dim - 1);
        return ExtractElementInst::Create(NCI, getUInt32(M, I), "",
                                          NCI->getNextNode()->getIterator());
      });
}

/// Remove trivial conversion functions
bool OCLToSPIRVBase::eraseUselessConvert(CallInst *CI, StringRef MangledName,
                                         StringRef DemangledName) {
  auto *TargetTy = CI->getType();
  auto *SrcTy = CI->getArgOperand(0)->getType();
  if (auto *VecTy = dyn_cast<VectorType>(TargetTy))
    TargetTy = VecTy->getElementType();
  if (auto *VecTy = dyn_cast<VectorType>(SrcTy))
    SrcTy = VecTy->getElementType();
  if (TargetTy == SrcTy) {
    if (isa<IntegerType>(TargetTy) &&
        DemangledName.find("_sat") != StringRef::npos &&
        isLastFuncParamSigned(MangledName) != (DemangledName[8] != 'u'))
      return false;
    CI->getArgOperand(0)->takeName(CI);
    SPIRVDBG(dbgs() << "[regularizeOCLConvert] " << *CI << " <- "
                    << *CI->getArgOperand(0) << '\n');
    CI->replaceAllUsesWith(CI->getArgOperand(0));
    ValuesToDelete.insert(CI);
    return true;
  }
  return false;
}

void OCLToSPIRVBase::visitCallBuiltinSimple(CallInst *CI, StringRef MangledName,
                                            StringRef DemangledName) {
  OCLBuiltinTransInfo Info;
  Info.MangledName = MangledName.str();
  Info.UniqName = DemangledName.str();
  transBuiltin(CI, Info);
}

void OCLToSPIRVBase::visitCallReadWriteImage(CallInst *CI,
                                             StringRef DemangledName) {
  OCLBuiltinTransInfo Info;
  if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0) {
    Info.UniqName = kOCLBuiltinName::ReadImage;
    unsigned ImgOpMask = getImageSignZeroExt(DemangledName);
    if (ImgOpMask) {
      Module *Mod = M;
      Info.PostProc = [ImgOpMask, Mod](BuiltinCallMutator &Mutator) {
        Mutator.appendArg(getInt32(Mod, ImgOpMask));
      };
    }
  }

  if (DemangledName.find(kOCLBuiltinName::WriteImage) == 0) {
    Info.UniqName = kOCLBuiltinName::WriteImage;
    Info.PostProc = [&](BuiltinCallMutator &Mutator) {
      unsigned ImgOpMask = getImageSignZeroExt(DemangledName);
      unsigned ImgOpMaskInsIndex = Mutator.arg_size();
      if (Mutator.arg_size() == 4) // write with lod
      {
        ImgOpMask |= ImageOperandsMask::ImageOperandsLodMask;
        ImgOpMaskInsIndex = Mutator.arg_size() - 1;
        Mutator.moveArg(2, Mutator.arg_size() - 1);
      }
      if (ImgOpMask) {
        Mutator.insertArg(ImgOpMaskInsIndex, getInt32(M, ImgOpMask));
      }
    };
  }

  transBuiltin(CI, Info);
}

void OCLToSPIRVBase::visitCallToAddr(CallInst *CI, StringRef DemangledName) {
  auto AddrSpace =
      static_cast<SPIRAddressSpace>(CI->getType()->getPointerAddressSpace());
  OCLBuiltinTransInfo Info;
  Info.UniqName = DemangledName.str();
  Info.Postfix = std::string(kSPIRVPostfix::Divider) + "To" +
                 SPIRAddrSpaceCapitalizedNameMap::map(AddrSpace);
  auto *StorageClass = addInt32(SPIRSPIRVAddrSpaceMap::map(AddrSpace));
  Info.RetTy = getInt8PtrTy(cast<PointerType>(CI->getType()));
  Info.PostProc = [=](BuiltinCallMutator &Mutator) {
    Mutator
        .mapArg(Mutator.arg_size() - 1,
                [&](Value *V) {
                  return std::make_pair(
                      castToInt8Ptr(V, CI->getIterator()),
                      TypedPointerType::get(Type::getInt8Ty(V->getContext()),
                                            SPIRAS_Generic));
                })
        .appendArg(StorageClass);
  };
  transBuiltin(CI, Info);
}

void OCLToSPIRVBase::visitCallRelational(CallInst *CI,
                                         StringRef DemangledName) {
  assert(CI->getCalledFunction() && "Unexpected indirect call");
  Op OC = OpNop;
  OCLSPIRVBuiltinMap::find(DemangledName.str(), &OC);
  // i1 or <i1 x N>, depending on whether it returns a vector type.
  Type *BoolTy = CI->getType()->getWithNewType(Type::getInt1Ty(*Ctx));
  mutateCallInst(CI, OC).changeReturnType(
      BoolTy, [=](IRBuilder<> &Builder, CallInst *NewCI) {
        Value *TrueOp = CI->getType()->isVectorTy()
                            ? Constant::getAllOnesValue(CI->getType())
                            : getInt32(M, 1);
        return Builder.CreateSelect(NewCI, TrueOp,
                                    Constant::getNullValue(CI->getType()));
      });
}

void OCLToSPIRVBase::visitCallVecLoadStore(CallInst *CI, StringRef MangledName,
                                           StringRef OrigDemangledName) {
  std::vector<int> PreOps;
  std::string DemangledName{OrigDemangledName};
  if (DemangledName.find(kOCLBuiltinName::VLoadPrefix) == 0 &&
      DemangledName != kOCLBuiltinName::VLoadHalf) {
    SPIRVWord Width = getVecLoadWidth(DemangledName);
    SPIRVDBG(spvdbgs() << "[visitCallVecLoadStore] DemangledName: "
                       << DemangledName << " Width: " << Width << '\n');
    PreOps.push_back(Width);
  } else if (DemangledName.find(kOCLBuiltinName::RoundingPrefix) !=
             std::string::npos) {
    auto R = SPIRSPIRVFPRoundingModeMap::map(DemangledName.substr(
        DemangledName.find(kOCLBuiltinName::RoundingPrefix) + 1, 3));
    PreOps.push_back(R);
  }

  if (DemangledName.find(kOCLBuiltinName::VLoadAPrefix) == 0)
    transVecLoadStoreName(DemangledName, kOCLBuiltinName::VLoadAPrefix, true);
  else
    transVecLoadStoreName(DemangledName, kOCLBuiltinName::VLoadPrefix, false);

  if (DemangledName.find(kOCLBuiltinName::VStoreAPrefix) == 0)
    transVecLoadStoreName(DemangledName, kOCLBuiltinName::VStoreAPrefix, true);
  else
    transVecLoadStoreName(DemangledName, kOCLBuiltinName::VStorePrefix, false);

  auto Consts = getInt32(M, PreOps);
  OCLBuiltinTransInfo Info;
  Info.MangledName = MangledName.str();
  Info.UniqName = DemangledName;
  if (DemangledName.find(kOCLBuiltinName::VLoadPrefix) == 0)
    Info.Postfix =
        std::string(kSPIRVPostfix::ExtDivider) + getPostfixForReturnType(CI);
  Info.PostProc = [=](BuiltinCallMutator &Mutator) {
    for (auto *Value : Consts)
      Mutator.appendArg(Value);
  };
  transBuiltin(CI, Info);
}

void OCLToSPIRVBase::visitCallGetFence(CallInst *CI, StringRef DemangledName) {
  Op OC = OpNop;
  OCLSPIRVBuiltinMap::find(DemangledName.str(), &OC);
  mutateCallInst(CI, OC).changeReturnType(
      CI->getType(), [](IRBuilder<> &Builder, CallInst *NewCI) {
        return Builder.CreateLShr(NewCI, Builder.getInt32(8));
      });
}

void OCLToSPIRVBase::visitCallDot(CallInst *CI) {
  IRBuilder<> Builder(CI);
  Value *FMulVal = Builder.CreateFMul(CI->getOperand(0), CI->getOperand(1));
  CI->replaceAllUsesWith(FMulVal);
  CI->eraseFromParent();
}

void OCLToSPIRVBase::visitCallDot(CallInst *CI, StringRef MangledName,
                                  StringRef DemangledName) {
  // translation for dot function calls,
  // to differentiate between integer dot products

  bool IsFirstSigned, IsSecondSigned;
  bool IsDot = DemangledName == kOCLBuiltinName::Dot;
  bool IsAccSat = DemangledName.contains(kOCLBuiltinName::DotAccSat);
  bool IsPacked = CI->getOperand(0)->getType()->isIntegerTy();
  if (!IsPacked) {
    if (IsDot) {
      // dot(char4, char4) _Z3dotDv4_cS_
      // dot(char4, uchar4) _Z3dotDv4_cDv4_h
      // dot(uchar4, char4) _Z3dotDv4_hDv4_c
      // dot(uchar4, uchar4) _Z3dotDv4_hS_
      // or
      // dot(short2, short2) _Z3dotDv2_sS_
      // dot(short2, ushort2) _Z3dotDv2_sDv2_t
      // dot(ushort2, short2) _Z3dotDv2_tDv2_s
      // dot(ushort2, ushort2) _Z3dotDv2_tS_
      assert(MangledName.starts_with("_Z3dotDv"));
      if (MangledName[MangledName.size() - 1] == '_') {
        IsFirstSigned = ((MangledName[MangledName.size() - 3] == 'c') ||
                         (MangledName[MangledName.size() - 3] == 's'));
        IsSecondSigned = IsFirstSigned;
      } else {
        IsFirstSigned = ((MangledName[MangledName.size() - 6] == 'c') ||
                         (MangledName[MangledName.size() - 6] == 's'));
        IsSecondSigned = ((MangledName[MangledName.size() - 1] == 'c') ||
                          (MangledName[MangledName.size() - 1] == 's'));
      }
    } else {
      // dot_acc_sat(char4, char4, int) _Z11dot_acc_satDv4_cS_i
      // dot_acc_sat(char4, uchar4, int) _Z11dot_acc_satDv4_cDv4_hi
      // dot_acc_sat(uchar4, char4, int) _Z11dot_acc_satDv4_hDv4_ci
      // dot_acc_sat(uchar4, uchar4, uint) _Z11dot_acc_satDv4_hS_j
      // or
      // dot_acc_sat(short2, short2, int) _Z11dot_acc_satDv4_sS_i
      // dot_acc_sat(short2, ushort2, int) _Z11dot_acc_satDv4_sDv4_ti
      // dot_acc_sat(ushort2, short2, int) _Z11dot_acc_satDv4_tDv4_si
      // dot_acc_sat(ushort2, ushort2, uint) _Z11dot_acc_satDv4_tS_j
      assert(MangledName.starts_with("_Z11dot_acc_satDv"));
      IsFirstSigned = ((MangledName[19] == 'c') || (MangledName[19] == 's'));
      IsSecondSigned = (MangledName[20] == 'S'
                            ? IsFirstSigned
                            : ((MangledName[MangledName.size() - 2] == 'c') ||
                               (MangledName[MangledName.size() - 2] == 's')));
    }
  } else {
    // for packed format
    // dot_4x8packed_ss_int(uint, uint) _Z20dot_4x8packed_ss_intjj
    // dot_4x8packed_su_int(uint, uint) _Z20dot_4x8packed_su_intjj
    // dot_4x8packed_us_int(uint, uint) _Z20dot_4x8packed_us_intjj
    // dot_4x8packed_uu_uint(uint, uint) _Z21dot_4x8packed_uu_uintjj
    // or
    // dot_acc_sat_4x8packed_ss_int(uint, uint, int)
    // _Z28dot_acc_sat_4x8packed_ss_intjji
    // dot_acc_sat_4x8packed_su_int(uint, uint, int)
    // _Z28dot_acc_sat_4x8packed_su_intjji
    // dot_acc_sat_4x8packed_us_int(uint, uint, int)
    // _Z28dot_acc_sat_4x8packed_us_intjji
    // dot_acc_sat_4x8packed_uu_uint(uint, uint, uint)
    // _Z29dot_acc_sat_4x8packed_uu_uintjjj
    assert(MangledName.starts_with("_Z20dot_4x8packed") ||
           MangledName.starts_with("_Z21dot_4x8packed") ||
           MangledName.starts_with("_Z28dot_acc_sat_4x8packed") ||
           MangledName.starts_with("_Z29dot_acc_sat_4x8packed"));
    size_t SignIndex = IsAccSat
                           ? strlen(kOCLBuiltinName::DotAccSat4x8PackedPrefix)
                           : strlen(kOCLBuiltinName::Dot4x8PackedPrefix);
    IsFirstSigned = DemangledName[SignIndex] == 's';
    IsSecondSigned = DemangledName[SignIndex + 1] == 's';
  }
  Op OC;
  if (!IsAccSat) {
    OC =
        (IsFirstSigned != IsSecondSigned ? OpSUDot
                                         : ((IsFirstSigned) ? OpSDot : OpUDot));
  } else {
    OC = (IsFirstSigned != IsSecondSigned
              ? OpSUDotAccSat
              : ((IsFirstSigned) ? OpSDotAccSat : OpUDotAccSat));
  }
  auto Mutator = mutateCallInst(CI, OC);
  // If arguments are in order unsigned -> signed
  // then the translator should swap them,
  // so that the OpSUDotKHR can be used properly
  if (IsFirstSigned == false && IsSecondSigned == true) {
    Mutator.moveArg(1, 0);
  }
  if (IsPacked) {
    // As per SPIRV specification the dot OpCodes
    // which use scalar integers to represent
    // packed vectors need additional argument
    // specified - the Packed Vector Format
    Mutator.appendArg(
        getInt32(M, PackedVectorFormatPackedVectorFormat4x8BitKHR));
  }
}

void OCLToSPIRVBase::visitCallClockRead(CallInst *CI, StringRef MangledName,
                                        StringRef DemangledName) {
  // The builtin returns i64 or <2 x i32>, but both variants are mapped to the
  // same instruction; hence include the return type.
  std::string OpName = getSPIRVFuncName(OpReadClockKHR, CI->getType());

  // Scope is part of the OpenCL builtin name.
  Scope ScopeArg = StringSwitch<Scope>(DemangledName)
                       .EndsWith("device", ScopeDevice)
                       .EndsWith("work_group", ScopeWorkgroup)
                       .EndsWith("sub_group", ScopeSubgroup)
                       .Default(ScopeMax);

  auto Mutator = mutateCallInst(CI, OpName);
  Mutator.appendArg(getInt32(M, ScopeArg));
}

void OCLToSPIRVBase::visitCallScalToVec(CallInst *CI, StringRef MangledName,
                                        StringRef DemangledName) {
  // Check if all arguments have the same type - it's simple case.
  auto Uniform = true;
  auto IsArg0Vector = isa<VectorType>(CI->getOperand(0)->getType());
  for (unsigned I = 1, E = CI->arg_size(); Uniform && (I != E); ++I) {
    Uniform = isa<VectorType>(CI->getOperand(I)->getType()) == IsArg0Vector;
  }
  if (Uniform) {
    visitCallBuiltinSimple(CI, MangledName, DemangledName);
    return;
  }

  std::vector<unsigned int> VecPos;
  std::vector<unsigned int> ScalarPos;
  if (DemangledName == kOCLBuiltinName::FMin ||
      DemangledName == kOCLBuiltinName::FMax ||
      DemangledName == kOCLBuiltinName::Min ||
      DemangledName == kOCLBuiltinName::Max) {
    VecPos.push_back(0);
    ScalarPos.push_back(1);
  } else if (DemangledName == kOCLBuiltinName::Clamp) {
    VecPos.push_back(0);
    ScalarPos.push_back(1);
    ScalarPos.push_back(2);
  } else if (DemangledName == kOCLBuiltinName::Mix) {
    VecPos.push_back(0);
    VecPos.push_back(1);
    ScalarPos.push_back(2);
  } else if (DemangledName == kOCLBuiltinName::Step) {
    VecPos.push_back(1);
    ScalarPos.push_back(0);
  } else if (DemangledName == kOCLBuiltinName::SmoothStep) {
    VecPos.push_back(2);
    ScalarPos.push_back(0);
    ScalarPos.push_back(1);
  }

  assert(CI->arg_size() == VecPos.size() + ScalarPos.size() &&
         "Argument counts do not match up.");

  Type *VecTy = CI->getOperand(VecPos[0])->getType();
  auto VecElemCount = cast<VectorType>(VecTy)->getElementCount();
  auto Mutator = mutateCallInst(
      CI, getSPIRVExtFuncName(SPIRVEIS_OpenCL,
                              getExtOp(MangledName, DemangledName)));
  for (auto I : ScalarPos)
    Mutator.mapArg(I, [&](Value *V) {
      Instruction *Inst = InsertElementInst::Create(
          PoisonValue::get(VecTy), V, getInt32(M, 0), "", CI->getIterator());
      return new ShuffleVectorInst(
          Inst, PoisonValue::get(VecTy),
          ConstantVector::getSplat(VecElemCount, getInt32(M, 0)), "",
          CI->getIterator());
    });
}

namespace {
// Return true if any users of the CallInst use any of the constants
// introduced by the SPV_EXT_image_raw10_raw12 extension.
bool usesSpvExtImageRaw10Raw12Constants(const CallInst *CI) {
  const std::array ExtConstants{
      OCLImageChannelDataTypeOffset + ImageChannelDataTypeUnsignedIntRaw10EXT,
      OCLImageChannelDataTypeOffset + ImageChannelDataTypeUnsignedIntRaw12EXT};

  // The return values for `OpImageQueryFormat` added by the extension are
  // integer constants that may appear anywhere in LLVM IR.  Only detect some
  // common use patterns here.
  for (auto *U : CI->users()) {
    for (auto C : ExtConstants) {
      CmpPredicate Pred;
      if (match(U, m_c_ICmp(Pred, m_Value(), m_SpecificInt(C)))) {
        return true;
      }
      if (auto *Switch = dyn_cast<SwitchInst>(U)) {
        if (any_of(Switch->cases(), [C](const auto &Case) {
              return Case.getCaseValue()->equalsInt(C);
            })) {
          return true;
        }
      }
    }
  }
  return false;
}
} // anonymous namespace

void OCLToSPIRVBase::visitCallGetImageChannel(CallInst *CI,
                                              StringRef DemangledName,
                                              unsigned int Offset) {
  assert(CI->getCalledFunction() && "Unexpected indirect call");

  if (Offset == OCLImageChannelDataTypeOffset) {
    // See if any of the SPV_EXT_image_raw10_raw12 constants are used, and
    // add the extension if not already there.
    if (usesSpvExtImageRaw10Raw12Constants(CI)) {
      const char *ExtStr = "SPV_EXT_image_raw10_raw12";
      NamedMDNode *NMD = M->getOrInsertNamedMetadata(kSPIRVMD::Extension);
      if (none_of(NMD->operands(), [ExtStr](MDNode *N) {
            return N->getOperand(0).equalsStr(ExtStr);
          })) {
        MDString *MDS = MDString::get(*Ctx, ExtStr);
        NMD->addOperand(MDNode::get(*Ctx, MDS));
      }
    }
  }

  Op OC = OpNop;
  OCLSPIRVBuiltinMap::find(DemangledName.str(), &OC);
  mutateCallInst(CI, OC).changeReturnType(
      CI->getType(), [=](IRBuilder<> &Builder, CallInst *NewCI) {
        return Builder.CreateAdd(NewCI, Builder.getInt32(Offset));
      });
}
void OCLToSPIRVBase::visitCallEnqueueKernel(CallInst *CI,
                                            StringRef DemangledName) {
  const DataLayout &DL = M->getDataLayout();
  bool HasEvents = DemangledName.find("events") != StringRef::npos;

  // SPIRV OpEnqueueKernel instruction has 10+ arguments.
  SmallVector<Value *, 16> Args;

  // Copy all arguments before block invoke function pointer
  // which match with what Clang 6.0 produced
  const unsigned BlockFIdx = HasEvents ? 6 : 3;
  Args.assign(CI->arg_begin(), CI->arg_begin() + BlockFIdx);

  // If no event arguments in original call, add dummy ones
  if (!HasEvents) {
    Args.push_back(getInt32(M, 0)); // dummy num events
    Value *Null = Constant::getNullValue(
        PointerType::get(CI->getContext(), SPIRAS_Generic));
    Args.push_back(Null); // dummy wait events
    Args.push_back(Null); // dummy ret event
  }

  // Invoke: Pointer to invoke function
  Value *BlockFunc = CI->getArgOperand(BlockFIdx);
  Args.push_back(cast<Function>(getUnderlyingObject(BlockFunc)));

  // Param: Pointer to block literal
  Value *BlockLiteral = CI->getArgOperand(BlockFIdx + 1);
  Args.push_back(BlockLiteral);
  BlockLiteral = BlockLiteral->stripPointerCasts();
  if (auto *GV = dyn_cast<GlobalVariable>(BlockLiteral)) {
    assert(GV->hasInitializer() && "Block literal should have an initializer");
    Constant *Init = GV->getInitializer();
    if (auto *Struct = dyn_cast<ConstantStruct>(Init)) {
      Constant *SizeConst = Struct->getOperand(0);
      Constant *AlignConst = Struct->getOperand(1);
      auto *SizeVal = dyn_cast<ConstantInt>(SizeConst);
      auto *AlignVal = dyn_cast<ConstantInt>(AlignConst);
      Args.push_back(SizeVal);
      Args.push_back(AlignVal);
    }
  } else {

    Value *Base = getUnderlyingObject(BlockLiteral);
    Value *ParamSizeVal = nullptr;
    Value *ParamAlignVal = nullptr;

    for (User *U : Base->users()) {
      if (auto *GEP = dyn_cast<GetElementPtrInst>(U)) {
        if (GEP->getNumIndices() < 2)
          continue;
        auto *CI1 = dyn_cast<ConstantInt>(GEP->getOperand(2));
        if (!CI1)
          continue;
        uint64_t FieldIndex = CI1->getZExtValue();
        for (User *GEPUser : GEP->users()) {
          if (auto *Store = dyn_cast<StoreInst>(GEPUser)) {
            if (FieldIndex == 0)
              ParamSizeVal = dyn_cast<ConstantInt>(Store->getValueOperand());
            else if (FieldIndex == 1)
              ParamAlignVal = dyn_cast<ConstantInt>(Store->getValueOperand());
          }
        }
      }
      if (ParamSizeVal && ParamAlignVal)
        break;
    }
    Type *ParamType = getBlockStructType(BlockLiteral);
    // Fallback to default if not found
    if (!ParamSizeVal)
      ParamSizeVal = getInt32(M, DL.getTypeStoreSize(ParamType));
    if (!ParamAlignVal)
      ParamAlignVal = getInt32(M, DL.getPrefTypeAlign(ParamType).value());

    Args.push_back(ParamSizeVal);
    Args.push_back(ParamAlignVal);
  }
  // Local sizes arguments: Sizes of block invoke arguments
  // Clang 6.0 and higher generates local size operands as an array,
  // so we need to unpack them
  if (DemangledName.find("_varargs") != StringRef::npos) {
    const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
    auto *LocalSizeArray =
        cast<GetElementPtrInst>(CI->getArgOperand(LocalSizeArrayIdx));
    auto *LocalSizeArrayTy =
        cast<ArrayType>(LocalSizeArray->getSourceElementType());
    const uint64_t LocalSizeNum = LocalSizeArrayTy->getNumElements();
    for (unsigned I = 0; I < LocalSizeNum; ++I)
      Args.push_back(GetElementPtrInst::Create(
          LocalSizeArray->getSourceElementType(), // Pointee type
          LocalSizeArray->getPointerOperand(),    // Alloca
          {getInt32(M, 0), getInt32(M, I)},       // Indices
          "", CI->getIterator()));
  }

  StringRef NewName = "__spirv_EnqueueKernel__";
  FunctionType *FT = FunctionType::get(
      CI->getType(), getTypes(ArrayRef<Value *>(Args)), false /*isVarArg*/);
  Function *NewF =
      Function::Create(FT, GlobalValue::ExternalLinkage, NewName, M);
  NewF->setCallingConv(CallingConv::SPIR_FUNC);
  CallInst *NewCall = CallInst::Create(NewF, Args, "", CI->getIterator());
  NewCall->setCallingConv(NewF->getCallingConv());
  CI->replaceAllUsesWith(NewCall);
  CI->eraseFromParent();
}

void OCLToSPIRVBase::visitCallKernelQuery(CallInst *CI,
                                          StringRef DemangledName) {
  const DataLayout &DL = M->getDataLayout();
  bool HasNDRange = DemangledName.find("_for_ndrange_impl") != StringRef::npos;
  // BIs with "_for_ndrange_impl" suffix has NDRange argument first, and
  // Invoke argument following. For other BIs Invoke function is the first arg
  const unsigned BlockFIdx = HasNDRange ? 1 : 0;
  Value *BlockFVal = CI->getArgOperand(BlockFIdx)->stripPointerCasts();

  auto *BlockF = cast<Function>(getUnderlyingObject(BlockFVal));

  AttributeList Attrs = CI->getCalledFunction()->getAttributes();
  ::mutateCallInst(
      M, CI,
      [=](CallInst *CI, std::vector<Value *> &Args) {
        Value *Param = *Args.rbegin();
        Type *ParamType = getBlockStructType(Param);
        // Last arg corresponds to SPIRV Param operand.
        // Insert Invoke in front of Param.
        // Add Param Size and Param Align at the end.
        Args[BlockFIdx] = BlockF;
        Args.push_back(getInt32(M, DL.getTypeStoreSize(ParamType)));
        Args.push_back(getInt32(M, DL.getPrefTypeAlign(ParamType).value()));

        Op Opcode = OCLSPIRVBuiltinMap::map(DemangledName.str());
        // Adding "__" postfix, so in case we have multiple such
        // functions and their names will have numerical postfix,
        // then the numerical postfix will be droped and we will get
        // correct function name.
        return getSPIRVFuncName(Opcode, kSPIRVName::Postfix);
      },
      /*BuiltinFuncMangleInfo*/ nullptr, &Attrs);
}

// Add postfix to overloaded intel subgroup block read/write builtins
// so new functions can be distinguished.
void OCLToSPIRVBase::processSubgroupBlockReadWriteINTEL(
    CallInst *CI, OCLBuiltinTransInfo &Info, const Type *DataTy) {
  unsigned VectorNumElements = 1;
  if (auto *VecTy = dyn_cast<FixedVectorType>(DataTy))
    VectorNumElements = VecTy->getNumElements();
  unsigned ElementBitSize = DataTy->getScalarSizeInBits();
  Info.Postfix = "_";
  Info.Postfix +=
      getIntelSubgroupBlockDataPostfix(ElementBitSize, VectorNumElements);
  assert(CI->getCalledFunction() && "Unexpected indirect call");
  mutateCallInst(CI, Info.UniqName + Info.Postfix);
}

// The intel_sub_group_block_read built-ins are overloaded to support both
// buffers and images, but need to be mapped to distinct SPIR-V instructions.
// Additionally, for block reads, need to distinguish between scalar block
// reads and vector block reads.
void OCLToSPIRVBase::visitSubgroupBlockReadINTEL(CallInst *CI) {
  OCLBuiltinTransInfo Info;
  if (isOCLImageType(getCallValueType(CI, 0)))
    Info.UniqName = getSPIRVFuncName(spv::OpSubgroupImageBlockReadINTEL);
  else
    Info.UniqName = getSPIRVFuncName(spv::OpSubgroupBlockReadINTEL);
  Type *DataTy = CI->getType();
  processSubgroupBlockReadWriteINTEL(CI, Info, DataTy);
}

// The intel_sub_group_block_write built-ins are similarly overloaded to support
// both buffers and images but need to be mapped to distinct SPIR-V
// instructions.
void OCLToSPIRVBase::visitSubgroupBlockWriteINTEL(CallInst *CI) {
  OCLBuiltinTransInfo Info;
  if (isOCLImageType(getCallValueType(CI, 0)))
    Info.UniqName = getSPIRVFuncName(spv::OpSubgroupImageBlockWriteINTEL);
  else
    Info.UniqName = getSPIRVFuncName(spv::OpSubgroupBlockWriteINTEL);
  assert(!CI->arg_empty() &&
         "Intel subgroup block write should have arguments");
  unsigned DataArg = CI->arg_size() - 1;
  Type *DataTy = CI->getArgOperand(DataArg)->getType();
  processSubgroupBlockReadWriteINTEL(CI, Info, DataTy);
}

void OCLToSPIRVBase::visitSubgroupImageMediaBlockINTEL(
    CallInst *CI, StringRef DemangledName) {
  spv::Op OpCode = DemangledName.rfind("read") != StringRef::npos
                       ? spv::OpSubgroupImageMediaBlockReadINTEL
                       : spv::OpSubgroupImageMediaBlockWriteINTEL;
  // Move the last argument to the beginning.
  mutateCallInst(CI, getSPIRVFuncName(OpCode, CI->getType()))
      .moveArg(CI->arg_size() - 1, 0);
}

static const char *getSubgroupAVCIntelOpKind(StringRef Name) {
  return StringSwitch<const char *>(Name.data())
      .StartsWith(kOCLSubgroupsAVCIntel::IMEPrefix, "ime")
      .StartsWith(kOCLSubgroupsAVCIntel::REFPrefix, "ref")
      .StartsWith(kOCLSubgroupsAVCIntel::SICPrefix, "sic");
}

static const char *getSubgroupAVCIntelTyKind(StringRef MangledName) {
  // We're looking for the type name of the last parameter, which will be at the
  // very end of the mangled name. Since we only care about the ending of the
  // name, we don't need to be any more clever than this.
  return MangledName.ends_with("_payload_t") ? "payload" : "result";
}

static Type *getSubgroupAVCIntelMCEType(Module *M, std::string &TName) {
  auto *Ty = StructType::getTypeByName(M->getContext(), TName);
  if (Ty)
    return Ty;

  return StructType::create(M->getContext(), TName);
}

static Op getSubgroupAVCIntelMCEOpCodeForWrapper(StringRef DemangledName) {
  if (DemangledName.size() <= strlen(kOCLSubgroupsAVCIntel::MCEPrefix))
    return OpNop; // this is not a VME built-in

  std::string MCEName{DemangledName};
  MCEName.replace(0, strlen(kOCLSubgroupsAVCIntel::MCEPrefix),
                  kOCLSubgroupsAVCIntel::MCEPrefix);
  Op MCEOC = OpNop;
  OCLSPIRVSubgroupAVCIntelBuiltinMap::find(MCEName, &MCEOC);
  return MCEOC;
}

// Handles Subgroup AVC Intel extension generic built-ins.
void OCLToSPIRVBase::visitSubgroupAVCBuiltinCall(CallInst *CI,
                                                 StringRef DemangledName) {
  Op OC = OpNop;
  std::string FName{DemangledName};
  std::string Prefix = kOCLSubgroupsAVCIntel::Prefix;

  // Update names for built-ins mapped on two or more SPIRV instructions
  if (FName.find(Prefix + "ime_get_streamout_major_shape_") == 0) {
    // _single_reference functions have 2 arguments, _dual_reference have 3
    // arguments.
    FName += (CI->arg_size() == 2) ? "_single_reference" : "_dual_reference";
  } else if (FName.find(Prefix + "sic_configure_ipe") == 0) {
    FName += (CI->arg_size() == 8) ? "_luma" : "_luma_chroma";
  }

  OCLSPIRVSubgroupAVCIntelBuiltinMap::find(FName, &OC);
  if (OC == OpNop) {
    if (Op MCEOC = getSubgroupAVCIntelMCEOpCodeForWrapper(DemangledName))
      // The called function is a VME wrapper built-in
      return visitSubgroupAVCWrapperBuiltinCall(CI, MCEOC, DemangledName);
    else
      // The called function isn't a VME built-in
      return;
  }

  mutateCallInst(CI, OC);
}

// Handles Subgroup AVC Intel extension wrapper built-ins.
// 'IME', 'REF' and 'SIC' sets contain wrapper built-ins which don't have
// corresponded instructions in SPIRV and should be translated to a
// conterpart from 'MCE' with conversion for an argument and result (if needed).
void OCLToSPIRVBase::visitSubgroupAVCWrapperBuiltinCall(
    CallInst *CI, Op WrappedOC, StringRef DemangledName) {
  std::string Prefix = kOCLSubgroupsAVCIntel::Prefix;

  // Find 'to_mce' conversion function.
  // The operand required conversion is always the last one.
  const char *OpKind = getSubgroupAVCIntelOpKind(DemangledName);
  const char *TyKind =
      getSubgroupAVCIntelTyKind(CI->getCalledFunction()->getName());
  std::string MCETName =
      std::string(kOCLSubgroupsAVCIntel::TypePrefix) + "mce_" + TyKind + "_t";
  auto *MCESTy = getSubgroupAVCIntelMCEType(M, MCETName);
  auto *MCETy = TypedPointerType::get(MCESTy, SPIRAS_Private);
  std::string ToMCEFName = Prefix + OpKind + "_convert_to_mce_" + TyKind;
  Op ToMCEOC = OpNop;
  OCLSPIRVSubgroupAVCIntelBuiltinMap::find(ToMCEFName, &ToMCEOC);
  assert(ToMCEOC != OpNop && "Invalid Subgroup AVC Intel built-in call");

  if (std::strcmp(TyKind, "payload") == 0) {
    // Wrapper built-ins which take the 'payload_t' argument return it as
    // the result: two conversion calls required.
    std::string FromMCEFName =
        Prefix + "mce_convert_to_" + OpKind + "_" + TyKind;
    Op FromMCEOC = OpNop;
    OCLSPIRVSubgroupAVCIntelBuiltinMap::find(FromMCEFName, &FromMCEOC);
    assert(FromMCEOC != OpNop && "Invalid Subgroup AVC Intel built-in call");

    mutateCallInst(CI, WrappedOC)
        .mapArg(CI->arg_size() - 1,
                [&](IRBuilder<> &Builder, Value *Arg, Type *ParamTy) {
                  // Create conversion function call for the last operand
                  return addSPIRVCallPair(Builder, ToMCEOC, MCETy, {Arg},
                                          {ParamTy});
                })
        .changeReturnType(MCETy, [&](IRBuilder<> &Builder, CallInst *NewCI) {
          // Create conversion function call for the return result
          return addSPIRVCall(Builder, FromMCEOC, CI->getType(), {NewCI},
                              {MCETy});
        });
  } else {
    // Wrapper built-ins which take the 'result_t' argument requires only one
    // conversion for the argument
    mutateCallInst(CI, WrappedOC)
        .mapArg(CI->arg_size() - 1, [&](IRBuilder<> &Builder, Value *Arg,
                                        Type *ParamTy) {
          // Create conversion function call for the last operand
          return addSPIRVCallPair(Builder, ToMCEOC, MCETy, {Arg}, {ParamTy});
        });
  }
}

// Handles Subgroup AVC Intel extension built-ins which take sampler as
// an argument (their SPIR-V counterparts take OpTypeVmeImageIntel instead)
void OCLToSPIRVBase::visitSubgroupAVCBuiltinCallWithSampler(
    CallInst *CI, StringRef DemangledName) {
  std::string FName{DemangledName};
  std::string Prefix = kOCLSubgroupsAVCIntel::Prefix;

  // Update names for built-ins mapped on two or more SPIRV instructions
  if (FName.find(Prefix + "ref_evaluate_with_multi_reference") == 0 ||
      FName.find(Prefix + "sic_evaluate_with_multi_reference") == 0) {
    FName += (CI->arg_size() == 5) ? "_interlaced" : "";
  }

  Op OC = OpNop;
  OCLSPIRVSubgroupAVCIntelBuiltinMap::find(FName, &OC);
  if (OC == OpNop)
    return; // this is not a VME built-in

  SmallVector<Type *, 4> ParamTys;
  [[maybe_unused]] bool DidDemangle =
      getParameterTypes(CI->getCalledFunction(), ParamTys);
  assert(DidDemangle && "Expected SPIR-V builtins to be properly mangled");
  auto *TyIt = std::find_if(ParamTys.begin(), ParamTys.end(), isSamplerTy);
  assert(TyIt != ParamTys.end() && "Invalid Subgroup AVC Intel built-in call");
  unsigned SamplerIndex = TyIt - ParamTys.begin();
  Value *SamplerVal = CI->getOperand(SamplerIndex);
  Type *SamplerTy = ParamTys[SamplerIndex];

  SmallVector<Type *, 4> AdaptedTys;
  for (unsigned I = 0; I < CI->arg_size(); I++)
    AdaptedTys.push_back(
        OCLTypeToSPIRVPtr->getAdaptedArgumentType(CI->getCalledFunction(), I));
  auto *AdaptedIter = AdaptedTys.begin();

  mutateCallInst(CI, OC)
      .mapArgs([&](IRBuilder<> &Builder, Value *Arg, Type *ArgTy) {
        if (!isOCLImageType(ArgTy))
          return BuiltinCallMutator::ValueTypePair(Arg, ArgTy);

        auto *ImageTy = *AdaptedIter++;
        if (!ImageTy)
          ImageTy = ArgTy;
        auto *SampledImgTy = adjustImageType(ImageTy, kSPIRVTypeName::Image,
                                             kSPIRVTypeName::VmeImageINTEL);

        Value *SampledImgArgs[] = {Arg, SamplerVal};
        return addSPIRVCallPair(Builder, OpVmeImageINTEL, SampledImgTy,
                                SampledImgArgs, {ArgTy, SamplerTy},
                                kSPIRVName::TempSampledImage);
      })
      .removeArg(SamplerIndex);
}

void OCLToSPIRVBase::visitCallSplitBarrierINTEL(CallInst *CI,
                                                StringRef DemangledName) {
  auto Lit = getBarrierLiterals(CI);
  Op OpCode =
      StringSwitch<Op>(DemangledName)
          .Case("intel_work_group_barrier_arrive", OpControlBarrierArriveINTEL)
          .Case("intel_work_group_barrier_wait", OpControlBarrierWaitINTEL)
          .Default(OpNop);

  // Map memory semantics as follows:
  // OpControlBarrierArriveINTEL -> Release,
  // OpControlBarrierWaitINTEL -> Acquire
  unsigned MemFenceFlag = std::get<0>(Lit);
  OCLMemOrderKind MemOrder =
      OpCode == OpControlBarrierArriveINTEL ? OCLMO_release : OCLMO_acquire;
  mutateCallInst(CI, OpCode)
      .removeArgs(0, CI->arg_size())
      // Execution scope
      .appendArg(addInt32(map<Scope>(std::get<2>(Lit))))
      // Memory scope
      .appendArg(addInt32(map<Scope>(std::get<1>(Lit))))
      // Memory semantics
      .appendArg(addInt32(mapOCLMemSemanticToSPIRV(MemFenceFlag, MemOrder)));
}

void OCLToSPIRVBase::visitCallLdexp(CallInst *CI, StringRef MangledName,
                                    StringRef DemangledName) {
  auto Args = getArguments(CI);
  if (Args.size() == 2) {
    Type *Type0 = Args[0]->getType();
    Type *Type1 = Args[1]->getType();
    // For OpenCL built-in math functions 'halfn ldexp(halfn x, int k)',
    // 'floatn ldexp(floatn x, int k)' and 'doublen ldexp (doublen x, int k)',
    // convert scalar arg to vector to keep consistency with SPIRV spec.
    // Regarding to SPIRV OpenCL Extended Instruction set, k operand must have
    // the same component count as Result Type and x operands
    if (auto *FixedVecType0 = dyn_cast<FixedVectorType>(Type0)) {
      auto ScalarTypeID = Type0->getScalarType()->getTypeID();
      if ((ScalarTypeID == llvm::Type::FloatTyID ||
           ScalarTypeID == llvm::Type::DoubleTyID ||
           ScalarTypeID == llvm::Type::HalfTyID) &&
          Type1->isIntegerTy()) {
        IRBuilder<> IRB(CI);
        unsigned Width = FixedVecType0->getNumElements();
        CI->setOperand(1, IRB.CreateVectorSplat(Width, CI->getArgOperand(1)));
      }
    }
  }
  visitCallBuiltinSimple(CI, MangledName, DemangledName);
}

void OCLToSPIRVBase::visitCallConvertBFloat16AsUshort(CallInst *CI,
                                                      StringRef DemangledName) {
  Type *RetTy = CI->getType();
  Type *ArgTy = CI->getOperand(0)->getType();
  if (DemangledName == kOCLBuiltinName::ConvertBFloat16AsUShort) {
    if (!RetTy->isIntegerTy(16U) || !ArgTy->isFloatTy())
      report_fatal_error(
          "OpConvertBFloat16AsUShort must be of i16 and take float");
  } else {
    FixedVectorType *RetTyVec = cast<FixedVectorType>(RetTy);
    FixedVectorType *ArgTyVec = cast<FixedVectorType>(ArgTy);
    if (!RetTyVec || !RetTyVec->getElementType()->isIntegerTy(16U) ||
        !ArgTyVec || !ArgTyVec->getElementType()->isFloatTy())
      report_fatal_error("OpConvertBFloat16NAsUShortN must be of <N x i16> and "
                         "take <N x float>");
    unsigned RetTyVecSize = RetTyVec->getNumElements();
    unsigned ArgTyVecSize = ArgTyVec->getNumElements();
    if (DemangledName == kOCLBuiltinName::ConvertBFloat162AsUShort2) {
      if (RetTyVecSize != 2 || ArgTyVecSize != 2)
        report_fatal_error("ConvertBFloat162AsUShort2 must be of <2 x i16> and "
                           "take <2 x float>");
    } else if (DemangledName == kOCLBuiltinName::ConvertBFloat163AsUShort3) {
      if (RetTyVecSize != 3 || ArgTyVecSize != 3)
        report_fatal_error("ConvertBFloat163AsUShort3 must be of <3 x i16> and "
                           "take <3 x float>");
    } else if (DemangledName == kOCLBuiltinName::ConvertBFloat164AsUShort4) {
      if (RetTyVecSize != 4 || ArgTyVecSize != 4)
        report_fatal_error("ConvertBFloat164AsUShort4 must be of <4 x i16> and "
                           "take <4 x float>");
    } else if (DemangledName == kOCLBuiltinName::ConvertBFloat168AsUShort8) {
      if (RetTyVecSize != 8 || ArgTyVecSize != 8)
        report_fatal_error("ConvertBFloat168AsUShort8 must be of <8 x i16> and "
                           "take <8 x float>");
    } else if (DemangledName == kOCLBuiltinName::ConvertBFloat1616AsUShort16) {
      if (RetTyVecSize != 16 || ArgTyVecSize != 16)
        report_fatal_error("ConvertBFloat1616AsUShort16 must be of <16 x i16> "
                           "and take <16 x float>");
    }
  }

  mutateCallInst(CI, OpConvertFToBF16INTEL);
}

void OCLToSPIRVBase::visitCallConvertAsBFloat16Float(CallInst *CI,
                                                     StringRef DemangledName) {
  Type *RetTy = CI->getType();
  Type *ArgTy = CI->getOperand(0)->getType();
  if (DemangledName == kOCLBuiltinName::ConvertAsBFloat16Float) {
    if (!RetTy->isFloatTy() || !ArgTy->isIntegerTy(16U))
      report_fatal_error(
          "OpConvertAsBFloat16Float must be of float and take i16");
  } else {
    FixedVectorType *RetTyVec = cast<FixedVectorType>(RetTy);
    FixedVectorType *ArgTyVec = cast<FixedVectorType>(ArgTy);
    if (!RetTyVec || !RetTyVec->getElementType()->isFloatTy() || !ArgTyVec ||
        !ArgTyVec->getElementType()->isIntegerTy(16U))
      report_fatal_error("OpConvertAsBFloat16NFloatN must be of <N x float> "
                         "and take <N x i16>");
    unsigned RetTyVecSize = RetTyVec->getNumElements();
    unsigned ArgTyVecSize = ArgTyVec->getNumElements();
    if (DemangledName == kOCLBuiltinName::ConvertAsBFloat162Float2) {
      if (RetTyVecSize != 2 || ArgTyVecSize != 2)
        report_fatal_error("ConvertAsBFloat162Float2 must be of <2 x float> "
                           "and take <2 x i16>");
    } else if (DemangledName == kOCLBuiltinName::ConvertAsBFloat163Float3) {
      if (RetTyVecSize != 3 || ArgTyVecSize != 3)
        report_fatal_error("ConvertAsBFloat163Float3 must be of <3 x float> "
                           "and take <3 x i16>");
    } else if (DemangledName == kOCLBuiltinName::ConvertAsBFloat164Float4) {
      if (RetTyVecSize != 4 || ArgTyVecSize != 4)
        report_fatal_error("ConvertAsBFloat164Float4 must be of <4 x float> "
                           "and take <4 x i16>");
    } else if (DemangledName == kOCLBuiltinName::ConvertAsBFloat168Float8) {
      if (RetTyVecSize != 8 || ArgTyVecSize != 8)
        report_fatal_error("ConvertAsBFloat168Float8 must be of <8 x float> "
                           "and take <8 x i16>");
    } else if (DemangledName == kOCLBuiltinName::ConvertAsBFloat1616Float16) {
      if (RetTyVecSize != 16 || ArgTyVecSize != 16)
        report_fatal_error("ConvertAsBFloat1616Float16 must be of <16 x float> "
                           "and take <16 x i16>");
    }
  }

  mutateCallInst(CI, OpConvertBF16ToFINTEL);
}
} // namespace SPIRV

INITIALIZE_PASS_BEGIN(OCLToSPIRVLegacy, "ocl-to-spv",
                      "Transform OCL 2.0 to SPIR-V", false, false)
INITIALIZE_PASS_DEPENDENCY(OCLTypeToSPIRVLegacy)
INITIALIZE_PASS_END(OCLToSPIRVLegacy, "ocl-to-spv",
                    "Transform OCL 2.0 to SPIR-V", false, false)

ModulePass *llvm::createOCLToSPIRVLegacy() { return new OCLToSPIRVLegacy(); }