Skip to content

Commit 09fb54a

Browse files
AlexeySachkovAlexeySachkov
committed
[SYCL] Add support for composites in SpecConstantsPass
1 parent 1c1c4c3 commit 09fb54a

File tree

6 files changed

+486
-50
lines changed

6 files changed

+486
-50
lines changed

‎llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
; RUN: sycl-post-link -spec-const=default --ir-output-only %s -S -o - \
2+
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
3+
;
4+
; TODO: consider adding a test case with vector type: the pass itself already
5+
; supports this, but at the moment, sycl::vec type is not a POD type, which
6+
; means we can't have it within a spec constant, i.e. we can't generate LLVM IR
7+
; from a real-life application to use it as a test here.
8+
;
9+
; CHECK: %[[#CAST:]] = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i
10+
; CHECK: store %struct._ZTS3POD.POD zeroinitializer, %struct._ZTS3POD.POD {{.*}}* %[[#CAST]]
11+
12+
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"
13+
target triple = "spir64-unknown-unknown-sycldevice"
14+
15+
%struct._ZTS3POD.POD = type { [2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" }
16+
%struct._ZTS1A.A = type { i32, float }
17+
%"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" = type { <2 x i32> }
18+
%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
19+
%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
20+
%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
21+
22+
$_ZTS4Test = comdat any
23+
24+
@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3POD\00", align 1
25+
26+
; Function Attrs: convergent norecurse uwtable
27+
define weak_odr dso_local spir_kernel void @_ZTS4Test(%struct._ZTS3POD.POD addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
28+
entry:
29+
%ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8
30+
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
31+
%1 = load i64, i64* %0, align 8
32+
%add.ptr.i = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %_arg_, i64 %1
33+
%2 = bitcast %struct._ZTS3POD.POD* %ref.tmp.i to i8*
34+
call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %2) #3
35+
%3 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)*
36+
call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8 %3, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4
37+
%4 = bitcast %struct._ZTS3POD.POD addrspace(1)* %add.ptr.i to i8 addrspace(1)*
38+
%5 = addrspacecast i8 addrspace(1)* %4 to i8 addrspace(4)*
39+
call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 dereferenceable(24) %5, i8* nonnull align 8 dereferenceable(24) %2, i64 24, i1 false), !tbaa.struct !5
40+
call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2) #3
41+
ret void
42+
}
43+
44+
; Function Attrs: argmemonly nounwind willreturn
45+
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
46+
47+
; Function Attrs: argmemonly nounwind willreturn
48+
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
49+
50+
; Function Attrs: argmemonly nounwind willreturn
51+
declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #1
52+
53+
; Function Attrs: convergent
54+
declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8, i8 addrspace(4)*) local_unnamed_addr #2
55+
56+
attributes #0 = { convergent norecurse uwtable "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="../sycl/test/spec_const/composite.cpp" "tune-cpu"="generic" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
57+
attributes #1 = { argmemonly nounwind willreturn }
58+
attributes #2 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "tune-cpu"="generic" "unsafe-fp-math"="false" "use-soft-float"="false" }
59+
attributes #3 = { nounwind }
60+
attributes #4 = { convergent }
61+
62+
!llvm.module.flags = !{!0}
63+
!opencl.spir.version = !{!1}
64+
!spirv.Source = !{!2}
65+
!llvm.ident = !{!3}
66+
67+
!0 = !{i32 1, !"wchar_size", i32 4}
68+
!1 = !{i32 1, i32 2}
69+
!2 = !{i32 4, i32 100000}
70+
!3 = !{!"clang version 12.0.0 (/data/github.com/intel/llvm/clang 56ee5b054b5a1f2f703722fc414fcb05af18b40a)"}
71+
!4 = !{i32 -1, i32 -1, i32 -1, i32 -1}
72+
!5 = !{i64 0, i64 16, !6, i64 16, i64 8, !6}
73+
!6 = !{!7, !7, i64 0}
74+
!7 = !{!"omnipotent char", !8, i64 0}
75+
!8 = !{!"Simple C++ TBAA"}

‎llvm/test/tools/sycl-post-link/composite-spec-constant.ll

Lines changed: 91 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,91 @@
1+
; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
2+
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
3+
;
4+
; This test is intended to check that sycl-post-link tool is capable of handling
5+
; composite specialization constants by lowering them into a set of SPIR-V
6+
; friendly IR operations representing those constants.
7+
;
8+
; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32
9+
; CHECK: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 1]], float
10+
; CHECK: %[[#NA0:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS0]], float %[[#NS1]])
11+
;
12+
; CHECK: %[[#NS2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 3]], i32
13+
; CHECK: %[[#NS3:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 4]], float
14+
; CHECK: %[[#NA1:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS2]], float %[[#NS3]])
15+
;
16+
; CHECK: %[[#NA:]] = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A %[[#NA0]], %struct._ZTS1A.A %[[#NA1]])
17+
;
18+
; CHECK: %[[#B0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 7]], i32{{.*}})
19+
; CHECK: %[[#B1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 8]], i32{{.*}})
20+
; CHECK: %[[#BV:]] = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32 %[[#B0]], i32 %[[#B1]])
21+
; CHECK: %[[#B:]] = call %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" @_Z29__spirv_SpecConstantCompositeDv2_i(<2 x i32> %[[#BV]])
22+
;
23+
; 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:]]
24+
; CHECK: store %struct._ZTS3POD.POD %[[#POD]]
25+
;
26+
; CHECK: ![[#MD]] = !{!"_ZTS3POD", i32 [[#ID]], i32 [[#ID + 1]], i32 [[#ID + 3]], i32 [[#ID + 4]], i32 [[#ID + 7]], i32 [[#ID + 8]]}
27+
28+
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"
29+
target triple = "spir64-unknown-unknown-sycldevice"
30+
31+
%struct._ZTS3POD.POD = type { [2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" }
32+
%struct._ZTS1A.A = type { i32, float }
33+
%"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" = type { <2 x i32> }
34+
%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
35+
%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
36+
%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
37+
38+
$_ZTS4Test = comdat any
39+
40+
@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3POD\00", align 1
41+
42+
; Function Attrs: convergent norecurse uwtable
43+
define weak_odr dso_local spir_kernel void @_ZTS4Test(%struct._ZTS3POD.POD addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
44+
entry:
45+
%ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8
46+
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
47+
%1 = load i64, i64* %0, align 8
48+
%add.ptr.i = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %_arg_, i64 %1
49+
%2 = bitcast %struct._ZTS3POD.POD* %ref.tmp.i to i8*
50+
call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %2) #3
51+
%3 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)*
52+
call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8 %3, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4
53+
%4 = bitcast %struct._ZTS3POD.POD addrspace(1)* %add.ptr.i to i8 addrspace(1)*
54+
%5 = addrspacecast i8 addrspace(1)* %4 to i8 addrspace(4)*
55+
call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 dereferenceable(24) %5, i8* nonnull align 8 dereferenceable(24) %2, i64 24, i1 false), !tbaa.struct !5
56+
call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2) #3
57+
ret void
58+
}
59+
60+
; Function Attrs: argmemonly nounwind willreturn
61+
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
62+
63+
; Function Attrs: argmemonly nounwind willreturn
64+
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
65+
66+
; Function Attrs: argmemonly nounwind willreturn
67+
declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #1
68+
69+
; Function Attrs: convergent
70+
declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8, i8 addrspace(4)*) local_unnamed_addr #2
71+
72+
attributes #0 = { convergent norecurse uwtable "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="../sycl/test/spec_const/composite.cpp" "tune-cpu"="generic" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
73+
attributes #1 = { argmemonly nounwind willreturn }
74+
attributes #2 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "tune-cpu"="generic" "unsafe-fp-math"="false" "use-soft-float"="false" }
75+
attributes #3 = { nounwind }
76+
attributes #4 = { convergent }
77+
78+
!llvm.module.flags = !{!0}
79+
!opencl.spir.version = !{!1}
80+
!spirv.Source = !{!2}
81+
!llvm.ident = !{!3}
82+
83+
!0 = !{i32 1, !"wchar_size", i32 4}
84+
!1 = !{i32 1, i32 2}
85+
!2 = !{i32 4, i32 100000}
86+
!3 = !{!"clang version 12.0.0 (/data/github.com/intel/llvm/clang 56ee5b054b5a1f2f703722fc414fcb05af18b40a)"}
87+
!4 = !{i32 -1, i32 -1, i32 -1, i32 -1}
88+
!5 = !{i64 0, i64 16, !6, i64 16, i64 8, !6}
89+
!6 = !{!7, !7, i64 0}
90+
!7 = !{!"omnipotent char", !8, i64 0}
91+
!8 = !{!"Simple C++ TBAA"}

‎llvm/test/tools/sycl-post-link/multiple-spec-const-usages.ll

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
; This test checks that IDs assigned to spec constants are correct, i.e. if some
2+
; spec constant is accessed twice, then metadata for both accesses should point
3+
; to the same ID
4+
5+
; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
6+
; RUN: | FileCheck %s
7+
8+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
9+
target triple = "spir64-unknown-unknown-sycldevice"
10+
11+
%"spec_constant" = type { i8 }
12+
13+
@SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1
14+
@SCSymID1 = private unnamed_addr constant [11 x i8] c"SpecConst1\00", align 1
15+
16+
declare dso_local spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)*)
17+
18+
; Function Attrs: norecurse
19+
define weak_odr dso_local spir_kernel void @Kernel() {
20+
%1 = call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
21+
; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID0:[0-9]+]]
22+
ret void
23+
}
24+
25+
; Function Attrs: norecurse
26+
define dso_local spir_func float @foo_float(%"spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
27+
%2 = tail call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([11 x i8], [11 x i8]* @SCSymID1, i64 0, i64 0) to i8 addrspace(4)*))
28+
; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID1:[0-9]+]]
29+
ret float %2
30+
}
31+
32+
; Function Attrs: norecurse
33+
define dso_local spir_func float @foo_float2(%"spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
34+
%2 = tail call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
35+
; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID0]]
36+
ret float %2
37+
}
38+
39+
; CHECK: ![[ID0]] = !{!"SpecConst", i32 0}
40+
; CHECK: ![[ID1]] = !{!"SpecConst1", i32 1}

0 commit comments

Comments
 (0)