[SYCL][Matrix] Add support for tf32 type using the unified interface #8702

dkhaldi · 2023-03-20T18:17:26Z

No description provided.

JackAKirk · 2023-03-20T18:42:36Z

sycl/test/matrix/matrix-tf32-test.cpp

+             auto wi_data_a =
+                 sycl::ext::intel::experimental::matrix::get_wi_data(sg, sub_a);
+             for (int i = 0; i < wi_data_a.length(); i++) {
+               wi_data_a[i] = round_to_tf32(wi_data_a[i]);


nit, but shouldn't this type of usage of round_to_tf32 really use joint_matrix_apply?

Yes, I will use joint_matrix_apply for A and leave the loop way for B to show both usages.

JackAKirk · 2023-03-20T18:46:30Z

I think this should be fine from NVPTX point of view: will wait for confirmation from the cuda llvm-test-suite run.

JackAKirk · 2023-03-20T18:47:02Z

sycl/test/matrix/matrix-tf32-test.cpp

@@ -0,0 +1,164 @@
+// RUN: %clangxx -fsycl -O2 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 %s -o %t.out
+#include <iostream>


Shouldn't this kind of runtime test be placed in llvm-test-suite/SYCL/Matrix instead?

I will update the header to only check compilation

If we use -emit-llvm and FileCheck shouldn't we add checks for the actual IR generated by the compiler?

MrSidims

Changes are LGTM

MrSidims · 2023-03-22T13:24:25Z

sycl/test/matrix/matrix-tf32-test.cpp

@@ -0,0 +1,164 @@
+// RUN: %clangxx -fsycl -O2 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 %s -o %t.out
+#include <iostream>


If we use -emit-llvm and FileCheck shouldn't we add checks for the actual IR generated by the compiler?

dkhaldi · 2023-03-22T14:16:25Z

sycl/test/matrix/matrix-tf32-test.cpp

+           joint_matrix_apply(sg, sub_a, [=](float x) { x *= 2; });
+           for (int i = 0; i < wi_slice_a.length(); i++) {
+             float elem = wi_slice_a[i];
+             wi_slice_a[i] *= 2;


move this ewo line to ewo test in llvm-test-suite

tests are added here: intel/llvm-test-suite#1677

dkhaldi · 2023-03-22T19:04:15Z

@MrSidims, I am changing the name to __spirv_RoundFToTF32INTEL
Is this what you will be using?

yubingex007-a11y

LGTM

yubingex007-a11y · 2023-04-03T08:34:59Z

e2e testcases got approval in intel/llvm-test-suite#1677

yubingex007-a11y · 2023-04-03T10:53:30Z

ping?

dkhaldi · 2023-04-03T13:09:06Z

@yubingex007-a11y the build is failing here because the round function is not available yet in intel/llvm.
Will #8846 solve the build issue?

yubingex007-a11y · 2023-04-04T07:55:14Z

the testcases fail are not related with this pr i think.

yubingex007-a11y · 2023-04-04T07:55:20Z

ping?

steffenlarsen

Small naming change request, but otherwise LGTM.

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

dkhaldi · 2023-04-05T13:08:13Z

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

+template <typename T> struct helper_traits {
+  using element_type = T;
+  using storage_element_type = T;
+  using fill_argument_type = T;


remove fill_argument_type as it is currently unused

steffenlarsen

LGTM!

yubingex007-a11y · 2023-04-06T05:14:12Z

@steffenlarsen Hi, can you help us merge this patch

steffenlarsen · 2023-04-06T09:59:25Z

@yubingex007-a11y - Could you please push another merge commit to rerun pre-commit tests? CI issues should hopefully have been addressed.

Nuullll · 2023-04-10T03:23:12Z

sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp

+//===----------------------------------------------------------------------===//
+// REQUIRES: matrix
+
+// RUN: %clangxx -fsycl %s -o %t.out -Dsycl/test-e2e_EXT_ONEAPI_MATRIX_VERSION=4


Should be -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4

Nuullll · 2023-04-10T03:23:21Z

sycl/test-e2e/Matrix/joint_matrix_tf32.cpp

+//===----------------------------------------------------------------------===//
+// REQUIRES: matrix
+
+// RUN: %clangxx -fsycl %s -o %t.out -Dsycl/test-e2e_EXT_ONEAPI_MATRIX_VERSION=4


Should be -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4

Nuullll · 2023-04-10T03:24:02Z

sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp

+             auto wi_data_b =
+                 sycl::ext::intel::experimental::matrix::get_wi_data(sg, sub_b);
+             for (int i = 0; i < wi_data_b.length(); i++) {
+               wi_data_b[i] = round_to_tf32(wi_data_b[i]);


No known conversion from wi_element<> to float &

[SYCL][Matrix] Add support for tf32 type using the unified interface

de3d230

dkhaldi marked this pull request as ready for review March 20, 2023 18:19

dkhaldi requested a review from a team as a code owner March 20, 2023 18:19

dkhaldi requested a review from bso-intel March 20, 2023 18:19

dkhaldi mentioned this pull request Mar 20, 2023

[SYCL][Matrix] Add support for tf32 type using the unified interface #8151

Closed

dkhaldi requested review from yubingex007-a11y, JackAKirk and MrSidims March 20, 2023 18:20

JackAKirk reviewed Mar 20, 2023

View reviewed changes

Address Jack's comments

88aad2c

dkhaldi temporarily deployed to aws March 20, 2023 20:50 — with GitHub Actions Inactive

MrSidims reviewed Mar 22, 2023

View reviewed changes

dkhaldi commented Mar 22, 2023

View reviewed changes

change convert to round in the name

a56084d

dkhaldi temporarily deployed to aws March 22, 2023 19:25 — with GitHub Actions Inactive

Merge remote-tracking branch 'intel_llvm/sycl' into tf32-joint-matrix

f7d7c4e

dkhaldi temporarily deployed to aws March 23, 2023 18:12 — with GitHub Actions Inactive

yubingex007-a11y approved these changes Mar 24, 2023

View reviewed changes

set xfail to test as SPIRV changes are not part of intel/llvm yet

0d23b8a

dkhaldi temporarily deployed to aws March 24, 2023 14:12 — with GitHub Actions Inactive

correct SYCL_EXTERNAL with the new naming __DPCPP_SYCL_EXTERNAL

7d05e6a

dkhaldi temporarily deployed to aws March 28, 2023 13:13 — with GitHub Actions Inactive

yubingex007-a11y self-assigned this Apr 3, 2023

yubingex007-a11y added 2 commits April 3, 2023 16:20

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

c162139

move e2e testcases from llvm-test-suite's pr

1f84b04

yubingex007-a11y requested a review from steffenlarsen April 3, 2023 11:05

yubingex007-a11y temporarily deployed to aws April 3, 2023 12:20 — with GitHub Actions Inactive

yubingex007-a11y added 2 commits April 4, 2023 14:02

rm XFAIL from sycl/test/matrix/matrix-tf32-test.cpp

93193e9

fix clang-format issue

8caee53

yubingex007-a11y temporarily deployed to aws April 4, 2023 06:51 — with GitHub Actions Inactive

yubingex007-a11y temporarily deployed to aws April 4, 2023 07:24 — with GitHub Actions Inactive

steffenlarsen reviewed Apr 4, 2023

View reviewed changes

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

dkhaldi commented Apr 5, 2023

View reviewed changes

address comments

1557f34

steffenlarsen approved these changes Apr 5, 2023

View reviewed changes

JackAKirk self-requested a review April 5, 2023 13:45

JackAKirk approved these changes Apr 5, 2023

View reviewed changes

yubingex007-a11y temporarily deployed to aws April 5, 2023 14:08 — with GitHub Actions Inactive

yubingex007-a11y temporarily deployed to aws April 5, 2023 14:40 — with GitHub Actions Inactive

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

19ada00

yubingex007-a11y temporarily deployed to aws April 6, 2023 10:39 — with GitHub Actions Inactive

yubingex007-a11y temporarily deployed to aws April 6, 2023 11:11 — with GitHub Actions Inactive

steffenlarsen merged commit aba6d85 into intel:sycl Apr 6, 2023

Nuullll reviewed Apr 10, 2023

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL][Matrix] Add support for tf32 type using the unified interface #8702

[SYCL][Matrix] Add support for tf32 type using the unified interface #8702

dkhaldi commented Mar 20, 2023

JackAKirk Mar 20, 2023

dkhaldi Mar 20, 2023

JackAKirk commented Mar 20, 2023

JackAKirk Mar 20, 2023

dkhaldi Mar 20, 2023

MrSidims Mar 22, 2023

MrSidims left a comment

MrSidims Mar 22, 2023

dkhaldi Mar 22, 2023 •

edited

Loading

dkhaldi Mar 22, 2023

dkhaldi commented Mar 22, 2023

yubingex007-a11y left a comment

yubingex007-a11y commented Apr 3, 2023

yubingex007-a11y commented Apr 3, 2023

dkhaldi commented Apr 3, 2023

yubingex007-a11y commented Apr 4, 2023

yubingex007-a11y commented Apr 4, 2023

steffenlarsen left a comment

dkhaldi Apr 5, 2023

steffenlarsen left a comment

yubingex007-a11y commented Apr 6, 2023

steffenlarsen commented Apr 6, 2023

Nuullll Apr 10, 2023

Nuullll Apr 10, 2023

Nuullll Apr 10, 2023

		@@ -0,0 +1,164 @@
		// RUN: %clangxx -fsycl -O2 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 %s -o %t.out
		#include <iostream>

[SYCL][Matrix] Add support for tf32 type using the unified interface #8702

[SYCL][Matrix] Add support for tf32 type using the unified interface #8702

Conversation

dkhaldi commented Mar 20, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

JackAKirk commented Mar 20, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MrSidims left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dkhaldi Mar 22, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dkhaldi commented Mar 22, 2023

yubingex007-a11y left a comment

Choose a reason for hiding this comment

yubingex007-a11y commented Apr 3, 2023

yubingex007-a11y commented Apr 3, 2023

dkhaldi commented Apr 3, 2023

yubingex007-a11y commented Apr 4, 2023

yubingex007-a11y commented Apr 4, 2023

steffenlarsen left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

steffenlarsen left a comment

Choose a reason for hiding this comment

yubingex007-a11y commented Apr 6, 2023

steffenlarsen commented Apr 6, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dkhaldi Mar 22, 2023 •

edited

Loading