From ebfa066523e12642068ba35a830cfadd57b91eeb Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 23 Sep 2025 14:53:49 -0700 Subject: [PATCH 1/2] [NFC][SYCL] Drop `bfloat16::Bfloat16StorageT` It's not part of the specification and should have never been a public type alias inside `bfloat16`. There aren't too many uses of it (`bfloat16` itself and `convertToOpenCLType`/`vec::convert`) so I don't see much value in creating a named type alias. --- .../sycl/detail/generic_type_traits.hpp | 3 +-- sycl/include/sycl/detail/vector_convert.hpp | 3 +-- sycl/include/sycl/ext/oneapi/bfloat16.hpp | 22 +++++++++++-------- sycl/test-e2e/BFloat16/bfloat_hw.cpp | 6 ++--- 4 files changed, 18 insertions(+), 16 deletions(-) diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp index d1f2caad67a97..250a4326298cb 100644 --- a/sycl/include/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/sycl/detail/generic_type_traits.hpp @@ -175,8 +175,7 @@ template auto convertToOpenCLType(T &&x) { } else if constexpr (std::is_same_v) { // On host, don't interpret BF16 as uint16. #ifdef __SYCL_DEVICE_ONLY__ - using OpenCLType = typename no_ref::Bfloat16StorageT; - return sycl::bit_cast(x); + return sycl::bit_cast(x); #else return std::forward(x); #endif diff --git a/sycl/include/sycl/detail/vector_convert.hpp b/sycl/include/sycl/detail/vector_convert.hpp index 544250aff82ac..8fe6e5c726342 100644 --- a/sycl/include/sycl/detail/vector_convert.hpp +++ b/sycl/include/sycl/detail/vector_convert.hpp @@ -895,8 +895,7 @@ vec vec::convert() const { #endif bool, /*->*/ std::uint8_t, // sycl::half, /*->*/ sycl::detail::half_impl::StorageT, // - sycl::ext::oneapi::bfloat16, - /*->*/ sycl::ext::oneapi::bfloat16::Bfloat16StorageT, // + sycl::ext::oneapi::bfloat16, /*->*/ uint16_t, // char, /*->*/ detail::ConvertToOpenCLType_t, // DataT, /*->*/ DataT // >::type diff --git a/sycl/include/sycl/ext/oneapi/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/bfloat16.hpp index bb47ac7667556..67111fad1bc7d 100644 --- a/sycl/include/sycl/ext/oneapi/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/bfloat16.hpp @@ -21,7 +21,11 @@ namespace ext::oneapi { class bfloat16 { public: - using Bfloat16StorageT = uint16_t; +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + using Bfloat16StorageT + __SYCL_DEPRECATED("bfloat16::Bfloat16StorageT is non-standard and has " + "been deprecated.") = uint16_t; +#endif bfloat16() = default; ~bfloat16() = default; @@ -58,7 +62,7 @@ class bfloat16 { friend bfloat16 operator-(const bfloat16 &lhs) { #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__) && \ (__SYCL_CUDA_ARCH__ >= 800) - Bfloat16StorageT res; + uint16_t res; asm("neg.bf16 %0, %1;" : "=h"(res) : "h"(lhs.value)); return bit_cast(res); #else @@ -146,18 +150,18 @@ class bfloat16 { #endif private: - Bfloat16StorageT value; + uint16_t value; // Private tag used to avoid constructor ambiguity. struct private_tag { explicit private_tag() = default; }; - constexpr bfloat16(Bfloat16StorageT Value, private_tag) : value{Value} {} + constexpr bfloat16(uint16_t Value, private_tag) : value{Value} {} // Explicit conversion functions - static float to_float(const Bfloat16StorageT &a); - static Bfloat16StorageT from_float(const float &a); + static float to_float(const uint16_t &a); + static uint16_t from_float(const float &a); // Friend traits. friend std::numeric_limits; @@ -178,7 +182,7 @@ class bfloat16 { extern "C" __DPCPP_SYCL_EXTERNAL float __devicelib_ConvertBF16ToFINTEL(const uint16_t &) noexcept; #endif -inline float bfloat16::to_float(const bfloat16::Bfloat16StorageT &a) { +inline float bfloat16::to_float(const uint16_t &a) { #if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__)) return __devicelib_ConvertBF16ToFINTEL(a); #else @@ -213,11 +217,11 @@ inline uint16_t from_float_to_uint16_t(const float &a) { extern "C" __DPCPP_SYCL_EXTERNAL uint16_t __devicelib_ConvertFToBF16INTEL(const float &) noexcept; #endif -inline bfloat16::Bfloat16StorageT bfloat16::from_float(const float &a) { +inline uint16_t bfloat16::from_float(const float &a) { #if defined(__SYCL_DEVICE_ONLY__) #if defined(__NVPTX__) #if (__SYCL_CUDA_ARCH__ >= 800) - Bfloat16StorageT res; + uint16_t res; asm("cvt.rn.bf16.f32 %0, %1;" : "=h"(res) : "f"(a)); return res; #else diff --git a/sycl/test-e2e/BFloat16/bfloat_hw.cpp b/sycl/test-e2e/BFloat16/bfloat_hw.cpp index 0154d21156fd1..98f4a15b0df31 100644 --- a/sycl/test-e2e/BFloat16/bfloat_hw.cpp +++ b/sycl/test-e2e/BFloat16/bfloat_hw.cpp @@ -17,15 +17,15 @@ using get_uint_type_of_size = typename std::conditional_t< std::conditional_t>>>; using bfloat16 = sycl::ext::oneapi::bfloat16; -using Bfloat16StorageT = get_uint_type_of_size; +static_assert(sizeof(bfloat16) == size(uint16_t)); -bool test(float Val, Bfloat16StorageT Bits) { +bool test(float Val, uint16_t Bits) { std::cout << "Value: " << Val << " Bits: " << std::hex << "0x" << Bits << std::dec << "...\n"; bool Passed = true; { std::cout << " float -> bfloat16 conversion ..."; - Bfloat16StorageT RawVal = sycl::bit_cast(bfloat16(Val)); + auto RawVal = sycl::bit_cast(bfloat16(Val)); bool Res = (RawVal == Bits); Passed &= Res; From 124e0d609f576191cf8d029d1384c42556defcec Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 23 Sep 2025 15:54:59 -0700 Subject: [PATCH 2/2] Update sycl/test-e2e/BFloat16/bfloat_hw.cpp --- sycl/test-e2e/BFloat16/bfloat_hw.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/BFloat16/bfloat_hw.cpp b/sycl/test-e2e/BFloat16/bfloat_hw.cpp index 98f4a15b0df31..79a5832c8fed3 100644 --- a/sycl/test-e2e/BFloat16/bfloat_hw.cpp +++ b/sycl/test-e2e/BFloat16/bfloat_hw.cpp @@ -17,7 +17,7 @@ using get_uint_type_of_size = typename std::conditional_t< std::conditional_t>>>; using bfloat16 = sycl::ext::oneapi::bfloat16; -static_assert(sizeof(bfloat16) == size(uint16_t)); +static_assert(sizeof(bfloat16) == sizeof(uint16_t)); bool test(float Val, uint16_t Bits) { std::cout << "Value: " << Val << " Bits: " << std::hex << "0x" << Bits