From 14d41bbda85a97d7c363a70432989300d11d95a7 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Wed, 2 Jun 2021 19:42:56 +0300 Subject: [PATCH 1/8] [sycl-post-link] Add device image property for assert feature A property with the name of the kernel is added whenever the kernel uses assert. Details: https://github.com/intel/llvm/blob/sycl/sycl/doc/Assert.md#online-linking-fallback-__devicelib_assert_fail --- llvm/include/llvm/Support/PropertySetIO.h | 1 + llvm/lib/Support/PropertySetIO.cpp | 1 + .../assert-property-with-split.ll | 145 ++++++++++++++++++ llvm/tools/sycl-post-link/sycl-post-link.cpp | 40 ++++- 4 files changed, 186 insertions(+), 1 deletion(-) create mode 100644 llvm/test/tools/sycl-post-link/assert-property-with-split.ll 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-with-split.ll b/llvm/test/tools/sycl-post-link/assert-property-with-split.ll new file mode 100644 index 0000000000000..8cdcb9378f2ea --- /dev/null +++ b/llvm/test/tools/sycl-post-link/assert-property-with-split.ll @@ -0,0 +1,145 @@ +; 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-DAG: _ZTSZ4mainE11TU0_kernel0 +; CHECK-DAG: _ZTSZ4mainE10TU1_kernel + +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 +} + +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 +} + +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 +define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 0 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 1 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 2 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 0 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 1 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 2 + ret i64 %1 +} + +; 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..298f860387d20 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -169,6 +169,7 @@ struct ImagePropSaveInfo { bool SpecConstsMet; bool EmitKernelParamInfo; bool IsEsimdKernel; + bool IsAssertEnabled; }; static void error(const Twine &Msg) { @@ -271,6 +272,28 @@ static void collectKernelModuleMap( } } +// Collect all the dependencies for the function. +static bool collectFunctionCallGraphNodes(llvm::Function *Func) { + std::vector Workqueue; + Workqueue.push_back(Func); + + while (!Workqueue.empty()) { + Function *F = &*Workqueue.back(); // To remove &* + Workqueue.pop_back(); + for (auto &I : instructions(F)) { + if (CallBase *CB = dyn_cast(&I)) + if (Function *CF = CB->getCalledFunction()) + if (!CF->isDeclaration()) { + if (CF->getName().startswith("__devicelib_assert_fail")) { + return true; + } + Workqueue.push_back(CF); + } + } + } + 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 +486,20 @@ static string_vector saveDeviceImageProperty( {"isEsimdImage", true}); } + if (ImgPSInfo.IsAssertEnabled) { + Module *M = ResultModules[I].get(); + std::vector SyclKernels; + for (auto &F : M->functions()) { + if (F.getCallingConv() == CallingConv::SPIR_KERNEL) { + if (collectFunctionCallGraphNodes(&F)) { + SyclKernels.push_back(&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_" : ""); @@ -609,7 +646,8 @@ static TableFiles processOneModule(std::unique_ptr M, bool IsEsimd, { ImagePropSaveInfo ImgPSInfo = { true, DoSpecConst, SetSpecConstAtRT, - SpecConstsMet, EmitKernelParamInfo, IsEsimd}; + SpecConstsMet, EmitKernelParamInfo, IsEsimd, + true}; string_vector Files = saveDeviceImageProperty(ResultModules, ImgPSInfo); std::copy(Files.begin(), Files.end(), std::back_inserter(TblFiles[COL_PROPS])); From e0251befd7627b16c7ae9d079dd5694c8212fc08 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Fri, 4 Jun 2021 20:35:58 +0300 Subject: [PATCH 2/8] Add test --- .../assert-property-with-split.ll | 2 + .../tools/sycl-post-link/assert-property.ll | 235 ++++++++++++++++++ 2 files changed, 237 insertions(+) create mode 100644 llvm/test/tools/sycl-post-link/assert-property.ll 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 index 8cdcb9378f2ea..46e9f5dac5e57 100644 --- a/llvm/test/tools/sycl-post-link/assert-property-with-split.ll +++ b/llvm/test/tools/sycl-post-link/assert-property-with-split.ll @@ -13,8 +13,10 @@ target triple = "spir64-unknown-linux-sycldevice" @_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-NOT: _ZTSZ4mainE11TU0_kernel1 ; CHECK-DAG: _ZTSZ4mainE11TU0_kernel0 ; CHECK-DAG: _ZTSZ4mainE10TU1_kernel +; CHECK-NOT: _ZTSZ4mainE11TU0_kernel1 define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 { entry: diff --git a/llvm/test/tools/sycl-post-link/assert-property.ll b/llvm/test/tools/sycl-post-link/assert-property.ll new file mode 100644 index 0000000000000..6d3680bb62803 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/assert-property.ll @@ -0,0 +1,235 @@ + +; 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) { +; bar(); +; baz(); +; }); +; }); +; 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] +; CHECK-DAG: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE9TeKernelEE +; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel +; CHECK-DAG: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel3EE +; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3 +; CHECK-NOT: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel2EE +; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2 + +; Function Attrs: convergent norecurse +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 { +entry: + %0 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* %_arg_, i64 0, i32 0, i32 0, i64 0 + %.sroa.0.0..sroa_cast9 = addrspacecast i64* %0 to i64 addrspace(4)* + %.sroa.0.0.copyload10 = load i64, i64 addrspace(4)* %.sroa.0.0..sroa_cast9, align 8 + %1 = 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 + %2 = extractelement <3 x i64> %1, i64 1 + %cmp.i.i = icmp ult i64 %2, 2147483648 + tail call void @llvm.assume(i1 %cmp.i.i) + %cmp.not.i = icmp ult i64 %2, %.sroa.0.0.copyload10 + br i1 %cmp.not.i, label %if.end.i, label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" + +if.end.i: ; preds = %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)) + br label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" + +"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit": ; preds = %entry, %if.end.i + ret void +} + +; Function Attrs: inaccessiblememonly nofree nosync nounwind willreturn +declare void @llvm.assume(i1 noundef) #1 + +; 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 +} + +; Function Attrs: convergent norecurse +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel"() #0 { +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 +} + +; Function Attrs: norecurse +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 { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 1 + %cmp.i.i = icmp ult i64 %1, 2147483648 + tail call void @llvm.assume(i1 %cmp.i.i) + ret void +} + +; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress +define dso_local spir_func void @_Z3barv() { +entry: + ret void +} + +; Function Attrs: norecurse +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2"() #1 { +entry: + ret void +} + +; Function Attrs: convergent norecurse +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 { +entry: + %0 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* %_arg_, i64 0, i32 0, i32 0, i64 0 + %.sroa.0.0..sroa_cast9 = addrspacecast i64* %0 to i64 addrspace(4)* + %.sroa.0.0.copyload10 = load i64, i64 addrspace(4)* %.sroa.0.0..sroa_cast9, align 8 + %1 = 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 + %2 = extractelement <3 x i64> %1, i64 1 + %cmp.i.i = icmp ult i64 %2, 2147483648 + tail call void @llvm.assume(i1 %cmp.i.i) + %cmp.not.i = icmp ult i64 %2, %.sroa.0.0.copyload10 + br i1 %cmp.not.i, label %if.end.i, label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E1_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" + +if.end.i: ; preds = %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)) + br label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E1_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" + +"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E1_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit": ; preds = %entry, %if.end.i + 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 +} + +; Function Attrs: convergent norecurse +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3"() #0 { +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 +} + +; 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 +define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 0 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 1 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 2 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 0 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 1 + ret i64 %1 +} + +; Function Attrs: inlinehint norecurse mustprogress +define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() { +entry: + %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 + %1 = extractelement <3 x i64> %0, i64 2 + ret i64 %1 +} + +; 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} From bda11f6d7e233d84efd9bce0accd210bcb7206a4 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Fri, 4 Jun 2021 20:51:51 +0300 Subject: [PATCH 3/8] Apply suggestions --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 19 +++++++++---------- 1 file changed, 9 insertions(+), 10 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 298f860387d20..0ca96a19d0545 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -272,23 +272,22 @@ static void collectKernelModuleMap( } } -// Collect all the dependencies for the function. -static bool collectFunctionCallGraphNodes(llvm::Function *Func) { +// Go through function call graph searching for assert call. +static bool hasAssertInFunctionCallGraph(llvm::Function *Func) { std::vector Workqueue; Workqueue.push_back(Func); while (!Workqueue.empty()) { - Function *F = &*Workqueue.back(); // To remove &* + Function *F = Workqueue.back(); Workqueue.pop_back(); for (auto &I : instructions(F)) { if (CallBase *CB = dyn_cast(&I)) - if (Function *CF = CB->getCalledFunction()) - if (!CF->isDeclaration()) { - if (CF->getName().startswith("__devicelib_assert_fail")) { - return true; - } + if (Function *CF = CB->getCalledFunction()) { + if (CF->getName().startswith("__devicelib_assert_fail")) + return true; + if (!CF->isDeclaration()) Workqueue.push_back(CF); - } + } } } return false; @@ -491,7 +490,7 @@ static string_vector saveDeviceImageProperty( std::vector SyclKernels; for (auto &F : M->functions()) { if (F.getCallingConv() == CallingConv::SPIR_KERNEL) { - if (collectFunctionCallGraphNodes(&F)) { + if (hasAssertInFunctionCallGraph(&F)) { SyclKernels.push_back(&F); PropSet[llvm::util::PropertySetRegistry::SYCL_ASSERT_USED].insert( {F.getName(), true}); From 2b7841260ba9d06a13f4c5111e3cb31aa04249a6 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Wed, 9 Jun 2021 13:52:53 +0300 Subject: [PATCH 4/8] Update tests and apply suggestions --- .../assert-property-with-split.ll | 53 +++++------------ .../tools/sycl-post-link/assert-property.ll | 57 +++++-------------- llvm/tools/sycl-post-link/sycl-post-link.cpp | 8 +-- 3 files changed, 32 insertions(+), 86 deletions(-) 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 index 46e9f5dac5e57..2d0f7e1bb5f30 100644 --- a/llvm/test/tools/sycl-post-link/assert-property-with-split.ll +++ b/llvm/test/tools/sycl-post-link/assert-property-with-split.ll @@ -1,3 +1,7 @@ +; 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 @@ -13,11 +17,8 @@ target triple = "spir64-unknown-linux-sycldevice" @_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-NOT: _ZTSZ4mainE11TU0_kernel1 -; CHECK-DAG: _ZTSZ4mainE11TU0_kernel0 -; CHECK-DAG: _ZTSZ4mainE10TU1_kernel -; CHECK-NOT: _ZTSZ4mainE11TU0_kernel1 +; CHECK: _ZTSZ4mainE11TU0_kernel0 define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 { entry: call spir_func void @_Z3foov() @@ -35,6 +36,7 @@ entry: ret void } +; CHECK-NOT: _ZTSZ4mainE11TU0_kernel1 define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel1() #0 { entry: call spir_func void @_Z4foo1v() @@ -49,6 +51,7 @@ entry: ret void } +; CHECK: _ZTSZ4mainE10TU1_kernel define dso_local spir_kernel void @_ZTSZ4mainE10TU1_kernel() #1 { entry: call spir_func void @_Z4foo2v() @@ -80,52 +83,22 @@ entry: } ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 0 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 1 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 2 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 0 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 1 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 2 - ret i64 %1 -} +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 { diff --git a/llvm/test/tools/sycl-post-link/assert-property.ll b/llvm/test/tools/sycl-post-link/assert-property.ll index 6d3680bb62803..f4b55458cc380 100644 --- a/llvm/test/tools/sycl-post-link/assert-property.ll +++ b/llvm/test/tools/sycl-post-link/assert-property.ll @@ -1,3 +1,6 @@ +; 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 @@ -46,13 +49,8 @@ target triple = "spir64_x86_64-unknown-unknown-sycldevice" @_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-DAG: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE9TeKernelEE -; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel -; CHECK-DAG: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel3EE -; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3 -; CHECK-NOT: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel2EE -; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2 +; CHECK: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE9TeKernelEE ; Function Attrs: convergent norecurse 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 { entry: @@ -84,6 +82,7 @@ entry: ret void } +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel ; Function Attrs: convergent norecurse define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel"() #0 { entry: @@ -91,6 +90,7 @@ entry: ret void } +; CHECK-NOT: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel2EE ; Function Attrs: norecurse 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 { entry: @@ -107,12 +107,14 @@ entry: ret void } +; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2 ; Function Attrs: norecurse define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2"() #1 { entry: ret void } +; CHECK: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel3EE ; Function Attrs: convergent norecurse 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 { entry: @@ -141,6 +143,7 @@ entry: ret void } +; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3 ; Function Attrs: convergent norecurse define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3"() #0 { entry: @@ -162,52 +165,22 @@ entry: } ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 0 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 1 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 2 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 0 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 1 - ret i64 %1 -} +declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr ; Function Attrs: inlinehint norecurse mustprogress -define weak dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 2 - ret i64 %1 -} +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) { diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 0ca96a19d0545..1554de5c29e3c 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -169,7 +169,6 @@ struct ImagePropSaveInfo { bool SpecConstsMet; bool EmitKernelParamInfo; bool IsEsimdKernel; - bool IsAssertEnabled; }; static void error(const Twine &Msg) { @@ -485,10 +484,12 @@ static string_vector saveDeviceImageProperty( {"isEsimdImage", true}); } - if (ImgPSInfo.IsAssertEnabled) { + { Module *M = ResultModules[I].get(); std::vector SyclKernels; 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)) { SyclKernels.push_back(&F); @@ -645,8 +646,7 @@ static TableFiles processOneModule(std::unique_ptr M, bool IsEsimd, { ImagePropSaveInfo ImgPSInfo = { true, DoSpecConst, SetSpecConstAtRT, - SpecConstsMet, EmitKernelParamInfo, IsEsimd, - true}; + SpecConstsMet, EmitKernelParamInfo, IsEsimd}; string_vector Files = saveDeviceImageProperty(ResultModules, ImgPSInfo); std::copy(Files.begin(), Files.end(), std::back_inserter(TblFiles[COL_PROPS])); From 19303047c0bf58552d76e84af698848b9a00adc0 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Thu, 10 Jun 2021 12:54:24 +0300 Subject: [PATCH 5/8] Re-work call graph traversal --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 39 +++++++++++++++++++- 1 file changed, 37 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 1554de5c29e3c..e8b00ddcc9261 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -273,21 +273,56 @@ 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 FuncList; + std::vector Workqueue; Workqueue.push_back(Func); while (!Workqueue.empty()) { Function *F = Workqueue.back(); Workqueue.pop_back(); + if (F != Func) + FuncList.push_back(F); + + bool IsVertex = true; for (auto &I : instructions(F)) { if (CallBase *CB = dyn_cast(&I)) if (Function *CF = CB->getCalledFunction()) { - if (CF->getName().startswith("__devicelib_assert_fail")) + // Return if we've already discovered if there are asserts in the + // function call graph. + if (hasAssertionInCallGraphMap.count(CF)) { + return hasAssertionInCallGraphMap[CF]; + } + + if (CF->getName().startswith("__devicelib_assert_fail")) { + // Mark all the functions above in call graph as ones that can call + // assert. + for (auto *It : FuncList) + hasAssertionInCallGraphMap[It] = true; + + hasAssertionInCallGraphMap[Func] = true; + hasAssertionInCallGraphMap[CF] = true; + return true; - if (!CF->isDeclaration()) + } + + if (!CF->isDeclaration()) { Workqueue.push_back(CF); + IsVertex = false; + } } } + + if (IsVertex) { + // Mark the above functions as ones that definetely do not call assert. + for (auto *It : FuncList) + hasAssertionInCallGraphMap[It] = false; + FuncList.clear(); + } } return false; } From fe84701b38dfb6f805589ad50161e68128da2a87 Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Sat, 12 Jun 2021 00:34:28 +0300 Subject: [PATCH 6/8] Tests update --- .../tools/sycl-post-link/assert-property-2.ll | 338 ++++++++++++++++++ .../tools/sycl-post-link/assert-property.ll | 65 +--- 2 files changed, 346 insertions(+), 57 deletions(-) create mode 100644 llvm/test/tools/sycl-post-link/assert-property-2.ll 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.ll b/llvm/test/tools/sycl-post-link/assert-property.ll index f4b55458cc380..40d0822883e6d 100644 --- a/llvm/test/tools/sycl-post-link/assert-property.ll +++ b/llvm/test/tools/sycl-post-link/assert-property.ll @@ -26,8 +26,8 @@ ; bar(); ; }); ; CGH.parallel_for(range<2>{2, 10}, [=](item<2> It) { -; bar(); ; baz(); +; bar(); ; }); ; }); ; Q.wait(); @@ -50,31 +50,6 @@ target triple = "spir64_x86_64-unknown-unknown-sycldevice" ; CHECK: [SYCL/assert used] -; CHECK: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE9TeKernelEE -; Function Attrs: convergent norecurse -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 { -entry: - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* %_arg_, i64 0, i32 0, i32 0, i64 0 - %.sroa.0.0..sroa_cast9 = addrspacecast i64* %0 to i64 addrspace(4)* - %.sroa.0.0.copyload10 = load i64, i64 addrspace(4)* %.sroa.0.0..sroa_cast9, align 8 - %1 = 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 - %2 = extractelement <3 x i64> %1, i64 1 - %cmp.i.i = icmp ult i64 %2, 2147483648 - tail call void @llvm.assume(i1 %cmp.i.i) - %cmp.not.i = icmp ult i64 %2, %.sroa.0.0.copyload10 - br i1 %cmp.not.i, label %if.end.i, label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" - -if.end.i: ; preds = %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)) - br label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" - -"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit": ; preds = %entry, %if.end.i - ret void -} - -; Function Attrs: inaccessiblememonly nofree nosync nounwind willreturn -declare void @llvm.assume(i1 noundef) #1 - ; Function Attrs: convergent norecurse nounwind mustprogress define dso_local spir_func void @_Z3foov() { entry: @@ -86,18 +61,7 @@ entry: ; Function Attrs: convergent norecurse define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel"() #0 { 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-NOT: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel2EE -; Function Attrs: norecurse -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 { -entry: - %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 - %1 = extractelement <3 x i64> %0, i64 1 - %cmp.i.i = icmp ult i64 %1, 2147483648 - tail call void @llvm.assume(i1 %cmp.i.i) + call spir_func void @_Z3foov() ret void } @@ -111,28 +75,15 @@ entry: ; Function Attrs: norecurse define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2"() #1 { entry: + call spir_func void @_Z3barv() ret void } -; CHECK: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperIZZ4mainENK3$_0clERNS0_7handlerEE10TheKernel3EE -; Function Attrs: convergent norecurse -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 { +; Function Attrs: convergent inlinehint norecurse nounwind mustprogress +define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_4itemILi2ELb1EEEE1_clES5_"() unnamed_addr #8 align 2 { entry: - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* %_arg_, i64 0, i32 0, i32 0, i64 0 - %.sroa.0.0..sroa_cast9 = addrspacecast i64* %0 to i64 addrspace(4)* - %.sroa.0.0.copyload10 = load i64, i64 addrspace(4)* %.sroa.0.0..sroa_cast9, align 8 - %1 = 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 - %2 = extractelement <3 x i64> %1, i64 1 - %cmp.i.i = icmp ult i64 %2, 2147483648 - tail call void @llvm.assume(i1 %cmp.i.i) - %cmp.not.i = icmp ult i64 %2, %.sroa.0.0.copyload10 - br i1 %cmp.not.i, label %if.end.i, label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E1_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" - -if.end.i: ; preds = %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)) - br label %"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E1_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit" - -"_ZZN2cl4sycl7handler27getRangeRoundedKernelLambdaINS0_4itemILi2ELb1EEELi2EZZ4mainENK3$_0clERS1_EUlS4_E1_LPv0EEEDaT1_NS0_5rangeIXT0_EEEENKUlS4_E_clES4_.exit": ; preds = %entry, %if.end.i + call spir_func void @_Z3bazv() + call spir_func void @_Z3barv() ret void } @@ -147,7 +98,7 @@ entry: ; Function Attrs: convergent norecurse define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3"() #0 { 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)) + call spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_4itemILi2ELb1EEEE1_clES5_"() ret void } From edee7c9c03d07a211e2f674d4f24ebda1ca4d6fd Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Sat, 12 Jun 2021 00:34:58 +0300 Subject: [PATCH 7/8] apply suggestions --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 73 +++++++++++--------- 1 file changed, 41 insertions(+), 32 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index e8b00ddcc9261..c36e528e9300a 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -277,51 +277,60 @@ static bool hasAssertInFunctionCallGraph(llvm::Function *Func) { // 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 FuncList; + std::vector FuncCallStack; - std::vector Workqueue; - Workqueue.push_back(Func); + std::vector Workstack; + Workstack.push_back(Func); - while (!Workqueue.empty()) { - Function *F = Workqueue.back(); - Workqueue.pop_back(); + while (!Workstack.empty()) { + Function *F = Workstack.back(); + Workstack.pop_back(); if (F != Func) - FuncList.push_back(F); + FuncCallStack.push_back(F); - bool IsVertex = true; + bool IsLeaf = true; for (auto &I : instructions(F)) { - if (CallBase *CB = dyn_cast(&I)) - if (Function *CF = CB->getCalledFunction()) { - // Return if we've already discovered if there are asserts in the - // function call graph. - if (hasAssertionInCallGraphMap.count(CF)) { - return hasAssertionInCallGraphMap[CF]; - } + 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 : FuncList) - hasAssertionInCallGraphMap[It] = 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; + hasAssertionInCallGraphMap[Func] = true; + hasAssertionInCallGraphMap[CF] = true; - return true; - } + return true; + } - if (!CF->isDeclaration()) { - Workqueue.push_back(CF); - IsVertex = false; - } - } + if (!CF->isDeclaration()) { + Workstack.push_back(CF); + IsLeaf = false; + } } - if (IsVertex) { + if (IsLeaf) { // Mark the above functions as ones that definetely do not call assert. - for (auto *It : FuncList) + for (auto *It : FuncCallStack) hasAssertionInCallGraphMap[It] = false; - FuncList.clear(); + FuncCallStack.clear(); } } return false; From d8606525b8dbe3648c63cadb9e5519510a4377bc Mon Sep 17 00:00:00 2001 From: Viktoria Maksimova Date: Tue, 15 Jun 2021 14:55:29 +0300 Subject: [PATCH 8/8] minor changes --- .../{assert-property.ll => assert-property-1.ll} | 0 llvm/tools/sycl-post-link/sycl-post-link.cpp | 5 +---- 2 files changed, 1 insertion(+), 4 deletions(-) rename llvm/test/tools/sycl-post-link/{assert-property.ll => assert-property-1.ll} (100%) diff --git a/llvm/test/tools/sycl-post-link/assert-property.ll b/llvm/test/tools/sycl-post-link/assert-property-1.ll similarity index 100% rename from llvm/test/tools/sycl-post-link/assert-property.ll rename to llvm/test/tools/sycl-post-link/assert-property-1.ll diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index c36e528e9300a..3858259e43c99 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -530,16 +530,13 @@ static string_vector saveDeviceImageProperty( { Module *M = ResultModules[I].get(); - std::vector SyclKernels; 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)) { - SyclKernels.push_back(&F); + if (hasAssertInFunctionCallGraph(&F)) PropSet[llvm::util::PropertySetRegistry::SYCL_ASSERT_USED].insert( {F.getName(), true}); - } } } }