Don't wrap kernels that are not being called in the module (#2119)

MrSidims · web-flow · commit 46285e43abcc · 2023-08-09T17:46:33.000+02:00
* Don't wrap kernels that are not being called in the module This patch is a result of a reflection about previously merged PR #1149 "add an entry point wrapper around functions (llvm pass)" and is enspired by various reported translator, clang (OpenCL) and Intel GPU drivers issues (see #2029 for reference). While SPIR-V spec states: === *OpName* --//--. This has nosemantic impact and can safely be removed from a module. === yet having EntryPoint function and a function that shares the name via OpName might be confusing by both (old) drivers and programmers, who read the SPIR-V file. This patch prevents generation of the wrapper function when it's not necessary to generate it aka if a kernel function is not called by other kernel. We can do better in other cases as well, for example I have experiments of renaming a wrapped function adding a previx, so it could be distinguished from the actual kernel/entry point, but for now it doesn't pass validation for E2E OpenCL tests. Signed-off-by: Sidorov, Dmitry <dmitry.sidorov@intel.com> * prevent a copy Signed-off-by: Sidorov, Dmitry <dmitry.sidorov@intel.com> This patch is a result of a reflection about previously merged PR #1149 "add an entry point wrapper around functions (llvm pass)" and is enspired by various reported translator, clang (OpenCL) and Intel GPU drivers issues (see #2029 for reference). While SPIR-V spec states: OpName --//--. This has nosemantic impact and can safely be removed from a module. yet having EntryPoint function and a function that shares the name via OpName might be confusing by both not-up-to-date drivers and programmers, who read the SPIR-V file. This patch prevents generation of the wrapper function when it's not necessary to generate it aka if a kernel function is not called by other kernel. Signed-off-by: Sidorov, Dmitry <dmitry.sidorov@intel.com>
diff --git a/lib/SPIRV/SPIRVRegularizeLLVM.cpp b/lib/SPIRV/SPIRVRegularizeLLVM.cpp
@@ -363,9 +363,10 @@ bool SPIRVRegularizeLLVMBase::runRegularizeLLVM(Module &Module) {
 /// Remove entities not representable by SPIR-V
 bool SPIRVRegularizeLLVMBase::regularize() {
   eraseUselessFunctions(M);
-  addKernelEntryPoint(M);
   expandSYCLTypeUsing(M);
 
+  // Kernels called by other kernels
+  std::vector<Function *> CalledKernels;
   for (auto I = M->begin(), E = M->end(); I != E;) {
     Function *F = &(*I++);
     if (F->isDeclaration() && F->use_empty()) {
@@ -379,7 +380,9 @@ bool SPIRVRegularizeLLVMBase::regularize() {
         if (auto *Call = dyn_cast<CallInst>(&II)) {
           Call->setTailCall(false);
           Function *CF = Call->getCalledFunction();
-          if (CF && CF->isIntrinsic()) {
+          if (CF && CF->getCallingConv() == CallingConv::SPIR_KERNEL) {
+            CalledKernels.push_back(CF);
+          } else if (CF && CF->isIntrinsic()) {
             removeFnAttr(Call, Attribute::NoUnwind);
             auto *II = cast<IntrinsicInst>(Call);
             if (II->getIntrinsicID() == Intrinsic::memset ||
@@ -497,19 +500,14 @@ bool SPIRVRegularizeLLVMBase::regularize() {
     }
   }
 
+  addKernelEntryPoint(CalledKernels);
   if (SPIRVDbgSaveRegularizedModule)
     saveLLVMModule(M, RegularizedModuleTmpFile);
   return true;
 }
 
-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);
-  }
+void SPIRVRegularizeLLVMBase::addKernelEntryPoint(
+    const std::vector<Function *> &Work) {
   for (auto &F : Work) {
     // for declarations just make them into SPIR functions.
     F->setCallingConv(CallingConv::SPIR_FUNC);
diff --git a/lib/SPIRV/SPIRVRegularizeLLVM.h b/lib/SPIRV/SPIRVRegularizeLLVM.h
@@ -51,10 +51,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(llvm::Module *M);
+  // SPIR-V disallows functions being entrypoints/kernels and called
+  // OpenCL doesn't. This adds a wrapper around the entry points that are called
+  // by other entry point that later SPIR-V writer renames. Caller of the
+  // function is responsible to fill 'CalledKernels' with such entry points.
+  void addKernelEntryPoint(const std::vector<Function *> &CalledKernels);
 
   /// Some LLVM intrinsics that have no SPIR-V counterpart may be wrapped in
   /// @spirv.llvm_intrinsic_* function. During reverse translation from SPIR-V
diff --git a/lib/SPIRV/SPIRVWriter.cpp b/lib/SPIRV/SPIRVWriter.cpp
@@ -788,10 +788,11 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) {
   BF->setFunctionControlMask(transFunctionControlMask(F));
   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()));
+      // If found, strip the entry point wrapper prefix as the runtime will
+      // be looking for this name
+      StringRef Name = F->getName();
+      (void) Name.consume_front(kSPIRVName::EntrypointPrefix);
+      BM->setName(BF, Name.str());
     } else {
       if (isUniformGroupOperation(F))
         BM->getErrorLog().checkError(
@@ -5329,8 +5330,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 Prefix = kSPIRVName::EntrypointPrefix;
-  std::string Name = F->getName().str().substr(Prefix.size());
+  StringRef FunName = F->getName();
+  (void) FunName.consume_front(kSPIRVName::EntrypointPrefix);
+  std::string Name = FunName.str();
   std::string KernelArgTypesMDStr = std::string(MDName) + "." + Name + ".";
   for (const auto &TyOp : MD->operands())
     KernelArgTypesMDStr += cast<MDString>(TyOp)->getString().str() + ",";
diff --git a/test/entry_point_func.ll b/test/entry_point_func.ll
@@ -12,10 +12,47 @@ define spir_kernel void @testfunction() {
    ret void
 }
 
+define spir_kernel void @callerfunction() {
+   call spir_kernel void @testfunction()
+   call spir_kernel void @testdeclaration()
+   ret void
+}
+
+declare spir_kernel void @testdeclaration()
+
 ; 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]]
+; CHECK-SPIRV: EntryPoint 6 [[#CallerEn:]] "callerfunction"
+; CHECK-SPIRV: EntryPoint 6 [[#TestEn:]] "testfunction"
+; CHECK-SPIRV: Name [[#TestDecl:]] "testdeclaration"
+; CHECK-SPIRV: Name [[#TestFn:]] "testfunction"
+; CHECK-SPIRV: Decorate [[#TestDecl]] LinkageAttributes "testdeclaration" Import
+; CHECK-SPIRV: Decorate [[#TestFn]] LinkageAttributes "testfunction" Export
+
+; CHECK-SPIRV: Function [[#]] [[#TestDecl]] [[#]] [[#]]
+; CHECK-SPIRV-EMPTY:
+; CHECK-SPIRV-NEXT: FunctionEnd
+
+; CHECK-SPIRV: Function [[#]] [[#TestFn]] [[#]] [[#]]
+; CHECK-SPIRV-EMPTY:
+; CHECK-SPIRV-NEXT: Label
+; CHECK-SPIRV-NEXT: Return
+; CHECK-SPIRV-EMPTY:
+; CHECK-SPIRV-NEXT: FunctionEnd
+
+; CHECK-SPIRV: Function [[#]] [[#CallerEn]] [[#]] [[#]]
+; CHECK-SPIRV-EMPTY:
+; CHECK-SPIRV-NEXT: Label
+; CHECK-SPIRV-NEXT: FunctionCall [[#]] [[#]] [[#TestFn]]
+; CHECK-SPIRV-NEXT: FunctionCall [[#]] [[#]] [[#TestDecl]]
+; CHECK-SPIRV-NEXT: Return
+; CHECK-SPIRV-EMPTY:
+; CHECK-SPIRV-NEXT: FunctionEnd
+
+
+; CHECK-SPIRV: Function [[#]] [[#TestEn]] [[#]] [[#]]
+; CHECK-SPIRV-EMPTY:
+; CHECK-SPIRV-NEXT: Label
+; CHECK-SPIRV-NEXT: FunctionCall [[#]] [[#]] [[#TestFn]]
+; CHECK-SPIRV-NEXT: Return
+; CHECK-SPIRV-EMPTY:
+; CHECK-SPIRV-NEXT: FunctionEnd
diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_argument_interfaces/sycl-kernel-arg-annotation.ll b/test/extensions/INTEL/SPV_INTEL_fpga_argument_interfaces/sycl-kernel-arg-annotation.ll
@@ -53,7 +53,6 @@ entry:
 ; CHECK-SPIRV: Capability FPGAArgumentInterfacesINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_fpga_argument_interfaces"
 ; CHECK-SPIRV: Extension "SPV_INTEL_fpga_buffer_location"
-; CHECK-SPIRV-DAG:  Name [[IDS:[0-9]+]] "_arg_p"
 ; CHECK-SPIRV-DAG:  Name [[ID:[0-9]+]] "_arg_p"
 ; CHECK-SPIRV:  Decorate [[ID]] Alignment 4
 ; CHECK-SPIRV:  Decorate [[ID]] MMHostInterfaceAddressWidthINTEL 32
diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/alias.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/alias.ll
@@ -10,11 +10,10 @@ target triple = "spir64-unknown-unknown"
 ; when used since they can't be translated directly.
 
 ; CHECK-SPIRV-DAG: Name [[#FOO:]] "foo"
-; CHECK-SPIRV-DAG: Name [[#BAR:]] "bar"
+; CHECK-SPIRV-DAG: EntryPoint [[#]] [[#BAR:]] "bar"
 ; CHECK-SPIRV-DAG: Name [[#Y:]] "y"
 ; CHECK-SPIRV-DAG: Name [[#FOOPTR:]] "foo.alias"
 ; CHECK-SPIRV-DAG: Decorate [[#FOO]] LinkageAttributes "foo" Export
-; CHECK-SPIRV-DAG: Decorate [[#BAR]] LinkageAttributes "bar" Export
 ; CHECK-SPIRV-DAG: TypeInt [[#I32:]] 32 0
 ; CHECK-SPIRV-DAG: TypeInt [[#I64:]] 64 0
 ; CHECK-SPIRV-DAG: TypeFunction [[#FOO_TYPE:]] [[#I32]] [[#I32]]
diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/fp-from-host.ll b/test/extensions/INTEL/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: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: EntryPoint [[#]] [[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]]
diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/function-pointer-as-function-arg.ll b/test/extensions/INTEL/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: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: EntryPoint [[#]] [[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]]
diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/function-pointer.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/function-pointer.ll
@@ -19,7 +19,7 @@
 ;
 ; CHECK-SPIRV: Capability FunctionPointersINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
-; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: EntryPoint [[#]] [[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]]
diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/non-uniform-function-pointer.ll b/test/extensions/INTEL/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: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: EntryPoint [[#]] [[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]]
diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/select.ll b/test/extensions/INTEL/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: Name [[#KERNEL_ID:]] "_ZTS6kernel"
+; CHECK-SPIRV-DAG: EntryPoint [[#]] [[#KERNEL_ID:]] "_ZTS6kernel"
 ; CHECK-SPIRV-DAG: Name [[#BAR:]] "_Z3barii"
 ; CHECK-SPIRV-DAG: Name [[#BAZ:]] "_Z3bazii"
 ; CHECK-SPIRV: TypeInt [[#INT32:]] 32
diff --git a/test/extensions/INTEL/SPV_INTEL_unstructured_loop_controls/FPGAUnstructuredLoopAttr.ll b/test/extensions/INTEL/SPV_INTEL_unstructured_loop_controls/FPGAUnstructuredLoopAttr.ll
@@ -5,34 +5,34 @@
 ; RUN: llvm-spirv -r -emit-opaque-pointers %t.spv -o %t.rev.bc
 ; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
 
-; CHECK-SPIRV: 2 Capability UnstructuredLoopControlsINTEL
-; 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: 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"
+; CHECK-SPIRV: Capability UnstructuredLoopControlsINTEL
+; CHECK-SPIRV: Capability FPGALoopControlsINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_fpga_loop_controls"
+; CHECK-SPIRV: Extension "SPV_INTEL_unstructured_loop_controls"
+; CHECK-SPIRV: EntryPoint [[#]] [[FOO:[0-9]+]] "foo"
+; CHECK-SPIRV: EntryPoint [[#]] [[BOO:[0-9]+]] "boo"
+; CHECK-SPIRV: Name [[ENTRY_1:[0-9]+]] "entry"
+; CHECK-SPIRV: Name [[FOR:[0-9]+]] "for.cond"
+; CHECK-SPIRV: Name [[ENTRY_2:[0-9]+]] "entry"
+; CHECK-SPIRV: Name [[WHILE:[0-9]+]] "while.body"
 
-; CHECK-SPIRV: 5 Function 2 [[FOO]] {{[0-9]+}} {{[0-9]+}}
-; CHECK-SPIRV: 2 Label [[ENTRY_1]]
-; CHECK-SPIRV: 2 Branch [[FOR]]
-; CHECK-SPIRV: 2 Label [[FOR]]
+; CHECK-SPIRV: Function [[#]] [[FOO]] {{[0-9]+}} {{[0-9]+}}
+; CHECK-SPIRV: Label [[ENTRY_1]]
+; CHECK-SPIRV: Branch [[FOR]]
+; CHECK-SPIRV: Label [[FOR]]
 ; Per SPIR-V spec extension INTEL/SPV_INTEL_fpga_loop_controls,
 ; LoopControlMaxConcurrencyINTELMask = 0x20000 (131072)
-; CHECK-SPIRV: 3 LoopControlINTEL 131072 2
-; CHECK-SPIRV-NEXT: 2 Branch [[FOR]]
+; CHECK-SPIRV: LoopControlINTEL 131072 2
+; CHECK-SPIRV-NEXT: Branch [[FOR]]
 
-; CHECK-SPIRV: 5 Function 2 [[BOO]] {{[0-9]+}} {{[0-9]+}}
-; CHECK-SPIRV: 2 Label [[ENTRY_2]]
-; CHECK-SPIRV: 2 Branch [[WHILE]]
-; CHECK-SPIRV: 2 Label [[WHILE]]
+; CHECK-SPIRV: Function [[#]] [[BOO]] {{[0-9]+}} {{[0-9]+}}
+; CHECK-SPIRV: Label [[ENTRY_2]]
+; CHECK-SPIRV: Branch [[WHILE]]
+; CHECK-SPIRV: Label [[WHILE]]
 ; Per SPIR-V spec extension INTEL/SPV_INTEL_fpga_loop_controls,
 ; LoopControlInitiationIntervalINTELMask = 0x10000 (65536)
-; CHECK-SPIRV: 3 LoopControlINTEL 65536 2
-; CHECK-SPIRV-NEXT: 2 Branch [[WHILE]]
+; CHECK-SPIRV: LoopControlINTEL 65536 2
+; CHECK-SPIRV-NEXT: Branch [[WHILE]]
 
 ; ModuleID = 'infinite.cl'
 source_filename = "infinite.cl"
diff --git a/test/mem2reg.cl b/test/mem2reg.cl
@@ -1,11 +1,10 @@
 // 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-WO
+// RUN: llvm-dis < %t.bc | FileCheck %s --check-prefixes=CHECK,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-W
-// CHECK-W-LABEL: spir_func void @foo
+// RUN: llvm-dis < %t.opt.bc | FileCheck %s --check-prefixes=CHECK,CHECK-W
+// CHECK-LABEL: spir_kernel void @foo
 // CHECK-W-NOT: alloca
-// CHECK-WO-LABEL: spir_kernel void @foo
 // CHECK-WO: alloca
 __kernel void foo(__global int *a) {
     *a = *a + 1;
diff --git a/test/transcoding/KernelArgTypeInOpString2.ll b/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 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-SPIRV-WORKAROUND: String [[#]] "kernel_arg_type.foo.cl::tt::vec<float, 4>*,"
+; CHECK-SPIRV-WORKAROUND-NEGATIVE-NOT: String [[#]] "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>*"}
diff --git a/test/transcoding/OpenCL/atomic_cmpxchg.cl b/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: Name [[TEST:[0-9]+]] "test_atomic_cmpxchg"
+// CHECK-SPIRV: EntryPoint [[#]] [[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]]
 //
diff --git a/test/transcoding/OpenCL/atomic_legacy.cl b/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: Name [[TEST:[0-9]+]] "test_legacy_atomics"
+// CHECK-SPIRV: EntryPoint [[#]] [[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]]
 //
diff --git a/test/transcoding/OpenCL/atomic_work_item_fence.cl b/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: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
+// CHECK-SPIRV: EntryPoint [[#]] [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
 // CHECK-SPIRV: TypeInt [[UINT:[0-9]+]] 32 0
 //
 // 0x0 Relaxed + 0x100 WorkgroupMemory
diff --git a/test/transcoding/OpenCL/barrier.cl b/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: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
+// CHECK-SPIRV: EntryPoint [[#]] [[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
diff --git a/test/transcoding/OpenCL/mem_fence.cl b/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: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_mem_fence_const_flags"
+// CHECK-SPIRV: EntryPoint [[#]] [[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
diff --git a/test/transcoding/OpenCL/sub_group_barrier.cl b/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: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
+// CHECK-SPIRV: EntryPoint [[#]] [[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
diff --git a/test/transcoding/OpenCL/work_group_barrier.cl b/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: Name [[TEST_CONST_FLAGS:[0-9]+]] "test_barrier_const_flags"
+// CHECK-SPIRV: EntryPoint [[#]] [[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
diff --git a/test/transcoding/SampledImage.cl b/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: Name [[sample_kernel_float:[0-9]+]] "sample_kernel_float"
-// CHECK-SPIRV: Name [[sample_kernel_int:[0-9]+]] "sample_kernel_int"
+// CHECK-SPIRV: EntryPoint [[#]] [[sample_kernel_float:[0-9]+]] "sample_kernel_float"
+// CHECK-SPIRV: EntryPoint [[#]] [[sample_kernel_int:[0-9]+]] "sample_kernel_int"
 
 // CHECK-SPIRV: TypeSampler [[TypeSampler:[0-9]+]]
 // CHECK-SPIRV: TypeSampledImage [[SampledImageTy:[0-9]+]]
diff --git a/test/transcoding/kernel_arg_type_qual.ll b/test/transcoding/kernel_arg_type_qual.ll
@@ -13,10 +13,10 @@ 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 18 "kernel_arg_type_qual.test.volatile,const,,"
-; CHECK-SPIRV: Name [[ARG:1[0-9]+]] "g"
+; CHECK-SPIRV: String [[#]] "kernel_arg_type_qual.test.volatile,const,,"
+; CHECK-SPIRV: Name [[ARG:[0-9]+]] "g"
 ; CHECK-SPIRV: Decorate [[ARG]] Volatile
-; CHECK-SPIRV-NEGATIVE-NOT: String 12 "kernel_arg_type_qual.test.volatile,const,,"
+; CHECK-SPIRV-NEGATIVE-NOT: String [[#]] "kernel_arg_type_qual.test.volatile,const,,"
 
 ; CHECK-LLVM-WORKAROUND: !kernel_arg_type_qual ![[QUAL:[0-9]+]]
 ; CHECK-LLVM-WORKAROUND: ![[QUAL]] = !{!"volatile", !"const", !""}
diff --git a/test/transcoding/kernel_query.ll b/test/transcoding/kernel_query.ll
diff --git a/test/transcoding/registerallocmode.ll b/test/transcoding/registerallocmode.ll
diff --git a/test/transcoding/spirv-target-types-buffer.ll b/test/transcoding/spirv-target-types-buffer.ll

Original file line number	Diff line number	Diff line change
`@@ -17,7 +17,7 @@`
`17`	`17`	`; CHECK-SPIRV: Capability FunctionPointersINTEL`
`18`	`18`	`; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"`
`19`	`19`	`;`
`20`		`-; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"`
	`20`	`+; CHECK-SPIRV: EntryPoint [[#]] [[KERNEL_ID:[0-9]+]] "test"`
`21`	`21`	`; CHECK-SPIRV: TypeInt [[INT32_TYPE_ID:[0-9]+]] 32`
`22`	`22`	`; CHECK-SPIRV: TypePointer [[INT_PTR:[0-9]+]] 5 [[INT32_TYPE_ID]]`
`23`	`23`	`; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[INT32_TYPE_ID]] [[INT32_TYPE_ID]]`
Original file line number	Diff line number	Diff line change
`@@ -33,7 +33,7 @@`
`33`	`33`	`; CHECK-SPIRV: Capability FunctionPointersINTEL`
`34`	`34`	`; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"`
`35`	`35`	`;`
`36`		`-; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"`
	`36`	`+; CHECK-SPIRV: EntryPoint [[#]] [[KERNEL_ID:[0-9]+]] "test"`
`37`	`37`	`; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9]+]] 32`
`38`	`38`	`; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]`
`39`	`39`	`; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]`
Original file line number	Diff line number	Diff line change
`@@ -19,7 +19,7 @@`
`19`	`19`	`;`
`20`	`20`	`; CHECK-SPIRV: Capability FunctionPointersINTEL`
`21`	`21`	`; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"`
`22`		`-; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"`
	`22`	`+; CHECK-SPIRV: EntryPoint [[#]] [[KERNEL_ID:[0-9]+]] "test"`
`23`	`23`	`; CHECK-SPIRV: TypeInt [[TYPE_INT_ID:[0-9]+]]`
`24`	`24`	`; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT_ID]] [[TYPE_INT_ID]]`
`25`	`25`	`; CHECK-SPIRV: TypePointer [[FOO_PTR_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]`
Original file line number	Diff line number	Diff line change
`@@ -29,7 +29,7 @@`
`29`	`29`	`; CHECK-SPIRV: Capability FunctionPointersINTEL`
`30`	`30`	`; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"`
`31`	`31`	`;`
`32`		`-; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"`
	`32`	`+; CHECK-SPIRV: EntryPoint [[#]] [[KERNEL_ID:[0-9]+]] "test"`
`33`	`33`	`; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9+]]] 32`
`34`	`34`	`; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]`
`35`	`35`	`; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]`