lib/SPIRV/OCL20ToSPIRV.cpp

//===- OCL20ToSPIRV.cpp - Transform OCL20 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 translation of OCL20 builtin functions.
//
//===----------------------------------------------------------------------===//
#define DEBUG_TYPE "cl20tospv"

#include "SPIRVInternal.h"
#include "OCLUtil.h"
#include "OCLTypeToSPIRV.h"

#include "llvm/ADT/StringSwitch.h"
#include "llvm/IR/InstVisitor.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Instruction.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Verifier.h"
#include "llvm/Pass.h"
#include "llvm/PassSupport.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/raw_ostream.h"

#include <set>

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

namespace SPIRV {
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);
}

class OCL20ToSPIRV: public ModulePass,
  public InstVisitor<OCL20ToSPIRV> {
public:
  OCL20ToSPIRV():ModulePass(ID), M(nullptr), Ctx(nullptr), CLVer(0) {
    initializeOCL20ToSPIRVPass(*PassRegistry::getPassRegistry());
  }
  virtual bool runOnModule(Module &M);

  void getAnalysisUsage(AnalysisUsage &AU) const {
    AU.addRequired<OCLTypeToSPIRV>();
  }

  virtual void visitCallInst(CallInst &CI);

  /// Transform barrier/work_group_barrier/sub_group_barrier
  ///     to __spirv_ControlBarrier.
  /// barrier(flag) =>
  ///   __spirv_ControlBarrier(workgroup, workgroup, map(flag))
  /// work_group_barrier(scope, flag) =>
  ///   __spirv_ControlBarrier(workgroup, map(scope), map(flag))
  /// sub_group_barrier(scope, flag) =>
  ///   __spirv_ControlBarrier(subgroup, map(scope), map(flag))
  void visitCallBarrier(CallInst *CI);

  /// Erase useless convert functions.
  /// \return true if the call instruction is erased.
  bool eraseUselessConvert(CallInst *Call, const std::string &MangledName,
      const std::string &DeMangledName);

  /// Transform convert_ to
  ///   __spirv_{CastOpName}_R{TargeTyName}{_sat}{_rt[p|n|z|e]}
  void visitCallConvert(CallInst *CI, StringRef MangledName,
    const std::string &DemangledName);

  /// Transform async_work_group{_strided}_copy.
  /// async_work_group_copy(dst, src, n, event)
  ///   => async_work_group_strided_copy(dst, src, n, 1, event)
  /// async_work_group_strided_copy(dst, src, n, stride, event)
  ///   => __spirv_AsyncGroupCopy(ScopeWorkGroup, dst, src, n, stride, event)
  void visitCallAsyncWorkGroupCopy(CallInst *CI,
      const std::string &DemangledName);

  /// Transform OCL builtin function to SPIR-V builtin function.
  void transBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info);

  /// Transform OCL work item builtin functions to SPIR-V builtin variables.
  void transWorkItemBuiltinsToVariables();

  /// Transform atomic_work_item_fence/mem_fence to __spirv_MemoryBarrier.
  /// func(flag, order, scope) =>
  ///   __spirv_MemoryBarrier(map(scope), map(flag)|map(order))
  void transMemoryBarrier(CallInst *CI, AtomicWorkItemFenceLiterals);

  /// Transform all to __spirv_Op(All|Any).  Note that the types mismatch so
  // some extra code is emitted to convert between the two.
  void visitCallAllAny(spv::Op OC, CallInst *CI);

  /// Transform atomic_* to __spirv_Atomic*.
  /// atomic_x(ptr_arg, args, order, scope) =>
  ///   __spirv_AtomicY(ptr_arg, map(order), map(scope), args)
  void transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info);

  /// Transform atomic_work_item_fence to __spirv_MemoryBarrier.
  /// atomic_work_item_fence(flag, order, scope) =>
  ///   __spirv_MemoryBarrier(map(scope), map(flag)|map(order))
  void visitCallAtomicWorkItemFence(CallInst *CI);

  /// Transform atomic_compare_exchange call.
  /// In atomic_compare_exchange, the expected value parameter is a pointer.
  /// However in SPIR-V it is a value. The transformation adds a load
  /// instruction, result of which is passed to atomic_compare_exchange as
  /// argument.
  /// The transformation adds a store instruction after the call, to update the
  /// value in expected with the value pointed to by object. Though, it is not
  /// necessary in case they are equal, this approach makes result code simpler.
  /// Also ICmp instruction is added, because the call must return result of
  /// comparison.
  /// \returns the call instruction of atomic_compare_exchange_strong.
  CallInst *visitCallAtomicCmpXchg(CallInst *CI,
      const std::string &DemangledName);

  /// Transform atomic_init.
  /// atomic_init(p, x) => store p, x
  void visitCallAtomicInit(CallInst *CI);

  /// Transform legacy OCL 1.x atomic builtins to SPIR-V builtins for extensions
  ///   cl_khr_int64_base_atomics
  ///   cl_khr_int64_extended_atomics
  /// Do nothing if the called function is not a legacy atomic builtin.
  void visitCallAtomicLegacy(CallInst *CI, StringRef MangledName,
    const std::string &DemangledName);

  /// Transform OCL 2.0 C++11 atomic builtins to SPIR-V builtins.
  /// Do nothing if the called function is not a C++11 atomic builtin.
  void visitCallAtomicCpp11(CallInst *CI, StringRef MangledName,
    const std::string &DemangledName);

  /// Transform OCL builtin function to SPIR-V builtin function.
  /// Assuming there is a simple name mapping without argument changes.
  /// Should be called at last.
  void visitCallBuiltinSimple(CallInst *CI, StringRef MangledName,
    const std::string &DemangledName);

  /// Transform get_image_{width|height|depth|dim}.
  /// get_image_xxx(...) =>
  ///   dimension = __spirv_ImageQuerySizeLod_R{ReturnType}(...);
  ///   return dimension.{x|y|z};
  void visitCallGetImageSize(CallInst *CI, StringRef MangledName,
    const std::string &DemangledName);

  /// Transform {work|sub}_group_x =>
  ///   __spirv_{OpName}
  ///
  /// Special handling of work_group_broadcast.
  ///   work_group_broadcast(a, x, y, z)
  ///     =>
  ///   __spirv_GroupBroadcast(a, vec3(x, y, z))

  void visitCallGroupBuiltin(CallInst *CI, StringRef MangledName,
    const std::string &DemangledName);

  /// Transform mem_fence to __spirv_MemoryBarrier.
  /// mem_fence(flag) => __spirv_MemoryBarrier(Workgroup, map(flag))
  void visitCallMemFence(CallInst *CI);

  void visitCallNDRange(CallInst *CI, const std::string &DemangledName);

  /// Transform OCL pipe builtin function to SPIR-V pipe builtin function.
  void visitCallPipeBuiltin(CallInst *CI, StringRef MangledName,
    const std::string &DemangledName);

  /// Transform read_image with sampler arguments.
  /// read_image(image, sampler, ...) =>
  ///   sampled_image = __spirv_SampledImage(image, sampler);
  ///   return __spirv_ImageSampleExplicitLod_R{ReturnType}(sampled_image, ...);
  void visitCallReadImageWithSampler(CallInst *CI, StringRef MangledName,
      const std::string &DemangledName);

  /// Transform read_image with msaa image arguments.
  /// Sample argument must be acoded as Image Operand.
  void visitCallReadImageMSAA(CallInst *CI, StringRef MangledName,
                              const std::string &DemangledName);

  /// Transform {read|write}_image without sampler arguments.
  void visitCallReadWriteImage(CallInst *CI, StringRef MangledName,
      const std::string &DemangledName);

  /// Transform to_{global|local|private}.
  ///
  /// T* a = ...;
  /// addr T* b = to_addr(a);
  ///   =>
  /// i8* x = cast<i8*>(a);
  /// addr i8* y = __spirv_GenericCastToPtr_ToAddr(x);
  /// addr T* b = cast<addr T*>(y);
  void visitCallToAddr(CallInst *CI, StringRef MangledName,
      const std::string &DemangledName);

  /// Transform return type of relatinal built-in functions like isnan, isfinite
  /// to boolean values.
  void visitCallRelational(CallInst *CI, const std::string &DemangledName);

  /// Transform vector load/store functions to SPIR-V extended builtin
  ///   functions
  /// {vload|vstore{a}}{_half}{n}{_rte|_rtz|_rtp|_rtn} =>
  ///   __spirv_ocl_{ExtendedInstructionOpCodeName}__R{ReturnType}
  void visitCallVecLoadStore(CallInst *CI, StringRef MangledName,
      const std::string &DemangledName);

  /// Transforms get_mem_fence built-in to SPIR-V function and aligns result values with SPIR 1.2.
  /// get_mem_fence(ptr) => __spirv_GenericPtrMemSemantics
  /// GenericPtrMemSemantics valid values are 0x100, 0x200 and 0x300, where is
  /// SPIR 1.2 defines them as 0x1, 0x2 and 0x3, so this function adjusts
  /// GenericPtrMemSemantics results to SPIR 1.2 values.
  void visitCallGetFence(CallInst *CI, StringRef MangledName, const std::string& DemangledName);

  /// Transforms OpDot instructions with a scalar type to a fmul instruction
  void visitCallDot(CallInst *CI);

  /// Fixes for built-in functions with vector+scalar arguments that are
  /// translated to the SPIR-V instructions where all arguments must have the
  /// same type.
  void visitCallScalToVec(CallInst *CI, StringRef MangledName,
                          const std::string &DemangledName);

  /// Transform get_image_channel_{order|data_type} built-in functions to
  ///   __spirv_ocl_{ImageQueryOrder|ImageQueryFormat}
  void visitCallGetImageChannel(CallInst *CI, StringRef MangledName,
                                const std::string &DemangledName,
                                unsigned int Offset);

  void visitDbgInfoIntrinsic(DbgInfoIntrinsic &I){
    I.dropAllReferences();
    I.eraseFromParent();
  }
  static char ID;
