DPC++ Runtime
Runtime libraries for oneAPI DPC++
Namespaces | Classes | Enumerations | Functions
sycl::_V1::ext::oneapi::experimental::matrix Namespace Reference

Namespaces

 precision
 

Classes

struct  spv_matrix_layout_traits
 
struct  spv_matrix_use_traits
 
struct  spv_scope_traits
 
struct  spv_scope_traits< sycl::sub_group >
 
struct  spv_scope_traits< sycl::group< D > >
 
struct  joint_matrix
 
struct  combination
 
struct  matrix_params
 
struct  matrix_params< architecture::intel_cpu_spr, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void >)>::type >
 
struct  matrix_params< architecture::intel_cpu_spr, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&sM !=0 &&sN !=0 &&sK !=0)>::type >
 
struct  matrix_params< architecture::intel_gpu_dg2_g10, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void >)>::type >
 
struct  matrix_params< architecture::intel_gpu_dg2_g10, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&sM !=0 &&sN !=0 &&sK !=0)>::type >
 
struct  matrix_params< architecture::intel_gpu_dg2_g11, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void >)>::type >
 
struct  matrix_params< architecture::intel_gpu_dg2_g11, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&sM !=0 &&sN !=0 &&sK !=0)>::type >
 
struct  matrix_params< architecture::intel_gpu_dg2_g12, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void >)>::type >
 
struct  matrix_params< architecture::intel_gpu_dg2_g12, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&sM !=0 &&sN !=0 &&sK !=0)>::type >
 
struct  matrix_params< architecture::intel_gpu_pvc, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void >)>::type >
 
struct  matrix_params< architecture::intel_gpu_pvc, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&sM !=0 &&sN !=0 &&sK !=0)>::type >
 
struct  matrix_params< architecture::amd_gpu_gfx90a, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if_t<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb > &&std::is_same_v< Tc, Td >)> >
 
struct  matrix_params< architecture::amd_gpu_gfx90a, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if_t<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb > &&std::is_same_v< Tc, Td > &&sM !=0 &&sN !=0 &&sK !=0)> >
 
struct  matrix_params< architecture::nvidia_gpu_sm_70, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if_t<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb >)> >
 
struct  matrix_params< architecture::nvidia_gpu_sm_72, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb >)>::type >
 
struct  matrix_params< architecture::nvidia_gpu_sm_80, Ta, Tb, Tc, Td, 0, 0, 0, typename std::enable_if_t<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb >)> >
 
struct  matrix_params< architecture::nvidia_gpu_sm_70, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if_t<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb > &&sM !=0 &&sN !=0 &&sK !=0)> >
 
struct  matrix_params< architecture::nvidia_gpu_sm_72, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if_t<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb > &&sM !=0 &&sN !=0 &&sK !=0)> >
 
struct  matrix_params< architecture::nvidia_gpu_sm_80, Ta, Tb, Tc, Td, sM, sN, sK, typename std::enable_if_t<(!std::is_same_v< Ta, void > &&!std::is_same_v< Tb, void > &&!std::is_same_v< Tc, void > &&!std::is_same_v< Td, void > &&std::is_same_v< Ta, Tb > &&sM !=0 &&sN !=0 &&sK !=0)> >
 

Enumerations

enum class  use { a , b , accumulator }
 
enum class  layout { row_major = 0 , col_major = 1 , ext_intel_packed = 2 , dynamic = 3 }
 
enum class  matrix_type {
  bf16 , fp16 , tf32 , fp32 ,
  fp64 , sint8 , sint16 , sint32 ,
  sint64 , uint8 , uint16 , uint32 ,
  uint64
}
 

Functions

template<typename Group , typename T , use Use, size_t M, size_t N, layout Layout, typename F >
__SYCL_ALWAYS_INLINE void joint_matrix_apply (Group sg, joint_matrix< Group, T, Use, M, N, Layout > &jm, F &&lambda)
 
