[Matrix][SYCL] Add bfloat16 support for joint_matrix #5566

yubingex007-a11y · 2022-02-14T03:47:58Z

Previously we use unsigned short to represent bfloat16, now we give formal support for sycl/s bfloat16 in joint_matrix implementation.

yubingex007-a11y · 2022-02-14T03:50:26Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit.hpp

@@ -11,6 +11,7 @@
 #include <CL/__spirv/spirv_ops.hpp>
 #include <CL/sycl/detail/defines_elementary.hpp>
 #include <CL/sycl/feature_test.hpp>
+#include <sycl/ext/intel/experimental/bfloat16.hpp>


Buildfail will happen since

class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16

and sycl_detail::uses_aspects is unsupported code in g++ which is used for intel/llvm's build

yubingex007-a11y · 2022-02-14T03:56:17Z

we are not able to let a end-to-end testcase pass since __spirv_ConvertBF16ToFINTEL is not supported in OCL CPU Backend

sycl/include/sycl/ext/oneapi/matrix/matrix-jit.hpp

AlexeySotkin · 2022-02-15T16:16:41Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit.hpp

@@ -737,6 +738,160 @@ class wi_element<uint16_t, NumRows, NumCols, Layout, Group> {
  }
 };

+template <size_t NumRows, size_t NumCols, matrix_layout Layout, typename Group>
+class wi_element<sycl::ext::intel::experimental::bfloat16, NumRows, NumCols,


Nit: probably alias like

using bfloat16 = sycl::ext::intel::experimental::bfloat16

could improve readability of this code.

we can't use "using" in header file

we can't use "using" in header file

What do you mean? I guess you can use them locally if it is in a scope not visible from the end-user.

dkhaldi

LGTM

yubingex007-a11y · 2022-04-13T13:22:27Z

ping? @intel/llvm-reviewers-runtime

steffenlarsen · 2022-04-13T13:39:15Z

Do we have tests for this?

yubingex007-a11y · 2022-04-21T13:52:49Z

I don't think the failed testcase has relationship with this PR.

dkhaldi · 2022-04-21T15:53:50Z

sycl/test/matrix/matrix-bfloat16-test.cpp

+  return *res;
+}
+
+unsigned short make_bf16(float x) {


Do we still need these functions?
There must be some equivalent DPC++/SPIRV ones, right?

in bfloat16 class, we don't have bf16<->float conversion in host device. so this make_bf16 is needed in line144

in bfloat16 class, we don't have bf16<->float conversion in host device. so this make_bf16 is needed in line144

I'd suggest to finish #5954 and reuse host version of the conversion.

sounds good.
@YetAnotherCompilerEngineer who are you? :)

in bfloat16 class, we don't have bf16<->float conversion in host device. so this make_bf16 is needed in line144

I'd suggest to finish #5954 and reuse host version of the conversion.

Could we merge the patch before #5954? i think we can create another PR for testcase after #5954

Could we merge the patch before #5954?

Sure, please merge them in an order you like the most :)

dkhaldi

LGTM

keryell · 2022-04-22T01:27:43Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit.hpp

+    (void)lhs;                                                                 \
+    (void)rhs;                                                                 \


Instead just remove the names from the op() declaration

i didn't get your point. actually this code is for host. the param list of friend type operator op should be the same as line555 which is for device code.

I believe what @keryell suggests is to change this to

friend type operator op( \ const wi_element<sycl::ext::oneapi::experimental::bfloat16, NumRows, \ NumCols, Layout, Group> &, \ const sycl::ext::oneapi::experimental::bfloat16 &) { \ throw runtime_error("joint matrix is not supported on host device.", \ PI_INVALID_DEVICE); \ }

which preserves the signature but doesn't have unused arguments as they are unnamed.

keryell · 2022-04-22T01:28:35Z

sycl/test/matrix/matrix-bfloat16-test.cpp

@@ -0,0 +1,192 @@
+// RUN: %clangxx -fsycl -O2 %s -o %t.out
+#include <CL/sycl.hpp>


Suggested change

#include <CL/sycl.hpp>

#include <sycl/sycl.hpp>

it is "include/sycl/CL/sycl.hpp", how could we change into "#include <sycl/sycl.hpp>"

keryell · 2022-04-22T01:29:21Z

sycl/test/matrix/matrix-bfloat16-test.cpp

+#if (SYCL_EXT_ONEAPI_MATRIX == 2)
+#include <iostream>
+
+using namespace sycl;


It is better to avoid.
If so, why the sycl:: on the 2 next lines? :-)

keryell · 2022-04-22T01:29:49Z

sycl/test/matrix/matrix-bfloat16-test.cpp

+using namespace sycl::ext::oneapi::experimental::matrix;
+using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
+
+#define TILE_SZ 16


constexpr everywhere instead of macros.

keryell · 2022-04-23T00:58:34Z

sycl/test/matrix/matrix-bfloat16-test.cpp

+using namespace sycl::ext::oneapi::experimental::matrix;
+using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
+
+static constexpr size_t TILE_SZ = 16;


Since your macro was not using any type at the first place, you could just use auto instead of size_t.
Because with the macro the numbers we interpreted as 32-bit int while here you are forcing 64-bit, which might be somewhat less efficient in some case (but not in this test I guess).

keryell

Thanks!

yubingex007-a11y · 2022-05-05T06:51:08Z

ping? @steffenlarsen @dkhaldi

dkhaldi

LGTM

steffenlarsen

LGTM!

yubingex007-a11y · 2022-05-05T16:08:18Z

@steffenlarsen could you help us merge it?

