diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp
index 2ffc49a4868b6..3dd5ed5bd161e 100644
--- a/sycl/include/sycl/detail/spirv.hpp
+++ b/sycl/include/sycl/detail/spirv.hpp
@@ -1225,22 +1225,26 @@ EnableIfGenericShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
 template <typename Group>
 typename std::enable_if_t<
     ext::oneapi::experimental::is_fixed_topology_group_v<Group>>
-ControlBarrier(Group, memory_scope FenceScope, memory_order Order) {
+ControlBarrier(Group, [[maybe_unused]] memory_scope FenceScope,
+               [[maybe_unused]] memory_order Order) {
+#ifdef __SYCL_DEVICE_ONLY__
   __spirv_ControlBarrier(group_scope<Group>::value, getScope(FenceScope),
                          getMemorySemanticsMask(Order) |
                              __spv::MemorySemanticsMask::SubgroupMemory |
                              __spv::MemorySemanticsMask::WorkgroupMemory |
                              __spv::MemorySemanticsMask::CrossWorkgroupMemory);
+#endif
 }
 
 template <typename Group>
 typename std::enable_if_t<
     ext::oneapi::experimental::is_user_constructed_group_v<Group>>
-ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
+ControlBarrier([[maybe_unused]] Group g,
+               [[maybe_unused]] memory_scope FenceScope,
+               [[maybe_unused]] memory_order Order) {
 #if defined(__NVPTX__)
   __nvvm_bar_warp_sync(detail::ExtractMask(detail::GetMask(g))[0]);
-#else
-  (void)g;
+#elif defined(__SYCL_DEVICE_ONLY__)
   // SPIR-V does not define an instruction to synchronize partial groups.
   // However, most (possibly all?) of the current SPIR-V targets execute
   // work-items in lockstep, so we can probably get away with a MemoryBarrier.
diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt
index 6035cbe030fa2..7b70618f76e23 100644
--- a/sycl/source/CMakeLists.txt
+++ b/sycl/source/CMakeLists.txt
@@ -330,7 +330,6 @@ set(SYCL_COMMON_SOURCES
     "queue.cpp"
     "sampler.cpp"
     "stream.cpp"
-    "spirv_ops.cpp"
     "virtual_mem.cpp"
     "detail/memory_pool_impl.cpp"
     "detail/async_alloc.cpp"
@@ -341,6 +340,7 @@ set(SYCL_COMMON_SOURCES
 )
 
 set(SYCL_NON_PREVIEW_SOURCES "${SYCL_COMMON_SOURCES}"
+  "spirv_ops.cpp"
 )
 
 
diff --git a/sycl/source/detail/platform_util.cpp b/sycl/source/detail/platform_util.cpp
index 68692ce14b96b..59d0f845b7ad9 100644
--- a/sycl/source/detail/platform_util.cpp
+++ b/sycl/source/detail/platform_util.cpp
@@ -131,26 +131,6 @@ uint32_t PlatformUtil::getNativeVectorWidth(PlatformUtil::TypeIndex TIndex) {
   return 0;
 }
 
-void PlatformUtil::prefetch(const char *Ptr, size_t NumBytes) {
-  if (!Ptr)
-    return;
-
-  const size_t CacheLineSize = PlatformUtil::getMemCacheLineSize();
-  const size_t CacheLineMask = ~(CacheLineSize - 1);
-  const char *PtrEnd = Ptr + NumBytes;
-
-  // Set the pointer to the beginning of the current cache line.
-  Ptr = reinterpret_cast<const char *>(reinterpret_cast<size_t>(Ptr) &
-                                       CacheLineMask);
-  for (; Ptr < PtrEnd; Ptr += CacheLineSize) {
-#if defined(__SYCL_RT_OS_LINUX)
-    __builtin_prefetch(Ptr);
-#elif defined(__SYCL_RT_OS_WINDOWS)
-    _mm_prefetch(Ptr, _MM_HINT_T0);
-#endif
-  }
-}
-
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/source/detail/platform_util.hpp b/sycl/source/detail/platform_util.hpp
index cd9353198d08a..85d88b4f63e4e 100644
--- a/sycl/source/detail/platform_util.hpp
+++ b/sycl/source/detail/platform_util.hpp
@@ -42,8 +42,6 @@ struct PlatformUtil {
   static uint32_t getMemCacheLineSize();
 
   static uint64_t getMemCacheSize();
-
-  static void prefetch(const char *Ptr, size_t NumBytes);
 };
 
 } // namespace detail
diff --git a/sycl/source/ld-version-script.txt b/sycl/source/ld-version-script.txt
index 61e2e6a874406..910e1eec1e6c5 100644
--- a/sycl/source/ld-version-script.txt
+++ b/sycl/source/ld-version-script.txt
@@ -18,10 +18,13 @@
     _ZN10__host_std*;
 
     /* Export SPIR-V built-ins for host device */
+    /* #ifndef __INTEL_PREVIEW_BREAKING_CHANGES */
+    /* TODO: drop those in the next ABI-breaking window */
     _Z23__spirv_GroupWaitEvents*;
     _Z22__spirv_ControlBarrier*;
     _Z21__spirv_MemoryBarrier*;
     _Z20__spirv_ocl_prefetch*;
+    /* #endif // __INTEL_PREVIEW_BREAKING_CHANGES */
 
     /* Export offload image hooks */
     __sycl_register_lib;
diff --git a/sycl/source/spirv_ops.cpp b/sycl/source/spirv_ops.cpp
index fdaa7e1834eae..231205b899389 100644
--- a/sycl/source/spirv_ops.cpp
+++ b/sycl/source/spirv_ops.cpp
@@ -45,7 +45,4 @@ __SYCL_EXPORT void __spirv_MemoryBarrier(__spv::Scope Memory,
   atomic_thread_fence(std::memory_order_seq_cst);
 }
 
-__SYCL_EXPORT void __spirv_ocl_prefetch(const char *Ptr,
-                                        size_t NumBytes) noexcept {
-  sycl::detail::PlatformUtil::prefetch(Ptr, NumBytes);
-}
+__SYCL_EXPORT void __spirv_ocl_prefetch(const char *, size_t) noexcept {}