[SYCL] half math is using float precision for multiply/add/etc operators.

[JackAKirk]

sycl::half operators like

llvm/sycl/include/sycl/half_type.hpp

Line 155 in b6faefb

*this = operator float() + static_cast<float>(rhs);

are converting to float for operations.

We can specialize the cuda backend case in order to call appropriate instructions for f16 precision operators instead, but I opened this issue to raise awareness for other backends/ to get potentially a general solution that lowers to backend specific instructions.

[0x12CC]

llvm/sycl/include/sycl/half_type.hpp

Line 229 in 7515da8

Data += rhs.Data;

@JackAKirk, the conversions aren't there anymore. Is this issue still relevant or can we close it?

[AlexeySachkov]

@JackAKirk, the code you are pointing to is host implementation for half. It does indeed makes a fallback to float math, because we have no half in standard C++.

However, on device side we should be using _Float16 for math operations and that has been the case for more than 4 years now, see #1089

[JackAKirk]

Thanks for the clarification, my mistake.