template<typename Group , typename T , use Use, size_t M, size_t N, layout Layout, typename F >
__SYCL_ALWAYS_INLINE void joint_matrix_apply (Group sg, joint_matrix< Group, T, Use, M, N, Layout > &jmsrc, joint_matrix< Group, T, Use, M, N, Layout > &jmdest, F &&lambda)
 
template<typename Group , typename T , size_t NumRows, size_t NumCols, use Use, layout Layout, typename T2 >
__SYCL_ALWAYS_INLINE void joint_matrix_fill (Group, joint_matrix< Group, T, Use, NumRows, NumCols, Layout > &res, const T2 &v)
 
template<typename Group , typename S , typename T , size_t NumRows, size_t NumCols, access::address_space Space, access::decorated IsDecorated, std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value, bool > = true>
__SYCL_ALWAYS_INLINE void joint_matrix_load (Group sg, joint_matrix< Group, S, use::accumulator, NumRows, NumCols, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &res, multi_ptr< T, Space, IsDecorated > src, size_t stride, sycl::ext::oneapi::experimental::matrix::layout Layout)
 
template<typename Group , typename S , typename T , use Use, size_t NumRows, size_t NumCols, matrix::layout Layout, access::address_space Space, access::decorated IsDecorated, std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value||(std::is_same< S, precision::tf32 >::value &&std::is_same< std::remove_const_t< T >, float >::value), bool > = true>
__SYCL_ALWAYS_INLINE void joint_matrix_load (Group sg, joint_matrix< Group, S, Use, NumRows, NumCols, Layout > &res, multi_ptr< T, Space, IsDecorated > src, size_t stride)
 
template<typename Group , typename S , typename T , size_t NumRows, size_t NumCols, typename PropertyListT , std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value, bool > = true>
__SYCL_ALWAYS_INLINE void joint_matrix_load (Group sg, joint_matrix< Group, S, use::accumulator, NumRows, NumCols, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &res, ext::oneapi::experimental::annotated_ptr< T, PropertyListT > src, size_t stride, sycl::ext::oneapi::experimental::matrix::layout Layout)
 
template<typename Group , typename S , typename T , use Use, size_t NumRows, size_t NumCols, matrix::layout Layout, typename PropertyListT , std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value||(std::is_same< S, precision::tf32 >::value &&std::is_same< std::remove_const_t< T >, float >::value), bool > = true>
__SYCL_ALWAYS_INLINE void joint_matrix_load (Group sg, joint_matrix< Group, S, Use, NumRows, NumCols, Layout > &res, ext::oneapi::experimental::annotated_ptr< T, PropertyListT > src, size_t stride)
 
template<typename Group , typename T , size_t NumRows, size_t NumCols, access::address_space Space, access::decorated IsDecorated>
__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)
 
template<typename Group , typename T , size_t NumRows, size_t NumCols, typename PropertyListT >
__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, ext::oneapi::experimental::annotated_ptr< T, PropertyListT > dst, size_t stride, sycl::ext::oneapi::experimental::matrix::layout Layout)
 
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>
__SYCL_ALWAYS_INLINE void joint_matrix_mad (Group, 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)
 
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)
 
__SYCL_ALWAYS_INLINE float round_to_tf32 (const float &a)
 
template<size_t NumRows, size_t NumCols, typename Group , typename T , typename Properties = ext::oneapi::experimental::empty_properties_t>
__SYCL_ALWAYS_INLINE void joint_matrix_prefetch (Group sg, T *Ptr, size_t stride, sycl::ext::oneapi::experimental::matrix::layout Layout, Properties properties={})
 
template<typename Ta , typename Tb , typename Tc >
constexpr bool is_combination_valid_amx (size_t sM, size_t sN, size_t sK)
 
template<typename Ta , typename Tb , typename Tc >
constexpr bool are_types_valid_amx ()
 
template<typename Ta , typename Tb , typename Tc >
constexpr bool is_combination_valid_xmx8 (size_t sM, size_t sN, size_t sK)
 
template<typename Ta , typename Tb , typename Tc >
constexpr bool are_types_valid_xmx8 ()
 
template<typename Ta , typename Tb , typename Tc >
constexpr bool is_combination_valid_xmx16 (size_t sM, size_t sN, size_t sK)
 
