diff --git a/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp b/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp
index 2c9b999309731..8d6aa5ec432c1 100644
--- a/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp
@@ -94,16 +94,16 @@ struct bit_equal<T, typename detail::enable_if_t<std::is_integral<T>::value>> {
 
 template <> struct bit_equal<float> {
   bool operator()(const float &lhs, const float &rhs) {
-    auto LhsInt = detail::bit_cast<uint32_t>(lhs);
-    auto RhsInt = detail::bit_cast<uint32_t>(rhs);
+    auto LhsInt = sycl::bit_cast<uint32_t>(lhs);
+    auto RhsInt = sycl::bit_cast<uint32_t>(rhs);
     return LhsInt == RhsInt;
   }
 };
 
 template <> struct bit_equal<double> {
   bool operator()(const double &lhs, const double &rhs) {
-    auto LhsInt = detail::bit_cast<uint64_t>(lhs);
-    auto RhsInt = detail::bit_cast<uint64_t>(rhs);
+    auto LhsInt = sycl::bit_cast<uint64_t>(lhs);
+    auto RhsInt = sycl::bit_cast<uint64_t>(rhs);
     return LhsInt == RhsInt;
   }
 };
diff --git a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp
index 121d94a5141a2..11d09c114bf81 100644
--- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp
@@ -55,7 +55,7 @@ T load(const multi_ptr<T, Space> src) {
   BlockT Ret =
       __spirv_SubgroupBlockReadINTEL<BlockT>(reinterpret_cast<PtrT>(src.get()));
 
-  return sycl::detail::bit_cast<T>(Ret);
+  return sycl::bit_cast<T>(Ret);
 }
 
 template <int N, typename T, access::address_space Space>
@@ -68,7 +68,7 @@ vec<T, N> load(const multi_ptr<T, Space> src) {
   VecT Ret =
       __spirv_SubgroupBlockReadINTEL<VecT>(reinterpret_cast<PtrT>(src.get()));
 
-  return sycl::detail::bit_cast<typename vec<T, N>::vector_t>(Ret);
+  return sycl::bit_cast<typename vec<T, N>::vector_t>(Ret);
 }
 
 template <typename T, access::address_space Space>
@@ -77,7 +77,7 @@ void store(multi_ptr<T, Space> dst, const T &x) {
   using PtrT = sycl::detail::ConvertToOpenCLType_t<multi_ptr<BlockT, Space>>;
 
   __spirv_SubgroupBlockWriteINTEL(reinterpret_cast<PtrT>(dst.get()),
-                                  sycl::detail::bit_cast<BlockT>(x));
+                                  sycl::bit_cast<BlockT>(x));
 }
 
 template <int N, typename T, access::address_space Space>
@@ -88,7 +88,7 @@ void store(multi_ptr<T, Space> dst, const vec<T, N> &x) {
       sycl::detail::ConvertToOpenCLType_t<const multi_ptr<BlockT, Space>>;
 
   __spirv_SubgroupBlockWriteINTEL(reinterpret_cast<PtrT>(dst.get()),
-                                  sycl::detail::bit_cast<VecT>(x));
+                                  sycl::bit_cast<VecT>(x));
 }
 #endif // __SYCL_DEVICE_ONLY__
 
diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp
index 6733e9666ab86..cbf615f80bb54 100644
--- a/sycl/include/CL/sycl/atomic.hpp
+++ b/sycl/include/CL/sycl/atomic.hpp
@@ -237,7 +237,7 @@ class atomic {
             Ptr);
     cl_int TmpVal = __spirv_AtomicLoad(
         TmpPtr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order));
-    cl_float ResVal = detail::bit_cast<cl_float>(TmpVal);
+    cl_float ResVal = bit_cast<cl_float>(TmpVal);
     return ResVal;
   }
 #else
