diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
index 3687b2f029b9e..1f11db722b60a 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
@@ -99,15 +99,16 @@ template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
 template <typename CommandGroupFunc, typename PropertiesT>
 void submit_impl(queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc) {
-  Q.submit_without_event(Props, std::forward<CommandGroupFunc>(CGF), CodeLoc);
+  Q.submit_without_event<__SYCL_USE_FALLBACK_ASSERT>(
+      Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
 }
 
 template <typename CommandGroupFunc, typename PropertiesT>
 event submit_with_event_impl(queue &Q, PropertiesT Props,
                              CommandGroupFunc &&CGF,
                              const sycl::detail::code_location &CodeLoc) {
-  return Q.submit_with_event(Props, std::forward<CommandGroupFunc>(CGF),
-                             nullptr, CodeLoc);
+  return Q.submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
+      Props, detail::type_erased_cgfo_ty{CGF}, nullptr, CodeLoc);
 }
 } // namespace detail
 
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 758daa3a81a9b..b815869f4d707 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -162,6 +162,38 @@ class graph_impl;
 } // namespace ext::oneapi::experimental::detail
 namespace detail {
 
+class type_erased_cgfo_ty {
+  // From SYCL 2020,  command group function object:
+  // A type which is callable with operator() that takes a reference to a
+  // command group handler, that defines a command group which can be submitted
+  // by a queue. The function object can be a named type, lambda function or
+  // std::function.
+  template <typename T> struct invoker {
+    static void call(void *object, handler &cgh) {
+      (*static_cast<T *>(object))(cgh);
+    }
+  };
+  void *object;
+  using invoker_ty = void (*)(void *, handler &);
+  const invoker_ty invoker_f;
+
+public:
+  template <class T>
+  type_erased_cgfo_ty(T &f)
+      // NOTE: Even if `T` is a pointer to a function, `&f` is a pointer to a
+      // pointer to a function and as such can be casted to `void *` (pointer to
+      // a function cannot be casted).
+      : object(static_cast<void *>(&f)), invoker_f(&invoker<T>::call) {}
+  ~type_erased_cgfo_ty() = default;
+
+  type_erased_cgfo_ty(const type_erased_cgfo_ty &) = delete;
+  type_erased_cgfo_ty(type_erased_cgfo_ty &&) = delete;
+  type_erased_cgfo_ty &operator=(const type_erased_cgfo_ty &) = delete;
+  type_erased_cgfo_ty &operator=(type_erased_cgfo_ty &&) = delete;
+
+  void operator()(sycl::handler &cgh) const { invoker_f(object, cgh); }
+};
+
 class kernel_bundle_impl;
 class work_group_memory_impl;
 class handler_impl;
diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp
index e2208d452d100..a311d9a0c36a8 100644
--- a/sycl/include/sycl/queue.hpp
+++ b/sycl/include/sycl/queue.hpp
@@ -75,10 +75,8 @@ auto get_native(const SyclObjectT &Obj)
 namespace detail {
 class queue_impl;
 
-#if __SYCL_USE_FALLBACK_ASSERT
 inline event submitAssertCapture(queue &, event &, queue *,
                                  const detail::code_location &);
-#endif
 
 // Function to postprocess submitted command
 // Arguments:
@@ -375,8 +373,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
       T CGF,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
-    return submit_with_event(
-        sycl::ext::oneapi::experimental::empty_properties_t{}, CGF,
+    return submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
+        sycl::ext::oneapi::experimental::empty_properties_t{},
+        detail::type_erased_cgfo_ty{CGF},
         /*SecondaryQueuePtr=*/nullptr, CodeLoc);
   }
 
@@ -395,9 +394,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
       T CGF, queue &SecondaryQueue,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
-    return submit_with_event(
-        sycl::ext::oneapi::experimental::empty_properties_t{}, CGF,
-        &SecondaryQueue, CodeLoc);
+    return submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
+        sycl::ext::oneapi::experimental::empty_properties_t{},
+        detail::type_erased_cgfo_ty{CGF}, &SecondaryQueue, CodeLoc);
   }
 
   /// Prevents any commands submitted afterward to this queue from executing
