-
Notifications
You must be signed in to change notification settings - Fork 769
[SYCL][ESIMD]Limit bfloat16 operators to scalars to enable operations with simd vectors #12089
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Conversation
@@ -155,13 +155,14 @@ class bfloat16 { | |||
f op static_cast<float>(rhs); \ | |||
return lhs = f; \ | |||
} \ | |||
template <typename T> \ | |||
template <typename T, class = std::enable_if_t<std::is_scalar_v<T>>> \ |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I am not sure, but this might be a little too brave fix. Will it limit the operation for T
= sycl::vec
?
Also, in my opinion, the current implementation (not your patch) is wrong.
Even if T
is convertible to float
does not mean this operator must be used.
For example, let's say some struct my_2d_vec
may be convertible to float
. It does not necessary mean that this operator defined here must be used. User may want to have my_2d_vec::operator+(bfloat16 v) {this.x += v; this.y += v;}
Much better way of supporting arith operations for bfloat16 is already used for sycl::half.
sycl::half implements a number of operations for various trivial types (float
, int32_t
, and some other) and let T use implicit conversion to one of those trivial types.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It is probably too limiting but sycl::half is not much better as it contains pretty limited amount of types it supports which means the user will likely be forced to explicitly convert the argument to one of those types for example it doesn't support int16 but for some reason supports int64 which is likely to cause overflow. Probably better to get the feedback from the owners of this class to see which direction they want it to go.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Reimplemented the operators in a way similar to half
|
||
Q.single_task([=]() SYCL_ESIMD_KERNEL { | ||
simd<T, N> Vec(TOne); | ||
Vec = Vec + TTen; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Please test other operations too. +
, -
, *
, +=
, etc.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Added test operations for +, -, *,/. I will test +=, -=, *= and /= in the test for the other JIRA
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
ESIMD changes look good to me.
…erations with simd vectors (intel#12089)" This reverts commit 8c92df9.
The purpose of this change is to limit operators defined for bfloat16 to scalar types to allow arithmetic operations between bfloat16 scalars and simd vectors. This allows to use simd operators that are defined separately and support operations between vectors and scalars