Skip to content

[SYCL] Fix bfloat16::to_float() host implementation. #6675

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 4 commits into from
Sep 2, 2022
Merged

[SYCL] Fix bfloat16::to_float() host implementation. #6675

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
546ad16
[SYCL] Fix bfloat16::to_float() host implementation.
kbobrovs Aug 30, 2022
94f982a
Fix Linux strict alias mode compilation error.
kbobrovs Aug 31, 2022
f1b19ae
Improve bfloat16::to_float not to use std::memcpy, fix test.
kbobrovs Aug 31, 2022
85fb1da
Review comments: use sycl::bit_cast instead of union-based hack.
kbobrovs Aug 31, 2022
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
3 changes: 1 addition & 2 deletions sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,10 +63,9 @@ class bfloat16 {
return __spirv_ConvertBF16ToFINTEL(a);
#endif
#else
// Shift temporary variable to silence the warning
uint32_t bits = a;
bits <<= 16;
return static_cast<float>(bits);
return sycl::bit_cast<float>(bits);
#endif
}

Expand Down
10 changes: 6 additions & 4 deletions sycl/test/extensions/bfloat16_host.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,15 +73,17 @@ int main() {
Success &= check_bf16_from_float(std::numeric_limits<float>::quiet_NaN(),
std::stoi("1111111111000001", nullptr, 2));

// see https://float.exposed/b0xffff
Success &= check_bf16_to_float(
0, bitsToFloatConv(std::string("00000000000000000000000000000000")));
Success &= check_bf16_to_float(
1, bitsToFloatConv(std::string("01000111100000000000000000000000")));
Copy link
Contributor

Choose a reason for hiding this comment

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

Oh, that was a huge mistake, sorry about that.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

bugs happen, no problem )

1, bitsToFloatConv(std::string("00000000000000010000000000000000")));
Success &= check_bf16_to_float(
42, bitsToFloatConv(std::string("01001010001010000000000000000000")));
42, bitsToFloatConv(std::string("00000000001010100000000000000000")));
Success &= check_bf16_to_float(
std::numeric_limits<uint16_t>::max(),
bitsToFloatConv(std::string("01001111011111111111111100000000")));
// std::numeric_limits<uint16_t>::max() - 0xffff is bfloat16 -Nan and
// -Nan == -Nan check in check_bf16_to_float would fail, so use not Nan:
65407, bitsToFloatConv(std::string("11111111011111110000000000000000")));
if (!Success)
return -1;
return 0;
Expand Down