diff --git a/sycl/include/sycl/detail/helpers.hpp b/sycl/include/sycl/detail/helpers.hpp
index e4e3ad12bb8d3..1576ea7ed1679 100644
--- a/sycl/include/sycl/detail/helpers.hpp
+++ b/sycl/include/sycl/detail/helpers.hpp
@@ -38,14 +38,11 @@ template <typename Type, std::size_t NumElements> class marray;
 enum class memory_order;
 
 namespace detail {
-class CGExecKernel;
 class buffer_impl;
-class context_impl;
-class queue_impl;
-using QueueImplPtr = std::shared_ptr<sycl::detail::queue_impl>;
-class RTDeviceBinaryImage;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 __SYCL_EXPORT void waitEvents(std::vector<sycl::event> DepEvents);
+#endif
 
 __SYCL_EXPORT void
 markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl);
@@ -251,10 +248,6 @@ template <size_t count, class F> void loop(F &&f) {
   loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
 }
 inline constexpr bool is_power_of_two(int x) { return (x & (x - 1)) == 0; }
-
-std::tuple<const RTDeviceBinaryImage *, ur_program_handle_t>
-retrieveKernelBinary(const QueueImplPtr &, const char *KernelName,
-                     CGExecKernel *CGKernel = nullptr);
 } // namespace detail
 
 } // namespace _V1
diff --git a/sycl/include/sycl/memory_enums.hpp b/sycl/include/sycl/memory_enums.hpp
index 9ae997896f5e8..3a8886034c560 100644
--- a/sycl/include/sycl/memory_enums.hpp
+++ b/sycl/include/sycl/memory_enums.hpp
@@ -8,8 +8,6 @@
 
 #pragma once
 
-#include <ur_api.h> // for ur_memory_order_capability_flags_t
-
 #include <atomic> // for memory_order
 #include <vector> // for vector
 
@@ -48,38 +46,6 @@ inline constexpr auto memory_order_seq_cst = memory_order::seq_cst;
 
 namespace detail {
 
-inline std::vector<memory_order>
-readMemoryOrderBitfield(ur_memory_order_capability_flags_t bits) {
-  std::vector<memory_order> result;
-  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_RELAXED)
-    result.push_back(memory_order::relaxed);
-  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_ACQUIRE)
-    result.push_back(memory_order::acquire);
-  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_RELEASE)
-    result.push_back(memory_order::release);
-  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_ACQ_REL)
-    result.push_back(memory_order::acq_rel);
-  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_SEQ_CST)
-    result.push_back(memory_order::seq_cst);
-  return result;
-}
-
-inline std::vector<memory_scope>
-readMemoryScopeBitfield(ur_memory_scope_capability_flags_t bits) {
-  std::vector<memory_scope> result;
-  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_ITEM)
-    result.push_back(memory_scope::work_item);
-  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_SUB_GROUP)
-    result.push_back(memory_scope::sub_group);
-  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_GROUP)
-    result.push_back(memory_scope::work_group);
-  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_DEVICE)
-    result.push_back(memory_scope::device);
-  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_SYSTEM)
-    result.push_back(memory_scope::system);
-  return result;
-}
-
 #ifndef __SYCL_DEVICE_ONLY__
 static constexpr std::memory_order getStdMemoryOrder(sycl::memory_order order) {
   switch (order) {
diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp
index 11daf34ae14c8..b84bfff755abb 100644
--- a/sycl/source/detail/device_info.hpp
+++ b/sycl/source/detail/device_info.hpp
@@ -23,6 +23,7 @@
 #include <sycl/info/info_desc.hpp>
 #include <sycl/memory_enums.hpp>
 #include <sycl/platform.hpp>
+#include <ur_api.h> // for ur_memory_order_capability_flags_t
 
 #include <chrono>
 #include <sstream>
@@ -34,6 +35,38 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
+inline std::vector<memory_order>
+readMemoryOrderBitfield(ur_memory_order_capability_flags_t bits) {
+  std::vector<memory_order> result;
+  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_RELAXED)
+    result.push_back(memory_order::relaxed);
+  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_ACQUIRE)
+    result.push_back(memory_order::acquire);
+  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_RELEASE)
+    result.push_back(memory_order::release);
+  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_ACQ_REL)
+    result.push_back(memory_order::acq_rel);
+  if (bits & UR_MEMORY_ORDER_CAPABILITY_FLAG_SEQ_CST)
+    result.push_back(memory_order::seq_cst);
+  return result;
+}
+
+inline std::vector<memory_scope>
+readMemoryScopeBitfield(ur_memory_scope_capability_flags_t bits) {
+  std::vector<memory_scope> result;
+  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_ITEM)
+    result.push_back(memory_scope::work_item);
+  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_SUB_GROUP)
+    result.push_back(memory_scope::sub_group);
+  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_GROUP)
+    result.push_back(memory_scope::work_group);
+  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_DEVICE)
+    result.push_back(memory_scope::device);
+  if (bits & UR_MEMORY_SCOPE_CAPABILITY_FLAG_SYSTEM)
+    result.push_back(memory_scope::system);
+  return result;
+}
+
 inline std::vector<info::fp_config>
 read_fp_bitfield(ur_device_fp_capability_flags_t bits) {
   std::vector<info::fp_config> result;
diff --git a/sycl/source/detail/helpers.cpp b/sycl/source/detail/helpers.cpp
index 59f4c42b66040..8a6583d6dd741 100644
--- a/sycl/source/detail/helpers.cpp
+++ b/sycl/source/detail/helpers.cpp
@@ -6,6 +6,8 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <detail/helpers.hpp>
+
 #include <detail/scheduler/commands.hpp>
 #include <sycl/detail/helpers.hpp>
 
@@ -30,7 +32,8 @@ void waitEvents(std::vector<sycl::event> DepEvents) {
   }
 }
 
