[SYCL] Move bfloat support from experimental to supported. #6524

rdeodhar · 2022-08-03T22:10:31Z

This change makes bfloat16 a supported feature.

Signed-off-by: Rajiv Deodhar rajiv.deodhar@intel.com

Signed-off-by: Rajiv Deodhar <rajiv.deodhar@intel.com>

JackAKirk · 2022-08-04T14:34:56Z

Looks good, do we want to also move these bfloat16 math functions out of experimental also:

llvm/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp

Line 160 in bdd88e5

std::enable_if_t<std::is_same<T, bfloat16>::value, T> fmin(T x, T y) {

since they are defined in the same extension document as the main bfloat16 class?

btw there should be an accompanying PR to intel/llvm-test-suite updating corresponding tests, otherwise there will be lots of failures: For example in this test: https://github.com/intel/llvm-test-suite/blob/intel/SYCL/BFloat16/bfloat16_type.hpp.

JackAKirk · 2022-08-04T14:36:15Z

Also FYI there is another open PR updating the bfloat16 class that it might be good to consider merging first #6492.

rdeodhar · 2022-08-04T20:41:21Z

/verify with intel/llvm-test-suite#1129

steffenlarsen

Overall LGTM, though I think the comment should be addressed first.

Also, should we consider leaving a deprecated version of bfloat16.hpp in the experimental folder to warn old users? Since the feature was experimental I don't think we strictly have to, but we could.

steffenlarsen · 2022-08-19T14:50:42Z

sycl/doc/extensions/supported/sycl_ext_oneapi_bfloat16.asciidoc

@@ -408,4 +406,5 @@ Compute absolute value of a `bfloat16`.
 |3|2021-08-18|Alexey Sotkin |Remove `uint16_t` constructor
 |4|2022-03-07|Aidan Belton and Jack Kirk |Switch from Intel vendor specific to oneapi
 |5|2022-04-05|Jack Kirk | Added section for bfloat16 math builtins
+|6|2022-08-03|Alexey Sotkin |Add `operator sycl::half()`


I think it should be your name here. 😄

It is carried over from that author, but I agree that more changes have been made, so changing name.

gmlueck · 2022-08-23T16:12:02Z

I also have two global comments:

Please update the specification to use the template in (https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/template.asciidoc) and follow the instructions in (https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/README-process.md). I would like all new "supported" extensions to use the new template format.
There was a request from the ML team that the bfloat16 type should not be an "optional feature". Instead, they proposed that it should be allowed for any device. Devices without native support should use a fall-back routine for conversion. Couldn't we do this by following the pattern of the other math operations we have in "llvm/libdevice"? If we do this, there is no need for the aspect, because application code wouldn't need to check if a device supports bloat16 before using it.

gmlueck · 2022-08-23T14:53:40Z

sycl/doc/extensions/supported/sycl_ext_oneapi_bfloat16.asciidoc

  operator float() const;
+  operator sycl::half() const;


I like this conversion to sycl::half. However, we should also add the opposite conversion from sycl::half to bfloat16:

bfloat16(const sycl::half &a); bfloat16 &operator=(const sycl::half &a);

Do we also need conversion to / from double?

This PR is intended to move the current bfloat16 support out of experimental space. Any changes to the level of bfloat16 support can be done in future PRs.

On Intel platforms the bfloat16 to/from float is done using the __spirv_ConvertBF16ToFINTELoperator. I suspect a double version of that does not exist.
Float to double conversion can be made in the usual C++ way more efficiently in hardware. A direct version of bfloat16 to double conversion in software will involve more bit twiddling than the float conversion where only trailing 0 bits of fraction need to be inserted.

The sycl::half class includes conversions to/from float. Those kick in when bfloat16 is used with sycl::half, so conversions between bfloat16 and sycl::half are not needed.

Are you saying that we should remove this conversion from bfloat16 to sycl::half?

Yes, its not needed.

This item was revisited and it turns out that sycl::half <-> bfloat16 conversions are needed. They have been added.

Sorry for joining the discussion late. May be it's a nitpick, but should we tell, that conversion half <-> bfloat16 follows IEEE 754 float <-> half conversion? In other words, what happens, if bfloat16 value overflows half range? Also are we adding last 3 fraction bits stochastically or they are guarantied to be zero (or it's implementation detail)?

