diff --git a/libdevice/cmake/modules/SYCLLibdevice.cmake b/libdevice/cmake/modules/SYCLLibdevice.cmake
index 255cd2abcd617..fc2c81806150c 100644
--- a/libdevice/cmake/modules/SYCLLibdevice.cmake
+++ b/libdevice/cmake/modules/SYCLLibdevice.cmake
@@ -2,11 +2,32 @@ set(obj_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}")
 if (WIN32)
   set(lib-suffix obj)
   set(spv_binary_dir "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}")
+  set(host_offload_target "host-x86_64-pc-windows-msvc")
 else()
   set(lib-suffix o)
   set(spv_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}")
+  set(host_offload_target "host-x86_64-unknown-linux-gnu")
 endif()
 set(clang $<TARGET_FILE:clang>)
+set(bundler $<TARGET_FILE:clang-offload-bundler>)
+set(llvm-link $<TARGET_FILE:llvm-link>)
+set(llvm-spirv $<TARGET_FILE:llvm-spirv>)
+
+string(CONCAT bundler_targets_opt
+  "-targets="
+  "sycl-spir64_x86_64-unknown-unknown,"
+  "sycl-spir64_gen-unknown-unknown,"
+  "sycl-spir64_fpga-unknown-unknown,"
+  "sycl-spir64-unknown-unknown,"
+  ${host_offload_target})
+string(CONCAT bundler_inputs_opt
+  "-inputs="
+  ${obj_binary_dir} "/fallback-cassert_spir64_x86_64." ${lib-suffix} ","
+  ${obj_binary_dir} "/fallback-cassert_spir64_gen." ${lib-suffix} ","
+  ${obj_binary_dir} "/fallback-cassert_spir64_fpga." ${lib-suffix} ","
+  ${obj_binary_dir} "/fallback-cassert_spir64." ${lib-suffix} ","
+  ${obj_binary_dir} "/assert_no_read_host." ${lib-suffix})
+
 
 string(CONCAT sycl_targets_opt
   "-fsycl-targets="
@@ -81,6 +102,61 @@ add_custom_command(OUTPUT ${devicelib-obj-cmath-fp64}
                    DEPENDS device_math.h device.h sycl-compiler
                    VERBATIM)
 
+################################################################################
+#[[
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                   MAIN_DEPENDENCY fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/assert_read_spir64_spv.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${spv_binary_dir}/assert_read_spir64_spv.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   COMMAND ${llvm-link} -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                           ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                           ${spv_binary_dir}/assert_read_spir64_spv.bc
+                   DEPENDS ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                           ${spv_binary_dir}/assert_read_spir64_spv.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                   COMMAND ${llvm-spirv} -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   DEPENDS ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                           llvm-spirv
+                   VERBATIM)
+]]
+
+
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   MAIN_DEPENDENCY fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                   COMMAND ${llvm-spirv} -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   DEPENDS ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                           llvm-spirv
+                   VERBATIM)
+#[[
 add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
                    COMMAND ${clang} -fsycl-device-only -fno-sycl-use-bitcode
                            ${compile_opts}
@@ -89,6 +165,8 @@ add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
                    MAIN_DEPENDENCY fallback-cassert.cpp
                    DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
                    VERBATIM)
+]]
+################################################################################
 
 add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cstring.spv
                    COMMAND ${clang} -fsycl-device-only -fno-sycl-use-bitcode
@@ -99,6 +177,127 @@ add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cstring.spv
                    DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
                    VERBATIM)
 
+################################################################################
+#[[
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64_x86_64-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64_gen.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64_gen-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64_gen.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64_fpga-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64_x86_64-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64_gen-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64_fpga-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_host.o
+                   COMMAND ${clang} -c
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_host.o
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64_x86_64.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64_x86_64.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64_gen.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64_gen.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64_gen.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64_gen.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64_fpga.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64_fpga.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                           llvm-link
+                   VERBATIM)
+
+add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffix}
+                   COMMAND ${bundler} -type=o ${bundler_targets_opt}
+                           -outputs=${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffix}
+                           ${bundler_inputs_opt}
+                   DEPENDS ${obj_binary_dir}/fallback-cassert_spir64_x86_64.${lib-suffix}
+                           ${obj_binary_dir}/fallback-cassert_spir64_gen.${lib-suffix}
+                           ${obj_binary_dir}/fallback-cassert_spir64_fpga.${lib-suffix}
+                           ${obj_binary_dir}/fallback-cassert_spir64.${lib-suffix}
+                           ${obj_binary_dir}/assert_no_read_host.${lib-suffix}
+                   VERBATIM)
+]]
+
+
 add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffix}
                    COMMAND ${clang} -fsycl -c
                            ${compile_opts} ${sycl_targets_opt}
@@ -108,6 +307,8 @@ add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffi
                    DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
                    VERBATIM)
 