private:
  Module *M;
  LLVMContext *Ctx;
  unsigned CLVer;                   /// OpenCL version as major*10+minor
  std::set<Value *> ValuesToDelete;

  ConstantInt *addInt32(int I) {
    return getInt32(M, I);
  }
  ConstantInt *addSizet(uint64_t I) {
    return getSizet(M, I);
  }

  /// Get vector width from OpenCL vload* function name.
  SPIRVWord 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 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 OCL20ToSPIRV::ID = 0;

bool
OCL20ToSPIRV::runOnModule(Module& Module) {
  M = &Module;
  Ctx = &M->getContext();
  auto Src = getSPIRVSource(&Module);
  if (std::get<0>(Src) != spv::SourceLanguageOpenCL_C)
    return false;

  CLVer = std::get<1>(Src);
  if (CLVer > kOCLVer::CL20)
    return false;

  DEBUG(dbgs() << "Enter OCL20ToSPIRV:\n");

  transWorkItemBuiltinsToVariables();

  visit(*M);

  for (auto &I:ValuesToDelete)
    if (auto Inst = dyn_cast<Instruction>(I))
      Inst->eraseFromParent();
  for (auto &I:ValuesToDelete)
    if (auto GV = dyn_cast<GlobalValue>(I))
      GV->eraseFromParent();

  DEBUG(dbgs() << "After OCL20ToSPIRV:\n" << *M);

  std::string Err;
  raw_string_ostream ErrorOS(Err);
  if (verifyModule(*M, &ErrorOS)){
    DEBUG(errs() << "Fails to verify module: " << ErrorOS.str());
  }
  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
OCL20ToSPIRV::visitCallInst(CallInst& CI) {
  DEBUG(dbgs() << "[visistCallInst] " << CI << '\n');
  auto F = CI.getCalledFunction();
  if (!F)
    return;

  auto MangledName = F->getName();
  std::string DemangledName;
  if (!oclIsBuiltin(MangledName, &DemangledName))
    return;

  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) {
    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 && "Wrong version of OpenCL");
      PCI = visitCallAtomicCmpXchg(PCI, DemangledName);
    }
    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, MangledName, DemangledName);
    return;
  }
  if ((DemangledName.find(kOCLBuiltinName::WorkGroupPrefix) == 0 &&
      DemangledName != kOCLBuiltinName::WorkGroupBarrier) ||
      DemangledName == kOCLBuiltinName::WaitGroupEvent ||
      (DemangledName.find(kOCLBuiltinName::SubGroupPrefix) == 0 &&
       DemangledName != kOCLBuiltinName::SubGroupBarrier)) {
    visitCallGroupBuiltin(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName.find(kOCLBuiltinName::Pipe) != std::string::npos) {
    visitCallPipeBuiltin(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::MemFence) {
    visitCallMemFence(&CI);
    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, DemangledName);
      return;
    }
  }
  if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0 ||
      DemangledName.find(kOCLBuiltinName::WriteImage) == 0) {
    visitCallReadWriteImage(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::ToGlobal ||
      DemangledName == kOCLBuiltinName::ToLocal ||
      DemangledName == kOCLBuiltinName::ToPrivate) {
    visitCallToAddr(&CI, MangledName, 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) {
    visitCallBarrier(&CI);
    return;
  }
  if (DemangledName == kOCLBuiltinName::GetFence) {
    visitCallGetFence(&CI, MangledName, DemangledName);
    return;
  }
  if (DemangledName == kOCLBuiltinName::Dot &&
      !(CI.getOperand(0)->getType()->isVectorTy())) {
    visitCallDot(&CI);
    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, MangledName, DemangledName,
                             OCLImageChannelDataTypeOffset);
    return;
  }
  if (DemangledName == kOCLBuiltinName::GetImageChannelOrder) {
    visitCallGetImageChannel(&CI, MangledName, DemangledName,
                             OCLImageChannelOrderOffset);
    return;
  }
  visitCallBuiltinSimple(&CI, MangledName, DemangledName);
}