sycl/doc/extensions/supported/sycl_ext_oneapi_bfloat16.asciidoc

rdeodhar · 2022-08-24T22:44:58Z

/verify with intel/llvm-test-suite#1129

rdeodhar · 2022-08-25T21:54:24Z

/run with intel/llvm-test-suite#1129

rdeodhar · 2022-08-26T01:50:54Z

/verify with intel/llvm-test-suite#1129

rdeodhar · 2022-08-26T23:01:26Z

/verify with intel/llvm-test-suite#1129

gmlueck

This is looking a lot better, just a few more comments below.

sycl/doc/extensions/supported/sycl_ext_oneapi_bfloat16.asciidoc

rdeodhar · 2022-08-31T00:30:25Z

/verify with intel/llvm-test-suite#1129

gmlueck

Added some minor doc comments below. I think the main remaining issue is that the aspect needs to be enabled.

sycl/doc/extensions/supported/sycl_ext_oneapi_bfloat16.asciidoc

JackAKirk · 2022-11-29T10:10:00Z

Precommit CI seems to be failing after this as well. See for example #7563.

Summary of current failures:

1. esimd/regression/windows_build_test.cpp in Windows post-commit - Addressed in [[SYCL] Remove experimental from bfloat16 in Windows test #7569](https://github.com/intel/llvm/pull/7569).

2. extensions/bfloat16.cpp in no-assert Linux - No current PRs addressing this.

3. ESIMD test-suite failures in CI - Addressed in [[SYCL] Correct bfloat16 namespace in ESIMD and matrix tests llvm-test-suite#1422](https://github.com/intel/llvm-test-suite/pull/1422).

4. CUDA test-suite compilation failures - Addressed in [[SYCL] Reintroduce experimental bfloat16 math functions #7567](https://github.com/intel/llvm/pull/7567) and [[SYCL] Fix CUDA tests using bfloat16 llvm-test-suite#1421](https://github.com/intel/llvm-test-suite/pull/1421).

5. `PI_ERROR_INVALID_BINARY` in test-suite bfloat16_type.cpp on CUDA - No current PRs addressing this.

Line 5 of bloat16_type.cpp looks like it is missing the -fsycl-targets=%sycl_triple ?:

// RUN: %clangxx -fsycl %s -o %t.out

steffenlarsen · 2022-11-29T10:30:58Z

Line 5 of bloat16_type.cpp looks like it is missing the -fsycl-targets=%sycl_triple ?:

// RUN: %clangxx -fsycl %s -o %t.out

You're absolutely right, it seems like we have a separate CUDA test for this since it has additional requirements, so I think is better to just keep these separate for now. I have opened a PR for this in intel/llvm-test-suite#1423.

pvchupin · 2022-11-29T16:22:58Z

@steffenlarsen should we revert this change?

intel/llvm#6524 moved bfloat16 out of the experimental namespace. This commit removes the last remaining uses of the experimental namespace in bfloat16 for ESIMD and matrix tests. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen · 2022-11-29T16:28:02Z

@steffenlarsen should we revert this change?

Tests seem to be resolvable, but in order to make sure we need the next nightly. Once all the mentioned patches are merged we should only have a check-sycl failure on no-assert to also address, which should only affect post-commit so I suggest we keep it in unless other issues pop up.

pvchupin · 2022-11-29T16:33:02Z

@steffenlarsen, thanks a lot for handling bunch of these!!!
ping @rdeodhar for remaining one.

JackAKirk · 2022-11-29T16:35:37Z

@steffenlarsen, thanks a lot for handling bunch of these!!! ping @rdeodhar for remaining one.

@steffenlarsen has also dealt with this:

" 5. PI_ERROR_INVALID_BINARY in test-suite bfloat16_type.cpp on CUDA - No current PRs addressing this.

"

in
intel/llvm-test-suite#1423

#6524 accidentally removed the experimental bfloat16 math functions while moving bfloat16 out of the experimental namespace. This commit reintroduces these in the bfloat16_math.hpp header file. Changes to sub_group.hpp are to resolve detail namespace ambiguities are are NFC. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Test was modified at intel#6524 Change fixes post-commit issue in no-asserts mode

Test was modified at #6524 Change fixes post-commit issue in no-asserts mode

…ntel#9143) The error in LIT test esimd/intel_fp16_converts.cpp is caused by intel#6524 which - moved 'bfloat16' out of 'experimental' namespace - created a wrapper __devicelib_ConvertBF16ToFINTEL() which simply calls __spirv_ConvertBF16ToFINTEL() The fix - create a test for bfloat16 conversions. - allows __devicelib_ConvertBF16ToFINTEL() and __devicelib_ConvertFToBF16INTEL() for ESIMD context. Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

#7981) …9143) The error in LIT test esimd/intel_fp16_converts.cpp is caused by #6524 which - moved 'bfloat16' out of 'experimental' namespace - created a wrapper __devicelib_ConvertBF16ToFINTEL() which simply calls __spirv_ConvertBF16ToFINTEL() The fix - creates a test for bfloat16 conversions. - allows __devicelib_ConvertBF16ToFINTEL() and __devicelib_ConvertFToBF16INTEL() for ESIMD context. Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

…YCL (#8257) This PR addresses an issue where if we use `__CUDA_ARCH__` causes intrinsics not to be defined in the CUDA include files. - Replace `__CUDA_ARCH__` with `__SYCL_CUDA_ARCH__` for SYCL device - Update the `sycl-macro.cpp` test to check the appropriate macro. --- As far as I could find the original issue was introduced from PR [#6524](7b47ebb) for enabling the bfloat16 support moving it from the experimental extension, and it breaks some codebases with CUDA interop calls. Current reports include github issues [#7722](#7722), [#8133](#8133) and [oneapi-src/oneMKL#257](oneapi-src/oneMKL#257). For that reason we define a unique `__SYCL_CUDA_ARCH__` macro and use it instead for SYCL device targets and leave `__CUDA_ARCH__` as before for CUDA targets.

…ntal status. (intel#1129) Tests changes for intel#6524 Signed-off-by: Rajiv Deodhar <rajiv.deodhar@intel.com> Co-authored-by: JackAKirk <jack.kirk@codeplay.com>

…ntal status. (intel/llvm-test-suite#1129) Tests changes for intel#6524 Signed-off-by: Rajiv Deodhar <rajiv.deodhar@intel.com> Co-authored-by: JackAKirk <jack.kirk@codeplay.com>

…vm-test-suite#1422) intel#6524 moved bfloat16 out of the experimental namespace. This commit removes the last remaining uses of the experimental namespace in bfloat16 for ESIMD and matrix tests. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

rdeodhar added 2 commits August 3, 2022 15:09

[SYCL] Move bfloat support from experimental to supported.

6014cef

Signed-off-by: Rajiv Deodhar <rajiv.deodhar@intel.com>

Corrections to tests.

bdd88e5

rdeodhar mentioned this pull request Aug 4, 2022

[SYCL] Test corrections after moving bfloat16 support out of experimental status. intel/llvm-test-suite#1129

Merged

rdeodhar marked this pull request as ready for review August 9, 2022 16:43

rdeodhar requested review from a team as code owners August 9, 2022 16:43

rdeodhar requested a review from steffenlarsen August 9, 2022 16:43

steffenlarsen reviewed Aug 19, 2022

View reviewed changes

gmlueck reviewed Aug 23, 2022

View reviewed changes

rdeodhar added 2 commits August 24, 2022 09:24

Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16

73ed541

Moved another file out of experimental space.

0fe1884

Responses to review comments.

feb9d5f

Removed unneeded sycl::half conversion and updated doc.

129f53f

rdeodhar requested a review from gmlueck August 26, 2022 23:02

Added conversion from sycl::half to bfloat16.

2115f09

gmlueck reviewed Aug 30, 2022

View reviewed changes

sycl/doc/extensions/supported/sycl_ext_oneapi_bfloat16.asciidoc Outdated Show resolved Hide resolved

Cleanup of documentation.

3c2eb80

rdeodhar requested a review from gmlueck August 31, 2022 00:15

gmlueck reviewed Aug 31, 2022

View reviewed changes

Hooked up bfloat16 aspect within OpenCL plugin.

74aa175

steffenlarsen mentioned this pull request Nov 29, 2022

[SYCL] Support __builtin_printf for SYCL device #7483

Merged

JackAKirk mentioned this pull request Nov 29, 2022

[SYCL][CUDA] Add SM version check to bfloat16 CUDA test intel/llvm-test-suite#1423

Merged

pvchupin mentioned this pull request Dec 1, 2022

[SYCL] Update bfloat16.cpp test to pass in no-assert mode #7602

Merged

pvchupin pushed a commit to pvchupin/llvm that referenced this pull request Dec 1, 2022

[SYCL] Update bfloat16.cpp test to pass in no-assert mode

b89376d

Test was modified at intel#6524 Change fixes post-commit issue in no-asserts mode

steffenlarsen pushed a commit that referenced this pull request Dec 1, 2022

[SYCL] Update bfloat16.cpp test to pass in no-assert mode (#7602)

9ff6045

Test was modified at #6524 Change fixes post-commit issue in no-asserts mode

AuroraPerego mentioned this pull request Dec 9, 2022

[SYCL][CUDA] __CUDA_ARCH__ defined when compiling for CUDA backend since sycl-nightly/20221129 #7722

Closed

yubingex007-a11y added a commit to yubingex007-a11y/llvm that referenced this pull request Jan 3, 2023

[Matrix] Fix the testcase issue brought by intel#6524

3b3fa8b

bader pushed a commit that referenced this pull request Jan 3, 2023

[Matrix] Fix the testcase issue brought by #6524 (#7902)

55e6f3d

v-klochkov mentioned this pull request Jan 11, 2023

[ESIMD] Fix errors caused by move of bfloat16 from experimental ns (#… #7981

Merged

JackAKirk mentioned this pull request Oct 9, 2023

Use sycl::bfloat16 class and functions instead of float casts. oneapi-src/SYCLomatic#1341

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Move bfloat support from experimental to supported. #6524

[SYCL] Move bfloat support from experimental to supported. #6524

rdeodhar commented Aug 3, 2022

JackAKirk commented Aug 4, 2022

JackAKirk commented Aug 4, 2022

rdeodhar commented Aug 4, 2022

steffenlarsen left a comment

steffenlarsen Aug 19, 2022

rdeodhar Aug 25, 2022

gmlueck commented Aug 23, 2022

gmlueck Aug 23, 2022

rdeodhar Aug 24, 2022

rdeodhar Aug 25, 2022

rdeodhar Aug 26, 2022

gmlueck Aug 26, 2022

rdeodhar Aug 26, 2022

rdeodhar Aug 30, 2022

MrSidims Sep 20, 2022 •

edited

Loading

rdeodhar commented Aug 24, 2022

rdeodhar commented Aug 25, 2022

rdeodhar commented Aug 26, 2022

rdeodhar commented Aug 26, 2022

gmlueck left a comment

rdeodhar commented Aug 31, 2022

gmlueck left a comment

JackAKirk commented Nov 29, 2022

steffenlarsen commented Nov 29, 2022

pvchupin commented Nov 29, 2022

steffenlarsen commented Nov 29, 2022

pvchupin commented Nov 29, 2022

JackAKirk commented Nov 29, 2022

[SYCL] Move bfloat support from experimental to supported. #6524

[SYCL] Move bfloat support from experimental to supported. #6524

Conversation

rdeodhar commented Aug 3, 2022

JackAKirk commented Aug 4, 2022

JackAKirk commented Aug 4, 2022

rdeodhar commented Aug 4, 2022

steffenlarsen left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

gmlueck commented Aug 23, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MrSidims Sep 20, 2022 • edited Loading

Choose a reason for hiding this comment

rdeodhar commented Aug 24, 2022

rdeodhar commented Aug 25, 2022

rdeodhar commented Aug 26, 2022

rdeodhar commented Aug 26, 2022

gmlueck left a comment

Choose a reason for hiding this comment

rdeodhar commented Aug 31, 2022

gmlueck left a comment

Choose a reason for hiding this comment

JackAKirk commented Nov 29, 2022

steffenlarsen commented Nov 29, 2022

pvchupin commented Nov 29, 2022

steffenlarsen commented Nov 29, 2022

pvchupin commented Nov 29, 2022

JackAKirk commented Nov 29, 2022

MrSidims Sep 20, 2022 •

edited

Loading