Skip to content

[SYCL] Speed up device compilation for half and bfloat16 #20050

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Sep 11, 2025
Merged

[SYCL] Speed up device compilation for half and bfloat16 #20050

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
2cda53a
[SYCL] Speed up device compilation for `half` and `bfloat16`
aelovikov-intel Sep 10, 2025
436c842
Try to fix win
aelovikov-intel Sep 10, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions sycl/include/sycl/detail/vector_arith.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,8 @@
#include <sycl/detail/type_traits/vec_marray_traits.hpp>
#include <sycl/ext/oneapi/bfloat16.hpp>

#include <functional>

namespace sycl {
inline namespace _V1 {
namespace detail {
Expand Down
13 changes: 0 additions & 13 deletions sycl/include/sycl/ext/intel/esimd/detail/half_type_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,19 +115,6 @@ struct is_esimd_arithmetic_type<half_raw_type, void> : std::true_type {};
template <>
struct is_esimd_arithmetic_type<sycl::half, void> : std::true_type {};

// Misc
inline std::ostream &operator<<(std::ostream &O, sycl::half const &rhs) {
O << static_cast<float>(rhs);
return O;
}

inline std::istream &operator>>(std::istream &I, sycl::half &rhs) {
float ValFloat = 0.0f;
I >> ValFloat;
rhs = ValFloat;
return I;
}
Comment on lines -118 to -129
Copy link
Contributor Author

@aelovikov-intel aelovikov-intel Sep 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

half's hidden friends already do this, no idea why ESIMD needed this in the first place.

Copy link
Contributor

@sarnex sarnex Sep 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

i think we had an implementation of half in esimd before sycl::half was implemented, probably we forgot to remove it

Copy link
Contributor

@bader bader Sep 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

i think we had an implementation of half in esimd before sycl::half was implemented, probably we forgot to remove it

I see stream operators implemented before ESIMD.

ESIMD implementation added by: f34ba2c

Here is sycl::half implementation for the same SHA:

llvm/sycl/include/CL/sycl/half_type.hpp

Lines 518 to 528 in f34ba2c

inline std::ostream &operator<<(std::ostream &O, cl::sycl::half const &rhs) {
O << static_cast<float>(rhs);
return O;
}
inline std::istream &operator>>(std::istream &I, cl::sycl::half &rhs) {
float ValFloat = 0.0f;
I >> ValFloat;
rhs = ValFloat;
return I;
}

Copy link
Contributor

@sarnex sarnex Sep 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

esimd doesn't work on the host now, im not sure about at the time, but there is no point in having this code for the host


} // namespace ext::intel::esimd::detail
} // namespace _V1
} // namespace sycl
Expand Down
17 changes: 12 additions & 5 deletions sycl/include/sycl/ext/oneapi/bfloat16.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,12 +8,12 @@

#pragma once

#include <sycl/aliases.hpp> // for half
#include <sycl/bit_cast.hpp> // for bit_cast
#include <sycl/detail/defines_elementary.hpp> // for __DPCPP_SYCL_EXTERNAL
#include <sycl/half_type.hpp> // for half
#include <sycl/aliases.hpp>
#include <sycl/bit_cast.hpp>
#include <sycl/detail/defines_elementary.hpp>
#include <sycl/half_type.hpp>

#include <cstdint> // for uint16_t, uint32_t
#include <cstdint>

namespace sycl {
inline namespace _V1 {
Expand Down Expand Up @@ -126,6 +126,12 @@ class bfloat16 {
// for floating-point types.

// Stream Operator << and >>
#ifdef __SYCL_DEVICE_ONLY__
// std::istream/std::ostream aren't usable on device, so don't provide a
// definition to save compile time by using lightweight `<iosfwd>`.
inline friend std::ostream &operator<<(std::ostream &O, bfloat16 const &rhs);
inline friend std::istream &operator>>(std::istream &I, bfloat16 &rhs);
#else
inline friend std::ostream &operator<<(std::ostream &O, bfloat16 const &rhs) {
O << static_cast<float>(rhs);
return O;
Expand All @@ -137,6 +143,7 @@ class bfloat16 {
rhs = ValFloat;
return I;
}
#endif

private:
Bfloat16StorageT value;
Expand Down
27 changes: 21 additions & 6 deletions sycl/include/sycl/half_type.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,18 +10,25 @@

#include <sycl/bit_cast.hpp> // for bit_cast
#include <sycl/detail/export.hpp> // for __SYCL_EXPORT
#include <sycl/detail/iostream_proxy.hpp> // for istream, ostream

#ifdef __SYCL_DEVICE_ONLY__
#include <sycl/aspects.hpp>
#endif

#include <cstddef> // for size_t
#include <cstdint> // for uint16_t, uint32_t, uint8_t
#include <functional> // for hash
#include <cstddef>
#include <cstdint>
#include <limits> // for float_denorm_style, float_r...
#include <string_view> // for hash
#include <type_traits> // for enable_if_t
#include <type_traits>

// For std::hash, seems to be the most lightweight header provide it under
// C++17:
#include <optional>

#ifdef __SYCL_DEVICE_ONLY__
#include <iosfwd>
#else
#include <sycl/detail/iostream_proxy.hpp>
#endif

#if !defined(__has_builtin) || !__has_builtin(__builtin_expect)
#define __builtin_expect(a, b) (a)
Expand Down Expand Up @@ -478,6 +485,13 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half {
#endif // __SYCL_DEVICE_ONLY__

// Operator << and >>
#ifdef __SYCL_DEVICE_ONLY__
// std::istream/std::ostream aren't usable on device, so don't provide a
// definition to save compile time by using lightweight `<iosfwd>`.
inline friend std::ostream &operator<<(std::ostream &O,
sycl::half const &rhs);
inline friend std::istream &operator>>(std::istream &I, sycl::half &rhs);
#else
inline friend std::ostream &operator<<(std::ostream &O,
sycl::half const &rhs) {
O << static_cast<float>(rhs);
Expand All @@ -490,6 +504,7 @@ class [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] half {
rhs = ValFloat;
return I;
}
#endif

template <typename Key> friend struct std::hash;

Expand Down