void
OCL20ToSPIRV::visitCallNDRange(CallInst *CI,
    const std::string &DemangledName) {
  assert(DemangledName.find(kOCLBuiltinName::NDRangePrefix) == 0);
  std::string lenStr = DemangledName.substr(8, 1);
  auto Len = atoi(lenStr.c_str());
  assert (Len >= 1 && Len <= 3);
  // 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.
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
    for (size_t I = 1, E = Args.size(); I != E; ++I)
      Args[I] = getScalarOrArray(Args[I], Len, CI);
    switch (Args.size()) {
    case 2: {
      // Has global work size.
      auto T = Args[1]->getType();
      auto C = getScalarOrArrayConstantInt(CI, T, Len, 0);
      Args.push_back(C);
      Args.push_back(C);
    }
      break;
    case 3: {
      // Has global and local work size.
      auto T = Args[1]->getType();
      Args.push_back(getScalarOrArrayConstantInt(CI, T, Len, 0));
    }
      break;
    case 4: {
      // Move offset arg to the end
      auto OffsetPos = Args.begin() + 1;
      Value* OffsetVal = *OffsetPos;
      Args.erase(OffsetPos);
      Args.push_back(OffsetVal);
    }
      break;
    default:
      assert(0 && "Invalid number of arguments");
    }
    // Translate ndrange_ND into differently named SPIR-V decorated functions because
    // they have array arugments of different dimension which mangled the same way.
    return getSPIRVFuncName(OpBuildNDRange, "_" + lenStr + "D");
  }, &Attrs);
}

void
OCL20ToSPIRV::visitCallAsyncWorkGroupCopy(CallInst* CI,
    const std::string &DemangledName) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
    if (DemangledName == OCLUtil::kOCLBuiltinName::AsyncWorkGroupCopy) {
      Args.insert(Args.begin()+3, addSizet(1));
    }
    Args.insert(Args.begin(), addInt32(ScopeWorkgroup));
    return getSPIRVFuncName(OpGroupAsyncCopy);
  }, &Attrs);
}

CallInst *
OCL20ToSPIRV::visitCallAtomicCmpXchg(CallInst* CI,
    const std::string& DemangledName) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  Value *Expected = nullptr;
  CallInst *NewCI = nullptr;
  mutateCallInstOCL(M, CI, [&](CallInst * CI, std::vector<Value *> &Args,
      Type *&RetTy){
    Expected = Args[1]; // temporary save second argument.
    Args[1] = new LoadInst(Args[1], "exp", false, CI);
    RetTy = Args[2]->getType();
    assert(Args[0]->getType()->getPointerElementType()->isIntegerTy() &&
      Args[1]->getType()->isIntegerTy() && Args[2]->getType()->isIntegerTy() &&
      "In SPIR-V 1.0 arguments of OpAtomicCompareExchange must be "
      "an integer type scalars");
    return kOCLBuiltinName::AtomicCmpXchgStrong;
  },
  [&](CallInst *NCI)->Instruction * {
    NewCI = NCI;
    Instruction* Store = new StoreInst(NCI, Expected, NCI->getNextNode());
    return new ICmpInst(Store->getNextNode(), CmpInst::ICMP_EQ, NCI,
                        NCI->getArgOperand(1));
  },
  &Attrs);
  return NewCI;
}

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

