[SPIRV] add IR regularization pass

[iliya-diyachkov]

Summary

The patch adds the regularization pass that prepare LLVM IR for the IR translation. It also implements a few small improvements and fixes some bugs:

• reduce code indentation, make getNonParametrizedType, getSamplerType, getPipeType, getImageType, getSampledImageType static (SPIRVBuiltins.cpp),
• rename mayBeOclOrSpirvBuiltin to getOclOrSpirvBuiltinDemangledName (SPIRVUtils.cpp),
• move isOpenCLBuiltinType, isSPIRVBuiltinType, isSpecialType from SPIRVGlobalRegistry.cpp to SPIRVUtils.cpp, renaming isSpecialType to isSpecialOpaqueType,
• implment getTgtMemIntrinsic() in SPIRVISelLowering,
• add hasSideEffects = 0 in Pseudo (SPIRVInstrFormats.td),
• add legalization rule for G_MEMSET, correct G_BRCOND rule (SPIRVLegalizerInfo.cpp),
• add capability processing for OpBuildNDRange (SPIRVModuleAnalysis.cpp),
• don't correct type of register holding constant and used in G_ADDRSPACE_CAST (SPIRVPreLegalizer.cpp),
• lower memset/bswap intrinsics to functions (SPIRVPrepareFunctions.cpp),
• change TLOF from TargetLoweringObjectFileELF to SPIRVTargetObjectFile (SPIRVTargetMachine.cpp),
• correct comments.

5 LIT tests are added to show the improvement.

[nlopes]

llvm/lib/Target/SPIRV/SPIRVLowerConstExpr.cpp

150 ↗

