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

MrSidims · jsji · commit 30e6487f4fae · 2023-08-10T19:31:38.000-07: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 KhronosGroup/SPIRV-LLVM-Translator#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 KhronosGroup/SPIRV-LLVM-Translator#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 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> Original commit: KhronosGroup/SPIRV-LLVM-Translator@46285e4
diff --git a/llvm-spirv/lib/SPIRV/SPIRVRegularizeLLVM.cpp b/llvm-spirv/lib/SPIRV/SPIRVRegularizeLLVM.cpp
@@ -427,6 +427,8 @@ bool SPIRVRegularizeLLVMBase::regularize() {
       simplifyBuiltinVarAccesses(&GV);
   }
 
+  // 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()) {
@@ -440,7 +442,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 ||
diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_function_pointers/alias.ll b/llvm-spirv/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/llvm-spirv/test/extensions/INTEL/SPV_INTEL_function_pointers/fp-from-host.ll b/llvm-spirv/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: EntryPoint {{[0-9]+}} [[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/llvm-spirv/test/extensions/INTEL/SPV_INTEL_function_pointers/function-pointer-as-function-arg.ll b/llvm-spirv/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: EntryPoint 6 [[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/llvm-spirv/test/extensions/INTEL/SPV_INTEL_function_pointers/function-pointer.ll b/llvm-spirv/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: EntryPoint 6 [[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/llvm-spirv/test/extensions/INTEL/SPV_INTEL_function_pointers/non-uniform-function-pointer.ll b/llvm-spirv/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: EntryPoint 6 [[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/llvm-spirv/test/extensions/INTEL/SPV_INTEL_function_pointers/select.ll b/llvm-spirv/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: EntryPoint 6 [[#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/llvm-spirv/test/extensions/INTEL/SPV_INTEL_unstructured_loop_controls/FPGAUnstructuredLoopAttr.ll b/llvm-spirv/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: 4 EntryPoint 6 [[FOO:[0-9]+]] "foo"
-; CHECK-SPIRV: 4 EntryPoint 6 [[BOO:[0-9]+]] "boo"
-; CHECK-SPIRV: 4 Name [[ENTRY_1:[0-9]+]] "entry"
-; CHECK-SPIRV: 5 Name [[FOR:[0-9]+]] "for.cond"
-; 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/llvm-spirv/test/mem2reg.cl b/llvm-spirv/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_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/llvm-spirv/test/transcoding/KernelArgTypeInOpString2.ll b/llvm-spirv/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 [[#]] "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/llvm-spirv/test/transcoding/OpenCL/atomic_cmpxchg.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/OpenCL/atomic_legacy.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/OpenCL/atomic_work_item_fence.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/OpenCL/barrier.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/OpenCL/mem_fence.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/OpenCL/sub_group_barrier.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/OpenCL/work_group_barrier.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/SampledImage.cl b/llvm-spirv/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: 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/llvm-spirv/test/transcoding/kernel_arg_type_qual.ll b/llvm-spirv/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 12 "kernel_arg_type_qual.test.volatile,const,,"
+; 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/llvm-spirv/test/transcoding/kernel_query.ll b/llvm-spirv/test/transcoding/kernel_query.ll
@@ -29,14 +29,14 @@ target triple = "spir-unknown-unknown"
 
 %struct.ndrange_t = type { i32 }
 
-; CHECK-SPIRV: EntryPoint {{.*}} [[BlockKer1:[0-9]+]] "__device_side_enqueue_block_invoke_kernel"
-; CHECK-SPIRV: EntryPoint {{.*}} [[BlockKer2:[0-9]+]] "__device_side_enqueue_block_invoke_2_kernel"
-; CHECK-SPIRV: EntryPoint {{.*}} [[BlockKer3:[0-9]+]] "__device_side_enqueue_block_invoke_3_kernel"
-; CHECK-SPIRV: EntryPoint {{.*}} [[BlockKer4:[0-9]+]] "__device_side_enqueue_block_invoke_4_kernel"
-; CHECK-SPIRV: Name [[BlockGlb1:[0-9]+]] "__block_literal_global"
-; CHECK-SPIRV: Name [[BlockGlb2:[0-9]+]] "__block_literal_global.1"
-; CHECK-SPIRV: Name [[BlockGlb3:[0-9]+]] "__block_literal_global.2"
-; CHECK-SPIRV: Name [[BlockGlb4:[0-9]+]] "__block_literal_global.3"
+; CHECK-SPIRV-DAG: Name [[BlockGlb1:[0-9]+]] "__block_literal_global"
+; CHECK-SPIRV-DAG: Name [[BlockGlb2:[0-9]+]] "__block_literal_global.1"
+; CHECK-SPIRV-DAG: Name [[BlockGlb3:[0-9]+]] "__block_literal_global.2"
+; CHECK-SPIRV-DAG: Name [[BlockGlb4:[0-9]+]] "__block_literal_global.3"
+; CHECK-SPIRV-DAG: EntryPoint [[#]] [[BlockKer1:[0-9]+]] "__device_side_enqueue_block_invoke_kernel"
+; CHECK-SPIRV-DAG: EntryPoint [[#]] [[BlockKer2:[0-9]+]] "__device_side_enqueue_block_invoke_2_kernel"
+; CHECK-SPIRV-DAG: EntryPoint [[#]] [[BlockKer3:[0-9]+]] "__device_side_enqueue_block_invoke_3_kernel"
+; CHECK-SPIRV-DAG: EntryPoint [[#]] [[BlockKer4:[0-9]+]] "__device_side_enqueue_block_invoke_4_kernel"
 
 ; CHECK-LLVM: [[BlockTy:%[0-9a-z\.]+]] = type { i32, i32 }
 %1 = type <{ i32, i32 }>
diff --git a/llvm-spirv/test/transcoding/registerallocmode.ll b/llvm-spirv/test/transcoding/registerallocmode.ll
@@ -4,15 +4,15 @@
 ; RUN: spirv-val %t.spv
 ; RUN: llvm-spirv -r -emit-opaque-pointers %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
 
-; CHECK-SPIRV: EntryPoint [[#TMP:]] [[#FUNC0:]] "main_l3"
-; CHECK-SPIRV: EntryPoint [[#TMP:]] [[#FUNC1:]] "main_l6"
-; CHECK-SPIRV: EntryPoint [[#TMP:]] [[#FUNC2:]] "main_l9"
-; CHECK-SPIRV: EntryPoint [[#TMP:]] [[#FUNC3:]] "main_l13"
-; CHECK-SPIRV: EntryPoint [[#TMP:]] [[#FUNC4:]] "main_l19"
+; CHECK-SPIRV: EntryPoint [[#]] [[#FUNC0:]] "main_l3"
+; CHECK-SPIRV: EntryPoint [[#]] [[#FUNC1:]] "main_l6"
+; CHECK-SPIRV: EntryPoint [[#]] [[#FUNC2:]] "main_l9"
+; CHECK-SPIRV: EntryPoint [[#]] [[#FUNC3:]] "main_l13"
+; CHECK-SPIRV: EntryPoint [[#]] [[#FUNC4:]] "main_l19"
 
 ; CHECK-SPIRV: Decorate [[#FUNC0]] UserSemantic "num-thread-per-eu 4"
 ; CHECK-SPIRV: Decorate [[#FUNC1]] UserSemantic "num-thread-per-eu 8"
-; CHECK-SPIRV:  Decorate [[#FUNC2]] UserSemantic "num-thread-per-eu 0"
+; CHECK-SPIRV: Decorate [[#FUNC2]] UserSemantic "num-thread-per-eu 0"
 ; CHECK-SPIRV-NOT: Decorate [[#FUNC3]] UserSemantic
 ; CHECK-SPIRV-NOT: Decorate [[#FUNC4]] UserSemantic
 
diff --git a/llvm-spirv/test/transcoding/spirv-target-types-buffer.ll b/llvm-spirv/test/transcoding/spirv-target-types-buffer.ll
@@ -8,7 +8,7 @@ target triple = "spir-unknown-unknown"
 
 ; CHECK-SPIRV: Capability VectorComputeINTEL
 ; CHECK-SPIRV: Extension "SPV_INTEL_vector_compute"
-; CHECK-SPIRV: EntryPoint {{[0-9]+}} [[#FuncName:]] "foo"
+; CHECK-SPIRV: EntryPoint [[#]] [[#FuncName:]] "foo"
 ; CHECK-SPIRV: Name [[#ParamName:]] "a"
 ; CHECK-SPIRV: TypeVoid  [[#VoidT:]]
 ; CHECK-SPIRV: TypeBufferSurfaceINTEL [[#BufferID:]]
@@ -20,4 +20,4 @@ define spir_kernel void @foo(target("spirv.BufferSurfaceINTEL", 0) %a) #0 {
   ret void
  }
 
-attributes #0 = { noinline norecurse nounwind readnone "VCFunction"}
+attributes #0 = { noinline norecurse nounwind readnone "VCFunction"}

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: EntryPoint {{[0-9]+}} [[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: EntryPoint 6 [[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: EntryPoint 6 [[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: EntryPoint 6 [[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]]`