void
OCL20ToSPIRV::visitCallAllAny(spv::Op OC, CallInst* CI) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();

  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);

  if (!isa<VectorType>(ArgTy)) {
    auto *Cast = CastInst::CreateZExtOrBitCast(Cmp, Type::getInt32Ty(*Ctx),
                                                "", Cmp->getNextNode());
    CI->replaceAllUsesWith(Cast);
    CI->eraseFromParent();
  } else {
    mutateCallInstSPIRV(
        M, CI,
        [&](CallInst *, std::vector<Value *> &Args, Type *&Ret) {
          Args[0] = Cmp;
          Ret = Type::getInt1Ty(*Ctx);

          return getSPIRVFuncName(OC);
        },
        [&](CallInst *CI) -> Instruction * {
          return CastInst::CreateZExtOrBitCast(CI, Type::getInt32Ty(*Ctx), "",
                                               CI->getNextNode());
        },
        &Attrs);
  }
}

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

void
OCL20ToSPIRV::visitCallMemFence(CallInst* CI) {
  transMemoryBarrier(CI, std::make_tuple(
      cast<ConstantInt>(CI->getArgOperand(0))->getZExtValue(),
      OCLMO_relaxed,
      OCLMS_work_group));
}

void OCL20ToSPIRV::transMemoryBarrier(CallInst* CI,
    AtomicWorkItemFenceLiterals Lit) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
    Args.resize(2);
    Args[0] = addInt32(map<Scope>(std::get<2>(Lit)));
    Args[1] = addInt32(mapOCLMemSemanticToSPIRV(std::get<0>(Lit),
        std::get<1>(Lit)));
    return getSPIRVFuncName(OpMemoryBarrier);
  }, &Attrs);
}

void
OCL20ToSPIRV::visitCallAtomicLegacy(CallInst* CI,
    StringRef MangledName, const std::string& DemangledName) {
  StringRef Stem = DemangledName;
  if (Stem.startswith("atom_"))
    Stem = Stem.drop_front(strlen("atom_"));
  else if (Stem.startswith("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.startswith("compare_exchange"))
    PostOps.push_back(OCLLegacyAtomicMemOrder);
  PostOps.push_back(OCLLegacyAtomicMemScope);

  Info.PostProc = [=](std::vector<Value *> &Ops){
    for (auto &I:PostOps){
      Ops.push_back(addInt32(I));
    }
  };
  transAtomicBuiltin(CI, Info);
}

void
OCL20ToSPIRV::visitCallAtomicCpp11(CallInst* CI,
    StringRef MangledName, const std::string& DemangledName) {
  StringRef Stem = DemangledName;
  if (Stem.startswith("atomic_"))
    Stem = Stem.drop_front(strlen("atomic_"));
  else
    return;

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

    if (!Stem.endswith("_explicit")) {
      NewStem = NewStem + "_explicit";
      PostOps.push_back(OCLMO_seq_cst);
      if (Stem.startswith("compare_exchange"))
        PostOps.push_back(OCLMO_seq_cst);
      PostOps.push_back(OCLMS_device);
    } else {
      auto MaxOps = getOCLCpp11AtomicMaxNumOps(
          Stem.drop_back(strlen("_explicit")));
      if (CI->getNumArgOperands() < 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 = [=](std::vector<Value *> &Ops){
    for (auto &I:PostOps){
      Ops.push_back(addInt32(I));
    }
  };

  transAtomicBuiltin(CI, Info);
}

void
OCL20ToSPIRV::transAtomicBuiltin(CallInst* CI,
    OCLBuiltinTransInfo& Info) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(M, CI, [=](CallInst * CI, std::vector<Value *> &Args){
    Info.PostProc(Args);
    // 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 = Args.size();
    const size_t ScopeIdx = ArgsCount - 1;
    const size_t OrderIdx = ScopeIdx - NumOrder;
    Args[ScopeIdx] = mapUInt(M, cast<ConstantInt>(Args[ScopeIdx]),
        [](unsigned I){
      return map<Scope>(static_cast<OCLScopeKind>(I));
    });
    for (size_t I = 0; I < NumOrder; ++I)
      Args[OrderIdx + I] = mapUInt(M, cast<ConstantInt>(Args[OrderIdx + I]),
          [](unsigned Ord) {
      return mapOCLMemSemanticToSPIRV(0, static_cast<OCLMemOrderKind>(Ord));
    });
    // Order of args in SPIR-V:
    // object, scope, 1-2 order, 0-2 other args
    std::swap(Args[1], Args[ScopeIdx]);
    if(OrderIdx > 2) {
      // For atomic_compare_exchange the swap above puts Comparator/Expected
      // argument just where it should be, so don't move the last argument then.
      int offset = Info.UniqName.find("atomic_compare_exchange") == 0 ? 1 : 0;
      std::rotate(Args.begin() + 2, Args.begin() + OrderIdx,
                  Args.end() - offset);
    }
    return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
  }, &Attrs);
}

void
OCL20ToSPIRV::visitCallBarrier(CallInst* CI) {
  auto Lit = getBarrierLiterals(CI);
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
    Args.resize(3);
    Args[0] = addInt32(map<Scope>(std::get<2>(Lit)));
    Args[1] = addInt32(map<Scope>(std::get<1>(Lit)));
    Args[2] = addInt32(mapOCLMemFenceFlagToSPIRV(std::get<0>(Lit)));
    return getSPIRVFuncName(OpControlBarrier);
  }, &Attrs);
}

void OCL20ToSPIRV::visitCallConvert(CallInst* CI,
    StringRef MangledName, const std::string& DemangledName) {
  if (eraseUselessConvert(CI, MangledName, DemangledName))
    return;
  Op OC = OpNop;
  auto TargetTy = CI->getType();
  auto SrcTy = CI->getArgOperand(0)->getType();
  if (isa<VectorType>(TargetTy))
    TargetTy = TargetTy->getVectorElementType();
  if (isa<VectorType>(SrcTy))
    SrcTy = SrcTy->getVectorElementType();
  auto IsTargetInt = isa<IntegerType>(TargetTy);

  std::string TargetTyName = DemangledName.substr(
      strlen(kOCLBuiltinName::ConvertPrefix));
  auto FirstUnderscoreLoc = TargetTyName.find('_');
  if (FirstUnderscoreLoc != std::string::npos)
    TargetTyName = TargetTyName.substr(0, FirstUnderscoreLoc);
  TargetTyName = std::string("_R") + TargetTyName;

  std::string Sat = DemangledName.find("_sat") != std::string::npos ?
      "_sat" : "";
  auto TargetSigned = DemangledName[8] != '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;
  }
  auto Loc = DemangledName.find("_rt");
  std::string Rounding;
  if (Loc != std::string::npos &&
      !(isa<IntegerType>(SrcTy) && IsTargetInt)) {
    Rounding = DemangledName.substr(Loc, 4);
  }
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
    return getSPIRVFuncName(OC, TargetTyName + Sat + Rounding);
  }, &Attrs);
}

