diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index cd7688d10f2db..3c332e3ca1f0b 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -189,6 +189,7 @@ class PropertySetRegistry { static constexpr char SYCL_DEVICELIB_REQ_MASK[] = "SYCL/devicelib req mask"; static constexpr char SYCL_KERNEL_PARAM_OPT_INFO[] = "SYCL/kernel param opt"; static constexpr char SYCL_MISC_PROP[] = "SYCL/misc properties"; + static constexpr char SYCL_ASSERT_USED[] = "SYCL/assert used"; // Function for bulk addition of an entire property set under given category // (property set name). diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index 97016364e0d0a..03ad6bb0b4cf9 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -198,6 +198,7 @@ constexpr char PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK[]; constexpr char PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[]; constexpr char PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO[]; constexpr char PropertySetRegistry::SYCL_MISC_PROP[]; +constexpr char PropertySetRegistry::SYCL_ASSERT_USED[]; } // namespace util } // namespace llvm diff --git a/llvm/test/tools/sycl-post-link/assert-property-1.ll b/llvm/test/tools/sycl-post-link/assert-property-1.ll new file mode 100644 index 0000000000000..40d0822883e6d --- /dev/null +++ b/llvm/test/tools/sycl-post-link/assert-property-1.ll @@ -0,0 +1,159 @@ +; This test checks that the post-link tool properly generates "assert used" +; property - it should include only kernels that call assertions in their call +; graph. + +; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: FileCheck %s -input-file=%t_0.prop + +; SYCL source: +; void foo() { +; assert(0); +; } +; void bar() { +; assert(1); +; } +; void baz() { +; foo(); +; } +; +; int main() { +; queue Q; +; Q.submit([&] (handler& CGH) { +; CGH.parallel_for(range<2>{2, 10}, [=](item<2> It) { +; foo(); +; }); +; CGH.parallel_for(range<2>{2, 10}, [=](item<2> It) { +; bar(); +; }); +; CGH.parallel_for(range<2>{2, 10}, [=](item<2> It) { +; baz(); +; bar(); +; }); +; }); +; Q.wait(); +; return 0; +; } + +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_x86_64-unknown-unknown-sycldevice" + +%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" } +%"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" = type { [2 x i64] } +%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon" = type { i8 } + +@.str = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", align 1 +@.str.1 = private unnamed_addr addrspace(1) constant [11 x i8] c"assert.cpp\00", align 1 +@__PRETTY_FUNCTION__._Z3foov = private unnamed_addr addrspace(1) constant [11 x i8] c"void foo()\00", align 1 +@__spirv_BuiltInGlobalInvocationId = external dso_local addrspace(1) constant <3 x i64>, align 32 +@__spirv_BuiltInLocalInvocationId = external dso_local addrspace(1) constant <3 x i64>, align 32 +@_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 + +; CHECK: [SYCL/assert used] + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z3foov() { +entry: + 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)) + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel +; Function Attrs: convergent norecurse +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel"() #0 { +entry: + call spir_func void @_Z3foov() + ret void +} + +; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress +define dso_local spir_func void @_Z3barv() { +entry: + ret void +} + +; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2 +; Function Attrs: norecurse +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2"() #1 { +entry: + call spir_func void @_Z3barv() + ret void +} + +; Function Attrs: convergent inlinehint norecurse nounwind mustprogress +define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_4itemILi2ELb1EEEE1_clES5_"() unnamed_addr #8 align 2 { +entry: + call spir_func void @_Z3bazv() + call spir_func void @_Z3barv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z3bazv() { +entry: + 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)) + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3 +; Function Attrs: convergent norecurse +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3"() #0 { +entry: + call spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_4itemILi2ELb1EEEE1_clES5_"() + ret void +} + +; Function Attrs: convergent norecurse mustprogress +define weak dso_local spir_func void @__assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func) { +entry: + %call = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() + %call1 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() + %call2 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() + %call3 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_xv() + %call4 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_yv() + %call5 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_zv() + 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) + ret void +} + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr + +; Function Attrs: convergent norecurse mustprogress +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) { +entry: + %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) + ret void +} + +; Function Attrs: convergent +declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...) + +attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="assert.cpp" "uniform-work-group-size"="true" } +attributes #1 = { norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="assert.cpp" "uniform-work-group-size"="true" } + +!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0} +!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1} +!llvm.ident = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2} +!llvm.module.flags = !{!3, !4} + +!0 = !{i32 1, i32 2} +!1 = !{i32 4, i32 100000} +!2 = !{!"clang version 13.0.0 (https://github.com/intel/llvm)"} +!3 = !{i32 1, !"wchar_size", i32 4} +!4 = !{i32 7, !"frame-pointer", i32 2} +!5 = !{i32 -1, i32 -1} diff --git a/llvm/test/tools/sycl-post-link/assert-property-2.ll b/llvm/test/tools/sycl-post-link/assert-property-2.ll new file mode 100644 index 0000000000000..8b707c1a4ef4c --- /dev/null +++ b/llvm/test/tools/sycl-post-link/assert-property-2.ll @@ -0,0 +1,338 @@ +; This test checks that the post-link tool properly generates "assert used" +; property - it should include only kernels that call assertions in their call +; graph. + +; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: FileCheck %s -input-file=%t_0.prop + +; SYCL source: +; void assert_func() { +; assert(0); +; } +; +; void A_excl() {} +; void B_incl() { assert_func(); } +; +; void A_incl() { assert_func(); } +; void B_excl() {} +; +; void C_excl() {} +; void D_incl() { assert_func(); } +; void common() { +; C_excl(); +; D_incl(); +; } +; +; void C_incl() { assert_func(); } +; void D_excl() {} +; void common2() { +; C_incl(); +; D_excl(); +; } +; +; void E_excl() {} +; void F_incl() { assert_func(); } +; +; void I_incl() { assert_func(); } +; void common3() { I_incl();} +; void G() { common3(); } +; void H() { common3(); } +; +; int main() { +; queue Q; +; Q.submit([&] (handler& CGH) { +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; A_excl(); +; B_incl(); +; }); +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; A_incl(); +; B_excl(); +; }); +; +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; common(); +; }); +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; common2(); +; }); +; +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; B_incl(); +; A_excl(); +; }); +; +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; E_excl(); +; E_excl(); +; }); +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; F_incl(); +; F_incl(); +; }); +; +; CGH.parallel_for(range<1>{1}, [=](id<1> i) { +; G(); +; H(); +; }); +; }); +; Q.wait(); +; return 0; +; } + +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_x86_64-unknown-unknown-sycldevice" + +@.str = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", align 1 +@.str.1 = private unnamed_addr addrspace(1) constant [16 x i8] c"assert_test.cpp\00", align 1 +@__PRETTY_FUNCTION__._Z11assert_funcv = private unnamed_addr addrspace(1) constant [19 x i8] c"void assert_func()\00", align 1 +@_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 + +; CHECK: [SYCL/assert used] + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z6B_inclv() local_unnamed_addr { +entry: + call spir_func void @_Z11assert_funcv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z11assert_funcv() local_unnamed_addr { +entry: + 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)) + ret void +} + +; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress +define dso_local spir_func void @_Z6A_exclv() local_unnamed_addr { +entry: + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE6Kernel +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE6Kernel"() local_unnamed_addr #0 { +entry: + call spir_func void @_Z6A_exclv() + call spir_func void @_Z6B_inclv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z6A_inclv() local_unnamed_addr { +entry: + call spir_func void @_Z11assert_funcv() + ret void +} + +; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress +define dso_local spir_func void @_Z6B_exclv() local_unnamed_addr { +entry: + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel2 +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel2"() local_unnamed_addr #0 { +entry: + call spir_func void @_Z6A_inclv() + call spir_func void @_Z6B_exclv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z6commonv() local_unnamed_addr { +entry: + call spir_func void @_Z6C_exclv() + call spir_func void @_Z6D_inclv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z6D_inclv() local_unnamed_addr { +entry: + call spir_func void @_Z11assert_funcv() + ret void +} + +; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress +define dso_local spir_func void @_Z6C_exclv() local_unnamed_addr { +entry: + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel3 +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel3"() local_unnamed_addr #0 { +entry: + call spir_func void @_Z6commonv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z7common2v() local_unnamed_addr { +entry: + call spir_func void @_Z6C_inclv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z6C_inclv() local_unnamed_addr { +entry: + call spir_func void @_Z11assert_funcv() + ret void +} + +; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress +define dso_local spir_func void @_Z6D_exclv() local_unnamed_addr { +entry: + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel4 +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel4"() local_unnamed_addr #0 { +entry: + call spir_func void @_Z7common2v() + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel5 +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel5"() local_unnamed_addr #0 { +entry: + call spir_func void @_Z6B_inclv() + call spir_func void @_Z6A_exclv() + ret void +} + +; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress +define dso_local spir_func void @_Z6E_exclv() local_unnamed_addr { +entry: + ret void +} + +; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel6"() local_unnamed_addr #0 { +entry: + call spir_func void @_Z6E_exclv() + call spir_func void @_Z6E_exclv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z6F_inclv() local_unnamed_addr { +entry: + call spir_func void @_Z11assert_funcv() + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel7 +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel7"() local_unnamed_addr #0 { +entry: + call spir_func void @_Z6F_inclv() + call spir_func void @_Z6F_inclv() + ret void +} + +; Function Attrs: convergent inlinehint norecurse nounwind mustprogress +define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_2idILi1EEEE6_clES5_"() unnamed_addr align 2 { +entry: + call spir_func void @_Z1Gv() + call spir_func void @_Z1Hv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z1Gv() local_unnamed_addr { +entry: + call spir_func void @_Z7common3v() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z1Hv() local_unnamed_addr { +entry: + call spir_func void @_Z7common3v() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z7common3v() local_unnamed_addr { +entry: + call spir_func void @_Z6I_inclv() + ret void +} + +; Function Attrs: convergent norecurse nounwind mustprogress +define dso_local spir_func void @_Z6I_inclv() local_unnamed_addr { +entry: + call spir_func void @_Z11assert_funcv() + ret void +} + +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel8 +; Function Attrs: convergent norecurse mustprogress +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel8"() local_unnamed_addr #0 { + call spir_func void @_Z1Gv() + call spir_func void @_Z1Hv() + ret void +} + +; Function Attrs: convergent norecurse mustprogress +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 { +entry: + %call = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() + %call1 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() + %call2 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() + %call3 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_xv() + %call4 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_yv() + %call5 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_zv() + 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) + ret void +} + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr + +; Function Attrs: convergent norecurse mustprogress +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) { +entry: + %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) + ret void +} + +; Function Attrs: convergent +declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...) local_unnamed_addr + +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" } + +!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0} +!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1} +!llvm.ident = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2} +!llvm.module.flags = !{!3, !4} + +!0 = !{i32 1, i32 2} +!1 = !{i32 4, i32 100000} +!2 = !{!"clang version 13.0.0 (https://github.com/intel/llvm)"} +!3 = !{i32 1, !"wchar_size", i32 4} +!4 = !{i32 7, !"frame-pointer", i32 2} +!5 = !{i32 -1, i32 -1} diff --git a/llvm/test/tools/sycl-post-link/assert-property-with-split.ll b/llvm/test/tools/sycl-post-link/assert-property-with-split.ll new file mode 100644 index 0000000000000..2d0f7e1bb5f30 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/assert-property-with-split.ll @@ -0,0 +1,120 @@ +; This test checks that the post-link tool properly generates "assert used" +; property in split mode - it should include only kernels that call assertions +; in their call graph. + +; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: FileCheck %s -input-file=%t_0.prop + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-linux-sycldevice" + +@_ZL2GV = internal addrspace(1) constant [1 x i32] [i32 42], align 4 +@.str = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", align 1 +@.str.1 = private unnamed_addr addrspace(1) constant [11 x i8] c"assert.cpp\00", align 1 +@__PRETTY_FUNCTION__._Z3foov = private unnamed_addr addrspace(1) constant [11 x i8] c"void foo()\00", align 1 +@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 +@__spirv_BuiltInLocalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 +@_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 + +; CHECK: [SYCL/assert used] + +; CHECK: _ZTSZ4mainE11TU0_kernel0 +define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 { +entry: + call spir_func void @_Z3foov() + ret void +} + +define dso_local spir_func void @_Z3foov() { +entry: + %a = alloca i32, align 4 + %ptr = bitcast i32* %a to i32 (i32)* + %call = call spir_func i32 %ptr(i32 1) + %add = add nsw i32 2, %call + store i32 %add, i32* %a, align 4 + 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)) + ret void +} + +; CHECK-NOT: _ZTSZ4mainE11TU0_kernel1 +define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel1() #0 { +entry: + call spir_func void @_Z4foo1v() + ret void +} + +; Function Attrs: nounwind +define dso_local spir_func void @_Z4foo1v() { +entry: + %a = alloca i32, align 4 + store i32 2, i32* %a, align 4 + ret void +} + +; CHECK: _ZTSZ4mainE10TU1_kernel +define dso_local spir_kernel void @_ZTSZ4mainE10TU1_kernel() #1 { +entry: + call spir_func void @_Z4foo2v() + ret void +} + +; Function Attrs: nounwind +define dso_local spir_func void @_Z4foo2v() { +entry: + %a = alloca i32, align 4 + %0 = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @_ZL2GV to [1 x i32] addrspace(4)*), i64 0, i64 0), align 4 + %add = add nsw i32 4, %0 + store i32 %add, i32* %a, align 4 + 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)) + ret void +} + +; Function Attrs: convergent norecurse mustprogress +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 { +entry: + %call = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() + %call1 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() + %call2 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() + %call3 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_xv() + %call4 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_yv() + %call5 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_zv() + 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) + ret void +} + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr + +; Function Attrs: inlinehint norecurse mustprogress +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr + +; Function Attrs: convergent norecurse mustprogress +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 { +entry: + %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) + ret void +} + +; Function Attrs: convergent +declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...) local_unnamed_addr + +attributes #0 = { "sycl-module-id"="TU1.cpp" } +attributes #1 = { "sycl-module-id"="TU2.cpp" } + +!opencl.spir.version = !{!0, !0} +!spirv.Source = !{!1, !1} + +!0 = !{i32 1, i32 2} +!1 = !{i32 4, i32 100000} diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 304ce183d44c8..3858259e43c99 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -271,6 +271,71 @@ static void collectKernelModuleMap( } } +// Go through function call graph searching for assert call. +static bool hasAssertInFunctionCallGraph(llvm::Function *Func) { + // Map holds the info about assertions in already examined functions: + // true - if there is an assertion in underlying functions, + // false - if there are definetely no assertions in underlying functions. + static std::map hasAssertionInCallGraphMap; + std::vector FuncCallStack; + + std::vector Workstack; + Workstack.push_back(Func); + + while (!Workstack.empty()) { + Function *F = Workstack.back(); + Workstack.pop_back(); + if (F != Func) + FuncCallStack.push_back(F); + + bool IsLeaf = true; + for (auto &I : instructions(F)) { + if (!isa(&I)) + continue; + + Function *CF = cast(&I)->getCalledFunction(); + if (!CF) + continue; + + // Return if we've already discovered if there are asserts in the + // function call graph. + if (hasAssertionInCallGraphMap.count(CF)) { + // If we know, that this function does not contain assert, we still + // should investigate another instructions in the function. + if (!hasAssertionInCallGraphMap[CF]) + continue; + + return true; + } + + if (CF->getName().startswith("__devicelib_assert_fail")) { + // Mark all the functions above in call graph as ones that can call + // assert. + for (auto *It : FuncCallStack) + hasAssertionInCallGraphMap[It] = true; + + hasAssertionInCallGraphMap[Func] = true; + hasAssertionInCallGraphMap[CF] = true; + + return true; + } + + if (!CF->isDeclaration()) { + Workstack.push_back(CF); + IsLeaf = false; + } + } + + if (IsLeaf) { + // Mark the above functions as ones that definetely do not call assert. + for (auto *It : FuncCallStack) + hasAssertionInCallGraphMap[It] = false; + FuncCallStack.clear(); + } + } + return false; +} + // Input parameter KernelModuleMap is a map containing groups of kernels with // same values of the sycl-module-id attribute. ResSymbolsLists is a vector of // kernel name lists. Each vector element is a string with kernel names from the @@ -463,6 +528,19 @@ static string_vector saveDeviceImageProperty( {"isEsimdImage", true}); } + { + Module *M = ResultModules[I].get(); + for (auto &F : M->functions()) { + // TODO: handle SYCL_EXTERNAL functions for dynamic linkage. + // TODO: handle function pointers. + if (F.getCallingConv() == CallingConv::SPIR_KERNEL) { + if (hasAssertInFunctionCallGraph(&F)) + PropSet[llvm::util::PropertySetRegistry::SYCL_ASSERT_USED].insert( + {F.getName(), true}); + } + } + } + std::error_code EC; std::string SCFile = makeResultFileName(".prop", I, ImgPSInfo.IsEsimdKernel ? "esimd_" : "");