29 #include <type_traits>
32 inline namespace _V1 {
35 namespace experimental {
42 #define SPV_MATRIX_LAYOUT_TRAITS(LAYOUT, SPV_LAYOUT) \
43 template <> struct spv_matrix_layout_traits<LAYOUT> { \
44 static constexpr __spv::MatrixLayout value = SPV_LAYOUT; \
56 #define SPV_MATRIX_USE_TRAITS(USE, SPV_USE) \
57 template <> struct spv_matrix_use_traits<USE> { \
58 static constexpr __spv::MatrixUse value = SPV_USE; \
74 template <
typename Group,
typename T,
use Use,
size_t Rows,
size_t Cols,
95 using namespace sycl::ext::oneapi::experimental::matrix;
98 template <
typename T,
size_t NumRows,
size_t NumCols,
101 sycl::ext::oneapi::experimental::matrix::layout::dynamic,
113 Group, T, Use, NumRows, NumCols, Layout> &Mat,
118 #if defined(__SYCL_DEVICE_ONLY__)
119 __ocl_vec_t<uint32_t, 2> coord =
120 __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx);
121 const size_t row = coord[0];
122 const size_t col = coord[1];
125 throw runtime_error(
"joint matrix is not supported on host device.",
126 PI_ERROR_INVALID_DEVICE);
131 #ifdef __SYCL_DEVICE_ONLY__
134 spv_matrix_use_traits<Use>::value,
135 spv_matrix_layout_traits<Layout>::value,
136 spv_scope_traits<Group>::value>(M.spvm,
140 throw runtime_error(
"joint matrix is not supported on host device.",
141 PI_ERROR_INVALID_DEVICE);
145 explicit operator bool() {
146 #ifdef __SYCL_DEVICE_ONLY__
149 spv_matrix_use_traits<Use>::value,
150 spv_matrix_layout_traits<Layout>::value,
151 spv_scope_traits<Group>::value>(
154 throw runtime_error(
"joint matrix is not supported on host device.",
155 PI_ERROR_INVALID_DEVICE);
160 #ifdef __SYCL_DEVICE_ONLY__
161 M.spvm = __spirv_VectorInsertDynamic(
166 throw runtime_error(
"joint matrix is not supported on host device.",
167 PI_ERROR_INVALID_DEVICE);
173 #ifdef __SYCL_DEVICE_ONLY__
174 M.spvm = __spirv_VectorInsertDynamic(
177 spv_matrix_use_traits<Use>::value,
178 spv_matrix_layout_traits<Layout>::value,
179 spv_scope_traits<Group>::value>(rhs.M.spvm,
185 throw runtime_error(
"joint matrix is not supported on host device.",
186 PI_ERROR_INVALID_DEVICE);
190 #if __SYCL_DEVICE_ONLY__
192 template <typename T2> wi_element &operator op##=(const T2 & rhs) { \
193 M.spvm = __spirv_VectorInsertDynamic( \
195 static_cast<storage_element_type>( \
196 __spirv_VectorExtractDynamic< \
197 storage_element_type, T, NumRows, NumCols, \
198 spv_matrix_use_traits<Use>::value, \
199 spv_matrix_layout_traits<Layout>::value, \
200 spv_scope_traits<Group>::value>(M.spvm, idx) \
201 op static_cast<storage_element_type>(rhs)), \
207 template <typename T2> wi_element &operator op##=(const T2 & rhs) { \
209 throw runtime_error("joint matrix is not supported on host device.", \
210 PI_ERROR_INVALID_DEVICE); \
220 template <
size_t NumRows,
size_t NumCols,
238 #if defined(__SYCL_DEVICE_ONLY__)
239 __ocl_vec_t<uint32_t, 2> coord =
240 __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx);
241 const uint32_t row = coord[0];
242 const uint32_t col = coord[1];
245 throw runtime_error(
"joint matrix is not supported on host device.",
246 PI_ERROR_INVALID_DEVICE);
251 #ifdef __SYCL_DEVICE_ONLY__
252 return __spirv_VectorExtractDynamic<
254 NumCols, spv_matrix_use_traits<Use>::value,
255 spv_matrix_layout_traits<Layout>::value,
256 spv_scope_traits<Group>::value>(M.spvm, idx);
258 throw runtime_error(
"joint matrix is not supported on host device.",
259 PI_ERROR_INVALID_DEVICE);
263 explicit operator bool() {
264 #ifdef __SYCL_DEVICE_ONLY__
266 __spirv_VectorExtractDynamic<
268 NumRows, NumCols, spv_matrix_use_traits<Use>::value,
269 spv_matrix_layout_traits<Layout>::value,
270 spv_scope_traits<Group>::value>(M.spvm, idx))) >=
271 std::numeric_limits<float>::epsilon();
273 throw runtime_error(
"joint matrix is not supported on host device.",
274 PI_ERROR_INVALID_DEVICE);
279 #ifdef __SYCL_DEVICE_ONLY__
280 M.spvm = __spirv_VectorInsertDynamic(M.spvm, rhs, idx);
284 throw runtime_error(
"joint matrix is not supported on host device.",
285 PI_ERROR_INVALID_DEVICE);
290 NumCols, Use, Layout, Group> &rhs) {
291 #ifdef __SYCL_DEVICE_ONLY__
292 M.spvm = __spirv_VectorInsertDynamic(
296 NumCols, spv_matrix_use_traits<Use>::value,
297 spv_matrix_layout_traits<Layout>::value,
298 spv_scope_traits<Group>::value>(rhs.M.spvm,
304 throw runtime_error(
"joint matrix is not supported on host device.",
305 PI_ERROR_INVALID_DEVICE);
309 #if __SYCL_DEVICE_ONLY__
310 #define OP(opassign, op) \
311 wi_element &operator opassign(const sycl::ext::oneapi::bfloat16 & rhs) { \
312 M.spvm = __spirv_VectorInsertDynamic( \
314 __spirv_VectorExtractDynamic< \
315 sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \
316 NumCols, spv_matrix_use_traits<Use>::value, \
317 spv_matrix_layout_traits<Layout>::value, \
318 spv_scope_traits<Group>::value>(M.spvm, idx) op rhs, \
323 #define OP(opassign, op) \
324 wi_element &operator opassign(const sycl::ext::oneapi::bfloat16 & rhs) { \
326 throw runtime_error("joint matrix is not supported on host device.", \
327 PI_ERROR_INVALID_DEVICE); \
336 #if __SYCL_DEVICE_ONLY__
337 #define OP(type, op) \
338 friend type operator op( \
339 const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, \
340 Layout, Group> &lhs, \
341 const sycl::ext::oneapi::bfloat16 &rhs) { \
342 return __spirv_VectorExtractDynamic< \
343 sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \
344 NumCols, spv_matrix_use_traits<Use>::value, \
345 spv_matrix_layout_traits<Layout>::value, \
346 spv_scope_traits<Group>::value>(lhs.M.spvm, lhs.idx) op rhs; \
348 friend type operator op( \
349 const sycl::ext::oneapi::bfloat16 &lhs, \
350 const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, \
351 Layout, Group> &rhs) { \
352 return __spirv_VectorExtractDynamic< \
353 sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \
354 NumCols, spv_matrix_use_traits<Use>::value, \
355 spv_matrix_layout_traits<Layout>::value, \
356 spv_scope_traits<Group>::value>(rhs.M.spvm, rhs.idx) op lhs; \
363 #define OP(type, op) \
364 friend type operator op( \
365 const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, \
366 Layout, Group> &lhs, \
367 const sycl::ext::oneapi::bfloat16 &rhs) { \
368 return type{static_cast<float>( \
369 __spirv_VectorExtractDynamic< \
370 sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \
371 NumCols, spv_matrix_use_traits<Use>::value, \
372 spv_matrix_layout_traits<Layout>::value, \
373 spv_scope_traits<Group>::value>(lhs.M.spvm, lhs.idx)) \
374 op static_cast<float>(rhs)}; \
376 friend type operator op( \
377 const sycl::ext::oneapi::bfloat16 &lhs, \
378 const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, \
379 Layout, Group> &rhs) { \
380 return type{static_cast<float>( \
381 __spirv_VectorExtractDynamic< \
382 sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \
383 NumCols, spv_matrix_use_traits<Use>::value, \
384 spv_matrix_layout_traits<Layout>::value, \
385 spv_scope_traits<Group>::value>(rhs.M.spvm, rhs.idx)) \
386 op static_cast<float>(lhs)}; \
396 #define OP(type, op) \
397 friend type operator op( \
398 const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, \
400 const sycl::ext::oneapi::bfloat16 &) { \
401 throw runtime_error("joint matrix is not supported on host device.", \
402 PI_ERROR_INVALID_DEVICE); \
404 friend type operator op( \
405 const sycl::ext::oneapi::bfloat16 &, \
406 const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, \
407 Layout, Group> &) { \
408 throw runtime_error("joint matrix is not supported on host device.", \
409 PI_ERROR_INVALID_DEVICE); \
412 OP(
sycl::ext::oneapi::bfloat16, -)
413 OP(
sycl::ext::oneapi::bfloat16, *)
414 OP(
sycl::ext::oneapi::bfloat16, /)
429 template <
typename Group,
typename T,
438 Group, T, Use, Rows, Cols, Layout> &_jm)
441 template <
typename Grp,
typename Type,
445 friend decltype(
auto)
447 Grp, Type, UseJm, NumRows, NumCols, LayoutJm> &);
451 #if __SYCL_DEVICE_ONLY__
452 return __spirv_JointMatrixWorkItemLengthINTEL(jm.spvm);
454 throw runtime_error(
"joint matrix is not supported on host device.",
455 PI_ERROR_INVALID_DEVICE);
459 decltype(
auto) operator[](
size_t i) {
464 template <
typename Group,
typename T,
469 Group, T, Use, Rows, Cols, Layout> &jm) {
478 namespace intel::experimental::matrix {
480 typename Group,
typename T,
typename Tp,
490 Group, Tp, Use, NumRows, NumCols, Layout> &src,
492 #if defined(__SYCL_DEVICE_ONLY__)
494 "Joint Matrix doesn't support store to private memory!");
495 #if defined(__NVPTX__)
498 std::ignore = stride;
500 "This version of the matrix extension is only currently supported on "
502 PI_ERROR_INVALID_DEVICE);
505 using DecorT =
typename sycl::detail::DecoratedType<T, Space>::type;
506 DecorT *Ptr = sycl::detail::getDecorated<DecorT>(dst);
507 __spirv_JointMatrixStoreINTEL<DecorT, Tp, NumRows, NumCols,
512 Ptr, src.spvm, stride,
520 std::ignore = stride;
521 throw runtime_error(
"joint matrix is not supported on host device.",
522 PI_ERROR_INVALID_DEVICE);
527 typename Group,
typename T,
typename Tp,
530 typename PropertyListT,
537 Group, Tp, Use, NumRows, NumCols, Layout> &src,
540 #if defined(__SYCL_DEVICE_ONLY__)
541 #if defined(__NVPTX__)
544 std::ignore = stride;
546 "This version of the matrix extension is only currently supported on "
548 PI_ERROR_INVALID_DEVICE);
552 __spirv_JointMatrixStoreINTEL<T, Tp, NumRows, NumCols,
557 Ptr, src.spvm, stride,
565 std::ignore = stride;
566 throw runtime_error(
"joint matrix is not supported on host device.",
567 PI_ERROR_INVALID_DEVICE);
571 template <
typename Group,
typename T,
580 #if defined(__SYCL_DEVICE_ONLY__)
581 #if defined(__NVPTX__)
583 for (
int i = 0; i < jm.matrix_impl.wi_marray.size(); i++) {
584 lambda(jm.matrix_impl.wi_marray[i]);
587 using storage_element_type =
589 T>::storage_element_type;
591 for (
int i = 0; i < wi_data_c.length(); i++) {
592 storage_element_type element = wi_data_c[i];
593 auto [row, col] = wi_data_c[i].get_coord();
594 lambda(element, row, col);
595 wi_data_c[i] = element;
601 std::ignore = lambda;
602 throw runtime_error(
"joint matrix is not supported on host device.",
603 PI_ERROR_INVALID_DEVICE);
607 using namespace sycl::ext::oneapi::experimental::matrix;
611 template <
typename Group,
typename T,
size_t NumRows,
size_t NumCols,
use Use,
612 layout Layout,
typename T2>
614 Group, joint_matrix<Group, T, Use, NumRows, NumCols, Layout> &Res,
615 const T2 &Value,
size_t Stride,
size_t Height,
size_t Width,
size_t CoordX,
617 #if defined(__SYCL_DEVICE_ONLY__)
618 using storage_element_type =
620 T>::storage_element_type;
621 Res.spvm = __spirv_CooperativeMatrixConstructCheckedINTEL<
622 storage_element_type, T, NumRows, NumCols,
623 spv_matrix_use_traits<Use>::value,
624 spv_matrix_layout_traits<Layout>::value>(
625 static_cast<storage_element_type
>(Value), Stride, Height, Width, CoordX,
630 std::ignore = Stride;
631 std::ignore = Height;
633 std::ignore = CoordX;
634 std::ignore = CoordY;
635 throw runtime_error(
"joint matrix is not supported on host device.",
636 PI_ERROR_INVALID_DEVICE);
641 typename Group,
typename S,
typename T,
size_t NumRows,
size_t NumCols,
643 std::enable_if_t<std::is_same<S, std::remove_const_t<T>>::value,
bool> =
647 joint_matrix<Group, S, use::accumulator, NumRows, NumCols, layout::dynamic>
650 size_t Height,
size_t Width,
size_t CoordX,
size_t CoordY) {
651 #if defined(__SYCL_DEVICE_ONLY__)
653 "Joint Matrix doesn't support load from private memory!");
655 using DecorT =
typename sycl::detail::DecoratedType<T, Space>::type;
656 DecorT *Ptr = sycl::detail::getDecorated<DecorT>(Src);
657 Res.spvm = __spirv_JointMatrixLoadCheckedINTEL<
658 DecorT, S, NumRows, NumCols,
659 spv_matrix_use_traits<use::accumulator>::value,
660 spv_matrix_layout_traits<layout::dynamic>::value>(
661 Ptr, Stride, Height, Width, CoordX, CoordY,
663 spv_scope_traits<Group>::value);
668 std::ignore = Stride;
669 std::ignore = Height;
671 std::ignore = Layout;
672 std::ignore = CoordX;
673 std::ignore = CoordY;
674 throw runtime_error(
"joint matrix is not supported on host device.",
675 PI_ERROR_INVALID_DEVICE);
680 typename Group,
typename S,
typename T,
use Use,
size_t NumRows,
683 std::enable_if_t<std::is_same<S, std::remove_const_t<T>>::value ||
684 (std::is_same<S, precision::tf32>::value &&
685 std::is_same<std::remove_const_t<T>,
float>::value),
688 Group sg, joint_matrix<Group, S, Use, NumRows, NumCols, Layout> &Res,
690 size_t Width,
size_t CoordX,
size_t CoordY) {
691 #if defined(__SYCL_DEVICE_ONLY__)
693 "Joint Matrix doesn't support load from private memory!");
695 using DecorT =
typename sycl::detail::DecoratedType<T, Space>::type;
696 DecorT *Ptr = sycl::detail::getDecorated<DecorT>(Src);
697 Res.spvm = __spirv_JointMatrixLoadCheckedINTEL<
698 DecorT, S, NumRows, NumCols, spv_matrix_use_traits<Use>::value,
699 spv_matrix_layout_traits<Layout>::value>(
700 Ptr, Stride, Height, Width, CoordX, CoordY,
701 spv_matrix_layout_traits<Layout>::value, spv_scope_traits<Group>::value);
706 std::ignore = Stride;
707 std::ignore = Height;
709 std::ignore = CoordX;
710 std::ignore = CoordY;
711 throw runtime_error(
"joint matrix is not supported on host device.",
712 PI_ERROR_INVALID_DEVICE);
716 template <
typename Group,
typename T,
size_t NumRows,
size_t NumCols,
720 joint_matrix<Group, T, use::accumulator, NumRows, NumCols, layout::dynamic>
723 size_t Height,
size_t Width,
size_t CoordX,
size_t CoordY) {
724 #if defined(__SYCL_DEVICE_ONLY__)
726 "Joint Matrix doesn't support store to private memory!");
728 using DecorT =
typename sycl::detail::DecoratedType<T, Space>::type;
729 DecorT *Ptr = sycl::detail::getDecorated<DecorT>(Dst);
730 __spirv_JointMatrixStoreCheckedINTEL<
731 DecorT, T, NumRows, NumCols,
732 spv_matrix_use_traits<use::accumulator>::value,
733 spv_matrix_layout_traits<layout::dynamic>::value>(
734 Ptr, Src.spvm, Stride, Height, Width, CoordX, CoordY,
736 spv_scope_traits<Group>::value);
741 std::ignore = Stride;
742 std::ignore = Height;
744 std::ignore = Layout;
745 std::ignore = CoordX;
746 std::ignore = CoordY;
747 throw runtime_error(
"joint matrix is not supported on host device.",
748 PI_ERROR_INVALID_DEVICE);
752 template <
typename Group,
typename T,
typename Tp,
use Use,
size_t NumRows,
755 std::enable_if_t<Use == use::a || Use == use::b, bool> =
true>
757 Group sg,
const joint_matrix<Group, Tp, Use, NumRows, NumCols, Layout> &Src,
759 size_t Width,
size_t CoordX,
size_t CoordY) {
760 #if defined(__SYCL_DEVICE_ONLY__)
762 "Joint Matrix doesn't support store to private memory!");
764 using DecorT =
typename sycl::detail::DecoratedType<T, Space>::type;
765 DecorT *Ptr = sycl::detail::getDecorated<DecorT>(Dst);
766 __spirv_JointMatrixStoreCheckedINTEL<DecorT, Tp, NumRows, NumCols,
767 spv_matrix_use_traits<Use>::value,
768 spv_matrix_layout_traits<Layout>::value>(
769 Ptr, Src.spvm, Stride, Height, Width, CoordX, CoordY,
770 spv_matrix_layout_traits<Layout>::value, spv_scope_traits<Group>::value);
775 std::ignore = Stride;
776 std::ignore = Height;
778 std::ignore = CoordX;
779 std::ignore = CoordY;
780 throw runtime_error(
"joint matrix is not supported on host device.",
781 PI_ERROR_INVALID_DEVICE);
786 template <
typename Group,
typename S,
typename T,
size_t NumRows,
787 size_t NumCols,
typename PropertyListT,
788 std::enable_if_t<std::is_same<S, std::remove_const_t<T>>::value,
792 joint_matrix<Group, S, use::accumulator, NumRows, NumCols, layout::dynamic>
795 size_t Stride,
layout Layout,
size_t Height,
size_t Width,
size_t CoordX,
797 #if defined(__SYCL_DEVICE_ONLY__)
800 Res.spvm = __spirv_JointMatrixLoadCheckedINTEL<
801 T, S, NumRows, NumCols, spv_matrix_use_traits<use::accumulator>::value,
802 spv_matrix_layout_traits<layout::dynamic>::value>(
803 Ptr, Stride, Height, Width, CoordX, CoordY,
805 spv_scope_traits<Group>::value);
810 std::ignore = Stride;
811 std::ignore = Height;
813 std::ignore = Layout;
814 std::ignore = CoordX;
815 std::ignore = CoordY;
816 throw runtime_error(
"joint matrix is not supported on host device.",
817 PI_ERROR_INVALID_DEVICE);
822 typename Group,
typename S,
typename T,
use Use,
size_t NumRows,
823 size_t NumCols,
layout Layout,
typename PropertyListT,
824 std::enable_if_t<std::is_same<S, std::remove_const_t<T>>::value ||
825 (std::is_same<S, precision::tf32>::value &&
826 std::is_same<std::remove_const_t<T>,
float>::value),
829 Group sg, joint_matrix<Group, S, Use, NumRows, NumCols, Layout> &Res,
831 size_t Stride,
size_t Height,
size_t Width,
size_t CoordX,
size_t CoordY) {
832 #if defined(__SYCL_DEVICE_ONLY__)
835 Res.spvm = __spirv_JointMatrixLoadCheckedINTEL<
836 T, S, NumRows, NumCols, spv_matrix_use_traits<Use>::value,
837 spv_matrix_layout_traits<Layout>::value>(
838 Ptr, Stride, Height, Width, CoordX, CoordY,
839 spv_matrix_layout_traits<Layout>::value, spv_scope_traits<Group>::value);
844 std::ignore = Stride;
845 std::ignore = Height;
847 std::ignore = CoordX;
848 std::ignore = CoordY;
849 throw runtime_error(
"joint matrix is not supported on host device.",
850 PI_ERROR_INVALID_DEVICE);
854 template <
typename Group,
typename T,
size_t NumRows,
size_t NumCols,
855 typename PropertyListT>
858 joint_matrix<Group, T, use::accumulator, NumRows, NumCols, layout::dynamic>
861 size_t Stride,
layout Layout,
size_t Height,
size_t Width,
size_t CoordX,
863 #if defined(__SYCL_DEVICE_ONLY__)
866 __spirv_JointMatrixStoreCheckedINTEL<
867 T, T, NumRows, NumCols, spv_matrix_use_traits<use::accumulator>::value,
868 spv_matrix_layout_traits<layout::dynamic>::value>(
869 Ptr, Src.spvm, Stride, Height, Width, CoordX, CoordY,
871 spv_scope_traits<Group>::value);
876 std::ignore = Stride;
877 std::ignore = Height;
879 std::ignore = Layout;
880 std::ignore = CoordX;
881 std::ignore = CoordY;
882 throw runtime_error(
"joint matrix is not supported on host device.",
883 PI_ERROR_INVALID_DEVICE);
887 template <
typename Group,
typename T,
typename Tp,
use Use,
size_t NumRows,
888 size_t NumCols,
layout Layout,
typename PropertyListT,
889 std::enable_if_t<Use == use::a || Use == use::b, bool> =
true>
891 Group sg,
const joint_matrix<Group, Tp, Use, NumRows, NumCols, Layout> &Src,
893 size_t Stride,
size_t Height,
size_t Width,
size_t CoordX,
size_t CoordY) {
894 #if defined(__SYCL_DEVICE_ONLY__)
897 __spirv_JointMatrixStoreCheckedINTEL<T, Tp, NumRows, NumCols,
898 spv_matrix_use_traits<Use>::value,
899 spv_matrix_layout_traits<Layout>::value>(
900 Ptr, Src.spvm, Stride, Height, Width, CoordX, CoordY,
901 spv_matrix_layout_traits<Layout>::value, spv_scope_traits<Group>::value);
906 std::ignore = Stride;
907 std::ignore = Height;
909 std::ignore = CoordX;
910 std::ignore = CoordY;
911 throw runtime_error(
"joint matrix is not supported on host device.",
912 PI_ERROR_INVALID_DEVICE);
wi_element(sycl::ext::oneapi::experimental::matrix::joint_matrix< Group, sycl::ext::oneapi::bfloat16, Use, NumRows, NumCols, Layout > &Mat, std::size_t i)
wi_element & operator=(const wi_element< sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout, Group > &rhs)
__SYCL_ALWAYS_INLINE std::tuple< uint32_t, uint32_t > get_coord()
wi_element & operator=(const sycl::ext::oneapi::bfloat16 &rhs)
wi_element & operator=(const wi_element< T, NumRows, NumCols, Use, Layout, Group > &rhs)
wi_element(sycl::ext::oneapi::experimental::matrix::joint_matrix< Group, T, Use, NumRows, NumCols, Layout > &Mat, std::size_t i)
wi_element & operator=(const T2 &rhs)
typename oneapi::detail::jm_type_interpretation_helper_trait< T >::storage_element_type storage_element_type
__SYCL_ALWAYS_INLINE std::tuple< size_t, size_t > get_coord()
#define __SYCL_ALWAYS_INLINE
#define SPV_MATRIX_USE_TRAITS(USE, SPV_USE)
#define SPV_MATRIX_LAYOUT_TRAITS(LAYOUT, SPV_LAYOUT)
__SYCL_ALWAYS_INLINE __spv::MatrixLayout joint_matrix_layout_to_spv(sycl::ext::oneapi::experimental::matrix::layout Layout)
constexpr tuple< Ts... > make_tuple(Ts... Args)
__SYCL_ALWAYS_INLINE void joint_matrix_fill_checked(Group, joint_matrix< Group, T, Use, NumRows, NumCols, Layout > &Res, const T2 &Value, size_t Stride, size_t Height, size_t Width, size_t CoordX, size_t CoordY)
__SYCL_ALWAYS_INLINE void joint_matrix_store(Group, const sycl::ext::oneapi::experimental::matrix::joint_matrix< Group, Tp, Use, NumRows, NumCols, Layout > &src, ext::oneapi::experimental::annotated_ptr< T, PropertyListT > dst, size_t stride)
__SYCL_ALWAYS_INLINE void joint_matrix_load_checked(Group sg, joint_matrix< Group, S, Use, NumRows, NumCols, Layout > &Res, ext::oneapi::experimental::annotated_ptr< T, PropertyListT > Src, size_t Stride, size_t Height, size_t Width, size_t CoordX, size_t CoordY)
__SYCL_ALWAYS_INLINE void joint_matrix_apply(Group sg, sycl::ext::oneapi::experimental::matrix::joint_matrix< Group, T, Use, Rows, Cols, Layout > &jm, F &&lambda)
__SYCL_ALWAYS_INLINE void joint_matrix_store_checked(Group sg, const joint_matrix< Group, Tp, Use, NumRows, NumCols, Layout > &Src, ext::oneapi::experimental::annotated_ptr< T, PropertyListT > Dst, size_t Stride, size_t Height, size_t Width, size_t CoordX, size_t CoordY)
decltype(auto) __SYCL_ALWAYS_INLINE get_wi_data(Group sg, sycl::ext::oneapi::experimental::matrix::joint_matrix< Group, T, Use, Rows, Cols, Layout > &jm)
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fabs(T x)
float storage_element_type
static constexpr __spv::MatrixLayout value
static constexpr __spv::MatrixUse value