+################################################################################
+
 add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cstring.${lib-suffix}
                    COMMAND ${clang} -fsycl -c
                            ${compile_opts} ${sycl_targets_opt}
diff --git a/libdevice/fallback-cassert.cl b/libdevice/fallback-cassert.cl
new file mode 100644
index 0000000000000..b3d1228d86f07
--- /dev/null
+++ b/libdevice/fallback-cassert.cl
@@ -0,0 +1,51 @@
+#pragma OPENCL EXTENSION cl_khr_int64_base_atomics:enable
+#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics:enable
+
+// NOTE Align these definitions with fallback-cassert.cpp
+#define ASSERT_NONE 0
+#define ASSERT_START 1
+#define ASSERT_FINISH 2
+
+// NOTE Layout of this structure should be aligned with the one in
+// sycl/include/CL/sycl/detail/assert_happened.hpp
+struct AssertHappened {
+  int Flag;
+  char Expr[256 + 1];
+  char File[256 + 1];
+  char Func[128 + 1];
+
+  int Line;
+
+  unsigned long GID0;
+  unsigned long GID1;
+  unsigned long GID2;
+
+  unsigned long LID0;
+  unsigned long LID1;
+  unsigned long LID2;
+};
+
+typedef struct AssertHappened AssertHappenedT;
+
+extern __global AssertHappenedT SPIR_AssertHappenedMem;
+
+__kernel void __devicelib_assert_read(__global void *_Dst) {
+  if (!_Dst)
+    return;
+
+  AssertHappenedT *Dst = (AssertHappenedT *)_Dst;
+
+  __global int *FlagPtr = &SPIR_AssertHappenedMem.Flag;
+  int Flag = atomic_add(FlagPtr, 0);
+
+  if (ASSERT_NONE == Flag) {
+    Dst->Flag = ASSERT_NONE;
+    return;
+  }
+
+  if (Flag != ASSERT_FINISH)
+    while (ASSERT_START == atomic_add(FlagPtr, 0))
+      ;
+
+  *Dst = SPIR_AssertHappenedMem;
+}
diff --git a/libdevice/fallback-cassert.cpp b/libdevice/fallback-cassert.cpp
index b03a3409b7bf8..41c76831df650 100644
--- a/libdevice/fallback-cassert.cpp
+++ b/libdevice/fallback-cassert.cpp
@@ -12,6 +12,7 @@
 
 #ifdef __SPIR__
 
+// NOTE Align these definitions with fallback-cassert.cl
 #define ASSERT_NONE 0
 #define ASSERT_START 1
 #define ASSERT_FINISH 2
@@ -19,7 +20,7 @@
 // definition
 SPIR_GLOBAL AssertHappened SPIR_AssertHappenedMem;
 