diff --git a/sycl/include/CL/sycl/bit_cast.hpp b/sycl/include/CL/sycl/bit_cast.hpp
new file mode 100644
index 0000000000000..2a042c92811bd
--- /dev/null
+++ b/sycl/include/CL/sycl/bit_cast.hpp
@@ -0,0 +1,67 @@
+//==---------------- bit_cast.hpp - SYCL bit_cast --------------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <type_traits>
+
+#if __cpp_lib_bit_cast
+#include <bit>
+#endif
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+
+// forward decl
+namespace detail {
+inline void memcpy(void *Dst, const void *Src, std::size_t Size);
+}
+
+// sycl::bit_cast ( no longer sycl::detail::bit_cast )
+template <typename To, typename From>
+#if __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast)
+constexpr
+#endif
+    To
+    bit_cast(const From &from) noexcept {
+  static_assert(sizeof(To) == sizeof(From),
+                "Sizes of To and From must be equal");
+  static_assert(std::is_trivially_copyable<From>::value,
+                "From must be trivially copyable");
+  static_assert(std::is_trivially_copyable<To>::value,
+                "To must be trivially copyable");
+#if __cpp_lib_bit_cast
+  return std::bit_cast<To>(from);
+#else // __cpp_lib_bit_cast
+
+#if __has_builtin(__builtin_bit_cast)
+  return __builtin_bit_cast(To, from);
+#else  // __has_builtin(__builtin_bit_cast)
+  static_assert(std::is_trivially_default_constructible<To>::value,
+                "To must be trivially default constructible");
+  To to;
+  sycl::detail::memcpy(&to, &from, sizeof(To));
+  return to;
+#endif // __has_builtin(__builtin_bit_cast)
+
+#endif // __cpp_lib_bit_cast
+}
+
+namespace detail {
+template <typename To, typename From>
+#if __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast)
+constexpr
+#endif
+    To
+    bit_cast(const From &from) noexcept {
+  return sycl::bit_cast<To>(from);
+}
+} // namespace detail
+
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/detail/helpers.hpp b/sycl/include/CL/sycl/detail/helpers.hpp
index 652719abd107d..118271a35bab5 100644
--- a/sycl/include/CL/sycl/detail/helpers.hpp
+++ b/sycl/include/CL/sycl/detail/helpers.hpp
@@ -16,9 +16,6 @@
 #include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 
-#if __cpp_lib_bit_cast
-#include <bit>
-#endif
 #include <memory>
 #include <stdexcept>
 #include <type_traits>
@@ -45,34 +42,6 @@ inline void memcpy(void *Dst, const void *Src, size_t Size) {
   }
 }
 
-template <typename To, typename From>
-#if __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast)
-constexpr
-#endif
-    To
-    bit_cast(const From &from) noexcept {
-  static_assert(sizeof(To) == sizeof(From),
-                "Sizes of To and From must be equal");
-  static_assert(std::is_trivially_copyable<From>::value,
-                "From must be trivially copyable");
-  static_assert(std::is_trivially_copyable<To>::value,
-                "To must be trivially copyable");
-#if __cpp_lib_bit_cast
-  return std::bit_cast<To>(from);
-#else // __cpp_lib_bit_cast
-
-#if __has_builtin(__builtin_bit_cast)
-  return __builtin_bit_cast(To, from);
-#else  // __has_builtin(__builtin_bit_cast)
-  static_assert(std::is_trivially_default_constructible<To>::value,
-                "To must be trivially default constructible");
-  To to;
-  sycl::detail::memcpy(&to, &from, sizeof(To));
-  return to;
-#endif // __has_builtin(__builtin_bit_cast)
-
-#endif // __cpp_lib_bit_cast
-}
 
 class context_impl;
 // The function returns list of events that can be passed to OpenCL API as
@@ -272,5 +241,6 @@ getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace,
 }
 
 } // namespace detail
