OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 2
Code Issues Packages Projects Releases Wiki Activity

[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@gmail.com>
Co-authored-by: Michal Paszkowski <michal.paszkowski@outlook.com>
Co-authored-by: Andrey Tretyakov <andrey1.tretyakov@intel.com>
Co-authored-by: Konrad Trifunovic <konrad.trifunovic@intel.com>
This commit is contained in:
Ilia Diachkov 2022-09-14 11:51:03 +03:00
parent f0c234d2a6
commit 3544d200d9
21 changed files with 719 additions and 108 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
llvm/lib/Target/SPIRV/CMakeLists.txt
View File

@ -30,6 +30,7 @@ add_llvm_target(SPIRVCodeGen
SPIRVPrepareFunctions.cpp
SPIRVRegisterBankInfo.cpp
SPIRVRegisterInfo.cpp
SPIRVRegularizer.cpp
SPIRVSubtarget.cpp
SPIRVTargetMachine.cpp
SPIRVUtils.cpp

1
llvm/lib/Target/SPIRV/SPIRV.h
View File

@ -20,6 +20,7 @@ class InstructionSelector;
class RegisterBankInfo;
ModulePass *createSPIRVPrepareFunctionsPass();
FunctionPass *createSPIRVRegularizerPass();
FunctionPass *createSPIRVPreLegalizerPass();
FunctionPass *createSPIRVEmitIntrinsicsPass(SPIRVTargetMachine *TM);
InstructionSelector *

141
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
View File

@ -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");
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.");
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");
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();

2
llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
View File

@ -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() &&

28
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
View File

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

29
llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
View File

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

3
llvm/lib/Target/SPIRV/SPIRVISelLowering.h
View File

@ -41,6 +41,9 @@ public:
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

1
llvm/lib/Target/SPIRV/SPIRVInstrFormats.td
View File

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

7
llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
View File

@ -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,

1
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
View File

@ -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:

10
llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
View File

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

71
llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
View File

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

249
llvm/lib/Target/SPIRV/SPIRVRegularizer.cpp Normal file
View File

@ -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();
}

8
llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp
View File

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

22
llvm/lib/Target/SPIRV/SPIRVUtils.cpp
View File

@ -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

9
llvm/lib/Target/SPIRV/SPIRVUtils.h
View File

@ -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

74
llvm/test/CodeGen/SPIRV/llvm-intrinsics/bswap.ll Normal file
View File

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

83
llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll Normal file
View File

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

18
llvm/test/CodeGen/SPIRV/lshr-constexpr.ll Normal file
View File

@ -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
}

53
llvm/test/CodeGen/SPIRV/opencl/get_global_id.ll Normal file
View File

@ -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]]

16
llvm/test/CodeGen/SPIRV/transcoding/OpMin.ll Normal file
View File

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