diff --git a/llvm/lib/Target/SPIRV/CMakeLists.txt b/llvm/lib/Target/SPIRV/CMakeLists.txt
--- a/llvm/lib/Target/SPIRV/CMakeLists.txt
+++ b/llvm/lib/Target/SPIRV/CMakeLists.txt
@@ -30,6 +30,7 @@
   SPIRVPrepareFunctions.cpp
   SPIRVRegisterBankInfo.cpp
   SPIRVRegisterInfo.cpp
+  SPIRVRegularizer.cpp
   SPIRVSubtarget.cpp
   SPIRVTargetMachine.cpp
   SPIRVUtils.cpp
diff --git a/llvm/lib/Target/SPIRV/SPIRV.h b/llvm/lib/Target/SPIRV/SPIRV.h
--- a/llvm/lib/Target/SPIRV/SPIRV.h
+++ b/llvm/lib/Target/SPIRV/SPIRV.h
@@ -20,6 +20,7 @@
 class RegisterBankInfo;
 
 ModulePass *createSPIRVPrepareFunctionsPass();
+FunctionPass *createSPIRVRegularizerPass();
 FunctionPass *createSPIRVPreLegalizerPass();
 FunctionPass *createSPIRVEmitIntrinsicsPass(SPIRVTargetMachine *TM);
 InstructionSelector *
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -809,7 +809,7 @@
 }
 
 // These queries ask for a single size_t result for a given dimension index, e.g
-// size_t get_global_id(uintt dimindex). In SPIR-V, the builtins corresonding to
+// size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
 // these values are all vec3 types, so we need to extract the correct index or
 // return defaultVal (0 or 1 depending on the query). We also handle extending
 // or tuncating in case size_t does not match the expected result type's
@@ -1655,16 +1655,15 @@
 static const SPIRV::DemangledType *findBuiltinType(StringRef Name) {
   if (Name.startswith("opencl."))
     return SPIRV::lookupBuiltinType(Name);
-  if (Name.startswith("spirv.")) {
-    // Some SPIR-V builtin types have a complex list of parameters as part of
-    // their name (e.g. spirv.Image._void_1_0_0_0_0_0_0). Those parameters often
-    // are numeric literals which cannot be easily represented by TableGen
-    // records and should be parsed instead.
-    unsigned BaseTypeNameLength =
-        Name.contains('_') ? Name.find('_') - 1 : Name.size();
-    return SPIRV::lookupBuiltinType(Name.substr(0, BaseTypeNameLength).str());
-  }
-  return nullptr;
+  if (!Name.startswith("spirv."))
+    return nullptr;
+  // Some SPIR-V builtin types have a complex list of parameters as part of
+  // their name (e.g. spirv.Image._void_1_0_0_0_0_0_0). Those parameters often
+  // are numeric literals which cannot be easily represented by TableGen
+  // records and should be parsed instead.
+  unsigned BaseTypeNameLength =
+      Name.contains('_') ? Name.find('_') - 1 : Name.size();
+  return SPIRV::lookupBuiltinType(Name.substr(0, BaseTypeNameLength).str());
 }
 
 static std::unique_ptr<const SPIRV::ImageType>
@@ -1674,37 +1673,36 @@
     const SPIRV::ImageType *Record = SPIRV::lookupImageType(Name);
     return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType(*Record));
   }
-  if (Name.startswith("spirv.")) {
-    // Parse the literals of SPIR-V image builtin parameters. The name should
-    // have the following format:
-    // spirv.Image._Type_Dim_Depth_Arrayed_MS_Sampled_ImageFormat_AccessQualifier
-    // e.g. %spirv.Image._void_1_0_0_0_0_0_0
-    StringRef TypeParametersString = Name.substr(strlen("spirv.Image."));
-    SmallVector<StringRef> TypeParameters;
-    SplitString(TypeParametersString, TypeParameters, "_");
-    assert(TypeParameters.size() == 8 &&
-           "Wrong number of literals in SPIR-V builtin image type");
-
-    StringRef SampledType = TypeParameters[0];
-    unsigned Dim, Depth, Arrayed, Multisampled, Sampled, Format, AccessQual;
-    bool AreParameterLiteralsValid =
-        !(TypeParameters[1].getAsInteger(10, Dim) ||
-          TypeParameters[2].getAsInteger(10, Depth) ||
-          TypeParameters[3].getAsInteger(10, Arrayed) ||
-          TypeParameters[4].getAsInteger(10, Multisampled) ||
-          TypeParameters[5].getAsInteger(10, Sampled) ||
-          TypeParameters[6].getAsInteger(10, Format) ||
-          TypeParameters[7].getAsInteger(10, AccessQual));
-    assert(AreParameterLiteralsValid &&
-           "Invalid format of SPIR-V image type parameter literals.");
-
-    return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType{
-        Name, SampledType, SPIRV::AccessQualifier::AccessQualifier(AccessQual),
-        SPIRV::Dim::Dim(Dim), static_cast<bool>(Arrayed),
-        static_cast<bool>(Depth), static_cast<bool>(Multisampled),
-        static_cast<bool>(Sampled), SPIRV::ImageFormat::ImageFormat(Format)});
-  }
-  llvm_unreachable("Unknown builtin image type name/literal");
+  if (!Name.startswith("spirv."))
+    llvm_unreachable("Unknown builtin image type name/literal");
+  // Parse the literals of SPIR-V image builtin parameters. The name should
+  // have the following format:
+  // spirv.Image._Type_Dim_Depth_Arrayed_MS_Sampled_ImageFormat_AccessQualifier
+  // e.g. %spirv.Image._void_1_0_0_0_0_0_0
+  StringRef TypeParametersString = Name.substr(strlen("spirv.Image."));
+  SmallVector<StringRef> TypeParameters;
+  SplitString(TypeParametersString, TypeParameters, "_");
+  assert(TypeParameters.size() == 8 &&
+         "Wrong number of literals in SPIR-V builtin image type");
+
+  StringRef SampledType = TypeParameters[0];
+  unsigned Dim, Depth, Arrayed, Multisampled, Sampled, Format, AccessQual;
+  bool AreParameterLiteralsValid =
+      !(TypeParameters[1].getAsInteger(10, Dim) ||
+        TypeParameters[2].getAsInteger(10, Depth) ||
+        TypeParameters[3].getAsInteger(10, Arrayed) ||
+        TypeParameters[4].getAsInteger(10, Multisampled) ||
+        TypeParameters[5].getAsInteger(10, Sampled) ||
+        TypeParameters[6].getAsInteger(10, Format) ||
+        TypeParameters[7].getAsInteger(10, AccessQual));
+  assert(AreParameterLiteralsValid &&
+         "Invalid format of SPIR-V image type parameter literals.");
+
+  return std::unique_ptr<SPIRV::ImageType>(new SPIRV::ImageType{
+      Name, SampledType, SPIRV::AccessQualifier::AccessQualifier(AccessQual),
+      SPIRV::Dim::Dim(Dim), static_cast<bool>(Arrayed),
+      static_cast<bool>(Depth), static_cast<bool>(Multisampled),
+      static_cast<bool>(Sampled), SPIRV::ImageFormat::ImageFormat(Format)});
 }
 
 static std::unique_ptr<const SPIRV::PipeType>