(On Diff #457784)

Please use PoisonValue as a placeholder.

[zuban32]

Could you pls specify a list of these small bug fixes, at the very high level at least? An ideal variant would be to not mix them up with this regularization stuff at all.

llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp

118

Alright, this doesn't look like a small change to me. Can this be moved into a separate patch, current description doesn't say anything specific about this anyway?

[iliya-diyachkov]

Two issues are resolved in the new patch.

Could you pls specify a list of these small bug fixes, at the very high level at least?

Sure, I will add it later today.

[andreytr]

Suggest small style test fixes.

llvm/test/CodeGen/SPIRV/lshr-constexpr.ll

6–12

llvm/test/CodeGen/SPIRV/opencl/get_global_id.ll

42–43

46–47

54

llvm/test/CodeGen/SPIRV/transcoding/OpMin.ll

5

7–8

[iliya-diyachkov]

Fix tests.

[mpaszkowski]

Thanks for the patch and fixing the issues/comments! LGTM! I think this could be committed now.

[zuban32]

llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp

1738

Pls remove tabs

1747

Tabs

1784

Tabs

[iliya-diyachkov]

These are not tabs, but spaces. Phabricator shows in this way that the patch contains an indentation change. In this particular case, I inserted "static" and needed to add spaces before parameters.

[zuban32]

OK, LGTM then

[nlopes]

You committed the patch with undefvalue. Could you please use poison value instead? Thank you!

llvm/lib/Target/SPIRV/CMakeLists.txt

add_llvm_target(SPIRVCodeGen
   SPIRVISelLowering.cpp
   SPIRVLegalizerInfo.cpp
   SPIRVMCInstLower.cpp
   SPIRVModuleAnalysis.cpp
   SPIRVPreLegalizer.cpp
   SPIRVPrepareFunctions.cpp
   SPIRVRegisterBankInfo.cpp
   SPIRVRegisterInfo.cpp
+  SPIRVRegularizer.cpp
   SPIRVSubtarget.cpp
   SPIRVTargetMachine.cpp
   SPIRVUtils.cpp
 
   LINK_COMPONENTS
   Analysis
   AsmPrinter
   CodeGen

llvm/lib/Target/SPIRV/SPIRV.h

 
 namespace llvm {
 class SPIRVTargetMachine;
 class SPIRVSubtarget;
 class InstructionSelector;
 class RegisterBankInfo;
 
 ModulePass *createSPIRVPrepareFunctionsPass();
+FunctionPass *createSPIRVRegularizerPass();
 FunctionPass *createSPIRVPreLegalizerPass();
 FunctionPass *createSPIRVEmitIntrinsicsPass(SPIRVTargetMachine *TM);
 InstructionSelector *
 createSPIRVInstructionSelector(const SPIRVTargetMachine &TM,
                                const SPIRVSubtarget &Subtarget,
                                const RegisterBankInfo &RBI);
 
 void initializeSPIRVModuleAnalysisPass(PassRegistry &);
 void initializeSPIRVPreLegalizerPass(PassRegistry &);
 void initializeSPIRVEmitIntrinsicsPass(PassRegistry &);
 } // namespace llvm
 
 #endif // LLVM_LIB_TARGET_SPIRV_SPIRV_H

llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp

static bool generateGroupInst(const SPIRV::IncomingCall *Call,
   // Build select instruction.
   if (HasBoolReturnTy)
     buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
                     Call->ReturnType, GR);
   return true;
 }
 
 // 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
 // bitwidth.
 //
 // For a constant index >= 3 we generate:
 //  %res = OpConstant %SizeT 0
 //

 //===----------------------------------------------------------------------===//
 // Misc functions for parsing builtin types and looking up implementation
 // details in TableGenerated tables.
 //===----------------------------------------------------------------------===//
 
 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;
 
 static std::unique_ptr<const SPIRV::ImageType>
 lookupOrParseBuiltinImageType(StringRef Name) {
   if (Name.startswith("opencl.")) {
     // Lookup OpenCL builtin image type lowering details in TableGen records.
     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");
+  if (!Name.startswith("spirv."))
+    llvm_unreachable("Unknown builtin image type name/literal");
 
-    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");
-}
 
 static std::unique_ptr<const SPIRV::PipeType>
 lookupOrParseBuiltinPipeType(StringRef Name) {
   if (Name.startswith("opencl.")) {
     // Lookup OpenCL builtin pipe type lowering details in TableGen records.
     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");
 
 //===----------------------------------------------------------------------===//
 // Implementation functions for builtin types.
 //===----------------------------------------------------------------------===//
 
-SPIRVType *getNonParametrizedType(const StructType *OpaqueType,
+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,
+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 =
       lookupOrParseBuiltinPipeType(OpaqueType->getName());
   // Create or get an existing type from GlobalRegistry.
   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,
   // Lookup image type lowering details in TableGen records or parse the
   // name/literal for details.
   std::unique_ptr<const SPIRV::ImageType> Record =
       lookupOrParseBuiltinImageType(OpaqueType->getName());
 
   SPIRVType *SampledType =
       GR->getOrCreateSPIRVTypeByName(Record.get()->SampledType, MIRBuilder);
   return GR->getOrCreateOpTypeImage(
       MIRBuilder, SampledType, Record.get()->Dimensionality,
       Record.get()->Depth, Record.get()->Arrayed, Record.get()->Multisampled,
       Record.get()->Sampled, Record.get()->Format,
       AccessQual == SPIRV::AccessQualifier::WriteOnly
           ? SPIRV::AccessQualifier::WriteOnly
           : Record.get()->Qualifier);
 }
 
-SPIRVType *getSampledImageType(const StructType *OpaqueType,
+static SPIRVType *getSampledImageType(const StructType *OpaqueType,
                                       MachineIRBuilder &MIRBuilder,
                                       SPIRVGlobalRegistry *GR) {
   StringRef TypeParametersString =
       OpaqueType->getName().substr(strlen("spirv.SampledImage."));
   LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
   Type *ImageOpaqueType = StructType::getTypeByName(
       Context, "spirv.Image." + TypeParametersString.str());
   SPIRVType *TargetImageType =
       GR->getOrCreateSPIRVType(ImageOpaqueType, MIRBuilder);
   return GR->getOrCreateOpTypeSampledImage(TargetImageType, MIRBuilder);

llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp

if (Const) {
       else
         ArgTypes[Const->getSExtValue()] = CMeta->getType();
     }
   }
 
   return FunctionType::get(RetTy, ArgTypes, F.isVarArg());
 }
 
 bool SPIRVCallLowering::lowerFormalArguments(MachineIRBuilder &MIRBuilder,
                                              const Function &F,
                                              ArrayRef<ArrayRef<Register>> VRegs,
                                              FunctionLoweringInfo &FLI) const {
   assert(GR && "Must initialize the SPIRV type registry before lowering args.");
   GR->setCurrentFunc(MIRBuilder.getMF());
 
   // Assign types and names to all args, and store their types for later.
   FunctionType *FTy = getOriginalFunctionType(F);
if (Info.Callee.isGlobal()) {
     if (CF == nullptr)
       return false;
     FTy = getOriginalFunctionType(*CF);
   }
 
   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() &&
       ST->canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
     const Type *OrigRetTy = Info.OrigRet.Ty;
     if (FTy)
       OrigRetTy = FTy->getReturnType();
     SmallVector<Register, 8> ArgVRegs;

llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp

for (const auto &Ty : FieldTypes)
     MIB.addUse(Ty);
   if (Ty->hasName())
     buildOpName(ResVReg, Ty->getName(), MIRBuilder);
   if (Ty->isPacked())
     buildOpDecorate(ResVReg, MIRBuilder, SPIRV::Decoration::CPacked, {});
   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) {
   // Some OpenCL and SPIRV builtins like image2d_t are passed in as
   // pointers, but should be treated as custom types like OpTypeImage.
   if (auto PType = dyn_cast<PointerType>(Ty)) {
     assert(!PType->isOpaque());
     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);
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOpTypePointer(
     SPIRV::StorageClass::StorageClass SC, SPIRVType *ElemType,
     MachineIRBuilder &MIRBuilder, Register Reg) {
   if (!Reg.isValid())
     Reg = createTypeVReg(MIRBuilder);
Register SPIRVGlobalRegistry::getSPIRVTypeID(const SPIRVType *SpirvType) const {
   if (SpirvType->getOpcode() == SPIRV::OpTypeForwardPointer)
     return SpirvType->uses().begin()->getReg();
   return SpirvType->defs().begin()->getReg();
 }
 
 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);
   if (t != TypeToSPIRVTypeMap.end()) {
     auto tt = t->second.find(&MIRBuilder.getMF());
     if (tt != t->second.end())
       return getSPIRVTypeForVReg(tt->second);
   }
SPIRVType *SPIRVGlobalRegistry::restOfCreateSPIRVType(
   SPIRVType *SpirvType = createSPIRVType(Ty, MIRBuilder, AccessQual, EmitIR);
   TypesInProcessing.erase(Ty);
   VRegToTypeMap[&MIRBuilder.getMF()][getSPIRVTypeID(SpirvType)] = SpirvType;
   SPIRVToLLVMType[SpirvType] = Ty;
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   // 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;
 }
 
 SPIRVType *SPIRVGlobalRegistry::getSPIRVTypeForVReg(Register VReg) const {
   auto t = VRegToTypeMap.find(CurMF);
   if (t != VRegToTypeMap.end()) {
     auto tt = t->second.find(VReg);
     if (tt != t->second.end())
       return tt->second;
   }
   return nullptr;
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVType(
     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);
   // Create normal pointer types for the corresponding OpTypeForwardPointers.
   for (auto &CU : ForwardPointerTypes) {
     const Type *Ty2 = CU.first;
     SPIRVType *STy2 = CU.second;
     if ((Reg = DT.find(Ty2, &MIRBuilder.getMF())).isValid())

llvm/lib/Target/SPIRV/SPIRVISelLowering.h

public:
   MVT getVectorIdxTy(const DataLayout &DL) const override {
     return MVT::getIntegerVT(32);
   }
   unsigned getNumRegistersForCallingConv(LLVMContext &Context,
                                          CallingConv::ID CC,
                                          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
 
 #endif // LLVM_LIB_TARGET_SPIRV_SPIRVISELLOWERING_H

llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp

 //===- SPIRVISelLowering.cpp - SPIR-V DAG Lowering Impl ---------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the SPIRVTargetLowering class.
 //
 //===----------------------------------------------------------------------===//
 
 #include "SPIRVISelLowering.h"
 #include "SPIRV.h"
+#include "llvm/IR/IntrinsicsSPIRV.h"
 
 #define DEBUG_TYPE "spirv-lower"
 
 using namespace llvm;
 
 unsigned SPIRVTargetLowering::getNumRegistersForCallingConv(
     LLVMContext &Context, CallingConv::ID CC, EVT VT) const {
   // This code avoids CallLowering fail inside getVectorTypeBreakdown
MVT SPIRVTargetLowering::getRegisterTypeForCallingConv(LLVMContext &Context,
   if (VT.isVector() && VT.getVectorNumElements() == 3) {
     if (VT.getVectorElementType() == MVT::i1)
       return MVT::v4i1;
     else if (VT.getVectorElementType() == MVT::i8)
       return MVT::v4i8;
   }
   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;
+}

llvm/lib/Target/SPIRV/SPIRVInstrFormats.td

class Op<bits<16> Opcode, dag outs, dag ins, string asmstr, list<dag> pattern = []>
   dag InOperandList = ins;
   let AsmString = asmstr;
   let Pattern = pattern;
 }
 
 // Pseudo instructions
 class Pseudo<dag outs, dag ins> : Op<0, outs, ins, ""> {
   let isPseudo = 1;
+  let hasSideEffects = 0;
 }

llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp

SPIRVLegalizerInfo::SPIRVLegalizerInfo(const SPIRVSubtarget &ST) {
   getActionDefinitionsBuilder(G_GLOBAL_VALUE).alwaysLegal();
 
   // TODO: add proper rules for vectors legalization.
   getActionDefinitionsBuilder({G_BUILD_VECTOR, G_SHUFFLE_VECTOR}).alwaysLegal();
 
   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);
 
   getActionDefinitionsBuilder({G_LOAD, G_STORE}).legalIf(typeInSet(1, allPtrs));
 
   getActionDefinitionsBuilder(G_BITREVERSE).legalFor(allFloatScalarsAndVectors);
 
   getActionDefinitionsBuilder(G_FMA).legalFor(allFloatScalarsAndVectors);
SPIRVLegalizerInfo::SPIRVLegalizerInfo(const SPIRVSubtarget &ST) {
 
   // FP conversions.
   getActionDefinitionsBuilder({G_FPTRUNC, G_FPEXT})
       .legalForCartesianProduct(allFloatScalarsAndVectors);
 
   // 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,
                                G_FEXP2,
                                G_FLOG,
                                G_FLOG2,
                                G_FABS,
                                G_FMINNUM,

llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp

case SPIRV::OpTypeEvent:
     Reqs.addCapability(SPIRV::Capability::Kernel);
     break;
   case SPIRV::OpTypePipe:
   case SPIRV::OpTypeReserveId:
     Reqs.addCapability(SPIRV::Capability::Pipes);
     break;
   case SPIRV::OpTypeDeviceEvent:
   case SPIRV::OpTypeQueue:
+  case SPIRV::OpBuildNDRange:
     Reqs.addCapability(SPIRV::Capability::DeviceEnqueue);
     break;
   case SPIRV::OpDecorate:
   case SPIRV::OpDecorateId:
   case SPIRV::OpDecorateString:
     addOpDecorateReqs(MI, 1, Reqs, ST);
     break;
   case SPIRV::OpMemberDecorate:

llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp

static void processInstrsWithTypeFolding(MachineFunction &MF,
   for (MachineBasicBlock &MBB : MF) {
     for (MachineInstr &MI : MBB) {
       // We need to rewrite dst types for ASSIGN_TYPE instrs to be able
       // to perform tblgen'erated selection and we can't do that on Legalizer
       // as it operates on gMIR only.
       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));
     }
   }
 }
 
 static void processSwitches(MachineFunction &MF, SPIRVGlobalRegistry *GR,
                             MachineIRBuilder MIB) {
   DenseMap<Register, SmallDenseMap<uint64_t, MachineBasicBlock *>>

llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp

 // NOTE: this pass is a module-level one due to the necessity to modify
 // GVs/functions.
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 #include "llvm/Transforms/Utils/LowerMemIntrinsics.h"
 
 using namespace llvm;
 
 namespace llvm {
if (F && F->getFunctionType() == FT)
     return F;
   Function *NewF = Function::Create(FT, GlobalValue::ExternalLinkage, Name, M);
   if (F)
     NewF->setDSOLocal(F->isDSOLocal());
   NewF->setCallingConv(CallingConv::SPIR_FUNC);
   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
   // function.
   // Generate LLVM IR for  i* @spirv.llvm_fsh?_i* (i* %a, i* %b, i* %c)
   FunctionType *FSHFuncTy = FSHIntrinsic->getFunctionType();
   Type *FSHRetTy = FSHFuncTy->getReturnType();
   const std::string FuncName = lowerLLVMIntrinsicName(FSHIntrinsic);
for (Instruction &I : BB) {
       auto Call = dyn_cast<CallInst>(&I);
       if (!Call)
         continue;
       Call->setTailCall(false);
       Function *CF = Call->getCalledFunction();
       if (!CF || !CF->isIntrinsic())
         continue;
       auto *II = cast<IntrinsicInst>(Call);
-      if (II->getIntrinsicID() == Intrinsic::fshl ||
+      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);
     }
   }
 }
 
 bool SPIRVPrepareFunctions::runOnModule(Module &M) {

llvm/lib/Target/SPIRV/SPIRVRegularizer.cpp

+//===-- 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();
+}

llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp

SPIRVTargetMachine::SPIRVTargetMachine(const Target &T, const Triple &TT,
                                        StringRef CPU, StringRef FS,
                                        const TargetOptions &Options,
                                        Optional<Reloc::Model> RM,
                                        Optional<CodeModel::Model> CM,
                                        CodeGenOpt::Level OL, bool JIT)
     : 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);
   setFastISel(false);
   setO0WantsFastISel(false);
   setRequiresStructuredCFG(false);
 }
 

 }
 
 TargetPassConfig *SPIRVTargetMachine::createPassConfig(PassManagerBase &PM) {
   return new SPIRVPassConfig(*this, PM);
 }
 
 void SPIRVPassConfig::addIRPasses() {
   TargetPassConfig::addIRPasses();
+  addPass(createSPIRVRegularizerPass());
   addPass(createSPIRVPrepareFunctionsPass());
 }
 
 void SPIRVPassConfig::addISelPrepare() {
   addPass(createSPIRVEmitIntrinsicsPass(&getTM<SPIRVTargetMachine>()));
   TargetPassConfig::addISelPrepare();
 }
 
 bool SPIRVPassConfig::addIRTranslator() {
   addPass(new IRTranslator(getOptLevel()));
   return false;
 }
 
 void SPIRVPassConfig::addPreLegalizeMachineIR() {
   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;
 }
 
 namespace {
 // A custom subclass of InstructionSelect, which is mostly the same except from
 // not requiring RegBankSelect to occur previously.
 class SPIRVInstructionSelect : public InstructionSelect {
   // We don't use register banks, so unset the requirement for them
   MachineFunctionProperties getRequiredProperties() const override {
     return InstructionSelect::getRequiredProperties().reset(
         MachineFunctionProperties::Property::RegBankSelected);
   }
 };
 } // namespace
 
+// Add the custom SPIRVInstructionSelect from above.
 bool SPIRVPassConfig::addGlobalInstructionSelect() {
   addPass(new SPIRVInstructionSelect());
   return false;
 }

llvm/lib/Target/SPIRV/SPIRVUtils.h

 uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI);
 
 // Check if MI is a SPIR-V specific intrinsic call.
 bool isSpvIntrinsic(MachineInstr &MI, Intrinsic::ID IntrinsicID);
 
 // 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

llvm/lib/Target/SPIRV/SPIRVUtils.cpp

static bool isNonMangledOCLBuiltin(StringRef Name) {
   if (!Name.startswith("__"))
     return false;
 
   return isEnqueueKernelBI(Name) || isKernelQueryBI(Name) ||
          isPipeOrAddressSpaceCastBI(Name.drop_front(2)) ||
          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");
 
   if (!IsNonMangledOCL && !IsNonMangledSPIRV && !IsMangled)
     return std::string();
 
   // Try to use the itanium demangler.
if (Name.substr(NameSpaceStart, 11) != "2cl7__spirv")
       return std::string();
     DemangledNameLenStart = NameSpaceStart + 11;
   }
   Start = Name.find_first_not_of("0123456789", DemangledNameLenStart);
   Name.substr(DemangledNameLenStart, Start - DemangledNameLenStart)
       .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

llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll

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

llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll

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

llvm/test/CodeGen/SPIRV/lshr-constexpr.ll

+; 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
+}

llvm/test/CodeGen/SPIRV/opencl/get_global_id.ll

+; 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]]

llvm/test/CodeGen/SPIRV/transcoding/OpMin.ll

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