[SYCL][COMPAT] Add bfloat16 support to several maths ops #15572

joeatodd · 2024-10-02T11:28:33Z

Adds support for sycl::ext::oneapi::bfloat16 to:

relu
clamp
fmax_nan
fmin_nan
min
max
compare_mask
unordered_compare_mask

bfloat16 can't be constexpr constructed

This includes sycl::half and sycl::ext::oneapi::bfloat16

- Adding bfloat to type lists - Making test fixtures work with sycl::vec<type,N> and sycl::marray<type,N>

No new functionality here

Mixed type cases aren't caught otherwise

- Added support for bfloat16 - Added tests for `sycl::vec<T,2>` & `sycl::marray<T,2>`

& a bit of tidying

& more tests

marray version never tested!

Function casts to sycl::complex<float> (no native bfloat16 support)

i.e. sycl::marray<bfloat16, N>, sycl::vec<bfloat16, N>

Also reimplement `fmax_nan` & `fmin_nan` in terms of improved max/min

- compare_mask & unordered_compare_mask support sycl::marray - sycl::marray tests for all `compare` APIs

This reverts commit 55e95ad.

ziranzha · 2024-10-08T06:02:42Z

sycl/include/syclcompat/math.hpp

-  return sycl::ext::oneapi::experimental::isnan(a);
+  if constexpr (detail::support_bfloat16_math &&
+                std::is_same_v<ValueT, sycl::ext::oneapi::bfloat16>) {
+    return sycl::ext::oneapi::experimental::isnan(a);


If not supported, how about throwing an exception

I added a static_assert since we'll know at compile time 👍

ziranzha · 2024-10-08T06:03:38Z

sycl/include/syclcompat/math.hpp

 min(ValueT a, ValueU b) {
-  return sycl::fmin(static_cast<std::common_type_t<ValueT, ValueU>>(a),
-                    static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  if constexpr (detail::support_bfloat16_math &&


How about throwing an exception if bfloat16 not supported

ziranzha · 2024-10-08T06:04:06Z

sycl/include/syclcompat/math.hpp

 max(ValueT a, ValueU b) {
-  return sycl::fmax(static_cast<std::common_type_t<ValueT, ValueU>>(a),
-                    static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  if constexpr (detail::support_bfloat16_math &&


Same as above.

ziranzha · 2024-10-08T08:31:05Z

sycl/include/syclcompat/math.hpp

+};
+
+template <typename T> struct common_type<sycl::ext::oneapi::bfloat16, T> {
+  using type = sycl::ext::oneapi::bfloat16; // std::common_type_t<float, T>;  //


I don't understand the comments. std::commone_type_t<float, sycl::ext::oneapi::bfloat16> will be bfloat16, and won't be float, why do the comments say bfloat16 will be promoted to float?

Woops! Thanks for catching that. Stale comment from earlier impl.

Alcpz

Thanks a lot for tackling this @joeatodd , It´ s quite a good PR

sycl/include/syclcompat/traits.hpp

sycl/include/syclcompat/math.hpp

Alcpz · 2024-10-09T09:20:41Z

sycl/include/syclcompat/math.hpp

+inline std::enable_if_t<syclcompat::is_floating_point_v<ValueT>, ValueT>
 cbrt(ValueT val) {
-  return sycl::cbrt(static_cast<ValueT>(val));
+  if constexpr (std::is_same_v<sycl::ext::oneapi::bfloat16, ValueT>) {


Is there a reason to not check for detail::support_bfloat16_math here?

Well, since we're casting to float & doing cbrt<float>, it's not needed. But I'm not confident casting is the right approach actually... The bfloat16 extension doesn't have cbrt.

I have reverted the cbrt change in this PR. The math op isn't currently implemented by the bfloat16 math extension, and it wasn't actually part of the request from DCPT.

sycl/include/syclcompat/math.hpp

sycl/test-e2e/syclcompat/common.hpp

sycl/test-e2e/syclcompat/math/math_compare.cpp

sycl/test-e2e/syclcompat/math/math_fixt.hpp

sycl/doc/syclcompat/README.md

and remove stale comments

joeatodd

Review comments addressed. Thanks @Alcpz and @ziranzha for your input 🙏

sycl/include/syclcompat/math.hpp

joeatodd · 2024-10-14T19:27:11Z

sycl/include/syclcompat/math.hpp

+inline std::enable_if_t<syclcompat::is_floating_point_v<ValueT>, ValueT>
 cbrt(ValueT val) {
-  return sycl::cbrt(static_cast<ValueT>(val));
+  if constexpr (std::is_same_v<sycl::ext::oneapi::bfloat16, ValueT>) {


Well, since we're casting to float & doing cbrt<float>, it's not needed. But I'm not confident casting is the right approach actually... The bfloat16 extension doesn't have cbrt.

joeatodd · 2024-10-14T19:31:18Z

sycl/include/syclcompat/math.hpp

+};
+
+template <typename T> struct common_type<sycl::ext::oneapi::bfloat16, T> {
+  using type = sycl::ext::oneapi::bfloat16; // std::common_type_t<float, T>;  //


Woops! Thanks for catching that. Stale comment from earlier impl.

sycl/include/syclcompat/math.hpp

joeatodd · 2024-10-16T11:41:21Z

sycl/include/syclcompat/math.hpp

-  return sycl::ext::oneapi::experimental::isnan(a);
+  if constexpr (detail::support_bfloat16_math &&
+                std::is_same_v<ValueT, sycl::ext::oneapi::bfloat16>) {
+    return sycl::ext::oneapi::experimental::isnan(a);


I added a static_assert since we'll know at compile time 👍

joeatodd · 2024-10-17T11:54:39Z

@intel/llvm-gatekeepers this is ready to merge 🙏

joeatodd added 20 commits September 25, 2024 15:17

No constexpr construction in math tests

5a000ac

bfloat16 can't be constexpr constructed

Add syclcompat::is_floating_point_v

c32b263

This includes sycl::half and sycl::ext::oneapi::bfloat16

Add bfloat16 to tests and generalize to container types

9f28f62

- Adding bfloat to type lists - Making test fixtures work with sycl::vec<type,N> and sycl::marray<type,N>

Add some guards & TODOs to math.hpp

8319654

No new functionality here

Generalize should_skip to variadic types

ec00197

Mixed type cases aren't caught otherwise

Specialize std::common_type for bfloat16

f8f62aa

Add bfloat16 support (& more extensive tests) to fmin/max_nan

dad728c

- Added support for bfloat16 - Added tests for `sycl::vec<T,2>` & `sycl::marray<T,2>`

Define constexpr bool for bfloat16 support

9681cc7

& a bit of tidying

Add marray fmin/fmax_nan to docs

a5e8e10

Extend isnan support to bfloat16

4497d82

& more tests

Formatting

1fe5980

Fix bug in cmul_add tests

847d45d

marray version never tested!

Add cmul_add<bfloat> & draft test

55e95ad

Function casts to sycl::complex<float> (no native bfloat16 support)

Enable relu for bfloat

5b0e0b5

Add bfloat16-container support to clamp

51860b3

i.e. sycl::marray<bfloat16, N>, sycl::vec<bfloat16, N>

Add max & min bfloat16 support & tests

d16dd56

Also reimplement `fmax_nan` & `fmin_nan` in terms of improved max/min

Generalize compare_mask & unordered_compare_mask

05b0748

- compare_mask & unordered_compare_mask support sycl::marray - sycl::marray tests for all `compare` APIs

Tidy a comment

fd9ebf4

Add bfloat16 support to cbrt

e5f9231

Revert "Add cmul_add<bfloat> & draft test"

f95dc85

This reverts commit 55e95ad.

joeatodd requested a review from a team as a code owner October 2, 2024 11:28

joeatodd temporarily deployed to WindowsCILock October 2, 2024 11:28 — with GitHub Actions Inactive

joeatodd had a problem deploying to WindowsCILock October 2, 2024 11:52 — with GitHub Actions Failure

ziranzha reviewed Oct 8, 2024

View reviewed changes

Alcpz suggested changes Oct 9, 2024

View reviewed changes

Merge branch 'sycl' into jtodd/bfloat16_support

d7b8676

joeatodd temporarily deployed to WindowsCILock October 14, 2024 11:43 — with GitHub Actions Inactive

joeatodd had a problem deploying to WindowsCILock October 14, 2024 12:15 — with GitHub Actions Failure

joeatodd added 2 commits October 15, 2024 11:59

Move std::common_type specialization to traits.hpp

4da9833

and remove stale comments

Document std::common_type_t<bfloat16,...>

ac2b4e9

joeatodd added 3 commits October 16, 2024 12:39

Revert unneeded cbrt<bfloat16> support

b0303bb

Review fix includes

fc11764

Assert bfloat16 math support in isnan, max, and min

f21f62e

joeatodd temporarily deployed to WindowsCILock October 16, 2024 11:41 — with GitHub Actions Inactive

joeatodd commented Oct 16, 2024

View reviewed changes

joeatodd requested a review from Alcpz October 16, 2024 11:45

joeatodd mentioned this pull request Oct 16, 2024

[SYCL][COMPAT] Add vectorized_ternary and vectorized_with_pred. #15550

Merged

joeatodd had a problem deploying to WindowsCILock October 16, 2024 12:05 — with GitHub Actions Failure

joeatodd added 2 commits October 16, 2024 15:15

Merge branch 'sycl' into jtodd/bfloat16_support

48e786a

Fix incorrect local memory usage in tests

20012b3

joeatodd temporarily deployed to WindowsCILock October 17, 2024 09:54 — with GitHub Actions Inactive

Alcpz approved these changes Oct 17, 2024

View reviewed changes

joeatodd temporarily deployed to WindowsCILock October 17, 2024 10:22 — with GitHub Actions Inactive

martygrant merged commit 1791115 into intel:sycl Oct 17, 2024
13 checks passed

[SYCL][COMPAT] Add bfloat16 support to several maths ops #15572

[SYCL][COMPAT] Add bfloat16 support to several maths ops #15572

Uh oh!

Conversation

joeatodd commented Oct 2, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Alcpz left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

joeatodd left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

joeatodd commented Oct 17, 2024

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

joeatodd commented Oct 2, 2024 •

edited

Loading

Alcpz left a comment •

edited

Loading