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