From c3f55a53e59ada8541df6438f151e7bb03960d3a Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 9 Sep 2021 19:49:50 +0300 Subject: [PATCH 1/3] Fix handling of composites in sycl-post-link Enchanced our mini-mangler in SpecConstants pass to avoid generating `call bitcast` construct when encountering almost the same functions For example, if you have a call to `%struct.A __spriv_SpecConstantComposite(i32, i32)` and you try to insert a call to `%struct.B __spirv_SpecConstantComposite(i32, i32)` you should be having two calls to two differently mangled functions instead of `call bitcast` construct, which confuses other toolchain components. --- .../spec-constants/SYCL-2020.ll | 64 ++++++++++++++++--- .../spec-constants/composite-O0.ll | 8 +-- .../spec-constants/composite-O2.ll | 14 ++-- .../spec-constants/composite-no-sret.ll | 4 +- .../multiple-composite-usages-2.ll | 4 +- .../multiple-composite-usages.ll | 6 +- llvm/tools/sycl-post-link/SpecConstants.cpp | 17 ++++- 7 files changed, 88 insertions(+), 29 deletions(-) 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 c57d493aa2e8b..bc0ddbe2a36d9 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 @@ -15,10 +15,17 @@ %class.specialization_id.2 = type { %struct.ComposConst2 } %struct.ComposConst2 = type { i8, %struct.myConst, double } +%struct.VectorConst = type { <2 x i32> } +%class.specialization_id.3 = type { %struct.VectorConst } +%struct.MArrayConst = type { [2 x i32] } +%class.specialization_id.4 = type { %struct.MArrayConst } + @id_half = dso_local global %class.specialization_id { half 0xH4000 }, align 8 @id_int = dso_local global %class.specialization_id.0 { i32 42 }, align 4 @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 +@id_vector = dso_local global %class.specialization_id.3 { %struct.VectorConst { <2 x i32> } }, align 8 +@id_marray = dso_local global %class.specialization_id.4 { %struct.MArrayConst { [2 x i32] [i32 1, i32 2] } }, align 8 ; 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 @@ -33,6 +40,8 @@ @__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 @__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 @__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 +@__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_vectorE17specialization_idI11VectorConstES1_ET1_v = private unnamed_addr constant [38 x i8] c"_ZTS14name_generatorIL_Z10id_vectorEE\00", align 1 +@__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_marrayE17specialization_idI11MArrayConstES1_ET1_v = private unnamed_addr constant [38 x i8] c"_ZTS14name_generatorIL_Z10id_marrayEE\00", align 1 ; CHECK-LABEL: define dso_local void @_Z4testv define dso_local void @_Z4testv() local_unnamed_addr #0 { @@ -77,8 +86,8 @@ entry: ; CHECK-RT: %[[#SE2:]] = call double @_Z20__spirv_SpecConstantid(i32 [[#SCID3:]], double 2.000000e+00) ; 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]]) +; CHECK-RT: %[[#CE1:]] = call %struct.myConst @_Z29__spirv_SpecConstantCompositeif_Rstruct.myConst(i32 %[[#SE3]], float %[[#SE4]]) +; CHECK-RT: %[[C1:[0-9a-z]+]] = call %struct.ComposConst @_Z29__spirv_SpecConstantCompositeidstruct.myConst_Rstruct.ComposConst(i32 %[[#SE1]], double %[[#SE2]], %struct.myConst %[[#CE1]]) ; ; CHECK: store %struct.ComposConst %[[C1]], %struct.ComposConst* @@ -93,9 +102,9 @@ entry: ; CHECK-RT: %[[#SE1:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID6:]], i8 1) ; CHECK-RT: %[[#SE2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID7:]], i32 52) ; 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: %[[#CE1:]] = call %struct.myConst @_Z29__spirv_SpecConstantCompositeif_Rstruct.myConst(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]]) +; CHECK-RT: %[[C2:[0-9a-z]+]] = call %struct.ComposConst2 @_Z29__spirv_SpecConstantCompositeastruct.myConstd_Rstruct.ComposConst2(i8 %[[#SE1]], %struct.myConst %[[#CE1]], double %[[#SE4]]) ; ; CHECK: store %struct.ComposConst2 %[[C2]], %struct.ComposConst2* @@ -110,13 +119,40 @@ entry: ; CHECK-RT: %[[#SE2:]] = call double @_Z20__spirv_SpecConstantid(i32 [[#SCID3]], double 2.000000e+00) ; 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]]) +; CHECK-RT: %[[#CE1:]] = call %struct.myConst @_Z29__spirv_SpecConstantCompositeif_Rstruct.myConst(i32 %[[#SE3]], float %[[#SE4]]) +; CHECK-RT: %[[C3:[0-9a-z]+]] = call %struct.ComposConst @_Z29__spirv_SpecConstantCompositeidstruct.myConst_Rstruct.ComposConst(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 ret void } + +define void @test3() { + %tmp = alloca %struct.VectorConst, align 8 + %tmp1 = alloca %struct.MArrayConst, align 8 + %1 = bitcast %struct.VectorConst* %tmp to i8* +; CHECK-DEF: %[[GEP1:[0-9a-z]+]] = getelementptr i8, i8* null, i32 54 +; CHECK-DEF: %[[BITCAST1:[0-9a-z]+]] = bitcast i8* %[[GEP1]] to %struct.VectorConst* +; CHECK-DEF: %[[C1:[0-9a-z]+]] = load %struct.VectorConst, %struct.VectorConst* %[[BITCAST1]], align 8 +; +; CHECK-RT: %[[#SE1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID10:]], i32 1) +; CHECK-RT: %[[#SE2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID11:]], i32 2) +; CHECK-RT: %[[#CE1:]] = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii_RDv2_i(i32 %[[#SE1]], i32 %[[#SE2]]) +; CHECK-RT: call %struct.VectorConst @_Z29__spirv_SpecConstantCompositeDv2_i_Rstruct.VectorConst(<2 x i32> %[[#CE1]]) + call void @_Z40__sycl_getComposite2020SpecConstantValueI11VectorConstET_PKcPvS4_(%struct.VectorConst* nonnull sret(%struct.VectorConst) align 8 %tmp, i8* getelementptr inbounds ([38 x i8], [38 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_vectorE17specialization_idI11VectorConstES1_ET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id.3* @id_vector to i8*), i8* null) + %2 = bitcast %struct.MArrayConst* %tmp1 to i8* +; CHECK-DEF: %[[GEP2:[0-9a-z]+]] = getelementptr i8, i8* null, i32 62 +; CHECK-DEF: %[[BITCAST2:[0-9a-z]+]] = bitcast i8* %[[GEP2]] to %struct.MArrayConst* +; CHECK-DEF: %[[C2:[0-9a-z]+]] = load %struct.MArrayConst, %struct.MArrayConst* %[[BITCAST2]], align 4 +; +; CHECK-RT: %[[#SE1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID12:]], i32 1) +; CHECK-RT: %[[#SE2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID13:]], i32 2) +; CHECK-RT: %[[#CE1:]] = call [2 x i32] @_Z29__spirv_SpecConstantCompositeii_RA2_i(i32 %[[#SE1]], i32 %[[#SE2]]) +; CHECK-RT: call %struct.MArrayConst @_Z29__spirv_SpecConstantCompositeA2_i_Rstruct.MArrayConst([2 x i32] %[[#CE1]]) + call void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConstET_PKcPvS4_(%struct.MArrayConst* nonnull sret(%struct.MArrayConst) align 8 %tmp1, i8* getelementptr inbounds ([38 x i8], [38 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_marrayE17specialization_idI11MArrayConstES1_ET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id.4* @id_marray to i8*), i8* null) + ret void +} + declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 @@ -131,14 +167,18 @@ declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11ComposCon declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI12ComposConst2ET_PKcPvS4_(%struct.ComposConst2* sret(%struct.ComposConst2) align 8, i8*, i8*, i8*) local_unnamed_addr #2 +declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11VectorConstET_PKcPvS4_(%struct.VectorConst* sret(%struct.VectorConst) align 8, i8*, i8*, i8*) local_unnamed_addr #2 + +declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConstET_PKcPvS4_(%struct.MArrayConst* sret(%struct.MArrayConst) align 8, i8*, i8*, i8*) local_unnamed_addr #2 + attributes #0 = { uwtable mustprogress "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-jump-tables"="false" "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 #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: !sycl.specialization-constants = !{![[#ID0:]], ![[#ID1:]], ![[#ID2:]], ![[#ID3:]]} +; CHECK: !sycl.specialization-constants = !{![[#ID0:]], ![[#ID1:]], ![[#ID2:]], ![[#ID3:]], ![[#ID4:]], ![[#ID5:]]} ; -; CHECK-DEF: !sycl.specialization-constants-default-values = !{![[#ID4:]], ![[#ID5:]], ![[#ID6:]], ![[#ID7:]]} +; CHECK-DEF: !sycl.specialization-constants-default-values = !{![[#ID4:]], ![[#ID5:]], ![[#ID6:]], ![[#ID7:]], ![[#ID8:]], ![[#ID9:]]} ; CHECK-RT-NOT: !sycl.specialization-constants-default-values ; ; CHECK: ![[#ID0]] = !{!"_ZTS14name_generatorIL_Z9id_halfEE", i32 0, i32 0, i32 2} @@ -162,3 +202,11 @@ attributes #3 = { nounwind } ; CHECK-DEF: ![[#ID5]] = !{i32 42} ; CHECK-DEF: ![[#ID6]] = !{%struct.ComposConst { i32 1, double 2.000000e+00, %struct.myConst { i32 13, float 0x4020666660000000 } }} ; CHECK-DEF: ![[#ID7]] = !{%struct.ComposConst2 { i8 1, %struct.myConst { i32 52, float 0x40479999A0000000 }, double 2.000000e+00 }} +; +; CHECK-DEF: ![[#ID8]] = !{%struct.VectorConst { <2 x i32> }} +; CHECK-DEF: ![[#ID9]] = !{%struct.MArrayConst { [2 x i32] [i32 1, i32 2] }} +; +; CHECK-RT: ![[#ID4]] = !{!"_ZTS14name_generatorIL_Z10id_vectorEE", i32 [[#SCID10]], i32 0, i32 4, +; CHECK-RT-SAME: i32 [[#SCID11]], i32 4, i32 4} +; CHECK-RT: ![[#ID5]] = !{!"_ZTS14name_generatorIL_Z10id_marrayEE", i32 [[#SCID12]], i32 0, i32 4, +; CHECK-RT-SAME: i32 [[#SCID13]], i32 4, i32 4} 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 355ea4a7dd157..13c8d3094eb91 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 @@ -10,17 +10,17 @@ ; ; 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]]) +; CHECK: %[[#NA0:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif_Rstruct._ZTS1A.A(i32 %[[#NS0]], float %[[#NS1]]) ; ; CHECK: %[[#NS2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 2]], i32 ; CHECK: %[[#NS3:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 3]], float -; CHECK: %[[#NA1:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS2]], float %[[#NS3]]) +; CHECK: %[[#NA1:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif_Rstruct._ZTS1A.A(i32 %[[#NS2]], float %[[#NS3]]) ; -; CHECK: %[[#NA:]] = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A %[[#NA0]], %struct._ZTS1A.A %[[#NA1]]) +; CHECK: %[[#NA:]] = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A_RA2_struct._ZTS1A.A(%struct._ZTS1A.A %[[#NA0]], %struct._ZTS1A.A %[[#NA1]]) ; ; 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]]) +; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @_Z29__spirv_SpecConstantCompositeA2_struct._ZTS1A.Ai_Rstruct._ZTS3POD.POD([2 x %struct._ZTS1A.A] %[[#NA]], i32 %[[#B]]) ; CHECK: store %struct._ZTS3POD.POD %[[#POD]] ; ; CHECK: !sycl.specialization-constants = !{![[#MD:]]} 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 9289c10e6f8c3..a1ca6da8e1e3a 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 @@ -8,26 +8,26 @@ ; 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]]) +; CHECK: %[[#NA0:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif_Rstruct._ZTS1A.A(i32 %[[#NS0]], float %[[#NS1]]) ; ; CHECK: %[[#NS2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 2]], i32 ; CHECK: %[[#NS3:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 3]], float -; CHECK: %[[#NA1:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS2]], float %[[#NS3]]) +; CHECK: %[[#NA1:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif_Rstruct._ZTS1A.A(i32 %[[#NS2]], float %[[#NS3]]) ; -; CHECK: %[[#NA:]] = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A %[[#NA0]], %struct._ZTS1A.A %[[#NA1]]) +; CHECK: %[[#NA:]] = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A_RA2_struct._ZTS1A.A(%struct._ZTS1A.A %[[#NA0]], %struct._ZTS1A.A %[[#NA1]]) ; ; CHECK: %[[#B0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 4]], i32{{.*}}) ; CHECK: %[[#B1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 5]], i32{{.*}}) -; 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: %[[#BV:]] = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii_RDv2_i(i32 %[[#B0]], i32 %[[#B1]]) +; CHECK: %[[#B:]] = call %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" @"_Z29__spirv_SpecConstantCompositeDv2_i_Rclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"(<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]]) +; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeA2_struct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec_Rstruct._ZTS3POD.POD"([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: %[[#CONST:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif_Rstruct._ZTS1A.A(i32 %[[#N0]], float %[[#N1]]) ; CHECK: %struct._ZTS1A.A %[[#CONST]] ; ; CHECK: !sycl.specialization-constants = !{![[#MD0:]], ![[#MD1:]]} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll index c173a2f79abeb..e7128b3ad97b0 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll @@ -3,10 +3,10 @@ ; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32 ; CHECK: %[[#NS1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 1]], i32 42) -; CHECK: %[[#NA0:]] = call %struct._ZTS10TestStruct.TestStruct @_Z29__spirv_SpecConstantCompositeii(i32 %[[#NS0]], i32 %[[#NS1]]) +; CHECK: %[[#NA0:]] = call %struct._ZTS10TestStruct.TestStruct @_Z29__spirv_SpecConstantCompositeii_Rstruct._ZTS10TestStruct.TestStruct(i32 %[[#NS0]], i32 %[[#NS1]]) ; CHECK: declare i32 @_Z20__spirv_SpecConstantii(i32, i32) -; CHECK: declare %struct._ZTS10TestStruct.TestStruct @_Z29__spirv_SpecConstantCompositeii(i32, i32) +; CHECK: declare %struct._ZTS10TestStruct.TestStruct @_Z29__spirv_SpecConstantCompositeii_Rstruct._ZTS10TestStruct.TestStruct(i32, i32) ; CHECK: !sycl.specialization-constants = !{![[#MD:]]} ; CHECK: ![[#MD]] = !{!"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL10SpecConst3EEE", i32 [[#ID]], i32 0, i32 4, 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 86738c7736c4e..4d36c451f91e2 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,11 +9,11 @@ ; 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]]) +; CHECK: call {{.*}} @"_Z29__spirv_SpecConstantCompositeff_Rstruct._ZTSN4test5pod_tE.test::pod_t"(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]]) +; CHECK: call {{.*}} @"_Z29__spirv_SpecConstantCompositeff_Rstruct._ZTSN4test5pod_tE.test::pod_t"(float %[[#X2]], float %[[#Y2]]) ; CHECK: !sycl.specialization-constants = !{![[#ID:]] 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 bcaf2392d3f9e..26213f8f80a74 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,11 +6,11 @@ ; once ; ; CHECK-LABEL: @foo1 -; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}) +; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeA2_struct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec_Rstruct._ZTS3POD.POD"({{.*}}) ; CHECK-LABEL: @_ZTS4Test -; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}) +; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeA2_struct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec_Rstruct._ZTS3POD.POD"({{.*}}) ; CHECK-LABEL: @foo2 -; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}) +; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeA2_struct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec_Rstruct._ZTS3POD.POD"({{.*}}) ; CHECK: !sycl.specialization-constants = !{![[#MD0:]], ![[#MD1:]] ; diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 414e171a311bb..858df441c48a2 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -192,8 +192,8 @@ std::string manglePrimitiveType(const Type *T) { llvm_unreachable("unsupported spec const integer type"); } } - // Mangling, which is generated below is not conformant with C++ ABI rules - // (https://itanium-cxx-abi.github.io/cxx-abi/abi.html#mangle.unqualified-name) + // Mangling, which is generated below is not fully conformant with C++ ABI + // rules (https://itanium-cxx-abi.github.io/cxx-abi/abi.html#mangle.unqualified-name) // But it should be more or less okay, because these declarations only // exists in the module between invocations of sycl-post-link and llvm-spirv, // llvm-spirv doesn't care about the mangling and the only intent here is to @@ -202,7 +202,9 @@ std::string manglePrimitiveType(const Type *T) { if (T->isStructTy()) return T->getStructName().str(); if (T->isArrayTy()) - return "A" + manglePrimitiveType(T->getArrayElementType()); + return "A" + std::to_string(T->getArrayNumElements()) + + "_" + manglePrimitiveType(T->getArrayElementType()); + if (auto *VecTy = dyn_cast(T)) return "Dv" + std::to_string(VecTy->getNumElements()) + "_" + manglePrimitiveType(VecTy->getElementType()); @@ -217,6 +219,15 @@ std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { (Twine("_Z") + Twine(BaseName.size()) + Twine(BaseName)).str(); for (unsigned I = 0; I < FT->getNumParams(); ++I) Res += manglePrimitiveType(FT->getParamType(I)); + if (FT->getReturnType()->isArrayTy() || + FT->getReturnType()->isStructTy() || + FT->getReturnType()->isVectorTy()) { + // It is possible that we need to generate several calls to + // __spirv_SpecConstantComposite, accepting the same argument types, but + // returning different types. Therefore, we incorporate the return type into + // the mangling name as well to distinguish between those functions + Res += "_R" + manglePrimitiveType(FT->getReturnType()); + } return Res; } From c88f6ffc433b324c3e3359d4887cc1c027ff2c64 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 10 Sep 2021 21:59:22 +0300 Subject: [PATCH 2/3] Apply comments --- .../spec-constants/SYCL-2020.ll | 25 +++++++++++++++++-- llvm/tools/sycl-post-link/SpecConstants.cpp | 20 ++++++++------- 2 files changed, 34 insertions(+), 11 deletions(-) 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 bc0ddbe2a36d9..903b2bca86733 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 @@ -19,6 +19,10 @@ %class.specialization_id.3 = type { %struct.VectorConst } %struct.MArrayConst = type { [2 x i32] } %class.specialization_id.4 = type { %struct.MArrayConst } +%struct.MArrayConst2 = type { [3 x i32] } +%class.specialization_id.5 = type { %struct.MArrayConst2 } +%struct.MArrayConst3 = type { [3 x i64] } +%class.specialization_id.6 = type { %struct.MArrayConst3 } @id_half = dso_local global %class.specialization_id { half 0xH4000 }, align 8 @id_int = dso_local global %class.specialization_id.0 { i32 42 }, align 4 @@ -26,6 +30,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 @id_vector = dso_local global %class.specialization_id.3 { %struct.VectorConst { <2 x i32> } }, align 8 @id_marray = dso_local global %class.specialization_id.4 { %struct.MArrayConst { [2 x i32] [i32 1, i32 2] } }, align 8 +@id_marray2 = dso_local global %class.specialization_id.5 { %struct.MArrayConst2 { [3 x i32] [i32 1, i32 2, i32 3] } }, align 8 +@id_marray3 = dso_local global %class.specialization_id.6 { %struct.MArrayConst3 { [3 x i64] [i64 1, i64 2, i64 3] } }, align 8 ; 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 @@ -42,6 +48,8 @@ @__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 @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_vectorE17specialization_idI11VectorConstES1_ET1_v = private unnamed_addr constant [38 x i8] c"_ZTS14name_generatorIL_Z10id_vectorEE\00", align 1 @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_marrayE17specialization_idI11MArrayConstES1_ET1_v = private unnamed_addr constant [38 x i8] c"_ZTS14name_generatorIL_Z10id_marrayEE\00", align 1 +@__builtin_unique_stable_name.id_marray2 = private unnamed_addr constant [39 x i8] c"_ZTS14name_generatorIL_Z10id_marray2EE\00", align 1 +@__builtin_unique_stable_name.id_marray3 = private unnamed_addr constant [39 x i8] c"_ZTS14name_generatorIL_Z10id_marray3EE\00", align 1 ; CHECK-LABEL: define dso_local void @_Z4testv define dso_local void @_Z4testv() local_unnamed_addr #0 { @@ -130,6 +138,8 @@ entry: define void @test3() { %tmp = alloca %struct.VectorConst, align 8 %tmp1 = alloca %struct.MArrayConst, align 8 + %tmp2 = alloca %struct.MArrayConst2, align 8 + %tmp3 = alloca %struct.MArrayConst3, align 8 %1 = bitcast %struct.VectorConst* %tmp to i8* ; CHECK-DEF: %[[GEP1:[0-9a-z]+]] = getelementptr i8, i8* null, i32 54 ; CHECK-DEF: %[[BITCAST1:[0-9a-z]+]] = bitcast i8* %[[GEP1]] to %struct.VectorConst* @@ -150,6 +160,13 @@ define void @test3() { ; CHECK-RT: %[[#CE1:]] = call [2 x i32] @_Z29__spirv_SpecConstantCompositeii_RA2_i(i32 %[[#SE1]], i32 %[[#SE2]]) ; CHECK-RT: call %struct.MArrayConst @_Z29__spirv_SpecConstantCompositeA2_i_Rstruct.MArrayConst([2 x i32] %[[#CE1]]) call void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConstET_PKcPvS4_(%struct.MArrayConst* nonnull sret(%struct.MArrayConst) align 8 %tmp1, i8* getelementptr inbounds ([38 x i8], [38 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_marrayE17specialization_idI11MArrayConstES1_ET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id.4* @id_marray to i8*), i8* null) +; Here we only check the mangling of generated __spirv_SpecConstantComposite function + %3 = bitcast %struct.MArrayConst2* %tmp2 to i8* +; CHECK-RT: call %struct.MArrayConst2 @_Z29__spirv_SpecConstantCompositeA3_i_Rstruct.MArrayConst2 + call void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst2ET_PKcPvS4_(%struct.MArrayConst2* nonnull sret(%struct.MArrayConst2) align 8 %tmp2, i8* getelementptr inbounds ([39 x i8], [39 x i8]* @__builtin_unique_stable_name.id_marray2, i64 0, i64 0), i8* bitcast (%class.specialization_id.5* @id_marray2 to i8*), i8* null) + %4 = bitcast %struct.MArrayConst3* %tmp3 to i8* +; CHECK-RT: call %struct.MArrayConst3 @_Z29__spirv_SpecConstantCompositeA3_x_Rstruct.MArrayConst3 + call void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst3ET_PKcPvS4_(%struct.MArrayConst3* nonnull sret(%struct.MArrayConst3) align 8 %tmp3, i8* getelementptr inbounds ([39 x i8], [39 x i8]* @__builtin_unique_stable_name.id_marray3, i64 0, i64 0), i8* bitcast (%class.specialization_id.6* @id_marray3 to i8*), i8* null) ret void } @@ -171,14 +188,18 @@ declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11VectorCon declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConstET_PKcPvS4_(%struct.MArrayConst* sret(%struct.MArrayConst) align 8, i8*, i8*, i8*) local_unnamed_addr #2 +declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst2ET_PKcPvS4_(%struct.MArrayConst2* sret(%struct.MArrayConst2) align 8, i8*, i8*, i8*) local_unnamed_addr #2 + +declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst3ET_PKcPvS4_(%struct.MArrayConst3* sret(%struct.MArrayConst3) align 8, i8*, i8*, i8*) local_unnamed_addr #2 + attributes #0 = { uwtable mustprogress "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-jump-tables"="false" "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 #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: !sycl.specialization-constants = !{![[#ID0:]], ![[#ID1:]], ![[#ID2:]], ![[#ID3:]], ![[#ID4:]], ![[#ID5:]]} +; CHECK: !sycl.specialization-constants = !{![[#ID0:]], ![[#ID1:]], ![[#ID2:]], ![[#ID3:]], ![[#ID4:]], ![[#ID5:]] ; -; CHECK-DEF: !sycl.specialization-constants-default-values = !{![[#ID4:]], ![[#ID5:]], ![[#ID6:]], ![[#ID7:]], ![[#ID8:]], ![[#ID9:]]} +; CHECK-DEF: !sycl.specialization-constants-default-values = !{![[#ID4:]], ![[#ID5:]], ![[#ID6:]], ![[#ID7:]], ![[#ID8:]], ![[#ID9:]] ; CHECK-RT-NOT: !sycl.specialization-constants-default-values ; ; CHECK: ![[#ID0]] = !{!"_ZTS14name_generatorIL_Z9id_halfEE", i32 0, i32 0, i32 2} diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 858df441c48a2..cb3275fe7749e 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -169,7 +169,7 @@ Value *getDefaultCPPValue(Type *T) { return nullptr; } -std::string manglePrimitiveType(const Type *T) { +std::string mangleType(const Type *T) { if (T->isFloatTy()) return "f"; if (T->isDoubleTy()) @@ -193,7 +193,8 @@ std::string manglePrimitiveType(const Type *T) { } } // Mangling, which is generated below is not fully conformant with C++ ABI - // rules (https://itanium-cxx-abi.github.io/cxx-abi/abi.html#mangle.unqualified-name) + // rules + // (https://itanium-cxx-abi.github.io/cxx-abi/abi.html#mangle.unqualified-name) // But it should be more or less okay, because these declarations only // exists in the module between invocations of sycl-post-link and llvm-spirv, // llvm-spirv doesn't care about the mangling and the only intent here is to @@ -202,31 +203,32 @@ std::string manglePrimitiveType(const Type *T) { if (T->isStructTy()) return T->getStructName().str(); if (T->isArrayTy()) - return "A" + std::to_string(T->getArrayNumElements()) + - "_" + manglePrimitiveType(T->getArrayElementType()); + return "A" + std::to_string(T->getArrayNumElements()) + "_" + + mangleType(T->getArrayElementType()); if (auto *VecTy = dyn_cast(T)) return "Dv" + std::to_string(VecTy->getNumElements()) + "_" + - manglePrimitiveType(VecTy->getElementType()); + mangleType(VecTy->getElementType()); llvm_unreachable("unsupported spec const type"); return ""; } // This is a very basic mangler which can mangle non-templated and non-member // functions with primitive types in the signature. +// FIXME: generated mangling is not always complies with C++ ABI rules and might +// not be demanglable. Consider fixing this. std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { std::string Res = (Twine("_Z") + Twine(BaseName.size()) + Twine(BaseName)).str(); for (unsigned I = 0; I < FT->getNumParams(); ++I) - Res += manglePrimitiveType(FT->getParamType(I)); - if (FT->getReturnType()->isArrayTy() || - FT->getReturnType()->isStructTy() || + Res += mangleType(FT->getParamType(I)); + if (FT->getReturnType()->isArrayTy() || FT->getReturnType()->isStructTy() || FT->getReturnType()->isVectorTy()) { // It is possible that we need to generate several calls to // __spirv_SpecConstantComposite, accepting the same argument types, but // returning different types. Therefore, we incorporate the return type into // the mangling name as well to distinguish between those functions - Res += "_R" + manglePrimitiveType(FT->getReturnType()); + Res += "_R" + mangleType(FT->getReturnType()); } return Res; } From bbcc8352a8571a17f2241025fa1bc8168e847813 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 10 Sep 2021 22:14:38 +0300 Subject: [PATCH 3/3] Add one more test --- .../tools/sycl-post-link/spec-constants/SYCL-2020.ll | 10 ++++++++++ 1 file changed, 10 insertions(+) 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 903b2bca86733..88975f11777b8 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 @@ -23,6 +23,8 @@ %class.specialization_id.5 = type { %struct.MArrayConst2 } %struct.MArrayConst3 = type { [3 x i64] } %class.specialization_id.6 = type { %struct.MArrayConst3 } +%struct.MArrayConst4 = type { [2 x [2 x [3 x i32]]] } +%class.specialization_id.7 = type { %struct.MArrayConst4 } @id_half = dso_local global %class.specialization_id { half 0xH4000 }, align 8 @id_int = dso_local global %class.specialization_id.0 { i32 42 }, align 4 @@ -32,6 +34,7 @@ @id_marray = dso_local global %class.specialization_id.4 { %struct.MArrayConst { [2 x i32] [i32 1, i32 2] } }, align 8 @id_marray2 = dso_local global %class.specialization_id.5 { %struct.MArrayConst2 { [3 x i32] [i32 1, i32 2, i32 3] } }, align 8 @id_marray3 = dso_local global %class.specialization_id.6 { %struct.MArrayConst3 { [3 x i64] [i64 1, i64 2, i64 3] } }, align 8 +@id_marray4 = dso_local global %class.specialization_id.7 { %struct.MArrayConst4 { [2 x [2 x [3 x i32]]] [[2 x [3 x i32]] [[3 x i32] [i32 1, i32 2, i32 3], [3 x i32] [i32 1, i32 2, i32 3]], [2 x [3 x i32]] [[3 x i32] [i32 1, i32 2, i32 3], [3 x i32] [i32 1, i32 2, i32 3]]] } }, align 8 ; 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 @@ -50,6 +53,7 @@ @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_marrayE17specialization_idI11MArrayConstES1_ET1_v = private unnamed_addr constant [38 x i8] c"_ZTS14name_generatorIL_Z10id_marrayEE\00", align 1 @__builtin_unique_stable_name.id_marray2 = private unnamed_addr constant [39 x i8] c"_ZTS14name_generatorIL_Z10id_marray2EE\00", align 1 @__builtin_unique_stable_name.id_marray3 = private unnamed_addr constant [39 x i8] c"_ZTS14name_generatorIL_Z10id_marray3EE\00", align 1 +@__builtin_unique_stable_name.id_marray4 = private unnamed_addr constant [39 x i8] c"_ZTS14name_generatorIL_Z10id_marray4EE\00", align 1 ; CHECK-LABEL: define dso_local void @_Z4testv define dso_local void @_Z4testv() local_unnamed_addr #0 { @@ -140,6 +144,7 @@ define void @test3() { %tmp1 = alloca %struct.MArrayConst, align 8 %tmp2 = alloca %struct.MArrayConst2, align 8 %tmp3 = alloca %struct.MArrayConst3, align 8 + %tmp4 = alloca %struct.MArrayConst4, align 8 %1 = bitcast %struct.VectorConst* %tmp to i8* ; CHECK-DEF: %[[GEP1:[0-9a-z]+]] = getelementptr i8, i8* null, i32 54 ; CHECK-DEF: %[[BITCAST1:[0-9a-z]+]] = bitcast i8* %[[GEP1]] to %struct.VectorConst* @@ -167,6 +172,9 @@ define void @test3() { %4 = bitcast %struct.MArrayConst3* %tmp3 to i8* ; CHECK-RT: call %struct.MArrayConst3 @_Z29__spirv_SpecConstantCompositeA3_x_Rstruct.MArrayConst3 call void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst3ET_PKcPvS4_(%struct.MArrayConst3* nonnull sret(%struct.MArrayConst3) align 8 %tmp3, i8* getelementptr inbounds ([39 x i8], [39 x i8]* @__builtin_unique_stable_name.id_marray3, i64 0, i64 0), i8* bitcast (%class.specialization_id.6* @id_marray3 to i8*), i8* null) + %5 = bitcast %struct.MArrayConst4* %tmp4 to i8* +; CHECK-RT: call %struct.MArrayConst4 @_Z29__spirv_SpecConstantCompositeA2_A2_A3_i_Rstruct.MArrayConst4 + call void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst4ET_PKcPvS4_(%struct.MArrayConst4* nonnull sret(%struct.MArrayConst4) align 8 %tmp4, i8* getelementptr inbounds ([39 x i8], [39 x i8]* @__builtin_unique_stable_name.id_marray4, i64 0, i64 0), i8* bitcast (%class.specialization_id.7* @id_marray4 to i8*), i8* null) ret void } @@ -192,6 +200,8 @@ declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayCon declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst3ET_PKcPvS4_(%struct.MArrayConst3* sret(%struct.MArrayConst3) align 8, i8*, i8*, i8*) local_unnamed_addr #2 +declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11MArrayConst4ET_PKcPvS4_(%struct.MArrayConst4* sret(%struct.MArrayConst4) align 8, i8*, i8*, i8*) local_unnamed_addr #2 + attributes #0 = { uwtable mustprogress "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-jump-tables"="false" "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 #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" }