lib/ReplaceOpenCLBuiltinPass.cpp

// Copyright 2017 The Clspv Authors. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//     http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.

#include <math.h>
#include <string>
#include <tuple>

#include "llvm/ADT/StringSwitch.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Module.h"
#include "llvm/IR/Operator.h"
#include "llvm/IR/ValueSymbolTable.h"
#include "llvm/Pass.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/Transforms/Utils/BasicBlockUtils.h"
#include "llvm/Transforms/Utils/Cloning.h"

#include "spirv/unified1/spirv.hpp"

#include "clspv/AddressSpace.h"
#include "clspv/Option.h"

#include "Builtins.h"
#include "Constants.h"
#include "ReplaceOpenCLBuiltinPass.h"
#include "SPIRVOp.h"
#include "Types.h"

using namespace clspv;
using namespace llvm;

namespace {

uint32_t clz(uint32_t v) {
  uint32_t r;
  uint32_t shift;

  r = (v > 0xFFFF) << 4;
  v >>= r;
  shift = (v > 0xFF) << 3;
  v >>= shift;
  r |= shift;
  shift = (v > 0xF) << 2;
  v >>= shift;
  r |= shift;
  shift = (v > 0x3) << 1;
  v >>= shift;
  r |= shift;
  r |= (v >> 1);

  return r;
}

Type *getIntOrIntVectorTyForCast(LLVMContext &C, Type *Ty) {
  Type *IntTy = Type::getIntNTy(C, Ty->getScalarSizeInBits());
  if (auto vec_ty = dyn_cast<VectorType>(Ty)) {
    IntTy = FixedVectorType::get(IntTy,
                                 vec_ty->getElementCount().getKnownMinValue());
  }
  return IntTy;
}

Value *MemoryOrderSemantics(Value *order, bool is_global,
                            Instruction *InsertBefore,
                            spv::MemorySemanticsMask base_semantics,
                            bool include_storage = true) {
  enum AtomicMemoryOrder : uint32_t {
    kMemoryOrderRelaxed = 0,
    kMemoryOrderAcquire = 2,
    kMemoryOrderRelease = 3,
    kMemoryOrderAcqRel = 4,
    kMemoryOrderSeqCst = 5
  };

  IRBuilder<> builder(InsertBefore);

  // Constants for OpenCL C 2.0 memory_order.
  const auto relaxed = builder.getInt32(AtomicMemoryOrder::kMemoryOrderRelaxed);
  const auto acquire = builder.getInt32(AtomicMemoryOrder::kMemoryOrderAcquire);
  const auto release = builder.getInt32(AtomicMemoryOrder::kMemoryOrderRelease);
  const auto acq_rel = builder.getInt32(AtomicMemoryOrder::kMemoryOrderAcqRel);

  // Constants for SPIR-V ordering memory semantics.
  const auto RelaxedSemantics = builder.getInt32(spv::MemorySemanticsMaskNone);
  const auto AcquireSemantics =
      builder.getInt32(spv::MemorySemanticsAcquireMask);
  const auto ReleaseSemantics =
      builder.getInt32(spv::MemorySemanticsReleaseMask);
  const auto AcqRelSemantics =
      builder.getInt32(spv::MemorySemanticsAcquireReleaseMask);

  // Constants for SPIR-V storage class semantics.
  const auto UniformSemantics =
      builder.getInt32(spv::MemorySemanticsUniformMemoryMask);
  const auto WorkgroupSemantics =
      builder.getInt32(spv::MemorySemanticsWorkgroupMemoryMask);

  // Instead of sequentially consistent, use acquire, release or acquire
  // release semantics.
  Value *base_order = nullptr;
  switch (base_semantics) {
  case spv::MemorySemanticsAcquireMask:
    base_order = AcquireSemantics;
    break;
  case spv::MemorySemanticsReleaseMask:
    base_order = ReleaseSemantics;
    break;
  default:
    base_order = AcqRelSemantics;
    break;
  }

  Value *storage = is_global ? UniformSemantics : WorkgroupSemantics;
  if (order == nullptr) {
    if (include_storage)
      return builder.CreateOr({storage, base_order});
    else
      return base_order;
  }

  auto is_relaxed = builder.CreateICmpEQ(order, relaxed);
  auto is_acquire = builder.CreateICmpEQ(order, acquire);
  auto is_release = builder.CreateICmpEQ(order, release);
  auto is_acq_rel = builder.CreateICmpEQ(order, acq_rel);
  auto semantics =
      builder.CreateSelect(is_relaxed, RelaxedSemantics, base_order);
  semantics = builder.CreateSelect(is_acquire, AcquireSemantics, semantics);
  semantics = builder.CreateSelect(is_release, ReleaseSemantics, semantics);
  semantics = builder.CreateSelect(is_acq_rel, AcqRelSemantics, semantics);
  if (include_storage)
    return builder.CreateOr({storage, semantics});
  else
    return semantics;
}

Value *MemoryScope(Value *scope, bool is_global, Instruction *InsertBefore) {
  enum AtomicMemoryScope : uint32_t {
    kMemoryScopeWorkItem = 0,
    kMemoryScopeWorkGroup = 1,
    kMemoryScopeDevice = 2,
    kMemoryScopeAllSVMDevices = 3, // not supported
    kMemoryScopeSubGroup = 4
  };

  IRBuilder<> builder(InsertBefore);

  // Constants for OpenCL C 2.0 memory_scope.
  const auto work_item =
      builder.getInt32(AtomicMemoryScope::kMemoryScopeWorkItem);
  const auto work_group =
      builder.getInt32(AtomicMemoryScope::kMemoryScopeWorkGroup);
  const auto sub_group =
      builder.getInt32(AtomicMemoryScope::kMemoryScopeSubGroup);
  const auto device = builder.getInt32(AtomicMemoryScope::kMemoryScopeDevice);

  // Constants for SPIR-V memory scopes.
  const auto InvocationScope = builder.getInt32(spv::ScopeInvocation);
  const auto WorkgroupScope = builder.getInt32(spv::ScopeWorkgroup);
  const auto DeviceScope = builder.getInt32(spv::ScopeDevice);
  const auto SubgroupScope = builder.getInt32(spv::ScopeSubgroup);

  auto base_scope = is_global ? DeviceScope : WorkgroupScope;
  if (scope == nullptr)
    return base_scope;

  auto is_work_item = builder.CreateICmpEQ(scope, work_item);
  auto is_work_group = builder.CreateICmpEQ(scope, work_group);
  auto is_sub_group = builder.CreateICmpEQ(scope, sub_group);
  auto is_device = builder.CreateICmpEQ(scope, device);

  scope = builder.CreateSelect(is_work_item, InvocationScope, base_scope);
  scope = builder.CreateSelect(is_work_group, WorkgroupScope, scope);
  scope = builder.CreateSelect(is_sub_group, SubgroupScope, scope);
  scope = builder.CreateSelect(is_device, DeviceScope, scope);

  return scope;
}

bool replaceCallsWithValue(Function &F,
                           std::function<Value *(CallInst *)> Replacer) {

  bool Changed = false;

  SmallVector<Instruction *, 4> ToRemoves;

  // Walk the users of the function.
  for (auto &U : F.uses()) {
    if (auto CI = dyn_cast<CallInst>(U.getUser())) {

      auto NewValue = Replacer(CI);

      if (NewValue != nullptr) {
        CI->replaceAllUsesWith(NewValue);

        // Lastly, remember to remove the user.
        ToRemoves.push_back(CI);
      }
    }
  }

  Changed = !ToRemoves.empty();

  // And cleanup the calls we don't use anymore.
  for (auto V : ToRemoves) {
    V->eraseFromParent();
  }

  return Changed;
}

} // namespace

PreservedAnalyses ReplaceOpenCLBuiltinPass::run(Module &M,
                                                ModuleAnalysisManager &MPM) {
  PreservedAnalyses PA;
  std::list<Function *> func_list;
  for (auto &F : M.getFunctionList()) {
    // process only function declarations
    if (F.isDeclaration() && runOnFunction(F)) {
      func_list.push_front(&F);
    }
  }
  if (func_list.size() != 0) {
    // recursively convert functions, but first remove dead
    for (auto *F : func_list) {
      if (F->use_empty()) {
        F->eraseFromParent();
      }
    }
    PA = run(M, MPM);
    return PA;
  }
  return PA;
}

bool ReplaceOpenCLBuiltinPass::runOnFunction(Function &F) {
  auto &FI = Builtins::Lookup(&F);
  switch (FI.getType()) {
  case Builtins::kAbs:
    if (!FI.getParameter(0).is_signed) {
      return replaceAbs(F);
    }
    break;
  case Builtins::kAbsDiff:
    return replaceAbsDiff(F, FI.getParameter(0).is_signed);

  case Builtins::kAddSat:
    return replaceAddSubSat(F, FI.getParameter(0).is_signed, true);

  case Builtins::kClz:
    return replaceCountZeroes(F, true);

  case Builtins::kCtz:
    return replaceCountZeroes(F, false);

  case Builtins::kHadd:
    return replaceHadd(F, FI.getParameter(0).is_signed, Instruction::And);
  case Builtins::kRhadd:
    return replaceHadd(F, FI.getParameter(0).is_signed, Instruction::Or);

  case Builtins::kCopysign:
    return replaceCopysign(F);

  case Builtins::kHalfRecip:
  case Builtins::kNativeRecip:
    return replaceRecip(F);

  case Builtins::kHalfDivide:
  case Builtins::kNativeDivide:
    return replaceDivide(F);

  case Builtins::kDot:
    return replaceDot(F);

  case Builtins::kExp10:
  case Builtins::kHalfExp10:
  case Builtins::kNativeExp10:
    return replaceExp10(F, FI.getName());

  case Builtins::kExpm1:
    return replaceExpm1(F);

  case Builtins::kLog10:
  case Builtins::kHalfLog10:
  case Builtins::kNativeLog10:
    return replaceLog10(F, FI.getName());

  case Builtins::kLog1p:
    return replaceLog1p(F);

  case Builtins::kFdim:
    return replaceFDim(F);

  case Builtins::kFmod:
    return replaceFmod(F);

  case Builtins::kPown:
    return replacePown(F);

  case Builtins::kRound:
    return replaceRound(F);

  case Builtins::kCospi:
  case Builtins::kSinpi:
  case Builtins::kTanpi:
    return replaceTrigPi(F, FI.getType());

  case Builtins::kSincos:
    return replaceSincos(F);

  case Builtins::kBarrier:
  case Builtins::kWorkGroupBarrier:
    return replaceBarrier(F);

  case Builtins::kSubGroupBarrier:
    return replaceBarrier(F, true);

  case Builtins::kAtomicWorkItemFence:
    return replaceMemFence(F, spv::MemorySemanticsMaskNone);
  case Builtins::kMemFence:
    return replaceMemFence(F, spv::MemorySemanticsAcquireReleaseMask);
  case Builtins::kReadMemFence:
    return replaceMemFence(F, spv::MemorySemanticsAcquireMask);
  case Builtins::kWriteMemFence:
    return replaceMemFence(F, spv::MemorySemanticsReleaseMask);

    // Relational
  case Builtins::kIsequal:
    return replaceRelational(F, CmpInst::FCMP_OEQ);
  case Builtins::kIsgreater:
    return replaceRelational(F, CmpInst::FCMP_OGT);
  case Builtins::kIsgreaterequal:
    return replaceRelational(F, CmpInst::FCMP_OGE);
  case Builtins::kIsless:
    return replaceRelational(F, CmpInst::FCMP_OLT);
  case Builtins::kIslessequal:
    return replaceRelational(F, CmpInst::FCMP_OLE);
  case Builtins::kIsnotequal:
    return replaceRelational(F, CmpInst::FCMP_UNE);
  case Builtins::kIslessgreater:
    return replaceRelational(F, CmpInst::FCMP_ONE);

  case Builtins::kIsordered:
    return replaceOrdered(F, true);

  case Builtins::kIsunordered:
    return replaceOrdered(F, false);

  case Builtins::kIsinf: {
    bool is_vec = FI.getParameter(0).vector_size != 0;
    return replaceIsInfAndIsNan(F, spv::OpIsInf, is_vec ? -1 : 1);
  }
  case Builtins::kIsnan: {
    bool is_vec = FI.getParameter(0).vector_size != 0;
    return replaceIsInfAndIsNan(F, spv::OpIsNan, is_vec ? -1 : 1);
  }

  case Builtins::kIsfinite:
    return replaceIsFinite(F);

  case Builtins::kAll: {
    bool is_vec = FI.getParameter(0).vector_size != 0;
    return replaceAllAndAny(F, !is_vec ? spv::OpNop : spv::OpAll);
  }
  case Builtins::kAny: {
    bool is_vec = FI.getParameter(0).vector_size != 0;
    return replaceAllAndAny(F, !is_vec ? spv::OpNop : spv::OpAny);
  }

  case Builtins::kIsnormal:
    return replaceIsNormal(F);

  case Builtins::kUpsample:
    return replaceUpsample(F);

  case Builtins::kRotate:
    return replaceRotate(F);

  case Builtins::kConvert:
    return replaceConvert(F, FI.getParameter(0).is_signed,
                          FI.getReturnType().is_signed);

  // OpenCL 2.0 explicit atomics have different default scopes and semantics
  // than legacy atomic functions.
  case Builtins::kAtomicLoad:
  case Builtins::kAtomicLoadExplicit:
    return replaceAtomicLoad(F);
  case Builtins::kAtomicStore:
  case Builtins::kAtomicStoreExplicit:
    return replaceExplicitAtomics(F, spv::OpAtomicStore,
                                  spv::MemorySemanticsReleaseMask);
  case Builtins::kAtomicExchange:
  case Builtins::kAtomicExchangeExplicit:
    return replaceExplicitAtomics(F, spv::OpAtomicExchange);
  case Builtins::kAtomicFetchAdd:
  case Builtins::kAtomicFetchAddExplicit:
    return replaceExplicitAtomics(F, spv::OpAtomicIAdd);
  case Builtins::kAtomicFetchSub:
  case Builtins::kAtomicFetchSubExplicit:
    return replaceExplicitAtomics(F, spv::OpAtomicISub);
  case Builtins::kAtomicFetchOr:
  case Builtins::kAtomicFetchOrExplicit:
    return replaceExplicitAtomics(F, spv::OpAtomicOr);
  case Builtins::kAtomicFetchXor:
  case Builtins::kAtomicFetchXorExplicit:
    return replaceExplicitAtomics(F, spv::OpAtomicXor);
  case Builtins::kAtomicFetchAnd:
  case Builtins::kAtomicFetchAndExplicit:
    return replaceExplicitAtomics(F, spv::OpAtomicAnd);
  case Builtins::kAtomicFetchMin:
  case Builtins::kAtomicFetchMinExplicit:
    return replaceExplicitAtomics(F, FI.getParameter(1).is_signed
                                         ? spv::OpAtomicSMin
                                         : spv::OpAtomicUMin);
  case Builtins::kAtomicFetchMax:
  case Builtins::kAtomicFetchMaxExplicit:
    return replaceExplicitAtomics(F, FI.getParameter(1).is_signed
                                         ? spv::OpAtomicSMax
                                         : spv::OpAtomicUMax);
  // Weak compare exchange is generated as strong compare exchange.
  case Builtins::kAtomicCompareExchangeWeak:
  case Builtins::kAtomicCompareExchangeWeakExplicit:
  case Builtins::kAtomicCompareExchangeStrong:
  case Builtins::kAtomicCompareExchangeStrongExplicit:
    return replaceAtomicCompareExchange(F);

  // Legacy atomic functions.
  case Builtins::kAtomicInc:
    return replaceAtomics(F, spv::OpAtomicIIncrement);
  case Builtins::kAtomicDec:
    return replaceAtomics(F, spv::OpAtomicIDecrement);
  case Builtins::kAtomicCmpxchg:
    return replaceAtomics(F, spv::OpAtomicCompareExchange);
  case Builtins::kAtomicAdd:
    return replaceAtomics(F, llvm::AtomicRMWInst::Add);
  case Builtins::kAtomicSub:
    return replaceAtomics(F, llvm::AtomicRMWInst::Sub);
  case Builtins::kAtomicXchg:
    return replaceAtomics(F, llvm::AtomicRMWInst::Xchg);
  case Builtins::kAtomicMin:
    return replaceAtomics(F, FI.getParameter(0).is_signed
                                 ? llvm::AtomicRMWInst::Min
                                 : llvm::AtomicRMWInst::UMin);
  case Builtins::kAtomicMax:
    return replaceAtomics(F, FI.getParameter(0).is_signed
                                 ? llvm::AtomicRMWInst::Max
                                 : llvm::AtomicRMWInst::UMax);
  case Builtins::kAtomicAnd:
    return replaceAtomics(F, llvm::AtomicRMWInst::And);
  case Builtins::kAtomicOr:
    return replaceAtomics(F, llvm::AtomicRMWInst::Or);
  case Builtins::kAtomicXor:
    return replaceAtomics(F, llvm::AtomicRMWInst::Xor);

  case Builtins::kCross:
    if (FI.getParameter(0).vector_size == 4) {
      return replaceCross(F);
    }
    break;

  case Builtins::kFract:
    if (FI.getParameterCount()) {
      return replaceFract(F, FI.getParameter(0).vector_size);
    }
    break;

  case Builtins::kMadHi:
    return replaceMulHi(F, FI.getParameter(0).is_signed, true);
  case Builtins::kMulHi:
    return replaceMulHi(F, FI.getParameter(0).is_signed, false);

  case Builtins::kMadSat:
    return replaceMadSat(F, FI.getParameter(0).is_signed);

  case Builtins::kMad:
  case Builtins::kMad24:
    return replaceMul(F, FI.getParameter(0).type_id == llvm::Type::FloatTyID,
                      true);
  case Builtins::kMul24:
    return replaceMul(F, FI.getParameter(0).type_id == llvm::Type::FloatTyID,
                      false);

  case Builtins::kSelect:
    return replaceSelect(F);

  case Builtins::kBitselect:
    return replaceBitSelect(F);

  case Builtins::kVload:
    return replaceVload(F);

  case Builtins::kVloadaHalf:
    return replaceVloadHalf(F, FI.getName(), FI.getParameter(0).vector_size,
                            true);
  case Builtins::kVloadHalf:
    return replaceVloadHalf(F, FI.getName(), FI.getParameter(0).vector_size,
                            false);

  case Builtins::kVstore:
    return replaceVstore(F);

  case Builtins::kVstoreaHalf:
    return replaceVstoreHalf(F, FI.getParameter(0).vector_size, true);
  case Builtins::kVstoreHalf:
    return replaceVstoreHalf(F, FI.getParameter(0).vector_size, false);

  case Builtins::kSmoothstep: {
    int vec_size = FI.getLastParameter().vector_size;
    if (FI.getParameter(0).vector_size == 0 && vec_size != 0) {
      return replaceStep(F, true);
    }
    break;
  }
  case Builtins::kStep: {
    int vec_size = FI.getLastParameter().vector_size;
    if (FI.getParameter(0).vector_size == 0 && vec_size != 0) {
      return replaceStep(F, false);
    }
    break;
  }

  case Builtins::kSignbit:
    return replaceSignbit(F, FI.getParameter(0).vector_size != 0);

  case Builtins::kSubSat:
    return replaceAddSubSat(F, FI.getParameter(0).is_signed, false);

  case Builtins::kReadImageh:
    return replaceHalfReadImage(F);
  case Builtins::kReadImagef:
  case Builtins::kReadImagei:
  case Builtins::kReadImageui: {
    if (FI.getParameter(1).isSampler() &&
        FI.getParameter(2).type_id == llvm::Type::IntegerTyID) {
      return replaceSampledReadImageWithIntCoords(F);
    }
    break;
  }

  case Builtins::kWriteImageh:
    return replaceHalfWriteImage(F);

  case Builtins::kPrefetch:
    return replacePrefetch(F);

  // Asynchronous copies
  case Builtins::kAsyncWorkGroupCopy:
    return replaceAsyncWorkGroupCopy(
        F, FI.getParameter(0).DataType(F.getParent()->getContext()));
  case Builtins::kAsyncWorkGroupStridedCopy:
    return replaceAsyncWorkGroupStridedCopy(
        F, FI.getParameter(0).DataType(F.getParent()->getContext()));
  case Builtins::kWaitGroupEvents:
    return replaceWaitGroupEvents(F);

  default:
    break;
  }

  return false;
}

