From 5d28c9377cba5dc233fd9b4e6f1305e1898079cd Mon Sep 17 00:00:00 2001 From: asudarsa Date: Wed, 2 Aug 2023 05:25:11 -0400 Subject: [PATCH] [llvm-spirv] Cherry pick Khronos changes to expand collection of entry point interfaces (PR #1334) (#10623) This PR pulls in the following PR from upstream Khronos SPIRV-LLVM-Translator repo: https://github.com/KhronosGroup/SPIRV-LLVM-Translator/pull/1334 ` This is a patch to expand the collection of entry point interfaces. In SPIR-V 1.4 and later OpEntryPoint must list all global variables in the interface. ` In addition, a couple of minor changes have been added to sync with latest code. This patch addresses https://github.com/intel/llvm/issues/9958 Updated the following tests to sync with upstream as well: llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl Thanks --------- Signed-off-by: Arvind Sudarsanam --- llvm-spirv/lib/SPIRV/SPIRVWriter.cpp | 30 ++++++----- llvm-spirv/lib/SPIRV/SPIRVWriter.h | 2 + llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h | 1 - llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp | 6 ++- llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h | 1 + llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp | 38 +++++--------- llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h | 7 ++- llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp | 1 + llvm-spirv/test/ExecutionMode.ll | 3 -- llvm-spirv/test/copy_object.spt | 2 +- llvm-spirv/test/entry-point-interfaces.ll | 52 +++++++++++++++++++ .../inline_asm_clobbers.cl | 8 +-- .../inline_asm_constraints.cl | 14 ++--- llvm-spirv/test/negative/unimplemented.spt | 2 +- llvm-spirv/test/right_shift.spt | 2 +- 15 files changed, 107 insertions(+), 62 deletions(-) create mode 100644 llvm-spirv/test/entry-point-interfaces.ll diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp index 76a599bf99fb0..f5525074086c7 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp @@ -874,9 +874,7 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) { BM->setName(BF, F->getName().str()); } - if (isKernel(F)) - BM->addEntryPoint(ExecutionModelKernel, BF->getId()); - else if (F->getLinkage() != GlobalValue::InternalLinkage) + if (!isKernel(F) && F->getLinkage() != GlobalValue::InternalLinkage) BF->setLinkageType(transLinkageType(F)); // Translate OpenCL/SYCL buffer_location metadata if it's attached to the @@ -4898,12 +4896,15 @@ bool LLVMToSPIRVBase::isAnyFunctionReachableFromFunction( return false; } -void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF, - Function *F) { +std::vector +LLVMToSPIRVBase::collectEntryPointInterfaces(SPIRVFunction *SF, Function *F) { + std::vector Interface; for (auto &GV : M->globals()) { const auto AS = GV.getAddressSpace(); - if (AS != SPIRAS_Input && AS != SPIRAS_Output) - continue; + SPIRVModule *BM = SF->getModule(); + if (!BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4)) + if (AS != SPIRAS_Input && AS != SPIRAS_Output) + continue; std::unordered_set Funcs; @@ -4915,9 +4916,14 @@ void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF, } if (isAnyFunctionReachableFromFunction(F, Funcs)) { - SF->addVariable(ValueMap[&GV]); + SPIRVWord ModuleVersion = static_cast(BM->getSPIRVVersion()); + if (AS != SPIRAS_Input && AS != SPIRAS_Output && + ModuleVersion < static_cast(VersionNumber::SPIRV_1_4)) + BM->setMinSPIRVVersion(VersionNumber::SPIRV_1_4); + Interface.push_back(ValueMap[&GV]->getId()); } } + return Interface; } void LLVMToSPIRVBase::mutateFuncArgType( @@ -5118,10 +5124,10 @@ void LLVMToSPIRVBase::transFunction(Function *I) { joinFPContract(I, FPContract::ENABLED); fpContractUpdateRecursive(I, getFPContract(I)); - bool IsKernelEntryPoint = isKernel(I); - - if (IsKernelEntryPoint) { - collectInputOutputVariables(BF, I); + if (isKernel(I)) { + auto Interface = collectEntryPointInterfaces(BF, I); + BM->addEntryPoint(ExecutionModelKernel, BF->getId(), I->getName().str(), + Interface); } } diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.h b/llvm-spirv/lib/SPIRV/SPIRVWriter.h index 6e082324451b4..4da4c03fbe570 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVWriter.h +++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.h @@ -248,6 +248,8 @@ class LLVMToSPIRVBase : protected BuiltinCallHelper { const Function *FS, const std::unordered_set Funcs) const; void collectInputOutputVariables(SPIRVFunction *SF, Function *F); + std::vector collectEntryPointInterfaces(SPIRVFunction *BF, + Function *F); }; class LLVMToSPIRVPass : public PassInfoMixin { diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h index d4b68d3522d79..c4c2a0cc1238b 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h @@ -268,7 +268,6 @@ class SPIRVDecorateLinkageAttr : public SPIRVDecorate { #ifdef _SPIRV_SUPPORT_TEXT_FMT if (SPIRVUseTextFormat) { Encoder << getString(Literals.cbegin(), Literals.cend() - 1); - Encoder.OS << " "; Encoder << (SPIRVLinkageTypeKind)Literals.back(); } else #endif diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp index 756f658203c3c..0ed8267eb59ed 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp @@ -621,9 +621,11 @@ void SPIRVEntryPoint::encode(spv_ostream &O) const { } void SPIRVEntryPoint::decode(std::istream &I) { - getDecoder(I) >> ExecModel >> Target >> Name >> Variables; + getDecoder(I) >> ExecModel >> Target >> Name; + Variables.resize(WordCount - FixedWC - getSizeInWords(Name) + 1); + getDecoder(I) >> Variables; Module->setName(getOrCreateTarget(), Name); - Module->addEntryPoint(ExecModel, Target); + Module->addEntryPoint(ExecModel, Target, Name, Variables); } void SPIRVExecutionMode::encode(spv_ostream &O) const { diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h index 2d5a03935d9dd..0b695c054c1cd 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h @@ -533,6 +533,7 @@ template class SPIRVAnnotation : public SPIRVAnnotationGeneric { class SPIRVEntryPoint : public SPIRVAnnotation { public: + static const SPIRVWord FixedWC = 4; SPIRVEntryPoint(SPIRVModule *TheModule, SPIRVExecutionModelKind, SPIRVId TheId, const std::string &TheName, std::vector Variables); diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp index 05d958de3b4cc..e4984d9df09ef 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp @@ -128,20 +128,6 @@ class SPIRVModuleImpl : public SPIRVModule { getValueTypes(const std::vector &) const override; SPIRVMemoryModelKind getMemoryModel() const override { return MemoryModel; } SPIRVConstant *getLiteralAsConstant(unsigned Literal) override; - unsigned getNumEntryPoints(SPIRVExecutionModelKind EM) const override { - auto Loc = EntryPointVec.find(EM); - if (Loc == EntryPointVec.end()) - return 0; - return Loc->second.size(); - } - SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind EM, - unsigned I) const override { - auto Loc = EntryPointVec.find(EM); - if (Loc == EntryPointVec.end()) - return nullptr; - assert(I < Loc->second.size()); - return get(Loc->second[I]); - } unsigned getNumFunctions() const override { return FuncVec.size(); } unsigned getNumVariables() const override { return VariableVec.size(); } SourceLanguage getSourceLanguage(SPIRVWord *Ver = nullptr) const override { @@ -225,8 +211,9 @@ class SPIRVModuleImpl : public SPIRVModule { SPIRVGroupMemberDecorate * addGroupMemberDecorate(SPIRVDecorationGroup *Group, const std::vector &Targets) override; - void addEntryPoint(SPIRVExecutionModelKind ExecModel, - SPIRVId EntryPoint) override; + void addEntryPoint(SPIRVExecutionModelKind ExecModel, SPIRVId EntryPoint, + const std::string &Name, + const std::vector &Variables) override; SPIRVForward *addForward(SPIRVType *Ty) override; SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) override; SPIRVFunction *addFunction(SPIRVFunction *) override; @@ -508,11 +495,11 @@ class SPIRVModuleImpl : public SPIRVModule { typedef std::vector SPIRVGroupDecVec; typedef std::vector SPIRVAsmTargetVector; typedef std::vector SPIRVAsmVector; + typedef std::vector SPIRVEntryPointVec; typedef std::map SPIRVIdToInstructionSetMap; std::map ExtInstSetIds; typedef std::map SPIRVIdToBuiltinSetMap; typedef std::map SPIRVExecModelIdSetMap; - typedef std::map SPIRVExecModelIdVecMap; typedef std::unordered_map SPIRVStringMap; typedef std::map>> SPIRVUnknownStructFieldMap; @@ -540,7 +527,7 @@ class SPIRVModuleImpl : public SPIRVModule { SPIRVAsmTargetVector AsmTargetVec; SPIRVAsmVector AsmVec; SPIRVExecModelIdSetMap EntryPointSet; - SPIRVExecModelIdVecMap EntryPointVec; + SPIRVEntryPointVec EntryPointVec; SPIRVStringMap StrMap; SPIRVCapMap CapMap; SPIRVUnknownStructFieldMap UnknownStructFieldMap; @@ -1086,11 +1073,14 @@ SPIRVModuleImpl::addDecorate(SPIRVDecorateGeneric *Dec) { } void SPIRVModuleImpl::addEntryPoint(SPIRVExecutionModelKind ExecModel, - SPIRVId EntryPoint) { + SPIRVId EntryPoint, const std::string &Name, + const std::vector &Variables) { assert(isValid(ExecModel) && "Invalid execution model"); assert(EntryPoint != SPIRVID_INVALID && "Invalid entry point"); + auto *EP = + add(new SPIRVEntryPoint(this, ExecModel, EntryPoint, Name, Variables)); + EntryPointVec.push_back(EP); EntryPointSet[ExecModel].insert(EntryPoint); - EntryPointVec[ExecModel].push_back(EntryPoint); addCapabilities(SPIRV::getCapability(ExecModel)); } @@ -1937,14 +1927,10 @@ spv_ostream &operator<<(spv_ostream &O, SPIRVModule &M) { O << SPIRVMemoryModel(&M); - for (auto &I : MI.EntryPointVec) - for (auto &II : I.second) - O << SPIRVEntryPoint(&M, I.first, II, M.get(II)->getName(), - M.get(II)->getVariables()); + O << MI.EntryPointVec; for (auto &I : MI.EntryPointVec) - for (auto &II : I.second) - MI.get(II)->encodeExecutionModes(O); + MI.get(I->getTargetId())->encodeExecutionModes(O); O << MI.StringVec; diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h index db2d475dfc2c7..b953f90c3a6fd 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h @@ -133,14 +133,11 @@ class SPIRVModule { virtual const SPIRVCapMap &getCapability() const = 0; virtual bool hasCapability(SPIRVCapabilityKind) const = 0; virtual SPIRVExtInstSetKind getBuiltinSet(SPIRVId) const = 0; - virtual SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind, - unsigned) const = 0; virtual std::set &getExtension() = 0; virtual SPIRVFunction *getFunction(unsigned) const = 0; virtual SPIRVVariable *getVariable(unsigned) const = 0; virtual SPIRVMemoryModelKind getMemoryModel() const = 0; virtual unsigned getNumFunctions() const = 0; - virtual unsigned getNumEntryPoints(SPIRVExecutionModelKind) const = 0; virtual unsigned getNumVariables() const = 0; virtual SourceLanguage getSourceLanguage(SPIRVWord *) const = 0; virtual std::set &getSourceExtension() = 0; @@ -223,7 +220,9 @@ class SPIRVModule { const std::vector &Targets) = 0; virtual SPIRVGroupDecorateGeneric * addGroupDecorateGeneric(SPIRVGroupDecorateGeneric *GDec) = 0; - virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId) = 0; + virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId, + const std::string &, + const std::vector &) = 0; virtual SPIRVForward *addForward(SPIRVType *Ty) = 0; virtual SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) = 0; virtual SPIRVFunction *addFunction(SPIRVFunction *) = 0; diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp index 7244676164e42..65ca5cbb61232 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp @@ -180,6 +180,7 @@ const SPIRVEncoder &operator<<(const SPIRVEncoder &O, const std::string &Str) { #ifdef _SPIRV_SUPPORT_TEXT_FMT if (SPIRVUseTextFormat) { writeQuotedString(O.OS, Str); + O.OS << " "; return O; } #endif diff --git a/llvm-spirv/test/ExecutionMode.ll b/llvm-spirv/test/ExecutionMode.ll index 9858342eb21c3..5f6aef54e69e2 100644 --- a/llvm-spirv/test/ExecutionMode.ll +++ b/llvm-spirv/test/ExecutionMode.ll @@ -1,9 +1,6 @@ ; RUN: llvm-as < %s | llvm-spirv -spirv-text -o %t ; RUN: FileCheck < %t %s -; check for magic number followed by version 1.1 -; CHECK: 119734787 65792 - ; CHECK-DAG: TypeVoid [[VOID:[0-9]+]] ; CHECK-DAG: EntryPoint 6 [[WORKER:[0-9]+]] "worker" diff --git a/llvm-spirv/test/copy_object.spt b/llvm-spirv/test/copy_object.spt index 8127dbd41b335..7ff03be1b2293 100644 --- a/llvm-spirv/test/copy_object.spt +++ b/llvm-spirv/test/copy_object.spt @@ -5,7 +5,7 @@ 2 Capability Int64 2 Capability Int8 3 MemoryModel 2 2 -8 EntryPoint 6 1 "copy_object" +6 EntryPoint 6 1 "copy_object" 3 Source 3 102000 3 Name 2 "in" 4 Decorate 3 BuiltIn 28 diff --git a/llvm-spirv/test/entry-point-interfaces.ll b/llvm-spirv/test/entry-point-interfaces.ll new file mode 100644 index 0000000000000..99934b397a72f --- /dev/null +++ b/llvm-spirv/test/entry-point-interfaces.ll @@ -0,0 +1,52 @@ +; RUN: llvm-as %s -o %t.bc + +; RUN: llvm-spirv %t.bc -o %t.spv +; RUN: spirv-val --target-env spv1.4 %t.spv +; RUN: llvm-spirv -to-text %t.spv -o %t.from.spv.spt +; RUN: FileCheck < %t.from.spv.spt %s --check-prefix=CHECK-SPIRV + +; RUN: llvm-spirv -spirv-text %t.bc -o %t.from.bc.spt +; RUN: FileCheck < %t.from.bc.spt %s --check-prefix=CHECK-SPIRV + +; CHECK-SPIRV: 7 EntryPoint 6 [[#]] "test" [[#Interface1:]] [[#Interface2:]] +; CHECK-SPIRV: TypeInt [[#TypeInt:]] 32 0 +; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant1:]] 1 +; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant2:]] 3 +; CHECK-SPIRV: Variable [[#]] [[#Interface1]] 0 [[#Constant1]] +; CHECK-SPIRV: Variable [[#]] [[#Interface2]] 0 [[#Constant2]] + +; ModuleID = 'source.cpp' +source_filename = "source.cpp" +target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spir" + +@var = dso_local addrspace(2) constant i32 1, align 4 +@var2 = dso_local addrspace(2) constant i32 3, align 4 +@var.const = private unnamed_addr addrspace(2) constant i32 1, align 4 +@var2.const = private unnamed_addr addrspace(2) constant i32 3, align 4 + +; Function Attrs: convergent noinline norecurse nounwind optnone +define dso_local spir_kernel void @test() #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !2 !kernel_arg_type !2 !kernel_arg_base_type !2 !kernel_arg_type_qual !2 !kernel_arg_host_accessible !2 !kernel_arg_pipe_depth !2 !kernel_arg_pipe_io !2 !kernel_arg_buffer_location !2 { +entry: + %0 = load i32, i32 addrspace(2)* @var.const, align 4 + %1 = load i32, i32 addrspace(2)* @var2.const, align 4 + %mul = mul nsw i32 %0, %1 + %mul1 = mul nsw i32 %mul, 2 + ret void +} + +attributes #0 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" } + +!opencl.enable.FP_CONTRACT = !{} +!opencl.ocl.version = !{!0} +!opencl.spir.version = !{!0} +!llvm.module.flags = !{!1} +!opencl.used.extensions = !{!2} +!opencl.used.optional.core.features = !{!2} +!opencl.compiler.options = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 2, i32 0} +!1 = !{i32 7, !"frame-pointer", i32 2} +!2 = !{} +!3 = !{!"Compiler"} diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl index ee07df4e37651..68680455b1da7 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl @@ -20,7 +20,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx); // 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-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "" "~{cc},~{memory}" // CHECK-LLVM: [[VALUE:%[0-9]+]] = load ptr addrspace(1), ptr // CHECK-LLVM-NEXT: getelementptr inbounds i32, ptr addrspace(1) [[VALUE]], i64 0 // CHECK-LLVM-NEXT: store i32 1, ptr addrspace(1) @@ -34,7 +34,7 @@ kernel void mem_clobber(global int *x) { } // 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-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]+}}, ptr [[VALUE:%[a-z0-9]+]], align 4 // CHECK-LLVM-NEXT: [[STOREVAL:%[a-z0-9]+]] = call i32 asm "earlyclobber_instruction_out $0", "=&r"() @@ -54,7 +54,7 @@ kernel void out_clobber(global int *x) { // 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-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]+}}, ptr [[LOADVAL:%[a-z0-9]+]], align 4 @@ -74,7 +74,7 @@ kernel void in_clobber(global int *x) { #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}" +// 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) { diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl index 4653776b6417c..a6c0fdc64b671 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl @@ -24,7 +24,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx); // 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-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]], ptr addrspace(1) @@ -34,7 +34,7 @@ kernel void test_int(global int *in, global int *out) { } // 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-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]], ptr addrspace(1) @@ -44,7 +44,7 @@ kernel void test_float(global float *in, global float *out) { } // 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-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]], ptr addrspace(1) @@ -55,7 +55,7 @@ kernel void test_mixed_integral(global uchar *A, global ushort *B, global uint * } // 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-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]], ptr addrspace(1) @@ -66,7 +66,7 @@ kernel void test_mixed_floating(global float *A, global half *B, global double * } // 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-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]], ptr addrspace(1) @@ -77,7 +77,7 @@ kernel void test_mixed_all(global uchar *A, global float *B, global uint *C, glo } // 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-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 @@ -90,7 +90,7 @@ kernel void test_multiple(global uchar *A, global float *B, global uint *C) { } // 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-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() { diff --git a/llvm-spirv/test/negative/unimplemented.spt b/llvm-spirv/test/negative/unimplemented.spt index fab5fe6105159..0f902155fcf39 100644 --- a/llvm-spirv/test/negative/unimplemented.spt +++ b/llvm-spirv/test/negative/unimplemented.spt @@ -2,7 +2,7 @@ 2 Capability Addresses 2 Capability Shader 3 MemoryModel 2 2 -6 EntryPoint 6 2 "foo" +4 EntryPoint 6 2 "foo" 3 Name 3 "res" 2 TypeVoid 12 3 TypeFloat 13 32 diff --git a/llvm-spirv/test/right_shift.spt b/llvm-spirv/test/right_shift.spt index c6ab2d336d1d4..66e04be0e295a 100644 --- a/llvm-spirv/test/right_shift.spt +++ b/llvm-spirv/test/right_shift.spt @@ -4,7 +4,7 @@ 2 Capability Kernel 2 Capability Int64 3 MemoryModel 2 2 -10 EntryPoint 6 1 "shift_right_arithmetic" +9 EntryPoint 6 1 "shift_right_arithmetic" 3 Source 3 102000 3 Name 2 "in" 4 Decorate 3 BuiltIn 28