Skip to content

Commit 2b78412

Browse files
vmaksimovmaksimo
committed
Update tests and apply suggestions
1 parent bda11f6 commit 2b78412

File tree

3 files changed

+32
-86
lines changed

3 files changed

+32
-86
lines changed

‎llvm/test/tools/sycl-post-link/assert-property-with-split.ll

Lines changed: 13 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,7 @@
1+
; This test checks that the post-link tool properly generates "assert used"
2+
; property in split mode - it should include only kernels that call assertions
3+
; in their call graph.
4+
15
; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
26
; RUN: FileCheck %s -input-file=%t_0.prop
37

@@ -13,11 +17,8 @@ target triple = "spir64-unknown-linux-sycldevice"
1317
@_ZL10assert_fmt = internal addrspace(2) constant [85 x i8] c"%s:%d: %s: global id: [%lu,%lu,%lu], local id: [%lu,%lu,%lu] Assertion `%s` failed.\0A\00", align 1
1418

1519
; CHECK: [SYCL/assert used]
16-
; CHECK-NOT: _ZTSZ4mainE11TU0_kernel1
17-
; CHECK-DAG: _ZTSZ4mainE11TU0_kernel0
18-
; CHECK-DAG: _ZTSZ4mainE10TU1_kernel
19-
; CHECK-NOT: _ZTSZ4mainE11TU0_kernel1
2020

21+
; CHECK: _ZTSZ4mainE11TU0_kernel0
2122
define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 {
2223
entry:
2324
call spir_func void @_Z3foov()
@@ -35,6 +36,7 @@ entry:
3536
ret void
3637
}
3738

39+
; CHECK-NOT: _ZTSZ4mainE11TU0_kernel1
3840
define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel1() #0 {
3941
entry:
4042
call spir_func void @_Z4foo1v()
@@ -49,6 +51,7 @@ entry:
4951
ret void
5052
}
5153

54+
; CHECK: _ZTSZ4mainE10TU1_kernel
5255
define dso_local spir_kernel void @_ZTSZ4mainE10TU1_kernel() #1 {
5356
entry:
5457
call spir_func void @_Z4foo2v()
@@ -80,52 +83,22 @@ entry:
8083
}
8184

8285
; Function Attrs: inlinehint norecurse mustprogress
83-
define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr {
84-
entry:
85-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
86-
%1 = extractelement <3 x i64> %0, i64 0
87-
ret i64 %1
88-
}
86+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr
8987

9088
; Function Attrs: inlinehint norecurse mustprogress
91-
define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr {
92-
entry:
93-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
94-
%1 = extractelement <3 x i64> %0, i64 1
95-
ret i64 %1
96-
}
89+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr
9790

9891
; Function Attrs: inlinehint norecurse mustprogress
99-
define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr {
100-
entry:
101-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
102-
%1 = extractelement <3 x i64> %0, i64 2
103-
ret i64 %1
104-
}
92+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr
10593

10694
; Function Attrs: inlinehint norecurse mustprogress
107-
define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr {
108-
entry:
109-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInLocalInvocationId to <3 x i64> addrspace(4)*), align 32
110-
%1 = extractelement <3 x i64> %0, i64 0
111-
ret i64 %1
112-
}
95+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr
11396

11497
; Function Attrs: inlinehint norecurse mustprogress
115-
define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr {
116-
entry:
117-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInLocalInvocationId to <3 x i64> addrspace(4)*), align 32
118-
%1 = extractelement <3 x i64> %0, i64 1
119-
ret i64 %1
120-
}
98+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr
12199

122100
; Function Attrs: inlinehint norecurse mustprogress
123-
define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr {
124-
entry:
125-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInLocalInvocationId to <3 x i64> addrspace(4)*), align 32
126-
%1 = extractelement <3 x i64> %0, i64 2
127-
ret i64 %1
128-
}
101+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr
129102