@@ -1714,46 +1712,46 @@
     const SPIRV::PipeType *Record = SPIRV::lookupPipeType(Name);
     return std::unique_ptr<SPIRV::PipeType>(new SPIRV::PipeType(*Record));
   }
-  if (Name.startswith("spirv.")) {
-    // Parse the access qualifier literal in the name of the SPIR-V pipe type.
-    // The name should have the following format:
-    // spirv.Pipe._AccessQualifier
-    // e.g. %spirv.Pipe._1
-    if (Name.endswith("_0"))
-      return std::unique_ptr<SPIRV::PipeType>(
-          new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadOnly});
-    if (Name.endswith("_1"))
-      return std::unique_ptr<SPIRV::PipeType>(
-          new SPIRV::PipeType{Name, SPIRV::AccessQualifier::WriteOnly});
-    if (Name.endswith("_2"))
-      return std::unique_ptr<SPIRV::PipeType>(
-          new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadWrite});
-    llvm_unreachable("Unknown pipe type access qualifier literal");
-  }
-  llvm_unreachable("Unknown builtin pipe type name/literal");
+  if (!Name.startswith("spirv."))
+    llvm_unreachable("Unknown builtin pipe type name/literal");
+  // Parse the access qualifier literal in the name of the SPIR-V pipe type.
+  // The name should have the following format:
+  // spirv.Pipe._AccessQualifier
+  // e.g. %spirv.Pipe._1
+  if (Name.endswith("_0"))
+    return std::unique_ptr<SPIRV::PipeType>(
+        new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadOnly});
+  if (Name.endswith("_1"))
+    return std::unique_ptr<SPIRV::PipeType>(
+        new SPIRV::PipeType{Name, SPIRV::AccessQualifier::WriteOnly});
+  if (Name.endswith("_2"))
+    return std::unique_ptr<SPIRV::PipeType>(
+        new SPIRV::PipeType{Name, SPIRV::AccessQualifier::ReadWrite});
+  llvm_unreachable("Unknown pipe type access qualifier literal");
 }
 
 //===----------------------------------------------------------------------===//
 // Implementation functions for builtin types.
 //===----------------------------------------------------------------------===//
 
-SPIRVType *getNonParametrizedType(const StructType *OpaqueType,
-                                  const SPIRV::DemangledType *TypeRecord,
-                                  MachineIRBuilder &MIRBuilder,
-                                  SPIRVGlobalRegistry *GR) {
+static SPIRVType *getNonParametrizedType(const StructType *OpaqueType,
+                                         const SPIRV::DemangledType *TypeRecord,
+                                         MachineIRBuilder &MIRBuilder,
+                                         SPIRVGlobalRegistry *GR) {
   unsigned Opcode = TypeRecord->Opcode;
   // Create or get an existing type from GlobalRegistry.
   return GR->getOrCreateOpTypeByOpcode(OpaqueType, MIRBuilder, Opcode);
 }
 
-SPIRVType *getSamplerType(MachineIRBuilder &MIRBuilder,
-                          SPIRVGlobalRegistry *GR) {
+static SPIRVType *getSamplerType(MachineIRBuilder &MIRBuilder,
+                                 SPIRVGlobalRegistry *GR) {
   // Create or get an existing type from GlobalRegistry.
   return GR->getOrCreateOpTypeSampler(MIRBuilder);
 }
 
-SPIRVType *getPipeType(const StructType *OpaqueType,
-                       MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
+static SPIRVType *getPipeType(const StructType *OpaqueType,
+                              MachineIRBuilder &MIRBuilder,
+                              SPIRVGlobalRegistry *GR) {
   // Lookup pipe type lowering details in TableGen records or parse the
   // name/literal for details.
   std::unique_ptr<const SPIRV::PipeType> Record =
@@ -1762,9 +1760,10 @@
   return GR->getOrCreateOpTypePipe(MIRBuilder, Record.get()->Qualifier);
 }
 
-SPIRVType *getImageType(const StructType *OpaqueType,
-                        SPIRV::AccessQualifier::AccessQualifier AccessQual,
-                        MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
+static SPIRVType *
+getImageType(const StructType *OpaqueType,
+             SPIRV::AccessQualifier::AccessQualifier AccessQual,
+             MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
   // Lookup image type lowering details in TableGen records or parse the
   // name/literal for details.
   std::unique_ptr<const SPIRV::ImageType> Record =
@@ -1781,9 +1780,9 @@
           : Record.get()->Qualifier);
 }
 
-SPIRVType *getSampledImageType(const StructType *OpaqueType,
-                               MachineIRBuilder &MIRBuilder,
-                               SPIRVGlobalRegistry *GR) {
+static SPIRVType *getSampledImageType(const StructType *OpaqueType,
+                                      MachineIRBuilder &MIRBuilder,
+                                      SPIRVGlobalRegistry *GR) {
   StringRef TypeParametersString =
       OpaqueType->getName().substr(strlen("spirv.SampledImage."));
   LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
diff --git a/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp b/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
@@ -286,7 +286,7 @@
   Register ResVReg =
       Info.OrigRet.Regs.empty() ? Register(0) : Info.OrigRet.Regs[0];
   std::string FuncName = Info.Callee.getGlobal()->getGlobalIdentifier();
-  std::string DemangledName = mayBeOclOrSpirvBuiltin(FuncName);
+  std::string DemangledName = getOclOrSpirvBuiltinDemangledName(FuncName);
   const auto *ST = static_cast<const SPIRVSubtarget *>(&MF.getSubtarget());
   // TODO: check that it's OCL builtin, then apply OpenCL_std.
   if (!DemangledName.empty() && CF && CF->isDeclaration() &&
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -544,26 +544,6 @@
   return MIB;
 }
 
-static bool isOpenCLBuiltinType(const StructType *SType) {
-  return SType->isOpaque() && SType->hasName() &&
-         SType->getName().startswith("opencl.");
-}
-
-static bool isSPIRVBuiltinType(const StructType *SType) {
-  return SType->isOpaque() && SType->hasName() &&
-         SType->getName().startswith("spirv.");
-}
-
-static bool isSpecialType(const Type *Ty) {
-  if (auto PType = dyn_cast<PointerType>(Ty)) {
-    if (!PType->isOpaque())
-      Ty = PType->getNonOpaquePointerElementType();
-  }
-  if (auto SType = dyn_cast<StructType>(Ty))
-    return isOpenCLBuiltinType(SType) || isSPIRVBuiltinType(SType);
-  return false;
-}
-
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSpecialType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccQual) {
@@ -574,7 +554,7 @@
     Ty = PType->getNonOpaquePointerElementType();
   }
   auto SType = cast<StructType>(Ty);
-  assert(isOpenCLBuiltinType(SType) || isSPIRVBuiltinType(SType));
+  assert(isSpecialOpaqueType(SType) && "Not a special opaque builtin type");
   return SPIRV::lowerBuiltinType(SType, AccQual, MIRBuilder, this);
 }
 
@@ -639,7 +619,7 @@
 SPIRVType *SPIRVGlobalRegistry::createSPIRVType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccQual, bool EmitIR) {
-  if (isSpecialType(Ty))
+  if (isSpecialOpaqueType(Ty))
     return getOrCreateSpecialType(Ty, MIRBuilder, AccQual);
   auto &TypeToSPIRVTypeMap = DT.getTypes()->getAllUses();
   auto t = TypeToSPIRVTypeMap.find(Ty);
@@ -725,7 +705,7 @@
   // Do not add OpTypeForwardPointer to DT, a corresponding normal pointer type
   // will be added later. For special types it is already added to DT.
   if (SpirvType->getOpcode() != SPIRV::OpTypeForwardPointer && !Reg.isValid() &&
-      !isSpecialType(Ty))
+      !isSpecialOpaqueType(Ty))
     DT.add(Ty, &MIRBuilder.getMF(), getSPIRVTypeID(SpirvType));
 
   return SpirvType;
@@ -745,7 +725,7 @@
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccessQual, bool EmitIR) {
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
-  if (Reg.isValid() && !isSpecialType(Ty))
+  if (Reg.isValid() && !isSpecialOpaqueType(Ty))
     return getSPIRVTypeForVReg(Reg);
   TypesInProcessing.clear();
   SPIRVType *STy = restOfCreateSPIRVType(Ty, MIRBuilder, AccessQual, EmitIR);
diff --git a/llvm/lib/Target/SPIRV/SPIRVISelLowering.h b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
--- a/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
+++ b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
@@ -41,6 +41,9 @@
                                          EVT VT) const override;
   MVT getRegisterTypeForCallingConv(LLVMContext &Context, CallingConv::ID CC,
                                     EVT VT) const override;
+  bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallInst &I,
+                          MachineFunction &MF,
+                          unsigned Intrinsic) const override;
 };
 } // namespace llvm
 