template<typename Ta , typename Tb , typename Tc >
constexpr bool are_types_valid_xmx16 ()
 
template<typename Ta , typename Tc >
constexpr bool is_combination_valid_amd_gfx90a (size_t sM, size_t sN, size_t sK)
 AMD Matrix Cores - GFX90A architecture ///. More...
 
template<typename Ta , typename Tc >
constexpr bool are_types_valid_amd_gfx90a ()
 
template<typename Ta , typename Tc , typename Td >
constexpr bool are_types_valid_cuda_sm70 ()
 CUDA Tensor Cores - sm70, sm72 and sm80 ///. More...
 
template<typename Ta , typename Tc , typename Td >
constexpr bool are_types_valid_cuda_sm72 ()
 
template<typename Ta , typename Tc , typename Td >
constexpr bool are_types_valid_cuda_sm80 ()
 
template<typename Ta , typename Tc , typename Td >
constexpr bool is_combination_valid_cuda_sm70 (size_t sM, size_t sN, size_t sK)
 
template<typename Ta , typename Tc , typename Td >
constexpr bool is_combination_valid_cuda_sm72 (size_t sM, size_t sN, size_t sK)
 
template<typename Ta , typename Tc , typename Td >
constexpr bool is_combination_valid_cuda_sm80 (size_t sM, size_t sN, size_t sK)
 

Enumeration Type Documentation

◆ layout

enum sycl::_V1::ext::oneapi::experimental::matrix::layout
strong
Enumerator
row_major 
col_major 
ext_intel_packed 
dynamic 

Definition at line 23 of file matrix-unified-utils.hpp.

◆ matrix_type

enum sycl::_V1::ext::oneapi::experimental::matrix::matrix_type
strong
Enumerator
bf16 
fp16 
tf32 
fp32 
fp64 
sint8 
sint16 
sint32 
sint64 
uint8 
uint16 
uint32 
uint64 

Definition at line 18 of file query-types.hpp.

◆ use

enum sycl::_V1::ext::oneapi::experimental::matrix::use
strong
Enumerator
a 
b 
accumulator 

Definition at line 21 of file matrix-unified-utils.hpp.

Function Documentation

◆ are_types_valid_amd_gfx90a()

template<typename Ta , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::are_types_valid_amd_gfx90a ( )
constexpr

Definition at line 507 of file static-query-use.hpp.

◆ are_types_valid_amx()

template<typename Ta , typename Tb , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::are_types_valid_amx ( )
constexpr

Definition at line 64 of file static-query-use.hpp.

◆ are_types_valid_cuda_sm70()

template<typename Ta , typename Tc , typename Td >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::are_types_valid_cuda_sm70 ( )
constexpr

CUDA Tensor Cores - sm70, sm72 and sm80 ///.

Definition at line 579 of file static-query-use.hpp.

◆ are_types_valid_cuda_sm72()

template<typename Ta , typename Tc , typename Td >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::are_types_valid_cuda_sm72 ( )
constexpr

Definition at line 591 of file static-query-use.hpp.

◆ are_types_valid_cuda_sm80()

template<typename Ta , typename Tc , typename Td >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::are_types_valid_cuda_sm80 ( )
constexpr

Definition at line 599 of file static-query-use.hpp.

◆ are_types_valid_xmx16()

template<typename Ta , typename Tb , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::are_types_valid_xmx16 ( )
constexpr

Definition at line 405 of file static-query-use.hpp.

◆ are_types_valid_xmx8()

template<typename Ta , typename Tb , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::are_types_valid_xmx8 ( )
constexpr

Definition at line 173 of file static-query-use.hpp.

◆ is_combination_valid_amd_gfx90a()

template<typename Ta , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::is_combination_valid_amd_gfx90a ( size_t  sM,
size_t  sN,
size_t  sK 
)
constexpr

AMD Matrix Cores - GFX90A architecture ///.

Definition at line 491 of file static-query-use.hpp.

◆ is_combination_valid_amx()