void OCL20ToSPIRV::visitCallGroupBuiltin(CallInst* CI,
    StringRef MangledName, const std::string& 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 GroupOp = DemangledName;
    GroupOp = GroupOp.drop_front(strlen(kSPIRVName::GroupPrefix));
    SPIRSPIRVGroupOperationMap::foreach_conditional([&](const std::string &S,
        SPIRVGroupOperationKind G){
      if (!GroupOp.startswith(S))
        return true; // continue
      PreOps.push_back(G);
      StringRef Op = GroupOp.drop_front(S.size() + 1);
      assert(!Op.empty() && "Invalid OpenCL group builtin function");
      char OpTyC = 0;
      auto NeedSign = Op == "max" || Op == "min";
      auto OpTy = F->getReturnType();
      if (OpTy->isFloatingPointTy())
        OpTyC = 'f';
      else if (OpTy->isIntegerTy()) {
        if (!NeedSign)
          OpTyC = 'i';
        else {
          if (isLastFuncParamSigned(F->getName()))
            OpTyC = 's';
          else
            OpTyC = 'u';
        }
      } else
        llvm_unreachable("Invalid OpenCL group builtin argument type");

      DemangledName = std::string(kSPIRVName::GroupPrefix) + OpTyC + Op.str();
      return false; // break out of loop
    });
  }

  bool IsGroupAllAny = (DemangledName.find("_all") != std::string::npos ||
                        DemangledName.find("_any") != std::string::npos);

  auto Consts = getInt32(M, PreOps);
  OCLBuiltinTransInfo Info;
  if (IsGroupAllAny)
    Info.RetTy = Type::getInt1Ty(*Ctx);
  Info.UniqName = DemangledName;
  Info.PostProc = [=](std::vector<Value *> &Ops) {
    if (IsGroupAllAny) {
      IRBuilder<> IRB(CI);
      Ops[0] =
          IRB.CreateICmpNE(Ops[0], ConstantInt::get(Type::getInt32Ty(*Ctx), 0));
    }
    size_t E = Ops.size();
    if (DemangledName == "group_broadcast" && E > 2) {
      assert(E == 3 || E == 4);
      makeVector(CI, Ops, std::make_pair(Ops.begin() + 1, Ops.end()));
    }
    Ops.insert(Ops.begin(), Consts.begin(), Consts.end());
  };
  transBuiltin(CI, Info);
}

void
OCL20ToSPIRV::transBuiltin(CallInst* CI,
    OCLBuiltinTransInfo& Info) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  Op OC = OpNop;
  unsigned ExtOp = ~0U;
  if (StringRef(Info.UniqName).startswith(kSPIRVName::Prefix))
      return;
  if (OCLSPIRVBuiltinMap::find(Info.UniqName, &OC))
    Info.UniqName = getSPIRVFuncName(OC);
  else if ((ExtOp = getExtOp(Info.MangledName, Info.UniqName)) != ~0U)
    Info.UniqName = getSPIRVExtFuncName(SPIRVEIS_OpenCL, ExtOp);
  else
    return;
  if (!Info.RetTy)
    mutateCallInstSPIRV(M, CI,
                        [=](CallInst *, std::vector<Value *> &Args) {
                          Info.PostProc(Args);
                          return Info.UniqName + Info.Postfix;
                        },
                        &Attrs);
  else
    mutateCallInstSPIRV(
        M, CI,
        [=](CallInst *, std::vector<Value *> &Args, Type *&RetTy) {
          Info.PostProc(Args);
          RetTy = Info.RetTy;
          return Info.UniqName + Info.Postfix;
        },
        [=](CallInst *NewCI) -> Instruction * {
          if (NewCI->getType()->isIntegerTy() && CI->getType()->isIntegerTy())
            return CastInst::CreateIntegerCast(NewCI, CI->getType(),
                                               Info.isRetSigned, "", CI);
          else
            return CastInst::CreatePointerBitCastOrAddrSpaceCast(
                NewCI, CI->getType(), "", CI);
        },
        &Attrs);
}