Type *ReplaceOpenCLBuiltinPass::GetPairStruct(Type *type) {
  auto iter = PairStructMap.find(type);
  if (iter != PairStructMap.end())
    return iter->second;

  auto new_struct = StructType::get(type->getContext(), {type, type});
  PairStructMap[type] = new_struct;
  return new_struct;
}

Value *ReplaceOpenCLBuiltinPass::InsertOpMulExtended(Instruction *InsertPoint,
                                                     Value *a, Value *b,
                                                     bool IsSigned,
                                                     bool Int64) {

  Type *Ty = a->getType();
  Type *RetTy = GetPairStruct(a->getType());
  assert(Ty == b->getType());

  if (!Option::HackMulExtended()) {
    spv::Op opcode = IsSigned ? spv::OpSMulExtended : spv::OpUMulExtended;

    return clspv::InsertSPIRVOp(InsertPoint, opcode, {Attribute::ReadNone},
                                RetTy, {a, b});
  }

  unsigned int ScalarSizeInBits = Ty->getScalarSizeInBits();
  bool IsVector = Ty->isVectorTy();

  IRBuilder<> Builder(InsertPoint);

  if (ScalarSizeInBits < 32 || (ScalarSizeInBits == 32 && Int64)) {
    /*
     * {mul_lo, mul_hi} = OpMulExtended(a, b, IsSigned) {
     *    S = SizeInBits(a)
     *    a_ext = ext2S(a, IsSigned)
     *    b_ext = ext2S(b, IsSigned)
     *    mul = a_ext * b_ext
     *    mul_lo = truncS(mul)
     *    mul_hi = truncS(mul >> S)
     *    return {mul_lo, mul_hi}
     * }
     */
    Type *TyTimes2 =
        Ty->getIntNTy(InsertPoint->getContext(), ScalarSizeInBits * 2);
    if (IsVector) {
      TyTimes2 = VectorType::get(TyTimes2, dyn_cast<VectorType>(Ty));
    }
    Value *aExtended, *bExtended;
    if (IsSigned) {
      aExtended = Builder.CreateSExt(a, TyTimes2);
      bExtended = Builder.CreateSExt(b, TyTimes2);
    } else {
      aExtended = Builder.CreateZExt(a, TyTimes2);
      bExtended = Builder.CreateZExt(b, TyTimes2);
    }
    auto mul = Builder.CreateMul(aExtended, bExtended);
    auto mul_lo = Builder.CreateTrunc(mul, Ty);
    auto mul_hi =
        Builder.CreateTrunc(Builder.CreateLShr(mul, ScalarSizeInBits), Ty);

    return Builder.CreateInsertValue(
        Builder.CreateInsertValue(UndefValue::get(RetTy), mul_lo, {0}), mul_hi,
        {1});
  } else if (ScalarSizeInBits == 64 || (ScalarSizeInBits == 32 && !Int64)) {
    /*
     * {mul_lo, mul_hi} = OpMulExtended(a, b, IsSigned) {
     *   S = SizeInBits(a)
     *   hS = S / 2
     *   if (IsSigned) {
     *     res_neg = (a > 0) ^ (b > 0) = (a ^ b) < 0
     *     a = abs(a)
     *     b = abs(b)
     *   }
     *   a0 = trunchS(a)
     *   a1 = trunchS(a >> hS)
     *   b0 = trunchS(b)
     *   b1 = trunchS(b >> hS)
     *   {a0b0_0, a0b0_1} = zextS(OpUMulExtended(a0, b0))
     *   {a1b0_0, a1b0_1} = zextS(OpUMulExtended(a1, b0))
     *   {a0b1_0, a0b1_1} = zextS(OpUMulExtended(a0, b1))
     *   {a1b1_0, a1b1_1} = zextS(OpUMulExtended(a1, b1))
     *
     *   mul_lo_hi = a0b0_1 + a1b0_0 + a0b1_0
     *   carry_mul_lo_hi = mul_lo_hi >> hS
     *   mul_hi_lo = a1b1_0 + a1b0_1 + a0b1_1 + carry_mul_lo_hi
     *   mul_lo = a0b0_0 + mul_lo_hi << hS
     *   mul_hi = mul_hi_lo + a1b1_1 << hS
     *
     *   if (IsSigned) {
     *     mul_lo_xor = mul_lo ^ -1
     *     {mul_lo_inv, carry} = OpIAddCarry(mul_lo_xor, 1)
     *     mul_hi_inv = mul_hi ^ -1 + carry
     *     mul_lo = res_neg ? mul_lo_inv : mul_lo
     *     mul_hi = res_neg ? mul_hi_inv : mul_hi
     *   }
     *   return {mul_lo, mul_hi}
     * }
     */
    Type *TyDiv2 =
        Ty->getIntNTy(InsertPoint->getContext(), ScalarSizeInBits / 2);
    if (IsVector) {
      TyDiv2 = VectorType::get(TyDiv2, dyn_cast<VectorType>(Ty));
    }

    Value *res_neg;
    if (IsSigned) {
      // We want to work with unsigned value.
      // Convert everything to unsigned and remember the signed of the end
      // result.
      auto a_b_xor = Builder.CreateXor(a, b);
      res_neg = Builder.CreateICmpSLT(a_b_xor, ConstantInt::get(Ty, 0, true));

      auto F = InsertPoint->getFunction();
      auto abs = Intrinsic::getDeclaration(F->getParent(), Intrinsic::abs, Ty);
      a = Builder.CreateCall(abs, {a, Builder.getInt1(false)});
      b = Builder.CreateCall(abs, {b, Builder.getInt1(false)});
    }

    auto a0 = Builder.CreateTrunc(a, TyDiv2);
    auto a1 = Builder.CreateTrunc(Builder.CreateLShr(a, ScalarSizeInBits / 2),
                                  TyDiv2);
    auto b0 = Builder.CreateTrunc(b, TyDiv2);
    auto b1 = Builder.CreateTrunc(Builder.CreateLShr(b, ScalarSizeInBits / 2),
                                  TyDiv2);

    auto a0b0 = InsertOpMulExtended(InsertPoint, a0, b0, false, true);
    auto a1b0 = InsertOpMulExtended(InsertPoint, a1, b0, false, true);
    auto a0b1 = InsertOpMulExtended(InsertPoint, a0, b1, false, true);
    auto a1b1 = InsertOpMulExtended(InsertPoint, a1, b1, false, true);
    auto a0b0_0 = Builder.CreateZExt(Builder.CreateExtractValue(a0b0, {0}), Ty);
    auto a0b0_1 = Builder.CreateZExt(Builder.CreateExtractValue(a0b0, {1}), Ty);
    auto a1b0_0 = Builder.CreateZExt(Builder.CreateExtractValue(a1b0, {0}), Ty);
    auto a1b0_1 = Builder.CreateZExt(Builder.CreateExtractValue(a1b0, {1}), Ty);
    auto a0b1_0 = Builder.CreateZExt(Builder.CreateExtractValue(a0b1, {0}), Ty);
    auto a0b1_1 = Builder.CreateZExt(Builder.CreateExtractValue(a0b1, {1}), Ty);
    auto a1b1_0 = Builder.CreateZExt(Builder.CreateExtractValue(a1b1, {0}), Ty);
    auto a1b1_1 = Builder.CreateZExt(Builder.CreateExtractValue(a1b1, {1}), Ty);

    auto mul_lo_hi =
        Builder.CreateAdd(Builder.CreateAdd(a0b0_1, a1b0_0), a0b1_0);
    auto carry_mul_lo_hi = Builder.CreateLShr(mul_lo_hi, ScalarSizeInBits / 2);
    auto mul_hi_lo = Builder.CreateAdd(
        Builder.CreateAdd(Builder.CreateAdd(a1b1_0, a1b0_1), a0b1_1),
        carry_mul_lo_hi);
    auto mul_lo = Builder.CreateAdd(
        a0b0_0, Builder.CreateShl(mul_lo_hi, ScalarSizeInBits / 2));
    auto mul_hi = Builder.CreateAdd(
        mul_hi_lo, Builder.CreateShl(a1b1_1, ScalarSizeInBits / 2));

    if (IsSigned) {
      // Apply the sign that we got from the previous if statement setting
      // res_neg.
      auto mul_lo_xor =
          Builder.CreateXor(mul_lo, Constant::getAllOnesValue(Ty));
      auto mul_lo_xor_add =
          InsertSPIRVOp(InsertPoint, spv::OpIAddCarry, {Attribute::ReadNone},
                        RetTy, {mul_lo_xor, ConstantInt::get(Ty, 1)});
      auto mul_lo_inv = Builder.CreateExtractValue(mul_lo_xor_add, {0});
      auto carry = Builder.CreateExtractValue(mul_lo_xor_add, {1});
      auto mul_hi_inv = Builder.CreateAdd(
          carry, Builder.CreateXor(mul_hi, Constant::getAllOnesValue(Ty)));
      mul_lo = Builder.CreateSelect(res_neg, mul_lo_inv, mul_lo);
      mul_hi = Builder.CreateSelect(res_neg, mul_hi_inv, mul_hi);
    }

    return Builder.CreateInsertValue(
        Builder.CreateInsertValue(UndefValue::get(RetTy), mul_lo, {0}), mul_hi,
        {1});
  } else {
    llvm_unreachable("Unexpected type for InsertOpMulExtended");
  }
}

bool ReplaceOpenCLBuiltinPass::replaceWaitGroupEvents(Function &F) {
  /* Simple implementation for wait_group_events to avoid dealing with the event
   * list:
   *
   * void wait_group_events(int num_events, event_t *event_list) {
   *   barrier(CLK_LOCAL_MEM_FENCE);
   * }
   *
   */

  enum {
    CLK_LOCAL_MEM_FENCE = 0x01,
    CLK_GLOBAL_MEM_FENCE = 0x02,
    CLK_IMAGE_MEM_FENCE = 0x04
  };

  return replaceCallsWithValue(F, [](CallInst *CI) {
    IRBuilder<> Builder(CI);

    const auto ConstantScopeWorkgroup = Builder.getInt32(spv::ScopeWorkgroup);
    const auto MemorySemanticsWorkgroup = BinaryOperator::Create(
        Instruction::Shl, Builder.getInt32(CLK_LOCAL_MEM_FENCE),
        Builder.getInt32(clz(spv::MemorySemanticsWorkgroupMemoryMask) -
                         clz(CLK_LOCAL_MEM_FENCE)),
        "", CI);
    auto MemorySemantics = BinaryOperator::Create(
        Instruction::Or, MemorySemanticsWorkgroup,
        ConstantInt::get(Builder.getInt32Ty(),
                         spv::MemorySemanticsAcquireReleaseMask),
        "", CI);

    return clspv::InsertSPIRVOp(
        CI, spv::OpControlBarrier,
        {Attribute::NoDuplicate, Attribute::Convergent}, Builder.getVoidTy(),
        {ConstantScopeWorkgroup, ConstantScopeWorkgroup, MemorySemantics});
  });
}

GlobalVariable *ReplaceOpenCLBuiltinPass::getOrCreateGlobalVariable(
    Module &M, std::string VariableName,
    AddressSpace::Type VariableAddressSpace) {
  GlobalVariable *GV = M.getGlobalVariable(VariableName);
  if (GV == nullptr) {
    IntegerType *IT = IntegerType::get(M.getContext(), 32);
    VectorType *VT = FixedVectorType::get(IT, 3);

    GV = new GlobalVariable(M, VT, false, GlobalValue::ExternalLinkage, nullptr,
                            VariableName, nullptr,
                            GlobalValue::ThreadLocalMode::NotThreadLocal,
                            VariableAddressSpace);
    GV->setInitializer(Constant::getNullValue(VT));
  }
  return GV;
}

Value *ReplaceOpenCLBuiltinPass::replaceAsyncWorkGroupCopies(
    Module &M, CallInst *CI, Value *Dst, Value *Src, Type *GenType,
    Value *NumGentypes, Value *Stride, Value *Event) {
  /*
   * event_t *async_work_group_strided_copy(T *dst, T *src, size_t num_gentypes,
   *                                        size_t stride, event_t event) {
   *   size_t start_id = ((get_local_id(2) * get_local_size(1))
   *                     + get_local_id(1)) * get_local_size(0)
   *                     + get_local_id(0);
   *   size_t incr = get_local_size(0) * get_local_size(1) * get_local_size(2);
   *   for (size_t it = start_id; it < num_gentypes; it += incr) {
   *     dst[it] = src[it * stride];
   *   }
   *   return event;
   * }
   */

  /* BB:
   *    before
   *    async_work_group_strided_copy
   *    after
   *
   * ================================
   *
   * BB:
   *    before
   *    start_id = f(get_local_ids, get_local_sizes)
   *    incr = g(get_local_sizes)
   *    br CmpBB
   *
   * CmpBB:
   *    it = PHI(start_id, it)
   *    cmp = it < NumGentypes
   *    condBr cmp, LoopBB, ExitBB
   *
   * LoopBB:
   *    dstI = dst[it]
   *    srcI = src[it * stride]
   *    OpCopyMemory dstI, srcI
   *    it += incr
   *    br CmpBB
   *
   * ExitBB:
   *    after
   */

  IRBuilder<> Builder(CI);

  auto Cst0 = Builder.getInt32(0);
  auto Cst1 = Builder.getInt32(1);
  auto Cst2 = Builder.getInt32(2);

  auto *IT = IntegerType::get(M.getContext(), 32);
  auto *VT = FixedVectorType::get(IT, 3);

  // get_local_id({0, 1, 2});
  GlobalVariable *GVId =
      getOrCreateGlobalVariable(M, clspv::LocalInvocationIdVariableName(),
                                clspv::LocalInvocationIdAddressSpace());
  Value *GEP0 = Builder.CreateGEP(VT, GVId, {Cst0, Cst0});
  Value *LocalId0 = Builder.CreateLoad(IT, GEP0);
  Value *GEP1 = Builder.CreateGEP(VT, GVId, {Cst0, Cst1});
  Value *LocalId1 = Builder.CreateLoad(IT, GEP1);
  Value *GEP2 = Builder.CreateGEP(VT, GVId, {Cst0, Cst2});
  Value *LocalId2 = Builder.CreateLoad(IT, GEP2);

  // get_local_size({0, 1, 2});
  GlobalVariable *GVSize =
      getOrCreateGlobalVariable(M, clspv::WorkgroupSizeVariableName(),
                                clspv::WorkgroupSizeAddressSpace());
  auto LocalSize = Builder.CreateLoad(VT, GVSize);
  auto LocalSize0 = Builder.CreateExtractElement(LocalSize, Cst0);
  auto LocalSize1 = Builder.CreateExtractElement(LocalSize, Cst1);
  auto LocalSize2 = Builder.CreateExtractElement(LocalSize, Cst2);

  // size_t start_id = ((get_local_id(2) * get_local_size(1))
  //                   + get_local_id(1)) * get_local_size(0)
  //                   + get_local_id(0);
  auto tmp0 = Builder.CreateMul(LocalId2, LocalSize1);
  auto tmp1 = Builder.CreateAdd(tmp0, LocalId1);
  auto tmp2 = Builder.CreateMul(tmp1, LocalSize0);
  auto StartId = Builder.CreateAdd(tmp2, LocalId0);

  // size_t incr = get_local_size(0) * get_local_size(1) * get_local_size(2);
  auto tmp3 = Builder.CreateMul(LocalSize0, LocalSize1);
  auto Incr = Builder.CreateMul(tmp3, LocalSize2);

  // Create BasicBlocks
  auto BB = CI->getParent();
  auto CmpBB = BasicBlock::Create(BB->getContext(), "", BB->getParent());
  auto LoopBB = BasicBlock::Create(BB->getContext(), "", BB->getParent());
  auto ExitBB = SplitBlock(BB, CI);

  // BB
  auto BrCmpBB = BranchInst::Create(CmpBB);
  ReplaceInstWithInst(BB->getTerminator(), BrCmpBB);

  // CmpBB
  Builder.SetInsertPoint(CmpBB);
  auto PHIIterator = Builder.CreatePHI(Builder.getInt32Ty(), 2);
  auto Cmp = Builder.CreateCmp(CmpInst::ICMP_ULT, PHIIterator, NumGentypes);
  Builder.CreateCondBr(Cmp, LoopBB, ExitBB);

  // LoopBB
  Builder.SetInsertPoint(LoopBB);

  // default values for non-strided copies
  Value *SrcIterator = PHIIterator;
  Value *DstIterator = PHIIterator;
  if (Stride != nullptr && (Dst->getType()->getPointerAddressSpace() ==
                            clspv::AddressSpace::Global)) {
    // async_work_group_strided_copy local to global case
    DstIterator = Builder.CreateMul(PHIIterator, Stride);
  } else if (Stride != nullptr && (Dst->getType()->getPointerAddressSpace() ==
                                   clspv::AddressSpace::Local)) {
    // async_work_group_strided_copy global to local case
    SrcIterator = Builder.CreateMul(PHIIterator, Stride);
  }
  auto DstI = Builder.CreateGEP(GenType, Dst, DstIterator);
  auto SrcI = Builder.CreateGEP(GenType, Src, SrcIterator);
  auto NewIterator = Builder.CreateAdd(PHIIterator, Incr);
  auto Br = Builder.CreateBr(CmpBB);
  clspv::InsertSPIRVOp(Br, spv::OpCopyMemory, {}, Builder.getVoidTy(),
                       {DstI, SrcI});

  // Set PHIIterator for CmpBB now that we have NewIterator
  PHIIterator->addIncoming(StartId, BB);
  PHIIterator->addIncoming(NewIterator, LoopBB);

  return Event;
}