diff --git a/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp b/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
@@ -12,6 +12,7 @@
 
 #include "SPIRVISelLowering.h"
 #include "SPIRV.h"
+#include "llvm/IR/IntrinsicsSPIRV.h"
 
 #define DEBUG_TYPE "spirv-lower"
 
@@ -43,3 +44,31 @@
   }
   return getRegisterType(Context, VT);
 }
+
+bool SPIRVTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info,
+                                             const CallInst &I,
+                                             MachineFunction &MF,
+                                             unsigned Intrinsic) const {
+  unsigned AlignIdx = 3;
+  switch (Intrinsic) {
+  case Intrinsic::spv_load:
+    AlignIdx = 2;
+    LLVM_FALLTHROUGH;
+  case Intrinsic::spv_store: {
+    if (I.getNumOperands() >= AlignIdx + 1) {
+      auto *AlignOp = cast<ConstantInt>(I.getOperand(AlignIdx));
+      Info.align = Align(AlignOp->getZExtValue());
+    }
+    Info.flags = static_cast<MachineMemOperand::Flags>(
+        cast<ConstantInt>(I.getOperand(AlignIdx - 1))->getZExtValue());
+    Info.memVT = MVT::i64;
+    // TODO: take into account opaque pointers (don't use getElementType).
+    // MVT::getVT(PtrTy->getElementType());
+    return true;
+    break;
+  }
+  default:
+    break;
+  }
+  return false;
+}
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstrFormats.td b/llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
--- a/llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
+++ b/llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
@@ -28,4 +28,5 @@
 // Pseudo instructions
 class Pseudo<dag outs, dag ins> : Op<0, outs, ins, ""> {
   let isPseudo = 1;
+  let hasSideEffects = 0;
 }
diff --git a/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp b/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
@@ -145,6 +145,9 @@
   getActionDefinitionsBuilder({G_MEMCPY, G_MEMMOVE})
       .legalIf(all(typeInSet(0, allWritablePtrs), typeInSet(1, allPtrs)));
 
+  getActionDefinitionsBuilder(G_MEMSET).legalIf(
+      all(typeInSet(0, allWritablePtrs), typeInSet(1, allIntScalars)));
+
   getActionDefinitionsBuilder(G_ADDRSPACE_CAST)
       .legalForCartesianProduct(allPtrs, allPtrs);
 
@@ -223,8 +226,8 @@
   // Pointer-handling.
   getActionDefinitionsBuilder(G_FRAME_INDEX).legalFor({p0});
 
-  // Control-flow.
-  getActionDefinitionsBuilder(G_BRCOND).legalFor({s1});
+  // Control-flow. In some cases (e.g. constants) s1 may be promoted to s32.
+  getActionDefinitionsBuilder(G_BRCOND).legalFor({s1, s32});
 
   getActionDefinitionsBuilder({G_FPOW,
                                G_FEXP,
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -751,6 +751,7 @@
     break;
   case SPIRV::OpTypeDeviceEvent:
   case SPIRV::OpTypeQueue:
+  case SPIRV::OpBuildNDRange:
     Reqs.addCapability(SPIRV::Capability::DeviceEnqueue);
     break;
   case SPIRV::OpDecorate:
diff --git a/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp b/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
@@ -369,11 +369,19 @@
       if (MI.getOpcode() != SPIRV::ASSIGN_TYPE)
         continue;
       Register SrcReg = MI.getOperand(1).getReg();
-      if (!isTypeFoldingSupported(MRI.getVRegDef(SrcReg)->getOpcode()))
+      unsigned Opcode = MRI.getVRegDef(SrcReg)->getOpcode();
+      if (!isTypeFoldingSupported(Opcode))
         continue;
       Register DstReg = MI.getOperand(0).getReg();
       if (MRI.getType(DstReg).isVector())
         MRI.setRegClass(DstReg, &SPIRV::IDRegClass);
+      // Don't need to reset type of register holding constant and used in
+      // G_ADDRSPACE_CAST, since it braaks legalizer.
+      if (Opcode == TargetOpcode::G_CONSTANT && MRI.hasOneUse(DstReg)) {
+        MachineInstr &UseMI = *MRI.use_instr_begin(DstReg);
+        if (UseMI.getOpcode() == TargetOpcode::G_ADDRSPACE_CAST)
+          continue;
+      }
       MRI.setType(DstReg, LLT::scalar(32));
     }
   }
diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
@@ -18,6 +18,7 @@
 #include "SPIRV.h"
 #include "SPIRVTargetMachine.h"
 #include "SPIRVUtils.h"
+#include "llvm/CodeGen/IntrinsicLowering.h"
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/Transforms/Utils/Cloning.h"
@@ -141,6 +142,69 @@
   return NewF;
 }
 
