DPC++ Runtime
Runtime libraries for oneAPI DPC++
Classes | Enumerations | Functions | Variables
cl::sycl::ext::intel::experimental::matrix Namespace Reference

Classes

struct  enable_if<!((NumRows<=tile_size) &&(NumCols *sizeof(T)/4<=tile_size) &&(Layout !=matrix_layout::col_major))>::type >
 
struct  enable_if<(NumRows<=tile_size) &&(NumCols *sizeof(T)/4<=tile_size)>::type >
 
struct  joint_matrix
 

Enumerations

enum  matrix_layout { matrix_layout::row_major, matrix_layout::col_major, matrix_layout::packed_a, matrix_layout::packed_b }
 

Functions

static _tile1024i tileloadd64_internal (short row, short col, char *buf, size_t stride)
 
static _tile1024i tdpbssd_internal (unsigned short m, unsigned short n, unsigned short k, _tile1024i dst, _tile1024i src1, _tile1024i src2)
 
static _tile1024i tdpbf16ps_internal (unsigned short m, unsigned short n, unsigned short k, _tile1024i dst, _tile1024i src1, _tile1024i src2)
 
static void tilestored64_internal (short row, short col, char *buf, size_t stride, _tile1024i tile)
 
template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if<(NumRows > tile_size)||(NumCols *sizeof(T)/4 > tile_size), void >::type joint_matrix_load (Group sg, joint_matrix< Group, T, NumRows, NumCols, Layout > &jm, multi_ptr< T, Space > src, size_t stride, matrix_layout layout)
 
template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if<(NumRows<=tile_size) &&(NumCols *sizeof(T)/4<=tile_size), void >::type joint_matrix_load (Group sg, joint_matrix< Group, T, NumRows, NumCols, Layout > &jm, multi_ptr< T, Space > src, size_t stride, matrix_layout layout)
 
template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if<(NumRows > tile_size)||(NumCols *sizeof(T)/4 > tile_size), void >::type joint_matrix_store (Group sg, joint_matrix< Group, T, NumRows, NumCols, Layout > &jm, multi_ptr< T, Space > dst, size_t stride, matrix_layout layout)
 
template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if<(NumRows<=tile_size) &&(NumCols *sizeof(T)/4<=tile_size), void >::type joint_matrix_store (Group sg, joint_matrix< Group, T, NumRows, NumCols, Layout > &jm, multi_ptr< T, Space > dst, size_t stride, matrix_layout layout)
 
template<typename Group , typename T1 , typename T2 , size_t NumRowsA, size_t NumColsA, size_t NumRowsB, size_t NumColsB, size_t NumRowsC, size_t NumColsC, matrix_layout LayoutA, matrix_layout LayoutB, matrix_layout LayoutC>
__SYCL_ALWAYS_INLINE std::enable_if<((std::is_same< T1, int8_t >::value &&std::is_same< T2, int32_t >::value)||(std::is_same< T1, unsigned short >::value &&std::is_same< T2, float >::value)) &&(LayoutA==matrix_layout::row_major) &&(LayoutB==matrix_layout::packed_b) &&(LayoutC==matrix_layout::row_major), joint_matrix< Group, T2, NumRowsC, NumColsC, LayoutC > >::type joint_matrix_mad (Group sg, joint_matrix< Group, T1, NumRowsA, NumColsA, LayoutA > &jmA, joint_matrix< Group, T1, NumRowsB, NumColsB, LayoutB > &jmB, joint_matrix< Group, T2, NumRowsC, NumColsC, LayoutC > &jmC)
 

Variables

constexpr size_t tile_size = 16
 

Enumeration Type Documentation

◆ matrix_layout

enum cl::sycl::ext::intel::experimental::matrix::matrix_layout
strong
Enumerator
row_major 
col_major 
packed_a 
packed_b 

Definition at line 118 of file matrix-aot-amx.hpp.

Function Documentation

◆ joint_matrix_load() [1/2]

template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if< (NumRows > tile_size) || (NumCols * sizeof(T) / 4 > tile_size), void>::type cl::sycl::ext::intel::experimental::matrix::joint_matrix_load ( Group  sg,
joint_matrix< Group, T, NumRows, NumCols, Layout > &  jm,
multi_ptr< T, Space >  src,
size_t  stride,
matrix_layout  layout 
)
inline

Definition at line 320 of file matrix-aot-amx.hpp.