void
OCL20ToSPIRV::visitCallPipeBuiltin(CallInst* CI,
    StringRef MangledName, const std::string& DemangledName) {
  std::string NewName = DemangledName;
  // Transform OpenCL read_pipe/write_pipe builtin function names
  // with reserve_id argument to reserved_read_pipe/reserved_write_pipe.
  if ((DemangledName.find(kOCLBuiltinName::ReadPipe) == 0 ||
      DemangledName.find(kOCLBuiltinName::WritePipe) == 0)
      && CI->getNumArgOperands() > 4)
    NewName = std::string(kSPIRVName::ReservedPrefix) + DemangledName;
  OCLBuiltinTransInfo Info;
  Info.UniqName = NewName;
  transBuiltin(CI, Info);
}

void OCL20ToSPIRV::visitCallReadImageMSAA(CallInst *CI, StringRef MangledName,
                                          const std::string &DemangledName) {
  assert(MangledName.find("msaa") != StringRef::npos);
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(
      M, CI,
      [=](CallInst *, std::vector<Value *> &Args) {
        Args.insert(Args.begin() + 2, getInt32(M, ImageOperandsSampleMask));
        return getSPIRVFuncName(OpImageRead,
                                std::string(kSPIRVPostfix::ExtDivider) +
                                    getPostfixForReturnType(CI));
      },
      &Attrs);
}

void OCL20ToSPIRV::visitCallReadImageWithSampler(
    CallInst *CI, StringRef MangledName, const std::string &DemangledName) {
  assert (MangledName.find(kMangledName::Sampler) != StringRef::npos);
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  bool isRetScalar = !CI->getType()->isVectorTy();
  mutateCallInstSPIRV(
      M, CI,
      [=](CallInst *, std::vector<Value *> &Args, Type *&Ret) {
        auto ImageTy = getAnalysis<OCLTypeToSPIRV>().getAdaptedType(Args[0]);
        if (isOCLImageType(ImageTy))
          ImageTy = getSPIRVImageTypeFromOCL(M, ImageTy);
        auto SampledImgTy = getSPIRVTypeByChangeBaseTypeName(
            M, ImageTy, kSPIRVTypeName::Image, kSPIRVTypeName::SampledImg);
        Value *SampledImgArgs[] = {Args[0], Args[1]};
        auto SampledImg = addCallInstSPIRV(
            M, getSPIRVFuncName(OpSampledImage), SampledImgTy, SampledImgArgs,
            nullptr, CI, kSPIRVName::TempSampledImage);

        Args[0] = SampledImg;
        Args.erase(Args.begin() + 1, Args.begin() + 2);

        switch (Args.size()) {
        case 2: // no lod
          Args.push_back(getInt32(M, ImageOperandsMask::ImageOperandsLodMask));
          Args.push_back(getFloat32(M, 0.f));
          break;
        case 3: // explicit lod
          Args.insert(Args.begin() + 2,
                      getInt32(M, ImageOperandsMask::ImageOperandsLodMask));
          break;
        case 4: // gradient
          Args.insert(Args.begin() + 2,
                      getInt32(M, ImageOperandsMask::ImageOperandsGradMask));
          break;
        default:
          assert(0 && "read_image* with unhandled number of args!");
        }

        // SPIR-V intruction always returns 4-element vector
        if (isRetScalar)
          Ret = VectorType::get(Ret, 4);
        return getSPIRVFuncName(OpImageSampleExplicitLod,
                                std::string(kSPIRVPostfix::ExtDivider) +
                                    getPostfixForReturnType(Ret));
      },
      [&](CallInst *CI) -> Instruction * {
        if (isRetScalar)
          return ExtractElementInst::Create(CI, getSizet(M, 0), "",
                                            CI->getNextNode());
        return CI;
      },
      &Attrs);
}

void
OCL20ToSPIRV::visitCallGetImageSize(CallInst* CI,
    StringRef MangledName, const std::string& DemangledName) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  StringRef TyName;
  SmallVector<StringRef, 4> SubStrs;
  auto IsImg = isOCLImageType(CI->getArgOperand(0)->getType(), &TyName);
  (void)IsImg;  // prevent warning about unused variable in NDEBUG build
  assert(IsImg);
  std::string ImageTyName = TyName.str();
  if (hasAccessQualifiedName(TyName))
    ImageTyName.erase(ImageTyName.size() - 5, 3);
  auto Desc = map<SPIRVTypeImageDescriptor>(ImageTyName);
  unsigned Dim = getImageDimension(Desc.Dim) + Desc.Arrayed;
  assert(Dim > 0 && "Invalid image dimension.");
  mutateCallInstSPIRV(M, CI,
    [&](CallInst *, std::vector<Value *> &Args, Type *&Ret){
      assert(Args.size() == 1);
      Ret = CI->getType()->isIntegerTy(64) ? Type::getInt64Ty(*Ctx)
                                           : Type::getInt32Ty(*Ctx);
      if (Dim > 1)
        Ret = VectorType::get(Ret, Dim);
      if (Desc.Dim == DimBuffer)
        return getSPIRVFuncName(OpImageQuerySize, CI->getType());
      else {
        Args.push_back(getInt32(M, 0));
        return getSPIRVFuncName(OpImageQuerySizeLod, CI->getType());
      }
    },
    [&](CallInst *NCI)->Instruction * {
      if (Dim == 1)
        return NCI;
      if (DemangledName == kOCLBuiltinName::GetImageDim) {
        if (Desc.Dim == Dim3D) {
          auto ZeroVec = ConstantVector::getSplat(3,
            Constant::getNullValue(NCI->getType()->getVectorElementType()));
          Constant *Index[] = {getInt32(M, 0), getInt32(M, 1),
              getInt32(M, 2), getInt32(M, 3)};
          return new ShuffleVectorInst(NCI, ZeroVec,
             ConstantVector::get(Index), "", CI);

        } else if (Desc.Dim == Dim2D && Desc.Arrayed) {
          Constant *Index[] = {getInt32(M, 0), getInt32(M, 1)};
          Constant *mask = ConstantVector::get(Index);
          return new ShuffleVectorInst(NCI, UndefValue::get(NCI->getType()),
                                       mask, NCI->getName(), CI);
        }
        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());
    },
  &Attrs);
}