steffenlarsen · 2022-05-05T16:32:58Z

Unexpectedly Passed Tests (1):
SYCL :: api/simd_view_select_2d_int.cpp

This was since fixed with intel/llvm-test-suite#1008

This reverts commit 80f29bd.

…#6108) This reverts commit 80f29bd.

yubingex007-a11y · 2022-05-06T08:10:08Z

i think i know what happen in https://github.com/intel/llvm/runs/6314339242?check_suite_focus=true.
Since unused arg caused the compfail and abort the build , llvm-spirv was not built so we can find it in bin/*. so tools/spirv-to-ir-wrapper/spirv-to-ir-wrapper.ll failed due to"llvm-spirv: command not found"

#6111 is to fix the unused variable issue.

[Matrix][SYCL] Add bfloat16 support for wi_slice

aa5cf45

yubingex007-a11y commented Feb 14, 2022

View reviewed changes

yubingex007-a11y requested a review from dkhaldi February 14, 2022 03:50

small fix

c310976

dkhaldi requested a review from AlexeySotkin February 14, 2022 14:23

AlexeySotkin reviewed Feb 15, 2022

View reviewed changes

sycl/include/sycl/ext/oneapi/matrix/matrix-jit.hpp Show resolved Hide resolved

AlexeySotkin reviewed Feb 15, 2022

View reviewed changes

yubingex007-a11y added 2 commits April 8, 2022 22:20

Merge remote-tracking branch 'intel_llvm/sycl' into addsyclbf16matrix

7b8b37e

small rebase

37d75f5

yubingex007-a11y changed the title ~~[Matrix][SYCL] Add bfloat16 support for wi_slice~~ [Matrix][SYCL] Add bfloat16 support for joint_matrix Apr 11, 2022

yubingex007-a11y marked this pull request as ready for review April 11, 2022 15:20

yubingex007-a11y requested a review from a team as a code owner April 11, 2022 15:20

yubingex007-a11y requested a review from steffenlarsen April 11, 2022 15:20

dkhaldi previously approved these changes Apr 11, 2022

View reviewed changes

add a testcase for bfloat16

a6a7c82

yubingex007-a11y dismissed dkhaldi’s stale review via a6a7c82 April 21, 2022 08:24

yubingex007-a11y requested a review from dkhaldi April 21, 2022 08:25

dkhaldi reviewed Apr 21, 2022

View reviewed changes

dkhaldi previously approved these changes Apr 21, 2022

View reviewed changes

keryell reviewed Apr 22, 2022

View reviewed changes

address Keryell's comments

6bc53d5

yubingex007-a11y dismissed dkhaldi’s stale review via 6bc53d5 April 22, 2022 09:33

change macro to constexpr

b5c1194

keryell reviewed Apr 23, 2022

View reviewed changes

small changes

a42e3aa

keryell approved these changes Apr 24, 2022

View reviewed changes

yubingex007-a11y requested a review from dkhaldi May 5, 2022 06:51

dkhaldi approved these changes May 5, 2022

View reviewed changes

steffenlarsen approved these changes May 5, 2022

View reviewed changes

steffenlarsen merged commit 80f29bd into intel:sycl May 5, 2022

pvchupin pushed a commit that referenced this pull request May 6, 2022

Revert "[Matrix][SYCL] Add bfloat16 support for joint_matrix (#5566)"

859da54

This reverts commit 80f29bd.

pvchupin mentioned this pull request May 6, 2022

Revert "[Matrix][SYCL] Add bfloat16 support for joint_matrix" #6108

Merged

pvchupin pushed a commit that referenced this pull request May 6, 2022

Revert "[Matrix][SYCL] Add bfloat16 support for joint_matrix (#5566)" (…

7461b63

…#6108) This reverts commit 80f29bd.

This was referenced May 6, 2022

[Matrix][SYCL] Add bfloat16 support for joint_matrix #6112

Closed

[Matrix][SYCL] Add bfloat16 support for joint_matrix #6113

Merged

		@@ -0,0 +1,192 @@
		// RUN: %clangxx -fsycl -O2 %s -o %t.out
		#include <CL/sycl.hpp>

[Matrix][SYCL] Add bfloat16 support for joint_matrix #5566

[Matrix][SYCL] Add bfloat16 support for joint_matrix #5566

Uh oh!

Conversation

yubingex007-a11y commented Feb 14, 2022 • 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!

yubingex007-a11y commented Feb 14, 2022

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

yubingex007-a11y Feb 17, 2022 • 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!

keryell Apr 24, 2022 • 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!

dkhaldi left a comment

Choose a reason for hiding this comment

Uh oh!

yubingex007-a11y commented Apr 13, 2022

Uh oh!

steffenlarsen commented Apr 13, 2022

Uh oh!

yubingex007-a11y commented Apr 21, 2022

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!

dkhaldi left a comment

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!

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!

keryell left a comment

Choose a reason for hiding this comment

Uh oh!

yubingex007-a11y commented May 5, 2022

Uh oh!

dkhaldi left a comment

Choose a reason for hiding this comment

Uh oh!

steffenlarsen left a comment

Choose a reason for hiding this comment

Uh oh!

yubingex007-a11y commented May 5, 2022

Uh oh!

steffenlarsen commented May 5, 2022

Uh oh!

yubingex007-a11y commented May 6, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

yubingex007-a11y commented Feb 14, 2022 •

edited

Loading

yubingex007-a11y Feb 17, 2022 •

edited

Loading

keryell Apr 24, 2022 •

edited

Loading

yubingex007-a11y commented May 6, 2022 •

edited

Loading