Skip to content

Commit fe84701

Browse files
vmaksimovmaksimo
committed
Tests update
1 parent 1930304 commit fe84701

File tree

2 files changed

+346
-57
lines changed

2 files changed

+346
-57
lines changed

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

Lines changed: 338 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,338 @@
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.
4+
5+
; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
6+
; RUN: FileCheck %s -input-file=%t_0.prop
7+
8+
; SYCL source:
9+
; void assert_func() {
10+
; assert(0);
11+
; }
12+
;
13+
; void A_excl() {}
14+
; void B_incl() { assert_func(); }
15+
;
16+
; void A_incl() { assert_func(); }
17+
; void B_excl() {}
18+
;
19+
; void C_excl() {}
20+
; void D_incl() { assert_func(); }
21+
; void common() {
22+
; C_excl();
23+
; D_incl();
24+
; }
25+
;
26+
; void C_incl() { assert_func(); }
27+
; void D_excl() {}
28+
; void common2() {
29+
; C_incl();
30+
; D_excl();
31+
; }
32+
;
33+
; void E_excl() {}
34+
; void F_incl() { assert_func(); }
35+
;
36+
; void I_incl() { assert_func(); }
37+
; void common3() { I_incl();}
38+
; void G() { common3(); }
39+
; void H() { common3(); }
40+
;
41+
; int main() {
42+
; queue Q;
43+
; Q.submit([&] (handler& CGH) {
44+
; CGH.parallel_for<class Kernel>(range<1>{1}, [=](id<1> i) {
45+
; A_excl();
46+
; B_incl();
47+
; });
48+
; CGH.parallel_for<class Kernel2>(range<1>{1}, [=](id<1> i) {
49+
; A_incl();
50+
; B_excl();
51+
; });
52+
;
53+
; CGH.parallel_for<class Kernel3>(range<1>{1}, [=](id<1> i) {
54+
; common();
55+
; });
56+
; CGH.parallel_for<class Kernel4>(range<1>{1}, [=](id<1> i) {
57+
; common2();
58+
; });
59+
;
60+
; CGH.parallel_for<class Kernel5>(range<1>{1}, [=](id<1> i) {
61+
; B_incl();
62+
; A_excl();
63+
; });
64+
;
65+
; CGH.parallel_for<class Kernel6>(range<1>{1}, [=](id<1> i) {
66+
; E_excl();
67+
; E_excl();
68+
; });
69+
; CGH.parallel_for<class Kernel7>(range<1>{1}, [=](id<1> i) {
70+
; F_incl();
71+
; F_incl();
72+
; });
73+
;
74+
; CGH.parallel_for<class Kernel8>(range<1>{1}, [=](id<1> i) {
75+
; G();
76+
; H();
77+
; });
78+
; });
79+
; Q.wait();
80+
; return 0;
81+
; }
82+
83+
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"
84+
target triple = "spir64_x86_64-unknown-unknown-sycldevice"
85+
86+
@.str = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", align 1
87+
@.str.1 = private unnamed_addr addrspace(1) constant [16 x i8] c"assert_test.cpp\00", align 1
88+
@__PRETTY_FUNCTION__._Z11assert_funcv = private unnamed_addr addrspace(1) constant [19 x i8] c"void assert_func()\00", align 1
89+
@_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
90+
91+
; CHECK: [SYCL/assert used]
92+
93+
; Function Attrs: convergent norecurse nounwind mustprogress
94+
define dso_local spir_func void @_Z6B_inclv() local_unnamed_addr {
95+
entry:
96+
call spir_func void @_Z11assert_funcv()
97+
ret void
98+
}
99+
100+
; Function Attrs: convergent norecurse nounwind mustprogress
101+
define dso_local spir_func void @_Z11assert_funcv() local_unnamed_addr {
102+
entry:
103+
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 ([16 x i8], [16 x i8] addrspace(4)* addrspacecast ([16 x i8] addrspace(1)* @.str.1 to [16 x i8] addrspace(4)*), i64 0, i64 0), i32 7, i8 addrspace(4)* getelementptr inbounds ([19 x i8], [19 x i8] addrspace(4)* addrspacecast ([19 x i8] addrspace(1)* @__PRETTY_FUNCTION__._Z11assert_funcv to [19 x i8] addrspace(4)*), i64 0, i64 0))
104+
ret void
105+
}
106+
107+
; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress
108+
define dso_local spir_func void @_Z6A_exclv() local_unnamed_addr {
109+
entry:
110+
ret void
111+
}
112+
113+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE6Kernel
114+
; Function Attrs: convergent norecurse mustprogress
115+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE6Kernel"() local_unnamed_addr #0 {
116+
entry:
117+
call spir_func void @_Z6A_exclv()
118+
call spir_func void @_Z6B_inclv()
119+
ret void
120+
}
121+
122+
; Function Attrs: convergent norecurse nounwind mustprogress
123+
define dso_local spir_func void @_Z6A_inclv() local_unnamed_addr {
124+
entry:
125+
call spir_func void @_Z11assert_funcv()
126+
ret void
127+
}
128+
129+
; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress
130+
define dso_local spir_func void @_Z6B_exclv() local_unnamed_addr {
131+
entry:
132+
ret void
133+
}
134+
135+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel2
136+
; Function Attrs: convergent norecurse mustprogress
137+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel2"() local_unnamed_addr #0 {
138+
entry:
139+
call spir_func void @_Z6A_inclv()
140+
call spir_func void @_Z6B_exclv()
141+
ret void
142+
}
143+
144+
; Function Attrs: convergent norecurse nounwind mustprogress
145+
define dso_local spir_func void @_Z6commonv() local_unnamed_addr {
146+
entry:
147+
call spir_func void @_Z6C_exclv()
148+
call spir_func void @_Z6D_inclv()
149+
ret void
150+
}
151+
152+
; Function Attrs: convergent norecurse nounwind mustprogress
153+
define dso_local spir_func void @_Z6D_inclv() local_unnamed_addr {
154+
entry:
155+
call spir_func void @_Z11assert_funcv()
156+
ret void
157+
}
158+
159+
; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress
160+
define dso_local spir_func void @_Z6C_exclv() local_unnamed_addr {
161+
entry:
162+
ret void
163+
}
164+
165+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel3
166+
; Function Attrs: convergent norecurse mustprogress
167+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel3"() local_unnamed_addr #0 {
168+
entry:
169+
call spir_func void @_Z6commonv()
170+
ret void
171+
}
172+
173+
; Function Attrs: convergent norecurse nounwind mustprogress
174+
define dso_local spir_func void @_Z7common2v() local_unnamed_addr {
175+
entry:
176+
call spir_func void @_Z6C_inclv()
177+
ret void
178+
}
179+
180+
; Function Attrs: convergent norecurse nounwind mustprogress
181+
define dso_local spir_func void @_Z6C_inclv() local_unnamed_addr {
182+
entry:
183+
call spir_func void @_Z11assert_funcv()
184+
ret void
185+
}
186+
187+
; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress
188+
define dso_local spir_func void @_Z6D_exclv() local_unnamed_addr {
189+
entry:
190+
ret void
191+
}
192+
193+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel4
194+
; Function Attrs: convergent norecurse mustprogress
195+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel4"() local_unnamed_addr #0 {
196+
entry:
197+
call spir_func void @_Z7common2v()
198+
ret void
199+
}
200+
201+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel5
202+
; Function Attrs: convergent norecurse mustprogress
203+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel5"() local_unnamed_addr #0 {
204+
entry:
205+
call spir_func void @_Z6B_inclv()
206+
call spir_func void @_Z6A_exclv()
207+
ret void
208+
}
209+
210+
; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress
211+
define dso_local spir_func void @_Z6E_exclv() local_unnamed_addr {
212+
entry:
213+
ret void
214+
}
215+
216+
; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel
217+
; Function Attrs: convergent norecurse mustprogress
218+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel6"() local_unnamed_addr #0 {
219+
entry:
220+
call spir_func void @_Z6E_exclv()
221+
call spir_func void @_Z6E_exclv()
222+
ret void
223+
}
224+
225+
; Function Attrs: convergent norecurse nounwind mustprogress
226+
define dso_local spir_func void @_Z6F_inclv() local_unnamed_addr {
227+
entry:
228+
call spir_func void @_Z11assert_funcv()
229+
ret void
230+
}
231+
232+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel7
233+
; Function Attrs: convergent norecurse mustprogress
234+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel7"() local_unnamed_addr #0 {
235+
entry:
236+
call spir_func void @_Z6F_inclv()
237+
call spir_func void @_Z6F_inclv()
238+
ret void
239+
}
240+
241+
; Function Attrs: convergent inlinehint norecurse nounwind mustprogress
242+
define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_2idILi1EEEE6_clES5_"() unnamed_addr align 2 {
243+
entry:
244+
call spir_func void @_Z1Gv()
245+
call spir_func void @_Z1Hv()
246+
ret void
247+
}
248+
249+
; Function Attrs: convergent norecurse nounwind mustprogress
250+
define dso_local spir_func void @_Z1Gv() local_unnamed_addr {
251+
entry:
252+
call spir_func void @_Z7common3v()
253+
ret void
254+
}
255+
256+
; Function Attrs: convergent norecurse nounwind mustprogress
257+
define dso_local spir_func void @_Z1Hv() local_unnamed_addr {
258+
entry:
259+
call spir_func void @_Z7common3v()
260+
ret void
261+
}
262+
263+
; Function Attrs: convergent norecurse nounwind mustprogress
264+
define dso_local spir_func void @_Z7common3v() local_unnamed_addr {
265+
entry:
266+
call spir_func void @_Z6I_inclv()
267+
ret void
268+
}
269+
270+
; Function Attrs: convergent norecurse nounwind mustprogress
271+
define dso_local spir_func void @_Z6I_inclv() local_unnamed_addr {
272+
entry:
273+
call spir_func void @_Z11assert_funcv()
274+
ret void
275+
}
276+
277+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel8
278+
; Function Attrs: convergent norecurse mustprogress
279+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel8"() local_unnamed_addr #0 {
280+
call spir_func void @_Z1Gv()
281+
call spir_func void @_Z1Hv()
282+
ret void
283+
}
284+
285+
; Function Attrs: convergent norecurse mustprogress
286+
define weak dso_local spir_func void @__assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func) local_unnamed_addr {
287+
entry:
288+
%call = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_xv()
289+
%call1 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_yv()
290+
%call2 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_zv()
291+
%call3 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_xv()
292+
%call4 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_yv()
293+
%call5 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_zv()
294+
tail call spir_func void @__devicelib_assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func, i64 %call, i64 %call1, i64 %call2, i64 %call3, i64 %call4, i64 %call5)
295+
ret void
296+
}
297+
298+
; Function Attrs: inlinehint norecurse mustprogress
299+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr
300+
301+
; Function Attrs: inlinehint norecurse mustprogress
302+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr
303+
304+
; Function Attrs: inlinehint norecurse mustprogress
305+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr
306+
307+
; Function Attrs: inlinehint norecurse mustprogress
308+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr
309+
310+
; Function Attrs: inlinehint norecurse mustprogress
311+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr
312+
313+
; Function Attrs: inlinehint norecurse mustprogress
314+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr
315+
316+
; Function Attrs: convergent norecurse mustprogress
317+
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) {
318+
entry:
319+
%call = tail call spir_func i32 (i8 addrspace(2)*, ...) @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)* getelementptr inbounds ([85 x i8], [85 x i8] addrspace(2)* @_ZL10assert_fmt, i64 0, i64 0), i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func, i64 %gid0, i64 %gid1, i64 %gid2, i64 %lid0, i64 %lid1, i64 %lid2, i8 addrspace(4)* %expr)
320+
ret void
321+
}
322+
323+
; Function Attrs: convergent
324+
declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...) local_unnamed_addr
325+
326+
attributes #0 = { convergent norecurse mustprogress "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="assert_test.cpp" "uniform-work-group-size"="true" }
327+
328+
!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0}
329+
!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1}
330+
!llvm.ident = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2}
331+
!llvm.module.flags = !{!3, !4}
332+
333+
!0 = !{i32 1, i32 2}
334+
!1 = !{i32 4, i32 100000}
335+
!2 = !{!"clang version 13.0.0 (https://github.com/intel/llvm)"}
336+
!3 = !{i32 1, !"wchar_size", i32 4}
337+
!4 = !{i32 7, !"frame-pointer", i32 2}
338+
!5 = !{i32 -1, i32 -1}

0 commit comments

Comments
 (0)