/// Remove trivial conversion functions
bool
OCL20ToSPIRV::eraseUselessConvert(CallInst *CI,
    const std::string &MangledName,
    const std::string &DemangledName) {
  auto TargetTy = CI->getType();
  auto SrcTy = CI->getArgOperand(0)->getType();
  if (isa<VectorType>(TargetTy))
    TargetTy = TargetTy->getVectorElementType();
  if (isa<VectorType>(SrcTy))
    SrcTy = SrcTy->getVectorElementType();
  if (TargetTy == SrcTy) {
    if (isa<IntegerType>(TargetTy) &&
        DemangledName.find("_sat") != std::string::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);
    ValuesToDelete.insert(CI->getCalledFunction());
    return true;
  }
  return false;
}

void
OCL20ToSPIRV::visitCallBuiltinSimple(CallInst* CI,
    StringRef MangledName, const std::string& DemangledName) {
  OCLBuiltinTransInfo Info;
  Info.MangledName = MangledName.str();
  Info.UniqName = DemangledName;
  transBuiltin(CI, Info);
}

/// Translates OCL work-item builtin functions to SPIRV builtin variables.
/// Function like get_global_id(i) -> x = load GlobalInvocationId; extract x, i
/// Function like get_work_dim() -> load WorkDim
void OCL20ToSPIRV::transWorkItemBuiltinsToVariables() {
  DEBUG(dbgs() << "Enter transWorkItemBuiltinsToVariables\n");
  std::vector<Function *> WorkList;
  for (auto I = M->begin(), E = M->end(); I != E; ++I) {
    std::string DemangledName;
    if (!oclIsBuiltin(I->getName(), &DemangledName))
      continue;
    DEBUG(dbgs() << "Function demangled name: " << DemangledName << '\n');
    std::string BuiltinVarName;
    SPIRVBuiltinVariableKind BVKind;
    if (!SPIRSPIRVBuiltinVariableMap::find(DemangledName, &BVKind))
      continue;
    BuiltinVarName = std::string(kSPIRVName::Prefix) +
        SPIRVBuiltInNameMap::map(BVKind);
    DEBUG(dbgs() << "builtin variable name: " << BuiltinVarName << '\n');
    bool IsVec = I->getFunctionType()->getNumParams() > 0;
    Type *GVType = IsVec ? VectorType::get(I->getReturnType(),3) :
        I->getReturnType();
    auto BV = new GlobalVariable(*M, GVType,
        true,
        GlobalValue::ExternalLinkage,
        nullptr, BuiltinVarName,
        0,
        GlobalVariable::NotThreadLocal,
        SPIRAS_Constant);
    std::vector<Instruction *> InstList;
    for (auto UI = I->user_begin(), UE = I->user_end(); UI != UE; ++UI) {
      auto CI = dyn_cast<CallInst>(*UI);
      assert(CI && "invalid instruction");
      Value * NewValue = new LoadInst(BV, "", CI);
      DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n');
      if (IsVec) {
        NewValue = ExtractElementInst::Create(NewValue,
          CI->getArgOperand(0),
          "", CI);
        DEBUG(dbgs() << *NewValue << '\n');
      }
      NewValue->takeName(CI);
      CI->replaceAllUsesWith(NewValue);
      InstList.push_back(CI);
    }
    for (auto &Inst:InstList) {
      Inst->dropAllReferences();
      Inst->removeFromParent();
    }
    WorkList.push_back(static_cast<Function*>(I));
  }
  for (auto &I:WorkList) {
    I->dropAllReferences();
    I->removeFromParent();
  }
}

void
OCL20ToSPIRV::visitCallReadWriteImage(CallInst* CI,
    StringRef MangledName, const std::string& DemangledName) {
  OCLBuiltinTransInfo Info;
  if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0)
    Info.UniqName = kOCLBuiltinName::ReadImage;

  if (DemangledName.find(kOCLBuiltinName::WriteImage) == 0)
  {
    Info.UniqName = kOCLBuiltinName::WriteImage;
    Info.PostProc = [&](std::vector<Value*> &Args) {
        if (Args.size() == 4) // write with lod
        {
            auto Lod = Args[2];
            Args.erase(Args.begin() + 2);
            Args.push_back(getInt32(M, ImageOperandsMask::ImageOperandsLodMask));
            Args.push_back(Lod);
        }
    };
  }

  transBuiltin(CI, Info);
}

