From c6f135714d3fc43710cf930377e4b127acb7c040 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Tue, 6 Feb 2024 07:45:29 -0800 Subject: [PATCH 1/2] [SYCL][Matrix] Correct Prefetch instruction name Signed-off-by: Sidorov, Dmitry --- sycl/include/CL/__spirv/spirv_ops.hpp | 2 +- sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index ea1a6580d30e6..f2dd855020e71 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -175,7 +175,7 @@ __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL *, Ts val, size_t i); template -extern __DPCPP_SYCL_EXTERNAL void __spirv_JointMatrixPrefetchINTEL( +extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL( T *Ptr, std::size_t coordX, std::size_t coordY, unsigned int CacheLevel, __spv::MatrixLayout Layout, std::size_t Stride); diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index 98aea6f04a48b..f19d179389b62 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -524,7 +524,7 @@ joint_matrix_prefetch(Group sg, T *Ptr, size_t stride, // Will be removed once SPIRV implementation also uses offsetpointer size_t coordX = 0; size_t coordY = 0; - __spirv_JointMatrixPrefetchINTEL( + __spirv_CooperativeMatrixPrefetchINTEL( Ptr, coordX, coordY, detail::PropertyMetaInfo::value, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); #endif // defined(__NVPTX__) From c9a699fd1f93ff2d1de8054d0a8a1c5ea0172dc3 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Tue, 6 Feb 2024 07:51:41 -0800 Subject: [PATCH 2/2] fix prefetch Signed-off-by: Sidorov, Dmitry --- sycl/include/CL/__spirv/spirv_ops.hpp | 7 ++++--- sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp | 5 +++-- 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index f2dd855020e71..9af5b7e75ae38 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -174,10 +174,11 @@ extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL * __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL *, Ts val, size_t i); -template +template extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL( - T *Ptr, std::size_t coordX, std::size_t coordY, unsigned int CacheLevel, - __spv::MatrixLayout Layout, std::size_t Stride); + T *Ptr, std::size_t coordX, std::size_t coordY, std::size_t NumRows, + std::size_t NumCols, unsigned int CacheLevel, __spv::MatrixLayout Layout, + std::size_t Stride); #ifndef __SPIRV_BUILTIN_DECLARATIONS__ #error \ diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index f19d179389b62..a07e9c144ba6a 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -524,8 +524,9 @@ joint_matrix_prefetch(Group sg, T *Ptr, size_t stride, // Will be removed once SPIRV implementation also uses offsetpointer size_t coordX = 0; size_t coordY = 0; - __spirv_CooperativeMatrixPrefetchINTEL( - Ptr, coordX, coordY, detail::PropertyMetaInfo::value, + __spirv_CooperativeMatrixPrefetchINTEL( + Ptr, coordX, coordY, NumRows, NumCols, + detail::PropertyMetaInfo::value, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); #endif // defined(__NVPTX__) #else