template<typename Ta , typename Tb , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::is_combination_valid_amx ( size_t  sM,
size_t  sN,
size_t  sK 
)
constexpr

Definition at line 44 of file static-query-use.hpp.

◆ is_combination_valid_cuda_sm70()

template<typename Ta , typename Tc , typename Td >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::is_combination_valid_cuda_sm70 ( size_t  sM,
size_t  sN,
size_t  sK 
)
constexpr

Definition at line 609 of file static-query-use.hpp.

◆ is_combination_valid_cuda_sm72()

template<typename Ta , typename Tc , typename Td >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::is_combination_valid_cuda_sm72 ( size_t  sM,
size_t  sN,
size_t  sK 
)
constexpr

Definition at line 617 of file static-query-use.hpp.

◆ is_combination_valid_cuda_sm80()

template<typename Ta , typename Tc , typename Td >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::is_combination_valid_cuda_sm80 ( size_t  sM,
size_t  sN,
size_t  sK 
)
constexpr

Definition at line 625 of file static-query-use.hpp.

◆ is_combination_valid_xmx16()

template<typename Ta , typename Tb , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::is_combination_valid_xmx16 ( size_t  sM,
size_t  sN,
size_t  sK 
)
constexpr

Definition at line 380 of file static-query-use.hpp.

◆ is_combination_valid_xmx8()

template<typename Ta , typename Tb , typename Tc >
constexpr bool sycl::_V1::ext::oneapi::experimental::matrix::is_combination_valid_xmx8 ( size_t  sM,
size_t  sN,
size_t  sK 
)
constexpr

Definition at line 148 of file static-query-use.hpp.

◆ joint_matrix_apply() [1/2]

template<typename Group , typename T , use Use, size_t M, size_t N, layout Layout, typename F >
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_apply ( Group  sg,
joint_matrix< Group, T, Use, M, N, Layout > &  jm,
F &&  lambda 
)
inline

Definition at line 87 of file matrix-unified.hpp.

References sycl::_V1::ext::oneapi::detail::get_wi_data().

◆ joint_matrix_apply() [2/2]

template<typename Group , typename T , use Use, size_t M, size_t N, layout Layout, typename F >
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_apply ( Group  sg,
joint_matrix< Group, T, Use, M, N, Layout > &  jmsrc,
joint_matrix< Group, T, Use, M, N, Layout > &  jmdest,
F &&  lambda 
)
inline

Definition at line 119 of file matrix-unified.hpp.

References sycl::_V1::ext::oneapi::detail::get_wi_data().

◆ joint_matrix_copy()

template<typename Group , typename T1 , typename T2 , size_t Rows, size_t Cols, use Use1, use Use2, layout Layout1, layout Layout2>
void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_copy ( Group  sg,
joint_matrix< Group, T1, Use1, Rows, Cols, Layout1 > &  src,
joint_matrix< Group, T2, Use2, Rows, Cols, Layout2 > &  dst 
)

Definition at line 493 of file matrix-unified.hpp.

References sycl::_V1::ext::oneapi::detail::get_wi_data().

◆ joint_matrix_fill()

template<typename Group , typename T , size_t NumRows, size_t NumCols, use Use, layout Layout, typename T2 >
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_fill ( Group  ,
joint_matrix< Group, T, Use, NumRows, NumCols, Layout > &  res,
const T2 &  v 
)
inline

Definition at line 155 of file matrix-unified.hpp.

◆ joint_matrix_load() [1/4]

template<typename Group , typename S , typename T , use Use, size_t NumRows, size_t NumCols, matrix::layout Layout, typename PropertyListT , std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value||(std::is_same< S, precision::tf32 >::value &&std::is_same< std::remove_const_t< T >, float >::value), bool > = true>
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_load ( Group  sg,
joint_matrix< Group, S, Use, NumRows, NumCols, Layout > &  res,
ext::oneapi::experimental::annotated_ptr< T, PropertyListT >  src,
size_t  stride 
)
inline

Definition at line 312 of file matrix-unified.hpp.

◆ joint_matrix_load() [2/4]

