Add an entry point wrapper around functions (llvm pass) (KhronosGroup…

…#1149)

SPIR-V spec states:
"It is invalid for any function to be targeted by both an OpEntryPoint instruction
and an OpFunctionCall instruction."

In order to satisfy SPIR-V that entrypoints and functions
must be different, this introduces an entrypoint wrapper around
functions at the LLVM IR level, then fixes up a few things like
naming at the SPIRV translation.

lib/SPIRV/SPIRVInternal.h

@@ -377,6 +377,7 @@ const static char TranslateOCLMemScope[] = "__translate_ocl_memory_scope";
 const static char TranslateSPIRVMemOrder[] = "__translate_spirv_memory_order";
 const static char TranslateSPIRVMemScope[] = "__translate_spirv_memory_scope";
 const static char TranslateSPIRVMemFence[] = "__translate_spirv_memory_fence";
+const static char EntrypointPrefix[] = "__spirv_entry_";
 } // namespace kSPIRVName
 
 namespace kSPIRVPostfix {

lib/SPIRV/SPIRVReader.cpp

@@ -2768,6 +2768,24 @@ Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) {
     return Loc->second;
 
   auto IsKernel = isKernel(BF);
+
+  if (IsKernel) {
+    // search for a previous function with the same name
+    // upgrade it to a kernel and drop this if it's found
+    for (auto &I : FuncMap) {
+      auto BFName = I.getFirst()->getName();
+      if (BF->getName() == BFName) {
+        auto *F = I.getSecond();
+        F->setCallingConv(CallingConv::SPIR_KERNEL);
+        F->setLinkage(GlobalValue::ExternalLinkage);
+        F->setDSOLocal(false);
+        F = cast<Function>(mapValue(BF, F));
+        mapFunction(BF, F);
+        return F;
+      }
+    }
+  }
+
   auto Linkage = IsKernel ? GlobalValue::ExternalLinkage : transLinkageType(BF);
   FunctionType *FT = dyn_cast<FunctionType>(transType(BF->getFunctionType()));
   std::string FuncName = BF->getName();

lib/SPIRV/SPIRVRegularizeLLVM.cpp

@@ -39,6 +39,7 @@
 
 #include "OCLUtil.h"
 #include "SPIRVInternal.h"
+#include "SPIRVMDWalker.h"
 #include "libSPIRV/SPIRVDebug.h"
 
 #include "llvm/ADT/StringExtras.h" // llvm::isDigit
@@ -72,6 +73,11 @@ class SPIRVRegularizeLLVMBase {
   // Lower functions
   bool regularize();
 
+  // SPIR-V disallows functions being entrypoints and called
+  // LLVM doesn't. This adds a wrapper around the entry point
+  // that later SPIR-V writer renames.
+  void addKernelEntryPoint(Module *M);
+
   /// Erase cast inst of function and replace with the function.
   /// Assuming F is a SPIR-V builtin function with op code \param OC.
   void lowerFuncPtr(Function *F, Op OC);
@@ -437,6 +443,7 @@ bool SPIRVRegularizeLLVMBase::runRegularizeLLVM(Module &Module) {
 bool SPIRVRegularizeLLVMBase::regularize() {
   eraseUselessFunctions(M);
   lowerFuncPtr(M);
+  addKernelEntryPoint(M);
 
   for (auto I = M->begin(), E = M->end(); I != E;) {
     Function *F = &(*I++);
@@ -605,6 +612,69 @@ void SPIRVRegularizeLLVMBase::lowerFuncPtr(Module *M) {
     lowerFuncPtr(I.first, I.second);
 }
 
+void SPIRVRegularizeLLVMBase::addKernelEntryPoint(Module *M) {
+  std::vector<Function *> Work;
+
+  // Get a list of all functions that have SPIR kernel calling conv
+  for (auto &F : *M) {
+    if (F.getCallingConv() == CallingConv::SPIR_KERNEL)
+      Work.push_back(&F);
+  }
+  for (auto &F : Work) {
+    // for declarations just make them into SPIR functions.
+    F->setCallingConv(CallingConv::SPIR_FUNC);
+    if (F->isDeclaration())
+      continue;
+
+    // Otherwise add a wrapper around the function to act as an entry point.
+    FunctionType *FType = F->getFunctionType();
+    std::string WrapName =
+        kSPIRVName::EntrypointPrefix + static_cast<std::string>(F->getName());
+    Function *WrapFn =
+        getOrCreateFunction(M, F->getReturnType(), FType->params(), WrapName);
+
+    auto *CallBB = BasicBlock::Create(M->getContext(), "", WrapFn);
+    IRBuilder<> Builder(CallBB);
+
+    Function::arg_iterator DestI = WrapFn->arg_begin();
+    for (const Argument &I : F->args()) {
+      DestI->setName(I.getName());
+      DestI++;
+    }
+    SmallVector<Value *, 1> Args;
+    for (Argument &I : WrapFn->args()) {
+      Args.emplace_back(&I);
+    }
+    auto *CI = CallInst::Create(F, ArrayRef<Value *>(Args), "", CallBB);
+    CI->setCallingConv(F->getCallingConv());
+    CI->setAttributes(F->getAttributes());
+
+    // copy over all the metadata (should it be removed from F?)
+    SmallVector<std::pair<unsigned, MDNode *>> MDs;
+    F->getAllMetadata(MDs);
+    WrapFn->setAttributes(F->getAttributes());
+    for (auto MD = MDs.begin(), End = MDs.end(); MD != End; ++MD) {
+      WrapFn->addMetadata(MD->first, *MD->second);
+    }
+    WrapFn->setCallingConv(CallingConv::SPIR_KERNEL);
+    WrapFn->setLinkage(llvm::GlobalValue::InternalLinkage);
+
+    Builder.CreateRet(F->getReturnType()->isVoidTy() ? nullptr : CI);
+
+    // Have to find the spir-v metadata for execution mode and transfer it to
+    // the wrapper.
+    if (auto NMD = SPIRVMDWalker(*M).getNamedMD(kSPIRVMD::ExecutionMode)) {
+      while (!NMD.atEnd()) {
+        Function *MDF = nullptr;
+        auto N = NMD.nextOp(); /* execution mode MDNode */
+        N.get(MDF);
+        if (MDF == F)
+          N.M->replaceOperandWith(0, ValueAsMetadata::get(WrapFn));
+      }
+    }
+  }
+}
+
 } // namespace SPIRV
 
 INITIALIZE_PASS(SPIRVRegularizeLLVMLegacy, "spvregular",

lib/SPIRV/SPIRVWriter.cpp

@@ -638,8 +638,15 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) {
   SPIRVFunction *BF =
       static_cast<SPIRVFunction *>(mapValue(F, BM->addFunction(BFT)));
   BF->setFunctionControlMask(transFunctionControlMask(F));
-  if (F->hasName())
-    BM->setName(BF, F->getName().str());
+  if (F->hasName()) {
+    if (isKernel(F)) {
+      /* strip the prefix as the runtime will be looking for this name */
+      std::string Prefix = kSPIRVName::EntrypointPrefix;
+      std::string Name = F->getName().str();
+      BM->setName(BF, Name.substr(Prefix.size()));
+    } else
+      BM->setName(BF, F->getName().str());
+  }
   if (!isKernel(F) && F->getLinkage() != GlobalValue::InternalLinkage)
     BF->setLinkageType(transLinkageType(F));
 
@@ -3735,7 +3742,7 @@ void LLVMToSPIRVBase::transFunction(Function *I) {
 
   if (isKernel(I)) {
     auto Interface = collectEntryPointInterfaces(BF, I);
-    BM->addEntryPoint(ExecutionModelKernel, BF->getId(), I->getName().str(),
+    BM->addEntryPoint(ExecutionModelKernel, BF->getId(), BF->getName(),
                       Interface);
   }
 }
@@ -4064,8 +4071,9 @@ bool LLVMToSPIRVBase::transMetadata() {
 // Work around to translate kernel_arg_type and kernel_arg_type_qual metadata
 static void transKernelArgTypeMD(SPIRVModule *BM, Function *F, MDNode *MD,
                                  std::string MDName) {
-  std::string KernelArgTypesMDStr =
-      std::string(MDName) + "." + F->getName().str() + ".";
+  std::string Prefix = kSPIRVName::EntrypointPrefix;
+  std::string Name = F->getName().str().substr(Prefix.size());
+  std::string KernelArgTypesMDStr = std::string(MDName) + "." + Name + ".";
   for (const auto &TyOp : MD->operands())
     KernelArgTypesMDStr += cast<MDString>(TyOp)->getString().str() + ",";
   BM->getString(KernelArgTypesMDStr);

test/entry_point_func.ll

@@ -0,0 +1,21 @@
+;; Test to check that an LLVM spir_kernel gets translated into an
+;; Entrypoint wrapper and Function with LinkageAttributes
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -o - -spirv-text | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val %t.spv
+
+target datalayout = "e-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 = "spir64-unknown-unknown"
+
+define spir_kernel void @testfunction() {
+   ret void
+}
+
+; Check there is an entrypoint and a function produced.
+; CHECK-SPIRV: EntryPoint 6 [[EP:[0-9]+]] "testfunction"
+; CHECK-SPIRV: Name [[FUNC:[0-9]+]] "testfunction"
+; CHECK-SPIRV: Decorate [[FUNC]] LinkageAttributes "testfunction" Export
+; CHECK-SPIRV: Function 2 [[FUNC]] 0 3
+; CHECK-SPIRV: Function 2 [[EP]] 0 3
+; CHECK-SPIRV: FunctionCall 2 8 [[FUNC]]

test/mem2reg.cl

@@ -1,10 +1,11 @@
 // RUN: %clang_cc1 -O0 -S -triple spir-unknown-unknown -cl-std=CL2.0 -x cl -disable-O0-optnone %s -emit-llvm-bc -o %t.bc
 // RUN: llvm-spirv -s %t.bc
-// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefixes=CHECK,CHECK-WO
+// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefixes=CHECK-WO
 // RUN: llvm-spirv -s -spirv-mem2reg %t.bc -o %t.opt.bc
-// RUN: llvm-dis < %t.opt.bc | FileCheck %s --check-prefixes=CHECK,CHECK-W
-// CHECK-LABEL: spir_kernel void @foo
+// RUN: llvm-dis < %t.opt.bc | FileCheck %s --check-prefixes=CHECK-W
+// CHECK-W-LABEL: spir_func void @foo
 // CHECK-W-NOT: alloca i32
+// CHECK-WO-LABEL: spir_kernel void @foo
 // CHECK-WO: alloca i32
 __kernel void foo(__global int *a) {
     *a = *a + 1;

test/transcoding/FPGAUnstructuredLoopAttr.ll

@@ -9,10 +9,10 @@
 ; CHECK-SPIRV: 2 Capability FPGALoopControlsINTEL
 ; CHECK-SPIRV: 9 Extension "SPV_INTEL_fpga_loop_controls"
 ; CHECK-SPIRV: 11 Extension "SPV_INTEL_unstructured_loop_controls"
-; CHECK-SPIRV: 4 EntryPoint 6 [[FOO:[0-9]+]] "foo"
-; CHECK-SPIRV: 4 EntryPoint 6 [[BOO:[0-9]+]] "boo"
+; CHECK-SPIRV: 3 Name [[FOO:[0-9]+]] "foo"
 ; CHECK-SPIRV: 4 Name [[ENTRY_1:[0-9]+]] "entry"
 ; CHECK-SPIRV: 5 Name [[FOR:[0-9]+]] "for.cond"
+; CHECK-SPIRV: 3 Name [[BOO:[0-9]+]] "boo"
 ; CHECK-SPIRV: 4 Name [[ENTRY_2:[0-9]+]] "entry"
 ; CHECK-SPIRV: 5 Name [[WHILE:[0-9]+]] "while.body"
 

test/transcoding/KernelArgTypeInOpString.ll

@@ -39,8 +39,8 @@
 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"
 target triple = "spir-unknown-unknown"
 
-; CHECK-SPIRV-WORKAROUND: String 14 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
-; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 14 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
+; CHECK-SPIRV-WORKAROUND: String 20 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
+; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 20 "kernel_arg_type.foo.image_kernel_data*,myInt,struct struct_name*,"
 
 ; CHECK-LLVM-WORKAROUND: !kernel_arg_type [[TYPE:![0-9]+]]
 ; CHECK-LLVM-WORKAROUND: [[TYPE]] = !{!"image_kernel_data*", !"myInt", !"struct struct_name*"}

test/transcoding/KernelArgTypeInOpString2.ll

@@ -41,8 +41,8 @@
 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"
 target triple = "spir"
 
-; CHECK-SPIRV-WORKAROUND: String 17 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
-; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 17 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
+; CHECK-SPIRV-WORKAROUND: String 21 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
+; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String 21 "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
 
 ; CHECK-LLVM-WORKAROUND: !kernel_arg_type [[TYPE:![0-9]+]]
 ; CHECK-LLVM-WORKAROUND: [[TYPE]] = !{!"cl::tt::vec<float, 4>*"}

test/transcoding/OpenCL/atomic_cmpxchg.cl

@@ -17,7 +17,7 @@ __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
   atomic_cmpxchg(up, ucmp, uval);
 }
 
-// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST:[0-9]+]] "test_atomic_cmpxchg"
+// CHECK-SPIRV: Name [[TEST:[0-9]+]] "test_atomic_cmpxchg"
 // CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
 // CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
 //

test/transcoding/OpenCL/atomic_legacy.cl

@@ -13,7 +13,7 @@ __kernel void test_legacy_atomics(__global int *p, int val) {
   atomic_add(p, val);   // from OpenCL C 1.1
 }
 
-// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST:[0-9]+]] "test_legacy_atomics"
+// CHECK-SPIRV: Name [[TEST:[0-9]+]] "test_legacy_atomics"
 // CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
 // CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
 //

test/transcoding/OpenCL/atomic_work_item_fence.cl

@@ -23,7 +23,7 @@ __kernel void test_mem_fence_non_const_flags(cl_mem_fence_flags flags, memory_or
   // atomic_work_item_fence(flags, order, scope);
 }
 
-// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
+// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
 // CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
 //
 // 0x0 Relaxed + 0x100 WorkgroupMemory

test/transcoding/OpenCL/barrier.cl

@@ -28,7 +28,7 @@ __kernel void test_barrier_non_const_flags(cl_mem_fence_flags flags) {
   // barrier(flags);
 }
 
-// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
+// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
 // CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
 //
 // In SPIR-V, barrier is represented as OpControlBarrier [3] and OpenCL

test/transcoding/OpenCL/mem_fence.cl

@@ -34,7 +34,7 @@ __kernel void test_mem_fence_non_const_flags(cl_mem_fence_flags flags) {
   // mem_fence(flags);
 }
 
-// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
+// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
 // CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
 //
 // In SPIR-V, mem_fence is represented as OpMemoryBarrier [2] and OpenCL

test/transcoding/OpenCL/sub_group_barrier.cl

@@ -31,7 +31,7 @@ __kernel void test_barrier_non_const_flags(cl_mem_fence_flags flags, memory_scop
   // sub_group_barrier(flags, scope);
 }
 
-// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
+// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
 // CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
 //
 // In SPIR-V, barrier is represented as OpControlBarrier [2] and OpenCL

test/transcoding/OpenCL/work_group_barrier.cl

@@ -33,7 +33,7 @@ __kernel void test_barrier_non_const_flags(cl_mem_fence_flags flags, memory_scop
   // work_group_barrier(flags, scope);
 }
 
-// CHECK-SPIRV: EntryPoint {{[0-9]+}} [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
+// CHECK-SPIRV: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
 // CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
 //
 // In SPIR-V, barrier is represented as OpControlBarrier [2] and OpenCL

test/transcoding/SPV_INTEL_function_pointers/fp-from-host.ll

@@ -17,7 +17,7 @@
 ; CHECK-SPIRV: Capability FunctionPointersINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
 ;
-; CHECK-SPIRV: EntryPoint {{[0-9]+}} [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
 ; CHECK-SPIRV: TypeInt [[INT32_TYPE_ID:[0-9]+]] 32
 ; CHECK-SPIRV: TypePointer [[INT_PTR:[0-9]+]] 5 [[INT32_TYPE_ID]]
 ; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[INT32_TYPE_ID]] [[INT32_TYPE_ID]]

test/transcoding/SPV_INTEL_function_pointers/function-pointer-as-function-arg.ll

@@ -33,7 +33,7 @@
 ; CHECK-SPIRV: Capability FunctionPointersINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
 ;
-; CHECK-SPIRV: EntryPoint 6 [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
 ; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9]+]] 32
 ; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]
 ; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]

test/transcoding/SPV_INTEL_function_pointers/function-pointer.ll

@@ -19,7 +19,7 @@
 ;
 ; CHECK-SPIRV: Capability FunctionPointersINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
-; CHECK-SPIRV: EntryPoint 6 [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
 ; CHECK-SPIRV: TypeInt [[TYPE_INT_ID:[0-9]+]]
 ; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT_ID]] [[TYPE_INT_ID]]
 ; CHECK-SPIRV: TypePointer [[FOO_PTR_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]

test/transcoding/SPV_INTEL_function_pointers/non-uniform-function-pointer.ll

@@ -29,7 +29,7 @@
 ; CHECK-SPIRV: Capability FunctionPointersINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
 ;
-; CHECK-SPIRV: EntryPoint 6 [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
 ; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9+]]] 32
 ; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]
 ; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]

