[SYCL][Matrix] Correct Prefetch instruction usage (#12623)

Signed-off-by: Sidorov, Dmitry <dmitry.sidorov@intel.com>
intel · Feb 8, 2024 · 0eac618 · 0eac618
1 parent 96073b9
commit 0eac618
Show file tree

Hide file tree

Showing 2 changed files with 8 additions and 6 deletions.
diff --git 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<T, R, C, L, S, U> *
 __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *,
  Ts val, size_t i);
 
-template <typename T, std::size_t NumRows, std::size_t NumCols>
-extern __DPCPP_SYCL_EXTERNAL void __spirv_JointMatrixPrefetchINTEL(
- T *Ptr, std::size_t coordX, std::size_t coordY, unsigned int CacheLevel,
- __spv::MatrixLayout Layout, std::size_t Stride);
+template <typename T>
+extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL(
+ 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
@@ -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_JointMatrixPrefetchINTEL<T, NumRows, NumCols>(
- Ptr, coordX, coordY, detail::PropertyMetaInfo<decltype(prop)>::value,
+ __spirv_CooperativeMatrixPrefetchINTEL<T>(
+ Ptr, coordX, coordY, NumRows, NumCols,
+ detail::PropertyMetaInfo<decltype(prop)>::value,
  sycl::detail::joint_matrix_layout_to_spv(Layout), stride);
 #endif // defined(__NVPTX__)
 #else