[SYCL] Add sycl::detail::bit_cast implementation

sycl/doc/extensions/README.md

@@ -7,7 +7,7 @@ DPC++ extensions status:
 
 |  Extension  |    Status   |   Comment   |
 |-------------|:------------|:------------|
-| [SYCL_INTEL_bitcast](Bitcast/SYCL_INTEL_bitcast.asciidoc)                                                                   | Proposal                                  | |
+| [SYCL_INTEL_bitcast](Bitcast/SYCL_INTEL_bitcast.asciidoc)                                                                   | Supported                                 | As sycl::detail::bit_cast |
 | [C and C++ Standard libraries support](C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst)                                     | Partially supported(OpenCL: CPU, GPU)     | |
 | [SYCL_INTEL_data_flow_pipes](DataFlowPipes/data_flow_pipes.asciidoc)                                                        | Partially supported(OpenCL: ACCELERATOR)  | kernel_host_pipe_support part is not implemented |
 | [SYCL_INTEL_deduction_guides](deduction_guides/SYCL_INTEL_deduction_guides.asciidoc)                                        | Supported                                 | |

sycl/include/CL/sycl/detail/helpers.hpp

@@ -17,6 +17,7 @@
 #include <CL/sycl/detail/type_traits.hpp>
 
 #include <memory>
+#include <numeric> // std::bit_cast
 #include <stdexcept>
 #include <type_traits>
 #include <vector>
@@ -42,6 +43,31 @@ inline void memcpy(void *Dst, const void *Src, size_t Size) {
   }
 }
 
+template <typename To, typename From>
+constexpr 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
 // dependency list and waits for others.

sycl/include/CL/sycl/intel/sub_group.hpp

@@ -21,7 +21,6 @@
 #include <CL/sycl/range.hpp>
 #include <CL/sycl/types.hpp>
 
-#include <numeric> // std::bit_cast
 #include <type_traits>
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -72,22 +71,6 @@ using AcceptableForLocalLoadStore =
     bool_constant<!std::is_same<void, SelectBlockT<T>>::value &&
                   Space == access::address_space::local_space>;
 
-// TODO: move this to public cl::sycl::bit_cast as extension?
-template <typename To, typename From> To bit_cast(const From &from) {
-#if __cpp_lib_bit_cast
-  return std::bit_cast<To>(from);
-#else
-
-#if __has_builtin(__builtin_bit_cast)
-  return __builtin_bit_cast(To, from);
-#else
-  To to;
-  sycl::detail::memcpy(&to, &from, sizeof(To));
-  return to;
-#endif // __has_builtin(__builtin_bit_cast)
-#endif // __cpp_lib_bit_cast
-}
-
 template <typename T, access::address_space Space>
 T load(const multi_ptr<T, Space> src) {
   using BlockT = SelectBlockT<T>;
@@ -97,7 +80,7 @@ T load(const multi_ptr<T, Space> src) {
   BlockT Ret =
       __spirv_SubgroupBlockReadINTEL<BlockT>(reinterpret_cast<PtrT>(src.get()));
 
-  return bit_cast<T>(Ret);
+  return sycl::detail::bit_cast<T>(Ret);
 }
 
 template <int N, typename T, access::address_space Space>
@@ -110,7 +93,7 @@ vec<T, N> load(const multi_ptr<T, Space> src) {
   VecT Ret =
       __spirv_SubgroupBlockReadINTEL<VecT>(reinterpret_cast<PtrT>(src.get()));
 
-  return bit_cast<typename vec<T, N>::vector_t>(Ret);
+  return sycl::detail::bit_cast<typename vec<T, N>::vector_t>(Ret);
 }
 
 template <typename T, access::address_space Space>
@@ -119,7 +102,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()),
-                                  bit_cast<BlockT>(x));
+                                  sycl::detail::bit_cast<BlockT>(x));
 }
 
 template <int N, typename T, access::address_space Space>
@@ -130,7 +113,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()),
-                                  bit_cast<VecT>(x));
+                                  sycl::detail::bit_cast<VecT>(x));
 }
 
 } // namespace sub_group

sycl/test/bit_cast/bit_cast.cpp

@@ -0,0 +1,84 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+#include <iostream>
+#include <math.h>
+#include <type_traits>
+
+constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write;
+
+template <typename To, typename From>
+class BitCastKernel;
+
+template <typename To, typename From>
+To doBitCast(const From &ValueToConvert) {
+  std::vector<To> Vec(1);
+  {
+    sycl::buffer<To, 1> Buf(Vec.data(), 1);
+    sycl::queue Queue;
+    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);
+      });
+    });
+  }
+  return Vec[0];
+}
+
+template <typename To, typename From>
+int test(const From &Value) {
+  auto ValueConvertedTwoTimes = doBitCast<From>(doBitCast<To>(Value));
+  bool isOriginalValueEqualsToConvertedTwoTimes = false;
+  if (std::is_integral<From>::value) {
+    isOriginalValueEqualsToConvertedTwoTimes = Value == ValueConvertedTwoTimes;
+  } else if ((std::is_floating_point<From>::value) || std::is_same<From, cl::sycl::half>::value) {
+    static const float Epsilon = 0.0000001f;
+    isOriginalValueEqualsToConvertedTwoTimes = fabs(Value - ValueConvertedTwoTimes) < Epsilon;
+  } else {
+    std::cerr << "Type " << typeid(From).name() << " neither integral nor floating point nor cl::sycl::half\n";
+    return 1;
+  }
+  if (!isOriginalValueEqualsToConvertedTwoTimes) {
+    std::cerr << "FAIL: Original value which is " << Value << " != value converted two times which is " << ValueConvertedTwoTimes << "\n";
+    return 1;
+  }
+  std::cout << "PASS\n";
+  return 0;
+}
+
+int main() {
+  int ReturnCode = 0;
+
+  std::cout << "cl::sycl::half to unsigned short ...\n";
+  ReturnCode += test<unsigned short>(cl::sycl::half(1.0f));
+
+  std::cout << "unsigned short to cl::sycl::half ...\n";
+  ReturnCode += test<cl::sycl::half>(static_cast<unsigned short>(16384));
+
+  std::cout << "cl::sycl::half to short ...\n";
+  ReturnCode += test<short>(cl::sycl::half(1.0f));
+
+  std::cout << "short to cl::sycl::half ...\n";
+  ReturnCode += test<cl::sycl::half>(static_cast<short>(16384));
+
+  std::cout << "int to float ...\n";
+  ReturnCode += test<float>(static_cast<int>(2));
+
+  std::cout << "float to int ...\n";
+  ReturnCode += test<int>(static_cast<float>(-2.4f));
+
+  std::cout << "unsigned int to float ...\n";
+  ReturnCode += test<float>(static_cast<unsigned int>(6));
+
+  std::cout << "float to unsigned int ...\n";
+  ReturnCode += test<unsigned int>(static_cast<float>(-2.4f));
+
+  return ReturnCode;
+}