diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index c0b3f7779c159..8338b894bd109 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -216,7 +216,6 @@ class PropertySetRegistry { static constexpr char SYCL_KERNEL_PARAM_OPT_INFO[] = "SYCL/kernel param opt"; static constexpr char SYCL_PROGRAM_METADATA[] = "SYCL/program metadata"; static constexpr char SYCL_MISC_PROP[] = "SYCL/misc properties"; - static constexpr char SYCL_ASSERT_USED[] = "SYCL/assert used"; static constexpr char SYCL_KERNEL_NAMES[] = "SYCL/kernel names"; static constexpr char SYCL_EXPORTED_SYMBOLS[] = "SYCL/exported symbols"; static constexpr char SYCL_IMPORTED_SYMBOLS[] = "SYCL/imported symbols"; diff --git a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp b/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp index c4bd7129fefe4..07afffbecf552 100644 --- a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp +++ b/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp @@ -57,85 +57,6 @@ bool isModuleUsingTsan(const Module &M) { return M.getNamedGlobal("__TsanKernelMetadata"); } -// This function traverses over reversed call graph by BFS algorithm. -// It means that an edge links some function @func with functions -// which contain call of function @func. It starts from -// @StartingFunction and lifts up until it reach all reachable functions, -// or it reaches some function containing "referenced-indirectly" attribute. -// If it reaches "referenced-indirectly" attribute than it returns an empty -// Optional. -// Otherwise, it returns an Optional containing a list of reached -// SPIR kernel function's names. -static std::optional> traverseCGToFindSPIRKernels( - const std::vector &StartingFunctionVec) { - std::queue FunctionsToVisit; - std::unordered_set VisitedFunctions; - for (const Function *FPtr : StartingFunctionVec) - FunctionsToVisit.push(FPtr); - std::vector KernelNames; - - while (!FunctionsToVisit.empty()) { - const Function *F = FunctionsToVisit.front(); - FunctionsToVisit.pop(); - - auto InsertionResult = VisitedFunctions.insert(F); - // It is possible that we insert some particular function several - // times in functionsToVisit queue. - if (!InsertionResult.second) - continue; - - for (const auto *U : F->users()) { - const CallInst *CI = dyn_cast(U); - if (!CI) - continue; - - const Function *ParentF = CI->getFunction(); - - if (VisitedFunctions.count(ParentF)) - continue; - - if (ParentF->hasFnAttribute("referenced-indirectly")) - return {}; - - if (ParentF->getCallingConv() == CallingConv::SPIR_KERNEL) - KernelNames.push_back(ParentF->getName()); - - FunctionsToVisit.push(ParentF); - } - } - - return {std::move(KernelNames)}; -} - -static std::vector -getKernelNamesUsingSpecialFunctions(const Module &M, - const std::vector &FNames) { - std::vector SpecialFunctionVec; - for (const auto Fn : FNames) { - Function *FPtr = M.getFunction(Fn); - if (FPtr) - SpecialFunctionVec.push_back(FPtr); - } - - if (SpecialFunctionVec.size() == 0) - return {}; - - auto TraverseResult = traverseCGToFindSPIRKernels(SpecialFunctionVec); - - if (TraverseResult.has_value()) - return std::move(*TraverseResult); - - // Here we reached "referenced-indirectly", so we need to find all kernels and - // return them. - std::vector SPIRKernelNames; - for (const Function &F : M) { - if (F.getCallingConv() == CallingConv::SPIR_KERNEL) - SPIRKernelNames.push_back(F.getName()); - } - - return SPIRKernelNames; -} - // Gets 1- to 3-dimension work-group related information for function Func. // Returns an empty vector if not present. template @@ -449,13 +370,6 @@ PropSetRegTy computeModuleProperties(const Module &M, if (OptLevel != -1) PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "optLevel", OptLevel); } - { - std::vector AssertFuncNames{"__devicelib_assert_fail"}; - std::vector FuncNames = - getKernelNamesUsingSpecialFunctions(M, AssertFuncNames); - for (const StringRef &FName : FuncNames) - PropSet.add(PropSetRegTy::SYCL_ASSERT_USED, FName, true); - } { std::vector> ArgPos = getKernelNamesUsingImplicitLocalMem(M); diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index d2030a8825fde..b562c67ff1eb3 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -200,7 +200,6 @@ constexpr char PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[]; constexpr char PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO[]; constexpr char PropertySetRegistry::SYCL_PROGRAM_METADATA[]; constexpr char PropertySetRegistry::SYCL_MISC_PROP[]; -constexpr char PropertySetRegistry::SYCL_ASSERT_USED[]; constexpr char PropertySetRegistry::SYCL_KERNEL_NAMES[]; constexpr char PropertySetRegistry::SYCL_EXPORTED_SYMBOLS[]; constexpr char PropertySetRegistry::SYCL_IMPORTED_SYMBOLS[]; diff --git a/llvm/test/tools/sycl-post-link/assert/indirect-with-split-2.ll b/llvm/test/tools/sycl-post-link/assert/indirect-with-split-2.ll deleted file mode 100644 index e7011ae6de141..0000000000000 --- a/llvm/test/tools/sycl-post-link/assert/indirect-with-split-2.ll +++ /dev/null @@ -1,165 +0,0 @@ -; This test checks that the post-link tool properly generates "assert used" -; property. This case validates that indirectly called function without assert -; does not cause all the module kernels to be marked as ones that can call -; assert indirectly. - -; Per design doc, if a callgraph for indirect callable function -; (marked with "referenced-indirectly" attribute in IR) has a call to -; __devicelib_assert_fail, then all kernels in the module are conservatively -; marked as using asserts. - -; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop -check-prefixes=CHECK,CHECK0 \ -; RUN: --implicit-check-not TU1 -; RUN: FileCheck %s -input-file=%t_1.prop -check-prefixes=CHECK,CHECK1 \ -; RUN: --implicit-check-not TU0 -; -; CHECK: [SYCL/assert used] -; CHECK0-DAG: main_TU1_kernel0 -; CHECK0-DAG: main_TU1_kernel1 -; -; CHECK1: main_TU0_kernel0 - -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" - -@_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 - -; PRESENCE-CHECK: [SYCL/assert used] - -; PRESENCE-CHECK-DAG: main_TU1_kernel1 -define dso_local spir_kernel void @main_TU1_kernel1() #2 { -entry: - call spir_func void @foo() - call spir_func void @bar() - ret void -} - -define dso_local spir_func void @foo() #2 { -entry: - call spir_func void @_Z4foo1v() - ret void -} - -; ABSENCE-CHECK-NOT: empty_kernel -define dso_local spir_kernel void @empty_kernel() #2 { - %1 = ptrtoint void ()* @bar to i64 - ret void -} - -define dso_local spir_func void @bar() #2 { -entry: - call spir_func void @_Z3foov() ; call assert - call spir_func void @_Z4foo2v() ; indirectly called - ret void -} - -; PRESENCE-CHECK-DAG: main_TU0_kernel0 -define dso_local spir_kernel void @main_TU0_kernel0() #0 { -entry: - call spir_func void @_Z3foov() ; call assert - 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 -} - -; 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 -} - -; PRESENCE-CHECK-DAG: main_TU1_kernel0 -define dso_local spir_kernel void @main_TU1_kernel0() #2 { -entry: - call spir_func void @_Z3foov() ; call assert - ret void -} - -; ABSENCE-CHECK-NOT: main_TU0_kernel1 -define dso_local spir_kernel void @main_TU0_kernel1() #0 { -entry: - call spir_func void @_Z4foo1v() - ret void -} - -; This function is marked with "referenced-indirectly", but it doesn't call an assert -; Function Attrs: nounwind -define dso_local spir_func void @_Z4foo2v() #1 { -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 - 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 = { "referenced-indirectly" "sycl-module-id"="TU2.cpp" } -attributes #2 = { "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/test/tools/sycl-post-link/assert/indirect-with-split.ll b/llvm/test/tools/sycl-post-link/assert/indirect-with-split.ll deleted file mode 100644 index 639dd73359c1b..0000000000000 --- a/llvm/test/tools/sycl-post-link/assert/indirect-with-split.ll +++ /dev/null @@ -1,143 +0,0 @@ -; This test checks that the post-link tool properly generates "assert used" -; property for indirectly called assertions - it should include all the kernels -; even they do not call assertions in their call graph. -; Per design doc, if a callgraph for indirect callable function -; (marked with "referenced-indirectly" attribute in IR) has a call to -; __devicelib_assert_fail, then all kernels in the module are conservatively -; marked as using asserts. - -; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes=CHECK,CHECK1 \ -; RUN: --implicit-check-not TU0 -; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes=CHECK,CHECK0 \ -; RUN: --implicit-check-not TU1 --implicit-check-not kernel1 -; -; With recent improvements to device code split, this file is actually being -; split to two modules and one of them does not contain "indirectly-referenced" -; function, meaning that only direct users of 'assert' will be mentioned in -; device image properties. -; -; CHECK: [SYCL/assert used] -; CHECK0: main_TU0_kernel0 -; -; CHECK1-DAG: main_TU1_kernel0 -; CHECK1-DAG: main_TU1_kernel1 - -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" - -@_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 - -define dso_local spir_kernel void @main_TU0_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 @main_TU0_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 @main_TU1_kernel0() #2 { -entry: - call spir_func void @_Z3foov() - ret void -} - -define dso_local spir_kernel void @main_TU1_kernel1() #2 { -entry: - call spir_func void @_Z4foo2v() - ret void -} - -; This function is marked with "referenced-indirectly" -; Function Attrs: nounwind -define dso_local spir_func void @_Z4foo2v() #1 { -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 = { "referenced-indirectly" "sycl-module-id"="TU2.cpp" } -attributes #2 = { "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/test/tools/sycl-post-link/assert/property-1.ll b/llvm/test/tools/sycl-post-link/assert/property-1.ll deleted file mode 100644 index 2df5e11dbbb07..0000000000000 --- a/llvm/test/tools/sycl-post-link/assert/property-1.ll +++ /dev/null @@ -1,176 +0,0 @@ -; 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 -properties -split=auto -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not TheKernel2 -; -; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not TheKernel2 -; -; RUN: sycl-post-link -properties -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not TheKernel2 -; -; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes=CHECK-K3 -; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes=CHECK-K1 -; RUN: FileCheck %s -input-file=%t_2.prop --check-prefixes=CHECK-K2 - -; SYCL source: -; void foo() { -; assert(0); -; } -; void bar() { -; -; } -; 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; -; } -; -; CHECK: [SYCL/assert used] -; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel=1|1 -; CHECK-DAG: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3=1|1 -; -; CHECK-K1: [SYCL/assert used] -; CHECK-K1: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3=1|1 -; CHECK-K2-NOT: [SYCL/assert used] -; CHECK-K3: [SYCL/assert used] -; CHECK-K3: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel=1|1 - - -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" - -%"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 - -; 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: - 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 -} - -; 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 -} - -; 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 deleted file mode 100644 index 4fafcde79e829..0000000000000 --- a/llvm/test/tools/sycl-post-link/assert/property-2.ll +++ /dev/null @@ -1,393 +0,0 @@ -; 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 -properties -split=auto -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop -check-prefix=PRESENCE-CHECK -; RUN: FileCheck %s -input-file=%t_0.prop -check-prefix=ABSENCE-CHECK - -; 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(); } -; -; void no_assert_func() { -; return; -; } -; void common4() { -; assert_func(); -; no_assert_func(); -; } -; void J() { -; common4(); -; } -; -; int main() { -; queue Q; -; Q.submit([&] (handler& CGH) { -; CGH.parallel_for(range<1>{1}, [=](id<1> i) { -; J(); -; }); -; CGH.parallel_for(range<1>{1}, [=](id<1> i) { -; common4(); -; }); -; 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" - -@.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 - -; PRESENCE-CHECK: [SYCL/assert used] - -; Function Attrs: convergent noinline norecurse optnone mustprogress -define dso_local spir_func void @_Z1Jv() #3 { -entry: - call spir_func void @_Z7common4v() - ret void -} - -; Function Attrs: convergent noinline norecurse optnone mustprogress -define dso_local spir_func void @_Z7common4v() #3 { -entry: - call spir_func void @_Z11assert_funcv() - call spir_func void @_Z14no_assert_funcv() - ret void -} - -; PRESENCE-CHECK-DAG: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E7Kernel9 -; Function Attrs: convergent noinline norecurse mustprogress -define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E7Kernel9() #0 { -entry: - call spir_func void @_Z1Jv() - ret void -} - -; PRESENCE-CHECK-DAG: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E8Kernel10 -; Function Attrs: convergent noinline norecurse optnone mustprogress -define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E8Kernel10() #0 { -entry: - call spir_func void @_Z7common4v() - ret void -} - -; Function Attrs: convergent noinline norecurse nounwind optnone mustprogress -define dso_local spir_func void @_Z14no_assert_funcv() #2 { -entry: - ret void -} - -; 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 -} - -; PRESENCE-CHECK-DAG: _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 -} - -; PRESENCE-CHECK-DAG: _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 -} - -; PRESENCE-CHECK-DAG: _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 -} - -; PRESENCE-CHECK-DAG: _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 -} - -; PRESENCE-CHECK-DAG: _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 -} - -; 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 -} - -; PRESENCE-CHECK-DAG: _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 -} - -; PRESENCE-CHECK-DAG: _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 -} - -; ABSENCE-CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel6 -; 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 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/sycl/doc/design/PropertySets.md b/sycl/doc/design/PropertySets.md index 1a55b544154fc..4cfdb552d6622 100644 --- a/sycl/doc/design/PropertySets.md +++ b/sycl/doc/design/PropertySets.md @@ -130,15 +130,6 @@ Miscellaneous properties: | "specConstsReplacedWithDefault" | 32 bit integer. ("1") | 1 if the specialization constants have been replaced by their default values and 0 or missing otherwise. | -### [SYCL/assert used] - -__Key:__ Kernel name. - -__Value type:__ 32 bit integer. ("1") - -__Value:__ 1 if the kernel uses assertions and 0 or missing otherwise. - - ### [SYCL/kernel names] __Key:__ Kernel name. diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index f7bff57c2df9a..826bc39982408 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -3817,13 +3817,6 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { friend auto get_native(const queue &Obj) -> backend_return_t; -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES -#if __SYCL_USE_FALLBACK_ASSERT - friend event detail::submitAssertCapture(const queue &, event &, - const detail::code_location &); -#endif -#endif - template friend void ext::oneapi::experimental::detail::submit_impl( const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF, diff --git a/sycl/source/detail/compiler.hpp b/sycl/source/detail/compiler.hpp index 2fc0b56135b20..670c8ce957ff4 100644 --- a/sycl/source/detail/compiler.hpp +++ b/sycl/source/detail/compiler.hpp @@ -55,8 +55,6 @@ #define __SYCL_PROPERTY_SET_PROGRAM_METADATA "SYCL/program metadata" /// PropertySetRegistry::SYCL_MISC_PROP defined in PropertySetIO.h #define __SYCL_PROPERTY_SET_SYCL_MISC_PROP "SYCL/misc properties" -/// PropertySetRegistry::SYCL_ASSERT_USED defined in PropertySetIO.h -#define __SYCL_PROPERTY_SET_SYCL_ASSERT_USED "SYCL/assert used" /// PropertySetRegistry::SYCL_KERNEL_NAMES defined in PropertySetIO.h #define __SYCL_PROPERTY_SET_SYCL_KERNEL_NAMES "SYCL/kernel names" /// PropertySetRegistry::SYCL_EXPORTED_SYMBOLS defined in PropertySetIO.h diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp index c8ff57631bb60..4ed4eea90fccc 100644 --- a/sycl/source/detail/device_binary_image.cpp +++ b/sycl/source/detail/device_binary_image.cpp @@ -191,7 +191,6 @@ RTDeviceBinaryImage::RTDeviceBinaryImage(sycl_device_binary Bin) { DeviceLibReqMask.init(Bin, __SYCL_PROPERTY_SET_DEVICELIB_REQ_MASK); DeviceLibMetadata.init(Bin, __SYCL_PROPERTY_SET_DEVICELIB_METADATA); KernelParamOptInfo.init(Bin, __SYCL_PROPERTY_SET_KERNEL_PARAM_OPT_INFO); - AssertUsed.init(Bin, __SYCL_PROPERTY_SET_SYCL_ASSERT_USED); ImplicitLocalArg.init(Bin, __SYCL_PROPERTY_SET_SYCL_IMPLICIT_LOCAL_ARG); ProgramMetadata.init(Bin, __SYCL_PROPERTY_SET_PROGRAM_METADATA); // Convert ProgramMetadata into the UR format @@ -517,8 +516,6 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage( naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) { return Img.getKernelParamOptInfo(); }); - auto MergedAssertUsed = naiveMergeBinaryProperties( - Imgs, [](const RTDeviceBinaryImage &Img) { return Img.getAssertUsed(); }); auto MergedDeviceGlobals = naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) { return Img.getDeviceGlobals(); @@ -546,13 +543,12 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage( return Img.getRegisteredKernels(); }); - std::array *, 11> MergedVecs{ + std::array *, 10> MergedVecs{ &MergedSpecConstants, &MergedSpecConstantsDefaultValues, - &MergedKernelParamOptInfo, &MergedAssertUsed, - &MergedDeviceGlobals, &MergedHostPipes, - &MergedVirtualFunctions, &MergedImplicitLocalArg, - &MergedKernelNames, &MergedExportedSymbols, - &MergedRegisteredKernels}; + &MergedKernelParamOptInfo, &MergedDeviceGlobals, + &MergedHostPipes, &MergedVirtualFunctions, + &MergedImplicitLocalArg, &MergedKernelNames, + &MergedExportedSymbols, &MergedRegisteredKernels}; // Exclusive merges. auto MergedDeviceLibReqMask = @@ -672,7 +668,6 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage( CopyPropertiesVec(MergedSpecConstantsDefaultValues, SpecConstDefaultValuesMap); CopyPropertiesVec(MergedKernelParamOptInfo, KernelParamOptInfo); - CopyPropertiesVec(MergedAssertUsed, AssertUsed); CopyPropertiesVec(MergedDeviceGlobals, DeviceGlobals); CopyPropertiesVec(MergedHostPipes, HostPipes); CopyPropertiesVec(MergedVirtualFunctions, VirtualFunctions); diff --git a/sycl/source/detail/device_binary_image.hpp b/sycl/source/detail/device_binary_image.hpp index 075229effb3ec..2cd380c91bd65 100644 --- a/sycl/source/detail/device_binary_image.hpp +++ b/sycl/source/detail/device_binary_image.hpp @@ -223,7 +223,6 @@ class RTDeviceBinaryImage { const PropertyRange &getKernelParamOptInfo() const { return KernelParamOptInfo; } - const PropertyRange &getAssertUsed() const { return AssertUsed; } const PropertyRange &getProgramMetadata() const { return ProgramMetadata; } const std::vector &getProgramMetadataUR() const { return ProgramMetadataUR; @@ -259,7 +258,6 @@ class RTDeviceBinaryImage { RTDeviceBinaryImage::PropertyRange DeviceLibReqMask; RTDeviceBinaryImage::PropertyRange DeviceLibMetadata; RTDeviceBinaryImage::PropertyRange KernelParamOptInfo; - RTDeviceBinaryImage::PropertyRange AssertUsed; RTDeviceBinaryImage::PropertyRange ProgramMetadata; RTDeviceBinaryImage::PropertyRange KernelNames; RTDeviceBinaryImage::PropertyRange ExportedSymbols; diff --git a/sycl/source/detail/device_kernel_info.cpp b/sycl/source/detail/device_kernel_info.cpp index 526f160c6596b..5c6dfad0d633d 100644 --- a/sycl/source/detail/device_kernel_info.cpp +++ b/sycl/source/detail/device_kernel_info.cpp @@ -24,7 +24,6 @@ DeviceKernelInfo::DeviceKernelInfo(const CompileTimeKernelInfoTy &Info) void DeviceKernelInfo::init(KernelNameStrRefT KernelName) { auto &PM = detail::ProgramManager::getInstance(); - MUsesAssert = PM.kernelUsesAssert(KernelName); MImplicitLocalArgPos = PM.kernelImplicitLocalArgPos(KernelName); #ifndef __INTEL_PREVIEW_BREAKING_CHANGES MInitialized.store(true); @@ -78,10 +77,7 @@ FastKernelSubcacheT &DeviceKernelInfo::getKernelSubcache() { assertInitialized(); return MFastKernelSubcache; } -bool DeviceKernelInfo::usesAssert() { - assertInitialized(); - return MUsesAssert; -} + const std::optional &DeviceKernelInfo::getImplicitLocalArgPos() { assertInitialized(); return MImplicitLocalArgPos; diff --git a/sycl/source/detail/device_kernel_info.hpp b/sycl/source/detail/device_kernel_info.hpp index 0ea4ff2d051e6..c76db34a3227d 100644 --- a/sycl/source/detail/device_kernel_info.hpp +++ b/sycl/source/detail/device_kernel_info.hpp @@ -108,7 +108,6 @@ class DeviceKernelInfo : public CompileTimeKernelInfoTy { void setCompileTimeInfoIfNeeded(const CompileTimeKernelInfoTy &Info); FastKernelSubcacheT &getKernelSubcache(); - bool usesAssert(); const std::optional &getImplicitLocalArgPos(); private: @@ -119,7 +118,6 @@ class DeviceKernelInfo : public CompileTimeKernelInfoTy { std::atomic MInitialized = false; #endif FastKernelSubcacheT MFastKernelSubcache; - bool MUsesAssert; std::optional MImplicitLocalArgPos; }; diff --git a/sycl/source/detail/kernel_data.hpp b/sycl/source/detail/kernel_data.hpp index 9f22fafa85aa4..aebf7e245c1a3 100644 --- a/sycl/source/detail/kernel_data.hpp +++ b/sycl/source/detail/kernel_data.hpp @@ -137,11 +137,6 @@ class KernelData { void setKernelFunc(void *KernelFuncPtr) { MKernelFuncPtr = KernelFuncPtr; } - bool usesAssert() const { - assert(MDeviceKernelInfoPtr); - return MDeviceKernelInfoPtr->usesAssert(); - } - // Kernel launch properties getter and setters. ur_kernel_cache_config_t getKernelCacheConfig() const { return MKernelCacheConfig; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 1d39a89a4dd45..d7c5525587076 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1794,14 +1794,6 @@ Managed ProgramManager::build( return LinkedProg; } -void ProgramManager::cacheKernelUsesAssertInfo(const RTDeviceBinaryImage &Img) { - const RTDeviceBinaryImage::PropertyRange &AssertUsedRange = - Img.getAssertUsed(); - if (AssertUsedRange.isAvailable()) - for (const auto &Prop : AssertUsedRange) - m_KernelUsesAssert.insert(Prop->Name); -} - void ProgramManager::cacheKernelImplicitLocalArg( const RTDeviceBinaryImage &Img) { const RTDeviceBinaryImage::PropertyRange &ImplicitLocalArgRange = @@ -2044,8 +2036,6 @@ void ProgramManager::addImage(sycl_device_binary RawImg, m_KernelNameRefCount[name]++; } - cacheKernelUsesAssertInfo(*Img); - // check if kernel uses sanitizer { sycl_device_binary_property SanProp = Img->getProperty("sanUsed"); @@ -2233,7 +2223,6 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) { if (--RefCount == 0) { // TODO aggregate all these maps into a single one since their entries // share lifetime. - m_KernelUsesAssert.erase(Name); m_KernelImplicitLocalArgPos.erase(Name); m_DeviceKernelInfoMap.erase(Name); m_KernelNameRefCount.erase(RefCountIt); diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 73eca2cd86e0a..8357dd2f81915 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -371,11 +371,6 @@ class ProgramManager { ProgramManager(); ~ProgramManager() = default; - template - bool kernelUsesAssert(const NameT &KernelName) const { - return m_KernelUsesAssert.find(KernelName) != m_KernelUsesAssert.end(); - } - SanitizerType kernelUsesSanitizer() const { return m_SanitizerFoundInImage; } std::optional @@ -412,9 +407,6 @@ class ProgramManager { /// Dumps image to current directory void dumpImage(const RTDeviceBinaryImage &Img, uint32_t SequenceID = 0) const; - /// Add info on kernels using assert into cache - void cacheKernelUsesAssertInfo(const RTDeviceBinaryImage &Img); - /// Add info on kernels using local arg into cache void cacheKernelImplicitLocalArg(const RTDeviceBinaryImage &Img); @@ -528,8 +520,6 @@ class ProgramManager { // different types without temporary key_type object creation. This includes // standard overloads, such as comparison between std::string and // std::string_view or just char*. - using KernelUsesAssertSet = std::set>; - KernelUsesAssertSet m_KernelUsesAssert; std::unordered_map m_KernelImplicitLocalArgPos; // Map for storing device kernel information. Runtime lookup should be avoided diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index 5b7bfb5e90fae..7f38c9266a37e 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -435,13 +435,6 @@ EventImplPtr queue_impl::submit_kernel_scheduler_bypass( } bool DiscardEvent = !EventNeeded && supportsDiscardingPiEvents(); - if (DiscardEvent) { - // Kernel only uses assert if it's non interop one - bool KernelUsesAssert = - !(KernelImplPtr && KernelImplPtr->isInterop()) && KData.usesAssert(); - DiscardEvent = !KernelUsesAssert; - } - std::shared_ptr ResultEvent = DiscardEvent ? nullptr : detail::event_impl::create_device_event(*this); diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 2285bbce42761..9de9d63ecd188 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -3348,16 +3348,6 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { const std::shared_ptr &SyclKernel = ExecKernel->MSyclKernel; KernelNameStrRefT KernelName = ExecKernel->MDeviceKernelInfo.Name; - - if (!EventImpl) { - // Kernel only uses assert if it's non interop one - bool KernelUsesAssert = (!SyclKernel || SyclKernel->hasSYCLMetadata()) && - ExecKernel->MDeviceKernelInfo.usesAssert(); - if (KernelUsesAssert) { - EventImpl = MEvent.get(); - } - } - const RTDeviceBinaryImage *BinImage = nullptr; if (detail::SYCLConfig::get()) { BinImage = retrieveKernelBinary(*MQueue, KernelName); diff --git a/sycl/unittests/helpers/MockDeviceImage.hpp b/sycl/unittests/helpers/MockDeviceImage.hpp index 343e8045b8ac5..40afada687c65 100644 --- a/sycl/unittests/helpers/MockDeviceImage.hpp +++ b/sycl/unittests/helpers/MockDeviceImage.hpp @@ -487,15 +487,6 @@ inline MockProperty makeSpecConstant(std::vector &ValData, return Prop; } -/// Utility function to mark kernel as the one using assert -inline void setKernelUsesAssert(const std::vector &Names, - MockPropertySet &Set) { - std::vector Value; - for (const std::string &N : Names) - Value.push_back({N, {0, 0, 0, 0}, SYCL_PROPERTY_TYPE_UINT32}); - Set.insert(__SYCL_PROPERTY_SET_SYCL_ASSERT_USED, std::move(Value)); -} - /// Utility function to add specialization constants to property set. /// /// This function overrides the default spec constant values. diff --git a/sycl/unittests/program_manager/Cleanup.cpp b/sycl/unittests/program_manager/Cleanup.cpp index 1bcbfa7676255..fa5b356b13eee 100644 --- a/sycl/unittests/program_manager/Cleanup.cpp +++ b/sycl/unittests/program_manager/Cleanup.cpp @@ -73,8 +73,6 @@ class ProgramManagerExposed : public sycl::detail::ProgramManager { return m_EliminatedKernelArgMasks; } - KernelUsesAssertSet &getKernelUsesAssert() { return m_KernelUsesAssert; } - std::unordered_map & getKernelImplicitLocalArgPos() { return m_KernelImplicitLocalArgPos; @@ -184,8 +182,6 @@ sycl::unittest::MockDeviceImage generateImage(const std::string &ImageId, PropSet.insert(__SYCL_PROPERTY_SET_SYCL_VIRTUAL_FUNCTIONS, createVFPropertySet(VirtualFunctions)); - setKernelUsesAssert(std::vector{KernelNames.begin()[0]}, - PropSet); PropSet.insert(__SYCL_PROPERTY_SET_SYCL_IMPLICIT_LOCAL_ARG, createPropertySet(ImplicitLocalArg)); @@ -311,9 +307,6 @@ void checkAllInvolvedContainers(ProgramManagerExposed &PM, "Kernel name reference count " + CommentPostfix); EXPECT_EQ(PM.getEliminatedKernelArgMask().size(), ExpectedImgCount) << "Eliminated kernel arg mask " + CommentPostfix; - checkContainer(PM.getKernelUsesAssert(), ExpectedEntryCount, - generateRefNames(ImgIds, "Kernel"), - "KernelUsesAssert " + CommentPostfix); EXPECT_EQ(PM.getKernelImplicitLocalArgPos().size(), ExpectedEntryCount) << "Kernel implicit local arg pos " + CommentPostfix;