From 42e2b171ce1af772fe960f7f533e942c4246f09f Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Wed, 10 Aug 2022 15:18:28 +0100 Subject: [PATCH] bmad doesn't use marray. Signed-off-by: JackAKirk --- .../sycl/ext/oneapi/matrix/matrix-tensorcore.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp index b1dc9a2036a94..b16e902c3f2e1 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp @@ -168,7 +168,7 @@ struct joint_matrix< "For the matrix_use::a case, matrix_layout::row_major must " "be used for Bitwise MAD"); }; - int32_t wi_marray[1]; + int32_t data; }; template @@ -181,7 +181,7 @@ struct joint_matrix< "For the matrix_use::b case, matrix_layout::col_major must " "be used for Bitwise MAD"); }; - int32_t wi_marray[1]; + int32_t data; }; #undef __SYCL_JOINT_MATRIX_OVERLOAD @@ -418,10 +418,10 @@ struct joint_matrix_load_impl< matrix::precision::b1>::value) { int32_t *tileptr = reinterpret_cast(src.get()); if constexpr (NumRows == 8 && NumCols == 128) { - __bmma_m8n8k128_ld_a_b1(res.wi_marray, tileptr, stride, + __bmma_m8n8k128_ld_a_b1(&res.data, tileptr, stride, get_layout_id()); } else if constexpr (NumRows == 128 && NumCols == 8) { - __bmma_m8n8k128_ld_b_b1(res.wi_marray, tileptr, stride, + __bmma_m8n8k128_ld_b_b1(&res.data, tileptr, stride, get_layout_id()); } } @@ -803,14 +803,14 @@ struct joint_matrix_bmad_impl< sycl::bit_and>::value) { __bmma_m8n8k128_mma_and_popc_b1( - reinterpret_cast(&D.wi_marray), A.wi_marray, B.wi_marray, + reinterpret_cast(&D.wi_marray), &A.data, &B.data, reinterpret_cast(&C.wi_marray), 1); } else if constexpr (std::is_same< BinaryOperation, sycl::bit_xor>::value) { __bmma_m8n8k128_mma_xor_popc_b1( - reinterpret_cast(&D.wi_marray), A.wi_marray, B.wi_marray, + reinterpret_cast(&D.wi_marray), &A.data, &B.data, reinterpret_cast(&C.wi_marray), 1); } return D;