◆ joint_matrix_load() [2/2]

template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if<(NumRows <= tile_size) && (NumCols * sizeof(T) / 4 <= tile_size), void>::type cl::sycl::ext::intel::experimental::matrix::joint_matrix_load ( Group  sg,
joint_matrix< Group, T, NumRows, NumCols, Layout > &  jm,
multi_ptr< T, Space >  src,
size_t  stride,
matrix_layout  layout 
)
inline

Definition at line 343 of file matrix-aot-amx.hpp.

◆ joint_matrix_mad()

template<typename Group , typename T1 , typename T2 , size_t NumRowsA, size_t NumColsA, size_t NumRowsB, size_t NumColsB, size_t NumRowsC, size_t NumColsC, matrix_layout LayoutA, matrix_layout LayoutB, matrix_layout LayoutC>
__SYCL_ALWAYS_INLINE std::enable_if< ((std::is_same<T1, int8_t>::value && std::is_same<T2, int32_t>::value) || (std::is_same<T1, unsigned short>::value && std::is_same<T2, float>::value)) && (LayoutA == matrix_layout::row_major) && (LayoutB == matrix_layout::packed_b) && (LayoutC == matrix_layout::row_major), joint_matrix<Group, T2, NumRowsC, NumColsC, LayoutC> >::type cl::sycl::ext::intel::experimental::matrix::joint_matrix_mad ( Group  sg,
joint_matrix< Group, T1, NumRowsA, NumColsA, LayoutA > &  jmA,
joint_matrix< Group, T1, NumRowsB, NumColsB, LayoutB > &  jmB,
joint_matrix< Group, T2, NumRowsC, NumColsC, LayoutC > &  jmC 
)
inline

Definition at line 406 of file matrix-aot-amx.hpp.

◆ joint_matrix_store() [1/2]

template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if< (NumRows > tile_size) || (NumCols * sizeof(T) / 4 > tile_size), void>::type cl::sycl::ext::intel::experimental::matrix::joint_matrix_store ( Group  sg,
joint_matrix< Group, T, NumRows, NumCols, Layout > &  jm,
multi_ptr< T, Space >  dst,
size_t  stride,
matrix_layout  layout 
)
inline

Definition at line 360 of file matrix-aot-amx.hpp.

◆ joint_matrix_store() [2/2]

template<typename Group , typename T , size_t NumRows, size_t NumCols, matrix_layout Layout, access::address_space Space>
__SYCL_ALWAYS_INLINE std::enable_if<(NumRows <= tile_size) && (NumCols * sizeof(T) / 4 <= tile_size), void>::type cl::sycl::ext::intel::experimental::matrix::joint_matrix_store ( Group  sg,
joint_matrix< Group, T, NumRows, NumCols, Layout > &  jm,
multi_ptr< T, Space >  dst,
size_t  stride,
matrix_layout  layout 
)
inline

Definition at line 382 of file matrix-aot-amx.hpp.

◆ tdpbf16ps_internal()

static _tile1024i cl::sycl::ext::intel::experimental::matrix::tdpbf16ps_internal ( unsigned short  m,
unsigned short  n,
unsigned short  k,
_tile1024i  dst,
_tile1024i  src1,
_tile1024i  src2 
)
static

Definition at line 107 of file matrix-aot-amx.hpp.

◆ tdpbssd_internal()

static _tile1024i cl::sycl::ext::intel::experimental::matrix::tdpbssd_internal ( unsigned short  m,
unsigned short  n,
unsigned short  k,
_tile1024i  dst,
_tile1024i  src1,
_tile1024i  src2 
)
static

Definition at line 102 of file matrix-aot-amx.hpp.

◆ tileloadd64_internal()

static _tile1024i cl::sycl::ext::intel::experimental::matrix::tileloadd64_internal ( short  row,
short  col,
char *  buf,
size_t  stride 
)
static

Definition at line 98 of file matrix-aot-amx.hpp.

◆ tilestored64_internal()

static void cl::sycl::ext::intel::experimental::matrix::tilestored64_internal ( short  row,
short  col,
char *  buf,
size_t  stride,
_tile1024i  tile 
)
static

Definition at line 112 of file matrix-aot-amx.hpp.

Variable Documentation

◆ tile_size

constexpr size_t cl::sycl::ext::intel::experimental::matrix::tile_size = 16
inlineconstexpr

Definition at line 120 of file matrix-aot-amx.hpp.