+static void lowerIntrinsicToFunction(Module *M, IntrinsicInst *Intrinsic) {
+  // For @llvm.memset.* intrinsic cases with constant value and length arguments
+  // are emulated via "storing" a constant array to the destination. For other
+  // cases we wrap the intrinsic in @spirv.llvm_memset_* function and expand the
+  // intrinsic to a loop via expandMemSetAsLoop().
+  if (auto *MSI = dyn_cast<MemSetInst>(Intrinsic))
+    if (isa<Constant>(MSI->getValue()) && isa<ConstantInt>(MSI->getLength()))
+      return; // It is handled later using OpCopyMemorySized.
+
+  std::string FuncName = lowerLLVMIntrinsicName(Intrinsic);
+  if (Intrinsic->isVolatile())
+    FuncName += ".volatile";
+  // Redirect @llvm.intrinsic.* call to @spirv.llvm_intrinsic_*
+  Function *F = M->getFunction(FuncName);
+  if (F) {
+    Intrinsic->setCalledFunction(F);
+    return;
+  }
+  // TODO copy arguments attributes: nocapture writeonly.
+  FunctionCallee FC =
+      M->getOrInsertFunction(FuncName, Intrinsic->getFunctionType());
+  auto IntrinsicID = Intrinsic->getIntrinsicID();
+  Intrinsic->setCalledFunction(FC);
+
+  F = dyn_cast<Function>(FC.getCallee());
+  assert(F && "Callee must be a function");
+
+  switch (IntrinsicID) {
+  case Intrinsic::memset: {
+    auto *MSI = static_cast<MemSetInst *>(Intrinsic);
+    Argument *Dest = F->getArg(0);
+    Argument *Val = F->getArg(1);
+    Argument *Len = F->getArg(2);
+    Argument *IsVolatile = F->getArg(3);
+    Dest->setName("dest");
+    Val->setName("val");
+    Len->setName("len");
+    IsVolatile->setName("isvolatile");
+    BasicBlock *EntryBB = BasicBlock::Create(M->getContext(), "entry", F);
+    IRBuilder<> IRB(EntryBB);
+    auto *MemSet = IRB.CreateMemSet(Dest, Val, Len, MSI->getDestAlign(),
+                                    MSI->isVolatile());
+    IRB.CreateRetVoid();
+    expandMemSetAsLoop(cast<MemSetInst>(MemSet));
+    MemSet->eraseFromParent();
+    break;
+  }
+  case Intrinsic::bswap: {
+    BasicBlock *EntryBB = BasicBlock::Create(M->getContext(), "entry", F);
+    IRBuilder<> IRB(EntryBB);
+    auto *BSwap = IRB.CreateIntrinsic(Intrinsic::bswap, Intrinsic->getType(),
+                                      F->getArg(0));
+    IRB.CreateRet(BSwap);
+    IntrinsicLowering IL(M->getDataLayout());
+    IL.LowerIntrinsicCall(BSwap);
+    break;
+  }
+  default:
+    break;
+  }
+  return;
+}
+
 static void lowerFunnelShifts(Module *M, IntrinsicInst *FSHIntrinsic) {
   // Get a separate function - otherwise, we'd have to rework the CFG of the
   // current one. Then simply replace the intrinsic uses with a call to the new
@@ -248,8 +312,11 @@
       if (!CF || !CF->isIntrinsic())
         continue;
       auto *II = cast<IntrinsicInst>(Call);
-      if (II->getIntrinsicID() == Intrinsic::fshl ||
-          II->getIntrinsicID() == Intrinsic::fshr)
+      if (II->getIntrinsicID() == Intrinsic::memset ||
+          II->getIntrinsicID() == Intrinsic::bswap)
+        lowerIntrinsicToFunction(M, II);
+      else if (II->getIntrinsicID() == Intrinsic::fshl ||
+               II->getIntrinsicID() == Intrinsic::fshr)
         lowerFunnelShifts(M, II);
       else if (II->getIntrinsicID() == Intrinsic::umul_with_overflow)
         lowerUMulWithOverflow(M, II);
diff --git a/llvm/lib/Target/SPIRV/SPIRVRegularizer.cpp b/llvm/lib/Target/SPIRV/SPIRVRegularizer.cpp
new file mode 100644
--- /dev/null
+++ b/llvm/lib/Target/SPIRV/SPIRVRegularizer.cpp
@@ -0,0 +1,249 @@
+//===-- SPIRVRegularizer.cpp - regularize IR for SPIR-V ---------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This pass implements regularization of LLVM IR for SPIR-V. The prototype of
+// the pass was taken from SPIRV-LLVM translator.
+//
+//===----------------------------------------------------------------------===//
+
+#include "SPIRV.h"
+#include "SPIRVTargetMachine.h"
+#include "llvm/Demangle/Demangle.h"
+#include "llvm/IR/InstIterator.h"
+#include "llvm/IR/InstVisitor.h"
+#include "llvm/IR/PassManager.h"
+#include "llvm/Transforms/Utils/Cloning.h"
+
+#include <list>
+
+#define DEBUG_TYPE "spirv-regularizer"
+
+using namespace llvm;
+
+namespace llvm {
+void initializeSPIRVRegularizerPass(PassRegistry &);
+}
+
+namespace {
+struct SPIRVRegularizer : public FunctionPass, InstVisitor<SPIRVRegularizer> {
+  DenseMap<Function *, Function *> Old2NewFuncs;
+
+public:
+  static char ID;
+  SPIRVRegularizer() : FunctionPass(ID) {
+    initializeSPIRVRegularizerPass(*PassRegistry::getPassRegistry());
+  }
+  bool runOnFunction(Function &F) override;
+  StringRef getPassName() const override { return "SPIR-V Regularizer"; }
+
+  void getAnalysisUsage(AnalysisUsage &AU) const override {
+    FunctionPass::getAnalysisUsage(AU);
+  }
+  void visitCallInst(CallInst &CI);
+
+private:
+  void visitCallScalToVec(CallInst *CI, StringRef MangledName,
+                          StringRef DemangledName);
+  void runLowerConstExpr(Function &F);
+};
+} // namespace
+
+char SPIRVRegularizer::ID = 0;
+
+INITIALIZE_PASS(SPIRVRegularizer, DEBUG_TYPE, "SPIR-V Regularizer", false,
+                false)
+
+// Since SPIR-V cannot represent constant expression, constant expressions
+// in LLVM IR need to be lowered to instructions. For each function,
+// the constant expressions used by instructions of the function are replaced
+// by instructions placed in the entry block since it dominates all other BBs.
+// Each constant expression only needs to be lowered once in each function
+// and all uses of it by instructions in that function are replaced by
+// one instruction.
+// TODO: remove redundant instructions for common subexpression.
+void SPIRVRegularizer::runLowerConstExpr(Function &F) {
+  LLVMContext &Ctx = F.getContext();
+  std::list<Instruction *> WorkList;
+  for (auto &II : instructions(F))
+    WorkList.push_back(&II);
+
+  auto FBegin = F.begin();
+  while (!WorkList.empty()) {
+    Instruction *II = WorkList.front();
+
+    auto LowerOp = [&II, &FBegin, &F](Value *V) -> Value * {
+      if (isa<Function>(V))
+        return V;
+      auto *CE = cast<ConstantExpr>(V);
+      LLVM_DEBUG(dbgs() << "[lowerConstantExpressions] " << *CE);
+      auto ReplInst = CE->getAsInstruction();
+      auto InsPoint = II->getParent() == &*FBegin ? II : &FBegin->back();
+      ReplInst->insertBefore(InsPoint);
+      LLVM_DEBUG(dbgs() << " -> " << *ReplInst << '\n');
+      std::vector<Instruction *> Users;
+      // Do not replace use during iteration of use. Do it in another loop.
+      for (auto U : CE->users()) {
+        LLVM_DEBUG(dbgs() << "[lowerConstantExpressions] Use: " << *U << '\n');
+        auto InstUser = dyn_cast<Instruction>(U);
+        // Only replace users in scope of current function.
+        if (InstUser && InstUser->getParent()->getParent() == &F)
+          Users.push_back(InstUser);
+      }
+      for (auto &User : Users) {
+        if (ReplInst->getParent() == User->getParent() &&
+            User->comesBefore(ReplInst))
+          ReplInst->moveBefore(User);
+        User->replaceUsesOfWith(CE, ReplInst);
+      }
+      return ReplInst;
+    };
+
+    WorkList.pop_front();
+    auto LowerConstantVec = [&II, &LowerOp, &WorkList,
+                             &Ctx](ConstantVector *Vec,
+                                   unsigned NumOfOp) -> Value * {
+      if (std::all_of(Vec->op_begin(), Vec->op_end(), [](Value *V) {
+            return isa<ConstantExpr>(V) || isa<Function>(V);
+          })) {
+        // Expand a vector of constexprs and construct it back with
+        // series of insertelement instructions.
+        std::list<Value *> OpList;
+        std::transform(Vec->op_begin(), Vec->op_end(),
+                       std::back_inserter(OpList),
+                       [LowerOp](Value *V) { return LowerOp(V); });
+        Value *Repl = nullptr;
+        unsigned Idx = 0;
+        auto *PhiII = dyn_cast<PHINode>(II);
+        Instruction *InsPoint =
+            PhiII ? &PhiII->getIncomingBlock(NumOfOp)->back() : II;
+        std::list<Instruction *> ReplList;
+        for (auto V : OpList) {
+          if (auto *Inst = dyn_cast<Instruction>(V))
+            ReplList.push_back(Inst);
+          Repl = InsertElementInst::Create(
+              (Repl ? Repl : PoisonValue::get(Vec->getType())), V,
+              ConstantInt::get(Type::getInt32Ty(Ctx), Idx++), "", InsPoint);
+        }
+        WorkList.splice(WorkList.begin(), ReplList);
+        return Repl;
+      }
+      return nullptr;
+    };
+    for (unsigned OI = 0, OE = II->getNumOperands(); OI != OE; ++OI) {
+      auto *Op = II->getOperand(OI);
+      if (auto *Vec = dyn_cast<ConstantVector>(Op)) {
+        Value *ReplInst = LowerConstantVec(Vec, OI);
+        if (ReplInst)
+          II->replaceUsesOfWith(Op, ReplInst);
+      } else if (auto CE = dyn_cast<ConstantExpr>(Op)) {
+        WorkList.push_front(cast<Instruction>(LowerOp(CE)));
+      } else if (auto MDAsVal = dyn_cast<MetadataAsValue>(Op)) {
+        auto ConstMD = dyn_cast<ConstantAsMetadata>(MDAsVal->getMetadata());
+        if (!ConstMD)
+          continue;
+        Constant *C = ConstMD->getValue();
+        Value *ReplInst = nullptr;
+        if (auto *Vec = dyn_cast<ConstantVector>(C))
+          ReplInst = LowerConstantVec(Vec, OI);
+        if (auto *CE = dyn_cast<ConstantExpr>(C))
+          ReplInst = LowerOp(CE);
+        if (!ReplInst)
+          continue;
+        Metadata *RepMD = ValueAsMetadata::get(ReplInst);
+        Value *RepMDVal = MetadataAsValue::get(Ctx, RepMD);
+        II->setOperand(OI, RepMDVal);
+        WorkList.push_front(cast<Instruction>(ReplInst));
+      }
+    }
+  }
+}
+
+// It fixes calls to OCL builtins that accept vector arguments and one of them
+// is actually a scalar splat.
+void SPIRVRegularizer::visitCallInst(CallInst &CI) {
+  auto F = CI.getCalledFunction();
+  if (!F)
+    return;
+
+  auto MangledName = F->getName();
+  size_t n;
+  int status;
+  char *NameStr = itaniumDemangle(F->getName().data(), nullptr, &n, &status);
+  StringRef DemangledName(NameStr);
+
+  // TODO: add support for other builtins.
+  if (DemangledName.startswith("fmin") || DemangledName.startswith("fmax") ||
+      DemangledName.startswith("min") || DemangledName.startswith("max"))
+    visitCallScalToVec(&CI, MangledName, DemangledName);
+  free(NameStr);
+}
+
+void SPIRVRegularizer::visitCallScalToVec(CallInst *CI, StringRef MangledName,
+                                          StringRef DemangledName) {
+  // Check if all arguments have the same type - it's simple case.
+  auto Uniform = true;
+  Type *Arg0Ty = CI->getOperand(0)->getType();
+  auto IsArg0Vector = isa<VectorType>(Arg0Ty);
+  for (unsigned I = 1, E = CI->arg_size(); Uniform && (I != E); ++I)
+    Uniform = isa<VectorType>(CI->getOperand(I)->getType()) == IsArg0Vector;
+  if (Uniform)
+    return;
+
+  auto *OldF = CI->getCalledFunction();
+  Function *NewF = nullptr;
+  if (!Old2NewFuncs.count(OldF)) {
+    AttributeList Attrs = CI->getCalledFunction()->getAttributes();
+    SmallVector<Type *, 2> ArgTypes = {OldF->getArg(0)->getType(), Arg0Ty};
+    auto *NewFTy =
+        FunctionType::get(OldF->getReturnType(), ArgTypes, OldF->isVarArg());
+    NewF = Function::Create(NewFTy, OldF->getLinkage(), OldF->getName(),
+                            *OldF->getParent());
+    ValueToValueMapTy VMap;
+    auto NewFArgIt = NewF->arg_begin();
+    for (auto &Arg : OldF->args()) {
+      auto ArgName = Arg.getName();
+      NewFArgIt->setName(ArgName);
+      VMap[&Arg] = &(*NewFArgIt++);
+    }
+    SmallVector<ReturnInst *, 8> Returns;
+    CloneFunctionInto(NewF, OldF, VMap,
+                      CloneFunctionChangeType::LocalChangesOnly, Returns);
+    NewF->setAttributes(Attrs);
+    Old2NewFuncs[OldF] = NewF;
+  } else {
+    NewF = Old2NewFuncs[OldF];
+  }
+  assert(NewF);
+
+  auto ConstInt = ConstantInt::get(IntegerType::get(CI->getContext(), 32), 0);
+  UndefValue *UndefVal = UndefValue::get(Arg0Ty);
+  Instruction *Inst =
+      InsertElementInst::Create(UndefVal, CI->getOperand(1), ConstInt, "", CI);
+  ElementCount VecElemCount = cast<VectorType>(Arg0Ty)->getElementCount();
+  Constant *ConstVec = ConstantVector::getSplat(VecElemCount, ConstInt);
+  Value *NewVec = new ShuffleVectorInst(Inst, UndefVal, ConstVec, "", CI);
+  CI->setOperand(1, NewVec);
+  CI->replaceUsesOfWith(OldF, NewF);
+  CI->mutateFunctionType(NewF->getFunctionType());
+}
+
+bool SPIRVRegularizer::runOnFunction(Function &F) {
+  runLowerConstExpr(F);
+  visit(F);
+  for (auto &OldNew : Old2NewFuncs) {
+    Function *OldF = OldNew.first;
+    Function *NewF = OldNew.second;
+    NewF->takeName(OldF);
+    OldF->eraseFromParent();
+  }
+  return true;
+}
+
+FunctionPass *llvm::createSPIRVRegularizerPass() {
+  return new SPIRVRegularizer();
+}
diff --git a/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp b/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
@@ -70,7 +70,7 @@
     : LLVMTargetMachine(T, computeDataLayout(TT), TT, CPU, FS, Options,
                         getEffectiveRelocModel(RM),
                         getEffectiveCodeModel(CM, CodeModel::Small), OL),
-      TLOF(std::make_unique<TargetLoweringObjectFileELF>()),
+      TLOF(std::make_unique<SPIRVTargetObjectFile>()),
       Subtarget(TT, CPU.str(), FS.str(), *this) {
   initAsmInfo();
   setGlobalISel(true);
@@ -142,6 +142,7 @@
 
 void SPIRVPassConfig::addIRPasses() {
   TargetPassConfig::addIRPasses();
+  addPass(createSPIRVRegularizerPass());
   addPass(createSPIRVPrepareFunctionsPass());
 }
 
@@ -159,13 +160,13 @@
   addPass(createSPIRVPreLegalizerPass());
 }
 
-// Use a default legalizer.
+// Use the default legalizer.
 bool SPIRVPassConfig::addLegalizeMachineIR() {
   addPass(new Legalizer());
   return false;
 }
 
-// Do not add a RegBankSelect pass, as we only ever need virtual registers.
+// Do not add the RegBankSelect pass, as we only ever need virtual registers.
 bool SPIRVPassConfig::addRegBankSelect() {
   disablePass(&RegBankSelect::ID);
   return false;
@@ -183,6 +184,7 @@
 };
 } // namespace
 
