[SYCL]Fix incorrect behaviors in some operations using sycl::vec<bool, N>
#8942
Conversation
Our implementation of sycl::vec<bool, N> has `signed char` rather than `bool` for the underlying data structure. This causes that our underlying implementation can have values other than `0/1`, which produces wrong results for some operations. This commit fixes that by doing appropriate casts and data conversions when dealing with bool type and operations that can lead to results other than `0/1`. Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
maarquitos14
changed the title
sycl::vec<bool, N>
maarquitos14
changed the title
sycl::vec<bool, N>sycl::vec<bool, N>
sycl/include/sycl/types.hpp
Outdated
| @@ -778,7 +792,8 @@ template <typename Type, int NumElements> class vec { | |||
|
|
|||
| template <typename Ty = DataT> | |||
| explicit constexpr vec(const EnableIfNotHostHalf<Ty> &arg) | |||
| : m_Data{(DataType)vec_data<Ty>::get(arg)} {} | |||
| : m_Data{DataType(vec_data<Ty>::get(arg))} { | |||
| } | |||
There was a problem hiding this comment.
I think this will still initialize a m_Data with -1 if given true as input - the argument given to initialize a bool ext_vector should be casted to something other than bool to avoid the special behavior of splatting a bool.
There was a problem hiding this comment.
Also, the operator= overload right below this function will still have the same bool splatting behavior.
There was a problem hiding this comment.
I think this will still initialize a
m_Datawith -1 if giventrueas input - the argument given to initialize aboolext_vector should be casted to something the special behavior of splatting a bool.
It looks like this change will first cast the bool value to the underlying std::intN_t and then splat. That should prevent the vector from being initialized to -1.
@maarquitos14, can you add a test case to show that the sycl::vec<bool, N> operations are correct with this change?
There was a problem hiding this comment.
For sycl::vec<bool, N>, If Ty = bool, then the following declarations suggest that vec_data<Ty>::get is a bool:
template <typename T> T vec_helper<T>::get();
template <typename T> using vec_data = detail::vec_helper<T>;There was a problem hiding this comment.
@maarquitos14, can you add a test case to show that the
sycl::vec<bool, N>operations are correct with this change?
I have just added a file with many tests, but please, feel free to suggest new ones if you think they should be added.
Also, the
operator=overload right below this function will still have the same bool splatting behavior.
I believe it's not necessary to change the operator= since the rhs should have correct values already. I might be wrong, though, so if you find a case where this could be an issue, please, share it with me so that I can add it to the test file and also update the code accordingly.
There was a problem hiding this comment.
Right, that makes sense. And I realized I didn't see/forgot about the specialization added for vec_helper so that's why I thought this wouldn't work.
There was a problem hiding this comment.
Right, that makes sense. And I realized I didn't see/forgot about the specialization added for vec_helper so that's why I thought this wouldn't work.
Ah, right, that is indeed the key to make this work.
Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>
|
Friendly ping @bso-intel @intel/llvm-reviewers-runtime. |
Update a `sycl::vec` constructor to ensure that element values are first converted to `DataT` before being stored as `vec_data_t<DataT>`. Add a test case that uses the updated constructor. Fixes bug from intel#8942. Signed-off-by: Michael Aziz <michael.aziz@intel.com>
Update a `sycl::vec` constructor to ensure that element values are first converted to `DataT` before being stored as `vec_data_t<DataT>`. Add a test case that uses the updated constructor. Fixes bug from #8942. Signed-off-by: Michael Aziz <michael.aziz@intel.com>
Our implementation of sycl::vec<bool, N> has
signed charrather thanboolfor the underlying data structure. This causes that our underlying implementation can have values other than0/1, which produces wrong results for some operations. This PR fixes that by doing appropriate casts and data conversions when dealing with bool type and operations that can lead to results other than0/1.