Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 19 commits into from
Apr 29, 2021
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
19 commits
Select commit Hold shift + click to select a range
7ab8773
Emit metadata for non-native spec constants
AlexeySachkov Apr 16, 2021
cc33b86
Refactor metadata generation for native spec constants
AlexeySachkov Apr 16, 2021
9518a8c
Refactor metadata collection in SpecConstantsPass
AlexeySachkov Apr 16, 2021
9353371
Cleanup dead code
AlexeySachkov Apr 16, 2021
f8d49b3
Outline common code
AlexeySachkov Apr 16, 2021
4ab1409
Adjust existing tests to new metadata
AlexeySachkov Apr 16, 2021
70a7900
Small refactoring/cleanup
AlexeySachkov Apr 16, 2021
af1f93b
Fix missing vector initialization
AlexeySachkov Apr 16, 2021
9e5fdd3
Refactor SpecConstants.h header
AlexeySachkov Apr 16, 2021
682ec2f
Refactoring
AlexeySachkov Apr 16, 2021
70b51b6
Restore necessary cleanup
AlexeySachkov Apr 19, 2021
dcf1567
Fix out-of-bounds access in SpecConstantsPass
AlexeySachkov Apr 19, 2021
4f48b28
Add test for new changes
AlexeySachkov Apr 19, 2021
7358c05
Emit spec constants property set even in emulation mode
AlexeySachkov Apr 21, 2021
f3e322e
Merge remote-tracking branch 'sycl' into private/asachkov/spec-consta…
AlexeySachkov Apr 21, 2021
c1f0fd8
Fix one more out-of-bounds access to vector
AlexeySachkov Apr 21, 2021
5b5602f
Add integration test for new device image properties
AlexeySachkov Apr 21, 2021
e7c93de
Apply comments to test
AlexeySachkov Apr 27, 2021
92649af
Use MapVector for storing spec constant metadata
AlexeySachkov Apr 27, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
34 changes: 24 additions & 10 deletions llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
}

Expand All @@ -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*

Expand All @@ -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*

Expand All @@ -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
Expand All @@ -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}
9 changes: 7 additions & 2 deletions llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down
27 changes: 20 additions & 7 deletions llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]])
Expand All @@ -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"

Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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}
Expand Down
9 changes: 6 additions & 3 deletions llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down
24 changes: 17 additions & 7 deletions llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down
11 changes: 6 additions & 5 deletions llvm/test/tools/sycl-post-link/spec-constants/multiple-scalar-usages.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
6 changes: 4 additions & 2 deletions llvm/test/tools/sycl-post-link/spec-constants/scalar-O0.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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}
Loading