+// Add the custom SPIRVInstructionSelect from above.
 bool SPIRVPassConfig::addGlobalInstructionSelect() {
   addPass(new SPIRVInstructionSelect());
   return false;
diff --git a/llvm/lib/Target/SPIRV/SPIRVUtils.h b/llvm/lib/Target/SPIRV/SPIRVUtils.h
--- a/llvm/lib/Target/SPIRV/SPIRVUtils.h
+++ b/llvm/lib/Target/SPIRV/SPIRVUtils.h
@@ -84,8 +84,11 @@
 // Get type of i-th operand of the metadata node.
 Type *getMDOperandAsType(const MDNode *N, unsigned I);
 
-// Return a demangled name with arg type info by itaniumDemangle().
-// If the parser fails, return only function name.
-std::string mayBeOclOrSpirvBuiltin(StringRef Name);
+// If OpenCL or SPIR-V builtin function name is recognized, return a demangled
+// name, otherwise return an empty string.
+std::string getOclOrSpirvBuiltinDemangledName(StringRef Name);
+
+// Check if given LLVM type is a special opaque builtin type.
+bool isSpecialOpaqueType(const Type *Ty);
 } // namespace llvm
 #endif // LLVM_LIB_TARGET_SPIRV_SPIRVUTILS_H
diff --git a/llvm/lib/Target/SPIRV/SPIRVUtils.cpp b/llvm/lib/Target/SPIRV/SPIRVUtils.cpp
--- a/llvm/lib/Target/SPIRV/SPIRVUtils.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVUtils.cpp
@@ -289,7 +289,7 @@
          Name == "__translate_sampler_initializer";
 }
 
