[SYCL][Matrix] syntax changes as preparation before moving joint matrix from experimental namespace #11215
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -413,7 +413,7 @@ void store_layoutT( | |
| template <typename T, size_t NumRows, size_t NumCols, | ||
| access::address_space Space, access::decorated IsDecorated> | ||
| void joint_matrix_store_cuda( | ||
| const joint_matrix_cuda< | ||
| T, sycl::ext::oneapi::experimental::matrix::use::accumulator, NumRows, | ||
| NumCols, sycl::ext::oneapi::experimental::matrix::layout::dynamic> &src, | ||
| multi_ptr<T, Space, IsDecorated> dst, size_t stride, | ||
|
|
@@ -482,11 +482,11 @@ void joint_matrix_mad_cuda( | |
| joint_matrix_cuda< | ||
| Tc, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, | ||
| sycl::ext::oneapi::experimental::matrix::layout::dynamic> &D, | ||
| const joint_matrix_cuda<Tm, sycl::ext::oneapi::experimental::matrix::use::a, | ||
| M, K, LayoutA> &A, | ||
| const joint_matrix_cuda<Tm, sycl::ext::oneapi::experimental::matrix::use::b, | ||
| K, N, LayoutB> &B, | ||
| const joint_matrix_cuda< | ||
| Tc, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, | ||
| sycl::ext::oneapi::experimental::matrix::layout::dynamic> &C) { | ||
| if constexpr (M == 16 && N == 16 && K == 16) { | ||
|
|
||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -16,7 +16,12 @@ namespace matrix { | |
|
|
||
| enum class use { a, b, accumulator }; | ||
|
|
||
| enum class layout { | ||
| row_major = 0, | ||
| col_major = 1, | ||
| ext_intel_packed = 2, | ||
| dynamic = 3 | ||
| }; | ||
|
|
||
| namespace precision { | ||
| class tf32 { | ||
|
|
||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -61,19 +61,8 @@ struct joint_matrix { | |
| } | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| #if defined(__SPIR__) | ||
| joint_matrix(const joint_matrix &other) = delete; | ||
| joint_matrix &operator=(const joint_matrix &rhs) = delete; | ||
| #endif // defined(__SPIR__) | ||
| #endif | ||
| }; | ||
|
|
@@ -98,19 +87,21 @@ class wi_data { | |
| #if defined(__NVPTX__) | ||
| return jm.cuda_impl.wi_marray.size(); | ||
| #else | ||
| throw runtime_error( | ||
|
||
| "get_wi_data is available using: ext::oneapi::detail::get_wi_data, but " | ||
| "intel users are expected to use joint_matrix_copy instead.", | ||
| PI_ERROR_INVALID_DEVICE); | ||
| #endif | ||
| }; | ||
|
|
||
| decltype(auto) operator[](size_t i) { | ||
| #if defined(__NVPTX__) | ||
| return (jm.cuda_impl.wi_marray[i]); | ||
| #else | ||
| throw runtime_error( | ||
| "get_wi_data is available using: ext::oneapi::detail::get_wi_data, but " | ||
|
||
| "intel users are expected to use joint_matrix_copy instead.", | ||
| PI_ERROR_INVALID_DEVICE); | ||
| #endif | ||
| }; | ||
| }; | ||
YuriPlyakhin marked this conversation as resolved.
Show resolved
Hide resolved
|
||
|
|
@@ -138,9 +129,9 @@ template <typename Group, typename T, use Use, size_t Rows, size_t Cols, | |
| __SYCL2020_DEPRECATED("get_wi_data() is deprecated for CUDA backend. Please " | ||
|
There was a problem hiding this comment. We should remove this. This is not really deprecated as joint_matrix is experimental so we can just remove APIs. Deprecated means they still exist and implementations maintain them. In the case of get_wi_data. it is replaced by joint_matrix_apply There was a problem hiding this comment. This will be addressed by @JackAKirk among other CUDA changes in a separate PR. There was a problem hiding this comment. Yes I will make this change as soon as this PR is merged. |
||
| "use joint_matrix_apply() instead.") | ||
| #else | ||
| __attribute__(( | ||
|
||
| unavailable("get_wi_data can't be used on intel device, please use " | ||
| "joint_matrix_apply instead!"))) | ||
| #endif | ||
dkhaldi marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| #endif | ||
| inline __SYCL_ALWAYS_INLINE decltype(auto) | ||
|
|
@@ -176,7 +167,7 @@ joint_matrix_apply(Group sg, joint_matrix<Group, T, Use, M, N, Layout> &jm, | |
| using storage_element_type = | ||
| typename oneapi::detail::jm_type_interpretation_helper_trait< | ||
| T>::storage_element_type; | ||
| auto wi_data_c = sycl::ext::oneapi::detail::get_wi_data(sg, jm); | ||
| for (int i = 0; i < wi_data_c.length(); i++) { | ||
| storage_element_type element = wi_data_c[i]; | ||
| lambda(element); | ||
|
|
@@ -262,7 +253,7 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load( | |
| Ptr, stride, __spv::MatrixLayout::ColumnMajor, | ||
| spv_scope_traits<Group>::value); | ||
| break; | ||
| case layout::ext_intel_packed: | ||
| res.spvm = __spirv_JointMatrixLoadINTEL< | ||
| DecorT, S, NumRows, NumCols, | ||
| spv_matrix_use_traits<use::accumulator>::value, | ||
|
|
@@ -327,8 +318,9 @@ template <typename Group, typename T, size_t NumRows, size_t NumCols, | |
| access::address_space Space, access::decorated IsDecorated> | ||
| inline __SYCL_ALWAYS_INLINE void joint_matrix_store( | ||
| Group sg, | ||
| const joint_matrix<Group, T, use::accumulator, NumRows, NumCols, | ||
| sycl::ext::oneapi::experimental::matrix::layout::dynamic> | ||
| &src, | ||
| multi_ptr<T, Space, IsDecorated> dst, size_t stride, | ||
| sycl::ext::oneapi::experimental::matrix::layout Layout) { | ||
| #if defined(__SYCL_DEVICE_ONLY__) | ||
|
|
@@ -361,7 +353,7 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( | |
| Ptr, src.spvm, stride, __spv::MatrixLayout::ColumnMajor, | ||
| spv_scope_traits<Group>::value); | ||
| break; | ||
| case layout::ext_intel_packed: | ||
| __spirv_JointMatrixStoreINTEL< | ||
| DecorT, T, NumRows, NumCols, | ||
| spv_matrix_use_traits<use::accumulator>::value, | ||
|
|
@@ -382,53 +374,79 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( | |
| #endif // defined(__SYCL_DEVICE_ONLY__) | ||
| } | ||
|
|
||
| template <typename Group, typename Ta, typename Tb, typename Tc, typename Td, | ||
| std::size_t M, std::size_t K, std::size_t N, layout LayoutA, | ||
| layout LayoutB> | ||
| inline __SYCL_ALWAYS_INLINE void joint_matrix_mad( | ||
| Group sg, | ||
| joint_matrix<Group, Td, use::accumulator, M, N, | ||
| sycl::ext::oneapi::experimental::matrix::layout::dynamic> &D, | ||
| const joint_matrix<Group, Ta, use::a, M, K, LayoutA> &A, | ||
| const joint_matrix<Group, Tb, use::b, K, N, LayoutB> &B, | ||
| const joint_matrix<Group, Tc, use::accumulator, M, N, | ||
| sycl::ext::oneapi::experimental::matrix::layout::dynamic> | ||
| &C) { | ||
| #if defined(__SYCL_DEVICE_ONLY__) | ||
| #if defined(__NVPTX__) | ||
| std::ignore = sg; | ||
| if constexpr (std::is_same<Ta, Tb>::value) { | ||
| sycl::ext::oneapi::detail::joint_matrix_mad_cuda<Ta, Tc, M, K, N, LayoutA, | ||
| LayoutB>( | ||
| D.cuda_impl, A.cuda_impl, B.cuda_impl, C.cuda_impl); | ||
| } else { | ||
| assert(false && "Ta != Tb : In the CUDA backend joint_matrix_mad " | ||
| "requires that joint_matrix data types Ta and Tb match"); | ||
| } | ||
| #else | ||
| if constexpr (std::is_same<Ta, uint16_t>::value && | ||
| std::is_same<Tb, uint16_t>::value && | ||
| std::is_same<Tc, float>::value) | ||
| D.spvm = __spirv_JointMatrixMadINTEL(A.spvm, B.spvm, C.spvm); | ||
| else if constexpr (std::is_unsigned<Ta>::value && std::is_unsigned<Tb>::value) | ||
| D.spvm = __spirv_JointMatrixUUMadINTEL(A.spvm, B.spvm, C.spvm); | ||
| else if constexpr (std::is_signed<Ta>::value && std::is_unsigned<Tb>::value) | ||
| D.spvm = __spirv_JointMatrixSUMadINTEL(A.spvm, B.spvm, C.spvm); | ||
| else if constexpr (std::is_unsigned<Ta>::value && std::is_signed<Tb>::value) | ||
| D.spvm = __spirv_JointMatrixUSMadINTEL(A.spvm, B.spvm, C.spvm); | ||
| else | ||
| D.spvm = __spirv_JointMatrixMadINTEL(A.spvm, B.spvm, C.spvm); | ||
| #endif // defined(__NVPTX__) | ||
| #else | ||
| std::ignore = sg; | ||
| std::ignore = A; | ||
| std::ignore = B; | ||
| std::ignore = C; | ||
| std::ignore = D; | ||
| throw runtime_error("joint matrix is not supported on host device.", | ||
| PI_ERROR_INVALID_DEVICE); | ||
| #endif // defined(__SYCL_DEVICE_ONLY__) | ||
| } | ||
|
|
||
| template <typename Group, typename T1, typename T2, size_t Rows, size_t Cols, | ||
| use Use1, use Use2, layout Layout1, layout Layout2> | ||
| void joint_matrix_copy( | ||
| Group sg, joint_matrix<Group, T1, Use1, Rows, Cols, Layout1> &src, | ||
| joint_matrix<Group, T2, Use2, Rows, Cols, Layout2> &dst) { | ||
| #if defined(__SYCL_DEVICE_ONLY__) | ||
| #if defined(__NVPTX__) | ||
| std::ignore = sg; | ||
| for (int i = 0; i < src.cuda_impl.wi_marray.size(); i++) { | ||
| dst.cuda_impl.wi_marray[i] = src.cuda_impl.wi_marray[i]; | ||
| } | ||
| #else | ||
| using storage_element_type = | ||
| typename oneapi::detail::jm_type_interpretation_helper_trait< | ||
| T2>::storage_element_type; | ||
| auto wi_data_c = sycl::ext::oneapi::detail::get_wi_data(sg, src); | ||
| auto wi_data_dst = sycl::ext::oneapi::detail::get_wi_data(sg, dst); | ||
| for (int i = 0; i < wi_data_c.length(); i++) { | ||
| wi_data_dst[i] = static_cast<storage_element_type>(wi_data_c[i]); | ||
| } | ||
| #endif // defined(__NVPTX__) | ||
| #else | ||
| std::ignore = sg; | ||
| std::ignore = dst; | ||
| std::ignore = src; | ||
| throw runtime_error("joint matrix is not supported on host device.", | ||
| PI_ERROR_INVALID_DEVICE); | ||
| #endif // defined(__SYCL_DEVICE_ONLY__) | ||
|
|
||
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
to be consistent with spec. Also depending on matrix instead of M and N, it can be N and K, etc..., so to avoid ambiguity, I suggest not to use M and N for API that can be applied to matrices with different use