diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll index 267fe0d8ffa9e..6661b7f0b861c 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll +++ b/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} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll index 94dfeb5abf7e0..695d2ac0da895 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll +++ b/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" diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll index a0b56bdc2fd39..1d5ae29c2f7c2 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll +++ b/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} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll b/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll index f07c9b5c47a1c..b8605b5f22cc7 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll +++ b/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" diff --git a/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages.ll b/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages.ll index 7eb8a4154c553..f6c0b1dc97ac2 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages.ll +++ b/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" diff --git a/llvm/test/tools/sycl-post-link/spec-constants/multiple-scalar-usages.ll b/llvm/test/tools/sycl-post-link/spec-constants/multiple-scalar-usages.ll index 0fc5b97ef56a7..88a947cdd320b 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/multiple-scalar-usages.ll +++ b/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} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/scalar-O0.ll b/llvm/test/tools/sycl-post-link/spec-constants/scalar-O0.ll index 14dad52ab3cf1..531949069937f 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/scalar-O0.ll +++ b/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} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/scalar-O2.ll b/llvm/test/tools/sycl-post-link/spec-constants/scalar-O2.ll index c82c199fd0f73..dd1f231ccba6c 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/scalar-O2.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/scalar-O2.ll @@ -51,7 +51,7 @@ define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrsp ;;;;;;;;;;;;;; check that __sycl* intrinsic goes away: ; CHECK-NOT: %{{[0-9]+}} ={{.*}} call {{.*}}@_Z33__sycl_getScalarSpecConstantValueIbET_PKc ;;;;;;;;;;;;;; check that with -spec-const=rt __spirv* intrinsic is generated: -; CHECK-RT: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[ID0:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false) ;;;;;;;;;;;;;; check that with -spec-const=default __spirv* intrinsic is not ;;;;;;;;;;;;;; generated: ; CHECK-DEF-NOT: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstant @@ -60,43 +60,43 @@ define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrsp ;;;;;;;;;;;;;; intrinsics are replaced with constants: ; CHECK-DEF: %[[VAL0:[0-9]+]] = zext i1 false to i32 %10 = tail call spir_func signext i8 @_Z33__sycl_getScalarSpecConstantValueIaET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([18 x i8], [18 x i8]* @__unique_stable_name.SC_Ia11MyInt8ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i8 @_Z20__spirv_SpecConstantia(i32 1, i8 0), !SYCL_SPEC_CONST_SYM_ID ![[ID1:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i8 @_Z20__spirv_SpecConstantia(i32 1, i8 0) %11 = sext i8 %10 to i32 ; CHECK-DEF: %[[VAL1:[0-9]+]] = sext i8 0 to i32 %12 = add nsw i32 %11, %9 ; CHECK-DEF: %[[SUM0:[0-9]+]] = add nsw i32 %[[VAL1]], %[[VAL0]] %13 = tail call spir_func zeroext i8 @_Z33__sycl_getScalarSpecConstantValueIhET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([19 x i8], [19 x i8]* @__unique_stable_name.SC_Ih12MyUInt8ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i8 @_Z20__spirv_SpecConstantia(i32 2, i8 0), !SYCL_SPEC_CONST_SYM_ID ![[ID2:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i8 @_Z20__spirv_SpecConstantia(i32 2, i8 0) %14 = zext i8 %13 to i32 ; CHECK-DEF: %[[VAL2:[0-9]+]] = zext i8 0 to i32 %15 = add nsw i32 %12, %14 ; CHECK-DEF: %[[SUM1:[0-9]+]] = add nsw i32 %[[SUM0]], %[[VAL2]] %16 = tail call spir_func signext i16 @_Z33__sycl_getScalarSpecConstantValueIsET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([19 x i8], [19 x i8]* @__unique_stable_name.SC_Is12MyInt16ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i16 @_Z20__spirv_SpecConstantis(i32 3, i16 0), !SYCL_SPEC_CONST_SYM_ID ![[ID3:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i16 @_Z20__spirv_SpecConstantis(i32 3, i16 0) %17 = sext i16 %16 to i32 ; CHECK-DEF: %[[VAL3:[0-9]+]] = sext i16 0 to i32 %18 = add nsw i32 %15, %17 ; CHECK-DEF: %[[SUM2:[0-9]+]] = add nsw i32 %[[SUM1]], %[[VAL3]] %19 = tail call spir_func zeroext i16 @_Z33__sycl_getScalarSpecConstantValueItET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([20 x i8], [20 x i8]* @__unique_stable_name.SC_It13MyUInt16ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i16 @_Z20__spirv_SpecConstantis(i32 4, i16 0), !SYCL_SPEC_CONST_SYM_ID ![[ID4:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i16 @_Z20__spirv_SpecConstantis(i32 4, i16 0) %20 = zext i16 %19 to i32 ; CHECK-DEF: %[[VAL4:[0-9]+]] = zext i16 0 to i32 %21 = add nsw i32 %18, %20 ; CHECK-DEF: %[[SUM3:[0-9]+]] = add nsw i32 %[[SUM2]], %[[VAL4]] %22 = tail call spir_func i32 @_Z33__sycl_getScalarSpecConstantValueIiET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([19 x i8], [19 x i8]* @__unique_stable_name.SC_Ii12MyInt32ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i32 @_Z20__spirv_SpecConstantii(i32 5, i32 0), !SYCL_SPEC_CONST_SYM_ID ![[ID5:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i32 @_Z20__spirv_SpecConstantii(i32 5, i32 0) %23 = add nsw i32 %21, %22 ; CHECK-DEF: %[[SUM4:[0-9]+]] = add nsw i32 %[[SUM3]], 0 %24 = tail call spir_func i32 @_Z33__sycl_getScalarSpecConstantValueIjET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([20 x i8], [20 x i8]* @__unique_stable_name.SC_Ij13MyUInt32ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i32 @_Z20__spirv_SpecConstantii(i32 6, i32 0), !SYCL_SPEC_CONST_SYM_ID ![[ID6:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i32 @_Z20__spirv_SpecConstantii(i32 6, i32 0) %25 = add i32 %23, %24 ; CHECK-DEF: %[[SUM5:[0-9]+]] = add i32 %[[SUM4]], 0 %26 = zext i32 %25 to i64 ; CHECK-DEF: %[[VAL5:[0-9]+]] = zext i32 %[[SUM5]] to i64 %27 = tail call spir_func i64 @_Z33__sycl_getScalarSpecConstantValueIlET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([19 x i8], [19 x i8]* @__unique_stable_name.SC_Il12MyInt64ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i64 @_Z20__spirv_SpecConstantix(i32 7, i64 0), !SYCL_SPEC_CONST_SYM_ID ![[ID7:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i64 @_Z20__spirv_SpecConstantix(i32 7, i64 0) %28 = tail call spir_func i64 @_Z33__sycl_getScalarSpecConstantValueImET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([20 x i8], [20 x i8]* @__unique_stable_name.SC_Im13MyUInt64ConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call i64 @_Z20__spirv_SpecConstantix(i32 8, i64 0), !SYCL_SPEC_CONST_SYM_ID ![[ID8:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call i64 @_Z20__spirv_SpecConstantix(i32 8, i64 0) %29 = add i64 %28, %27 ; CHECK-DEF: %[[SUM6:[0-9]+]] = add i64 0, 0 %30 = add i64 %29, %26 @@ -104,17 +104,17 @@ define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrsp %31 = uitofp i64 %30 to float ; CHECK-DEF: %[[VAL6:[0-9]+]] = uitofp i64 %[[SUM7]] to float %32 = tail call spir_func float @_Z33__sycl_getScalarSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([19 x i8], [19 x i8]* @__unique_stable_name.SC_If12MyFloatConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call float @_Z20__spirv_SpecConstantif(i32 9, float 0.000000e+00), !SYCL_SPEC_CONST_SYM_ID ![[ID9:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call float @_Z20__spirv_SpecConstantif(i32 9, float 0.000000e+00) %33 = fadd float %32, %31 ; CHECK-DEF: %[[SUM8:[0-9]+]] = fadd float 0.000000e+00, %[[VAL6]] %34 = fpext float %33 to double ; CHECK-DEF: %[[VAL7:[0-9]+]] = fpext float %[[SUM8]] to double %35 = tail call spir_func double @_Z33__sycl_getScalarSpecConstantValueIdET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([20 x i8], [20 x i8]* @__unique_stable_name.SC_Id13MyDoubleConstE3getEv, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK-RT: %{{[0-9]+}} = call double @_Z20__spirv_SpecConstantid(i32 10, double 0.000000e+00), !SYCL_SPEC_CONST_SYM_ID ![[ID10:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call double @_Z20__spirv_SpecConstantid(i32 10, double 0.000000e+00) %36 = fadd double %35, %34 ; CHECK-DEF: %[[SUM9:[0-9]+]] = fadd double 0.000000e+00, %[[VAL7]] %37 = tail call spir_func double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([21 x i8], [21 x i8]* @__unique_stable_name.SC_Id14MyDoubleConst2E3getEv, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* null, i8 addrspace(4)* null) -; CHECK-RT: %{{[0-9]+}} = call double @_Z20__spirv_SpecConstantid(i32 11, double 0.000000e+00), !SYCL_SPEC_CONST_SYM_ID ![[ID11:[0-9]+]] +; CHECK-RT: %{{[0-9]+}} = call double @_Z20__spirv_SpecConstantid(i32 11, double 0.000000e+00) %38 = fadd double %37, %36 ; CHECK-DEF: %[[GEP:[0-9a-z]+]] = getelementptr i8, i8 addrspace(4)* null, i32 0 ; CHECK-DEF: %[[BITCAST:[0-9a-z]+]] = bitcast i8 addrspace(4)* %[[GEP]] to double addrspace(4)* @@ -172,15 +172,17 @@ attributes #1 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail- !10 = !{!"omnipotent char", !11, i64 0} !11 = !{!"Simple C++ TBAA"} -; CHECK-RT: ![[ID0]] = !{!"_ZTS11MyBoolConst", i32 0} -; CHECK-RT: ![[ID1]] = !{!"_ZTS11MyInt8Const", i32 1} -; CHECK-RT: ![[ID2]] = !{!"_ZTS12MyUInt8Const", i32 2} -; CHECK-RT: ![[ID3]] = !{!"_ZTS12MyInt16Const", i32 3} -; CHECK-RT: ![[ID4]] = !{!"_ZTS13MyUInt16Const", i32 4} -; CHECK-RT: ![[ID5]] = !{!"_ZTS12MyInt32Const", i32 5} -; CHECK-RT: ![[ID6]] = !{!"_ZTS13MyUInt32Const", i32 6} -; CHECK-RT: ![[ID7]] = !{!"_ZTS12MyInt64Const", i32 7} -; CHECK-RT: ![[ID8]] = !{!"_ZTS13MyUInt64Const", i32 8} -; CHECK-RT: ![[ID9]] = !{!"_ZTS12MyFloatConst", i32 9} -; CHECK-RT: ![[ID10]] = !{!"_ZTS13MyDoubleConst", i32 10} -; CHECK-RT: ![[ID11]] = !{!"_ZTS14MyDoubleConst2", i32 11} +; CHECK-RT: !sycl.specialization-constants = !{![[#ID0:]], ![[#ID1:]], ![[#ID2:]], ![[#ID3:]], ![[#ID4:]], ![[#ID5:]], ![[#ID6:]], ![[#ID7:]], ![[#ID8:]], ![[#ID9:]], ![[#ID10:]], ![[#ID11:]]} + +; CHECK-RT: ![[#ID0:]] = !{!"_ZTS11MyBoolConst", i32 0, i32 0, i32 1} +; CHECK-RT: ![[#ID1:]] = !{!"_ZTS11MyInt8Const", i32 1, i32 0, i32 1} +; CHECK-RT: ![[#ID2:]] = !{!"_ZTS12MyUInt8Const", i32 2, i32 0, i32 1} +; CHECK-RT: ![[#ID3:]] = !{!"_ZTS12MyInt16Const", i32 3, i32 0, i32 2} +; CHECK-RT: ![[#ID4:]] = !{!"_ZTS13MyUInt16Const", i32 4, i32 0, i32 2} +; CHECK-RT: ![[#ID5:]] = !{!"_ZTS12MyInt32Const", i32 5, i32 0, i32 4} +; CHECK-RT: ![[#ID6:]] = !{!"_ZTS13MyUInt32Const", i32 6, i32 0, i32 4} +; CHECK-RT: ![[#ID7:]] = !{!"_ZTS12MyInt64Const", i32 7, i32 0, i32 8} +; CHECK-RT: ![[#ID8:]] = !{!"_ZTS13MyUInt64Const", i32 8, i32 0, i32 8} +; CHECK-RT: ![[#ID9:]] = !{!"_ZTS12MyFloatConst", i32 9, i32 0, i32 4} +; CHECK-RT: ![[#ID10:]] = !{!"_ZTS13MyDoubleConst", i32 10, i32 0, i32 8} +; CHECK-RT: ![[#ID11:]] = !{!"_ZTS14MyDoubleConst2", i32 11, i32 0, i32 8} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/spec_const_and_split.ll b/llvm/test/tools/sycl-post-link/spec-constants/spec_const_and_split.ll index a5a8649f9924b..b49a5dc91e821 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/spec_const_and_split.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/spec_const_and_split.ll @@ -13,14 +13,16 @@ declare dso_local spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbE define dso_local spir_kernel void @KERNEL_AAA() { %1 = call spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK0: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[MD_ID:[0-9]+]] +; CHECK0: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false) ret void } define dso_local spir_kernel void @KERNEL_BBB() { %1 = call spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*)) -; CHECK1: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[MD_ID:[0-9]+]] +; CHECK1: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false) ret void } -; CHECK: ![[MD_ID]] = !{!"SpecConst", i32 0} +; CHECK: !sycl.specialization-constants = !{![[#MD:]]} + +; CHECK: ![[#MD:]] = !{!"SpecConst", i32 0, i32 0, i32 1} diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 532af4203d565..ac1a505dd095e 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -11,6 +11,7 @@ #include "SpecConstants.h" #include "llvm/ADT/APInt.h" +#include "llvm/ADT/MapVector.h" #include "llvm/ADT/StringMap.h" #include "llvm/ADT/StringRef.h" #include "llvm/IR/InstIterator.h" @@ -43,10 +44,9 @@ constexpr char SPIRV_GET_SPEC_CONST_VAL[] = "__spirv_SpecConstant"; constexpr char SPIRV_GET_SPEC_CONST_COMPOSITE[] = "__spirv_SpecConstantComposite"; -// Metadata ID string added to calls to __spirv_SpecConstant to record the -// original symbolic spec constant ID. For composite spec constants it contains -// IDs of all scalar spec constants included into a composite -constexpr char SPEC_CONST_SYM_ID_MD_STRING[] = "SYCL_SPEC_CONST_SYM_ID"; +// Name of the metadata which holds a list of all specialization constants (with +// associated information) encountered in the module +constexpr char SPEC_CONST_MD_STRING[] = "sycl.specialization-constants"; void AssertRelease(bool Cond, const char *Msg) { if (!Cond) @@ -215,41 +215,10 @@ std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { return Res; } -void setSpecConstSymIDMetadata(Instruction *I, StringRef SymID, - ArrayRef IntIDs) { - LLVMContext &Ctx = I->getContext(); - SmallVector MDOperands; - MDOperands.push_back(MDString::get(Ctx, SymID)); - for (unsigned ID : IntIDs) - MDOperands.push_back( - ConstantAsMetadata::get(ConstantInt::get(Ctx, APInt(32, ID)))); - MDNode *Entry = MDNode::get(Ctx, MDOperands); - I->setMetadata(SPEC_CONST_SYM_ID_MD_STRING, Entry); -} - -std::pair> -getScalarSpecConstMetadata(const Instruction *I) { - const MDNode *N = I->getMetadata(SPEC_CONST_SYM_ID_MD_STRING); - if (!N) - return std::make_pair("", std::vector{}); - const auto *MDSym = cast(N->getOperand(0)); - const auto *MDInt = cast(N->getOperand(1)); - unsigned ID = static_cast( - cast(MDInt->getValue())->getValue().getZExtValue()); - std::vector Res(1); - Res[0].ID = ID; - // We need to add an additional byte if the type size is not evenly - // divisible by eight, which might be the case for i1, i.e. booleans - Res[0].Size = I->getType()->getPrimitiveSizeInBits() / 8 + - (I->getType()->getPrimitiveSizeInBits() % 8 != 0); - Res[0].Offset = 0; - return std::make_pair(MDSym->getString(), Res); -} - /// Recursively iterates over a composite type in order to collect information /// about its scalar elements. void collectCompositeElementsInfoRecursive( - const Module *M, Type *Ty, unsigned &Index, unsigned &Offset, + const Module &M, Type *Ty, unsigned &Index, unsigned &Offset, std::vector &Result) { if (auto *ArrTy = dyn_cast(Ty)) { for (size_t I = 0; I < ArrTy->getNumElements(); ++I) { @@ -260,7 +229,7 @@ void collectCompositeElementsInfoRecursive( Offset, Result); } } else if (auto *StructTy = dyn_cast(Ty)) { - const StructLayout *SL = M->getDataLayout().getStructLayout(StructTy); + const StructLayout *SL = M.getDataLayout().getStructLayout(StructTy); for (size_t I = 0, E = StructTy->getNumElements(); I < E; ++I) { auto *ElTy = StructTy->getElementType(I); // When handling elements of a structure, we do not use manually @@ -296,25 +265,50 @@ void collectCompositeElementsInfoRecursive( } } -std::pair> -getCompositeSpecConstMetadata(const Instruction *I) { - const MDNode *N = I->getMetadata(SPEC_CONST_SYM_ID_MD_STRING); - if (!N) - return std::make_pair("", std::vector{}); - const auto *MDSym = cast(N->getOperand(0)); - - std::vector Result(N->getNumOperands() - 1); - unsigned Index = 0, Offset = 0; - collectCompositeElementsInfoRecursive(I->getModule(), I->getType(), Index, - Offset, Result); - - for (unsigned I = 1; I < N->getNumOperands(); ++I) { - const auto *MDInt = cast(N->getOperand(I)); - unsigned ID = static_cast( - cast(MDInt->getValue())->getValue().getZExtValue()); - Result[I - 1].ID = ID; +MDNode *generateSpecConstantMetadata(const Module &M, StringRef SymbolicID, + Type *SCTy, ArrayRef IDs, + bool IsNativeSpecConstant) { + SmallVector MDOps; + LLVMContext &Ctx = M.getContext(); + auto *Int32Ty = Type::getInt32Ty(Ctx); + + // First element is always Symbolic ID + MDOps.push_back(MDString::get(Ctx, SymbolicID)); + + if (IsNativeSpecConstant) { + std::vector Result(IDs.size()); + unsigned Index = 0, Offset = 0; + collectCompositeElementsInfoRecursive(M, SCTy, Index, Offset, Result); + + for (unsigned I = 0; I < Result.size(); ++I) { + MDOps.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Int32Ty, APInt(32, IDs[I])))); + MDOps.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Int32Ty, APInt(32, Result[I].Offset)))); + MDOps.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Int32Ty, APInt(32, Result[I].Size)))); + } + } else { + assert(IDs.size() == 1 && + "There must be a single ID for emulated spec constant"); + MDOps.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Int32Ty, APInt(32, IDs[0])))); + // Second element is always zero here + MDOps.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Int32Ty, APInt(32, 0)))); + + unsigned Size = 0; + if (auto *StructTy = dyn_cast(SCTy)) { + const auto *SL = M.getDataLayout().getStructLayout(StructTy); + Size = SL->getSizeInBytes(); + } else + Size = SCTy->getScalarSizeInBits() / CHAR_BIT; + + MDOps.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Int32Ty, APInt(32, Size)))); } - return std::make_pair(MDSym->getString(), Result); + + return MDNode::get(Ctx, MDOps); } Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName, @@ -440,9 +434,10 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, unsigned NextOffset = 0; StringMap> IDMap; StringMap OffsetMap; + MapVector SCMetadata; // Iterate through all declarations of instances of function template - // template T __sycl_getSpecConstantValue(const char *ID) + // template T __sycl_get*SpecConstantValue(const char *ID) // intrinsic to find its calls and lower them depending on the SetValAtRT // setting (see below). bool IRModified = false; @@ -466,7 +461,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, IRModified = IRModified || (SCIntrCalls.size() > 0); for (auto *CI : SCIntrCalls) { - // 1. Find the symbolic ID (string literal) passed as the actual argument + // 1. Find the Symbolic ID (string literal) passed as the actual argument // to the intrinsic - this should always be possible, as only string // literals are passed to it in the SYCL RT source code, and application // code can't use this intrinsic directly. @@ -489,12 +484,12 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, SCTy = PtrTy->getElementType(); } StringRef SymID = getStringLiteralArg(CI, NameArgNo, DelInsts); + Value *Replacement = nullptr; if (SetValAtRT) { // 2. Spec constant value will be set at run time - then add the literal - // to a "spec const string literal ID" -> "integer ID" map or - // "composite spec const string literal ID" -> "vector of integer IDs" - // map, uniquing the integer IDs if this is new literal + // to a "spec const string literal ID" -> "vector of integer IDs" map, + // uniquing the integer IDs if this is a new literal auto Ins = IDMap.insert(std::make_pair(SymID, SmallVector{})); bool IsNewSpecConstant = Ins.second; @@ -524,44 +519,25 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // 3. Transform to spirv intrinsic _Z*__spirv_SpecConstant* or // _Z*__spirv_SpecConstantComposite - auto *SPIRVCall = - emitSpecConstantRecursive(SCTy, CI, IDs, DefaultValue); + Replacement = emitSpecConstantRecursive(SCTy, CI, IDs, DefaultValue); if (IsNewSpecConstant) { // emitSpecConstantRecursive might emit more than one spec constant - // (because of composite types) and therefore, we need to ajudst + // (because of composite types) and therefore, we need to adjust // NextID according to the actual amount of emitted spec constants. NextID += IDs.size(); - } - if (IsComposite) { - // __sycl_getCompositeSpecConstant returns through argument, so, the - // only thing we need to do here is to store into a memory pointed by - // that argument - new StoreInst(SPIRVCall, CI->getArgOperand(0), CI); - } else { - CI->replaceAllUsesWith(SPIRVCall); + // Generate necessary metadata which later will be pulled by + // sycl-post-link and transformed into device image properties + SCMetadata[SymID] = generateSpecConstantMetadata( + M, SymID, SCTy, IDs, /* is native spec constant */ true); } - - // Mark the instruction with list for later - // recollection by collectSpecConstantMetadata method. - setSpecConstSymIDMetadata(SPIRVCall, SymID, IDs); - // Example of the emitted call when spec constant is integer: - // %6 = call i32 @_Z20__spirv_SpecConstantii(i32 0, i32 0), \ - // !SYCL_SPEC_CONST_SYM_ID !22 - // !22 = {!"string-id", i32 0} - // Example of the emitted call when spec constant is vector consisting - // of two integers: - // %1 = call i32 @_Z20__spirv_SpecConstantii(i32 3, i32 0) - // %2 = call i32 @_Z20__spirv_SpecConstantii(i32 4, i32 0) - // %3 = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32 \ - // %1, i32 %2), !SYCL_SPEC_CONST_SYM_ID !23 - // !23 = {!"string-id-2", i32 3, i32 4} } else { - // 2a. Spec constant must be resolved at compile time - replace the - // intrinsic with the actual value for spec constant. - Value *Val = nullptr; + // 2a. For SYCL 2020: spec constant will be passed as kernel argument; + // For older proposal against SYCL 1.2.1 spec constant must be resolved + // at compile time - replace the intrinsic with the actual value for + // spec constant. if (Is2020Intrinsic) { - // Handle SYCL2020 version of intrinsic - replace it with a load from + // Handle SYCL 2020 version of intrinsic - replace it with a load from // the pointer to the specialization constant value. // A pointer to a single RT-buffer with all the values of // specialization constants is passed as a 3rd argument of intrinsic. @@ -573,23 +549,28 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // literal. auto Ins = OffsetMap.insert(std::make_pair(SymID, NextOffset)); bool IsNewSpecConstant = Ins.second; - auto CurrentOffset = Ins.first->second; + unsigned CurrentOffset = Ins.first->second; if (IsNewSpecConstant) { + unsigned Size = 0; if (IsComposite) { // When handling elements of a structure, we do not use manually // calculated offsets (which are sum of sizes of all previously // encountered elements), but instead rely on data provided for us // by DataLayout, because the structure can be unpacked, i.e. // padded in order to ensure particular alignment of its elements. - auto *StructTy = cast( - CI->getArgOperand(0)->getType()->getPointerElementType()); // We rely on the fact that the StructLayout of spec constant RT // values is the same for the host and the device. const StructLayout *SL = - M.getDataLayout().getStructLayout(StructTy); - NextOffset += SL->getSizeInBytes(); + M.getDataLayout().getStructLayout(cast(SCTy)); + Size = SL->getSizeInBytes(); } else - NextOffset += SCTy->getScalarSizeInBits() / CHAR_BIT; + Size = SCTy->getScalarSizeInBits() / CHAR_BIT; + + SCMetadata[SymID] = generateSpecConstantMetadata( + M, SymID, SCTy, NextID, /* is native spec constant */ false); + + ++NextID; + NextOffset += Size; } Type *Int8Ty = Type::getInt8Ty(CI->getContext()); @@ -601,62 +582,73 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, BitCastInst *BitCast = new BitCastInst( GEP, PointerType::get(SCTy, GEP->getAddressSpace()), "bc", CI); - LoadInst *Load = new LoadInst(SCTy, BitCast, "load", CI); - Val = Load; + Replacement = new LoadInst(SCTy, BitCast, "load", CI); } else { // Replace the intrinsic with default C++ value for the spec constant // type. - Val = getDefaultCPPValue(SCTy); + Replacement = getDefaultCPPValue(SCTy); } + } - if (IsComposite) { - // __sycl_getCompositeSpecConstant returns through argument, so, the - // only thing we need to do here is to store into a memory pointed - // by that argument - new StoreInst(Val, CI->getArgOperand(0), CI); - } else { - CI->replaceAllUsesWith(Val); - } + if (IsComposite) { + // __sycl_getCompositeSpecConstant returns through argument, so, the + // only thing we need to do here is to store into a memory pointed + // by that argument + new StoreInst(Replacement, CI->getArgOperand(0), CI); + } else { + CI->replaceAllUsesWith(Replacement); } for (auto *I : DelInsts) { - assert(I->getNumUses() == 0 && "removing live instruction"); I->removeFromParent(); I->deleteValue(); } } } + + // Emit metadata about encountered specializaiton constants. This metadata + // is later queried by sycl-post-link in order to be converted into device + // image properties. + // Generated metadata looks like: + // !sycl.specialization-constants = !{!1, !2, ... for each spec constant} + // !1 = !{!"SymbolicID1", i32 1, i32 0, i32 4, i32 2, i32 4, i32 8} + // !2 = !{!"SymbolicID2", i32 3, i32 0, i32 4} + // The format is [Symbolic ID, list of triplets: numeric ID, offset, size] + // For more infor about meaning of those triplets see comments about + // SpecConstantDescriptor structure in SpecConstants.h + NamedMDNode *MD = M.getOrInsertNamedMetadata(SPEC_CONST_MD_STRING); + for (const auto &P : SCMetadata) + MD->addOperand(P.second); + return IRModified ? PreservedAnalyses::none() : PreservedAnalyses::all(); } bool SpecConstantsPass::collectSpecConstantMetadata(Module &M, SpecIDMapTy &IDMap) { - bool Met = false; - - for (Function &F : M) { - if (F.isDeclaration()) - continue; - SmallVector SCIntrCalls; - - for (Instruction &I : instructions(F)) { - auto *CI = dyn_cast(&I); - Function *Callee = nullptr; - if (!CI || CI->isIndirectCall() || !(Callee = CI->getCalledFunction())) - continue; - - std::pair> Res; - if (Callee->getName().contains(SPIRV_GET_SPEC_CONST_COMPOSITE)) { - Res = getCompositeSpecConstMetadata(CI); - } else if (Callee->getName().contains(SPIRV_GET_SPEC_CONST_VAL)) { - Res = getScalarSpecConstMetadata(CI); - } + NamedMDNode *MD = M.getOrInsertNamedMetadata(SPEC_CONST_MD_STRING); + if (!MD) + return false; + + auto ExtractIntegerFromMDNodeOperand = [=](const MDNode *N, + unsigned OpNo) -> unsigned { + Constant *C = + cast(N->getOperand(OpNo).get())->getValue(); + return static_cast(C->getUniqueInteger().getZExtValue()); + }; - if (!Res.first.empty()) { - IDMap[Res.first] = Res.second; - Met = true; - } + for (const auto *Node : MD->operands()) { + StringRef ID = cast(Node->getOperand(0).get())->getString(); + assert((Node->getNumOperands() - 1) % 3 == 0 && + "Unexpected amount of operands"); + std::vector Descs((Node->getNumOperands() - 1) / 3); + for (unsigned NI = 1, I = 0; NI < Node->getNumOperands(); NI += 3, ++I) { + Descs[I].ID = ExtractIntegerFromMDNodeOperand(Node, NI + 0); + Descs[I].Offset = ExtractIntegerFromMDNodeOperand(Node, NI + 1); + Descs[I].Size = ExtractIntegerFromMDNodeOperand(Node, NI + 2); } + + IDMap[ID] = Descs; } - return Met; + return true; } diff --git a/llvm/tools/sycl-post-link/SpecConstants.h b/llvm/tools/sycl-post-link/SpecConstants.h index 88215a099d02d..505b37c3c97a7 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.h +++ b/llvm/tools/sycl-post-link/SpecConstants.h @@ -11,20 +11,21 @@ // constant operations. The spec constant IDs are symbolic before linkage to // make separate compilation possible. After linkage all spec constants are // available to the pass, and it can assign consistent integer IDs. -// -// The pass is used w/o a pass manager currently, but the interface is based on -// the standard Module pass interface to move it around easier in future. //===----------------------------------------------------------------------===// #pragma once -#include "llvm/ADT/StringMap.h" #include "llvm/IR/Module.h" #include "llvm/IR/PassManager.h" -using namespace llvm; +#include +#include + +namespace llvm { +class StringRef; +} -// Represents either an element of a composite speciailization constant or a +// Represents either an element of a composite specialization constant or a // single scalar specialization constant - at SYCL RT level composite // specialization constants are being represented as a single byte-array, while // at SPIR-V level they are represented by a number of scalar specialization @@ -44,23 +45,22 @@ struct SpecConstantDescriptor { // Encodes size of scalar specialization constant. unsigned Size; }; -using SpecIDMapTy = std::map>; +using SpecIDMapTy = + std::map>; -class SpecConstantsPass : public PassInfoMixin { +class SpecConstantsPass : public llvm::PassInfoMixin { public: // SetValAtRT parameter controls spec constant lowering mode: // - if true, it is lowered to SPIRV intrinsic which retrieves constant value // - if false, it is replaced with C++ default (used for AOT compilers) SpecConstantsPass(bool SetValAtRT = true) : SetValAtRT(SetValAtRT) {} - PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM); + llvm::PreservedAnalyses run(llvm::Module &M, + llvm::ModuleAnalysisManager &MAM); - // Searches given module for occurences of specialization constant-specific - // metadata at call instructions and builds a - // "spec constant name" -> "spec constant int ID" map for scalar spec - // constants and - // "spec constant name" -> vector<"spec constant int ID"> map for composite - // spec constants - static bool collectSpecConstantMetadata(Module &M, SpecIDMapTy &IDMap); + // Searches given module for occurrences of specialization constant-specific + // metadata and builds "spec constant name" -> vector<"spec constant int ID"> + // map + static bool collectSpecConstantMetadata(llvm::Module &M, SpecIDMapTy &IDMap); private: bool SetValAtRT; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 9bb3a4739ef11..ce3f85185a7c9 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -410,7 +410,7 @@ static string_vector saveDeviceImageProperty( PropSet.add(llvm::util::PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK, RMEntry); } - if (ImgPSInfo.DoSpecConst && ImgPSInfo.SetSpecConstAtRT) { + if (ImgPSInfo.DoSpecConst) { if (ImgPSInfo.SpecConstsMet) { // extract spec constant maps per each module SpecIDMapTy TmpSpecIDMap; diff --git a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp new file mode 100644 index 0000000000000..8bfb21a2f11d3 --- /dev/null +++ b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp @@ -0,0 +1,102 @@ +// RUN: %clangxx -fsycl -fsycl-device-only -c -o %t.bc %s +// RUN: sycl-post-link %t.bc -spec-const=rt -o %t-split1.txt +// RUN: cat %t-split1_0.prop | FileCheck %s +// RUN: sycl-post-link %t.bc -spec-const=default -o %t-split2.txt +// RUN: cat %t-split2_0.prop | FileCheck %s +// RUN: llvm-spirv -o %t-split1_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t-split1_0.bc +// RUN: llvm-spirv -o %t-split2_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t-split2_0.bc +// +//==----------- SYCL-2020-spec-constants.cpp -------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// The test checks that the tool chain correctly identifies all specialization +// constants, emits correct specialization constats map file and can properly +// translate the resulting bitcode to SPIR-V. + +#include + +#include +#include + +// FIXME: there is a bug with generation SPIR-V friendly IR for boolean spec +// constants, which was likely interoduced by https://github.com/intel/llvm/pull/3513 +// constexpr sycl::specialization_id bool_id; +constexpr sycl::specialization_id int8_id(1); +constexpr sycl::specialization_id uint8_id(2); +constexpr sycl::specialization_id int16_id(3); +constexpr sycl::specialization_id uint16_id(4); +constexpr sycl::specialization_id int32_id(5); +constexpr sycl::specialization_id uint32_id(6); +constexpr sycl::specialization_id int64_id(7); +constexpr sycl::specialization_id uint64_id(8); +// FIXME: enable once we support constexpr constructor for half +// constexpr sycl::specialization_id half_id(9.0); +constexpr sycl::specialization_id float_id(10.0); +constexpr sycl::specialization_id double_id(11.0); + +struct composite { + int a; + int b; + constexpr composite(int a, int b): a(a), b(b) {} +}; + +constexpr sycl::specialization_id composite_id(12, 13); + +class SpecializedKernel; + +int main() { + sycl::queue queue; + + std::vector vec(1); + { + sycl::buffer buf(vec.data(), vec.size()); + queue.submit([&](sycl::handler &h) { + auto acc = buf.get_access(h); + + h.single_task([=](sycl::kernel_handler kh) { + // see FIXME above about bool type support + // auto i1 = kh.get_specialization_constant(); + auto i8 = kh.get_specialization_constant(); + auto u8 = kh.get_specialization_constant(); + auto i16 = kh.get_specialization_constant(); + auto u16 = kh.get_specialization_constant(); + auto i32 = kh.get_specialization_constant(); + auto u32 = kh.get_specialization_constant(); + auto i64 = kh.get_specialization_constant(); + auto u64 = kh.get_specialization_constant(); + // see FIXME above about half type support + // auto f16 = kh.get_specialization_constant(); + auto f32 = kh.get_specialization_constant(); + auto f64 = kh.get_specialization_constant(); + auto c = kh.get_specialization_constant(); + + // see FIXMEs above about bool and half types support + acc[0] = /*i1 +*/ i8 + u8 + i16 + u16 + i32 + u32 + i64 + u64 + + //f16.get() + + f32 + f64 + c.a + c.b; + }); + }); + } + + return 0; +} + +// CHECK: [SYCL/specialization constants] +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL12composite_idEEE=2| +// See FIXME above about bool type support +// CHECK-disabled: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL7bool_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL7int8_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL8float_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL8int16_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL8int32_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL8int64_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL8uint8_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL9double_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL9uint16_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL9uint32_idEEE=2| +// CHECK-DAG: _ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL9uint64_idEEE=2| +// FIXME: check line for half constant