-std::string mayBeOclOrSpirvBuiltin(StringRef Name) {
+std::string getOclOrSpirvBuiltinDemangledName(StringRef Name) {
   bool IsNonMangledOCL = isNonMangledOCLBuiltin(Name);
   bool IsNonMangledSPIRV = Name.startswith("__spirv_");
   bool IsMangled = Name.startswith("_Z");
@@ -331,4 +331,24 @@
       .getAsInteger(10, Len);
   return Name.substr(Start, Len).str();
 }
+
+static bool isOpenCLBuiltinType(const StructType *SType) {
+  return SType->isOpaque() && SType->hasName() &&
+         SType->getName().startswith("opencl.");
+}
+
+static bool isSPIRVBuiltinType(const StructType *SType) {
+  return SType->isOpaque() && SType->hasName() &&
+         SType->getName().startswith("spirv.");
+}
+
+bool isSpecialOpaqueType(const Type *Ty) {
+  if (auto PType = dyn_cast<PointerType>(Ty)) {
+    if (!PType->isOpaque())
+      Ty = PType->getNonOpaquePointerElementType();
+  }
+  if (auto SType = dyn_cast<StructType>(Ty))
+    return isOpenCLBuiltinType(SType) || isSPIRVBuiltinType(SType);
+  return false;
+}
 } // namespace llvm
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll
new file mode 100644
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll
@@ -0,0 +1,74 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#FuncNameInt16:]] "spirv.llvm_bswap_i16"
+; CHECK-SPIRV: OpName %[[#FuncNameInt32:]] "spirv.llvm_bswap_i32"
+; CHECK-SPIRV: OpName %[[#FuncNameInt64:]] "spirv.llvm_bswap_i64"
+
+; CHECK-SPIRV: %[[#TypeInt32:]] = OpTypeInt 32 0
+; CHECK-SPIRV: %[[#TypeInt16:]] = OpTypeInt 16 0
+; CHECK-SPIRV: %[[#TypeInt64:]] = OpTypeInt 64 0
+
+; CHECK-SPIRV: %[[#FuncNameInt16]] = OpFunction %[[#TypeInt16]]
+; CHECK-SPIRV: %[[#FuncParameter:]] = OpFunctionParameter %[[#TypeInt16]]
+; CHECK-SPIRV: %[[#]] = OpShiftLeftLogical %[[#TypeInt16]] %[[#FuncParameter]]
+; CHECK-SPIRV: %[[#]] = OpShiftRightLogical %[[#TypeInt16]] %[[#FuncParameter]]
+; CHECK-SPIRV: %[[#RetVal:]] = OpBitwiseOr %[[#TypeInt16]]
+; CHECK-SPIRV: OpReturnValue %[[#RetVal]]
+; CHECK-SPIRV: OpFunctionEnd
+
+; CHECK-SPIRV: %[[#FuncNameInt32]] = OpFunction %[[#TypeInt32]]
+; CHECK-SPIRV: %[[#FuncParameter:]] = OpFunctionParameter %[[#TypeInt32]]
+; CHECK-SPIRV-COUNT-2: %[[#]] = OpShiftLeftLogical %[[#TypeInt32]] %[[#FuncParameter]]
+; CHECK-SPIRV-COUNT-2: %[[#]] = OpShiftRightLogical %[[#TypeInt32]] %[[#FuncParameter]]
+; CHECK-SPIRV-COUNT-2: OpBitwiseAnd %[[#TypeInt32]]
+; CHECK-SPIRV-COUNT-2: OpBitwiseOr %[[#TypeInt32]]
+; CHECK-SPIRV: %[[#RetVal:]] = OpBitwiseOr %[[#TypeInt32]]
+; CHECK-SPIRV: OpReturnValue %[[#RetVal:]]
+; CHECK-SPIRV: OpFunctionEnd
+
+; CHECK-SPIRV: %[[#FuncNameInt64]]  = OpFunction %[[#TypeInt64]]
+; CHECK-SPIRV: %[[#FuncParameter:]]  = OpFunctionParameter %[[#TypeInt64]]
+; CHECK-SPIRV-COUNT-4: %[[#]] = OpShiftLeftLogical %[[#TypeInt64]] %[[#FuncParameter]] %[[#]]
+; CHECK-SPIRV-COUNT-4: %[[#]] = OpShiftRightLogical %[[#TypeInt64]] %[[#FuncParameter]] %[[#]]
+; CHECK-SPIRV-COUNT-6: OpBitwiseAnd %[[#TypeInt64]]
+; CHECK-SPIRV-COUNT-6: OpBitwiseOr %[[#TypeInt64]]
+; CHECK-SPIRV: %[[#RetVal:]] = OpBitwiseOr %[[#TypeInt64]]
+; CHECK-SPIRV: OpReturnValue %[[#RetVal]]
+; CHECK-SPIRV: OpFunctionEnd
+
+define dso_local i32 @main() {
+entry:
+  %retval = alloca i32, align 4
+  %a = alloca i16, align 2
+  %b = alloca i16, align 2
+  %h = alloca i16, align 2
+  %i = alloca i16, align 2
+  %c = alloca i32, align 4
+  %d = alloca i32, align 4
+  %e = alloca i64, align 8
+  %f = alloca i64, align 8
+  store i32 0, i32* %retval, align 4
+  store i16 258, i16* %a, align 2
+  %0 = load i16, i16* %a, align 2
+  %1 = call i16 @llvm.bswap.i16(i16 %0)
+  store i16 %1, i16* %b, align 2
+  store i16 234, i16* %h, align 2
+  %2 = load i16, i16* %h, align 2
+  %3 = call i16 @llvm.bswap.i16(i16 %2)
+  store i16 %3, i16* %i, align 2
+  store i32 566, i32* %c, align 4
+  %4 = load i32, i32* %c, align 4
+  %5 = call i32 @llvm.bswap.i32(i32 %4)
+  store i32 %5, i32* %d, align 4
+  store i64 12587, i64* %e, align 8
+  %6 = load i64, i64* %e, align 8
+  %7 = call i64 @llvm.bswap.i64(i64 %6)
+  store i64 %7, i64* %f, align 8
+  ret i32 0
+}
+
+declare i16 @llvm.bswap.i16(i16)
+
+declare i32 @llvm.bswap.i32(i32)
+
+declare i64 @llvm.bswap.i64(i64)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll
new file mode 100644
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll
@@ -0,0 +1,83 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpDecorate %[[#NonConstMemset:]] LinkageAttributes "spirv.llvm_memset_p3i8_i32"
+; CHECK-SPIRV: %[[#Int32:]] = OpTypeInt 32 0
+; CHECK-SPIRV: %[[#Int8:]] = OpTypeInt 8 0
+; CHECK-SPIRV: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]]
+; CHECK-SPIRV: %[[#Lenmemset21:]] = OpConstant %[[#]] 4
+; CHECK-SPIRV: %[[#Int8x4:]] = OpTypeArray %[[#Int8]] %[[#Lenmemset21]]
+; CHECK-SPIRV: %[[#Int8PtrConst:]] = OpTypePointer UniformConstant %[[#Int8]]
+; CHECK-SPIRV: %[[#Lenmemset0:]] = OpConstant %[[#Int32]] 12
+; CHECK-SPIRV: %[[#Int8x12:]] = OpTypeArray %[[#Int8]] %[[#Lenmemset0]]
+; CHECK-SPIRV: %[[#Const21:]] = OpConstant %[[#]] 21
+; CHECK-SPIRV: %[[#False:]] = OpConstantFalse %[[#]]
+; CHECK-SPIRV: %[[#InitComp:]] = OpConstantComposite %[[#Int8x4]] %[[#Const21]] %[[#Const21]] %[[#Const21]] %[[#Const21]]
+; CHECK-SPIRV: %[[#Init:]] = OpConstantNull %[[#Int8x12]]
+; CHECK-SPIRV: %[[#ValComp:]] = OpVariable %[[#]] UniformConstant %[[#InitComp]]
+; CHECK-SPIRV: %[[#Val:]] = OpVariable %[[#]] UniformConstant %[[#Init]]
+
+; CHECK-SPIRV: %[[#Target:]] = OpBitcast %[[#Int8Ptr]] %[[#]]
+; CHECK-SPIRV: %[[#Source:]] = OpBitcast %[[#Int8PtrConst]] %[[#Val]]
+; CHECK-SPIRV: OpCopyMemorySized %[[#Target]] %[[#Source]] %[[#Lenmemset0]] Aligned 4
+
+; CHECK-SPIRV: %[[#SourceComp:]] = OpBitcast %[[#Int8PtrConst]] %[[#ValComp]]
+; CHECK-SPIRV: OpCopyMemorySized %[[#]] %[[#SourceComp]] %[[#Lenmemset21]] Aligned 4
+
+; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#]] %[[#NonConstMemset]] %[[#]] %[[#]] %[[#]] %[[#False]]
+
+; CHECK-SPIRV: %[[#NonConstMemset]] = OpFunction %[[#]]
+; CHECK-SPIRV: %[[#Dest:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV: %[[#Value:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV: %[[#Len:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV: %[[#Volatile:]] = OpFunctionParameter %[[#]]
+
+; CHECK-SPIRV: %[[#Entry:]] = OpLabel
+; CHECK-SPIRV: %[[#IsZeroLen:]] = OpIEqual %[[#]] %[[#Zero:]] %[[#Len]]
+; CHECK-SPIRV: OpBranchConditional %[[#IsZeroLen]] %[[#End:]] %[[#WhileBody:]]
+
+; CHECK-SPIRV: %[[#WhileBody]] = OpLabel
+; CHECK-SPIRV: %[[#Offset:]] = OpPhi %[[#]] %[[#Zero]] %[[#Entry]] %[[#OffsetInc:]] %[[#WhileBody]]
+; CHECK-SPIRV: %[[#Ptr:]] = OpInBoundsPtrAccessChain %[[#]] %[[#Dest]] %[[#Offset]]
+; CHECK-SPIRV: OpStore %[[#Ptr]] %[[#Value]] Aligned 1
+; CHECK-SPIRV: %[[#OffsetInc]] = OpIAdd %[[#]] %[[#Offset]] %[[#One:]]
+; CHECK-SPIRV: %[[#NotEnd:]] = OpULessThan %[[#]] %[[#OffsetInc]] %[[#Len]]
+; CHECK-SPIRV: OpBranchConditional %[[#NotEnd]] %[[#WhileBody]] %[[#End]]
+
+; CHECK-SPIRV: %[[#End]] = OpLabel
+; CHECK-SPIRV: OpReturn
+
+; CHECK-SPIRV: OpFunctionEnd
+
+%struct.S1 = type { i32, i32, i32 }
+
+define spir_func void @_Z5foo11v(%struct.S1 addrspace(4)* noalias nocapture sret(%struct.S1 addrspace(4)*) %agg.result, i32 %s1, i64 %s2, i8 %v) {
+  %x = alloca [4 x i8]
+  %x.bc = bitcast [4 x i8]* %x to i8*
+  %1 = bitcast %struct.S1 addrspace(4)* %agg.result to i8 addrspace(4)*
+  tail call void @llvm.memset.p4i8.i32(i8 addrspace(4)* align 4 %1, i8 0, i32 12, i1 false)
+  tail call void @llvm.memset.p0i8.i32(i8* align 4 %x.bc, i8 21, i32 4, i1 false)
+
+  ;; non-const value
+  tail call void @llvm.memset.p0i8.i32(i8* align 4 %x.bc, i8 %v, i32 3, i1 false)
+
+  ;; non-const value and size
+  tail call void @llvm.memset.p0i8.i32(i8*  align 4 %x.bc, i8 %v, i32 %s1, i1 false)
+
+  ;; Address spaces, non-const value and size
+  %a = addrspacecast i8 addrspace(4)* %1 to i8 addrspace(3)*
+  tail call void @llvm.memset.p3i8.i32(i8 addrspace(3)* align 4 %a, i8 %v, i32 %s1, i1 false)
+  %b = addrspacecast i8 addrspace(4)* %1 to i8 addrspace(1)*
+  tail call void @llvm.memset.p1i8.i64(i8 addrspace(1)* align 4 %b, i8 %v, i64 %s2, i1 false)
+
+  ;; Volatile
+  tail call void @llvm.memset.p1i8.i64(i8 addrspace(1)* align 4 %b, i8 %v, i64 %s2, i1 true)
+  ret void
+}
+
+declare void @llvm.memset.p4i8.i32(i8 addrspace(4)* nocapture, i8, i32, i1)
+
+declare void @llvm.memset.p0i8.i32(i8* nocapture, i8, i32, i1)
+
+declare void @llvm.memset.p3i8.i32(i8 addrspace(3)*, i8, i32, i1)
+
+declare void @llvm.memset.p1i8.i64(i8 addrspace(1)*, i8, i64, i1)
diff --git a/llvm/test/CodeGen/SPIRV/lshr-constexpr.ll b/llvm/test/CodeGen/SPIRV/lshr-constexpr.ll
new file mode 100644
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/lshr-constexpr.ll
@@ -0,0 +1,18 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: %[[#type_int32:]] = OpTypeInt 32 0
+; CHECK-SPIRV-DAG: %[[#type_int64:]] = OpTypeInt 64 0
+; CHECK-SPIRV:     %[[#type_vec:]] = OpTypeVector %[[#type_int32]] 2
+; CHECK-SPIRV:     %[[#const1:]] = OpConstant %[[#type_int32]] 1
+; CHECK-SPIRV:     %[[#vec_const:]] = OpConstantComposite %[[#type_vec]] %[[#const1]] %[[#const1]]
+; CHECK-SPIRV:     %[[#const32:]] = OpConstant %[[#type_int64]] 32 0
+
+; CHECK-SPIRV:     %[[#bitcast_res:]] = OpBitcast %[[#type_int64]] %[[#vec_const]]
+; CHECK-SPIRV:     %[[#shift_res:]] = OpShiftRightLogical %[[#type_int64]] %[[#bitcast_res]] %[[#const32]]
+
+define void @foo(i64* %arg) {
+entry:
+  %0 = lshr i64 bitcast (<2 x i32> <i32 1, i32 1> to i64), 32
+  store i64 %0, i64* %arg
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/opencl/get_global_id.ll b/llvm/test/CodeGen/SPIRV/opencl/get_global_id.ll
new file mode 100644
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/opencl/get_global_id.ll
@@ -0,0 +1,53 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown -opaque-pointers=0 %s -o - | FileCheck %s
+
+;; The set of valid inputs for get_global_id depends on the runtime NDRange,
+;; but inputs outside of [0, 2] always return 0.
+;; Here we assume Itanium mangling for function name.
+declare i64 @_Z13get_global_idj(i32)
+
+define i64 @foo(i32 %dim) {
+  %x = call i64 @_Z13get_global_idj(i32 0)
+  %zero = call i64 @_Z13get_global_idj(i32 5)
+  %unknown = call i64 @_Z13get_global_idj(i32 %dim)
+  %acc = add i64 %x, %zero
+  %ret = add i64 %acc, %unknown
+  ret i64 %ret
+}
+
+;; Capabilities:
+; CHECK-DAG: OpCapability Kernel
+; CHECK-DAG: OpCapability Int64
+
+; CHECK-NOT: DAG-FENCE
+
+;; Decorations:
+; CHECK-DAG: OpDecorate %[[#GET_GLOBAL_ID:]] BuiltIn GlobalInvocationId
+; CHECK-DAG: OpDecorate %[[#GET_GLOBAL_ID]] Constant
+
+; CHECK-NOT: DAG-FENCE
+
+;; Types, Constants and Variables:
+; CHECK-DAG: %[[#BOOL:]] = OpTypeBool
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#I64:]] = OpTypeInt 64 0
+; CHECK-DAG: %[[#VEC:]] = OpTypeVector %[[#I64]] 3
+; CHECK-DAG: %[[#PTR:]] = OpTypePointer Input %[[#VEC]]
+; CHECK-DAG: %[[#FN:]] = OpTypeFunction %[[#I64]] %[[#I32]]
+; CHECK-DAG: %[[#GET_GLOBAL_ID]] = OpVariable %[[#PTR]] Input
+; CHECK-DAG: %[[#ZERO:]] = OpConstantNull %[[#I64]]
+; CHECK-DAG: %[[#THREE:]] = OpConstant %[[#I32]] 3
+
+;; Functions:
+; CHECK:     OpFunction %[[#I64]] None %[[#FN]]
+; CHECK:     %[[#DIM:]] = OpFunctionParameter %[[#I32]]
+
+;; get_global_id(0): OpLoad + OpCompositeExtract.
+; CHECK:     %[[#TMP1:]] = OpLoad %[[#VEC]] %[[#GET_GLOBAL_ID]]
+; CHECK:     %[[#X:]] = OpCompositeExtract %[[#I64]] %[[#TMP1]] 0
+
+;; get_global_id(5): OpConstant (above) of zero.
+;; get_global_id(dim): Here we assume a specific implementation using select.
+; CHECK-DAG: %[[#TMP2:]] = OpLoad %[[#VEC]] %[[#GET_GLOBAL_ID]]
+; CHECK-DAG: %[[#TMP3:]] = OpVectorExtractDynamic %[[#I64]] %[[#TMP2]] %[[#DIM]]
+; CHECK-DAG: %[[#COND:]] = OpULessThan %[[#BOOL]] %[[#DIM]] %[[#THREE]]
+; CHECK:     %[[#UNKNOWN:]] = OpSelect %[[#I64]] %[[#COND]] %[[#TMP3]] %[[#ZERO]]
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpMin.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpMin.ll
new file mode 100644
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpMin.ll
@@ -0,0 +1,16 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: %[[#SetInstID:]] = OpExtInstImport "OpenCL.std"
+; CHECK-SPIRV: %[[#IntTypeID:]] = OpTypeInt 32 [[#]]
+; CHECK-SPIRV: %[[#Int2TypeID:]] = OpTypeVector %[[#IntTypeID]] 2
+; CHECK-SPIRV: %[[#CompositeID:]] = OpCompositeInsert %[[#Int2TypeID]] %[[#]] %[[#]] [[#]]
+; CHECK-SPIRV: %[[#ShuffleID:]] = OpVectorShuffle %[[#Int2TypeID]] %[[#CompositeID]] %[[#]] [[#]] [[#]]
+; CHECK-SPIRV: %[[#]] = OpExtInst %[[#Int2TypeID]] %[[#SetInstID]] s_min %[[#]] %[[#ShuffleID]]
+
+define spir_kernel void @test() {
+entry:
+  %call = tail call spir_func <2 x i32> @_Z3minDv2_ii(<2 x i32> <i32 1, i32 10>, i32 5) #2
+  ret void
+}
+
+declare spir_func <2 x i32> @_Z3minDv2_ii(<2 x i32>, i32)