[sycl-post-link] Do not launch GlobalDCE for spec constants cleanup (#3851)

The pass removes `specialization_id` global variables, because they are
unused in a module after `SpecConstantsPass` finishes its work, but they
are still referenced by debug metadata (if it is present) and such
crippled module can't be later handled by llvm-dis or llvm-spirv.

I wasn't able to quickly find a simple way to detect debug info presence
and also I think that the effect of `GlobalDCE` here is not so crucial, because
the amount of cleaned up code is small and can that code can easily be
removed by backends

Author: AlexeySachkov

llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll

@@ -20,15 +20,19 @@
 @id_compos = dso_local global %class.specialization_id.1 { %struct.ComposConst { i32 1, double 2.000000e+00, %struct.myConst { i32 13, float 0x4020666660000000 } } }, align 8
 @id_compos2 = dso_local global %class.specialization_id.2 { %struct.ComposConst2 { i8 1, %struct.myConst { i32 52, float 0x40479999A0000000 }, double 2.000000e+00 } }, align 8
 
-; check that no longer used globals are removed:
+; check that the following globals are preserved: even though they are won't be
+; used in the module anymore, they could still be referenced by debug info
+; metadata (specialization_id objects are used as template arguments in SYCL
+; specialization constant APIs)
+; CHECK: @id_double
+; CHECK: @id_int
+; CHECK: @id_compos
+; CHECK: @id_compos2
+
 @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_doubleE17specialization_idIdEdET1_v = private unnamed_addr constant [37 x i8] c"_ZTS14name_generatorIL_Z9id_doubleEE\00", align 1
-; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_doubleE17specialization_idIdEdET1_v
 @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z6id_intE17specialization_idIiEiET1_v = private unnamed_addr constant [34 x i8] c"_ZTS14name_generatorIL_Z6id_intEE\00", align 1
-; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z6id_intE17specialization_idIiEiET1_v
 @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_composE17specialization_idI11ComposConstES1_ET1_v = private unnamed_addr constant [37 x i8] c"_ZTS14name_generatorIL_Z9id_composEE\00", align 1
-; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_composE17specialization_idI11ComposConstES1_ET1_v
 @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_compos2E17specialization_idI12ComposConst2ES1_ET1_v = private unnamed_addr constant [39 x i8] c"_ZTS14name_generatorIL_Z10id_compos2EE\00", align 1
-; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_compos2E17specialization_idI12ComposConst2ES1_ET1_v
 
 ; CHECK-LABEL: define dso_local void @_Z4testv
 define dso_local void @_Z4testv() local_unnamed_addr #0 {

llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll

@@ -1,5 +1,5 @@
 ; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
-; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
+; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue"
 ;
 ; This test is intended to check that sycl-post-link tool is capable of handling
 ; composite specialization constants by lowering them into a set of SPIR-V

llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll

@@ -1,5 +1,5 @@
 ; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
-; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue --implicit-check-not __sycl_getComposite2020SpecConstantValue
+; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue" --implicit-check-not "call {{.*}} __sycl_getComposite2020SpecConstantValue"
 ;
 ; This test is intended to check that sycl-post-link tool is capable of handling
 ; composite specialization constants by lowering them into a set of SPIR-V

llvm/test/tools/sycl-post-link/spec-constants/composite-default-value.ll

@@ -1,5 +1,5 @@
 ; RUN: sycl-post-link -spec-const=default --ir-output-only %s -S -o - \
-; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue --implicit-check-not __sycl_getComposite2020SpecConstantValue
+; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue" --implicit-check-not "call {{.*}} __sycl_getComposite2020SpecConstantValue"
 ;
 ; This test checks that composite specialization constants can be correctly
 ; initialized by sycl-post-link tool for AOT use-case (default initialization

llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll

@@ -1,5 +1,5 @@
 ; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
-; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue --implicit-check-not __sycl_getComposite2020SpecConstantValue
+; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue" --implicit-check-not "call {{.*}} __sycl_getComposite2020SpecConstantValue"
 
 ; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32
 ; CHECK: %[[#NS1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 1]], i32 42)

llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll

@@ -1,5 +1,5 @@
 ; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
-; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
+; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue"
 ;
 ; This test is intended to check that sycl-post-link tool is capable of handling
 ; situations when the same composite specialization constants is used more than

llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages.ll

@@ -1,5 +1,5 @@
 ; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
-; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
+; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue"
 ;
 ; This test is intended to check that sycl-post-link tool is capable of handling
 ; situations when the same composite specialization constants is used more than

llvm/test/tools/sycl-post-link/spec-constants/multiple-sym-id-usages.ll

@@ -1,16 +1,14 @@
 ; This test checks that the tool does not crash and removes the unused spec
 ; constant global symbol when it is referenced more than once.
 
-; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
-; RUN: | FileCheck %s 
+; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o %t.ll
 
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
 target triple = "spir64-unknown-unknown-sycldevice"
 
 %"spec_constant" = type { i8 }
 
 @SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1
-; CHECK-NOT: @SCSymID
 
 declare dso_local spir_func float @_Z33__sycl_getScalarSpecConstantValueIfET_PKc(i8 addrspace(4)*)
 

llvm/test/tools/sycl-post-link/spec-constants/scalar-O0.ll

@@ -14,8 +14,6 @@ target triple = "spir64-unknown-linux-sycldevice"
 $FOO = comdat any
 
 @__unique_stable_name.FOO = private unnamed_addr constant [18 x i8] c"_ZTS11MyBoolConst\00", align 1
-; check that post link removes the no longer used global
-; CHECK-NOT: @__unique_stable_name.FOO
 
 declare dso_local spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbET_PKc(i8 addrspace(4)*) #7
 

llvm/test/tools/sycl-post-link/spec-constants/scalar-O2.ll

@@ -17,30 +17,17 @@ target triple = "spir64-unknown-linux-sycldevice"
 $_ZTS17SpecializedKernel = comdat any
 
 @__unique_stable_name.SC_Ib11MyBoolConstE3getEv = private unnamed_addr constant [18 x i8] c"_ZTS11MyBoolConst\00", align 1
-;;;;;;;;;;;;;; check that no longer used globals are removed:
-; CHECK-NOT: @__unique_stable_name.SC_Ib11MyBoolConstE3getEv
 @__unique_stable_name.SC_Ia11MyInt8ConstE3getEv = private unnamed_addr constant [18 x i8] c"_ZTS11MyInt8Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Ia11MyInt8ConstE3getEv
 @__unique_stable_name.SC_Ih12MyUInt8ConstE3getEv = private unnamed_addr constant [19 x i8] c"_ZTS12MyUInt8Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Ih12MyUInt8ConstE3getEv
 @__unique_stable_name.SC_Is12MyInt16ConstE3getEv = private unnamed_addr constant [19 x i8] c"_ZTS12MyInt16Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Is12MyInt16ConstE3getEv
 @__unique_stable_name.SC_It13MyUInt16ConstE3getEv = private unnamed_addr constant [20 x i8] c"_ZTS13MyUInt16Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_It13MyUInt16ConstE3getEv
 @__unique_stable_name.SC_Ii12MyInt32ConstE3getEv = private unnamed_addr constant [19 x i8] c"_ZTS12MyInt32Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Ii12MyInt32ConstE3getEv
 @__unique_stable_name.SC_Ij13MyUInt32ConstE3getEv = private unnamed_addr constant [20 x i8] c"_ZTS13MyUInt32Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Ij13MyUInt32ConstE3getEv
 @__unique_stable_name.SC_Il12MyInt64ConstE3getEv = private unnamed_addr constant [19 x i8] c"_ZTS12MyInt64Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Il12MyInt64ConstE3getEv
 @__unique_stable_name.SC_Im13MyUInt64ConstE3getEv = private unnamed_addr constant [20 x i8] c"_ZTS13MyUInt64Const\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Im13MyUInt64ConstE3getEv
 @__unique_stable_name.SC_If12MyFloatConstE3getEv = private unnamed_addr constant [19 x i8] c"_ZTS12MyFloatConst\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_If12MyFloatConstE3getEv
 @__unique_stable_name.SC_Id13MyDoubleConstE3getEv = private unnamed_addr constant [20 x i8] c"_ZTS13MyDoubleConst\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Id13MyDoubleConstE3getEv
 @__unique_stable_name.SC_Id14MyDoubleConst2E3getEv = private unnamed_addr constant [21 x i8] c"_ZTS14MyDoubleConst2\00", align 1
-; CHECK-NOT: @__unique_stable_name.SC_Id14MyDoubleConst2E3getEv
 
 ; Function Attrs: norecurse
 define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrspace(1)* %0, %"cl::sycl::range"* byval(%"cl::sycl::range") align 8 %1, %"cl::sycl::range"* byval(%"cl::sycl::range") align 8 %2, %"cl::sycl::id"* byval(%"cl::sycl::id") align 8 %3) local_unnamed_addr #0 comdat !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {

llvm/test/tools/sycl-post-link/spec-constants/spec_const_and_split.ll

@@ -7,17 +7,15 @@
 ; into properties of each device image
 
 ; RUN: sycl-post-link -split=kernel -spec-const=rt -S %s -o %t.files.table
-; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefixes CHECK-IR0,CHECK-IR
-; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefixes CHECK-IR1,CHECK-IR
-; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefixes CHECK-IR2,CHECK-IR
+; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefixes CHECK-IR0
+; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefixes CHECK-IR1
+; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefixes CHECK-IR2
 ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefixes CHECK-PROP0
 ; RUN: FileCheck %s -input-file=%t.files_1.prop --check-prefixes CHECK-PROP1
 ; RUN: FileCheck %s -input-file=%t.files_2.prop --check-prefixes CHECK-PROP2
 
 @SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1
 @SCSymID2 = private unnamed_addr constant [11 x i8] c"SpecConst2\00", align 1
-; CHECK-IR-NOT: @SCSymID
-; CHECK-IR-NOT: @SCSymID2
 
 declare dso_local spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbET_PKc(i8 addrspace(4)*)
 

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -573,8 +573,6 @@ static TableFiles processOneModule(std::unique_ptr<Module> M, bool IsEsimd,
     // Register required analysis
     MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
     RunSpecConst.addPass(SCP);
-    // This pass deletes unreachable globals.
-    RunSpecConst.addPass(GlobalDCEPass());
 
     for (auto &MPtr : ResultModules) {
       // perform the spec constant intrinsics transformation on each resulting

sycl/test/basic_tests/SYCL-2020-spec-const-ids-order.cpp

@@ -1,7 +1,7 @@
 // RUN: %clangxx -fsycl -fsycl-device-only -c -o %t.bc %s
 // RUN: sycl-post-link %t.bc -spec-const=default -S -o %t-split1.txt
-// RUN: cat %t-split1_0.ll | FileCheck %s -check-prefixes=CHECK,CHECK-IR
-// RUN: cat %t-split1_0.prop | FileCheck %s -check-prefixes=CHECK,CHECK-PROP
+// RUN: cat %t-split1_0.ll | FileCheck %s -check-prefixes=CHECK-IR
+// RUN: cat %t-split1_0.prop | FileCheck %s -check-prefixes=CHECK-PROP
 //
 //==----------- SYCL-2020-spec-const-ids-order.cpp -------------------------==//
 //
@@ -46,11 +46,13 @@ int main() {
 }
 
 // CHECK-PROP: [SYCL/specialization constants]
-// CHECK: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL5Val23EEE
-// CHECK-IR-SAME: i32 [[#ID:]]
-// CHECK-NEXT: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL10ConstantIdEEE
-// CHECK-IR-SAME: i32 [[#ID+1]]
-// CHECK-NEXT: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL11SecondValueEEE
-// CHECK-IR-SAME: i32 [[#ID+2]]
-// CHECK-NEXT: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL11SpecConst42EEE
-// CHECK-IR-SAME: i32 [[#ID+3]]
+// CHECK-PROP-NEXT: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL5Val23EEE
+// CHECK-PROP-NEXT: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL10ConstantIdEEE
+// CHECK-PROP-NEXT: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL11SecondValueEEE
+// CHECK-PROP-NEXT: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL11SpecConst42EEE
+//
+// CHECK-IR: !sycl.specialization-constants = !{![[#MD0:]], ![[#MD1:]], ![[#MD2:]], ![[#MD3:]]}
+// CHECK-IR: ![[#MD0]] = !{!"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL5Val23EEE", i32 [[#ID:]]
+// CHECK-IR: ![[#MD1]] = !{!"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL10ConstantIdEEE", i32 [[#ID+1]]
+// CHECK-IR: ![[#MD2]] = !{!"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL11SecondValueEEE", i32 [[#ID+2]]
+// CHECK-IR: ![[#MD3]] = !{!"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL11SpecConst42EEE", i32 [[#ID+3]]