Fix handling of composites in sycl-post-link (intel#4530)

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.
That is achieved by including a mangled return type substring into
the resulting mangling name as `_R{ret-type-mangling}` suffix.

Author: AlexeySachkov

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

@@ -15,10 +15,26 @@
 %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 }
+%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 }
+%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
 @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> <i32 1, i32 2>  } }, 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
+@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
@@ -33,6 +49,11 @@
 @__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
+@__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 {
@@ -77,8 +98,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 +114,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 +131,53 @@ 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
+  %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*
+; 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)
+; 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)
+  %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
+}
+
 declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
 
 declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
@@ -131,14 +192,24 @@ 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
+
+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
+
+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" }
 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 +233,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> <i32 1, i32 2> }}
+; 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}

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:]]}

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:]]}