Skip to content

[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

Merged
merged 8 commits into from
Jan 3, 2024

Conversation

fineg74
Copy link
Contributor

@fineg74 fineg74 commented Dec 6, 2023

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

@fineg74 fineg74 requested review from a team as code owners December 6, 2023 00:23
@fineg74 fineg74 requested a review from bso-intel December 6, 2023 00:23
@@ -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>>> \
Copy link
Contributor

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.

Copy link
Contributor Author

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.

Copy link
Contributor Author

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;
Copy link
Contributor

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.

Copy link
Contributor Author

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

@fineg74 fineg74 marked this pull request as draft December 12, 2023 18:45
@fineg74 fineg74 marked this pull request as ready for review December 13, 2023 20:06
Copy link
Contributor

@v-klochkov v-klochkov left a 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.

@againull againull merged commit 8c92df9 into intel:sycl Jan 3, 2024
@fineg74 fineg74 deleted the bfloat16Operator branch January 3, 2024 18:21
fineg74 added a commit to fineg74/llvm that referenced this pull request Jan 23, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants