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