-DEVICE_EXTERN_C void __devicelib_assert_read(void *_Dst) {
+void __devicelib_assert_read(__SYCL_GLOBAL__ void *_Dst) {
   AssertHappened *Dst = (AssertHappened *)_Dst;
   int Flag = atomicLoad(&SPIR_AssertHappenedMem.Flag);
 
@@ -35,11 +36,58 @@ DEVICE_EXTERN_C void __devicelib_assert_read(void *_Dst) {
   *Dst = SPIR_AssertHappenedMem;
 }
 
+#ifdef SYCL_LANGUAGE_VERSION
+#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel]]
+#else
+#define __SYCL_KERNEL_ATTR__
+#endif
+
+// Create a kernel entry point for __devicelib_assert_read
+template <typename KernelName, typename KernelType>
+__SYCL_KERNEL_ATTR__ __attribute__((noinline)) void
+kernel_caller(KernelType KernelFunc) {
+  KernelFunc();
+}
+
+namespace detail {
+class DevicelibAssertReadKernel {
+  __SYCL_GLOBAL__ void *MDst;
+public:
+  DevicelibAssertReadKernel(__SYCL_GLOBAL__ void *Dst) : MDst{Dst} {}
+  __attribute__((noinline))
+  void operator()() const {
+    __devicelib_assert_read(MDst);
+  }
+};
+} // namespace detail
+
+
+using DevicelibAssertReadT = detail::DevicelibAssertReadKernel;//void (*)(__SYCL_GLOBAL__ void *);
+
+__attribute__((noinline))
+struct DevicelibAssertReadKernelName;
+
+// a stub function to enforce assert info reader kernel in devicelib image
+__attribute__((noinline))
+void __devicelib_stub() {
+  kernel_caller<DevicelibAssertReadKernelName, DevicelibAssertReadT>(DevicelibAssertReadT{NULL});
+}
+
 DEVICE_EXTERN_C void __devicelib_assert_fail(const char *expr, const char *file,
                                              int32_t line, const char *func,
                                              uint64_t gid0, uint64_t gid1,
                                              uint64_t gid2, uint64_t lid0,
                                              uint64_t lid1, uint64_t lid2) {
+  // FIXME make offline linking against __devicelib_assert_fail enforce linking
+  // against __devicelib_assert_read also
+  {
+    __devicelib_stub();
+//    kernel_caller<DevicelibAssertReadKernelName, DevicelibAssertReadT>(DevicelibAssertReadT{NULL});
+//    __devicelib_stub();
+//    __devicelib_assert_read(NULL);
+  }
+
+
   int Expected = ASSERT_NONE;
   int Desired = ASSERT_START;
 
diff --git a/libdevice/include/assert-happened.hpp b/libdevice/include/assert-happened.hpp
index 8b50f5ef216b3..3720a6a2ac3da 100644
--- a/libdevice/include/assert-happened.hpp
+++ b/libdevice/include/assert-happened.hpp
@@ -43,6 +43,6 @@ struct AssertHappened {
 #define __SYCL_GLOBAL__ __attribute__((opencl_global))
 
 // declaration
-extern SPIR_GLOBAL_VAR __SYCL_GLOBAL__ AssertHappened SPIR_AssertHappenedMem;
+extern "C" SPIR_GLOBAL_VAR __SYCL_GLOBAL__ AssertHappened SPIR_AssertHappenedMem;
 
 #endif
diff --git a/sycl/include/CL/sycl/detail/assert_happened.hpp b/sycl/include/CL/sycl/detail/assert_happened.hpp
index 7c5f45c3fc4e5..d84dbcf50118b 100644
--- a/sycl/include/CL/sycl/detail/assert_happened.hpp
+++ b/sycl/include/CL/sycl/detail/assert_happened.hpp
@@ -12,12 +12,6 @@
 
 #include <cstdint>
 
-#ifdef __SYCL_DEVICE_ONLY__
-// Reads Flag of AssertHappened on device
-SYCL_EXTERNAL __attribute__((weak)) extern "C" void
-__devicelib_assert_read(void *);
-#endif
-
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp
index f21ffec77230c..5bc94d2020eb7 100644
--- a/sycl/include/CL/sycl/queue.hpp
+++ b/sycl/include/CL/sycl/queue.hpp
@@ -83,10 +83,12 @@ class queue;
 
 namespace detail {
 class queue_impl;
+#if 0
 #if __SYCL_USE_FALLBACK_ASSERT
 static event submitAssertCapture(queue &, event &, queue *,
                                  const detail::code_location &);
 #endif
+#endif
 } // namespace detail
 
 /// Encapsulates a single SYCL queue which schedules kernels on a SYCL device.
@@ -261,8 +263,8 @@ class __SYCL_EXPORT queue {
           // Linking against fallback impl of __devicelib_assert_fail is
           // performed by program manager class
           // Fallback assert isn't supported for FPGA
-          submitAssertCapture(*this, E, /* SecondaryQueue = */ nullptr,
-                              CodeLoc);
+          //submitAssertCapture(*this, E, /* SecondaryQueue = */ nullptr,
+          //                    CodeLoc);
         }
       };
 
@@ -307,7 +309,7 @@ class __SYCL_EXPORT queue {
           // Linking against fallback impl of __devicelib_assert_fail is
           // performed by program manager class
           // Fallback assert isn't supported for FPGA
-          submitAssertCapture(*this, E, DeviceSecondaryQueue, CodeLoc);
+          //submitAssertCapture(*this, E, DeviceSecondaryQueue, CodeLoc);
         }
       };
 
@@ -1073,9 +1075,11 @@ class __SYCL_EXPORT queue {
   template <class T>
   friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
 
+#if 0
 #if __SYCL_USE_FALLBACK_ASSERT
   friend event detail::submitAssertCapture(queue &, event &, queue *,
                                            const detail::code_location &);
+#endif
 #endif
 
   /// A template-free version of submit.
@@ -1183,6 +1187,7 @@ class __SYCL_EXPORT queue {
 };
 
 namespace detail {
+#if 0
 #if __SYCL_USE_FALLBACK_ASSERT
 #define __SYCL_ASSERT_START 1
 /**
@@ -1263,6 +1268,7 @@ event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
 }
 #undef __SYCL_ASSERT_START
 #endif // __SYCL_USE_FALLBACK_ASSERT
+#endif
 } // namespace detail
 
 } // namespace sycl
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 7ba084f315890..b2c643aa68fa7 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -43,7 +43,7 @@ namespace detail {
 
 using ContextImplPtr = std::shared_ptr<cl::sycl::detail::context_impl>;
 
-static constexpr int DbgProgMgr = 0;
+static constexpr int DbgProgMgr = 999;
 
 enum BuildState { BS_InProgress, BS_Done, BS_Failed };
 
diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp
index ed2ee3e6f78dc..530e31b11bbd9 100644
--- a/sycl/source/detail/scheduler/graph_builder.cpp
+++ b/sycl/source/detail/scheduler/graph_builder.cpp
@@ -1018,6 +1018,8 @@ Scheduler::GraphBuilder::addCG(std::unique_ptr<detail::CG> CommandGroup,
     NewCmd->MEmptyCmd =
         addEmptyCmd(NewCmd.get(), NewCmd->getCG().MRequirements, Queue,
                     Command::BlockReason::HostTask, ToEnqueue);
+  else if (CGType == CG::CGTYPE::Kernel)
+    addAssertInfoCheckerCGs(NewCmd.get());
 
   if (MPrintOptionsArray[AfterAddCG])
     printGraphAsDot("after_addCG");
@@ -1361,6 +1363,73 @@ Command *Scheduler::GraphBuilder::connectDepEvent(
   return ConnectCmd;
 }
 
+Command *Scheduler::GraphBuilder::addAssertInfoCheckerCGs(Command *Cmd) {
+  assert(Cmd->getType() == Command::RUN_CG &&
+         static_cast<ExecCGCommand *>(Cmd)->getCG().getType() == CG::Kernel &&
+         "Only kernel commands are allowed to be appended with assert info "
+         "copier/checker");
+
+  ExecCGCommand *KernelCmd = static_cast<ExecCGCommand *>(Cmd);
+  CGExecKernel &KernelCG = static_cast<CGExecKernel &>(KernelCmd->getCG());
+
+  (void)KernelCG;
+
+  bool FallbackAssertDisabled = false; // TODO check queue for compile-time value
+                                       // check env var also
+  bool KernelUsesAssert = true; // TODO check program manager.
+                                 // imply interop kernel doesn't use assert
+
+  // Don't enqueue assert info copier and checker in case of either of:
+  //  * kernel is launched on host
+  //  * fallback assert is disabled
+  //  * kernel doesn't use assert
+  if (KernelCmd->getWorkerQueue()->is_host() || FallbackAssertDisabled ||
+      !KernelUsesAssert)
+    return nullptr;
+
+  // 1. create CG for copier kernel, depend it on user's kernel => CopierEv
+  NDRDescT NDRDesc;
+  NDRDesc.set(range<1>{1});
+/*
+  NDRDescT NDRDesc,
+  std::unique_ptr<HostKernelBase> HKernel = nullptr,
+  std::shared_ptr<detail::kernel_impl> SyclKernel = get kernel from devicelib,
+  std::vector<std::vector<char>> ArgsStorage,
+  std::vector<detail::AccessorImplPtr> AccStorage,
+  std::vector<std::shared_ptr<const void>> SharedPtrStorage,
+  std::vector<Requirement *> Requirements,
+  std::vector<detail::EventImplPtr> Events,
+  std::vector<ArgDesc> Args, std::string KernelName,
+  detail::OSModuleHandle OSModuleHandle,
+  std::vector<std::shared_ptr<detail::stream_impl>> Streams,
+  CGTYPE Type, detail::code_location loc = {}
+*/
+  {
+  RT::PiKernel Kernel;
+  std::mutex *Mtx;
+  RT::PiProgram Prg;
+  std::tie(Kernel, Mtx, Prg) = ProgramManager::getInstance().getOrCreateKernel(
+      (OSModuleHandle)(-1),
+      KernelCmd->getWorkerQueue()->getContextImplPtr(),
+      KernelCmd->getWorkerQueue()->getDeviceImplPtr(),
+      "_ZTS29DevicelibAssertReadKernelName",
+      nullptr);
+  fprintf(stderr, "Kernel: %p, Mtx: %p, Prg: %p\n",
+          (const void *)Kernel, (const void *)Mtx, (const void *)Prg);
+  }
+#if 0
+  std::unique_ptr<detail::CG> CopierCG(new detail::CGExecKernel(
+      NDRDesc, HostKernel, SyclKernel,
+      // TODO
+  ));
+#endif
+  // 2. create CG for host task, depend on CopierEv
+
+  // TODO
+
+  return nullptr;
+}
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/source/detail/scheduler/scheduler.hpp b/sycl/source/detail/scheduler/scheduler.hpp
index 18ed2f5004c06..cc91aef092bf2 100644
--- a/sycl/source/detail/scheduler/scheduler.hpp
+++ b/sycl/source/detail/scheduler/scheduler.hpp
@@ -609,6 +609,8 @@ class Scheduler {
                 std::vector<Command *> &ToEnqueue,
                 const bool AddDepsToLeaves = true);
 
+    Command *addAssertInfoCheckerCGs(Command *KernelCmd);
+
   protected:
     /// Finds a command dependency corresponding to the record.
     DepDesc findDepForRecord(Command *Cmd, MemObjRecord *Record);