diff --git a/sycl/doc/extensions/README.md b/sycl/doc/extensions/README.md index c9f1914a449a0..4938fa36b9673 100644 --- a/sycl/doc/extensions/README.md +++ b/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 | | diff --git a/sycl/include/CL/sycl/detail/helpers.hpp b/sycl/include/CL/sycl/detail/helpers.hpp index 024801274b006..d71c5be9e6185 100644 --- a/sycl/include/CL/sycl/detail/helpers.hpp +++ b/sycl/include/CL/sycl/detail/helpers.hpp @@ -17,6 +17,7 @@ #include #include +#include // std::bit_cast #include #include #include @@ -42,6 +43,31 @@ inline void memcpy(void *Dst, const void *Src, size_t Size) { } } +template +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::value, + "From must be trivially copyable"); + static_assert(std::is_trivially_copyable::value, + "To must be trivially copyable"); +#if __cpp_lib_bit_cast + return std::bit_cast(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::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. diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 268e9b5d23a96..a8de9028ad557 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -21,7 +21,6 @@ #include #include -#include // std::bit_cast #include #ifdef __SYCL_DEVICE_ONLY__ @@ -72,22 +71,6 @@ using AcceptableForLocalLoadStore = bool_constant>::value && Space == access::address_space::local_space>; -// TODO: move this to public cl::sycl::bit_cast as extension? -template To bit_cast(const From &from) { -#if __cpp_lib_bit_cast - return std::bit_cast(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 T load(const multi_ptr src) { using BlockT = SelectBlockT; @@ -97,7 +80,7 @@ T load(const multi_ptr src) { BlockT Ret = __spirv_SubgroupBlockReadINTEL(reinterpret_cast(src.get())); - return bit_cast(Ret); + return sycl::detail::bit_cast(Ret); } template @@ -110,7 +93,7 @@ vec load(const multi_ptr src) { VecT Ret = __spirv_SubgroupBlockReadINTEL(reinterpret_cast(src.get())); - return bit_cast::vector_t>(Ret); + return sycl::detail::bit_cast::vector_t>(Ret); } template @@ -119,7 +102,7 @@ void store(multi_ptr dst, const T &x) { using PtrT = sycl::detail::ConvertToOpenCLType_t>; __spirv_SubgroupBlockWriteINTEL(reinterpret_cast(dst.get()), - bit_cast(x)); + sycl::detail::bit_cast(x)); } template @@ -130,7 +113,7 @@ void store(multi_ptr dst, const vec &x) { sycl::detail::ConvertToOpenCLType_t>; __spirv_SubgroupBlockWriteINTEL(reinterpret_cast(dst.get()), - bit_cast(x)); + sycl::detail::bit_cast(x)); } } // namespace sub_group diff --git a/sycl/test/bit_cast/bit_cast.cpp b/sycl/test/bit_cast/bit_cast.cpp new file mode 100644 index 0000000000000..fcca0a8b2bfd9 --- /dev/null +++ b/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 + +#include +#include +#include + +constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write; + +template +class BitCastKernel; + +template +To doBitCast(const From &ValueToConvert) { + std::vector Vec(1); + { + sycl::buffer Buf(Vec.data(), 1); + sycl::queue Queue; + Queue.submit([&](sycl::handler &cgh) { + auto acc = Buf.template get_access(cgh); + cgh.single_task>([=]() { + // TODO: change to sycl::bit_cast in the future + acc[0] = sycl::detail::bit_cast(ValueToConvert); + }); + }); + } + return Vec[0]; +} + +template +int test(const From &Value) { + auto ValueConvertedTwoTimes = doBitCast(doBitCast(Value)); + bool isOriginalValueEqualsToConvertedTwoTimes = false; + if (std::is_integral::value) { + isOriginalValueEqualsToConvertedTwoTimes = Value == ValueConvertedTwoTimes; + } else if ((std::is_floating_point::value) || std::is_same::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(cl::sycl::half(1.0f)); + + std::cout << "unsigned short to cl::sycl::half ...\n"; + ReturnCode += test(static_cast(16384)); + + std::cout << "cl::sycl::half to short ...\n"; + ReturnCode += test(cl::sycl::half(1.0f)); + + std::cout << "short to cl::sycl::half ...\n"; + ReturnCode += test(static_cast(16384)); + + std::cout << "int to float ...\n"; + ReturnCode += test(static_cast(2)); + + std::cout << "float to int ...\n"; + ReturnCode += test(static_cast(-2.4f)); + + std::cout << "unsigned int to float ...\n"; + ReturnCode += test(static_cast(6)); + + std::cout << "float to unsigned int ...\n"; + ReturnCode += test(static_cast(-2.4f)); + + return ReturnCode; +}