- 
                Notifications
    
You must be signed in to change notification settings  - Fork 794
 
          [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.
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
Conversation
* `std::istream`/`std::ostream` aren't usable on device, so limit `operator<<`/`operator>>` to declaration only to limit device includes to `<iosfwd>` instead of much heavier `<istream>`/`<ostream>`. * Use "lighter" `<optional>` to get `std::hash`
| // 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; | ||
| } | 
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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
There was a problem hiding this comment.
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; | |
| } | 
There was a problem hiding this comment.
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
| // 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; | ||
| } | 
There was a problem hiding this comment.
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
std::istream/std::ostreamaren't usable on device, so limitoperator<</operator>>to declaration only to limit device includes to<iosfwd>instead of much heavier<istream>/<ostream>.<optional>to getstd::hash