Skip to content

Commit a7ef268

Browse files
maksimsabmaksimsab
authored
[SYCL] Add generation of device image with specialization constants replaced by default values (#10115)
1 parent 970a2df commit a7ef268

File tree

11 files changed

+472
-20
lines changed

11 files changed

+472
-20
lines changed

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/bool.ll

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
; Test checks handling of bool specialization constant.
2+
3+
; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts
4+
; RUN: FileCheck %s -input-file %t_1.ll --implicit-check-not="SpecConst"
5+
6+
; CHECK: %bool1 = trunc i8 1 to i1
7+
; CHECK: %frombool = zext i1 %bool1 to i8
8+
9+
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"
10+
target triple = "spir64-unknown-unknown"
11+
12+
%"class.sycl::_V1::specialization_id" = type { i8 }
13+
%struct.A = type { i8 }
14+
15+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
16+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { i8 1 }, align 4
17+
18+
; Function Attrs: convergent nounwind
19+
declare dso_local spir_func noundef zeroext i1 @_Z37__sycl_getScalar2020SpecConstantValueIbET_PKcPKvS4_(i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
20+
21+
define spir_kernel void @kernel() {
22+
entry:
23+
%bool = call spir_func noundef zeroext i1 @_Z37__sycl_getScalar2020SpecConstantValueIbET_PKcPKvS4_(i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds (%"class.sycl::_V1::specialization_id", %"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c, i64 0, i32 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef null) #4
24+
%frombool = zext i1 %bool to i8
25+
ret void
26+
}

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll

Lines changed: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,90 @@
1+
; Test checks the content of simple generated device image.
2+
; It checks scalar, sret and "return by value" versions of SpecConstant functions.
3+
; Also test checks generated symbols.
4+
5+
; RUN: sycl-post-link -split=auto -spec-const=native -symbols -S -o %t.table %s -generate-device-image-default-spec-consts
6+
; RUN: FileCheck %s -input-file %t.table -check-prefix=CHECK-TABLE
7+
; RUN: FileCheck %s -input-file %t_0.prop -check-prefix=CHECK-PROP0
8+
; RUN: FileCheck %s -input-file %t_1.prop -check-prefix=CHECK-PROP1
9+
; RUN: FileCheck %s -input-file %t_0.ll -check-prefix=CHECK-IR0
10+
; RUN: FileCheck %s -input-file %t_1.ll -check-prefix=CHECK-IR1 --implicit-check-not "SpecConstant"
11+
; RUN: FileCheck %s -input-file %t_0.sym -check-prefix=CHECK-SYM0
12+
; RUN: FileCheck %s -input-file %t_1.sym -check-prefix=CHECK-SYM1
13+
14+
; CHECK-TABLE: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
15+
; CHECK-TABLE: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
16+
17+
; CHECK-PROP0-NOT: specConstsReplacedWithDefault=1|1
18+
19+
; CHECK-PROP1: specConstsReplacedWithDefault=1|1
20+
21+
; CHECK-IR0: call i32 @_Z20__spirv_SpecConstantii
22+
; CHECK-IR0: call %struct.B @_Z29__spirv_SpecConstantCompositeiii_Rstruct.B
23+
; CHECK-IR0: call %struct.A @_Z29__spirv_SpecConstantCompositeistruct.B_Rstruct.A
24+
25+
; CHECK-IR1: store %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } }, %struct.A addrspace(4)* %a.ascast.i, align 4
26+
27+
; Check that %scalar value has been replaced by global value.
28+
; CHECK-IR1-NOT: %scalar = call
29+
; CHECK-IR1: %scalar2 = add i32 123, 1
30+
31+
; Check that %returned_spec_const value has been replaced by global value.
32+
; CHECK-IR1-NOT: %returned_spec_const = call
33+
; CHECK-IR1: %sc.e = extractvalue %struct.C { i32 1 }, 0
34+
35+
; CHECK-SYM0: kernel
36+
; CHECK-SYM0-EMPTY:
37+
38+
; CHECK-SYM1: kernel
39+
; CHECK-SYM1-EMPTY:
40+
41+
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"
42+
target triple = "spir64-unknown-unknown"
43+
44+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
45+
%"class.sycl::_V1::specialization_id.2" = type { i32 }
46+
%"class.sycl::_V1::specialization_id.3" = type { %struct.C }
47+
%struct.A = type { i32, %struct.B }
48+
%struct.B = type { i32, i32, i32 }
49+
%struct.C = type { i32 }
50+
51+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
52+
@__usid_str.1 = private unnamed_addr constant [33 x i8] c"uidcac21ed8fab7d507____ZL6valueS\00", align 1
53+
@__usid_str.2 = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8f____ZL1b\00", align 1
54+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
55+
@_ZL6valueS = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.2" { i32 123 }, align 4
56+
@_ZL1b = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.3" { %struct.C { i32 1 } }, align 4
57+
58+
; This constant checks `zeroinitializer` field.
59+
@_ZL1d = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B zeroinitializer } }, align 4
60+
61+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
62+
63+
declare spir_func %struct.C @_Z40__sycl_getComposite2020SpecConstantValueI1CET_PKcPKvS5_(i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
64+
65+
declare dso_local spir_func noundef i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
66+
67+
; Function for testing symbol generation
68+
define spir_func void @func() {
69+
ret void
70+
}
71+
72+
define spir_kernel void @kernel() {
73+
entry:
74+
%a.i = alloca %struct.A, align 4
75+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
76+
%0 = bitcast %struct.A* %a.i to i8*
77+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
78+
%scalar = call spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([33 x i8], [33 x i8]* @__usid_str.1, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id.2" addrspace(1)* @_ZL6valueS to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
79+
%scalar2 = add i32 %scalar, 1
80+
81+
%returned_spec_const = call spir_func %struct.C @_Z40__sycl_getComposite2020SpecConstantValueI1CET_PKcPKvS5_(i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str.2, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id.3" addrspace(1)* @_ZL1b to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
82+
%sc.e = extractvalue %struct.C %returned_spec_const, 0
83+
%scalar3 = add i32 %sc.e, 1
84+
85+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1d to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
86+
87+
88+
call void @func()
89+
ret void
90+
}

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/esimd.ll

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
; Test checks generation of device image of esimd kernel.
2+
3+
; RUN: sycl-post-link -split=auto -split-esimd -lower-esimd -O2 -spec-const=native -o %t.table %s -generate-device-image-default-spec-consts
4+
; RUN: FileCheck %s -input-file=%t.table -check-prefix=CHECK-TABLE
5+
; RUN: FileCheck %s -input-file=%t_1.prop -check-prefix=CHECK-PROP
6+
; RUN: FileCheck %s -input-file=%t_esimd_1.prop -check-prefix=CHECK-ESIMD-PROP
7+
8+
; CHECK-TABLE: {{.*}}_esimd_0.bc|{{.*}}_esimd_0.prop
9+
; CHECK-TABLE: {{.*}}_0.bc|{{.*}}_0.prop
10+
; CHECK-TABLE: {{.*}}_esimd_1.bc|{{.*}}_esimd_1.prop
11+
; CHECK-TABLE: {{.*}}_1.bc|{{.*}}_1.prop
12+
13+
; CHECK-PROP: specConstsReplacedWithDefault=1|1
14+
15+
; CHECK-ESIMD-PROP: isEsimdImage=1|1
16+
; CHECK-ESIMD-PROP: specConstsReplacedWithDefault=1|1
17+
18+
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"
19+
target triple = "spir64-unknown-unknown"
20+
21+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
22+
%struct.A = type { i32, %struct.B }
23+
%struct.B = type { i32, i32, i32 }
24+
25+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
26+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
27+
28+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
29+
30+
define spir_kernel void @func1() !kernel_arg_buffer_location !7 !sycl_kernel_omit_args !8 {
31+
entry:
32+
%a.i = alloca %struct.A, align 4
33+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
34+
%0 = bitcast %struct.A* %a.i to i8*
35+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
36+
ret void
37+
}
38+
39+
define spir_kernel void @func2(i8 addrspace(1)* noundef align 1 %_arg__specialization_constants_buffer) !sycl_explicit_simd !1 !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 !kernel_arg_accessor_ptr !6 {
40+
entry:
41+
%a.i = alloca %struct.A, align 4
42+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
43+
%0 = bitcast %struct.A* %a.i to i8*
44+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
45+
ret void
46+
}
47+
48+
!1 = !{}
49+
!2 = !{i32 1}
50+
!3 = !{!"none"}
51+
!4 = !{!"char*"}
52+
!5 = !{!""}
53+
!6 = !{i1 false}
54+
!7 = !{i32 -1}
55+
!8 = !{i1 true}

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-kernel.ll

Lines changed: 51 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,51 @@
1+
; Test checks generation of device images for splitted kernels.
2+
3+
; RUN: sycl-post-link -split=kernel -o %t.table %s -generate-device-image-default-spec-consts
4+
; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE
5+
; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0
6+
; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
7+
; RUN: cat %t_2.prop | FileCheck %s -check-prefix=CHECK-PROP2
8+
; RUN: cat %t_3.prop | FileCheck %s -check-prefix=CHECK-PROP3
9+
10+
; CHECK-TABLE: {{.*}}_0.bc|{{.*}}_0.prop
11+
; CHECK-TABLE: {{.*}}_1.bc|{{.*}}_1.prop
12+
; CHECK-TABLE: {{.*}}_2.bc|{{.*}}_2.prop
13+
; CHECK-TABLE: {{.*}}_3.bc|{{.*}}_3.prop
14+
15+
; CHECK-PROP0-NOT: specConstsReplacedWithDefault=1|1
16+
17+
; CHECK-PROP1: specConstsReplacedWithDefault=1|1
18+
19+
; CHECK-PROP2-NOT: specConstsReplacedWithDefault=1|1
20+
21+
; CHECK-PROP3: specConstsReplacedWithDefault=1|1
22+
23+
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"
24+
target triple = "spir64-unknown-unknown"
25+
26+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
27+
%struct.A = type { i32, %struct.B }
28+
%struct.B = type { i32, i32, i32 }
29+
30+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
31+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
32+
33+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
34+
35+
define spir_kernel void @kernel1() {
36+
entry:
37+
%a.i = alloca %struct.A, align 4
38+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
39+
%0 = bitcast %struct.A* %a.i to i8*
40+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
41+
ret void
42+
}
43+
44+
define spir_kernel void @kernel2() {
45+
entry:
46+
%a.i = alloca %struct.A, align 4
47+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
48+
%0 = bitcast %struct.A* %a.i to i8*
49+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
50+
ret void
51+
}

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-source.ll

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
; Test checks generation of device images for splitted kernels by source.
2+
3+
; RUN: sycl-post-link -split=source -o %t.table %s -generate-device-image-default-spec-consts
4+
; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE
5+
; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0
6+
; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
7+
; RUN: cat %t_2.prop | FileCheck %s -check-prefix=CHECK-PROP2
8+
; RUN: cat %t_3.prop | FileCheck %s -check-prefix=CHECK-PROP3
9+
10+
; CHECK-TABLE: {{.*}}_0.bc|{{.*}}_0.prop
11+
; CHECK-TABLE: {{.*}}_1.bc|{{.*}}_1.prop
12+
; CHECK-TABLE: {{.*}}_2.bc|{{.*}}_2.prop
13+
; CHECK-TABLE: {{.*}}_3.bc|{{.*}}_3.prop
14+
15+
; CHECK-PROP0-NOT: specConstsReplacedWithDefault=1|1
16+
17+
; CHECK-PROP1: specConstsReplacedWithDefault=1|1
18+
19+
; CHECK-PROP2-NOT: specConstsReplacedWithDefault=1|1
20+
21+
; CHECK-PROP3: specConstsReplacedWithDefault=1|1
22+
23+
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"
24+
target triple = "spir64-unknown-unknown"
25+
26+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
27+
%struct.A = type { i32, %struct.B }
28+
%struct.B = type { i32, i32, i32 }
29+
30+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
31+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
32+
33+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
34+
35+
define spir_kernel void @kernel1() #0 {
36+
entry:
37+
%a.i = alloca %struct.A, align 4
38+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
39+
%0 = bitcast %struct.A* %a.i to i8*
40+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
41+
ret void
42+
}
43+
44+
define spir_kernel void @kernel2() #1 {
45+
entry:
46+
%a.i = alloca %struct.A, align 4
47+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
48+
%0 = bitcast %struct.A* %a.i to i8*
49+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
50+
ret void
51+
}
52+
53+
attributes #0 = { "sycl-module-id"="TU1.cpp" }
54+
attributes #1 = { "sycl-module-id"="TU2.cpp" }

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/struct-with-padding.ll

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
; Test checks that struct with padding is handled correctly.
2+
3+
; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts
4+
; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE
5+
; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
6+
; RUN: cat %t_1.ll | FileCheck %s -check-prefix=CHECK-IR1 --implicit-check-not SpecConstant
7+
8+
; CHECK-TABLE: {{.*}}_0.ll|{{.*}}_0.prop
9+
; CHECK-TABLE: {{.*}}_1.ll|{{.*}}_1.prop
10+
11+
; CHECK-PROP1: specConstsReplacedWithDefault=1|1
12+
13+
; CHECK-IR1: store { float, i32, i8 } { float 0x40091EB860000000, i32 42, i8 8 }, { float, i32, i8 } addrspace(4)* %1, align 4
14+
15+
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"
16+
target triple = "spir64-unknown-unknown"
17+
18+
%struct.A = type <{ float, i32, i8, [3 x i8] }>
19+
20+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
21+
@_ZL1c = internal addrspace(1) constant { { float, i32, i8 } } { { float, i32, i8 } { float 0x40091EB860000000, i32 42, i8 8 } }, align 4
22+
23+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
24+
25+
define spir_kernel void @func1() {
26+
entry:
27+
%a.i = alloca %struct.A, align 4
28+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
29+
%0 = bitcast %struct.A* %a.i to i8*
30+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast ({ { float, i32, i8 } } addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
31+
ret void
32+
}

‎llvm/tools/sycl-post-link/ModuleSplitter.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -563,6 +563,21 @@ void ModuleDesc::cleanup() {
563563
MPM.run(*M, MAM);
564564
}
565565

566+
bool ModuleDesc::isSpecConstantDefault() const {
567+
return Props.IsSpecConstantDefault;
568+
}
569+
570+
void ModuleDesc::setSpecConstantDefault(bool Value) {
571+
Props.IsSpecConstantDefault = Value;
572+
}
573+
574+
ModuleDesc ModuleDesc::clone() const {
575+
std::unique_ptr<Module> NewModule = CloneModule(getModule());
576+
ModuleDesc NewMD(std::move(NewModule));
577+
NewMD.EntryPoints.Props = EntryPoints.Props;
578+
return NewMD;
579+
}
580+
566581
#ifndef NDEBUG
567582
void ModuleDesc::verifyESIMDProperty() const {
568583
if (EntryPoints.Props.HasESIMD == SyclEsimdSplitStatus::SYCL_AND_ESIMD) {
@@ -630,6 +645,12 @@ void EntryPointGroup::rebuildFromNames(const std::vector<std::string> &Names,
630645
});
631646
}
632647

648+
void EntryPointGroup::rebuild(const Module &M) {
649+
for (const Function &F : M.functions())
650+
if (F.getCallingConv() == CallingConv::SPIR_KERNEL)
651+
Functions.insert(const_cast<Function *>(&F));
652+
}
653+
633654
namespace {
634655
// This is a helper class, which allows to group/categorize function based on
635656
// provided rules. It is intended to be used in device code split

0 commit comments

Comments
 (0)