test/transcoding/SPV_INTEL_function_pointers/select.ll

@@ -6,7 +6,7 @@
 ; RUN: llvm-dis %t.r.bc -o %t.r.ll
 ; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
 
-; CHECK-SPIRV: EntryPoint 6 [[#KERNEL_ID:]] "_ZTS6kernel"
+; CHECK-SPIRV: Name [[#KERNEL_ID:]] "_ZTS6kernel"
 ; CHECK-SPIRV-DAG: Name [[#BAR:]] "_Z3barii"
 ; CHECK-SPIRV-DAG: Name [[#BAZ:]] "_Z3bazii"
 ; CHECK-SPIRV: TypeInt [[#INT32:]] 32

test/transcoding/SPV_INTEL_joint_matrix/joint_matrix.ll

@@ -13,7 +13,7 @@
 
 ; CHECK-SPIRV: Capability JointMatrixINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_joint_matrix"
-; CHECK-SPIRV: EntryPoint 6 [[#Kernel:]]
+; CHECK-SPIRV: Name [[#Kernel:]] "_ZTSZ4mainE11matrix_test"
 
 ; CHECK-SPIRV-DAG: TypeInt [[#ShortTy:]] 16 0
 ; CHECK-SPIRV-DAG: TypeInt [[#CharTy:]] 8 0

test/transcoding/SampledImage.cl

@@ -27,8 +27,8 @@ void sample_kernel_int(image2d_t input, float2 coords, global int4 *results, sam
 }
 
 // CHECK-SPIRV: Capability LiteralSampler
-// CHECK-SPIRV: EntryPoint 6 [[sample_kernel_float:[0-9]+]] "sample_kernel_float"
-// CHECK-SPIRV: EntryPoint 6 [[sample_kernel_int:[0-9]+]] "sample_kernel_int"
+// CHECK-SPIRV: Name [[sample_kernel_float:[0-9]+]] "sample_kernel_float"
+// CHECK-SPIRV: Name [[sample_kernel_int:[0-9]+]] "sample_kernel_int"
 
 // CHECK-SPIRV: TypeSampler [[TypeSampler:[0-9]+]]
 // CHECK-SPIRV: TypeSampledImage [[SampledImageTy:[0-9]+]]
@@ -81,4 +81,4 @@ void sample_kernel_int(image2d_t input, float2 coords, global int4 *results, sam
 // CHECK-SPIRV: ImageSampleExplicitLod {{.*}} [[SampledImage6]]
 // CHECK-LLVM: call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %1, <2 x float> %coords)
 // CHECK-SPV-IR: call spir_func %spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS133__spirv_Image__void_1_0_0_0_0_0_0PU3AS215__spirv_Sampler(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %input, %spirv.Sampler addrspace(2)* %1)
-// CHECK-SPV-IR: call spir_func <4 x i32> @_Z36__spirv_ImageSampleExplicitLod_Rint4PU3AS140__spirv_SampledImage__void_1_0_0_0_0_0_0Dv2_fif(%spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* %TempSampledImage6, <2 x float> %coords, i32 2, float 0.000000e+00)
+// CHECK-SPV-IR: call spir_func <4 x i32> @_Z36__spirv_ImageSampleExplicitLod_Rint4PU3AS140__spirv_SampledImage__void_1_0_0_0_0_0_0Dv2_fif(%spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* %TempSampledImage6, <2 x float> %coords, i32 2, float 0.000000e+00)

test/transcoding/kernel_arg_type_qual.ll

@@ -13,8 +13,8 @@ source_filename = "test.cl"
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "spir64-unknown-unknown."
 
-; CHECK-SPIRV: String 12 "kernel_arg_type_qual.test.volatile,const,,"
-; CHECK-SPIRV: Name [[ARG:[0-9]+]] "g"
+; CHECK-SPIRV: String 18 "kernel_arg_type_qual.test.volatile,const,,"
+; CHECK-SPIRV: Name [[ARG:1[0-9]+]] "g"
 ; CHECK-SPIRV: Decorate [[ARG]] Volatile
 ; CHECK-SPIRV-NEGATIVE-NOT: String 12 "kernel_arg_type_qual.test.volatile,const,,"