Skip to content

Commit

Permalink
[SYCL] Fix float-to-half conversion for minimum subnormal on host (#7401
Browse files Browse the repository at this point in the history
)

Currently when converting a floating point value representing the
minimum subnormal value for half to a half on host will result in 0
rather than the expected minimum value. This is due to a faulty check
for the minimum supported single-precision exponent when converting.
This commit fixes this check.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
  • Loading branch information
steffenlarsen committed Nov 18, 2022
1 parent 45d516c commit 514708b
Show file tree
Hide file tree
Showing 2 changed files with 26 additions and 1 deletion.
2 changes: 1 addition & 1 deletion sycl/include/sycl/half_type.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -78,7 +78,7 @@ inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) {
// Tie to even.
else if (roundBits == halfway)
Frac16 += Frac16 & 1;
} else if (__builtin_expect(Exp32Diff > -24, 0)) {
} else if (__builtin_expect(Exp32Diff > -25, 0)) {
// subnormals
Frac16 = (Frac32 | (uint32_t(1) << 23)) >> (-Exp32Diff - 1);
}
Expand Down
25 changes: 25 additions & 0 deletions sycl/test/regression/half_host_subnormal_min.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,25 @@
// RUN: %clangxx -fsycl %s -o %t.out
// RUN: %t.out
//
// Checks that sycl::half on host can correctly cast its minimum subnormal value
// to and from a floating point value.

#include <sycl/sycl.hpp>

int main() {
sycl::half SubnormalMin =
sycl::bit_cast<sycl::half>((uint16_t)0b0000000000000001u);
sycl::half ConvertedSubnormalMin =
static_cast<sycl::half>(static_cast<float>(SubnormalMin));

if (SubnormalMin != ConvertedSubnormalMin) {
std::cout << "Failed! (0x" << std::hex
<< sycl::bit_cast<uint16_t>(SubnormalMin) << " != 0x"
<< sycl::bit_cast<uint16_t>(ConvertedSubnormalMin) << ")"
<< std::endl;
return 1;
}

std::cout << "Passed!" << std::endl;
return 0;
}

0 comments on commit 514708b

Please sign in to comment.