130103
; Function Attrs: convergent norecurse mustprogress
131104
define weak dso_local spir_func void @__devicelib_assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func, i64 %gid0, i64 %gid1, i64 %gid2, i64 %lid0, i64 %lid1, i64 %lid2) local_unnamed_addr {

‎llvm/test/tools/sycl-post-link/assert-property.ll

Lines changed: 15 additions & 42 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,6 @@
1+
; This test checks that the post-link tool properly generates "assert used"
2+
; property - it should include only kernels that call assertions in their call
3+
; graph.
14

25
; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
36
; RUN: FileCheck %s -input-file=%t_0.prop
@@ -46,13 +49,8 @@ target triple = "spir64_x86_64-unknown-unknown-sycldevice"
4649
@_ZL10assert_fmt = internal addrspace(2) constant [85 x i8] c"%s:%d: %s: global id: [%lu,%lu,%lu], local id: [%lu,%lu,%lu] Assertion `%s` failed.\0A\00", align 1
4750

4851
; CHECK: [SYCL/assert used]
49-
; CHECK-DAG: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE9TeKernelEE
50-
; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel
51-
; CHECK-DAG: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel3EE
52-
; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3
53-
; CHECK-NOT: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel2EE
54-
; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2
5552

53+
; CHECK: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE9TeKernelEE
5654
; Function Attrs: convergent norecurse
5755
define weak_odr dso_local spir_kernel void @"_ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE9TeKernelEE"(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range") align 8 %_arg_, %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon"* byval(%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon") align 1 %_arg_1) #0 {
5856
entry:
@@ -84,13 +82,15 @@ entry:
8482
ret void
8583
}
8684

85+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel
8786
; Function Attrs: convergent norecurse
8887
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel"() #0 {
8988
entry:
9089
tail call spir_func void @__assert_fail(i8 addrspace(4)* getelementptr inbounds ([2 x i8], [2 x i8] addrspace(4)* addrspacecast ([2 x i8] addrspace(1)* @.str to [2 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* getelementptr inbounds ([11 x i8], [11 x i8] addrspace(4)* addrspacecast ([11 x i8] addrspace(1)* @.str.1 to [11 x i8] addrspace(4)*), i64 0, i64 0), i32 8, i8 addrspace(4)* getelementptr inbounds ([11 x i8], [11 x i8] addrspace(4)* addrspacecast ([11 x i8] addrspace(1)* @__PRETTY_FUNCTION__._Z3foov to [11 x i8] addrspace(4)*), i64 0, i64 0))
9190
ret void
9291
}
9392

93+
; CHECK-NOT: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel2EE
9494
; Function Attrs: norecurse
9595
define weak_odr dso_local spir_kernel void @"_ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel2EE"(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range") align 8 %_arg_, %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon"* byval(%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon") align 1 %_arg_1) #1 {
9696
entry:
@@ -107,12 +107,14 @@ entry:
107107
ret void
108108
}
109109

110+
; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2
110111
; Function Attrs: norecurse
111112
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2"() #1 {
112113
entry:
113114
ret void
114115
}
115116

117+
; CHECK: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel3EE
116118
; Function Attrs: convergent norecurse
117119
define weak_odr dso_local spir_kernel void @"_ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel3EE"(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range") align 8 %_arg_, %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon"* byval(%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon") align 1 %_arg_1) #0 {
118120
entry:
@@ -141,6 +143,7 @@ entry:
141143
ret void
142144
}
143145

146+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3
144147
; Function Attrs: convergent norecurse
145148
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3"() #0 {
146149
entry:
@@ -162,52 +165,22 @@ entry:
162165
}
163166

164167
; Function Attrs: inlinehint norecurse mustprogress
165-
define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() {
166-
entry:
167-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
168-
%1 = extractelement <3 x i64> %0, i64 0
169-
ret i64 %1
170-
}
168+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr
171169

172170
; Function Attrs: inlinehint norecurse mustprogress
173-
define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() {
174-
entry:
175-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
176-
%1 = extractelement <3 x i64> %0, i64 1
177-
ret i64 %1
178-
}
171+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr
179172

180173
; Function Attrs: inlinehint norecurse mustprogress
181-
define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() {
182-
entry:
183-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
184-
%1 = extractelement <3 x i64> %0, i64 2
185-
ret i64 %1
186-
}
174+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr
187175

188176
; Function Attrs: inlinehint norecurse mustprogress
189-
define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() {
190-
entry:
191-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInLocalInvocationId to <3 x i64> addrspace(4)*), align 32
192-
%1 = extractelement <3 x i64> %0, i64 0
193-
ret i64 %1
194-
}
177+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr
195178

196179
; Function Attrs: inlinehint norecurse mustprogress
197-
define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() {
198-
entry:
199-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInLocalInvocationId to <3 x i64> addrspace(4)*), align 32
200-
%1 = extractelement <3 x i64> %0, i64 1
201-
ret i64 %1
202-
}
180+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr
203181

204182
; Function Attrs: inlinehint norecurse mustprogress
205-
define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() {
206-
entry:
207-
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInLocalInvocationId to <3 x i64> addrspace(4)*), align 32
208-
%1 = extractelement <3 x i64> %0, i64 2
209-
ret i64 %1
210-
}
183+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr
211184

212185
; Function Attrs: convergent norecurse mustprogress
213186
define weak dso_local spir_func void @__devicelib_assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func, i64 %gid0, i64 %gid1, i64 %gid2, i64 %lid0, i64 %lid1, i64 %lid2) {

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -169,7 +169,6 @@ struct ImagePropSaveInfo {
169169
bool SpecConstsMet;
170170
bool EmitKernelParamInfo;
171171
bool IsEsimdKernel;
172-
bool IsAssertEnabled;
173172
};
174173

175174
static void error(const Twine &Msg) {
@@ -485,10 +484,12 @@ static string_vector saveDeviceImageProperty(
485484
{"isEsimdImage", true});
486485
}
487486

488-
if (ImgPSInfo.IsAssertEnabled) {
487+
{
489488
Module *M = ResultModules[I].get();
490489
std::vector<Function *> SyclKernels;
491490
for (auto &F : M->functions()) {
491+
// TODO: handle SYCL_EXTERNAL functions for dynamic linkage.
492+
// TODO: handle function pointers.
492493
if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {
493494
if (hasAssertInFunctionCallGraph(&F)) {
494495
SyclKernels.push_back(&F);
@@ -645,8 +646,7 @@ static TableFiles processOneModule(std::unique_ptr<Module> M, bool IsEsimd,
645646
{
646647
ImagePropSaveInfo ImgPSInfo = {
647648
true, DoSpecConst, SetSpecConstAtRT,
648-
SpecConstsMet, EmitKernelParamInfo, IsEsimd,
649-
true};
649+
SpecConstsMet, EmitKernelParamInfo, IsEsimd};
650650
string_vector Files = saveDeviceImageProperty(ResultModules, ImgPSInfo);
651651
std::copy(Files.begin(), Files.end(),
652652
std::back_inserter(TblFiles[COL_PROPS]));

0 commit comments

Comments
 (0)