template<typename Group , typename S , typename T , use Use, size_t NumRows, size_t NumCols, matrix::layout Layout, access::address_space Space, access::decorated IsDecorated, std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value||(std::is_same< S, precision::tf32 >::value &&std::is_same< std::remove_const_t< T >, float >::value), bool > = true>
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_load ( Group  sg,
joint_matrix< Group, S, Use, NumRows, NumCols, Layout > &  res,
multi_ptr< T, Space, IsDecorated >  src,
size_t  stride 
)
inline

Definition at line 231 of file matrix-unified.hpp.

References sycl::_V1::ext::oneapi::detail::load_multiplicand_hip(), and sycl::_V1::access::private_space.

◆ joint_matrix_load() [3/4]

template<typename Group , typename S , typename T , size_t NumRows, size_t NumCols, typename PropertyListT , std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value, bool > = true>
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_load ( Group  sg,
joint_matrix< Group, S, use::accumulator, NumRows, NumCols, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &  res,
ext::oneapi::experimental::annotated_ptr< T, PropertyListT >  src,
size_t  stride,
sycl::ext::oneapi::experimental::matrix::layout  Layout 
)
inline

Definition at line 271 of file matrix-unified.hpp.

References sycl::_V1::detail::joint_matrix_layout_to_spv().

◆ joint_matrix_load() [4/4]

template<typename Group , typename S , typename T , size_t NumRows, size_t NumCols, access::address_space Space, access::decorated IsDecorated, std::enable_if_t< std::is_same< S, std::remove_const_t< T >>::value, bool > = true>
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_load ( Group  sg,
joint_matrix< Group, S, use::accumulator, NumRows, NumCols, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &  res,
multi_ptr< T, Space, IsDecorated >  src,
size_t  stride,
sycl::ext::oneapi::experimental::matrix::layout  Layout 
)
inline

Definition at line 184 of file matrix-unified.hpp.

References sycl::_V1::detail::joint_matrix_layout_to_spv(), sycl::_V1::ext::oneapi::detail::load_accumulator_hip(), and sycl::_V1::access::private_space.

◆ joint_matrix_mad()

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>
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_mad ( Group  ,
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 
)
inline

Definition at line 439 of file matrix-unified.hpp.

References sycl::_V1::ext::oneapi::detail::joint_matrix_mad_hip().

◆ joint_matrix_prefetch()

template<size_t NumRows, size_t NumCols, typename Group , typename T , typename Properties = ext::oneapi::experimental::empty_properties_t>
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_prefetch ( Group  sg,
T *  Ptr,
size_t  stride,
sycl::ext::oneapi::experimental::matrix::layout  Layout,
Properties  properties = {} 
)
inline

Definition at line 542 of file matrix-unified.hpp.

◆ joint_matrix_store() [1/2]

template<typename Group , typename T , size_t NumRows, size_t NumCols, typename PropertyListT >
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix_store ( Group  sg,
const joint_matrix< Group, T, use::accumulator, NumRows, NumCols, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &  src,
ext::oneapi::experimental::annotated_ptr< T, PropertyListT >  dst,
size_t  stride,
sycl::ext::oneapi::experimental::matrix::layout  Layout 
)
inline

Definition at line 389 of file matrix-unified.hpp.

References sycl::_V1::detail::joint_matrix_layout_to_spv().

◆ joint_matrix_store() [2/2]

template<typename Group , typename T , size_t NumRows, size_t NumCols, access::address_space Space, access::decorated IsDecorated>
__SYCL_ALWAYS_INLINE void sycl::_V1::ext::oneapi::experimental::matrix::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 
)
inline

Definition at line 346 of file matrix-unified.hpp.

References sycl::_V1::detail::joint_matrix_layout_to_spv(), sycl::_V1::ext::oneapi::detail::joint_matrix_store_hip(), and sycl::_V1::access::private_space.

◆ round_to_tf32()

__SYCL_ALWAYS_INLINE float sycl::_V1::ext::oneapi::experimental::matrix::round_to_tf32 ( const float &  a)
inline

Definition at line 521 of file matrix-unified.hpp.