+
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp
index ccf2cf3863304..cc89053783485 100644
--- a/sycl/include/CL/sycl/detail/spirv.hpp
+++ b/sycl/include/CL/sycl/detail/spirv.hpp
@@ -140,11 +140,11 @@ EnableIfBitcastBroadcast<T, IdT> GroupBroadcast(T x, IdT local_id) {
   GroupIdT GroupLocalId = static_cast<GroupIdT>(local_id);
   using BroadcastT = ConvertToNativeBroadcastType_t<T>;
   using OCLIdT = detail::ConvertToOpenCLType_t<GroupIdT>;
-  auto BroadcastX = detail::bit_cast<BroadcastT>(x);
+  auto BroadcastX = bit_cast<BroadcastT>(x);
   OCLIdT OCLId = detail::convertDataToType<GroupIdT, OCLIdT>(GroupLocalId);
   BroadcastT Result =
       __spirv_GroupBroadcast(group_scope<Group>::value, BroadcastX, OCLId);
-  return detail::bit_cast<T>(Result);
+  return bit_cast<T>(Result);
 }
 template <typename Group, typename T, typename IdT>
 EnableIfGenericBroadcast<T, IdT> GroupBroadcast(T x, IdT local_id) {
@@ -190,11 +190,11 @@ EnableIfBitcastBroadcast<T> GroupBroadcast(T x, id<Dimensions> local_id) {
   for (int i = 0; i < Dimensions; ++i) {
     VecId[i] = local_id[Dimensions - i - 1];
   }
-  auto BroadcastX = detail::bit_cast<BroadcastT>(x);
+  auto BroadcastX = bit_cast<BroadcastT>(x);
   OCLIdT OCLId = detail::convertDataToType<IdT, OCLIdT>(VecId);
   BroadcastT Result =
       __spirv_GroupBroadcast(group_scope<Group>::value, BroadcastX, OCLId);
-  return detail::bit_cast<T>(Result);
+  return bit_cast<T>(Result);
 }
 template <typename Group, typename T, int Dimensions>
 EnableIfGenericBroadcast<T> GroupBroadcast(T x, id<Dimensions> local_id) {
@@ -284,11 +284,11 @@ AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
   auto *PtrInt =
       reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
           MPtr.get());
-  I DesiredInt = detail::bit_cast<I>(Desired);
-  I ExpectedInt = detail::bit_cast<I>(Expected);
+  I DesiredInt = bit_cast<I>(Desired);
+  I ExpectedInt = bit_cast<I>(Expected);
   I ResultInt = __spirv_AtomicCompareExchange(
       PtrInt, SPIRVScope, SPIRVSuccess, SPIRVFailure, DesiredInt, ExpectedInt);
-  return detail::bit_cast<T>(ResultInt);
+  return bit_cast<T>(ResultInt);
 }
 
 template <typename T, access::address_space AddressSpace>
@@ -312,7 +312,7 @@ AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ONEAPI::memory_scope Scope,
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
   I ResultInt = __spirv_AtomicLoad(PtrInt, SPIRVScope, SPIRVOrder);
-  return detail::bit_cast<T>(ResultInt);
+  return bit_cast<T>(ResultInt);
 }
 
 template <typename T, access::address_space AddressSpace>
@@ -335,7 +335,7 @@ AtomicStore(multi_ptr<T, AddressSpace> MPtr, ONEAPI::memory_scope Scope,
           MPtr.get());
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
-  I ValueInt = detail::bit_cast<I>(Value);
+  I ValueInt = bit_cast<I>(Value);
   __spirv_AtomicStore(PtrInt, SPIRVScope, SPIRVOrder, ValueInt);
 }
 
@@ -359,10 +359,10 @@ AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ONEAPI::memory_scope Scope,
           MPtr.get());
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
-  I ValueInt = detail::bit_cast<I>(Value);
+  I ValueInt = bit_cast<I>(Value);
   I ResultInt =
       __spirv_AtomicExchange(PtrInt, SPIRVScope, SPIRVOrder, ValueInt);
-  return detail::bit_cast<T>(ResultInt);
+  return bit_cast<T>(ResultInt);
 }
 
 template <typename T, access::address_space AddressSpace>
