Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[SYCL][Matrix] Fix issues in bf16 test that were causing failed execu… #535

Merged
merged 4 commits into from
Nov 11, 2021
Merged

[SYCL][Matrix] Fix issues in bf16 test that were causing failed execu… #535

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
5467e1c
[SYCL][Matrix] Fix issues in bf16 test that were causing failed execu…
dkhaldi Oct 26, 2021
995b6fd
[SYCL][Matrix]fix clang-format issue
dkhaldi Nov 8, 2021
79b2665
[SYCL][Matrix] clang-format issue fix
dkhaldi Nov 8, 2021
8fd0f41
[SYCL][Matrix] Fix clang-format issue
dkhaldi Nov 8, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 7 additions & 11 deletions SYCL/Matrix/joint_matrix_bf16.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,18 +70,14 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
const auto sg_startx = global_idx - spmd_item.get_local_id(0);
const auto sg_starty = global_idy - spmd_item.get_local_id(1);

ext::oneapi::sub_group sg = spmd_item.get_sub_group();
joint_matrix<ext::oneapi::sub_group, unsigned short, TM, TK> sub_a(
sg);
sub_group sg = spmd_item.get_sub_group();
joint_matrix<unsigned short, TM, TK> sub_a(sg);
// For B, since current implementation does not support non-packed
// layout, users need to specify the updated VNNI sizes along with
// the packed_b layout. By default, the layout is row_major and size
// is (TK, TN).
joint_matrix<ext::oneapi::sub_group, unsigned short, TK, TN,
matrix_layout::packed_b>
sub_b(sg);
joint_matrix<ext::oneapi::sub_group, float, TM, TN> sub_c(sg);

// layout, users need to specify the packed_b layout.
// By default, the layout is row_major
joint_matrix<unsigned short, TK, TN, matrix_layout::packed_b> sub_b(
sg);
joint_matrix<float, TM, TN> sub_c(sg);
joint_matrix_load(sg, sub_c,
accC.get_pointer() + (sg_startx * TM) * N +
sg_starty / SG_SZ * TN,
Expand Down