[SYCL]Update tests for half operator #1012
Conversation
|
|
||
| using namespace cl::sycl; | ||
| using namespace sycl::ext::intel; | ||
| using namespace sycl::ext::intel::esimd; |
There was a problem hiding this comment.
I wonder if we can use something like this (not sure about the proper spelling)
template <class T>
using simd = esimd::simd<T>
to eliminate the changes below.
There was a problem hiding this comment.
I think using esimd::simd<T, N> in the code below is clearer compared to introducing simd alias.
There was a problem hiding this comment.
I agree with @kbobrovs I think esimd::simd<T, N> is more clear.
| int16_t cmp_val = reinterpret_cast<int16_t &>(val); | ||
| int16_t cmp_ref = reinterpret_cast<int16_t &>(ref); | ||
| if (std::abs(cmp_val - cmp_ref) > 1) { | ||
| if (std::abs(cmp_val - cmp_ref) > 2) { |
There was a problem hiding this comment.
Do you know why we need this at all? Other fp types seem to be checked for exact equality (line 112).
There was a problem hiding this comment.
This compares the results on the host vs device. The host does not have support for _Float16 so this is implemented with uint16_t data type and converting into 32-bit float for arithmetic and back to uint16_t.
The device has support for _Float16, so it is performing the same operations at 16-bit precision. So the difference in the precision/computation between host and device means there is a larger difference in output, compared to other data types.
This PR adds missing half operations +-*/. There are existing operations for the legacy class host_half_impl, however these were not extended to the half class. These operations were being performed as floating point operations via the implicit floating conversion. This results in the output being a float not half type. The template is limited to arithmetic types to prevent ambiguous templating. A minor change to built-in __fract is needed to ensure fmin is not ambiguous. llvm-test-suite PR: intel/llvm-test-suite#1012 Fixes: #6028
This PR makes some small changes due to the introduction of some missing operators. It changes the namespaces in ext_math.cpp as it was ambiguous as to if the sycl::abs() or sycl::ext::intel::esimd::abs() function should be used. There is also an adjustment of the load store's check tolerance. The checks tolerance is slightly increased, as previously these operations would have been performed as 32-bit floats. Depends on: intel/llvm#6061
This PR makes some small changes due to the introduction of some missing operators. It changes the namespaces in ext_math.cpp as it was ambiguous as to if the sycl::abs() or sycl::ext::intel::esimd::abs() function should be used. There is also an adjustment of the load store's check tolerance. The checks tolerance is slightly increased, as previously these operations would have been performed as 32-bit floats. Depends on: intel#6061
4 participants
This PR makes some small changes due to the introduction of some missing operators.
It changes the namespaces in
ext_math.cppas it was ambiguous as to if thesycl::abs()orsycl::ext::intel::esimd::abs()function should be used.There is also an adjustment of the load store's check tolerance. The checks tolerance is slightly increased, as previously these operations would have been performed as 32-bit floats.
Depends on: intel/llvm#6061