From 7a0767f2a83c49003d66edfbc70c9cc6cd374ab7 Mon Sep 17 00:00:00 2001 From: Konstantin Vladimirov Date: Wed, 18 Mar 2020 21:15:13 +0300 Subject: [PATCH] Implementation of SPV_INTEL_inline_assembly extension Extension is published as https://github.com/intel/llvm/pull/1290 Co-Authored-By: Nikita Rudenko Co-Authored-By: Anton Sidorenko Co-Authored-By: Alexey Sachkov --- include/LLVMSPIRVExtensions.inc | 1 + lib/SPIRV/SPIRVReader.cpp | 27 +++++ lib/SPIRV/SPIRVReader.h | 3 + lib/SPIRV/SPIRVWriter.cpp | 37 +++++++ lib/SPIRV/SPIRVWriter.h | 2 + lib/SPIRV/libSPIRV/SPIRVAsm.h | 142 ++++++++++++++++++++++++++ lib/SPIRV/libSPIRV/SPIRVEntry.cpp | 1 + lib/SPIRV/libSPIRV/SPIRVEnum.h | 1 + lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h | 3 + lib/SPIRV/libSPIRV/SPIRVModule.cpp | 56 +++++++++- lib/SPIRV/libSPIRV/SPIRVModule.h | 9 ++ lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 2 + lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h | 3 + lib/SPIRV/libSPIRV/spirv.hpp | 139 +++++++++++++------------ test/inline_asm_basic.cl | 17 +++ test/inline_asm_clobbers.cl | 94 +++++++++++++++++ test/inline_asm_constraints.cl | 100 ++++++++++++++++++ 17 files changed, 568 insertions(+), 69 deletions(-) create mode 100644 lib/SPIRV/libSPIRV/SPIRVAsm.h create mode 100644 test/inline_asm_basic.cl create mode 100644 test/inline_asm_clobbers.cl create mode 100644 test/inline_asm_constraints.cl diff --git a/include/LLVMSPIRVExtensions.inc b/include/LLVMSPIRVExtensions.inc index 7e6b8a7a55..6071892c8b 100644 --- a/include/LLVMSPIRVExtensions.inc +++ b/include/LLVMSPIRVExtensions.inc @@ -15,3 +15,4 @@ EXT(SPV_INTEL_blocking_pipes) EXT(SPV_INTEL_function_pointers) EXT(SPV_INTEL_kernel_attributes) EXT(SPV_INTEL_io_pipes) +EXT(SPV_INTEL_inline_assembly) diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp index 692bb7e010..3283646f30 100644 --- a/lib/SPIRV/SPIRVReader.cpp +++ b/lib/SPIRV/SPIRVReader.cpp @@ -38,6 +38,7 @@ //===----------------------------------------------------------------------===// #include "SPIRVReader.h" #include "OCLUtil.h" +#include "SPIRVAsm.h" #include "SPIRVBasicBlock.h" #include "SPIRVExtInst.h" #include "SPIRVFunction.h" @@ -57,6 +58,7 @@ #include "llvm/IR/DerivedTypes.h" #include "llvm/IR/Dominators.h" #include "llvm/IR/IRBuilder.h" +#include "llvm/IR/InlineAsm.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/IntrinsicInst.h" #include "llvm/IR/LegacyPassManager.h" @@ -1516,6 +1518,9 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, case OpFunction: return mapValue(BV, transFunction(static_cast(BV))); + case OpAsmINTEL: + return mapValue(BV, transAsmINTEL(static_cast(BV))); + case OpLabel: return mapValue(BV, BasicBlock::Create(*Context, BV->getName(), F)); @@ -2112,6 +2117,10 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, return mapValue(BV, Call); } + case OpAsmCallINTEL: + return mapValue( + BV, transAsmCallINTEL(static_cast(BV), F, BB)); + case OpFunctionPointerCallINTEL: { SPIRVFunctionPointerCallINTEL *BC = static_cast(BV); @@ -2380,6 +2389,24 @@ Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) { return F; } +Value *SPIRVToLLVM::transAsmINTEL(SPIRVAsmINTEL *BA) { + assert(BA); + bool HasSideEffect = BA->hasDecorate(DecorationSideEffectsINTEL); + return InlineAsm::get( + cast(transType(BA->getFunctionType())), + BA->getInstructions(), BA->getConstraints(), HasSideEffect, + /* IsAlignStack */ false, InlineAsm::AsmDialect::AD_ATT); +} + +CallInst *SPIRVToLLVM::transAsmCallINTEL(SPIRVAsmCallINTEL *BI, Function *F, + BasicBlock *BB) { + assert(BI); + auto *IA = cast(transValue(BI->getAsm(), F, BB)); + auto Args = transValue(BM->getValues(BI->getArguments()), F, BB); + return CallInst::Create(cast(IA->getFunctionType()), IA, Args, + BI->getName(), BB); +} + /// LLVM convert builtin functions is translated to two instructions: /// y = i32 islessgreater(float x, float z) -> /// y = i32 ZExt(bool LessOrGreater(float x, float z)) diff --git a/lib/SPIRV/SPIRVReader.h b/lib/SPIRV/SPIRVReader.h index 89d27c09d4..42ebdcf072 100644 --- a/lib/SPIRV/SPIRVReader.h +++ b/lib/SPIRV/SPIRVReader.h @@ -103,6 +103,9 @@ class SPIRVToLLVM { Instruction *transSGSizeQueryBI(SPIRVInstruction *BI, BasicBlock *BB); bool transFPContractMetadata(); bool transKernelMetadata(); + Value *transAsmINTEL(SPIRVAsmINTEL *BA); + CallInst *transAsmCallINTEL(SPIRVAsmCallINTEL *BI, Function *F, + BasicBlock *BB); bool transNonTemporalMetadata(Instruction *I); bool transSourceLanguage(); bool transSourceExtension(); diff --git a/lib/SPIRV/SPIRVWriter.cpp b/lib/SPIRV/SPIRVWriter.cpp index f7901ff541..8b613cd93b 100644 --- a/lib/SPIRV/SPIRVWriter.cpp +++ b/lib/SPIRV/SPIRVWriter.cpp @@ -40,6 +40,7 @@ #include "SPIRVWriter.h" #include "LLVMToSPIRVDbgTran.h" +#include "SPIRVAsm.h" #include "SPIRVBasicBlock.h" #include "SPIRVEntry.h" #include "SPIRVEnum.h" @@ -60,6 +61,7 @@ #include "llvm/IR/Constants.h" #include "llvm/IR/DerivedTypes.h" #include "llvm/IR/Function.h" +#include "llvm/IR/InlineAsm.h" #include "llvm/IR/InstrTypes.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/IntrinsicInst.h" @@ -1265,6 +1267,10 @@ SPIRVValue *LLVMToSPIRV::transValueWithoutDecoration(Value *V, return BV ? mapValue(V, BV) : nullptr; } + if (InlineAsm *IA = dyn_cast(V)) + if (BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_inline_assembly)) + return mapValue(V, transAsmINTEL(IA)); + if (CallInst *CI = dyn_cast(V)) return mapValue(V, transCallInst(CI, BB)); @@ -1774,6 +1780,11 @@ SPIRVValue *LLVMToSPIRV::transIntrinsicInst(IntrinsicInst *II, } SPIRVValue *LLVMToSPIRV::transCallInst(CallInst *CI, SPIRVBasicBlock *BB) { + assert(CI); + if (isa(CI->getCalledOperand()) && + BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_inline_assembly)) + return transAsmCallINTEL(CI, BB); + if (CI->isIndirectCall()) return transIndirectCallInst(CI, BB); return transDirectCallInst(CI, BB); @@ -1827,6 +1838,32 @@ SPIRVValue *LLVMToSPIRV::transIndirectCallInst(CallInst *CI, BB); } +SPIRVValue *LLVMToSPIRV::transAsmINTEL(InlineAsm *IA) { + assert(IA); + + // TODO: intention here is to provide information about actual target + // but in fact spir-64 is substituted as triple when translator works + // eventually we need to fix it (not urgent) + StringRef TripleStr(M->getTargetTriple()); + auto AsmTarget = static_cast( + BM->getOrAddAsmTargetINTEL(TripleStr.str())); + auto SIA = BM->addAsmINTEL( + static_cast(transType(IA->getFunctionType())), + AsmTarget, IA->getAsmString(), IA->getConstraintString()); + if (IA->hasSideEffects()) + SIA->addDecorate(DecorationSideEffectsINTEL); + return SIA; +} + +SPIRVValue *LLVMToSPIRV::transAsmCallINTEL(CallInst *CI, SPIRVBasicBlock *BB) { + assert(CI); + auto IA = cast(CI->getCalledOperand()); + return BM->addAsmCallINTELInst( + static_cast(transValue(IA, BB, false)), + transArguments(CI, BB, SPIRVEntry::createUnique(OpAsmCallINTEL).get()), + BB); +} + bool LLVMToSPIRV::transAddressingMode() { Triple TargetTriple(M->getTargetTriple()); diff --git a/lib/SPIRV/SPIRVWriter.h b/lib/SPIRV/SPIRVWriter.h index de4a499f67..f83cf98abf 100644 --- a/lib/SPIRV/SPIRVWriter.h +++ b/lib/SPIRV/SPIRVWriter.h @@ -100,6 +100,8 @@ class LLVMToSPIRV : public ModulePass { SPIRVValue *transCallInst(CallInst *Call, SPIRVBasicBlock *BB); SPIRVValue *transDirectCallInst(CallInst *Call, SPIRVBasicBlock *BB); SPIRVValue *transIndirectCallInst(CallInst *Call, SPIRVBasicBlock *BB); + SPIRVValue *transAsmINTEL(InlineAsm *Asm); + SPIRVValue *transAsmCallINTEL(CallInst *Call, SPIRVBasicBlock *BB); bool transDecoration(Value *V, SPIRVValue *BV); SPIRVWord transFunctionControlMask(Function *); SPIRVFunction *transFunctionDecl(Function *F); diff --git a/lib/SPIRV/libSPIRV/SPIRVAsm.h b/lib/SPIRV/libSPIRV/SPIRVAsm.h new file mode 100644 index 0000000000..d1ee9d8cb1 --- /dev/null +++ b/lib/SPIRV/libSPIRV/SPIRVAsm.h @@ -0,0 +1,142 @@ +//===- SPIRVAsm.h - --*- C++ -*-===// +// +// The LLVM/SPIRV Translator +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// +/// \file +/// +/// This file defines the inline assembler entries defined in SPIRV spec with op +/// codes. +/// +//===----------------------------------------------------------------------===// + +#ifndef SPIRV_LIBSPIRV_SPIRVASM_H +#define SPIRV_LIBSPIRV_SPIRVASM_H + +#include "SPIRVEntry.h" +#include "SPIRVInstruction.h" +#include "SPIRVValue.h" + +namespace SPIRV { + +class SPIRVAsmTargetINTEL : public SPIRVEntry { +public: + static const SPIRVWord FixedWC = 2; + static const Op OC = OpAsmTargetINTEL; + // Complete constructor + SPIRVAsmTargetINTEL(SPIRVModule *M, SPIRVId TheId, + const std::string &TheTarget) + : SPIRVEntry(M, FixedWC + getSizeInWords(TheTarget), OC, TheId), + Target(TheTarget) { + validate(); + } + // Incomplete constructor + SPIRVAsmTargetINTEL() : SPIRVEntry(OC) {} + SPIRVCapVec getRequiredCapability() const override { + return getVec(CapabilityAsmINTEL); + } + SPIRVExtSet getRequiredExtensions() const override { + return getSet(ExtensionID::SPV_INTEL_inline_assembly); + } + const std::string &getTarget() const { return Target; } + +protected: + void validate() const override { + SPIRVEntry::validate(); + assert(WordCount > FixedWC); + assert(OpCode == OC); + } + _SPIRV_DEF_ENCDEC2(Id, Target) + std::string Target; +}; + +class SPIRVAsmINTEL : public SPIRVValue { +public: + static const SPIRVWord FixedWC = 5; + static const Op OC = OpAsmINTEL; + // Complete constructor + SPIRVAsmINTEL(SPIRVModule *M, SPIRVTypeFunction *TheFunctionType, + SPIRVId TheId, SPIRVAsmTargetINTEL *TheTarget, + const std::string &TheInstructions, + const std::string &TheConstraints) + : SPIRVValue(M, + FixedWC + getSizeInWords(TheInstructions) + + getSizeInWords(TheConstraints), + OC, TheFunctionType->getReturnType(), TheId), + Target(TheTarget), FunctionType(TheFunctionType), + Instructions(TheInstructions), Constraints(TheConstraints) { + validate(); + } + // Incomplete constructor + SPIRVAsmINTEL() : SPIRVValue(OC) {} + SPIRVCapVec getRequiredCapability() const override { + return getVec(CapabilityAsmINTEL); + } + SPIRVExtSet getRequiredExtensions() const override { + return getSet(ExtensionID::SPV_INTEL_inline_assembly); + } + const std::string &getInstructions() const { return Instructions; } + const std::string &getConstraints() const { return Constraints; } + SPIRVTypeFunction *getFunctionType() const { return FunctionType; } + +protected: + _SPIRV_DEF_ENCDEC6(Type, Id, FunctionType, Target, Instructions, Constraints) + void validate() const override { + SPIRVValue::validate(); + assert(WordCount > FixedWC); + assert(OpCode == OC); + } + SPIRVAsmTargetINTEL *Target; + SPIRVTypeFunction *FunctionType; + std::string Instructions; + std::string Constraints; +}; + +class SPIRVAsmCallINTEL : public SPIRVInstruction { +public: + static const SPIRVWord FixedWC = 4; + static const Op OC = OpAsmCallINTEL; + // Complete constructor + SPIRVAsmCallINTEL(SPIRVId TheId, SPIRVAsmINTEL *TheAsm, + const std::vector &TheArgs, + SPIRVBasicBlock *TheBB) + : SPIRVInstruction(FixedWC + TheArgs.size(), OC, TheAsm->getType(), TheId, + TheBB), + Asm(TheAsm), Args(TheArgs) { + validate(); + } + // Incomplete constructor + SPIRVAsmCallINTEL() : SPIRVInstruction(OC) {} + SPIRVCapVec getRequiredCapability() const override { + return getVec(CapabilityAsmINTEL); + } + SPIRVExtSet getRequiredExtensions() const override { + return getSet(ExtensionID::SPV_INTEL_inline_assembly); + } + bool isOperandLiteral(unsigned int Index) const override { return false; } + void setWordCount(SPIRVWord TheWordCount) override { + SPIRVEntry::setWordCount(TheWordCount); + Args.resize(TheWordCount - FixedWC); + } + const std::vector &getArguments() const { return Args; } + + SPIRVAsmINTEL *getAsm() const { return Asm; } + +protected: + _SPIRV_DEF_ENCDEC4(Type, Id, Asm, Args) + void validate() const override { + SPIRVInstruction::validate(); + assert(WordCount >= FixedWC); + assert(OpCode == OC); + assert(getBasicBlock() && "Invalid BB"); + assert(getBasicBlock()->getModule() == Asm->getModule()); + } + SPIRVAsmINTEL *Asm; + std::vector Args; +}; + +} // namespace SPIRV +#endif // SPIRV_LIBSPIRV_SPIRVASM_H diff --git a/lib/SPIRV/libSPIRV/SPIRVEntry.cpp b/lib/SPIRV/libSPIRV/SPIRVEntry.cpp index 6bc701bf53..28bacfdb46 100644 --- a/lib/SPIRV/libSPIRV/SPIRVEntry.cpp +++ b/lib/SPIRV/libSPIRV/SPIRVEntry.cpp @@ -38,6 +38,7 @@ //===----------------------------------------------------------------------===// #include "SPIRVEntry.h" +#include "SPIRVAsm.h" #include "SPIRVBasicBlock.h" #include "SPIRVDebug.h" #include "SPIRVDecorate.h" diff --git a/lib/SPIRV/libSPIRV/SPIRVEnum.h b/lib/SPIRV/libSPIRV/SPIRVEnum.h index 33308a1df8..eed9a76cc0 100644 --- a/lib/SPIRV/libSPIRV/SPIRVEnum.h +++ b/lib/SPIRV/libSPIRV/SPIRVEnum.h @@ -366,6 +366,7 @@ template <> inline void SPIRVMap::init() { ADD_VEC_INIT(DecorationReferencedIndirectlyINTEL, {CapabilityIndirectReferencesINTEL}); ADD_VEC_INIT(DecorationIOPipeStorageINTEL, {CapabilityIOPipeINTEL}); + ADD_VEC_INIT(DecorationSideEffectsINTEL, {CapabilityAsmINTEL}); } template <> inline void SPIRVMap::init() { diff --git a/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h b/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h index d3ef9d22f7..5e59bd704b 100644 --- a/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h +++ b/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h @@ -893,6 +893,9 @@ inline bool isValid(spv::Op V) { case OpSubgroupImageBlockWriteINTEL: case OpSubgroupImageMediaBlockReadINTEL: case OpSubgroupImageMediaBlockWriteINTEL: + case OpAsmTargetINTEL: + case OpAsmINTEL: + case OpAsmCallINTEL: case OpVmeImageINTEL: case OpTypeVmeImageINTEL: case OpTypeAvcImePayloadINTEL: diff --git a/lib/SPIRV/libSPIRV/SPIRVModule.cpp b/lib/SPIRV/libSPIRV/SPIRVModule.cpp index 0cb669c104..d9cccd2d93 100644 --- a/lib/SPIRV/libSPIRV/SPIRVModule.cpp +++ b/lib/SPIRV/libSPIRV/SPIRVModule.cpp @@ -38,6 +38,7 @@ //===----------------------------------------------------------------------===// #include "SPIRVModule.h" +#include "SPIRVAsm.h" #include "SPIRVDebug.h" #include "SPIRVEntry.h" #include "SPIRVExtInst.h" @@ -297,6 +298,12 @@ class SPIRVModuleImpl : public SPIRVModule { SPIRVBasicBlock *) override; SPIRVInstruction *addFunctionPointerINTELInst(SPIRVType *, SPIRVFunction *, SPIRVBasicBlock *) override; + SPIRVEntry *getOrAddAsmTargetINTEL(const std::string &) override; + SPIRVValue *addAsmINTEL(SPIRVTypeFunction *, SPIRVAsmTargetINTEL *, + const std::string &, const std::string &) override; + SPIRVInstruction *addAsmCallINTELInst(SPIRVAsmINTEL *, + const std::vector &, + SPIRVBasicBlock *) override; SPIRVInstruction *addCmpInst(Op, SPIRVType *, SPIRVValue *, SPIRVValue *, SPIRVBasicBlock *) override; SPIRVInstruction *addLoadInst(SPIRVValue *, const std::vector &, @@ -444,6 +451,8 @@ class SPIRVModuleImpl : public SPIRVModule { typedef std::vector SPIRVMemberNameVec; typedef std::vector SPIRVDecGroupVec; typedef std::vector SPIRVGroupDecVec; + typedef std::vector SPIRVAsmTargetVector; + typedef std::vector SPIRVAsmVector; typedef std::map SPIRVIdToInstructionSetMap; std::map ExtInstSetIds; typedef std::map SPIRVIdToBuiltinSetMap; @@ -469,6 +478,8 @@ class SPIRVModuleImpl : public SPIRVModule { SPIRVDecorateSet DecorateSet; SPIRVDecGroupVec DecGroupVec; SPIRVGroupDecVec GroupDecVec; + SPIRVAsmTargetVector AsmTargetVec; + SPIRVAsmVector AsmVec; SPIRVExecModelIdSetMap EntryPointSet; SPIRVExecModelIdVecMap EntryPointVec; SPIRVStringMap StrMap; @@ -632,6 +643,14 @@ void SPIRVModuleImpl::layoutEntry(SPIRVEntry *E) { } break; } + case OpAsmTargetINTEL: { + addTo(AsmTargetVec, E); + break; + } + case OpAsmINTEL: { + addTo(AsmVec, E); + break; + } default: if (isTypeOpCode(OC)) TypeVec.push_back(static_cast(E)); @@ -1212,6 +1231,34 @@ SPIRVInstruction *SPIRVModuleImpl::addFunctionPointerINTELInst( new SPIRVFunctionPointerINTEL(getId(), TheType, TheFunction, BB), BB); } +SPIRVEntry * +SPIRVModuleImpl::getOrAddAsmTargetINTEL(const std::string &TheTarget) { + auto TargetIt = std::find_if(AsmTargetVec.begin(), AsmTargetVec.end(), + [&TheTarget](const SPIRVAsmTargetINTEL *Target) { + return Target->getTarget() == TheTarget; + }); + if (TargetIt == AsmTargetVec.end()) + return add(new SPIRVAsmTargetINTEL(this, getId(), TheTarget)); + return *TargetIt; +} + +SPIRVValue *SPIRVModuleImpl::addAsmINTEL(SPIRVTypeFunction *TheType, + SPIRVAsmTargetINTEL *TheTarget, + const std::string &TheInstructions, + const std::string &TheConstraints) { + auto Asm = new SPIRVAsmINTEL(this, TheType, getId(), TheTarget, + TheInstructions, TheConstraints); + return add(Asm); +} + +SPIRVInstruction * +SPIRVModuleImpl::addAsmCallINTELInst(SPIRVAsmINTEL *TheAsm, + const std::vector &TheArguments, + SPIRVBasicBlock *BB) { + return addInstruction( + new SPIRVAsmCallINTEL(getId(), TheAsm, TheArguments, BB), BB); +} + SPIRVInstruction *SPIRVModuleImpl::addBinaryInst(Op TheOpCode, SPIRVType *Type, SPIRVValue *Op1, SPIRVValue *Op2, @@ -1606,8 +1653,13 @@ spv_ostream &operator<<(spv_ostream &O, SPIRVModule &M) { O << MI.MemberNameVec << MI.DecGroupVec << MI.DecorateSet << MI.GroupDecVec << MI.ForwardPointerVec << TopologicalSort(MI.TypeVec, MI.ConstVec, MI.VariableVec, - MI.ForwardPointerVec) - << SPIRVNL() << MI.DebugInstVec << SPIRVNL() << MI.FuncVec; + MI.ForwardPointerVec); + + if (M.isAllowedToUseExtension(ExtensionID::SPV_INTEL_inline_assembly)) { + O << SPIRVNL() << MI.AsmTargetVec << MI.AsmVec; + } + + O << SPIRVNL() << MI.DebugInstVec << SPIRVNL() << MI.FuncVec; return O; } diff --git a/lib/SPIRV/libSPIRV/SPIRVModule.h b/lib/SPIRV/libSPIRV/SPIRVModule.h index a9a0ed0368..d753fbf623 100644 --- a/lib/SPIRV/libSPIRV/SPIRVModule.h +++ b/lib/SPIRV/libSPIRV/SPIRVModule.h @@ -86,6 +86,9 @@ class SPIRVGroupDecorate; class SPIRVGroupMemberDecorate; class SPIRVGroupDecorateGeneric; class SPIRVInstTemplateBase; +class SPIRVAsmTargetINTEL; +class SPIRVAsmINTEL; +class SPIRVAsmCallINTEL; typedef SPIRVBasicBlock SPIRVLabel; struct SPIRVTypeImageDescriptor; @@ -303,6 +306,12 @@ class SPIRVModule { virtual SPIRVInstruction *addFunctionPointerINTELInst(SPIRVType *, SPIRVFunction *, SPIRVBasicBlock *) = 0; + virtual SPIRVEntry *getOrAddAsmTargetINTEL(const std::string &) = 0; + virtual SPIRVValue *addAsmINTEL(SPIRVTypeFunction *, SPIRVAsmTargetINTEL *, + const std::string &, const std::string &) = 0; + virtual SPIRVInstruction *addAsmCallINTELInst(SPIRVAsmINTEL *, + const std::vector &, + SPIRVBasicBlock *) = 0; virtual SPIRVInstruction * addCompositeConstructInst(SPIRVType *, const std::vector &, SPIRVBasicBlock *) = 0; diff --git a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h index 1e3c6f09b9..41c22c5342 100644 --- a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h +++ b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h @@ -335,6 +335,7 @@ template <> inline void SPIRVMap::init() { add(DecorationMaxByteOffset, "MaxByteOffset"); add(DecorationNoSignedWrap, "NoSignedWrap"); add(DecorationNoUnsignedWrap, "NoUnsignedWrap"); + add(DecorationSideEffectsINTEL, "SideEffectsINTEL"); add(DecorationUserSemantic, "UserSemantic"); add(DecorationRegisterINTEL, "RegisterINTEL"); add(DecorationMemoryINTEL, "MemoryINTEL"); @@ -488,6 +489,7 @@ template <> inline void SPIRVMap::init() { add(CapabilitySubgroupImageBlockIOINTEL, "SubgroupImageBlockIOINTEL"); add(CapabilitySubgroupImageMediaBlockIOINTEL, "SubgroupImageMediaBlockIOINTEL"); + add(CapabilityAsmINTEL, "AsmINTEL"); add(CapabilitySubgroupAvcMotionEstimationINTEL, "SubgroupAvcMotionEstimationINTEL"); add(CapabilitySubgroupAvcMotionEstimationIntraINTEL, diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h b/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h index 1279603297..be24e40a5d 100644 --- a/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h +++ b/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h @@ -307,6 +307,9 @@ _SPIRV_OP(SubgroupImageMediaBlockReadINTEL, 5580) _SPIRV_OP(SubgroupImageMediaBlockWriteINTEL, 5581) _SPIRV_OP(FunctionPointerINTEL, 5600) _SPIRV_OP(FunctionPointerCallINTEL, 5601) +_SPIRV_OP(AsmTargetINTEL, 5609) +_SPIRV_OP(AsmINTEL, 5610) +_SPIRV_OP(AsmCallINTEL, 5611) _SPIRV_OP(VmeImageINTEL, 5699) _SPIRV_OP(TypeVmeImageINTEL, 5700) _SPIRV_OP(TypeAvcImePayloadINTEL, 5701) diff --git a/lib/SPIRV/libSPIRV/spirv.hpp b/lib/SPIRV/libSPIRV/spirv.hpp index 2a59f2cd90..b9e0fb3247 100644 --- a/lib/SPIRV/libSPIRV/spirv.hpp +++ b/lib/SPIRV/libSPIRV/spirv.hpp @@ -342,71 +342,72 @@ enum FunctionParameterAttribute { }; enum Decoration { - DecorationRelaxedPrecision = 0, - DecorationSpecId = 1, - DecorationBlock = 2, - DecorationBufferBlock = 3, - DecorationRowMajor = 4, - DecorationColMajor = 5, - DecorationArrayStride = 6, - DecorationMatrixStride = 7, - DecorationGLSLShared = 8, - DecorationGLSLPacked = 9, - DecorationCPacked = 10, - DecorationBuiltIn = 11, - DecorationNoPerspective = 13, - DecorationFlat = 14, - DecorationPatch = 15, - DecorationCentroid = 16, - DecorationSample = 17, - DecorationInvariant = 18, - DecorationRestrict = 19, - DecorationAliased = 20, - DecorationVolatile = 21, - DecorationConstant = 22, - DecorationCoherent = 23, - DecorationNonWritable = 24, - DecorationNonReadable = 25, - DecorationUniform = 26, - DecorationSaturatedConversion = 28, - DecorationStream = 29, - DecorationLocation = 30, - DecorationComponent = 31, - DecorationIndex = 32, - DecorationBinding = 33, - DecorationDescriptorSet = 34, - DecorationOffset = 35, - DecorationXfbBuffer = 36, - DecorationXfbStride = 37, - DecorationFuncParamAttr = 38, - DecorationFPRoundingMode = 39, - DecorationFPFastMathMode = 40, - DecorationLinkageAttributes = 41, - DecorationNoContraction = 42, - DecorationInputAttachmentIndex = 43, - DecorationAlignment = 44, - DecorationMaxByteOffset = 45, - DecorationNoSignedWrap = 4469, - DecorationNoUnsignedWrap = 4470, - DecorationOverrideCoverageNV = 5248, - DecorationPassthroughNV = 5250, - DecorationViewportRelativeNV = 5252, - DecorationSecondaryViewportRelativeNV = 5256, - DecorationReferencedIndirectlyINTEL = 5602, - DecorationUserSemantic = 5635, - DecorationRegisterINTEL = 5825, - DecorationMemoryINTEL = 5826, - DecorationNumbanksINTEL = 5827, - DecorationBankwidthINTEL = 5828, - DecorationMaxPrivateCopiesINTEL = 5829, - DecorationSinglepumpINTEL = 5830, - DecorationDoublepumpINTEL = 5831, - DecorationMaxReplicatesINTEL = 5832, - DecorationSimpleDualPortINTEL = 5833, - DecorationMergeINTEL = 5834, - DecorationBankBitsINTEL = 5835, - DecorationIOPipeStorageINTEL = 5944, - DecorationMax = 0x7fffffff, + DecorationRelaxedPrecision = 0, + DecorationSpecId = 1, + DecorationBlock = 2, + DecorationBufferBlock = 3, + DecorationRowMajor = 4, + DecorationColMajor = 5, + DecorationArrayStride = 6, + DecorationMatrixStride = 7, + DecorationGLSLShared = 8, + DecorationGLSLPacked = 9, + DecorationCPacked = 10, + DecorationBuiltIn = 11, + DecorationNoPerspective = 13, + DecorationFlat = 14, + DecorationPatch = 15, + DecorationCentroid = 16, + DecorationSample = 17, + DecorationInvariant = 18, + DecorationRestrict = 19, + DecorationAliased = 20, + DecorationVolatile = 21, + DecorationConstant = 22, + DecorationCoherent = 23, + DecorationNonWritable = 24, + DecorationNonReadable = 25, + DecorationUniform = 26, + DecorationSaturatedConversion = 28, + DecorationStream = 29, + DecorationLocation = 30, + DecorationComponent = 31, + DecorationIndex = 32, + DecorationBinding = 33, + DecorationDescriptorSet = 34, + DecorationOffset = 35, + DecorationXfbBuffer = 36, + DecorationXfbStride = 37, + DecorationFuncParamAttr = 38, + DecorationFPRoundingMode = 39, + DecorationFPFastMathMode = 40, + DecorationLinkageAttributes = 41, + DecorationNoContraction = 42, + DecorationInputAttachmentIndex = 43, + DecorationAlignment = 44, + DecorationMaxByteOffset = 45, + DecorationNoSignedWrap = 4469, + DecorationNoUnsignedWrap = 4470, + DecorationOverrideCoverageNV = 5248, + DecorationPassthroughNV = 5250, + DecorationViewportRelativeNV = 5252, + DecorationSecondaryViewportRelativeNV = 5256, + DecorationReferencedIndirectlyINTEL = 5602, + DecorationSideEffectsINTEL = 5608, + DecorationUserSemantic = 5635, + DecorationRegisterINTEL = 5825, + DecorationMemoryINTEL = 5826, + DecorationNumbanksINTEL = 5827, + DecorationBankwidthINTEL = 5828, + DecorationMaxPrivateCopiesINTEL = 5829, + DecorationSinglepumpINTEL = 5830, + DecorationDoublepumpINTEL = 5831, + DecorationMaxReplicatesINTEL = 5832, + DecorationSimpleDualPortINTEL = 5833, + DecorationMergeINTEL = 5834, + DecorationBankBitsINTEL = 5835, + DecorationIOPipeStorageINTEL = 5944, + DecorationMax = 0x7fffffff, }; enum BuiltIn { @@ -681,6 +682,7 @@ enum Capability { CapabilitySubgroupImageMediaBlockIOINTEL = 5579, CapabilityFunctionPointersINTEL = 5603, CapabilityIndirectReferencesINTEL = 5604, + CapabilityAsmINTEL = 5606, CapabilitySubgroupAvcMotionEstimationINTEL = 5696, CapabilitySubgroupAvcMotionEstimationIntraINTEL = 5697, CapabilitySubgroupAvcMotionEstimationChromaINTEL = 5698, @@ -689,8 +691,8 @@ enum Capability { CapabilityFPGALoopControlsINTEL = 5888, CapabilityBlockingPipesINTEL = 5945, CapabilityFPGARegINTEL = 5948, - CapabilityKernelAttributesINTEL= 5892, - CapabilityFPGAKernelAttributesINTEL= 5897, + CapabilityKernelAttributesINTEL = 5892, + CapabilityFPGAKernelAttributesINTEL = 5897, CapabilityIOPipeINTEL = 5943, CapabilityMax = 0x7fffffff, }; @@ -1019,6 +1021,9 @@ enum Op { OpSubgroupImageMediaBlockWriteINTEL = 5581, OpFunctionPointerINTEL = 5600, OpFunctionPointerCallINTEL = 5601, + OpAsmTargetINTEL = 5609, + OpAsmINTEL = 5610, + OpAsmCallINTEL = 5611, OpVmeImageINTEL = 5699, OpTypeVmeImageINTEL = 5700, OpTypeAvcImePayloadINTEL = 5701, diff --git a/test/inline_asm_basic.cl b/test/inline_asm_basic.cl new file mode 100644 index 0000000000..745a24fe96 --- /dev/null +++ b/test/inline_asm_basic.cl @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -triple spir64-unknown-unknown -x cl -cl-std=CL2.0 -O0 -emit-llvm-bc %s -o %t.bc +// RUN: llvm-spirv -spirv-ext=+SPV_INTEL_inline_assembly %t.bc -o %t.spv +// RUN: llvm-spirv %t.spv -to-text -o %t.spt +// RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +// RUN: llvm-spirv -r %t.spv -o %t.bc +// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM + +// CHECK-SPIRV: {{[0-9]+}} Capability AsmINTEL +// CHECK-SPIRV: {{[0-9]+}} Extension "SPV_INTEL_inline_assembly" +// CHECK-SPIRV: {{[0-9]+}} AsmTargetINTEL +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL + +// CHECK-LLVM: call void asm sideeffect + +kernel void foo() { + __asm__ volatile (""); +} diff --git a/test/inline_asm_clobbers.cl b/test/inline_asm_clobbers.cl new file mode 100644 index 0000000000..a640430ec9 --- /dev/null +++ b/test/inline_asm_clobbers.cl @@ -0,0 +1,94 @@ +// RUN: %clang_cc1 -triple spir64-unknown-unknown -x cl -cl-std=CL2.0 -O0 -emit-llvm-bc %s -o %t.bc +// RUN: llvm-spirv -spirv-ext=+SPV_INTEL_inline_assembly %t.bc -o %t.spv +// RUN: llvm-spirv %t.spv -to-text -o %t.spt +// RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +// RUN: llvm-spirv -r %t.spv -o %t.bc +// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM + +// Excerpt from opencl-c-base.h +typedef __SIZE_TYPE__ size_t; + +// Excerpt from opencl-c.h to speed up compilation. +#define __ovld __attribute__((overloadable)) +#define __cnfn __attribute__((const)) +size_t __ovld __cnfn get_global_id(unsigned int dimindx); + +// CHECK-SPIRV: {{[0-9]+}} Capability AsmINTEL +// CHECK-SPIRV: {{[0-9]+}} Extension "SPV_INTEL_inline_assembly" +// CHECK-SPIRV: {{[0-9]+}} AsmTargetINTEL + +// XCHECK-LLVM: [[STRUCTYPE:%[a-z0-9]+]] = type { i32, i32 } + +// CHECK-LLVM-LABEL: define spir_kernel void @mem_clobber +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} """~{cc},~{memory}" +// CHECK-LLVM: [[VALUE:%[0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** +// CHECK-LLVM-NEXT: getelementptr inbounds i32, i32 addrspace(1)* [[VALUE]], i64 0 +// CHECK-LLVM-NEXT: store i32 1, i32 addrspace(1)* +// CHECK-LLVM-NEXT: call void asm sideeffect "", "~{cc},~{memory}"() +// CHECK-LLVM-NEXT: load i32 addrspace(1)*, i32 addrspace(1)** + +kernel void mem_clobber(global int *x) { + x[0] = 1; + __asm__ ("":::"cc","memory"); + x[0] += 1; +} + +// CHECK-LLVM-LABEL: define spir_kernel void @out_clobber +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_out $0""=&r" +// CHECK-LLVM: barrier +// CHECK-LLVM: store i32 %{{[a-z0-9]+}}, i32* [[VALUE:%[a-z0-9]+]], align 4 +// CHECK-LLVM-NEXT: [[STOREVAL:%[a-z0-9]+]] = call i32 asm "earlyclobber_instruction_out $0", "=&r"() +// CHECK-LLVM: store i32 [[STOREVAL]], i32* [[VALUE]], align 4 + +kernel void out_clobber(global int *x) { + int i = get_global_id(0); + __asm__ ("barrier"); + int a = x[i]; + __asm__ ("earlyclobber_instruction_out %0":"=&r"(a)); + a += 1; + x[i] = a; +} + +// TODO: This fails on debug build with assert "function type not legal for constraints" +// Probably I am not completely understand what happens +// Or bug in clang FE. To investigate later, change xchecks to checks and enable + +// XCHECK-LLVM-LABEL: define spir_kernel void @in_clobber +// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_in $0""&r" +// XCHECK-LLVM: barrier +// XCHECK-LLVM: getelementptr +// XCHECK-LLVM: store i32 %{{[a-z0-9]+}}, i32* [[LOADVAL:%[a-z0-9]+]], align 4 +// XCHECK-LLVM-NEXT: [[VALUE:%[a-z0-9]+]] = load i32, i32* [[LOADVAL]], align 4 +// XCHECK-LLVM-NEXT: call void asm sideeffect "earlyclobber_instruction_in $0", "&r"(i32 [[VALUE]]) +// XCHECK-LLVM: %{{[a-z0-9]+}} = load i32, i32* [[LOADVAL]], align 4 + +#if 0 +kernel void in_clobber(global int *x) { + int i = get_global_id(0); + __asm__ ("barrier"); + int a = x[i]; + __asm__ ("earlyclobber_instruction_in %0"::"&r"(a)); + a += 1; + x[i] = a; +} +#endif + +// XCHECK-LLVM-LABEL: define spir_kernel void @mixed_clobber +// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixedclobber_instruction $0 $1 $2""=&r,=&r,&r,1,~{cc},~{memory}" + +#if 0 +kernel void mixed_clobber(global int *x, global int *y, global int *z) { + int i = get_global_id(0); + int a = x[i]; + int b = y[i]; + int c = z[i]; + __asm__ ("mixedclobber_instruction %0 %1 %2":"=&r"(a),"+&r"(b):"&r"(c):"cc","memory"); + a += 1; + b += 1; + c += 1; + x[i] = c; + y[i] = a; + z[i] = b; +} +#endif + diff --git a/test/inline_asm_constraints.cl b/test/inline_asm_constraints.cl new file mode 100644 index 0000000000..98bf490642 --- /dev/null +++ b/test/inline_asm_constraints.cl @@ -0,0 +1,100 @@ +// RUN: %clang_cc1 -triple spir64-unknown-unknown -x cl -cl-std=CL2.0 -O0 -emit-llvm-bc %s -o %t.bc +// RUN: llvm-spirv -spirv-ext=+SPV_INTEL_inline_assembly %t.bc -o %t.spv +// RUN: llvm-spirv %t.spv -to-text -o %t.spt +// RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +// RUN: llvm-spirv -r %t.spv -o %t.bc +// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM + +// Excerpt from opencl-c-base.h +typedef __SIZE_TYPE__ size_t; +typedef unsigned char uchar; +typedef unsigned short ushort; +typedef unsigned int uint; +typedef unsigned long ulong; + +// Excerpt from opencl-c.h to speed up compilation. +#define __ovld __attribute__((overloadable)) +#define __cnfn __attribute__((const)) +size_t __ovld __cnfn get_global_id(unsigned int dimindx); + +// CHECK-SPIRV: {{[0-9]+}} Capability AsmINTEL +// CHECK-SPIRV: {{[0-9]+}} Extension "SPV_INTEL_inline_assembly" +// CHECK-SPIRV-COUNT-1: {{[0-9]+}} AsmTargetINTEL + +// CHECK-LLVM: [[STRUCTYPE:%[a-z]+]] = type { i32, i8, float } + +// CHECK-LLVM-LABEL: define spir_kernel void @test_int +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "intcommand $0 $1""=r,r" +// CHECK-LLVM: [[VALUE:%[0-9]+]] = call i32 asm sideeffect "intcommand $0 $1", "=r,r"(i32 %{{[0-9]+}}) +// CHECK-LLVM-NEXT: store i32 [[VALUE]], i32 addrspace(1)* + +kernel void test_int(global int *in, global int *out) { + int i = get_global_id(0); + __asm__ volatile ("intcommand %0 %1" : "=r"(out[i]) : "r"(in[i])); +} + +// CHECK-LLVM-LABEL: define spir_kernel void @test_float +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "floatcommand $0 $1""=r,r" +// CHECK-LLVM: [[VALUE:%[0-9]+]] = call float asm sideeffect "floatcommand $0 $1", "=r,r"(float %{{[0-9]+}}) +// CHECK-LLVM-NEXT: store float [[VALUE]], float addrspace(1)* + +kernel void test_float(global float *in, global float *out) { + int i = get_global_id(0); + __asm__ volatile ("floatcommand %0 %1" : "=r"(out[i]) : "r"(in[i])); +} + +// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_integral +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_integral_command $0 $3 $1 $2""=r,r,r,r" +// CHECK-LLVM: [[VALUE:%[0-9]+]] = call i64 asm sideeffect "mixed_integral_command $0 $3 $1 $2", "=r,r,r,r"(i16 %{{[0-9]+}}, i32 %{{[0-9]+}}, i8 %{{[0-9]+}}) +// CHECK-LLVM-NEXT: store i64 [[VALUE]], i64 addrspace(1)* + +kernel void test_mixed_integral(global uchar *A, global ushort *B, global uint *C, global ulong *D) { + int wiId = get_global_id(0); + __asm__ volatile ("mixed_integral_command %0 %3 %1 %2" + : "=r"(D[wiId]) : "r"(B[wiId]), "r"(C[wiId]), "r"(A[wiId])); +} + +// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_floating +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_floating_command $0 $1 $2""=r,r,r" +// CHECK-LLVM: [[VALUE:%[0-9]+]] = call half asm sideeffect "mixed_floating_command $0 $1 $2", "=r,r,r"(double %{{[0-9]+}}, float %{{[0-9]+}}) +// CHECK-LLVM-NEXT: store half [[VALUE]], half addrspace(1)* + +kernel void test_mixed_floating(global float *A, global half *B, global double *C) { + int wiId = get_global_id(0); + __asm__ volatile ("mixed_floating_command %0 %1 %2" + : "=r"(B[wiId]) : "r"(C[wiId]), "r"(A[wiId])); +} + +// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_all +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_all_command $0 $3 $1 $2""=r,r,r,r" +// CHECK-LLVM: [[VALUE:%[0-9]+]] = call i8 asm sideeffect "mixed_all_command $0 $3 $1 $2", "=r,r,r,r"(float %{{[0-9]+}}, i32 %{{[0-9]+}}, i8 %{{[0-9]+}}) +// CHECK-LLVM-NEXT: store i8 [[VALUE]], i8 addrspace(1)* + +kernel void test_mixed_all(global uchar *A, global float *B, global uint *C, global bool *D) { + int wiId = get_global_id(0); + __asm__ volatile ("mixed_all_command %0 %3 %1 %2" + : "=r"(D[wiId]) : "r"(B[wiId]), "r"(C[wiId]), "r"(A[wiId])); +} + +// CHECK-LLVM-LABEL: define spir_kernel void @test_multiple +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "multiple_command $0 $0 $1 $1 $2 $2""=r,=r,=r,0,1,2" +// CHECK-LLVM: [[VALUE:%[0-9]+]] = call [[STRUCTYPE]] asm sideeffect "multiple_command $0 $0 $1 $1 $2 $2", "=r,=r,=r,0,1,2"(i32 %{{[0-9]+}}, i8 %{{[0-9]+}}, float %{{[0-9]+}}) +// CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 0 +// CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 1 +// CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 2 + +kernel void test_multiple(global uchar *A, global float *B, global uint *C) { + int wiId = get_global_id(0); + __asm__ volatile ("multiple_command %0 %0 %1 %1 %2 %2" + : "+r"(C[wiId]), "+r"(A[wiId]), "+r"(B[wiId])); +} + +// CHECK-LLVM-LABEL: define spir_kernel void @test_constants +// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "constcommand $0 $1""i,i" +// CHECK-LLVM: call void asm sideeffect "constcommand $0 $1", "i,i"(i32 1, double 2.000000e+00) + +kernel void test_constants() { + int i = get_global_id(0); + __asm__ volatile ("constcommand %0 %1" : : "i"(1), "i"(2.0)); +} +