@@ -2786,6 +2785,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   /// TODO: Unused. Remove these when ABI-break window is open.
+  /// Not using `type_erased_cgfo_ty` on purpose.
   event submit_impl(std::function<void(handler &)> CGH,
                     const detail::code_location &CodeLoc);
   event submit_impl(std::function<void(handler &)> CGH,
@@ -2815,16 +2815,28 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
       std::function<void(handler &)> CGH, queue secondQueue,
       const detail::code_location &CodeLoc,
       const detail::SubmitPostProcessF &PostProcess, bool IsTopCodeLoc);
+
+  // Old version when `std::function` was used in place of
+  // `std::function<void(handler &)>`.
+  event submit_with_event_impl(std::function<void(handler &)> CGH,
+                               const detail::SubmissionInfo &SubmitInfo,
+                               const detail::code_location &CodeLoc,
+                               bool IsTopCodeLoc);
+
+  void submit_without_event_impl(std::function<void(handler &)> CGH,
+                                 const detail::SubmissionInfo &SubmitInfo,
+                                 const detail::code_location &CodeLoc,
+                                 bool IsTopCodeLoc);
 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
 
   /// A template-free versions of submit.
-  event submit_with_event_impl(std::function<void(handler &)> CGH,
+  event submit_with_event_impl(const detail::type_erased_cgfo_ty &CGH,
                                const detail::SubmissionInfo &SubmitInfo,
                                const detail::code_location &CodeLoc,
                                bool IsTopCodeLoc);
 
   /// A template-free version of submit_without_event.
-  void submit_without_event_impl(std::function<void(handler &)> CGH,
+  void submit_without_event_impl(const detail::type_erased_cgfo_ty &CGH,
                                  const detail::SubmissionInfo &SubmitInfo,
                                  const detail::code_location &CodeLoc,
                                  bool IsTopCodeLoc);
@@ -2836,32 +2848,35 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   /// \param CGF is a function object containing command group.
   /// \param CodeLoc is the code location of the submit call (default argument)
   /// \return a SYCL event object for the submitted command group.
-  template <typename T, typename PropertiesT>
-  std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event>
-  submit_with_event(
-      PropertiesT Props, T CGF, queue *SecondaryQueuePtr,
+  //
+  // UseFallBackAssert as template param vs `#if` in function body is necessary
+  // to prevent ODR-violation between TUs built with different fallback assert
+  // modes.
+  template <bool UseFallbackAssert, typename PropertiesT>
+  event submit_with_event(
+      PropertiesT Props, const detail::type_erased_cgfo_ty &CGF,
+      queue *SecondaryQueuePtr,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
     detail::SubmissionInfo SI{};
     ProcessSubmitProperties(Props, SI);
     if (SecondaryQueuePtr)
       SI.SecondaryQueue() = detail::getSyclObjImpl(*SecondaryQueuePtr);
-#if __SYCL_USE_FALLBACK_ASSERT
-    SI.PostProcessorFunc() =
-        [this, &SecondaryQueuePtr,
-         &TlsCodeLocCapture](bool IsKernel, bool KernelUsesAssert, event &E) {
-          if (IsKernel && !device_has(aspect::ext_oneapi_native_assert) &&
-              KernelUsesAssert && !device_has(aspect::accelerator)) {
-            // __devicelib_assert_fail isn't supported by Device-side Runtime
-            // Linking against fallback impl of __devicelib_assert_fail is
-            // performed by program manager class
-            // Fallback assert isn't supported for FPGA
-            submitAssertCapture(*this, E, SecondaryQueuePtr,
-                                TlsCodeLocCapture.query());
-          }
-        };
-#endif // __SYCL_USE_FALLBACK_ASSERT
-    return submit_with_event_impl(std::move(CGF), SI, TlsCodeLocCapture.query(),
+    if constexpr (UseFallbackAssert)
+      SI.PostProcessorFunc() =
+          [this, &SecondaryQueuePtr,
+           &TlsCodeLocCapture](bool IsKernel, bool KernelUsesAssert, event &E) {
+            if (IsKernel && !device_has(aspect::ext_oneapi_native_assert) &&
+                KernelUsesAssert && !device_has(aspect::accelerator)) {
+              // __devicelib_assert_fail isn't supported by Device-side Runtime
+              // Linking against fallback impl of __devicelib_assert_fail is
+              // performed by program manager class
+              // Fallback assert isn't supported for FPGA
+              submitAssertCapture(*this, E, SecondaryQueuePtr,
+                                  TlsCodeLocCapture.query());
+            }
+          };
+    return submit_with_event_impl(CGF, SI, TlsCodeLocCapture.query(),
                                   TlsCodeLocCapture.isToplevel());
   }
 
@@ -2871,21 +2886,25 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   /// \param Props is a property list with submission properties.
   /// \param CGF is a function object containing command group.
   /// \param CodeLoc is the code location of the submit call (default argument)
-  template <typename T, typename PropertiesT>
-  std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, void>
-  submit_without_event(PropertiesT Props, T CGF,
-                       const detail::code_location &CodeLoc) {
-#if __SYCL_USE_FALLBACK_ASSERT
-    // If post-processing is needed, fall back to the regular submit.
-    // TODO: Revisit whether we can avoid this.
-    submit_with_event(Props, CGF, nullptr, CodeLoc);
-#else
-    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
-    detail::SubmissionInfo SI{};
-    ProcessSubmitProperties(Props, SI);
-    submit_without_event_impl(CGF, SI, TlsCodeLocCapture.query(),
-                              TlsCodeLocCapture.isToplevel());
-#endif // __SYCL_USE_FALLBACK_ASSERT
+  //
+  // UseFallBackAssert as template param vs `#if` in function body is necessary
+  // to prevent ODR-violation between TUs built with different fallback assert
+  // modes.
+  template <bool UseFallbackAssert, typename PropertiesT>
+  void submit_without_event(PropertiesT Props,
+                            const detail::type_erased_cgfo_ty &CGF,
+                            const detail::code_location &CodeLoc) {
+    if constexpr (UseFallbackAssert) {
+      // If post-processing is needed, fall back to the regular submit.
+      // TODO: Revisit whether we can avoid this.
+      submit_with_event<UseFallbackAssert>(Props, CGF, nullptr, CodeLoc);
+    } else {
+      detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+      detail::SubmissionInfo SI{};
+      ProcessSubmitProperties(Props, SI);
+      submit_without_event_impl(CGF, SI, TlsCodeLocCapture.query(),
+                                TlsCodeLocCapture.isToplevel());
+    }
   }
 
   /// parallel_for_impl with a kernel represented as a lambda + range that
@@ -3114,10 +3133,10 @@ event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
     });
   };
 
-  CopierEv = Self.submit_with_event(
+  CopierEv = Self.submit_with_event<true>(
       sycl::ext::oneapi::experimental::empty_properties_t{}, CopierCGF,
       SecondaryQueue, CodeLoc);
-  CheckerEv = Self.submit_with_event(
+  CheckerEv = Self.submit_with_event<true>(
       sycl::ext::oneapi::experimental::empty_properties_t{}, CheckerCGF,
       SecondaryQueue, CodeLoc);
 
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
index 8e82cc0f3082d..816c115295f87 100644
--- a/sycl/source/detail/queue_impl.cpp
+++ b/sycl/source/detail/queue_impl.cpp
@@ -349,7 +349,7 @@ void queue_impl::addSharedEvent(const event &Event) {
   MEventsShared.push_back(Event);
 }
 
-event queue_impl::submit_impl(const std::function<void(handler &)> &CGF,
+event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
                               const std::shared_ptr<queue_impl> &Self,
                               const std::shared_ptr<queue_impl> &PrimaryQueue,
                               const std::shared_ptr<queue_impl> &SecondaryQueue,
@@ -402,10 +402,13 @@ event queue_impl::submit_impl(const std::function<void(handler &)> &CGF,
     // We don't want stream flushing to be blocking operation that is why submit
     // a host task to print stream buffer. It will fire up as soon as the kernel
     // finishes execution.
-    event FlushEvent = submit_impl(
-        [&](handler &ServiceCGH) { Stream->generateFlushCommand(ServiceCGH); },
-        Self, PrimaryQueue, SecondaryQueue, /*CallerNeedsEvent*/ true, Loc,
-        IsTopCodeLoc, {});
+    auto L = [&](handler &ServiceCGH) {
+      Stream->generateFlushCommand(ServiceCGH);
+    };
+    detail::type_erased_cgfo_ty CGF{L};
+    event FlushEvent =
+        submit_impl(CGF, Self, PrimaryQueue, SecondaryQueue,
+                    /*CallerNeedsEvent*/ true, Loc, IsTopCodeLoc, {});
     EventImpl->attachEventToCompleteWeak(detail::getSyclObjImpl(FlushEvent));
     registerStreamServiceEvent(detail::getSyclObjImpl(FlushEvent));
   }
@@ -419,21 +422,19 @@ event queue_impl::submitWithHandler(const std::shared_ptr<queue_impl> &Self,
                                     bool CallerNeedsEvent,
                                     HandlerFuncT HandlerFunc) {
   SubmissionInfo SI{};
+  auto L = [&](handler &CGH) {
+    CGH.depends_on(DepEvents);
+    HandlerFunc(CGH);
+  };
+  detail::type_erased_cgfo_ty CGF{L};
+
   if (!CallerNeedsEvent && supportsDiscardingPiEvents()) {
-    submit_without_event(
-        [&](handler &CGH) {
-          CGH.depends_on(DepEvents);
-          HandlerFunc(CGH);
-        },
-        Self, SI, /*CodeLoc*/ {}, /*IsTopCodeLoc*/ true);
+    submit_without_event(CGF, Self, SI,
+                         /*CodeLoc*/ {}, /*IsTopCodeLoc*/ true);
     return createDiscardedEvent();
   }
-  return submit_with_event(
-      [&](handler &CGH) {
-        CGH.depends_on(DepEvents);
-        HandlerFunc(CGH);
-      },
-      Self, SI, /*CodeLoc*/ {}, /*IsTopCodeLoc*/ true);
+  return submit_with_event(CGF, Self, SI,
+                           /*CodeLoc*/ {}, /*IsTopCodeLoc*/ true);
 }
 
 template <typename HandlerFuncT, typename MemOpFuncT, typename... MemOpArgTs>
diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp
index 7a9128d75292c..bf18e97c50fca 100644
--- a/sycl/source/detail/queue_impl.hpp
+++ b/sycl/source/detail/queue_impl.hpp
@@ -340,7 +340,7 @@ class queue_impl {
   /// \param StoreAdditionalInfo makes additional info be stored in event_impl
   /// \return a SYCL event object, which corresponds to the queue the command
   /// group is being enqueued on.
-  event submit(const std::function<void(handler &)> &CGF,
+  event submit(const detail::type_erased_cgfo_ty &CGF,
                const std::shared_ptr<queue_impl> &Self,
                const std::shared_ptr<queue_impl> &SecondQueue,
                const detail::code_location &Loc, bool IsTopCodeLoc,
@@ -362,7 +362,7 @@ class queue_impl {
   /// \param Loc is the code location of the submit call (default argument)
   /// \param StoreAdditionalInfo makes additional info be stored in event_impl
   /// \return a SYCL event object for the submitted command group.
-  event submit_with_event(const std::function<void(handler &)> &CGF,
+  event submit_with_event(const detail::type_erased_cgfo_ty &CGF,
                           const std::shared_ptr<queue_impl> &Self,
                           const SubmissionInfo &SubmitInfo,
                           const detail::code_location &Loc, bool IsTopCodeLoc) {
@@ -387,7 +387,7 @@ class queue_impl {
     return discard_or_return(ResEvent);
   }
 
-  void submit_without_event(const std::function<void(handler &)> &CGF,
+  void submit_without_event(const detail::type_erased_cgfo_ty &CGF,
                             const std::shared_ptr<queue_impl> &Self,
                             const SubmissionInfo &SubmitInfo,
                             const detail::code_location &Loc,
@@ -855,7 +855,7 @@ class queue_impl {
   /// \param Loc is the code location of the submit call (default argument)
   /// \param SubmitInfo is additional optional information for the submission.
   /// \return a SYCL event representing submitted command group.
-  event submit_impl(const std::function<void(handler &)> &CGF,
+  event submit_impl(const detail::type_erased_cgfo_ty &CGF,
                     const std::shared_ptr<queue_impl> &Self,
                     const std::shared_ptr<queue_impl> &PrimaryQueue,
                     const std::shared_ptr<queue_impl> &SecondaryQueue,
diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp
index efd8f58cac5a3..6550677bb9854 100644
--- a/sycl/source/queue.cpp
+++ b/sycl/source/queue.cpp
@@ -255,7 +255,6 @@ event queue::submit_impl_and_postprocess(
   return impl->submit(CGH, impl, SecondQueue.impl, CodeLoc, IsTopCodeLoc,
                       &PostProcess);
 }
-#endif // __INTEL_PREVIEW_BREAKING_CHANGES
 
 event queue::submit_with_event_impl(std::function<void(handler &)> CGH,
                                     const detail::SubmissionInfo &SubmitInfo,
@@ -270,6 +269,21 @@ void queue::submit_without_event_impl(std::function<void(handler &)> CGH,
                                       bool IsTopCodeLoc) {
   impl->submit_without_event(CGH, impl, SubmitInfo, CodeLoc, IsTopCodeLoc);
 }
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
+
+event queue::submit_with_event_impl(const detail::type_erased_cgfo_ty &CGH,
+                                    const detail::SubmissionInfo &SubmitInfo,
+                                    const detail::code_location &CodeLoc,
+                                    bool IsTopCodeLoc) {
+  return impl->submit_with_event(CGH, impl, SubmitInfo, CodeLoc, IsTopCodeLoc);
+}
+
+void queue::submit_without_event_impl(const detail::type_erased_cgfo_ty &CGH,
+                                      const detail::SubmissionInfo &SubmitInfo,
+                                      const detail::code_location &CodeLoc,
+                                      bool IsTopCodeLoc) {
+  impl->submit_without_event(CGH, impl, SubmitInfo, CodeLoc, IsTopCodeLoc);
+}
 
 void queue::wait_proxy(const detail::code_location &CodeLoc) {
   impl->wait(CodeLoc);
diff --git a/sycl/test-e2e/Basic/submit_fn_ptr.cpp b/sycl/test-e2e/Basic/submit_fn_ptr.cpp
new file mode 100644
index 0000000000000..8b8d6e6ddcefe
--- /dev/null
+++ b/sycl/test-e2e/Basic/submit_fn_ptr.cpp
@@ -0,0 +1,22 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+int *p = nullptr;
+
+void foo(sycl::handler &cgh) {
+  auto *copy = p;
+  cgh.single_task([=]() { *copy = 42; });
+}
+
+int main() {
+  sycl::queue q;
+  p = sycl::malloc_shared<int>(1, q);
+  *p = 0;
+  q.submit(foo).wait();
+  assert(*p == 42);
+  sycl::free(p, q);
+  return 0;
+}
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index a8b45ebd93a4f..98dee24572890 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3137,9 +3137,11 @@ _ZN4sycl3_V15queue18throw_asynchronousEv
 _ZN4sycl3_V15queue20memcpyToDeviceGlobalEPvPKvbmmRKSt6vectorINS0_5eventESaIS6_EE
 _ZN4sycl3_V15queue20wait_and_throw_proxyERKNS0_6detail13code_locationE
 _ZN4sycl3_V15queue22memcpyFromDeviceGlobalEPvPKvbmmRKSt6vectorINS0_5eventESaIS6_EE
+_ZN4sycl3_V15queue22submit_with_event_implERKNS0_6detail19type_erased_cgfo_tyERKNS2_14SubmissionInfoERKNS2_13code_locationEb
 _ZN4sycl3_V15queue22submit_with_event_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail14SubmissionInfoERKNS7_13code_locationEb
 _ZN4sycl3_V15queue25ext_oneapi_submit_barrierERKNS0_6detail13code_locationE
 _ZN4sycl3_V15queue25ext_oneapi_submit_barrierERKSt6vectorINS0_5eventESaIS3_EERKNS0_6detail13code_locationE
+_ZN4sycl3_V15queue25submit_without_event_implERKNS0_6detail19type_erased_cgfo_tyERKNS2_14SubmissionInfoERKNS2_13code_locationEb
 _ZN4sycl3_V15queue25submit_without_event_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code_locationE
 _ZN4sycl3_V15queue25submit_without_event_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code_locationEb
 _ZN4sycl3_V15queue25submit_without_event_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail14SubmissionInfoERKNS7_13code_locationEb
@@ -3524,6 +3526,7 @@ _ZN4sycl3_V17handler20associateWithHandlerEPNS0_6detail30UnsampledImageAccessorB
 _ZN4sycl3_V17handler20memcpyToDeviceGlobalEPKvS3_bmm
 _ZN4sycl3_V17handler20setKernelCacheConfigENS1_23StableKernelCacheConfigE
 _ZN4sycl3_V17handler20setStateSpecConstSetEv
+_ZN4sycl3_V17handler21setKernelWorkGroupMemEm
 _ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE
 _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm
 _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm
@@ -3531,7 +3534,6 @@ _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE
 _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE
 _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi3EEEi
 _ZN4sycl3_V17handler22setKernelIsCooperativeEb
-_ZN4sycl3_V17handler21setKernelWorkGroupMemEm
 _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_
 _ZN4sycl3_V17handler24ext_intel_read_host_pipeENS0_6detail11string_viewEPvmb
 _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index 1519f0b0bf3a7..edc11b37a071f 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -4305,7 +4305,9 @@
 ?submit_impl_and_postprocess@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@AEBV?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@6@_N@Z
 ?submit_impl_and_postprocess@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V123@AEBUcode_location@detail@23@AEBV?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@6@@Z
 ?submit_impl_and_postprocess@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V123@AEBUcode_location@detail@23@AEBV?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@6@_N@Z
+?submit_with_event_impl@queue@_V1@sycl@@AEAA?AVevent@23@AEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@623@AEBUcode_location@623@_N@Z
 ?submit_with_event_impl@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBVSubmissionInfo@detail@23@AEBUcode_location@823@_N@Z
+?submit_without_event_impl@queue@_V1@sycl@@AEAAXAEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@523@AEBUcode_location@523@_N@Z
 ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@@Z
 ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@_N@Z
 ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBVSubmissionInfo@detail@23@AEBUcode_location@723@_N@Z
diff --git a/sycl/unittests/scheduler/HostTaskAndBarrier.cpp b/sycl/unittests/scheduler/HostTaskAndBarrier.cpp
index 2b07861ef28fd..9dbd9e0b34346 100644
--- a/sycl/unittests/scheduler/HostTaskAndBarrier.cpp
+++ b/sycl/unittests/scheduler/HostTaskAndBarrier.cpp
@@ -59,28 +59,28 @@ class BarrierHandlingWithHostTask : public ::testing::Test {
 
   sycl::event AddTask(TestCGType Type, bool BlockHostTask = true) {
     if (Type == TestCGType::HOST_TASK) {
-      return QueueDevImpl->submit(
-          [&](handler &CGH) {
-            CGH.host_task(BlockHostTask ? CustomHostLambda : [] {});
-          },
-          QueueDevImpl, nullptr, {}, true);
+      auto L = [&](handler &CGH) {
+        CGH.host_task(BlockHostTask ? CustomHostLambda : [] {});
+      };
+      return QueueDevImpl->submit(sycl::detail::type_erased_cgfo_ty{L},
+                                  QueueDevImpl, nullptr, {}, true);
     } else if (Type == TestCGType::KERNEL_TASK) {
-      return QueueDevImpl->submit(
-          [&](handler &CGH) { CGH.single_task<TestKernel<>>([] {}); },
-          QueueDevImpl, nullptr, {}, true);
+      auto L = [&](handler &CGH) { CGH.single_task<TestKernel<>>([] {}); };
+      return QueueDevImpl->submit(sycl::detail::type_erased_cgfo_ty{L},
+                                  QueueDevImpl, nullptr, {}, true);
     } else // (Type == TestCGType::BARRIER)
     {
-      return QueueDevImpl->submit(
-          [&](handler &CGH) { CGH.ext_oneapi_barrier(); }, QueueDevImpl,
-          nullptr, {}, true);
+      auto L = [&](handler &CGH) { CGH.ext_oneapi_barrier(); };
+      return QueueDevImpl->submit(sycl::detail::type_erased_cgfo_ty{L},
+                                  QueueDevImpl, nullptr, {}, true);
     }
   }
 
   sycl::event
   InsertBarrierWithWaitList(const std::vector<sycl::event> &WaitList) {
-    return QueueDevImpl->submit(
-        [&](handler &CGH) { CGH.ext_oneapi_barrier(WaitList); }, QueueDevImpl,
-        nullptr, {}, true);
+    auto L = [&](handler &CGH) { CGH.ext_oneapi_barrier(WaitList); };
+    return QueueDevImpl->submit(sycl::detail::type_erased_cgfo_ty{L},
+                                QueueDevImpl, nullptr, {}, true);
   }
 
   void BuildAndCheckInnerQueueState(std::vector<EventImplPtr> &Events) {