@@ -600,7 +600,7 @@ using ConvertToNativeShuffleType_t = select_cl_scalar_integral_unsigned_t<T>;
 template <typename T>
 EnableIfBitcastShuffle<T> SubgroupShuffle(T x, id<1> local_id) {
   using ShuffleT = ConvertToNativeShuffleType_t<T>;
-  auto ShuffleX = detail::bit_cast<ShuffleT>(x);
+  auto ShuffleX = bit_cast<ShuffleT>(x);
 #ifndef __NVPTX__
   ShuffleT Result = __spirv_SubgroupShuffleINTEL(
       ShuffleX, static_cast<uint32_t>(local_id.get(0)));
@@ -608,13 +608,13 @@ EnableIfBitcastShuffle<T> SubgroupShuffle(T x, id<1> local_id) {
   ShuffleT Result =
       __nvvm_shfl_sync_idx_i32(membermask(), ShuffleX, local_id.get(0), 0x1f);
 #endif
-  return detail::bit_cast<T>(Result);
+  return bit_cast<T>(Result);
 }
 
 template <typename T>
 EnableIfBitcastShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
   using ShuffleT = ConvertToNativeShuffleType_t<T>;
-  auto ShuffleX = detail::bit_cast<ShuffleT>(x);
+  auto ShuffleX = bit_cast<ShuffleT>(x);
 #ifndef __NVPTX__
   ShuffleT Result = __spirv_SubgroupShuffleXorINTEL(
       ShuffleX, static_cast<uint32_t>(local_id.get(0)));
@@ -622,13 +622,13 @@ EnableIfBitcastShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
   ShuffleT Result =
       __nvvm_shfl_sync_bfly_i32(membermask(), ShuffleX, local_id.get(0), 0x1f);
 #endif
-  return detail::bit_cast<T>(Result);
+  return bit_cast<T>(Result);
 }
 
 template <typename T>
 EnableIfBitcastShuffle<T> SubgroupShuffleDown(T x, id<1> local_id) {
   using ShuffleT = ConvertToNativeShuffleType_t<T>;
-  auto ShuffleX = detail::bit_cast<ShuffleT>(x);
+  auto ShuffleX = bit_cast<ShuffleT>(x);
 #ifndef __NVPTX__
   ShuffleT Result = __spirv_SubgroupShuffleDownINTEL(
       ShuffleX, ShuffleX, static_cast<uint32_t>(local_id.get(0)));
@@ -636,13 +636,13 @@ EnableIfBitcastShuffle<T> SubgroupShuffleDown(T x, id<1> local_id) {
   ShuffleT Result =
       __nvvm_shfl_sync_down_i32(membermask(), ShuffleX, local_id.get(0), 0x1f);
 #endif
-  return detail::bit_cast<T>(Result);
+  return bit_cast<T>(Result);
 }
 
 template <typename T>
 EnableIfBitcastShuffle<T> SubgroupShuffleUp(T x, id<1> local_id) {
   using ShuffleT = ConvertToNativeShuffleType_t<T>;
-  auto ShuffleX = detail::bit_cast<ShuffleT>(x);
+  auto ShuffleX = bit_cast<ShuffleT>(x);
 #ifndef __NVPTX__
   ShuffleT Result = __spirv_SubgroupShuffleUpINTEL(
       ShuffleX, ShuffleX, static_cast<uint32_t>(local_id.get(0)));
@@ -650,7 +650,7 @@ EnableIfBitcastShuffle<T> SubgroupShuffleUp(T x, id<1> local_id) {
   ShuffleT Result =
       __nvvm_shfl_sync_up_i32(membermask(), ShuffleX, local_id.get(0), 0);
 #endif
-  return detail::bit_cast<T>(Result);
+  return bit_cast<T>(Result);
 }
 
 // Generic shuffles may require multiple calls to SubgroupShuffle
diff --git a/sycl/include/CL/sycl/stl.hpp b/sycl/include/CL/sycl/stl.hpp
index 6a9b0b844d838..5b002aff3be0b 100644
--- a/sycl/include/CL/sycl/stl.hpp
+++ b/sycl/include/CL/sycl/stl.hpp
@@ -10,6 +10,7 @@
 
 // 4.5 C++ Standard library classes required for the interface
 
+#include <CL/sycl/bit_cast.hpp>
 #include <CL/sycl/detail/defines.hpp>
 
 #include <exception>
diff --git a/sycl/test/bit_cast/bit_cast.cpp b/sycl/test/bit_cast/bit_cast.cpp
index f383bf70ca313..bf9af037a5b0c 100644
--- a/sycl/test/bit_cast/bit_cast.cpp
+++ b/sycl/test/bit_cast/bit_cast.cpp
@@ -21,8 +21,7 @@ To doBitCast(const From &ValueToConvert) {
     Queue.submit([&](sycl::handler &cgh) {
       auto acc = Buf.template get_access<sycl_write>(cgh);
       cgh.single_task<class BitCastKernel<To, From>>([=]() {
-        // TODO: change to sycl::bit_cast in the future
-        acc[0] = sycl::detail::bit_cast<To>(ValueToConvert);
+        acc[0] = sycl::bit_cast<To>(ValueToConvert);
       });
     });
   }