[sycl-post-link] Generate spec constants device image property even for non-native spec constants (#3561)

The main purpose of this change is to emit the same device image
property SYCL/specialization constants not only when native spec
constants are supported, but also when they are emulated.

This is needed to support SYCL 2020 specialization constants, design
document can be found in: #3331.

The task is achieved by changing the way how SpecConstantsPass
works with metadata: in order to provide info which should be put
into device image properties, it was attached as metadata to
__spirv_SpecConstant intrinsics and they are not generated if native
spec constants are not available. New approach creates a single
named metadata sycl.specialization-constants which contains a list
of metadatas, corresponding to each specialization constant found
in the module.

Such new representation makes it easier to generate metadata and
more significantly, it should improve the speed of reading them,
because we don't need to look over each call to
__spirv_SpecConstant intrinsic.

Author: AlexeySachkov

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

@@ -38,21 +38,21 @@ entry:
 ; CHECK-DEF: %[[BITCAST:[0-9a-z]+]] = bitcast i8* %[[GEP]] to double*
 ; CHECK-DEF: %[[LOAD:[0-9a-z]+]] = load double, double* %[[BITCAST]], align 8
 ;
-; CHECK-RT: call double @_Z20__spirv_SpecConstantid(i32 [[#SCID0:]], double 3.140000e+00), !SYCL_SPEC_CONST_SYM_ID ![[#MID0:]]
+; CHECK-RT: call double @_Z20__spirv_SpecConstantid(i32 [[#SCID0:]], double 3.140000e+00)
 
   %call.i3 = tail call i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPvS3_(i8* getelementptr inbounds ([34 x i8], [34 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z6id_intE17specialization_idIiEiET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id.0* @id_int to i8*), i8* null)
 ; CHECK-DEF: %[[GEP1:[0-9a-z]+]] = getelementptr i8, i8* null, i32 8
 ; CHECK-DEF: %[[BITCAST1:[0-9a-z]+]] = bitcast i8* %[[GEP1]] to i32*
 ; CHECK-DEF: %[[LOAD1:[0-9a-z]+]] = load i32, i32* %[[BITCAST1]], align 4
 ;
-; CHECK-RT: call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID1:]], i32 42), !SYCL_SPEC_CONST_SYM_ID ![[#MID1:]]
+; CHECK-RT: call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID1:]], i32 42)
 
   %call.i4 = tail call fast double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8* getelementptr inbounds ([37 x i8], [37 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_doubleE17specialization_idIdEdET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id* @id_double to i8*), i8* null)
 ; CHECK-DEF: %[[GEP2:[0-9a-z]+]] = getelementptr i8, i8* null, i32 0
 ; CHECK-DEF: %[[BITCAST2:[0-9a-z]+]] = bitcast i8* %[[GEP2]] to double*
 ; CHECK-DEF: %[[LOAD2:[0-9a-z]+]] = load double, double* %[[BITCAST2]], align 8
 ;
-; CHECK-RT: call double @_Z20__spirv_SpecConstantid(i32 [[#SCID0]], double 3.140000e+00), !SYCL_SPEC_CONST_SYM_ID ![[#MID0]]
+; CHECK-RT: call double @_Z20__spirv_SpecConstantid(i32 [[#SCID0]], double 3.140000e+00)
   ret void
 }
 
@@ -74,7 +74,7 @@ entry:
 ; CHECK-RT: %[[#SE3:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID4:]], i32 13)
 ; CHECK-RT: %[[#SE4:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0x4020666660000000)
 ; CHECK-RT: %[[#CE1:]] = call %struct.myConst @_Z29__spirv_SpecConstantCompositeif(i32 %[[#SE3]], float %[[#SE4]])
-; CHECK-RT: %[[C1:[0-9a-z]+]] = call %struct.ComposConst @_Z29__spirv_SpecConstantCompositeidstruct.myConst(i32 %[[#SE1]], double %[[#SE2]], %struct.myConst %[[#CE1]]), !SYCL_SPEC_CONST_SYM_ID ![[#MID2:]]
+; CHECK-RT: %[[C1:[0-9a-z]+]] = call %struct.ComposConst @_Z29__spirv_SpecConstantCompositeidstruct.myConst(i32 %[[#SE1]], double %[[#SE2]], %struct.myConst %[[#CE1]])
 ;
 ; CHECK: store %struct.ComposConst %[[C1]], %struct.ComposConst*
 
@@ -91,7 +91,7 @@ entry:
 ; CHECK-RT: %[[#SE3:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID8:]], float 0x40479999A0000000)
 ; CHECK-RT: %[[#CE1:]] = call %struct.myConst @_Z29__spirv_SpecConstantCompositeif(i32 %[[#SE2]], float %[[#SE3]])
 ; CHECK-RT: %[[#SE4:]] = call double @_Z20__spirv_SpecConstantid(i32 [[#SCID9:]], double 2.000000e+00)
-; CHECK-RT: %[[C2:[0-9a-z]+]] = call %struct.ComposConst2 @_Z29__spirv_SpecConstantCompositeastruct.myConstd(i8 %[[#SE1]], %struct.myConst %[[#CE1]], double %[[#SE4]]), !SYCL_SPEC_CONST_SYM_ID ![[#MID3:]]
+; CHECK-RT: %[[C2:[0-9a-z]+]] = call %struct.ComposConst2 @_Z29__spirv_SpecConstantCompositeastruct.myConstd(i8 %[[#SE1]], %struct.myConst %[[#CE1]], double %[[#SE4]])
 ;
 ; CHECK: store %struct.ComposConst2 %[[C2]], %struct.ComposConst2*
 
@@ -107,7 +107,7 @@ entry:
 ; CHECK-RT: %[[#SE3:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID4]], i32 13)
 ; CHECK-RT: %[[#SE4:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5]], float 0x4020666660000000)
 ; CHECK-RT: %[[#CE1:]] = call %struct.myConst @_Z29__spirv_SpecConstantCompositeif(i32 %[[#SE3]], float %[[#SE4]])
-; CHECK-RT: %[[C3:[0-9a-z]+]] = call %struct.ComposConst @_Z29__spirv_SpecConstantCompositeidstruct.myConst(i32 %[[#SE1]], double %[[#SE2]], %struct.myConst %[[#CE1]]), !SYCL_SPEC_CONST_SYM_ID ![[#MID2]]
+; CHECK-RT: %[[C3:[0-9a-z]+]] = call %struct.ComposConst @_Z29__spirv_SpecConstantCompositeidstruct.myConst(i32 %[[#SE1]], double %[[#SE2]], %struct.myConst %[[#CE1]])
 ;
 ; CHECK: store %struct.ComposConst %[[C3]], %struct.ComposConst*
   call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2) #3
@@ -130,7 +130,21 @@ attributes #1 = { argmemonly nofree nosync nounwind willreturn }
 attributes #2 = { "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" "unsafe-fp-math"="true" "use-soft-float"="false" }
 attributes #3 = { nounwind }
 
-; CHECK-RT: ![[#MID0]] = !{!"_ZTS14name_generatorIL_Z9id_doubleEE", i32 [[#SCID0]]}
-; CHECK-RT: ![[#MID1]] = !{!"_ZTS14name_generatorIL_Z6id_intEE", i32 [[#SCID1]]}
-; CHECK-RT: ![[#MID2]] = !{!"_ZTS14name_generatorIL_Z9id_composEE", i32 [[#SCID2]], i32 [[#SCID3]], i32 [[#SCID4]], i32 [[#SCID5]]}
-; CHECK-RT: ![[#MID3]] = !{!"_ZTS14name_generatorIL_Z10id_compos2EE", i32 [[#SCID6]], i32 [[#SCID7]], i32 [[#SCID8]], i32 [[#SCID9]]}
+; CHECK: !sycl.specialization-constants = !{![[#ID0:]], ![[#ID1:]], ![[#ID2:]], ![[#ID3:]]}
+;
+; CHECK: ![[#ID0]] = !{!"_ZTS14name_generatorIL_Z9id_doubleEE", i32 0, i32 0, i32 8}
+; CHECK: ![[#ID1]] = !{!"_ZTS14name_generatorIL_Z6id_intEE", i32 1, i32 0, i32 4}
+;
+; For composite types, the amount of metadata is a bit different between native and emulated spec constants
+;
+; CHECK-DEF: ![[#ID2]] = !{!"_ZTS14name_generatorIL_Z9id_composEE", i32 2, i32 0, i32 24}
+; CHECK-DEF: ![[#ID3]] = !{!"_ZTS14name_generatorIL_Z10id_compos2EE", i32 3, i32 0, i32 24
+;
+; CHECK-RT: ![[#ID2]] = !{!"_ZTS14name_generatorIL_Z9id_composEE", i32 [[#SCID2]], i32 0, i32 4,
+; CHECK-RT-SAME: i32 [[#SCID3]], i32 8, i32 8,
+; CHECK-RT-SAME: i32 [[#SCID4]], i32 16, i32 4,
+; CHECK-RT-SAME: i32 [[#SCID5]], i32 20, i32 4}
+; CHECK-RT: ![[#ID3]] = !{!"_ZTS14name_generatorIL_Z10id_compos2EE", i32 [[#SCID6]], i32 0, i32 1,
+; CHECK-RT-SAME: i32 [[#SCID7]], i32 4, i32 4,
+; CHECK-RT-SAME: i32 [[#SCID8]], i32 8, i32 4,
+; CHECK-RT-SAME: i32 [[#SCID9]], i32 16, i32 8}

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

@@ -20,10 +20,15 @@
 ;
 ; CHECK: %[[#B:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 4]], i32{{.*}})
 ;
-; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Ai([2 x %struct._ZTS1A.A] %[[#NA]], i32 %[[#B]]), !SYCL_SPEC_CONST_SYM_ID ![[#MD:]]
+; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Ai([2 x %struct._ZTS1A.A] %[[#NA]], i32 %[[#B]])
 ; CHECK: store %struct._ZTS3POD.POD %[[#POD]]
 ;
-; CHECK: ![[#MD]] = !{!"_ZTS3POD", i32 [[#ID]], i32 [[#ID + 1]], i32 [[#ID + 2]], i32 [[#ID + 3]], i32 [[#ID + 4]]}
+; CHECK: !sycl.specialization-constants = !{![[#MD:]]}
+; CHECK: ![[#MD]] = !{!"_ZTS3POD", i32 [[#ID]], i32 0, i32 4,
+; CHECK-SAME: i32 [[#ID + 1]], i32 4, i32 4,
+; CHECK-SAME: i32 [[#ID + 2]], i32 8, i32 4,
+; CHECK-SAME: i32 [[#ID + 3]], i32 12, i32 4,
+; CHECK-SAME: i32 [[#ID + 4]], i32 16, i32 4}
 
 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-sycldevice"

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

@@ -5,6 +5,7 @@
 ; composite specialization constants by lowering them into a set of SPIR-V
 ; friendly IR operations representing those constants.
 ;
+; CHECK-LABEL: define {{.*}} spir_kernel void @_ZTS4Test
 ; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32
 ; CHECK: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 1]], float
 ; CHECK: %[[#NA0:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS0]], float %[[#NS1]])
@@ -20,9 +21,27 @@
 ; CHECK: %[[#BV:]] = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32 %[[#B0]], i32 %[[#B1]])
 ; CHECK: %[[#B:]] = call %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" @_Z29__spirv_SpecConstantCompositeDv2_i(<2 x i32> %[[#BV]])
 ;
-; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"([2 x %struct._ZTS1A.A] %[[#NA]], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" %[[#B]]), !SYCL_SPEC_CONST_SYM_ID ![[#MD:]]
+; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"([2 x %struct._ZTS1A.A] %[[#NA]], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" %[[#B]])
 ; CHECK: store %struct._ZTS3POD.POD %[[#POD]]
 
+; CHECK-LABEL: define {{.*}} spir_kernel void @_ZTS17SpecializedKernel
+; CHECK: %[[#N0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 6]], i32
+; CHECK: %[[#N1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 7]], float
+; CHECK: %[[#CONST:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#N0]], float %[[#N1]])
+; CHECK: %struct._ZTS1A.A %[[#CONST]]
+;
+; CHECK: !sycl.specialization-constants = !{![[#MD0:]], ![[#MD1:]]}
+;
+; CHECK: ![[#MD0]] = !{!"_ZTS3POD", i32 [[#ID]], i32 0, i32 4,
+; CHECK-SAME: i32 [[#ID + 1]], i32 4, i32 4,
+; CHECK-SAME: i32 [[#ID + 2]], i32 8, i32 4,
+; CHECK-SAME: i32 [[#ID + 3]], i32 12, i32 4,
+; CHECK-SAME: i32 [[#ID + 4]], i32 16, i32 4,
+; CHECK-SAME: i32 [[#ID + 5]], i32 20, i32 4}
+
+; CHECK: ![[#MD1]] = !{!"_ZTS13MyComposConst", i32 [[#ID + 6]], i32 0, i32 4,
+; CHECK-SAME: i32 [[#ID + 7]], i32 4, i32 4}
+
 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-sycldevice"
 
@@ -69,10 +88,6 @@ entry:
   %3 = bitcast %struct._ZTS1A.A* %c.i to i8*
   call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %3) #3
   call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4 %c.ascast.i, i8 addrspace(4)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(4)* addrspacecast ([20 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv to [20 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null, i8 addrspace(4)* null) #4
-; CHECK: %[[#N0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 6]], i32
-; CHECK: %[[#N1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 7]], float
-; CHECK: %[[#CONST:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#N0]], float %[[#N1]]), !SYCL_SPEC_CONST_SYM_ID ![[#MD1:]]
-; CHECK: %struct._ZTS1A.A %[[#CONST]]
   %a.i = getelementptr inbounds %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %c.ascast.i, i64 0, i32 0
   %4 = load i32, i32 addrspace(4)* %a.i, align 4
   %conv.i = sitofp i32 %4 to float
@@ -111,8 +126,6 @@ attributes #4 = { convergent }
 !spirv.Source = !{!2}
 !llvm.ident = !{!3}
 
-; CHECK: ![[#MD]] = !{!"_ZTS3POD", i32 [[#ID]], i32 [[#ID + 1]], i32 [[#ID + 2]], i32 [[#ID + 3]], i32 [[#ID + 4]], i32 [[#ID + 5]]}
-; CHECK: ![[#MD1]] = !{!"_ZTS13MyComposConst", i32 [[#ID + 6]], i32 [[#ID + 7]]}
 !0 = !{i32 1, !"wchar_size", i32 4}
 !1 = !{i32 1, i32 2}
 !2 = !{i32 4, i32 100000}

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

@@ -9,12 +9,15 @@
 ; CHECK-LABEL: @_ZTSN4test8kernel_tIfEE
 ; CHECK: %[[#X1:]] = call float @_Z20__spirv_SpecConstantif(i32 0, float 0
 ; CHECK: %[[#Y1:]] = call float @_Z20__spirv_SpecConstantif(i32 1, float 0
-; CHECK: call {{.*}} @_Z29__spirv_SpecConstantCompositeff(float %[[#X1]], float %[[#Y1]]), !SYCL_SPEC_CONST_SYM_ID ![[#ID:]]
+; CHECK: call {{.*}} @_Z29__spirv_SpecConstantCompositeff(float %[[#X1]], float %[[#Y1]])
 ; CHECK-LABEL: @_ZTSN4test8kernel_tIiEE
 ; CHECK: %[[#X2:]] = call float @_Z20__spirv_SpecConstantif(i32 0, float 0
 ; CHECK: %[[#Y2:]] = call float @_Z20__spirv_SpecConstantif(i32 1, float 0
-; CHECK: call {{.*}} @_Z29__spirv_SpecConstantCompositeff(float %[[#X2]], float %[[#Y2]]), !SYCL_SPEC_CONST_SYM_ID ![[#ID]]
-; CHECK: ![[#ID]] = !{!"_ZTS11sc_kernel_t", i32 0, i32 1}
+; CHECK: call {{.*}} @_Z29__spirv_SpecConstantCompositeff(float %[[#X2]], float %[[#Y2]])
+
+; CHECK: !sycl.specialization-constants = !{![[#ID:]]
+
+; CHECK: ![[#ID]] = !{!"_ZTS11sc_kernel_t", i32 0, i32 0, i32 4, i32 1, i32 4, i32 4}
 
 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-sycldevice"

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

@@ -6,16 +6,26 @@
 ; once
 ;
 ; CHECK-LABEL: @foo1
-; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[#MD0:]]
+; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}})
 ; CHECK-LABEL: @_ZTS4Test
-; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[#MD1:]]
+; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}})
 ; CHECK-LABEL: @foo2
-; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[#MD0:]]
+; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}})
+
+; CHECK: !sycl.specialization-constants = !{![[#MD0:]], ![[#MD1:]]
 ;
-; CHECK-DAG: ![[#MD0]] = !{!"_ZTS3PO2", i32 [[#ID:]],
-; CHECK-SAME: i32 [[#ID + 1]], i32 [[#ID + 2]], i32 [[#ID + 3]], i32 [[#ID + 4]], i32 [[#ID + 5]]}
-; CHECK-DAG: ![[#MD1]] = !{!"_ZTS3POD", i32 [[#ID1:]],
-; CHECK-SAME: i32 [[#ID1 + 1]], i32 [[#ID1 + 2]], i32 [[#ID1 + 3]], i32 [[#ID1 + 4]], i32 [[#ID1 + 5]]}
+; CHECK-DAG: ![[#MD0]] = !{!"_ZTS3PO2", i32 [[#ID:]], i32 0, i32 4,
+; CHECK-SAME: i32 [[#ID + 1]], i32 4, i32 4,
+; CHECK-SAME: i32 [[#ID + 2]], i32 8, i32 4,
+; CHECK-SAME: i32 [[#ID + 3]], i32 12, i32 4,
+; CHECK-SAME: i32 [[#ID + 4]], i32 16, i32 4,
+; CHECK-SAME: i32 [[#ID + 5]], i32 20, i32 4}
+; CHECK-DAG: ![[#MD1]] = !{!"_ZTS3POD", i32 [[#ID1:]], i32 0, i32 4,
+; CHECK-SAME: i32 [[#ID1 + 1]], i32 4, i32 4,
+; CHECK-SAME: i32 [[#ID1 + 2]], i32 8, i32 4,
+; CHECK-SAME: i32 [[#ID1 + 3]], i32 12, i32 4,
+; CHECK-SAME: i32 [[#ID1 + 4]], i32 16, i32 4,
+; CHECK-SAME: i32 [[#ID1 + 5]], i32 20, i32 4}
 
 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-sycldevice"

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

@@ -18,23 +18,24 @@ declare dso_local spir_func float @_Z33__sycl_getScalarSpecConstantValueIfET_PKc
 ; Function Attrs: norecurse
 define weak_odr dso_local spir_kernel void @Kernel() {
   %1 = call spir_func float @_Z33__sycl_getScalarSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
-; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID0:[0-9]+]]
+; CHECK: call float @_Z20__spirv_SpecConstantif(i32 0, {{.*}})
   ret void
 }
 
 ; Function Attrs: norecurse
 define dso_local spir_func float @foo_float(%"spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
   %2 = tail call spir_func float @_Z33__sycl_getScalarSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([11 x i8], [11 x i8]* @SCSymID1, i64 0, i64 0) to i8 addrspace(4)*))
-; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID1:[0-9]+]]
+; CHECK: call float @_Z20__spirv_SpecConstantif(i32 1, {{.*}})
   ret float %2
 }
 
 ; Function Attrs: norecurse
 define dso_local spir_func float @foo_float2(%"spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
   %2 = tail call spir_func float @_Z33__sycl_getScalarSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
-; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID0]]
+; CHECK: call float @_Z20__spirv_SpecConstantif(i32 0, {{.*}})
   ret float %2
 }
 
-; CHECK: ![[ID0]] = !{!"SpecConst", i32 0}
-; CHECK: ![[ID1]] = !{!"SpecConst1", i32 1}
+; CHECK: !sycl.specialization-constants = !{![[#MD1:]], ![[#MD2:]]}
+; CHECK: ![[#MD1]] = !{!"SpecConst", i32 0, i32 0, i32 4}
+; CHECK: ![[#MD2]] = !{!"SpecConst1", i32 1, i32 0, i32 4}

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

@@ -39,7 +39,7 @@ define spir_func zeroext i1 @FOO(%"UserSpecConstIDType" addrspace(4)* %0) comdat
   %7 = call spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbET_PKc(i8 addrspace(4)* %6)
 ; with -spec-const=rt the __sycl_getSpecConstantValue is replaced with
 ;  SPIRV intrinsic
-; CHECK-RT: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[MD_ID:[0-9]+]]
+; CHECK-RT: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false)
   %8 = bitcast i8 addrspace(4)** %3 to i8*
   call void @llvm.lifetime.end.p0i8(i64 8, i8* %8) #8
   ret i1 %7
@@ -65,4 +65,6 @@ attributes #8 = { nounwind }
 !9 = !{!"any pointer", !10, i64 0}
 !10 = !{!"omnipotent char", !11, i64 0}
 !11 = !{!"Simple C++ TBAA"}
-; CHECK-RT: ![[MD_ID]] = !{!"_ZTS11MyBoolConst", i32 0}
+
+; CHECK-RT: !sycl.specialization-constants = !{![[#MD:]]}
+; CHECK-RT: ![[#MD]] = !{!"_ZTS11MyBoolConst", i32 0, i32 0, i32 1}