bool ReplaceOpenCLBuiltinPass::replaceAsyncWorkGroupCopy(Function &F, Type *ty) {
  return replaceCallsWithValue(F, [&F, ty, this](CallInst *CI) {
    Module &M = *F.getParent();

    auto Dst = CI->getOperand(0);
    auto Src = CI->getOperand(1);
    auto NumGentypes = CI->getOperand(2);
    auto Event = CI->getOperand(3);

    return replaceAsyncWorkGroupCopies(M, CI, Dst, Src, ty, NumGentypes,
                                       nullptr, Event);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAsyncWorkGroupStridedCopy(Function &F, Type *ty) {
  return replaceCallsWithValue(F, [&F, ty, this](CallInst *CI) {
    Module &M = *F.getParent();

    auto Dst = CI->getOperand(0);
    auto Src = CI->getOperand(1);
    auto NumGentypes = CI->getOperand(2);
    auto Stride = CI->getOperand(3);
    auto Event = CI->getOperand(4);

    return replaceAsyncWorkGroupCopies(M, CI, Dst, Src, ty, NumGentypes, Stride,
                                       Event);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAbs(Function &F) {
  return replaceCallsWithValue(F,
                               [](CallInst *CI) { return CI->getOperand(0); });
}

bool ReplaceOpenCLBuiltinPass::replaceAbsDiff(Function &F, bool is_signed) {
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto XValue = CI->getOperand(0);
    auto YValue = CI->getOperand(1);

    IRBuilder<> Builder(CI);
    auto XmY = Builder.CreateSub(XValue, YValue);
    auto YmX = Builder.CreateSub(YValue, XValue);

    Value *Cmp = nullptr;
    if (is_signed) {
      Cmp = Builder.CreateICmpSGT(YValue, XValue);
    } else {
      Cmp = Builder.CreateICmpUGT(YValue, XValue);
    }

    return Builder.CreateSelect(Cmp, YmX, XmY);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceCopysign(Function &F) {
  return replaceCallsWithValue(F, [&F](CallInst *Call) {
    const auto x = Call->getArgOperand(0);
    const auto y = Call->getArgOperand(1);
    auto intrinsic = Intrinsic::getDeclaration(
        F.getParent(), Intrinsic::copysign, Call->getType());
    return CallInst::Create(intrinsic->getFunctionType(), intrinsic, {x, y}, "",
                            Call);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceRecip(Function &F) {
  return replaceCallsWithValue(F, [](CallInst *CI) {
    // Recip has one arg.
    auto Arg = CI->getOperand(0);
    auto Cst1 = ConstantFP::get(Arg->getType(), 1.0);
    return BinaryOperator::Create(Instruction::FDiv, Cst1, Arg, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceDivide(Function &F) {
  return replaceCallsWithValue(F, [](CallInst *CI) {
    auto Op0 = CI->getOperand(0);
    auto Op1 = CI->getOperand(1);
    return BinaryOperator::Create(Instruction::FDiv, Op0, Op1, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceDot(Function &F) {
  return replaceCallsWithValue(F, [](CallInst *CI) {
    auto Op0 = CI->getOperand(0);
    auto Op1 = CI->getOperand(1);

    Value *V = nullptr;
    if (Op0->getType()->isVectorTy()) {
      V = clspv::InsertSPIRVOp(CI, spv::OpDot, {Attribute::ReadNone},
                               CI->getType(), {Op0, Op1});
    } else {
      V = BinaryOperator::Create(Instruction::FMul, Op0, Op1, "", CI);
    }

    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceExp10(Function &F,
                                            const std::string &basename) {
  // convert to natural
  auto slen = basename.length() - 2;
  std::string NewFName = basename.substr(0, slen);
  NewFName =
      Builtins::GetMangledFunctionName(NewFName.c_str(), F.getFunctionType());

  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto NewF = M.getOrInsertFunction(NewFName, F.getFunctionType());

    auto Arg = CI->getOperand(0);

    // Constant of the natural log of 10 (ln(10)).
    const double Ln10 =
        2.302585092994045684017991454684364207601101488628772976033;

    auto Mul = BinaryOperator::Create(
        Instruction::FMul, ConstantFP::get(Arg->getType(), Ln10), Arg, "", CI);

    return CallInst::Create(NewF, Mul, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceFmod(Function &F) {
  // OpenCL fmod(x,y) is x - y * trunc(x/y)
  // The sign for a non-zero result is taken from x.
  // (Try an example.)
  // So translate to FRem
  return replaceCallsWithValue(F, [](CallInst *CI) {
    auto Op0 = CI->getOperand(0);
    auto Op1 = CI->getOperand(1);
    return BinaryOperator::Create(Instruction::FRem, Op0, Op1, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceLog10(Function &F,
                                            const std::string &basename) {
  // convert to natural
  auto slen = basename.length() - 2;
  std::string NewFName = basename.substr(0, slen);
  NewFName =
      Builtins::GetMangledFunctionName(NewFName.c_str(), F.getFunctionType());

  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto NewF = M.getOrInsertFunction(NewFName, F.getFunctionType());

    auto Arg = CI->getOperand(0);

    // Constant of the reciprocal of the natural log of 10 (ln(10)).
    const double Ln10 =
        0.434294481903251827651128918916605082294397005803666566114;

    auto NewCI = CallInst::Create(NewF, Arg, "", CI);

    return BinaryOperator::Create(Instruction::FMul,
                                  ConstantFP::get(Arg->getType(), Ln10), NewCI,
                                  "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceLog1p(Function &F) {
  // convert to natural
  return replaceCallsWithValue(F, [&F](CallInst *CI) {
    auto Arg = CI->getOperand(0);

    auto ArgP1 = BinaryOperator::Create(
        Instruction::FAdd, ConstantFP::get(Arg->getType(), 1.0), Arg, "", CI);

    auto log =
        Intrinsic::getDeclaration(F.getParent(), Intrinsic::log, CI->getType());
    return CallInst::Create(log, ArgP1, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceBarrier(Function &F, bool subgroup) {

  enum {
    CLK_LOCAL_MEM_FENCE = 0x01,
    CLK_GLOBAL_MEM_FENCE = 0x02,
    CLK_IMAGE_MEM_FENCE = 0x04
  };

  return replaceCallsWithValue(F, [subgroup](CallInst *CI) {
    auto Arg = CI->getOperand(0);

    // We need to map the OpenCL constants to the SPIR-V equivalents.
    const auto LocalMemFence =
        ConstantInt::get(Arg->getType(), CLK_LOCAL_MEM_FENCE);
    const auto GlobalMemFence =
        ConstantInt::get(Arg->getType(), CLK_GLOBAL_MEM_FENCE);
    const auto ImageMemFence =
        ConstantInt::get(Arg->getType(), CLK_IMAGE_MEM_FENCE);
    const auto ConstantAcquireRelease = ConstantInt::get(
        Arg->getType(), spv::MemorySemanticsAcquireReleaseMask);
    const auto ConstantScopeDevice =
        ConstantInt::get(Arg->getType(), spv::ScopeDevice);
    const auto ConstantScopeWorkgroup =
        ConstantInt::get(Arg->getType(), spv::ScopeWorkgroup);
    const auto ConstantScopeSubgroup =
        ConstantInt::get(Arg->getType(), spv::ScopeSubgroup);

    // Map CLK_LOCAL_MEM_FENCE to MemorySemanticsWorkgroupMemoryMask.
    const auto LocalMemFenceMask =
        BinaryOperator::Create(Instruction::And, LocalMemFence, Arg, "", CI);
    const auto WorkgroupShiftAmount =
        clz(spv::MemorySemanticsWorkgroupMemoryMask) - clz(CLK_LOCAL_MEM_FENCE);
    const auto MemorySemanticsWorkgroup = BinaryOperator::Create(
        Instruction::Shl, LocalMemFenceMask,
        ConstantInt::get(Arg->getType(), WorkgroupShiftAmount), "", CI);

    // Map CLK_GLOBAL_MEM_FENCE to MemorySemanticsUniformMemoryMask.
    const auto GlobalMemFenceMask =
        BinaryOperator::Create(Instruction::And, GlobalMemFence, Arg, "", CI);
    const auto UniformShiftAmount =
        clz(spv::MemorySemanticsUniformMemoryMask) - clz(CLK_GLOBAL_MEM_FENCE);
    const auto MemorySemanticsUniform = BinaryOperator::Create(
        Instruction::Shl, GlobalMemFenceMask,
        ConstantInt::get(Arg->getType(), UniformShiftAmount), "", CI);

    // OpenCL 2.0
    // Map CLK_IMAGE_MEM_FENCE to MemorySemanticsImageMemoryMask.
    const auto ImageMemFenceMask =
        BinaryOperator::Create(Instruction::And, ImageMemFence, Arg, "", CI);
    const auto ImageShiftAmount =
        clz(spv::MemorySemanticsImageMemoryMask) - clz(CLK_IMAGE_MEM_FENCE);
    const auto MemorySemanticsImage = BinaryOperator::Create(
        Instruction::Shl, ImageMemFenceMask,
        ConstantInt::get(Arg->getType(), ImageShiftAmount), "", CI);

    // And combine the above together, also adding in
    // MemorySemanticsSequentiallyConsistentMask.
    auto MemorySemantics1 =
        BinaryOperator::Create(Instruction::Or, MemorySemanticsWorkgroup,
                               ConstantAcquireRelease, "", CI);
    auto MemorySemantics2 = BinaryOperator::Create(
        Instruction::Or, MemorySemanticsUniform, MemorySemanticsImage, "", CI);
    auto MemorySemantics = BinaryOperator::Create(
        Instruction::Or, MemorySemantics1, MemorySemantics2, "", CI);

    // If the memory scope is not specified explicitly, it is either Subgroup
    // or Workgroup depending on the type of barrier.
    Value *MemoryScope =
        subgroup ? ConstantScopeSubgroup : ConstantScopeWorkgroup;
    if (CI->data_operands_size() > 1) {
      enum {
        CL_MEMORY_SCOPE_WORKGROUP = 0x1,
        CL_MEMORY_SCOPE_DEVICE = 0x2,
        CL_MEMORY_SCOPE_SUBGROUP = 0x4
      };
      // The call was given an explicit memory scope.
      const auto MemoryScopeSubgroup =
          ConstantInt::get(Arg->getType(), CL_MEMORY_SCOPE_SUBGROUP);
      const auto MemoryScopeDevice =
          ConstantInt::get(Arg->getType(), CL_MEMORY_SCOPE_DEVICE);

      auto Cmp =
          CmpInst::Create(Instruction::ICmp, CmpInst::ICMP_EQ,
                          MemoryScopeSubgroup, CI->getOperand(1), "", CI);
      MemoryScope = SelectInst::Create(Cmp, ConstantScopeSubgroup,
                                       ConstantScopeWorkgroup, "", CI);
      Cmp = CmpInst::Create(Instruction::ICmp, CmpInst::ICMP_EQ,
                            MemoryScopeDevice, CI->getOperand(1), "", CI);
      MemoryScope =
          SelectInst::Create(Cmp, ConstantScopeDevice, MemoryScope, "", CI);
    }

    // Lastly, the Execution Scope is either Workgroup or Subgroup depending on
    // the type of barrier;
    const auto ExecutionScope =
        subgroup ? ConstantScopeSubgroup : ConstantScopeWorkgroup;

    return clspv::InsertSPIRVOp(CI, spv::OpControlBarrier,
                                {Attribute::NoDuplicate, Attribute::Convergent},
                                CI->getType(),
                                {ExecutionScope, MemoryScope, MemorySemantics});
  });
}

bool ReplaceOpenCLBuiltinPass::replaceMemFence(
    Function &F, spv::MemorySemanticsMask semantics) {

  return replaceCallsWithValue(F, [&](CallInst *CI) {
    enum {
      CLK_LOCAL_MEM_FENCE = 0x01,
      CLK_GLOBAL_MEM_FENCE = 0x02,
      CLK_IMAGE_MEM_FENCE = 0x04,
    };

    auto Arg = CI->getOperand(0);

    // We need to map the OpenCL constants to the SPIR-V equivalents.
    const auto LocalMemFence =
        ConstantInt::get(Arg->getType(), CLK_LOCAL_MEM_FENCE);
    const auto GlobalMemFence =
        ConstantInt::get(Arg->getType(), CLK_GLOBAL_MEM_FENCE);
    const auto ImageMemFence =
        ConstantInt::get(Arg->getType(), CLK_IMAGE_MEM_FENCE);
    const auto ConstantMemorySemantics =
        ConstantInt::get(Arg->getType(), semantics);
    const auto ConstantScopeWorkgroup =
        ConstantInt::get(Arg->getType(), spv::ScopeWorkgroup);

    // Map CLK_LOCAL_MEM_FENCE to MemorySemanticsWorkgroupMemoryMask.
    const auto LocalMemFenceMask =
        BinaryOperator::Create(Instruction::And, LocalMemFence, Arg, "", CI);
    const auto WorkgroupShiftAmount =
        clz(spv::MemorySemanticsWorkgroupMemoryMask) - clz(CLK_LOCAL_MEM_FENCE);
    const auto MemorySemanticsWorkgroup = BinaryOperator::Create(
        Instruction::Shl, LocalMemFenceMask,
        ConstantInt::get(Arg->getType(), WorkgroupShiftAmount), "", CI);

    // Map CLK_GLOBAL_MEM_FENCE to MemorySemanticsUniformMemoryMask.
    const auto GlobalMemFenceMask =
        BinaryOperator::Create(Instruction::And, GlobalMemFence, Arg, "", CI);
    const auto UniformShiftAmount =
        clz(spv::MemorySemanticsUniformMemoryMask) - clz(CLK_GLOBAL_MEM_FENCE);
    const auto MemorySemanticsUniform = BinaryOperator::Create(
        Instruction::Shl, GlobalMemFenceMask,
        ConstantInt::get(Arg->getType(), UniformShiftAmount), "", CI);

    // OpenCL 2.0
    // Map CLK_IMAGE_MEM_FENCE to MemorySemanticsImageMemoryMask.
    const auto ImageMemFenceMask =
        BinaryOperator::Create(Instruction::And, ImageMemFence, Arg, "", CI);
    const auto ImageShiftAmount =
        clz(spv::MemorySemanticsImageMemoryMask) - clz(CLK_IMAGE_MEM_FENCE);
    const auto MemorySemanticsImage = BinaryOperator::Create(
        Instruction::Shl, ImageMemFenceMask,
        ConstantInt::get(Arg->getType(), ImageShiftAmount), "", CI);

    Value *MemOrder = ConstantMemorySemantics;
    Value *MemScope = ConstantScopeWorkgroup;
    IRBuilder<> builder(CI);
    if (CI->arg_size() > 1) {
      MemOrder = MemoryOrderSemantics(CI->getArgOperand(1), false, CI,
                                      semantics, false);
      MemScope = MemoryScope(CI->getArgOperand(2), false, CI);
    }
    // Join the storage semantics and the order semantics.
    auto MemorySemantics1 =
        builder.CreateOr({MemorySemanticsWorkgroup, MemorySemanticsUniform});
    auto MemorySemantics2 = builder.CreateOr({MemorySemanticsImage, MemOrder});
    auto MemorySemantics =
        builder.CreateOr({MemorySemantics1, MemorySemantics2});

    return clspv::InsertSPIRVOp(CI, spv::OpMemoryBarrier,
                                {Attribute::Convergent}, CI->getType(),
                                {MemScope, MemorySemantics});
  });
}

bool ReplaceOpenCLBuiltinPass::replacePrefetch(Function &F) {
  bool Changed = false;

  SmallVector<Instruction *, 4> ToRemoves;

  // Find all calls to the function
  for (auto &U : F.uses()) {
    if (auto CI = dyn_cast<CallInst>(U.getUser())) {
      ToRemoves.push_back(CI);
    }
  }

  Changed = !ToRemoves.empty();

  // Delete them
  for (auto V : ToRemoves) {
    V->eraseFromParent();
  }

  return Changed;
}

bool ReplaceOpenCLBuiltinPass::replaceRelational(Function &F,
                                                 CmpInst::Predicate P) {
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The predicate to use in the CmpInst.
    auto Predicate = P;

    auto Arg1 = CI->getOperand(0);
    auto Arg2 = CI->getOperand(1);

    const auto Cmp =
        CmpInst::Create(Instruction::FCmp, Predicate, Arg1, Arg2, "", CI);
    if (isa<VectorType>(F.getReturnType()))
      return CastInst::Create(Instruction::SExt, Cmp, CI->getType(), "", CI);
    return CastInst::Create(Instruction::ZExt, Cmp, CI->getType(), "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceIsInfAndIsNan(Function &F,
                                                    spv::Op SPIRVOp,
                                                    int32_t C) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    const auto CITy = CI->getType();

    // The value to return for true.
    auto TrueValue = ConstantInt::getSigned(CITy, C);

    // The value to return for false.
    auto FalseValue = Constant::getNullValue(CITy);

    Type *CorrespondingBoolTy = Type::getInt1Ty(M.getContext());
    if (auto CIVecTy = dyn_cast<VectorType>(CITy)) {
      CorrespondingBoolTy =
          FixedVectorType::get(Type::getInt1Ty(M.getContext()),
                               CIVecTy->getElementCount().getKnownMinValue());
    }

    auto NewCI = clspv::InsertSPIRVOp(CI, SPIRVOp, {Attribute::ReadNone},
                                      CorrespondingBoolTy, {CI->getOperand(0)});

    return SelectInst::Create(NewCI, TrueValue, FalseValue, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceIsFinite(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto &C = M.getContext();
    auto Val = CI->getOperand(0);
    auto ValTy = Val->getType();
    auto RetTy = CI->getType();

    // Get a suitable integer type to represent the number
    auto IntTy = getIntOrIntVectorTyForCast(C, ValTy);

    // Create Mask
    auto ScalarSize = ValTy->getScalarSizeInBits();
    Value *InfMask = nullptr;
    switch (ScalarSize) {
    case 16:
      InfMask = ConstantInt::get(IntTy, 0x7C00U);
      break;
    case 32:
      InfMask = ConstantInt::get(IntTy, 0x7F800000U);
      break;
    case 64:
      InfMask = ConstantInt::get(IntTy, 0x7FF0000000000000ULL);
      break;
    default:
      llvm_unreachable("Unsupported floating-point type");
    }

    IRBuilder<> Builder(CI);

    // Bitcast to int
    auto ValInt = Builder.CreateBitCast(Val, IntTy);

    // Mask and compare
    auto InfBits = Builder.CreateAnd(InfMask, ValInt);
    auto Cmp = Builder.CreateICmp(CmpInst::ICMP_EQ, InfBits, InfMask);

    auto RetFalse = ConstantInt::get(RetTy, 0);
    Value *RetTrue = nullptr;
    if (ValTy->isVectorTy()) {
      RetTrue = ConstantInt::getSigned(RetTy, -1);
    } else {
      RetTrue = ConstantInt::get(RetTy, 1);
    }
    return Builder.CreateSelect(Cmp, RetFalse, RetTrue);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAllAndAny(Function &F, spv::Op SPIRVOp) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto Arg = CI->getOperand(0);

    Value *V = nullptr;

    // If the argument is a 32-bit int, just use a shift
    if (Arg->getType() == Type::getInt32Ty(M.getContext())) {
      V = BinaryOperator::Create(Instruction::LShr, Arg,
                                 ConstantInt::get(Arg->getType(), 31), "", CI);
    } else {
      // The value for zero to compare against.
      const auto ZeroValue = Constant::getNullValue(Arg->getType());

      // The value to return for true.
      const auto TrueValue = ConstantInt::get(CI->getType(), 1);

      // The value to return for false.
      const auto FalseValue = Constant::getNullValue(CI->getType());

      const auto Cmp = CmpInst::Create(Instruction::ICmp, CmpInst::ICMP_SLT,
                                       Arg, ZeroValue, "", CI);

      Value *SelectSource = nullptr;

      // If we have a function to call, call it!
      if (SPIRVOp != spv::OpNop) {

        const auto BoolTy = Type::getInt1Ty(M.getContext());

        const auto NewCI = clspv::InsertSPIRVOp(
            CI, SPIRVOp, {Attribute::ReadNone}, BoolTy, {Cmp});
        SelectSource = NewCI;

      } else {
        SelectSource = Cmp;
      }

      V = SelectInst::Create(SelectSource, TrueValue, FalseValue, "", CI);
    }
    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceUpsample(Function &F) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    // Get arguments
    auto HiValue = CI->getOperand(0);
    auto LoValue = CI->getOperand(1);

    // Don't touch overloads that aren't in OpenCL C
    auto HiType = HiValue->getType();
    auto LoType = LoValue->getType();

    if (HiType != LoType) {
      return nullptr;
    }

    if (!HiType->isIntOrIntVectorTy()) {
      return nullptr;
    }

    if (HiType->getScalarSizeInBits() * 2 !=
        CI->getType()->getScalarSizeInBits()) {
      return nullptr;
    }

    if ((HiType->getScalarSizeInBits() != 8) &&
        (HiType->getScalarSizeInBits() != 16) &&
        (HiType->getScalarSizeInBits() != 32)) {
      return nullptr;
    }

    if (auto HiVecType = dyn_cast<VectorType>(HiType)) {
      unsigned NumElements = HiVecType->getElementCount().getKnownMinValue();
      if ((NumElements != 2) && (NumElements != 3) && (NumElements != 4) &&
          (NumElements != 8) && (NumElements != 16)) {
        return nullptr;
      }
    }

    // Convert both operands to the result type
    auto HiCast = CastInst::CreateZExtOrBitCast(HiValue, CI->getType(), "", CI);
    auto LoCast = CastInst::CreateZExtOrBitCast(LoValue, CI->getType(), "", CI);

    // Shift high operand
    auto ShiftAmount =
        ConstantInt::get(CI->getType(), HiType->getScalarSizeInBits());
    auto HiShifted =
        BinaryOperator::Create(Instruction::Shl, HiCast, ShiftAmount, "", CI);

    // OR both results
    return BinaryOperator::Create(Instruction::Or, HiShifted, LoCast, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceRotate(Function &F) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    // Get arguments
    auto SrcValue = CI->getOperand(0);
    auto RotAmount = CI->getOperand(1);

    // Don't touch overloads that aren't in OpenCL C
    auto SrcType = SrcValue->getType();
    auto RotType = RotAmount->getType();

    if ((SrcType != RotType) || (CI->getType() != SrcType)) {
      return nullptr;
    }

    if (!SrcType->isIntOrIntVectorTy()) {
      return nullptr;
    }

    if ((SrcType->getScalarSizeInBits() != 8) &&
        (SrcType->getScalarSizeInBits() != 16) &&
        (SrcType->getScalarSizeInBits() != 32) &&
        (SrcType->getScalarSizeInBits() != 64)) {
      return nullptr;
    }

    if (auto SrcVecType = dyn_cast<VectorType>(SrcType)) {
      unsigned NumElements = SrcVecType->getElementCount().getKnownMinValue();
      if ((NumElements != 2) && (NumElements != 3) && (NumElements != 4) &&
          (NumElements != 8) && (NumElements != 16)) {
        return nullptr;
      }
    }

    // Replace with LLVM's funnel shift left intrinsic because it is more
    // generic than rotate.
    Function *intrinsic =
        Intrinsic::getDeclaration(F.getParent(), Intrinsic::fshl, SrcType);
    return CallInst::Create(intrinsic->getFunctionType(), intrinsic,
                            {SrcValue, SrcValue, RotAmount}, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceConvert(Function &F, bool SrcIsSigned,
                                              bool DstIsSigned) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    Value *V = nullptr;
    // Get arguments
    auto SrcValue = CI->getOperand(0);

    // Don't touch overloads that aren't in OpenCL C
    auto SrcType = SrcValue->getType();
    auto DstType = CI->getType();

    if ((SrcType->isVectorTy() && !DstType->isVectorTy()) ||
        (!SrcType->isVectorTy() && DstType->isVectorTy())) {
      return V;
    }

    if (auto SrcVecType = dyn_cast<VectorType>(SrcType)) {
      unsigned SrcNumElements =
          SrcVecType->getElementCount().getKnownMinValue();
      unsigned DstNumElements =
          cast<VectorType>(DstType)->getElementCount().getKnownMinValue();
      if (SrcNumElements != DstNumElements) {
        return V;
      }

      if ((SrcNumElements != 2) && (SrcNumElements != 3) &&
          (SrcNumElements != 4) && (SrcNumElements != 8) &&
          (SrcNumElements != 16)) {
        return V;
      }
    }

    bool SrcIsFloat = SrcType->getScalarType()->isFloatingPointTy();
    bool DstIsFloat = DstType->getScalarType()->isFloatingPointTy();

    bool SrcIsInt = SrcType->isIntOrIntVectorTy();
    bool DstIsInt = DstType->isIntOrIntVectorTy();

    if (SrcType == DstType && DstIsSigned == SrcIsSigned) {
      // Unnecessary cast operation.
      V = SrcValue;
    } else if (SrcIsFloat && DstIsFloat) {
      V = CastInst::CreateFPCast(SrcValue, DstType, "", CI);
    } else if (SrcIsFloat && DstIsInt) {
      if (DstIsSigned) {
        V = CastInst::Create(Instruction::FPToSI, SrcValue, DstType, "", CI);
      } else {
        V = CastInst::Create(Instruction::FPToUI, SrcValue, DstType, "", CI);
      }
    } else if (SrcIsInt && DstIsFloat) {
      if (SrcIsSigned) {
        V = CastInst::Create(Instruction::SIToFP, SrcValue, DstType, "", CI);
      } else {
        V = CastInst::Create(Instruction::UIToFP, SrcValue, DstType, "", CI);
      }
    } else if (SrcIsInt && DstIsInt) {
      V = CastInst::CreateIntegerCast(SrcValue, DstType, SrcIsSigned, "", CI);
    } else {
      // Not something we're supposed to handle, just move on
    }

    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceMulHi(Function &F, bool is_signed,
                                            bool is_mad) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    Value *V = nullptr;
    // Get arguments
    auto AValue = CI->getOperand(0);
    auto BValue = CI->getOperand(1);
    auto CValue = CI->getOperand(2);

    // Don't touch overloads that aren't in OpenCL C
    auto AType = AValue->getType();
    auto BType = BValue->getType();
    auto CType = CValue->getType();

    if ((AType != BType) || (CI->getType() != AType) ||
        (is_mad && (AType != CType))) {
      return V;
    }

    if (!AType->isIntOrIntVectorTy()) {
      return V;
    }

    if ((AType->getScalarSizeInBits() != 8) &&
        (AType->getScalarSizeInBits() != 16) &&
        (AType->getScalarSizeInBits() != 32) &&
        (AType->getScalarSizeInBits() != 64)) {
      return V;
    }

    if (auto AVecType = dyn_cast<VectorType>(AType)) {
      unsigned NumElements = AVecType->getElementCount().getKnownMinValue();
      if ((NumElements != 2) && (NumElements != 3) && (NumElements != 4) &&
          (NumElements != 8) && (NumElements != 16)) {
        return V;
      }
    }

    auto Call = InsertOpMulExtended(CI, AValue, BValue, is_signed);

    // Get the high part of the result
    unsigned Idxs[] = {1};
    V = ExtractValueInst::Create(Call, Idxs, "", CI);

    // If we're handling a mad_hi, add the third argument to the result
    if (is_mad) {
      V = BinaryOperator::Create(Instruction::Add, V, CValue, "", CI);
    }

    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceSelect(Function &F) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    // Get arguments
    auto FalseValue = CI->getOperand(0);
    auto TrueValue = CI->getOperand(1);
    auto PredicateValue = CI->getOperand(2);

    // Don't touch overloads that aren't in OpenCL C
    auto FalseType = FalseValue->getType();
    auto TrueType = TrueValue->getType();
    auto PredicateType = PredicateValue->getType();

    if (FalseType != TrueType) {
      return nullptr;
    }

    if (!PredicateType->isIntOrIntVectorTy()) {
      return nullptr;
    }

    if (!FalseType->isIntOrIntVectorTy() &&
        !FalseType->getScalarType()->isFloatingPointTy()) {
      return nullptr;
    }

    if (FalseType->isVectorTy() && !PredicateType->isVectorTy()) {
      return nullptr;
    }

    if (FalseType->getScalarSizeInBits() !=
        PredicateType->getScalarSizeInBits()) {
      return nullptr;
    }

    if (auto FalseVecType = dyn_cast<VectorType>(FalseType)) {
      unsigned NumElements = FalseVecType->getElementCount().getKnownMinValue();
      if (NumElements != cast<VectorType>(PredicateType)
                             ->getElementCount()
                             .getKnownMinValue()) {
        return nullptr;
      }

      if ((NumElements != 2) && (NumElements != 3) && (NumElements != 4) &&
          (NumElements != 8) && (NumElements != 16)) {
        return nullptr;
      }
    }

    // Create constant
    const auto ZeroValue = Constant::getNullValue(PredicateType);

    // Scalar and vector are to be treated differently
    CmpInst::Predicate Pred;
    if (PredicateType->isVectorTy()) {
      Pred = CmpInst::ICMP_SLT;
    } else {
      Pred = CmpInst::ICMP_NE;
    }

    // Create comparison instruction
    auto Cmp = CmpInst::Create(Instruction::ICmp, Pred, PredicateValue,
                               ZeroValue, "", CI);

    // Create select
    return SelectInst::Create(Cmp, TrueValue, FalseValue, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceBitSelect(Function &F) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    Value *V = nullptr;
    if (CI->getNumOperands() != 4) {
      return V;
    }

    // Get arguments
    auto FalseValue = CI->getOperand(0);
    auto TrueValue = CI->getOperand(1);
    auto PredicateValue = CI->getOperand(2);

    // Don't touch overloads that aren't in OpenCL C
    auto FalseType = FalseValue->getType();
    auto TrueType = TrueValue->getType();
    auto PredicateType = PredicateValue->getType();

    if ((FalseType != TrueType) || (PredicateType != TrueType)) {
      return V;
    }

    if (auto TrueVecType = dyn_cast<VectorType>(TrueType)) {
      if (!TrueType->getScalarType()->isFloatingPointTy() &&
          !TrueType->getScalarType()->isIntegerTy()) {
        return V;
      }
      unsigned NumElements = TrueVecType->getElementCount().getKnownMinValue();
      if ((NumElements != 2) && (NumElements != 3) && (NumElements != 4) &&
          (NumElements != 8) && (NumElements != 16)) {
        return V;
      }
    }

    // Remember the type of the operands
    auto OpType = TrueType;

    // The actual bit selection will always be done on an integer type,
    // declare it here
    Type *BitType;

    // If the operands are float, then bitcast them to int
    if (OpType->getScalarType()->isFloatingPointTy()) {

      // First create the new type
      BitType = getIntOrIntVectorTyForCast(F.getContext(), OpType);

      // Then bitcast all operands
      PredicateValue =
          CastInst::CreateZExtOrBitCast(PredicateValue, BitType, "", CI);
      FalseValue = CastInst::CreateZExtOrBitCast(FalseValue, BitType, "", CI);
      TrueValue = CastInst::CreateZExtOrBitCast(TrueValue, BitType, "", CI);

    } else {
      // The operands have an integer type, use it directly
      BitType = OpType;
    }

    // All the operands are now always integers
    // implement as (c & b) | (~c & a)

    // Create our negated predicate value
    auto AllOnes = Constant::getAllOnesValue(BitType);
    auto NotPredicateValue = BinaryOperator::Create(
        Instruction::Xor, PredicateValue, AllOnes, "", CI);

    // Then put everything together
    auto BitsFalse = BinaryOperator::Create(Instruction::And, NotPredicateValue,
                                            FalseValue, "", CI);
    auto BitsTrue = BinaryOperator::Create(Instruction::And, PredicateValue,
                                           TrueValue, "", CI);

    V = BinaryOperator::Create(Instruction::Or, BitsFalse, BitsTrue, "", CI);

    // If we were dealing with a floating point type, we must bitcast
    // the result back to that
    if (OpType->getScalarType()->isFloatingPointTy()) {
      V = CastInst::CreateZExtOrBitCast(V, OpType, "", CI);
    }

    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceStep(Function &F, bool is_smooth) {
  // convert to vector versions
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    SmallVector<Value *, 2> ArgsToSplat = {CI->getOperand(0)};
    Value *VectorArg = nullptr;

    // First figure out which function we're dealing with
    if (is_smooth) {
      ArgsToSplat.push_back(CI->getOperand(1));
      VectorArg = CI->getOperand(2);
    } else {
      VectorArg = CI->getOperand(1);
    }

    // Splat arguments that need to be
    SmallVector<Value *, 2> SplatArgs;
    auto VecType = cast<VectorType>(VectorArg->getType());

    for (auto arg : ArgsToSplat) {
      Value *NewVectorArg = UndefValue::get(VecType);
      for (size_t i = 0; i < VecType->getElementCount().getKnownMinValue();
           i++) {
        auto index = ConstantInt::get(Type::getInt32Ty(M.getContext()), i);
        NewVectorArg =
            InsertElementInst::Create(NewVectorArg, arg, index, "", CI);
      }
      SplatArgs.push_back(NewVectorArg);
    }

    // Replace the call with the vector/vector flavour
    SmallVector<Type *, 3> NewArgTypes(ArgsToSplat.size() + 1, VecType);
    const auto NewFType = FunctionType::get(CI->getType(), NewArgTypes, false);

    std::string NewFName = Builtins::GetMangledFunctionName(
        is_smooth ? "smoothstep" : "step", NewFType);

    const auto NewF = M.getOrInsertFunction(NewFName, NewFType);

    SmallVector<Value *, 3> NewArgs;
    for (auto arg : SplatArgs) {
      NewArgs.push_back(arg);
    }
    NewArgs.push_back(VectorArg);

    return CallInst::Create(NewF, NewArgs, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceSignbit(Function &F, bool is_vec) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    auto Arg = CI->getOperand(0);
    auto Op = is_vec ? Instruction::AShr : Instruction::LShr;

    auto Bitcast = CastInst::CreateZExtOrBitCast(Arg, CI->getType(), "", CI);

    return BinaryOperator::Create(Op, Bitcast,
                                  ConstantInt::get(CI->getType(), 31), "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceMul(Function &F, bool is_float,
                                          bool is_mad) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    // The multiply instruction to use.
    auto MulInst = is_float ? Instruction::FMul : Instruction::Mul;

    SmallVector<Value *, 8> Args(CI->arg_begin(), CI->arg_end());

    Value *V = BinaryOperator::Create(MulInst, CI->getArgOperand(0),
                                      CI->getArgOperand(1), "", CI);

    if (is_mad) {
      // The add instruction to use.
      auto AddInst = is_float ? Instruction::FAdd : Instruction::Add;

      V = BinaryOperator::Create(AddInst, V, CI->getArgOperand(2), "", CI);
    }

    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstore(Function &F) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    Value *V = nullptr;
    auto data = CI->getOperand(0);

    auto data_type = data->getType();
    if (!data_type->isVectorTy())
      return V;

    auto vec_data_type = cast<VectorType>(data_type);

    auto elems = vec_data_type->getElementCount().getKnownMinValue();
    if (elems != 2 && elems != 3 && elems != 4 && elems != 8 && elems != 16)
      return V;

    auto offset = CI->getOperand(1);
    auto ptr = CI->getOperand(2);

    // Avoid pointer casts. Instead generate the correct number of stores
    // and rely on drivers to coalesce appropriately.
    IRBuilder<> builder(CI);
    auto elems_const = builder.getInt32(elems);
    auto adjust = builder.CreateMul(offset, elems_const);
    for (size_t i = 0; i < elems; ++i) {
      auto idx = builder.getInt32(i);
      auto add = builder.CreateAdd(adjust, idx);
      auto gep = builder.CreateGEP(vec_data_type->getScalarType(), ptr, add);
      auto extract = builder.CreateExtractElement(data, i);
      V = builder.CreateStore(extract, gep);
    }
    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVload(Function &F) {
  return replaceCallsWithValue(F, [&](CallInst *CI) -> llvm::Value * {
    Value *V = nullptr;
    auto ret_type = F.getReturnType();
    if (!ret_type->isVectorTy())
      return V;

    auto vec_ret_type = cast<VectorType>(ret_type);

    auto elems = vec_ret_type->getElementCount().getKnownMinValue();
    if (elems != 2 && elems != 3 && elems != 4 && elems != 8 && elems != 16)
      return V;

    auto offset = CI->getOperand(0);
    auto ptr = CI->getOperand(1);

    // Avoid pointer casts. Instead generate the correct number of loads
    // and rely on drivers to coalesce appropriately.
    IRBuilder<> builder(CI);
    auto elems_const = builder.getInt32(elems);
    V = UndefValue::get(ret_type);
    auto adjust = builder.CreateMul(offset, elems_const);
    for (unsigned i = 0; i < elems; ++i) {
      auto idx = builder.getInt32(i);
      auto add = builder.CreateAdd(adjust, idx);
      auto gep = builder.CreateGEP(vec_ret_type->getScalarType(), ptr, add);
      auto load = builder.CreateLoad(vec_ret_type->getScalarType(), gep);
      V = builder.CreateInsertElement(V, load, i);
    }
    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVloadHalf(Function &F,
                                                const std::string &name,
                                                int vec_size, bool aligned) {
  bool is_clspv_version = !name.compare(0, 8, "__clspv_");
  if (!vec_size) {
    // deduce vec_size from last characters of name (e.g. vload_half4)
    std::string half = "half";
    vec_size = std::atoi(
        name.substr(name.find(half) + half.size(), std::string::npos).c_str());
  }
  switch (vec_size) {
  case 2:
    return is_clspv_version ? replaceClspvVloadaHalf2(F) : replaceVloadHalf2(F);
  case 3:
    if (!is_clspv_version) {
      return aligned ? replaceVloadaHalf3(F) : replaceVloadHalf3(F);
    }
    break;
  case 4:
    return is_clspv_version ? replaceClspvVloadaHalf4(F) : replaceVloadHalf4(F);
  case 8:
    if (!is_clspv_version) {
      return replaceVloadHalf8(F);
    }
    break;
  case 16:
    if (!is_clspv_version) {
      return replaceVloadHalf16(F);
    }
    break;
  case 0:
    if (!is_clspv_version) {
      return replaceVloadHalf(F);
    }
    break;
  }
  llvm_unreachable("Unsupported vload_half vector size");
}

bool ReplaceOpenCLBuiltinPass::replaceVloadHalf(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The index argument from vload_half.
    auto Arg0 = CI->getOperand(0);

    // The pointer argument from vload_half.
    auto Arg1 = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    Value *V = nullptr;

    bool supports_16bit_storage = true;
    switch (Arg1->getType()->getPointerAddressSpace()) {
    case clspv::AddressSpace::Global:
      supports_16bit_storage = clspv::Option::Supports16BitStorageClass(
          clspv::Option::StorageClass::kSSBO);
      break;
    case clspv::AddressSpace::Constant:
      if (clspv::Option::ConstantArgsInUniformBuffer())
        supports_16bit_storage = clspv::Option::Supports16BitStorageClass(
            clspv::Option::StorageClass::kUBO);
      else
        supports_16bit_storage = clspv::Option::Supports16BitStorageClass(
            clspv::Option::StorageClass::kSSBO);
      break;
    default:
      // Clspv will emit the Float16 capability if the half type is
      // encountered. That capability covers private and local addressspaces.
      break;
    }

    if (supports_16bit_storage) {
      auto ShortTy = Type::getInt16Ty(M.getContext());
      auto ShortPointerTy =
          PointerType::get(ShortTy, Arg1->getType()->getPointerAddressSpace());

      // Cast the half* pointer to short*.
      auto Cast = CastInst::CreatePointerCast(Arg1, ShortPointerTy, "", CI);

      // Index into the correct address of the casted pointer.
      auto Index = GetElementPtrInst::Create(ShortTy, Cast, Arg0, "", CI);

      // Load from the short* we casted to.
      auto Load = new LoadInst(ShortTy, Index, "", CI);

      // ZExt the short -> int.
      auto ZExt = CastInst::CreateZExtOrBitCast(Load, IntTy, "", CI);

      // Get our float2.
      auto Call = CallInst::Create(NewF, ZExt, "", CI);

      // Extract out the bottom element which is our float result.
      V = ExtractElementInst::Create(Call, ConstantInt::get(IntTy, 0), "", CI);
    } else {
      // Assume the pointer argument points to storage aligned to 32bits
      // or more.
      // TODO(dneto): Do more analysis to make sure this is true?
      //
      // Replace call vstore_half(i32 %index, half addrspace(1) %base)
      // with:
      //
      //   %base_i32_ptr = bitcast half addrspace(1)* %base to i32
      //   addrspace(1)* %index_is_odd32 = and i32 %index, 1 %index_i32 =
      //   lshr i32 %index, 1 %in_ptr = getlementptr i32, i32
      //   addrspace(1)* %base_i32_ptr, %index_i32 %value_i32 = load i32,
      //   i32 addrspace(1)* %in_ptr %converted = call <2 x float>
      //   @spirv.unpack.v2f16(i32 %value_i32) %value = extractelement <2
      //   x float> %converted, %index_is_odd32

      auto IntPointerTy =
          PointerType::get(IntTy, Arg1->getType()->getPointerAddressSpace());

      // Cast the base pointer to int*.
      // In a valid call (according to assumptions), this should get
      // optimized away in the simplify GEP pass.
      auto Cast = CastInst::CreatePointerCast(Arg1, IntPointerTy, "", CI);

      auto One = ConstantInt::get(IntTy, 1);
      auto IndexIsOdd = BinaryOperator::CreateAnd(Arg0, One, "", CI);
      auto IndexIntoI32 = BinaryOperator::CreateLShr(Arg0, One, "", CI);

      // Index into the correct address of the casted pointer.
      auto Ptr = GetElementPtrInst::Create(IntTy, Cast, IndexIntoI32, "", CI);

      // Load from the int* we casted to.
      auto Load = new LoadInst(IntTy, Ptr, "", CI);

      // Get our float2.
      auto Call = CallInst::Create(NewF, Load, "", CI);

      // Extract out the float result, where the element number is
      // determined by whether the original index was even or odd.
      V = ExtractElementInst::Create(Call, IndexIsOdd, "", CI);
    }
    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVloadHalf2(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The index argument from vload_half.
    auto Arg0 = CI->getOperand(0);

    // The pointer argument from vload_half.
    auto Arg1 = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(IntTy, Arg1->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    // Cast the half* pointer to int*.
    auto Cast = CastInst::CreatePointerCast(Arg1, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Index = GetElementPtrInst::Create(IntTy, Cast, Arg0, "", CI);

    // Load from the int* we casted to.
    auto Load = new LoadInst(IntTy, Index, "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Get our float2.
    return CallInst::Create(NewF, Load, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVloadHalf3(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The index argument from vload_half.
    auto Arg0 = CI->getOperand(0);

    // The pointer argument from vload_half.
    auto Arg1 = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto ShortTy = Type::getInt16Ty(M.getContext());
    auto FloatTy = Type::getFloatTy(M.getContext());
    auto Float2Ty = FixedVectorType::get(FloatTy, 2);
    auto Float3Ty = FixedVectorType::get(FloatTy, 3);
    auto NewPointerTy =
        PointerType::get(ShortTy, Arg1->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    auto Int0 = ConstantInt::get(IntTy, 0);
    auto Int1 = ConstantInt::get(IntTy, 1);
    auto Int2 = ConstantInt::get(IntTy, 2);

    // Cast the half* pointer to short*.
    auto Cast = CastInst::CreatePointerCast(Arg1, NewPointerTy, "", CI);

    // Load the first element
    auto Index0 = BinaryOperator::Create(
        Instruction::Add,
        BinaryOperator::Create(Instruction::Shl, Arg0, Int1, "", CI), Arg0, "",
        CI);
    auto GEP0 = GetElementPtrInst::Create(ShortTy, Cast, Index0, "", CI);
    auto Load0 = new LoadInst(ShortTy, GEP0, "", CI);

    // Load the second element
    auto Index1 =
        BinaryOperator::Create(Instruction::Add, Index0, Int1, "", CI);
    auto GEP1 = GetElementPtrInst::Create(ShortTy, Cast, Index1, "", CI);
    auto Load1 = new LoadInst(ShortTy, GEP1, "", CI);

    // Load the third element
    auto Index2 =
        BinaryOperator::Create(Instruction::Add, Index1, Int1, "", CI);
    auto GEP2 = GetElementPtrInst::Create(ShortTy, Cast, Index2, "", CI);
    auto Load2 = new LoadInst(ShortTy, GEP2, "", CI);

    // Extend each short to int.
    auto X0 = CastInst::Create(Instruction::ZExt, Load0, IntTy, "", CI);
    auto X1 = CastInst::Create(Instruction::ZExt, Load1, IntTy, "", CI);
    auto X2 = CastInst::Create(Instruction::ZExt, Load2, IntTy, "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Convert int to float2 and extract the uniq meaningful float
    auto Y0 = ExtractElementInst::Create(CallInst::Create(NewF, X0, "", CI),
                                         Int0, "", CI);
    auto Y1 = ExtractElementInst::Create(CallInst::Create(NewF, X1, "", CI),
                                         Int0, "", CI);
    auto Y2 = ExtractElementInst::Create(CallInst::Create(NewF, X2, "", CI),
                                         Int0, "", CI);

    // Create the final float3 to be returned
    auto Combine =
        InsertElementInst::Create(UndefValue::get(Float3Ty), Y0, Int0, "", CI);
    Combine = InsertElementInst::Create(Combine, Y1, Int1, "", CI);
    Combine = InsertElementInst::Create(Combine, Y2, Int2, "", CI);

    return Combine;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVloadaHalf3(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The index argument from vload_half.
    auto Arg0 = CI->getOperand(0);

    // The pointer argument from vload_half.
    auto Arg1 = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int2Ty = FixedVectorType::get(IntTy, 2);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(Int2Ty, Arg1->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    // Cast the half* pointer to int2*.
    auto Cast = CastInst::CreatePointerCast(Arg1, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Index = GetElementPtrInst::Create(Int2Ty, Cast, Arg0, "", CI);

    // Load from the int2* we casted to.
    auto Load = new LoadInst(Int2Ty, Index, "", CI);

    // Extract each element from the loaded int2.
    auto X =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 0), "", CI);
    auto Y =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 1), "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Get the lower (x & y) components of our final float4.
    auto Lo = CallInst::Create(NewF, X, "", CI);

    // Get the higher (z & w) components of our final float4.
    auto Hi = CallInst::Create(NewF, Y, "", CI);

    Constant *ShuffleMask[3] = {ConstantInt::get(IntTy, 0),
                                ConstantInt::get(IntTy, 1),
                                ConstantInt::get(IntTy, 2)};

    // Combine our two float2's into one float4.
    return new ShuffleVectorInst(Lo, Hi, ConstantVector::get(ShuffleMask), "",
                                 CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVloadHalf4(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The index argument from vload_half.
    auto Arg0 = CI->getOperand(0);

    // The pointer argument from vload_half.
    auto Arg1 = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int2Ty = FixedVectorType::get(IntTy, 2);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(Int2Ty, Arg1->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    // Cast the half* pointer to int2*.
    auto Cast = CastInst::CreatePointerCast(Arg1, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Index = GetElementPtrInst::Create(Int2Ty, Cast, Arg0, "", CI);

    // Load from the int2* we casted to.
    auto Load = new LoadInst(Int2Ty, Index, "", CI);

    // Extract each element from the loaded int2.
    auto X =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 0), "", CI);
    auto Y =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 1), "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Get the lower (x & y) components of our final float4.
    auto Lo = CallInst::Create(NewF, X, "", CI);

    // Get the higher (z & w) components of our final float4.
    auto Hi = CallInst::Create(NewF, Y, "", CI);

    Constant *ShuffleMask[4] = {
        ConstantInt::get(IntTy, 0), ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2), ConstantInt::get(IntTy, 3)};

    // Combine our two float2's into one float4.
    return new ShuffleVectorInst(Lo, Hi, ConstantVector::get(ShuffleMask), "",
                                 CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVloadHalf8(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The index argument from vload_half.
    auto Arg0 = CI->getOperand(0);

    // The pointer argument from vload_half.
    auto Arg1 = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int4Ty = FixedVectorType::get(IntTy, 4);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(Int4Ty, Arg1->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    // Cast the half* pointer to int4*.
    auto Cast = CastInst::CreatePointerCast(Arg1, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Index = GetElementPtrInst::Create(Int4Ty, Cast, Arg0, "", CI);

    // Load from the int4* we casted to.
    auto Load = new LoadInst(Int4Ty, Index, "", CI);

    // Extract each element from the loaded int4.
    auto X1 =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 0), "", CI);
    auto X2 =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 1), "", CI);
    auto X3 =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 2), "", CI);
    auto X4 =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 3), "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Convert the 4 int into 4 float2
    auto Y1 = CallInst::Create(NewF, X1, "", CI);
    auto Y2 = CallInst::Create(NewF, X2, "", CI);
    auto Y3 = CallInst::Create(NewF, X3, "", CI);
    auto Y4 = CallInst::Create(NewF, X4, "", CI);

    Constant *ShuffleMask4[4] = {
        ConstantInt::get(IntTy, 0), ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2), ConstantInt::get(IntTy, 3)};

    // Combine our two float2's into one float4.
    auto Z1 = new ShuffleVectorInst(Y1, Y2, ConstantVector::get(ShuffleMask4),
                                    "", CI);
    auto Z2 = new ShuffleVectorInst(Y3, Y4, ConstantVector::get(ShuffleMask4),
                                    "", CI);

    Constant *ShuffleMask8[8] = {
        ConstantInt::get(IntTy, 0), ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2), ConstantInt::get(IntTy, 3),
        ConstantInt::get(IntTy, 4), ConstantInt::get(IntTy, 5),
        ConstantInt::get(IntTy, 6), ConstantInt::get(IntTy, 7)};

    // Combine our two float4's into one float8.
    return new ShuffleVectorInst(Z1, Z2, ConstantVector::get(ShuffleMask8), "",
                                 CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVloadHalf16(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The index argument from vload_half.
    auto Arg0 = CI->getOperand(0);

    // The pointer argument from vload_half.
    auto Arg1 = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int4Ty = FixedVectorType::get(IntTy, 4);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(Int4Ty, Arg1->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    // Cast the half* pointer to int4*.
    auto Cast = CastInst::CreatePointerCast(Arg1, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Arg0x2 = BinaryOperator::Create(Instruction::Shl, Arg0, ConstantInt::get(IntTy, 1), "", CI);
    auto Index1 = GetElementPtrInst::Create(Int4Ty, Cast, Arg0x2, "", CI);
    auto Arg0x2p1 = BinaryOperator::Create(Instruction::Add, Arg0x2, ConstantInt::get(IntTy, 1), "", CI);
    auto Index2 = GetElementPtrInst::Create(Int4Ty, Cast, Arg0x2p1, "", CI);

    // Load from the int4* we casted to.
    auto Load1 = new LoadInst(Int4Ty, Index1, "", CI);
    auto Load2 = new LoadInst(Int4Ty, Index2, "", CI);

    // Extract each element from the two loaded int4.
    auto X1 =
        ExtractElementInst::Create(Load1, ConstantInt::get(IntTy, 0), "", CI);
    auto X2 =
        ExtractElementInst::Create(Load1, ConstantInt::get(IntTy, 1), "", CI);
    auto X3 =
        ExtractElementInst::Create(Load1, ConstantInt::get(IntTy, 2), "", CI);
    auto X4 =
        ExtractElementInst::Create(Load1, ConstantInt::get(IntTy, 3), "", CI);
    auto X5 =
        ExtractElementInst::Create(Load2, ConstantInt::get(IntTy, 0), "", CI);
    auto X6 =
        ExtractElementInst::Create(Load2, ConstantInt::get(IntTy, 1), "", CI);
    auto X7 =
        ExtractElementInst::Create(Load2, ConstantInt::get(IntTy, 2), "", CI);
    auto X8 =
        ExtractElementInst::Create(Load2, ConstantInt::get(IntTy, 3), "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Convert the eight int into float2
    auto Y1 = CallInst::Create(NewF, X1, "", CI);
    auto Y2 = CallInst::Create(NewF, X2, "", CI);
    auto Y3 = CallInst::Create(NewF, X3, "", CI);
    auto Y4 = CallInst::Create(NewF, X4, "", CI);
    auto Y5 = CallInst::Create(NewF, X5, "", CI);
    auto Y6 = CallInst::Create(NewF, X6, "", CI);
    auto Y7 = CallInst::Create(NewF, X7, "", CI);
    auto Y8 = CallInst::Create(NewF, X8, "", CI);

    Constant *ShuffleMask4[4] = {
        ConstantInt::get(IntTy, 0), ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2), ConstantInt::get(IntTy, 3)};

    // Combine our two float2's into one float4.
    auto Z1 = new ShuffleVectorInst(Y1, Y2, ConstantVector::get(ShuffleMask4),
                                    "", CI);
    auto Z2 = new ShuffleVectorInst(Y3, Y4, ConstantVector::get(ShuffleMask4),
                                    "", CI);
    auto Z3 = new ShuffleVectorInst(Y5, Y6, ConstantVector::get(ShuffleMask4),
                                    "", CI);
    auto Z4 = new ShuffleVectorInst(Y7, Y8, ConstantVector::get(ShuffleMask4),
                                    "", CI);

    Constant *ShuffleMask8[8] = {
        ConstantInt::get(IntTy, 0), ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2), ConstantInt::get(IntTy, 3),
        ConstantInt::get(IntTy, 4), ConstantInt::get(IntTy, 5),
        ConstantInt::get(IntTy, 6), ConstantInt::get(IntTy, 7)};

    // Combine our two float4's into one float8.
    auto Z5 = new ShuffleVectorInst(Z1, Z2, ConstantVector::get(ShuffleMask8),
                                    "", CI);
    auto Z6 = new ShuffleVectorInst(Z3, Z4, ConstantVector::get(ShuffleMask8),
                                    "", CI);
    Constant *ShuffleMask16[16] = {
        ConstantInt::get(IntTy, 0),  ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2),  ConstantInt::get(IntTy, 3),
        ConstantInt::get(IntTy, 4),  ConstantInt::get(IntTy, 5),
        ConstantInt::get(IntTy, 6),  ConstantInt::get(IntTy, 7),
        ConstantInt::get(IntTy, 8),  ConstantInt::get(IntTy, 9),
        ConstantInt::get(IntTy, 10), ConstantInt::get(IntTy, 11),
        ConstantInt::get(IntTy, 12), ConstantInt::get(IntTy, 13),
        ConstantInt::get(IntTy, 14), ConstantInt::get(IntTy, 15)};
    // Combine our two float8's into one float16.
    return new ShuffleVectorInst(Z5, Z6, ConstantVector::get(ShuffleMask16), "",
                                 CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceClspvVloadaHalf2(Function &F) {

  // Replace __clspv_vloada_half2(uint Index, global uint* Ptr) with:
  //
  //    %u = load i32 %ptr
  //    %result = call <2 x float> Unpack2xHalf(u)
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto Index = CI->getOperand(0);
    auto Ptr = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    auto IndexedPtr = GetElementPtrInst::Create(IntTy, Ptr, Index, "", CI);
    auto Load = new LoadInst(IntTy, IndexedPtr, "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Get our final float2.
    return CallInst::Create(NewF, Load, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceClspvVloadaHalf4(Function &F) {

  // Replace __clspv_vloada_half4(uint Index, global uint2* Ptr) with:
  //
  //    %u2 = load <2 x i32> %ptr
  //    %u2xy = extractelement %u2, 0
  //    %u2zw = extractelement %u2, 1
  //    %fxy = call <2 x float> Unpack2xHalf(uint)
  //    %fzw = call <2 x float> Unpack2xHalf(uint)
  //    %result = shufflevector %fxy %fzw <4 x float> <0, 1, 2, 3>
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto Index = CI->getOperand(0);
    auto Ptr = CI->getOperand(1);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int2Ty = FixedVectorType::get(IntTy, 2);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewFType = FunctionType::get(Float2Ty, IntTy, false);

    auto IndexedPtr = GetElementPtrInst::Create(Int2Ty, Ptr, Index, "", CI);
    auto Load = new LoadInst(Int2Ty, IndexedPtr, "", CI);

    // Extract each element from the loaded int2.
    auto X =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 0), "", CI);
    auto Y =
        ExtractElementInst::Create(Load, ConstantInt::get(IntTy, 1), "", CI);

    // Our intrinsic to unpack a float2 from an int.
    auto SPIRVIntrinsic = clspv::UnpackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Get the lower (x & y) components of our final float4.
    auto Lo = CallInst::Create(NewF, X, "", CI);

    // Get the higher (z & w) components of our final float4.
    auto Hi = CallInst::Create(NewF, Y, "", CI);

    Constant *ShuffleMask[4] = {
        ConstantInt::get(IntTy, 0), ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2), ConstantInt::get(IntTy, 3)};

    // Combine our two float2's into one float4.
    return new ShuffleVectorInst(Lo, Hi, ConstantVector::get(ShuffleMask), "",
                                 CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreHalf(Function &F, int vec_size, bool aligned) {
  switch (vec_size) {
  case 0:
    return replaceVstoreHalf(F);
  case 2:
    return replaceVstoreHalf2(F);
  case 3:
    return aligned ? replaceVstoreaHalf3(F) : replaceVstoreHalf3(F);
  case 4:
    return replaceVstoreHalf4(F);
  case 8:
    return replaceVstoreHalf8(F);
  case 16:
    return replaceVstoreHalf16(F);
  default:
    llvm_unreachable("Unsupported vstore_half vector size");
    break;
  }
  return false;
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreHalf(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The value to store.
    auto Arg0 = CI->getOperand(0);

    // The index argument from vstore_half.
    auto Arg1 = CI->getOperand(1);

    // The pointer argument from vstore_half.
    auto Arg2 = CI->getOperand(2);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewFType = FunctionType::get(IntTy, Float2Ty, false);
    auto One = ConstantInt::get(IntTy, 1);

    // Our intrinsic to pack a float2 to an int.
    auto SPIRVIntrinsic = clspv::PackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Insert our value into a float2 so that we can pack it.
    auto TempVec = InsertElementInst::Create(
        UndefValue::get(Float2Ty), Arg0, ConstantInt::get(IntTy, 0), "", CI);

    // Pack the float2 -> half2 (in an int).
    auto X = CallInst::Create(NewF, TempVec, "", CI);

    bool supports_16bit_storage = true;
    switch (Arg2->getType()->getPointerAddressSpace()) {
    case clspv::AddressSpace::Global:
      supports_16bit_storage = clspv::Option::Supports16BitStorageClass(
          clspv::Option::StorageClass::kSSBO);
      break;
    case clspv::AddressSpace::Constant:
      if (clspv::Option::ConstantArgsInUniformBuffer())
        supports_16bit_storage = clspv::Option::Supports16BitStorageClass(
            clspv::Option::StorageClass::kUBO);
      else
        supports_16bit_storage = clspv::Option::Supports16BitStorageClass(
            clspv::Option::StorageClass::kSSBO);
      break;
    default:
      // Clspv will emit the Float16 capability if the half type is
      // encountered. That capability covers private and local addressspaces.
      break;
    }

    Value *V = nullptr;
    if (supports_16bit_storage) {
      auto ShortTy = Type::getInt16Ty(M.getContext());
      auto ShortPointerTy =
          PointerType::get(ShortTy, Arg2->getType()->getPointerAddressSpace());

      // Truncate our i32 to an i16.
      auto Trunc = CastInst::CreateTruncOrBitCast(X, ShortTy, "", CI);

      // Cast the half* pointer to short*.
      auto Cast = CastInst::CreatePointerCast(Arg2, ShortPointerTy, "", CI);

      // Index into the correct address of the casted pointer.
      auto Index = GetElementPtrInst::Create(ShortTy, Cast, Arg1, "", CI);

      // Store to the int* we casted to.
      V = new StoreInst(Trunc, Index, CI);
    } else {
      // We can only write to 32-bit aligned words.
      //
      // Assuming base is aligned to 32-bits, replace the equivalent of
      //   vstore_half(value, index, base)
      // with:
      //   uint32_t* target_ptr = (uint32_t*)(base) + index / 2;
      //   uint32_t write_to_upper_half = index & 1u;
      //   uint32_t shift = write_to_upper_half << 4;
      //
      //   // Pack the float value as a half number in bottom 16 bits
      //   // of an i32.
      //   uint32_t packed = spirv.pack.v2f16((float2)(value, undef));
      //
      //   uint32_t xor_value =   (*target_ptr & (0xffff << shift))
      //                        ^ ((packed & 0xffff) << shift)
      //   // We only need relaxed consistency, but OpenCL 1.2 only has
      //   // sequentially consistent atomics.
      //   // TODO(dneto): Use relaxed consistency.
      //   atomic_xor(target_ptr, xor_value)
      auto IntPointerTy =
          PointerType::get(IntTy, Arg2->getType()->getPointerAddressSpace());

      auto Four = ConstantInt::get(IntTy, 4);
      auto FFFF = ConstantInt::get(IntTy, 0xffff);

      auto IndexIsOdd =
          BinaryOperator::CreateAnd(Arg1, One, "index_is_odd_i32", CI);
      // Compute index / 2
      auto IndexIntoI32 =
          BinaryOperator::CreateLShr(Arg1, One, "index_into_i32", CI);
      auto BaseI32Ptr =
          CastInst::CreatePointerCast(Arg2, IntPointerTy, "base_i32_ptr", CI);
      auto OutPtr = GetElementPtrInst::Create(IntTy, BaseI32Ptr, IndexIntoI32,
                                              "base_i32_ptr", CI);
      auto CurrentValue = new LoadInst(IntTy, OutPtr, "current_value", CI);
      auto Shift = BinaryOperator::CreateShl(IndexIsOdd, Four, "shift", CI);
      auto MaskBitsToWrite =
          BinaryOperator::CreateShl(FFFF, Shift, "mask_bits_to_write", CI);
      auto MaskedCurrent = BinaryOperator::CreateAnd(
          MaskBitsToWrite, CurrentValue, "masked_current", CI);

      auto XLowerBits =
          BinaryOperator::CreateAnd(X, FFFF, "lower_bits_of_packed", CI);
      auto NewBitsToWrite =
          BinaryOperator::CreateShl(XLowerBits, Shift, "new_bits_to_write", CI);
      auto ValueToXor = BinaryOperator::CreateXor(MaskedCurrent, NewBitsToWrite,
                                                  "value_to_xor", CI);

      // Generate the call to atomi_xor.
      SmallVector<Type *, 5> ParamTypes;
      // The pointer type.
      ParamTypes.push_back(IntPointerTy);
      // The Types for memory scope, semantics, and value.
      ParamTypes.push_back(IntTy);
      ParamTypes.push_back(IntTy);
      ParamTypes.push_back(IntTy);
      auto NewFType = FunctionType::get(IntTy, ParamTypes, false);
      auto NewF = M.getOrInsertFunction("spirv.atomic_xor", NewFType);

      const auto ConstantScopeDevice =
          ConstantInt::get(IntTy, spv::ScopeDevice);
      // Assume the pointee is in OpenCL global (SPIR-V Uniform) or local
      // (SPIR-V Workgroup).
      const auto AddrSpaceSemanticsBits =
          IntPointerTy->getPointerAddressSpace() == 1
              ? spv::MemorySemanticsUniformMemoryMask
              : spv::MemorySemanticsWorkgroupMemoryMask;

      // We're using relaxed consistency here.
      const auto ConstantMemorySemantics =
          ConstantInt::get(IntTy, spv::MemorySemanticsUniformMemoryMask |
                                      AddrSpaceSemanticsBits);

      SmallVector<Value *, 5> Params{OutPtr, ConstantScopeDevice,
                                     ConstantMemorySemantics, ValueToXor};
      CallInst::Create(NewF, Params, "store_halfword_xor_trick", CI);

      // Return a Nop so the old Call is removed
      Function *donothing = Intrinsic::getDeclaration(&M, Intrinsic::donothing);
      V = CallInst::Create(donothing, {}, "", CI);
    }

    return V;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreHalf2(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The value to store.
    auto Arg0 = CI->getOperand(0);

    // The index argument from vstore_half.
    auto Arg1 = CI->getOperand(1);

    // The pointer argument from vstore_half.
    auto Arg2 = CI->getOperand(2);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(IntTy, Arg2->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(IntTy, Float2Ty, false);

    // Our intrinsic to pack a float2 to an int.
    auto SPIRVIntrinsic = clspv::PackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Turn the packed x & y into the final packing.
    auto X = CallInst::Create(NewF, Arg0, "", CI);

    // Cast the half* pointer to int*.
    auto Cast = CastInst::CreatePointerCast(Arg2, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Index = GetElementPtrInst::Create(IntTy, Cast, Arg1, "", CI);

    // Store to the int* we casted to.
    return new StoreInst(X, Index, CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreHalf3(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The value to store.
    auto Arg0 = CI->getOperand(0);

    // The index argument from vstore_half.
    auto Arg1 = CI->getOperand(1);

    // The pointer argument from vstore_half.
    auto Arg2 = CI->getOperand(2);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto ShortTy = Type::getInt16Ty(M.getContext());
    auto FloatTy = Type::getFloatTy(M.getContext());
    auto Float2Ty = FixedVectorType::get(FloatTy, 2);
    auto NewPointerTy =
        PointerType::get(ShortTy, Arg2->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(IntTy, Float2Ty, false);

    auto Int0 = ConstantInt::get(IntTy, 0);
    auto Int1 = ConstantInt::get(IntTy, 1);
    auto Int2 = ConstantInt::get(IntTy, 2);

    auto X0 = InsertElementInst::Create(
        UndefValue::get(Float2Ty),
        ExtractElementInst::Create(Arg0, Int0, "", CI), Int0, "", CI);
    auto X1 = InsertElementInst::Create(
        UndefValue::get(Float2Ty),
        ExtractElementInst::Create(Arg0, Int1, "", CI), Int0, "", CI);
    auto X2 = InsertElementInst::Create(
        UndefValue::get(Float2Ty),
        ExtractElementInst::Create(Arg0, Int2, "", CI), Int0, "", CI);

    // Our intrinsic to pack a float2 to an int.
    auto SPIRVIntrinsic = clspv::PackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Convert float2 into int and trunc to short to keep only the meaningful
    // part of it
    auto Y0 =
        CastInst::Create(Instruction::Trunc, CallInst::Create(NewF, X0, "", CI),
                         ShortTy, "", CI);
    auto Y1 =
        CastInst::Create(Instruction::Trunc, CallInst::Create(NewF, X1, "", CI),
                         ShortTy, "", CI);
    auto Y2 =
        CastInst::Create(Instruction::Trunc, CallInst::Create(NewF, X2, "", CI),
                         ShortTy, "", CI);

    // Cast the half* pointer to short*.
    auto Cast = CastInst::CreatePointerCast(Arg2, NewPointerTy, "", CI);

    auto Index0 = BinaryOperator::Create(
        Instruction::Add,
        BinaryOperator::Create(Instruction::Shl, Arg1, Int1, "", CI), Arg1, "",
        CI);
    auto GEP0 = GetElementPtrInst::Create(ShortTy, Cast, Index0, "", CI);
    new StoreInst(Y0, GEP0, CI);

    auto Index1 =
        BinaryOperator::Create(Instruction::Add, Index0, Int1, "", CI);
    auto GEP1 = GetElementPtrInst::Create(ShortTy, Cast, Index1, "", CI);
    new StoreInst(Y1, GEP1, CI);

    auto Index2 =
        BinaryOperator::Create(Instruction::Add, Index1, Int1, "", CI);
    auto GEP2 = GetElementPtrInst::Create(ShortTy, Cast, Index2, "", CI);
    return new StoreInst(Y2, GEP2, CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreaHalf3(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The value to store.
    auto Arg0 = CI->getOperand(0);

    // The index argument from vstore_half.
    auto Arg1 = CI->getOperand(1);

    // The pointer argument from vstore_half.
    auto Arg2 = CI->getOperand(2);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto ShortTy = Type::getInt16Ty(M.getContext());
    auto FloatTy = Type::getFloatTy(M.getContext());
    auto Float2Ty = FixedVectorType::get(FloatTy, 2);
    auto NewPointerTy =
        PointerType::get(ShortTy, Arg2->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(IntTy, Float2Ty, false);

    auto Int0 = ConstantInt::get(IntTy, 0);
    auto Int1 = ConstantInt::get(IntTy, 1);
    auto Int2 = ConstantInt::get(IntTy, 2);

    auto X0 = InsertElementInst::Create(
        UndefValue::get(Float2Ty),
        ExtractElementInst::Create(Arg0, Int0, "", CI), Int0, "", CI);
    auto X1 = InsertElementInst::Create(
        UndefValue::get(Float2Ty),
        ExtractElementInst::Create(Arg0, Int1, "", CI), Int0, "", CI);
    auto X2 = InsertElementInst::Create(
        UndefValue::get(Float2Ty),
        ExtractElementInst::Create(Arg0, Int2, "", CI), Int0, "", CI);

    // Our intrinsic to pack a float2 to an int.
    auto SPIRVIntrinsic = clspv::PackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Convert float2 into int and trunc to short to keep only the meaningful
    // part of it
    auto Y0 =
        CastInst::Create(Instruction::Trunc, CallInst::Create(NewF, X0, "", CI),
                         ShortTy, "", CI);
    auto Y1 =
        CastInst::Create(Instruction::Trunc, CallInst::Create(NewF, X1, "", CI),
                         ShortTy, "", CI);
    auto Y2 =
        CastInst::Create(Instruction::Trunc, CallInst::Create(NewF, X2, "", CI),
                         ShortTy, "", CI);

    // Cast the half* pointer to short*.
    auto Cast = CastInst::CreatePointerCast(Arg2, NewPointerTy, "", CI);

    auto Index0 = BinaryOperator::Create(Instruction::Shl, Arg1, Int2, "", CI);
    auto GEP0 = GetElementPtrInst::Create(ShortTy, Cast, Index0, "", CI);
    new StoreInst(Y0, GEP0, CI);

    auto Index1 =
        BinaryOperator::Create(Instruction::Add, Index0, Int1, "", CI);
    auto GEP1 = GetElementPtrInst::Create(ShortTy, Cast, Index1, "", CI);
    new StoreInst(Y1, GEP1, CI);

    auto Index2 =
        BinaryOperator::Create(Instruction::Add, Index1, Int1, "", CI);
    auto GEP2 = GetElementPtrInst::Create(ShortTy, Cast, Index2, "", CI);
    return new StoreInst(Y2, GEP2, CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreHalf4(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The value to store.
    auto Arg0 = CI->getOperand(0);

    // The index argument from vstore_half.
    auto Arg1 = CI->getOperand(1);

    // The pointer argument from vstore_half.
    auto Arg2 = CI->getOperand(2);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int2Ty = FixedVectorType::get(IntTy, 2);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(Int2Ty, Arg2->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(IntTy, Float2Ty, false);

    Constant *LoShuffleMask[2] = {ConstantInt::get(IntTy, 0),
                                  ConstantInt::get(IntTy, 1)};

    // Extract out the x & y components of our to store value.
    auto Lo = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(LoShuffleMask), "", CI);

    Constant *HiShuffleMask[2] = {ConstantInt::get(IntTy, 2),
                                  ConstantInt::get(IntTy, 3)};

    // Extract out the z & w components of our to store value.
    auto Hi = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(HiShuffleMask), "", CI);

    // Our intrinsic to pack a float2 to an int.
    auto SPIRVIntrinsic = clspv::PackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    // Turn the packed x & y into the final component of our int2.
    auto X = CallInst::Create(NewF, Lo, "", CI);

    // Turn the packed z & w into the final component of our int2.
    auto Y = CallInst::Create(NewF, Hi, "", CI);

    auto Combine = InsertElementInst::Create(
        UndefValue::get(Int2Ty), X, ConstantInt::get(IntTy, 0), "", CI);
    Combine = InsertElementInst::Create(Combine, Y, ConstantInt::get(IntTy, 1),
                                        "", CI);

    // Cast the half* pointer to int2*.
    auto Cast = CastInst::CreatePointerCast(Arg2, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Index = GetElementPtrInst::Create(Int2Ty, Cast, Arg1, "", CI);

    // Store to the int2* we casted to.
    return new StoreInst(Combine, Index, CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreHalf8(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The value to store.
    auto Arg0 = CI->getOperand(0);

    // The index argument from vstore_half.
    auto Arg1 = CI->getOperand(1);

    // The pointer argument from vstore_half.
    auto Arg2 = CI->getOperand(2);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int4Ty = FixedVectorType::get(IntTy, 4);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(Int4Ty, Arg2->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(IntTy, Float2Ty, false);

    Constant *ShuffleMask01[2] = {ConstantInt::get(IntTy, 0),
                                  ConstantInt::get(IntTy, 1)};
    auto X01 =
        new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                              ConstantVector::get(ShuffleMask01), "", CI);
    Constant *ShuffleMask23[2] = {ConstantInt::get(IntTy, 2),
                                  ConstantInt::get(IntTy, 3)};
    auto X23 =
        new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                              ConstantVector::get(ShuffleMask23), "", CI);
    Constant *ShuffleMask45[2] = {ConstantInt::get(IntTy, 4),
                                  ConstantInt::get(IntTy, 5)};
    auto X45 =
        new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                              ConstantVector::get(ShuffleMask45), "", CI);
    Constant *ShuffleMask67[2] = {ConstantInt::get(IntTy, 6),
                                  ConstantInt::get(IntTy, 7)};
    auto X67 =
        new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                              ConstantVector::get(ShuffleMask67), "", CI);

    // Our intrinsic to pack a float2 to an int.
    auto SPIRVIntrinsic = clspv::PackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    auto Y01 = CallInst::Create(NewF, X01, "", CI);
    auto Y23 = CallInst::Create(NewF, X23, "", CI);
    auto Y45 = CallInst::Create(NewF, X45, "", CI);
    auto Y67 = CallInst::Create(NewF, X67, "", CI);

    auto Combine = InsertElementInst::Create(
        UndefValue::get(Int4Ty), Y01, ConstantInt::get(IntTy, 0), "", CI);
    Combine = InsertElementInst::Create(Combine, Y23,
                                        ConstantInt::get(IntTy, 1), "", CI);
    Combine = InsertElementInst::Create(Combine, Y45,
                                        ConstantInt::get(IntTy, 2), "", CI);
    Combine = InsertElementInst::Create(Combine, Y67,
                                        ConstantInt::get(IntTy, 3), "", CI);

    // Cast the half* pointer to int4*.
    auto Cast = CastInst::CreatePointerCast(Arg2, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Index = GetElementPtrInst::Create(Int4Ty, Cast, Arg1, "", CI);

    // Store to the int4* we casted to.
    return new StoreInst(Combine, Index, CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceVstoreHalf16(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The value to store.
    auto Arg0 = CI->getOperand(0);

    // The index argument from vstore_half.
    auto Arg1 = CI->getOperand(1);

    // The pointer argument from vstore_half.
    auto Arg2 = CI->getOperand(2);

    auto IntTy = Type::getInt32Ty(M.getContext());
    auto Int4Ty = FixedVectorType::get(IntTy, 4);
    auto Float2Ty = FixedVectorType::get(Type::getFloatTy(M.getContext()), 2);
    auto NewPointerTy =
        PointerType::get(Int4Ty, Arg2->getType()->getPointerAddressSpace());
    auto NewFType = FunctionType::get(IntTy, Float2Ty, false);

    Constant *ShuffleMask0[2] = {ConstantInt::get(IntTy, 0),
                                 ConstantInt::get(IntTy, 1)};
    auto X0 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask0), "", CI);
    Constant *ShuffleMask1[2] = {ConstantInt::get(IntTy, 2),
                                 ConstantInt::get(IntTy, 3)};
    auto X1 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask1), "", CI);
    Constant *ShuffleMask2[2] = {ConstantInt::get(IntTy, 4),
                                 ConstantInt::get(IntTy, 5)};
    auto X2 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask2), "", CI);
    Constant *ShuffleMask3[2] = {ConstantInt::get(IntTy, 6),
                                 ConstantInt::get(IntTy, 7)};
    auto X3 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask3), "", CI);
    Constant *ShuffleMask4[2] = {ConstantInt::get(IntTy, 8),
                                 ConstantInt::get(IntTy, 9)};
    auto X4 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask4), "", CI);
    Constant *ShuffleMask5[2] = {ConstantInt::get(IntTy, 10),
                                 ConstantInt::get(IntTy, 11)};
    auto X5 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask5), "", CI);
    Constant *ShuffleMask6[2] = {ConstantInt::get(IntTy, 12),
                                 ConstantInt::get(IntTy, 13)};
    auto X6 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask6), "", CI);
    Constant *ShuffleMask7[2] = {ConstantInt::get(IntTy, 14),
                                 ConstantInt::get(IntTy, 15)};
    auto X7 = new ShuffleVectorInst(Arg0, UndefValue::get(Arg0->getType()),
                                    ConstantVector::get(ShuffleMask7), "", CI);

    // Our intrinsic to pack a float2 to an int.
    auto SPIRVIntrinsic = clspv::PackFunction();

    auto NewF = M.getOrInsertFunction(SPIRVIntrinsic, NewFType);

    auto Y0 = CallInst::Create(NewF, X0, "", CI);
    auto Y1 = CallInst::Create(NewF, X1, "", CI);
    auto Y2 = CallInst::Create(NewF, X2, "", CI);
    auto Y3 = CallInst::Create(NewF, X3, "", CI);
    auto Y4 = CallInst::Create(NewF, X4, "", CI);
    auto Y5 = CallInst::Create(NewF, X5, "", CI);
    auto Y6 = CallInst::Create(NewF, X6, "", CI);
    auto Y7 = CallInst::Create(NewF, X7, "", CI);

    auto Combine1 = InsertElementInst::Create(
        UndefValue::get(Int4Ty), Y0, ConstantInt::get(IntTy, 0), "", CI);
    Combine1 = InsertElementInst::Create(Combine1, Y1,
                                         ConstantInt::get(IntTy, 1), "", CI);
    Combine1 = InsertElementInst::Create(Combine1, Y2,
                                         ConstantInt::get(IntTy, 2), "", CI);
    Combine1 = InsertElementInst::Create(Combine1, Y3,
                                         ConstantInt::get(IntTy, 3), "", CI);

    auto Combine2 = InsertElementInst::Create(
        UndefValue::get(Int4Ty), Y4, ConstantInt::get(IntTy, 0), "", CI);
    Combine2 = InsertElementInst::Create(Combine2, Y5,
                                         ConstantInt::get(IntTy, 1), "", CI);
    Combine2 = InsertElementInst::Create(Combine2, Y6,
                                         ConstantInt::get(IntTy, 2), "", CI);
    Combine2 = InsertElementInst::Create(Combine2, Y7,
                                         ConstantInt::get(IntTy, 3), "", CI);

    // Cast the half* pointer to int4*.
    auto Cast = CastInst::CreatePointerCast(Arg2, NewPointerTy, "", CI);

    // Index into the correct address of the casted pointer.
    auto Arg1x2 = BinaryOperator::Create(Instruction::Shl, Arg1,
                                         ConstantInt::get(IntTy, 1), "", CI);
    auto Index1 = GetElementPtrInst::Create(Int4Ty, Cast, Arg1x2, "", CI);

    // Store to the int4* we casted to.
    new StoreInst(Combine1, Index1, CI);

    // Index into the correct address of the casted pointer.
    auto Arg1Plus1 = BinaryOperator::Create(Instruction::Add, Arg1x2,
                                            ConstantInt::get(IntTy, 1), "", CI);
    auto Index2 = GetElementPtrInst::Create(Int4Ty, Cast, Arg1Plus1, "", CI);

    // Store to the int4* we casted to.
    return new StoreInst(Combine2, Index2, CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceHalfReadImage(Function &F) {
  // convert half to float
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    SmallVector<Type *, 3> types;
    SmallVector<Value *, 3> args;
    for (size_t i = 0; i < CI->arg_size(); ++i) {
      types.push_back(CI->getArgOperand(i)->getType());
      args.push_back(CI->getArgOperand(i));
    }

    auto NewFType =
        FunctionType::get(FixedVectorType::get(Type::getFloatTy(M.getContext()),
                                               cast<VectorType>(CI->getType())
                                                   ->getElementCount()
                                                   .getKnownMinValue()),
                          types, false);

    std::string NewFName =
        Builtins::GetMangledFunctionName("read_imagef", NewFType);

    auto NewF = M.getOrInsertFunction(NewFName, NewFType);

    auto NewCI = CallInst::Create(NewF, args, "", CI);

    // Convert to the half type.
    return CastInst::CreateFPCast(NewCI, CI->getType(), "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceHalfWriteImage(Function &F) {
  // convert half to float
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    SmallVector<Type *, 3> types(3);
    SmallVector<Value *, 3> args(3);

    // Image
    types[0] = CI->getArgOperand(0)->getType();
    args[0] = CI->getArgOperand(0);

    // Coord
    types[1] = CI->getArgOperand(1)->getType();
    args[1] = CI->getArgOperand(1);

    // Data
    types[2] =
        FixedVectorType::get(Type::getFloatTy(M.getContext()),
                             cast<VectorType>(CI->getArgOperand(2)->getType())
                                 ->getElementCount()
                                 .getKnownMinValue());

    auto NewFType =
        FunctionType::get(Type::getVoidTy(M.getContext()), types, false);

    std::string NewFName =
        Builtins::GetMangledFunctionName("write_imagef", NewFType);

    auto NewF = M.getOrInsertFunction(NewFName, NewFType);

    // Convert data to the float type.
    auto Cast = CastInst::CreateFPCast(CI->getArgOperand(2), types[2], "", CI);
    args[2] = Cast;

    return CallInst::Create(NewF, args, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceSampledReadImageWithIntCoords(
    Function &F) {
  // convert read_image with int coords to float coords
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // The image.
    auto Arg0 = CI->getOperand(0);

    // The sampler.
    auto Arg1 = CI->getOperand(1);

    // The coordinate (integer type that we can't handle).
    auto Arg2 = CI->getOperand(2);

    uint32_t dim = clspv::ImageNumDimensions(Arg0->getType());
    uint32_t components =
        dim + (clspv::IsArrayImageType(Arg0->getType()) ? 1 : 0);
    Type *float_ty = nullptr;
    if (components == 1) {
      float_ty = Type::getFloatTy(M.getContext());
    } else {
      float_ty = FixedVectorType::get(Type::getFloatTy(M.getContext()),
                                      cast<VectorType>(Arg2->getType())
                                          ->getElementCount()
                                          .getKnownMinValue());
    }

    auto NewFType = FunctionType::get(
        CI->getType(), {Arg0->getType(), Arg1->getType(), float_ty}, false);

    std::string NewFName = F.getName().str();
    NewFName[NewFName.length() - 1] = 'f';

    auto NewF = M.getOrInsertFunction(NewFName, NewFType);

    auto Cast = CastInst::Create(Instruction::SIToFP, Arg2, float_ty, "", CI);

    return CallInst::Create(NewF, {Arg0, Arg1, Cast}, "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAtomics(Function &F, spv::Op Op) {
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto IntTy = Type::getInt32Ty(F.getContext());

    // We need to map the OpenCL constants to the SPIR-V equivalents.
    const auto ConstantScopeDevice = ConstantInt::get(IntTy, spv::ScopeDevice);
    const auto ConstantMemorySemantics = ConstantInt::get(
        IntTy, spv::MemorySemanticsUniformMemoryMask |
                   spv::MemorySemanticsSequentiallyConsistentMask);

    SmallVector<Value *, 5> Params;

    // The pointer.
    Params.push_back(CI->getArgOperand(0));

    // The memory scope.
    Params.push_back(ConstantScopeDevice);

    // The memory semantics.
    Params.push_back(ConstantMemorySemantics);

    if (2 < CI->arg_size()) {
      // The unequal memory semantics.
      Params.push_back(ConstantMemorySemantics);

      // The value.
      Params.push_back(CI->getArgOperand(2));

      // The comparator.
      Params.push_back(CI->getArgOperand(1));
    } else if (1 < CI->arg_size()) {
      // The value.
      Params.push_back(CI->getArgOperand(1));
    }

    return clspv::InsertSPIRVOp(CI, Op, {}, CI->getType(), Params);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAtomics(Function &F,
                                              llvm::AtomicRMWInst::BinOp Op) {
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto align = F.getParent()->getDataLayout().getABITypeAlign(
        CI->getArgOperand(1)->getType());
    return new AtomicRMWInst(Op, CI->getArgOperand(0), CI->getArgOperand(1),
                             align, AtomicOrdering::SequentiallyConsistent,
                             SyncScope::System, CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceCross(Function &F) {
  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    auto IntTy = Type::getInt32Ty(M.getContext());
    auto FloatTy = Type::getFloatTy(M.getContext());

    Constant *DownShuffleMask[3] = {ConstantInt::get(IntTy, 0),
                                    ConstantInt::get(IntTy, 1),
                                    ConstantInt::get(IntTy, 2)};

    Constant *UpShuffleMask[4] = {
        ConstantInt::get(IntTy, 0), ConstantInt::get(IntTy, 1),
        ConstantInt::get(IntTy, 2), ConstantInt::get(IntTy, 3)};

    Constant *FloatVec[3] = {ConstantFP::get(FloatTy, 0.0f),
                             UndefValue::get(FloatTy),
                             UndefValue::get(FloatTy)};

    auto Vec4Ty = CI->getArgOperand(0)->getType();
    auto Arg0 =
        new ShuffleVectorInst(CI->getArgOperand(0), UndefValue::get(Vec4Ty),
                              ConstantVector::get(DownShuffleMask), "", CI);
    auto Arg1 =
        new ShuffleVectorInst(CI->getArgOperand(1), UndefValue::get(Vec4Ty),
                              ConstantVector::get(DownShuffleMask), "", CI);
    auto Vec3Ty = Arg0->getType();

    auto NewFType = FunctionType::get(Vec3Ty, {Vec3Ty, Vec3Ty}, false);
    auto NewFName = Builtins::GetMangledFunctionName("cross", NewFType);

    auto Cross3Func = M.getOrInsertFunction(NewFName, NewFType);

    auto DownResult = CallInst::Create(Cross3Func, {Arg0, Arg1}, "", CI);

    return new ShuffleVectorInst(DownResult, ConstantVector::get(FloatVec),
                                 ConstantVector::get(UpShuffleMask), "", CI);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceFract(Function &F, int vec_size) {
  // OpenCL's   float result = fract(float x, float* ptr)
  //
  // In the LLVM domain:
  //
  //    %floor_result = call spir_func float @floor(float %x)
  //    store float %floor_result, float * %ptr
  //    %fract_intermediate = call spir_func float @clspv.fract(float %x)
  //    %result = call spir_func float
  //        @fmin(float %fract_intermediate, float 0x1.fffffep-1f)
  //
  // Becomes in the SPIR-V domain, where translations of floor, fmin,
  // and clspv.fract occur in the SPIR-V generator pass:
  //
  //    %glsl_ext = OpExtInstImport "GLSL.std.450"
  //    %just_under_1 = OpConstant %float 0x1.fffffep-1f
  //    ...
  //    %floor_result = OpExtInst %float %glsl_ext Floor %x
  //    OpStore %ptr %floor_result
  //    %fract_intermediate = OpExtInst %float %glsl_ext Fract %x
  //    %fract_result = OpExtInst %float
  //       %glsl_ext Nmin %fract_intermediate %just_under_1

  using std::string;

  // Mapping from the fract builtin to the floor, fmin, and clspv.fract builtins
  // we need.  The clspv.fract builtin is the same as GLSL.std.450 Fract.

  Module &M = *F.getParent();
  return replaceCallsWithValue(F, [&](CallInst *CI) {
    // This is either float or a float vector.  All the float-like
    // types are this type.
    auto result_ty = F.getReturnType();

    std::string fmin_name = Builtins::GetMangledFunctionName("fmin", result_ty);
    Function *fmin_fn = M.getFunction(fmin_name);
    if (!fmin_fn) {
      // Make the fmin function.
      FunctionType *fn_ty =
          FunctionType::get(result_ty, {result_ty, result_ty}, false);
      fmin_fn =
          cast<Function>(M.getOrInsertFunction(fmin_name, fn_ty).getCallee());
      fmin_fn->addFnAttr(Attribute::ReadNone);
      fmin_fn->setCallingConv(CallingConv::SPIR_FUNC);
    }

    std::string floor_name =
        Builtins::GetMangledFunctionName("floor", result_ty);
    Function *floor_fn = M.getFunction(floor_name);
    if (!floor_fn) {
      // Make the floor function.
      FunctionType *fn_ty = FunctionType::get(result_ty, {result_ty}, false);
      floor_fn =
          cast<Function>(M.getOrInsertFunction(floor_name, fn_ty).getCallee());
      floor_fn->addFnAttr(Attribute::ReadNone);
      floor_fn->setCallingConv(CallingConv::SPIR_FUNC);
    }

    std::string clspv_fract_name =
        Builtins::GetMangledFunctionName("clspv.fract", result_ty);
    Function *clspv_fract_fn = M.getFunction(clspv_fract_name);
    if (!clspv_fract_fn) {
      // Make the clspv_fract function.
      FunctionType *fn_ty = FunctionType::get(result_ty, {result_ty}, false);
      clspv_fract_fn = cast<Function>(
          M.getOrInsertFunction(clspv_fract_name, fn_ty).getCallee());
      clspv_fract_fn->addFnAttr(Attribute::ReadNone);
      clspv_fract_fn->setCallingConv(CallingConv::SPIR_FUNC);
    }

    // Number of significant significand bits, whether represented or not.
    unsigned num_significand_bits;
    switch (result_ty->getScalarType()->getTypeID()) {
    case Type::HalfTyID:
      num_significand_bits = 11;
      break;
    case Type::FloatTyID:
      num_significand_bits = 24;
      break;
    case Type::DoubleTyID:
      num_significand_bits = 53;
      break;
    default:
      llvm_unreachable("Unhandled float type when processing fract builtin");
      break;
    }
    // Beware that the disassembler displays this value as
    //   OpConstant %float 1
    // which is not quite right.
    const double kJustUnderOneScalar =
        ldexp(double((1 << num_significand_bits) - 1), -num_significand_bits);

    Constant *just_under_one =
        ConstantFP::get(result_ty->getScalarType(), kJustUnderOneScalar);
    if (result_ty->isVectorTy()) {
      just_under_one = ConstantVector::getSplat(
          cast<VectorType>(result_ty)->getElementCount(), just_under_one);
    }

    IRBuilder<> Builder(CI);

    auto arg = CI->getArgOperand(0);
    auto ptr = CI->getArgOperand(1);

    // Compute floor result and store it.
    auto floor = Builder.CreateCall(floor_fn, {arg});
    Builder.CreateStore(floor, ptr);

    auto fract_intermediate = Builder.CreateCall(clspv_fract_fn, arg);
    auto fract_result =
        Builder.CreateCall(fmin_fn, {fract_intermediate, just_under_one});

    return fract_result;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceHadd(Function &F, bool is_signed,
                                           Instruction::BinaryOps join_opcode) {
  return replaceCallsWithValue(F, [is_signed, join_opcode](CallInst *Call) {
    // a_shr = a >> 1
    // b_shr = b >> 1
    // add1 = a_shr + b_shr
    // join = a |join_opcode| b
    // and = join & 1
    // add = add1 + and
    const auto a = Call->getArgOperand(0);
    const auto b = Call->getArgOperand(1);
    IRBuilder<> builder(Call);
    Value *a_shift, *b_shift;
    if (is_signed) {
      a_shift = builder.CreateAShr(a, 1);
      b_shift = builder.CreateAShr(b, 1);
    } else {
      a_shift = builder.CreateLShr(a, 1);
      b_shift = builder.CreateLShr(b, 1);
    }
    auto add = builder.CreateAdd(a_shift, b_shift);
    auto join = BinaryOperator::Create(join_opcode, a, b, "", Call);
    auto constant_one = ConstantInt::get(a->getType(), 1);
    auto and_bit = builder.CreateAnd(join, constant_one);
    return builder.CreateAdd(add, and_bit);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAddSubSat(Function &F, bool is_signed,
                                                bool is_add) {
  return replaceCallsWithValue(F, [&F, this, is_signed,
                                   is_add](CallInst *Call) {
    auto intrinsic_type =
        is_signed ? (is_add ? Intrinsic::sadd_sat : Intrinsic::ssub_sat)
                  : (is_add ? Intrinsic::uadd_sat : Intrinsic::usub_sat);
    auto a = Call->getArgOperand(0);
    auto b = Call->getArgOperand(1);
    auto intrinsic = Intrinsic::getDeclaration(F.getParent(), intrinsic_type,
                                               Call->getType());
    return CallInst::Create(intrinsic->getFunctionType(), intrinsic, {a, b}, "",
                            Call);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAtomicLoad(Function &F) {
  return replaceCallsWithValue(F, [](CallInst *Call) {
    auto pointer = Call->getArgOperand(0);
    // Clang emits an address space cast to the generic address space. Skip the
    // cast and use the input directly.
    if (auto cast = dyn_cast<AddrSpaceCastOperator>(pointer)) {
      pointer = cast->getPointerOperand();
    }
    Value *order_arg = Call->arg_size() > 1 ? Call->getArgOperand(1) : nullptr;
    Value *scope_arg = Call->arg_size() > 2 ? Call->getArgOperand(2) : nullptr;
    bool is_global = pointer->getType()->getPointerAddressSpace() ==
                     clspv::AddressSpace::Global;
    auto order = MemoryOrderSemantics(order_arg, is_global, Call,
                                      spv::MemorySemanticsAcquireMask);
    auto scope = MemoryScope(scope_arg, is_global, Call);
    return InsertSPIRVOp(Call, spv::OpAtomicLoad, {Attribute::Convergent},
                         Call->getType(), {pointer, scope, order});
  });
}

bool ReplaceOpenCLBuiltinPass::replaceExplicitAtomics(
    Function &F, spv::Op Op, spv::MemorySemanticsMask semantics) {
  return replaceCallsWithValue(F, [Op, semantics](CallInst *Call) {
    auto pointer = Call->getArgOperand(0);
    // Clang emits an address space cast to the generic address space. Skip the
    // cast and use the input directly.
    if (auto cast = dyn_cast<AddrSpaceCastOperator>(pointer)) {
      pointer = cast->getPointerOperand();
    }
    Value *value = Call->getArgOperand(1);
    Value *order_arg = Call->arg_size() > 2 ? Call->getArgOperand(2) : nullptr;
    Value *scope_arg = Call->arg_size() > 3 ? Call->getArgOperand(3) : nullptr;
    bool is_global = pointer->getType()->getPointerAddressSpace() ==
                     clspv::AddressSpace::Global;
    auto scope = MemoryScope(scope_arg, is_global, Call);
    auto order = MemoryOrderSemantics(order_arg, is_global, Call, semantics);
    return InsertSPIRVOp(Call, Op, {Attribute::Convergent}, Call->getType(),
                         {pointer, scope, order, value});
  });
}

bool ReplaceOpenCLBuiltinPass::replaceAtomicCompareExchange(Function &F) {
  return replaceCallsWithValue(F, [](CallInst *Call) {
    auto pointer = Call->getArgOperand(0);
    // Clang emits an address space cast to the generic address space. Skip the
    // cast and use the input directly.
    if (auto cast = dyn_cast<AddrSpaceCastOperator>(pointer)) {
      pointer = cast->getPointerOperand();
    }
    auto expected = Call->getArgOperand(1);
    if (auto cast = dyn_cast<AddrSpaceCastOperator>(expected)) {
      expected = cast->getPointerOperand();
    }
    auto value = Call->getArgOperand(2);
    bool is_global = pointer->getType()->getPointerAddressSpace() ==
                     clspv::AddressSpace::Global;
    Value *success_arg =
        Call->arg_size() > 3 ? Call->getArgOperand(3) : nullptr;
    Value *failure_arg =
        Call->arg_size() > 4 ? Call->getArgOperand(4) : nullptr;
    Value *scope_arg = Call->arg_size() > 5 ? Call->getArgOperand(5) : nullptr;
    auto scope = MemoryScope(scope_arg, is_global, Call);
    auto success = MemoryOrderSemantics(success_arg, is_global, Call,
                                        spv::MemorySemanticsAcquireReleaseMask);
    auto failure = MemoryOrderSemantics(failure_arg, is_global, Call,
                                        spv::MemorySemanticsAcquireMask);

    // If the value pointed to by |expected| equals the value pointed to by
    // |pointer|, |value| is written into |pointer|, otherwise the value in
    // |pointer| is written into |expected|. In order to avoid extra stores,
    // the basic block with the original atomic is split and the store is
    // performed in the |then| block. The condition is the inversion of the
    // comparison result.
    IRBuilder<> builder(Call);
    auto load = builder.CreateLoad(value->getType(), expected);
    auto cmp_xchg = InsertSPIRVOp(
        Call, spv::OpAtomicCompareExchange, {Attribute::Convergent},
        value->getType(), {pointer, scope, success, failure, value, load});
    auto cmp = builder.CreateICmpEQ(cmp_xchg, load);
    auto not_cmp = builder.CreateNot(cmp);
    auto then_branch = SplitBlockAndInsertIfThen(not_cmp, Call, false);
    builder.SetInsertPoint(then_branch);
    builder.CreateStore(cmp_xchg, expected);
    return cmp;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceCountZeroes(Function &F, bool leading) {
  if (!isa<IntegerType>(F.getReturnType()->getScalarType()))
    return false;

  auto bitwidth = F.getReturnType()->getScalarSizeInBits();
  if (bitwidth > 64)
    return false;

  return replaceCallsWithValue(F, [&F, leading](CallInst *Call) {
    Function *intrinsic = Intrinsic::getDeclaration(
        F.getParent(), leading ? Intrinsic::ctlz : Intrinsic::cttz,
        Call->getType());
    const auto c_false = ConstantInt::getFalse(Call->getContext());
    return CallInst::Create(intrinsic->getFunctionType(), intrinsic,
                            {Call->getArgOperand(0), c_false}, "", Call);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceMadSat(Function &F, bool is_signed) {
  return replaceCallsWithValue(F, [&F, is_signed, this](CallInst *Call) {
    const auto ty = Call->getType();
    const auto a = Call->getArgOperand(0);
    const auto b = Call->getArgOperand(1);
    const auto c = Call->getArgOperand(2);
    IRBuilder<> builder(Call);
    if (is_signed) {
      unsigned bitwidth = Call->getType()->getScalarSizeInBits();
      if (bitwidth < 32) {
        // mul = sext(a) * sext(b)
        // add = mul + sext(c)
        // res = clamp(add, MIN, MAX)
        unsigned extended_width = bitwidth << 1;
        if (clspv::Option::HackClampWidth() && extended_width < 32) {
          extended_width = 32;
        }
        Type *extended_ty = IntegerType::get(F.getContext(), extended_width);
        if (auto vec_ty = dyn_cast<VectorType>(ty)) {
          extended_ty = VectorType::get(extended_ty, vec_ty->getElementCount());
        }
        auto a_sext = builder.CreateSExt(a, extended_ty);
        auto b_sext = builder.CreateSExt(b, extended_ty);
        auto c_sext = builder.CreateSExt(c, extended_ty);
        // Extended the size so no overflows occur.
        auto mul = builder.CreateMul(a_sext, b_sext, "", true, true);
        auto add = builder.CreateAdd(mul, c_sext, "", true, true);
        auto func_ty = FunctionType::get(
            extended_ty, {extended_ty, extended_ty, extended_ty}, false);
        // Don't use function type because we need signed parameters.
        std::string clamp_name = Builtins::GetMangledFunctionName("clamp");
        // The clamp values are the signed min and max of the original bitwidth
        // sign extended to the extended bitwidth.
        Constant *min = ConstantInt::get(
            Call->getContext(),
            APInt::getSignedMinValue(bitwidth).sext(extended_width));
        Constant *max = ConstantInt::get(
            Call->getContext(),
            APInt::getSignedMaxValue(bitwidth).sext(extended_width));
        if (auto vec_ty = dyn_cast<VectorType>(ty)) {
          min = ConstantVector::getSplat(vec_ty->getElementCount(), min);
          max = ConstantVector::getSplat(vec_ty->getElementCount(), max);
          unsigned vec_width = vec_ty->getElementCount().getKnownMinValue();
          if (extended_width == 32)
            clamp_name += "Dv" + std::to_string(vec_width) + "_iS_S_";
          else
            clamp_name += "Dv" + std::to_string(vec_width) + "_sS_S_";
        } else {
          if (extended_width == 32)
            clamp_name += "iii";
          else
            clamp_name += "sss";
        }
        auto callee = F.getParent()->getOrInsertFunction(clamp_name, func_ty);
        auto clamp = builder.CreateCall(callee, {add, min, max});
        return builder.CreateTrunc(clamp, ty);
      } else {
        // Compute
        // {hi, lo} = smul_extended(a, b)
        // add = lo + c
        auto mul_ext = InsertOpMulExtended(Call, a, b, true);

        auto mul_lo = builder.CreateExtractValue(mul_ext, {0});
        auto mul_hi = builder.CreateExtractValue(mul_ext, {1});
        auto add = builder.CreateAdd(mul_lo, c);

        // Constants for use in the calculation.
        Constant *min = ConstantInt::get(Call->getContext(),
                                         APInt::getSignedMinValue(bitwidth));
        Constant *max = ConstantInt::get(Call->getContext(),
                                         APInt::getSignedMaxValue(bitwidth));
        Constant *max_plus_1 = ConstantInt::get(
            Call->getContext(),
            APInt::getSignedMaxValue(bitwidth) + APInt(bitwidth, 1));
        if (auto vec_ty = dyn_cast<VectorType>(ty)) {
          min = ConstantVector::getSplat(vec_ty->getElementCount(), min);
          max = ConstantVector::getSplat(vec_ty->getElementCount(), max);
          max_plus_1 =
              ConstantVector::getSplat(vec_ty->getElementCount(), max_plus_1);
        }

        auto a_xor_b = builder.CreateXor(a, b);
        auto same_sign =
            builder.CreateICmpSGT(a_xor_b, Constant::getAllOnesValue(ty));
        auto different_sign = builder.CreateNot(same_sign);
        auto hi_eq_0 = builder.CreateICmpEQ(mul_hi, Constant::getNullValue(ty));
        auto hi_ne_0 = builder.CreateNot(hi_eq_0);
        auto lo_ge_max = builder.CreateICmpUGE(mul_lo, max);
        auto c_gt_0 = builder.CreateICmpSGT(c, Constant::getNullValue(ty));
        auto c_lt_0 = builder.CreateICmpSLT(c, Constant::getNullValue(ty));
        auto add_gt_max = builder.CreateICmpUGT(add, max);
        auto hi_eq_m1 =
            builder.CreateICmpEQ(mul_hi, Constant::getAllOnesValue(ty));
        auto hi_ne_m1 = builder.CreateNot(hi_eq_m1);
        auto lo_le_max_plus_1 = builder.CreateICmpULE(mul_lo, max_plus_1);
        auto max_sub_lo = builder.CreateSub(max, mul_lo);
        auto c_lt_max_sub_lo = builder.CreateICmpULT(c, max_sub_lo);

        // Equivalent to:
        // if (((x < 0) == (y < 0)) && mul_hi != 0)
        //   return MAX
        // if (mul_hi == 0 && mul_lo >= MAX && (z > 0 || add > MAX))
        //   return MAX
        // if (((x < 0) != (y < 0)) && mul_hi != -1)
        //   return MIN
        // if (hi == -1 && mul_lo <= (MAX + 1) && (z < 0 || z < (MAX - mul_lo))
        //   return MIN
        // return add
        auto max_clamp_1 = builder.CreateAnd(same_sign, hi_ne_0);
        auto max_clamp_2 = builder.CreateOr(c_gt_0, add_gt_max);
        auto tmp = builder.CreateAnd(hi_eq_0, lo_ge_max);
        max_clamp_2 = builder.CreateAnd(tmp, max_clamp_2);
        auto max_clamp = builder.CreateOr(max_clamp_1, max_clamp_2);
        auto min_clamp_1 = builder.CreateAnd(different_sign, hi_ne_m1);
        auto min_clamp_2 = builder.CreateOr(c_lt_0, c_lt_max_sub_lo);
        tmp = builder.CreateAnd(hi_eq_m1, lo_le_max_plus_1);
        min_clamp_2 = builder.CreateAnd(tmp, min_clamp_2);
        auto min_clamp = builder.CreateOr(min_clamp_1, min_clamp_2);
        auto sel = builder.CreateSelect(min_clamp, min, add);
        return builder.CreateSelect(max_clamp, max, sel);
      }
    } else {
      // {lo, hi} = mul_extended(a, b)
      // {add, carry} = add_carry(lo, c)
      // cmp = (mul_hi | carry) == 0
      // mad_sat = cmp ? add : MAX
      auto struct_ty = GetPairStruct(ty);
      auto mul_ext = InsertOpMulExtended(Call, a, b, false);
      auto mul_lo = builder.CreateExtractValue(mul_ext, {0});
      auto mul_hi = builder.CreateExtractValue(mul_ext, {1});
      auto add_carry =
          InsertSPIRVOp(Call, spv::OpIAddCarry, {Attribute::ReadNone},
                        struct_ty, {mul_lo, c});
      auto add = builder.CreateExtractValue(add_carry, {0});
      auto carry = builder.CreateExtractValue(add_carry, {1});
      auto or_value = builder.CreateOr(mul_hi, carry);
      auto cmp = builder.CreateICmpEQ(or_value, Constant::getNullValue(ty));
      return builder.CreateSelect(cmp, add, Constant::getAllOnesValue(ty));
    }
  });
}

bool ReplaceOpenCLBuiltinPass::replaceOrdered(Function &F, bool is_ordered) {
  if (!isa<IntegerType>(F.getReturnType()->getScalarType()))
    return false;

  if (F.getFunctionType()->getNumParams() != 2)
    return false;

  if (F.getFunctionType()->getParamType(0) !=
      F.getFunctionType()->getParamType(1)) {
    return false;
  }

  switch (F.getFunctionType()->getParamType(0)->getScalarType()->getTypeID()) {
  case Type::FloatTyID:
  case Type::HalfTyID:
  case Type::DoubleTyID:
    break;
  default:
    return false;
  }

  // Scalar versions all return an int, while vector versions return a vector
  // of an equally sized integer types (e.g. short, int or long).
  if (isa<VectorType>(F.getReturnType())) {
    if (F.getReturnType()->getScalarSizeInBits() !=
        F.getFunctionType()->getParamType(0)->getScalarSizeInBits()) {
      return false;
    }
  } else {
    if (F.getReturnType()->getScalarSizeInBits() != 32)
      return false;
  }

  return replaceCallsWithValue(F, [is_ordered](CallInst *Call) {
    // Replace with a floating point [un]ordered comparison followed by an
    // extension.
    auto x = Call->getArgOperand(0);
    auto y = Call->getArgOperand(1);
    IRBuilder<> builder(Call);
    Value *tmp = nullptr;
    if (is_ordered) {
      // This leads to a slight inefficiency in the SPIR-V that is easy for
      // drivers to optimize where the SPIR-V for the comparison and the
      // extension could be fused to drop the inversion of the OpIsNan.
      tmp = builder.CreateFCmpORD(x, y);
    } else {
      tmp = builder.CreateFCmpUNO(x, y);
    }
    // OpenCL CTS requires that vector versions use sign extension, but scalar
    // versions use zero extension.
    if (isa<VectorType>(Call->getType()))
      return builder.CreateSExt(tmp, Call->getType());
    return builder.CreateZExt(tmp, Call->getType());
  });
}

bool ReplaceOpenCLBuiltinPass::replaceIsNormal(Function &F) {
  return replaceCallsWithValue(F, [](CallInst *Call) {
    auto ty = Call->getType();
    auto x = Call->getArgOperand(0);
    unsigned width = x->getType()->getScalarSizeInBits();
    Type *int_ty = IntegerType::get(Call->getContext(), width);
    uint64_t abs_mask = 0x7fffffff;
    uint64_t exp_mask = 0x7f800000;
    uint64_t min_mask = 0x00800000;
    if (width == 16) {
      abs_mask = 0x7fff;
      exp_mask = 0x7c00;
      min_mask = 0x0400;
    } else if (width == 64) {
      abs_mask = 0x7fffffffffffffff;
      exp_mask = 0x7ff0000000000000;
      min_mask = 0x0010000000000000;
    }
    Constant *abs_const = ConstantInt::get(int_ty, APInt(width, abs_mask));
    Constant *exp_const = ConstantInt::get(int_ty, APInt(width, exp_mask));
    Constant *min_const = ConstantInt::get(int_ty, APInt(width, min_mask));
    if (auto vec_ty = dyn_cast<VectorType>(ty)) {
      int_ty = VectorType::get(int_ty, vec_ty->getElementCount());
      abs_const =
          ConstantVector::getSplat(vec_ty->getElementCount(), abs_const);
      exp_const =
          ConstantVector::getSplat(vec_ty->getElementCount(), exp_const);
      min_const =
          ConstantVector::getSplat(vec_ty->getElementCount(), min_const);
    }
    // Drop the sign bit and then check that the number is between
    // (exclusive) the min and max exponent values for the bit width.
    IRBuilder<> builder(Call);
    auto bitcast = builder.CreateBitCast(x, int_ty);
    auto abs = builder.CreateAnd(bitcast, abs_const);
    auto lt = builder.CreateICmpULT(abs, exp_const);
    auto ge = builder.CreateICmpUGE(abs, min_const);
    auto tmp = builder.CreateAnd(lt, ge);
    // OpenCL CTS requires that vector versions use sign extension, but scalar
    // versions use zero extension.
    if (isa<VectorType>(ty))
      return builder.CreateSExt(tmp, ty);
    return builder.CreateZExt(tmp, ty);
  });
}

bool ReplaceOpenCLBuiltinPass::replaceFDim(Function &F) {
  return replaceCallsWithValue(F, [](CallInst *Call) {
    const auto x = Call->getArgOperand(0);
    const auto y = Call->getArgOperand(1);
    IRBuilder<> builder(Call);
    auto sub = builder.CreateFSub(x, y);
    auto cmp = builder.CreateFCmpUGT(x, y);
    return builder.CreateSelect(cmp, sub,
                                Constant::getNullValue(Call->getType()));
  });
}

bool ReplaceOpenCLBuiltinPass::replaceRound(Function &F) {
  return replaceCallsWithValue(F, [&F](CallInst *Call) {
    const auto x = Call->getArgOperand(0);
    const double c_halfway = 0.5;
    auto halfway = ConstantFP::get(Call->getType(), c_halfway);

    const auto clspv_fract_name =
        Builtins::GetMangledFunctionName("clspv.fract", F.getFunctionType());
    Function *clspv_fract_fn = F.getParent()->getFunction(clspv_fract_name);
    if (!clspv_fract_fn) {
      // Make the clspv_fract function.
      clspv_fract_fn = cast<Function>(
          F.getParent()
              ->getOrInsertFunction(clspv_fract_name, F.getFunctionType())
              .getCallee());
      clspv_fract_fn->addFnAttr(Attribute::ReadNone);
      clspv_fract_fn->setCallingConv(CallingConv::SPIR_FUNC);
    }

    auto ceil = Intrinsic::getDeclaration(F.getParent(), Intrinsic::ceil,
                                          Call->getType());
    auto floor = Intrinsic::getDeclaration(F.getParent(), Intrinsic::floor,
                                           Call->getType());
    auto fabs = Intrinsic::getDeclaration(F.getParent(), Intrinsic::fabs,
                                          Call->getType());
    auto copysign = Intrinsic::getDeclaration(
        F.getParent(), Intrinsic::copysign, {Call->getType(), Call->getType()});

    IRBuilder<> builder(Call);

    auto fabs_call = builder.CreateCall(F.getFunctionType(), fabs, {x});
    auto ceil_call = builder.CreateCall(F.getFunctionType(), ceil, {fabs_call});
    auto floor_call =
        builder.CreateCall(F.getFunctionType(), floor, {fabs_call});
    auto fract_call =
        builder.CreateCall(F.getFunctionType(), clspv_fract_fn, {fabs_call});
    auto cmp = builder.CreateFCmpOGE(fract_call, halfway);
    auto sel = builder.CreateSelect(cmp, ceil_call, floor_call);
    return builder.CreateCall(copysign->getFunctionType(), copysign, {sel, x});
  });
}

bool ReplaceOpenCLBuiltinPass::replaceTrigPi(Function &F,
                                             Builtins::BuiltinType type) {
  return replaceCallsWithValue(F, [&F, type](CallInst *Call) -> Value * {
    const auto x = Call->getArgOperand(0);
    const double k_pi = 0x1.921fb54442d18p+1;
    Constant *pi = ConstantFP::get(x->getType(), k_pi);

    IRBuilder<> builder(Call);
    auto mul = builder.CreateFMul(x, pi);
    switch (type) {
    case Builtins::kSinpi: {
      auto func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin,
                                            x->getType());
      return builder.CreateCall(func->getFunctionType(), func, {mul});
    }
    case Builtins::kCospi: {
      auto func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos,
                                            x->getType());
      return builder.CreateCall(func->getFunctionType(), func, {mul});
    }
    case Builtins::kTanpi: {
      auto sin = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin,
                                           x->getType());
      auto sin_call = builder.CreateCall(sin->getFunctionType(), sin, {mul});
      auto cos = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos,
                                           x->getType());
      auto cos_call = builder.CreateCall(cos->getFunctionType(), cos, {mul});
      return builder.CreateFDiv(sin_call, cos_call);
    }
    default:
      llvm_unreachable("unexpected builtin");
      break;
    }
    return nullptr;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceSincos(Function &F) {
  return replaceCallsWithValue(F, [&F](CallInst *Call) {
    auto sin_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin,
                                              Call->getType());
    auto cos_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos,
                                              Call->getType());

    IRBuilder<> builder(Call);
    auto sin = builder.CreateCall(sin_func->getFunctionType(), sin_func,
                                  {Call->getArgOperand(0)});
    auto cos = builder.CreateCall(cos_func->getFunctionType(), cos_func,
                                  {Call->getArgOperand(0)});
    builder.CreateStore(cos, Call->getArgOperand(1));
    return sin;
  });
}

bool ReplaceOpenCLBuiltinPass::replaceExpm1(Function &F) {
  return replaceCallsWithValue(F, [&F](CallInst *Call) {
    auto exp_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::exp,
                                              Call->getType());

    IRBuilder<> builder(Call);
    auto exp = builder.CreateCall(exp_func->getFunctionType(), exp_func,
                                  {Call->getArgOperand(0)});
    return builder.CreateFSub(exp, ConstantFP::get(Call->getType(), 1.0));
  });
}

bool ReplaceOpenCLBuiltinPass::replacePown(Function &F) {
  return replaceCallsWithValue(F, [&F](CallInst *Call) {
    auto pow_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::pow,
                                              Call->getType());

    IRBuilder<> builder(Call);
    auto conv = builder.CreateSIToFP(Call->getArgOperand(1), Call->getType());
    return builder.CreateCall(pow_func->getFunctionType(), pow_func,
                              {Call->getArgOperand(0), conv});
  });
}