[llvm] 3544d20 - [SPIRV] add IR regularization pass
Ilia Diachkov via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 15 05:10:22 PDT 2022
Author: Ilia Diachkov
Date: 2022-09-15T15:53:44+03:00
New Revision: 3544d200d982cb6ce9654130751f0e24d2a61a49
URL: https://github.com/llvm/llvm-project/commit/3544d200d982cb6ce9654130751f0e24d2a61a49
DIFF: https://github.com/llvm/llvm-project/commit/3544d200d982cb6ce9654130751f0e24d2a61a49.diff
LOG: [SPIRV] add IR regularization pass
The patch adds the regularization pass that prepare LLVM IR for
the IR translation. It also contains following changes:
- reduce indentation, make getNonParametrizedType, getSamplerType,
getPipeType, getImageType, getSampledImageType static in SPIRVBuiltins,
- rename mayBeOclOrSpirvBuiltin to getOclOrSpirvBuiltinDemangledName,
- 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,
- add capability processing for OpBuildNDRange in SPIRVModuleAnalysis,
- don't correct types of registers holding constants and used in
G_ADDRSPACE_CAST (SPIRVPreLegalizer.cpp),
- lower memset/bswap intrinsics to functions in SPIRVPrepareFunctions,
- change TargetLoweringObjectFileELF to SPIRVTargetObjectFile
in SPIRVTargetMachine.cpp,
- correct comments.
5 LIT tests are added to show the improvement.
Differential Revision: https://reviews.llvm.org/D133253
Co-authored-by: Aleksandr Bezzubikov <zuban32s at gmail.com>
Co-authored-by: Michal Paszkowski <michal.paszkowski at outlook.com>
Co-authored-by: Andrey Tretyakov <andrey1.tretyakov at intel.com>
Co-authored-by: Konrad Trifunovic <konrad.trifunovic at intel.com>
Added:
llvm/lib/Target/SPIRV/SPIRVRegularizer.cpp
llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll
llvm/test/CodeGen/SPIRV/lshr-constexpr.ll
llvm/test/CodeGen/SPIRV/opencl/get_global_id.ll
llvm/test/CodeGen/SPIRV/transcoding/OpMin.ll
Modified:
llvm/lib/Target/SPIRV/CMakeLists.txt
llvm/lib/Target/SPIRV/SPIRV.h
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
llvm/lib/Target/SPIRV/SPIRVISelLowering.h
llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
llvm/lib/Target/SPIRV/SPIRVUtils.cpp
llvm/lib/Target/SPIRV/SPIRVUtils.h
Removed:
################################################################################
diff --git a/llvm/lib/Target/SPIRV/CMakeLists.txt b/llvm/lib/Target/SPIRV/CMakeLists.txt
index 4e0ac874b80e4..039b6f84805ad 100644
--- a/llvm/lib/Target/SPIRV/CMakeLists.txt
+++ b/llvm/lib/Target/SPIRV/CMakeLists.txt
@@ -30,6 +30,7 @@ add_llvm_target(SPIRVCodeGen
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
index 5a7f2e51afb8c..20834c5476468 100644
--- a/llvm/lib/Target/SPIRV/SPIRV.h
+++ b/llvm/lib/Target/SPIRV/SPIRV.h
@@ -20,6 +20,7 @@ class InstructionSelector;
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
index 5655662bd766d..c4ccfee7aaa5e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -809,7 +809,7 @@ static bool generateGroupInst(const SPIRV::IncomingCall *Call,
}
// 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 @@ using namespace ImageFormat;
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 @@ lookupOrParseBuiltinImageType(StringRef Name) {
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 @@ lookupOrParseBuiltinPipeType(StringRef Name) {
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 @@ SPIRVType *getPipeType(const StructType *OpaqueType,
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 @@ SPIRVType *getImageType(const StructType *OpaqueType,
: 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
index ce26a9a497626..774941d1f17ea 100644
--- a/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
@@ -286,7 +286,7 @@ bool SPIRVCallLowering::lowerCall(MachineIRBuilder &MIRBuilder,
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
index 3131794eed2d0..bbb86ce5595a5 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -544,26 +544,6 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeStruct(const StructType *Ty,
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 @@ SPIRVType *SPIRVGlobalRegistry::getOrCreateSpecialType(
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 @@ Register SPIRVGlobalRegistry::getSPIRVTypeID(const SPIRVType *SpirvType) const {
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 @@ SPIRVType *SPIRVGlobalRegistry::restOfCreateSPIRVType(
// 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 @@ 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);
diff --git a/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp b/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
index 66ff51c912b0a..c226b1ac7528f 100644
--- 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 @@ MVT SPIRVTargetLowering::getRegisterTypeForCallingConv(LLVMContext &Context,
}
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/SPIRVISelLowering.h b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
index bee9220f52486..f317b26207195 100644
--- a/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
+++ b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
@@ -41,6 +41,9 @@ class SPIRVTargetLowering : public TargetLowering {
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/SPIRVInstrFormats.td b/llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
index c78c8ee115901..9451583a5fa85 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
+++ b/llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
@@ -28,4 +28,5 @@ class Op<bits<16> Opcode, dag outs, dag ins, string asmstr, list<dag> pattern =
// 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
index 3d12077cd2c03..b0028f8c80a40 100644
--- a/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
@@ -145,6 +145,9 @@ SPIRVLegalizerInfo::SPIRVLegalizerInfo(const SPIRVSubtarget &ST) {
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 @@ SPIRVLegalizerInfo::SPIRVLegalizerInfo(const SPIRVSubtarget &ST) {
// 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
index 1d57398e9b6f3..bc3f234960f5b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -751,6 +751,7 @@ void addInstrRequirements(const MachineInstr &MI,
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
index 370a2e2aef06d..7c24d9557711f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
@@ -369,11 +369,19 @@ static void processInstrsWithTypeFolding(MachineFunction &MF,
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
index 13c3c12c1b417..262d60ac2cad3 100644
--- 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 @@ static Function *getOrCreateFunction(Module *M, Type *RetTy,
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 @@ static void substituteIntrinsicCalls(Module *M, Function *F) {
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
index 0000000000000..d7c66addf25cb
--- /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
index 7f5f14dc3ce8d..3d5d58134c89f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
@@ -70,7 +70,7 @@ SPIRVTargetMachine::SPIRVTargetMachine(const Target &T, const Triple &TT,
: 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 @@ TargetPassConfig *SPIRVTargetMachine::createPassConfig(PassManagerBase &PM) {
void SPIRVPassConfig::addIRPasses() {
TargetPassConfig::addIRPasses();
+ addPass(createSPIRVRegularizerPass());
addPass(createSPIRVPrepareFunctionsPass());
}
@@ -159,13 +160,13 @@ 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;
@@ -183,6 +184,7 @@ class SPIRVInstructionSelect : public InstructionSelect {
};
} // namespace
+// Add the custom SPIRVInstructionSelect from above.
bool SPIRVPassConfig::addGlobalInstructionSelect() {
addPass(new SPIRVInstructionSelect());
return false;
diff --git a/llvm/lib/Target/SPIRV/SPIRVUtils.cpp b/llvm/lib/Target/SPIRV/SPIRVUtils.cpp
index 85ce9c3b67847..9eeccb5a39370 100644
--- a/llvm/lib/Target/SPIRV/SPIRVUtils.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVUtils.cpp
@@ -289,7 +289,7 @@ static bool isNonMangledOCLBuiltin(StringRef Name) {
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 @@ std::string mayBeOclOrSpirvBuiltin(StringRef Name) {
.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/lib/Target/SPIRV/SPIRVUtils.h b/llvm/lib/Target/SPIRV/SPIRVUtils.h
index ffec7d78c0adb..5006caedd81d3 100644
--- a/llvm/lib/Target/SPIRV/SPIRVUtils.h
+++ b/llvm/lib/Target/SPIRV/SPIRVUtils.h
@@ -84,8 +84,11 @@ 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
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll
new file mode 100644
index 0000000000000..3f2ab9fa7190b
--- /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
index 0000000000000..0ccace31781ca
--- /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
index 0000000000000..c435cb25ef5c2
--- /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
index 0000000000000..7512f29233722
--- /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
index 0000000000000..5cc3ea01e5191
--- /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)
More information about the llvm-commits
mailing list