Fix nan check on 16bit precision in triangular solve. #1860
Conversation
c38d506 to
0d97768
Compare
| while (is_nan_exact( | ||
| val = load_relaxed(x + dependency * x_stride + rhs))) { | ||
| output_val = load_relaxed(x + dependency * x_stride + rhs))) { |
There was a problem hiding this comment.
Since the conversions are the issue here, I would propose a slightly different approach:
- Make sure that
xstores the bit-wise exact nan value ofSharedValueType, viamemcpy. - Use a function
bitwise_equal(a, b), whereais the loaded value andbthe expected nan value.
I think this could make it a bit clearer what is happening, and not use the gko<half|bfloat> types internally.
There was a problem hiding this comment.
Maybe I get the first point wrong. x is stored exact nan value because we fill it by gko::nan<ValueType> and x_shared is stored by gko::nan<SharedValueType>. We do not have the casting during assignment.
I think bitwise_equal(a, b) indeed makes the comparison value more clear. (or we can make it more clear/consistent what the exact nan among different backend/compiler)
Unfortunately, we do not have the host_type map currently to do bitwise_eqaul(a, nan<host_type<T>>) and I would not like to introduce another type mapping in short time of notice.
There was a problem hiding this comment.
You're right, I guess it was getting a bit late for me. Could an alternative be to fill the x with gko::nan<device_type<ValueType>> already? I will still approve this, just curious.
There was a problem hiding this comment.
no, at least not a small change. Postponing it to later such that we do not rush with a mistake.
gko::nan will call std::numeric_limits<>::quiet_NaN, which is only a default value on device type (usually 0).
Thus, it needs to specialize that on the device type or specialize gko::nan on device/math.
I think ensuring that quiet_NaN in device_numeric_limits on device type can give the same value as std::numeric_limits on host type might be more consistent.
Filling some value with device_type value is a bit weird to me unless we use another kernel to fill the value in the kernel directly not set the value from host.
0d97768 to
ed2fac8
Compare
We use the nan to check the status of dependence by checking the same bit representation as
gko::nan(quiet_NaN) in #1665This PR contains two parts:
half,bfloat16are checked bygko::nan<gko::half>andgko::nan<gko::bfloat16>, which is the value we initialize in the x vector.is_nan_exactcheck on the orignial type because the casting can give not the same kind of quiet NaN or different precision can use different kind of NaN.There is no specification for quiet NaN in IEEE more than exponent field is filled by 1 and the significant field != 0
Note. the first one can also be done by add quiet_NaN into device_numeric_limit and add test ensuring they are identical between device and host. However, the current value of quiet NaN of gko::half might be conflict with one of sycl::half. Or different compiler might also give different quiet_NaN. It needs more t 8000 ime to check the consistence, so I put the fix only under
is_nan_exactnow