[llvm-spirv] Cherry pick Khronos changes to expand collection of entr…

…y point interfaces (PR #1334) (#10623)

This PR pulls in the following PR from upstream Khronos
SPIRV-LLVM-Translator repo:
KhronosGroup/SPIRV-LLVM-Translator#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 #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 <[email protected]>

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<SPIRVId>
+LLVMToSPIRVBase::collectEntryPointInterfaces(SPIRVFunction *SF, Function *F) {
+  std::vector<SPIRVId> 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<const Function *> Funcs;
 
@@ -4915,9 +4916,14 @@ void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF,
     }
 
     if (isAnyFunctionReachableFromFunction(F, Funcs)) {
-      SF->addVariable(ValueMap[&GV]);
+      SPIRVWord ModuleVersion = static_cast<SPIRVWord>(BM->getSPIRVVersion());
+      if (AS != SPIRAS_Input && AS != SPIRAS_Output &&
+          ModuleVersion < static_cast<SPIRVWord>(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);
   }
 }
 

llvm-spirv/lib/SPIRV/SPIRVWriter.h

@@ -248,6 +248,8 @@ class LLVMToSPIRVBase : protected BuiltinCallHelper {
       const Function *FS,
       const std::unordered_set<const Function *> Funcs) const;
   void collectInputOutputVariables(SPIRVFunction *SF, Function *F);
+  std::vector<SPIRVId> collectEntryPointInterfaces(SPIRVFunction *BF,
+                                                   Function *F);
 };
 
 class LLVMToSPIRVPass : public PassInfoMixin<LLVMToSPIRVPass> {

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

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 {

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h

@@ -533,6 +533,7 @@ template <Op OC> class SPIRVAnnotation : public SPIRVAnnotationGeneric {
 
 class SPIRVEntryPoint : public SPIRVAnnotation<OpEntryPoint> {
 public:
+  static const SPIRVWord FixedWC = 4;
   SPIRVEntryPoint(SPIRVModule *TheModule, SPIRVExecutionModelKind,
                   SPIRVId TheId, const std::string &TheName,
                   std::vector<SPIRVId> Variables);

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp

@@ -128,20 +128,6 @@ class SPIRVModuleImpl : public SPIRVModule {
   getValueTypes(const std::vector<SPIRVId> &) 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<SPIRVFunction>(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<SPIRVEntry *> &Targets) override;
-  void addEntryPoint(SPIRVExecutionModelKind ExecModel,
-                     SPIRVId EntryPoint) override;
+  void addEntryPoint(SPIRVExecutionModelKind ExecModel, SPIRVId EntryPoint,
+                     const std::string &Name,
+                     const std::vector<SPIRVId> &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<SPIRVGroupDecorateGeneric *> SPIRVGroupDecVec;
   typedef std::vector<SPIRVAsmTargetINTEL *> SPIRVAsmTargetVector;
   typedef std::vector<SPIRVAsmINTEL *> SPIRVAsmVector;
+  typedef std::vector<SPIRVEntryPoint *> SPIRVEntryPointVec;
   typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToInstructionSetMap;
   std::map<SPIRVExtInstSetKind, SPIRVId> ExtInstSetIds;
   typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToBuiltinSetMap;
   typedef std::map<SPIRVExecutionModelKind, SPIRVIdSet> SPIRVExecModelIdSetMap;
-  typedef std::map<SPIRVExecutionModelKind, SPIRVIdVec> SPIRVExecModelIdVecMap;
   typedef std::unordered_map<std::string, SPIRVString *> SPIRVStringMap;
   typedef std::map<SPIRVTypeStruct *, std::vector<std::pair<unsigned, SPIRVId>>>
       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<SPIRVId> &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<SPIRVFunction>(II)->getName(),
-                           M.get<SPIRVFunction>(II)->getVariables());
+  O << MI.EntryPointVec;
 
   for (auto &I : MI.EntryPointVec)
-    for (auto &II : I.second)
-      MI.get<SPIRVFunction>(II)->encodeExecutionModes(O);
+    MI.get<SPIRVFunction>(I->getTargetId())->encodeExecutionModes(O);
 
   O << MI.StringVec;
 

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<std::string> &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<std::string> &getSourceExtension() = 0;
@@ -223,7 +220,9 @@ class SPIRVModule {
                          const std::vector<SPIRVEntry *> &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<SPIRVId> &) = 0;
   virtual SPIRVForward *addForward(SPIRVType *Ty) = 0;
   virtual SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) = 0;
   virtual SPIRVFunction *addFunction(SPIRVFunction *) = 0;

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

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"

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

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

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

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() {