-void markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl) {
+__SYCL_EXPORT void
+markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl) {
   BufImpl->markAsInternal();
 }
 
diff --git a/sycl/source/detail/helpers.hpp b/sycl/source/detail/helpers.hpp
new file mode 100644
index 0000000000000..57287959b7575
--- /dev/null
+++ b/sycl/source/detail/helpers.hpp
@@ -0,0 +1,34 @@
+//==---------------- helpers.cpp - SYCL helpers ---------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <ur_api.h>
+
+#include <memory>
+#include <tuple>
+#include <vector>
+
+namespace sycl {
+inline namespace _V1 {
+class event;
+
+namespace detail {
+class CGExecKernel;
+class queue_impl;
+using QueueImplPtr = std::shared_ptr<sycl::detail::queue_impl>;
+class RTDeviceBinaryImage;
+
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+void waitEvents(std::vector<sycl::event> DepEvents);
+#endif
+
+std::tuple<const RTDeviceBinaryImage *, ur_program_handle_t>
+retrieveKernelBinary(const QueueImplPtr &, const char *KernelName,
+                     CGExecKernel *CGKernel = nullptr);
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp
index 8abb9ddabfd7d..c48ba4c15b802 100644
--- a/sycl/source/detail/jit_compiler.cpp
+++ b/sycl/source/detail/jit_compiler.cpp
@@ -9,6 +9,7 @@
 #if SYCL_EXT_JIT_ENABLE
 #include <KernelFusion.h>
 #include <detail/device_image_impl.hpp>
+#include <detail/helpers.hpp>
 #include <detail/jit_compiler.hpp>
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_impl.hpp>
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 133f55992f356..acc61ab39d77d 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -10,6 +10,7 @@
 
 #include <detail/context_impl.hpp>
 #include <detail/event_impl.hpp>
+#include <detail/helpers.hpp>
 #include <detail/host_pipe_map_entry.hpp>
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_impl.hpp>
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index a7b66235b6653..3cf52728f4d06 100644
--- a/sycl/source/handler.cpp
+++ b/sycl/source/handler.cpp
@@ -14,6 +14,7 @@
 #include <detail/global_handler.hpp>
 #include <detail/graph_impl.hpp>
 #include <detail/handler_impl.hpp>
+#include <detail/helpers.hpp>
 #include <detail/host_task.hpp>
 #include <detail/image_impl.hpp>
 #include <detail/kernel_bundle_impl.hpp>
diff --git a/sycl/test/include_deps/sycl_accessor.hpp.cpp b/sycl/test/include_deps/sycl_accessor.hpp.cpp
index 5e51278562b34..ae896a3987e85 100644
--- a/sycl/test/include_deps/sycl_accessor.hpp.cpp
+++ b/sycl/test/include_deps/sycl_accessor.hpp.cpp
@@ -14,7 +14,6 @@
 // CHECK-NEXT: atomic.hpp
 // CHECK-NEXT: detail/helpers.hpp
 // CHECK-NEXT: memory_enums.hpp
-// CHECK-NEXT: ur_api.h
 // CHECK-NEXT: CL/__spirv/spirv_vars.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: aliases.hpp
@@ -66,6 +65,7 @@
 // CHECK-NEXT: CL/cl_platform.h
 // CHECK-NEXT: CL/cl_ext.h
 // CHECK-NEXT: detail/string.hpp
+// CHECK-NEXT: ur_api.h
 // CHECK-NEXT: detail/common.hpp
 // CHECK-NEXT: detail/is_device_copyable.hpp
 // CHECK-NEXT: detail/owner_less_base.hpp
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
index ad97f74c140c2..6ed467969333d 100644
--- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp
+++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
@@ -15,7 +15,6 @@
 // CHECK-NEXT: atomic.hpp
 // CHECK-NEXT: detail/helpers.hpp
 // CHECK-NEXT: memory_enums.hpp
-// CHECK-NEXT: ur_api.h
 // CHECK-NEXT: CL/__spirv/spirv_vars.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: aliases.hpp
@@ -67,6 +66,7 @@
 // CHECK-NEXT: CL/cl_platform.h
 // CHECK-NEXT: CL/cl_ext.h
 // CHECK-NEXT: detail/string.hpp
+// CHECK-NEXT: ur_api.h
 // CHECK-NEXT: detail/common.hpp
 // CHECK-NEXT: detail/is_device_copyable.hpp
 // CHECK-NEXT: detail/owner_less_base.hpp