void
OCL20ToSPIRV::visitCallToAddr(CallInst* CI, StringRef MangledName,
    const std::string &DemangledName) {
  auto AddrSpace = static_cast<SPIRAddressSpace>(
      CI->getType()->getPointerAddressSpace());
  OCLBuiltinTransInfo Info;
  Info.UniqName = DemangledName;
  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 = [=](std::vector<Value *> &Ops){
    auto P = Ops.back();
    Ops.pop_back();
    Ops.push_back(castToInt8Ptr(P, CI));
    Ops.push_back(StorageClass);
  };
  transBuiltin(CI, Info);
}

void OCL20ToSPIRV::visitCallRelational(CallInst *CI,
                                       const std::string &DemangledName) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  Op OC = OpNop;
  OCLSPIRVBuiltinMap::find(DemangledName, &OC);
  std::string SPIRVName = getSPIRVFuncName(OC);
  mutateCallInstSPIRV(
      M, CI,
      [=](CallInst *, std::vector<Value *> &Args, Type *&Ret) {
        Ret = Type::getInt1Ty(*Ctx);
        if (CI->getOperand(0)->getType()->isVectorTy())
          Ret = VectorType::get(
              Type::getInt1Ty(*Ctx),
              CI->getOperand(0)->getType()->getVectorNumElements());
        return SPIRVName;
      },
      [=](CallInst *NewCI) -> Instruction * {
        Value *False = nullptr, *True = nullptr;
        if (NewCI->getType()->isVectorTy()) {
          Type *IntTy = Type::getInt32Ty(*Ctx);
          if (cast<VectorType>(NewCI->getOperand(0)->getType())
                  ->getElementType()
                  ->isDoubleTy())
            IntTy = Type::getInt64Ty(*Ctx);
          if (cast<VectorType>(NewCI->getOperand(0)->getType())
                  ->getElementType()
                  ->isHalfTy())
            IntTy = Type::getInt16Ty(*Ctx);
          Type *VTy = VectorType::get(IntTy,
                                      NewCI->getType()->getVectorNumElements());
          False = Constant::getNullValue(VTy);
          True = Constant::getAllOnesValue(VTy);
        } else {
          False = getInt32(M, 0);
          True = getInt32(M, 1);
        }
        return SelectInst::Create(NewCI, True, False, "", NewCI->getNextNode());
      },
      &Attrs);
}

void
OCL20ToSPIRV::visitCallVecLoadStore(CallInst* CI,
    StringRef MangledName, const std::string& 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;
  Info.UniqName = DemangledName;
  if (DemangledName.find(kOCLBuiltinName::VLoadPrefix) == 0)
    Info.Postfix = std::string(kSPIRVPostfix::ExtDivider) +
      getPostfixForReturnType(CI);
  Info.PostProc = [=](std::vector<Value *> &Ops){
    Ops.insert(Ops.end(), Consts.begin(), Consts.end());
  };
  transBuiltin(CI, Info);
}

void OCL20ToSPIRV::visitCallGetFence(CallInst *CI, StringRef MangledName,
                                     const std::string &DemangledName) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  Op OC = OpNop;
  OCLSPIRVBuiltinMap::find(DemangledName, &OC);
  std::string SPIRVName = getSPIRVFuncName(OC);
  mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args,
                                 Type *&Ret) { return SPIRVName; },
            [=](CallInst *NewCI) -> Instruction * {
              return BinaryOperator::CreateLShr(NewCI, getInt32(M, 8), "", CI);
            },
            &Attrs);
}

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

void OCL20ToSPIRV::visitCallScalToVec(CallInst *CI, StringRef MangledName,
                                      const std::string &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->getNumArgOperands(); 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);
  }

  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  mutateCallInstSPIRV(
      M, CI,
      [=](CallInst *, std::vector<Value *> &Args) {
        Args.resize(VecPos.size() + ScalarPos.size());
        for (auto I : VecPos) {
          Args[I] = CI->getOperand(I);
        }
        auto VecArgWidth =
            CI->getOperand(VecPos[0])->getType()->getVectorNumElements();
        for (auto I : ScalarPos) {
          Instruction *Inst = InsertElementInst::Create(
              UndefValue::get(CI->getOperand(VecPos[0])->getType()),
              CI->getOperand(I), getInt32(M, 0), "", CI);
          Value *NewVec = new ShuffleVectorInst(
              Inst, UndefValue::get(CI->getOperand(VecPos[0])->getType()),
              ConstantVector::getSplat(VecArgWidth, getInt32(M, 0)), "", CI);

          Args[I] = NewVec;
        }
        return getSPIRVExtFuncName(SPIRVEIS_OpenCL,
                                   getExtOp(MangledName, DemangledName));
      },
      &Attrs);
}

void OCL20ToSPIRV::visitCallGetImageChannel(CallInst *CI, StringRef MangledName,
                                            const std::string &DemangledName,
                                            unsigned int Offset) {
  AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
  Op OC = OpNop;
  OCLSPIRVBuiltinMap::find(DemangledName, &OC);
  std::string SPIRVName = getSPIRVFuncName(OC);
  mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args,
                                 Type *&Ret) { return SPIRVName; },
                      [=](CallInst *NewCI) -> Instruction * {
                        return BinaryOperator::CreateAdd(
                            NewCI, getInt32(M, Offset), "", CI);
                      },
                      &Attrs);
}
}

INITIALIZE_PASS_BEGIN(OCL20ToSPIRV, "cl20tospv", "Transform OCL 2.0 to SPIR-V",
    false, false)
INITIALIZE_PASS_DEPENDENCY(OCLTypeToSPIRV)
INITIALIZE_PASS_END(OCL20ToSPIRV, "cl20tospv", "Transform OCL 2.0 to SPIR-V",
    false, false)

ModulePass *llvm::createOCL20ToSPIRV() {
  return new OCL20ToSPIRV();
}