DPC++ Runtime
Runtime libraries for oneAPI DPC++
matrix-hip.hpp File Reference
Include dependency graph for matrix-hip.hpp:

Go to the source code of this file.

Classes

struct  sycl::_V1::ext::oneapi::detail::to_hip_type< T >
 
struct  sycl::_V1::ext::oneapi::detail::to_hip_type< bfloat16 >
 
struct  sycl::_V1::ext::oneapi::detail::to_hip_type< half >
 
struct  sycl::_V1::ext::oneapi::detail::to_hip_type< int8_t >
 

Namespaces

 sycl
 
 sycl::_V1
 
 sycl::_V1::ext
 
 sycl::_V1::ext::oneapi
 
 sycl::_V1::ext::oneapi::detail
 

Macros

#define __HIP_PLATFORM_AMD_MFMA__
 
#define __SYCL_JOINT_MATRIX_OVERLOAD_ARR(TYPE, USE, M, N, SIZE)
 
#define __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(TYPE, M, N)
 

Typedefs

using sycl::_V1::ext::oneapi::detail::bfloat16x4 = __attribute__((__vector_size__(4 *sizeof(__bf16)))) __fp16
 
using sycl::_V1::ext::oneapi::detail::float16x4 = __attribute__((__vector_size__(4 *sizeof(__fp16)))) __fp16
 
using sycl::_V1::ext::oneapi::detail::floatx4 = __attribute__((__vector_size__(4 *sizeof(float)))) float
 
using sycl::_V1::ext::oneapi::detail::floatx16 = __attribute__((__vector_size__(16 *sizeof(float)))) float
 
using sycl::_V1::ext::oneapi::detail::int32x4 = __attribute__((__vector_size__(4 *sizeof(int32_t)))) int
 
using sycl::_V1::ext::oneapi::detail::int32x16 = __attribute__((__vector_size__(16 *sizeof(int32_t)))) int
 
using sycl::_V1::ext::oneapi::detail::doublex4 = __attribute__((__vector_size__(4 *sizeof(double)))) double
 

Functions

template<sycl::ext::oneapi::experimental::matrix::layout Layout, typename S , typename T , size_t M, size_t N, access::address_space Space, access::decorated IsDecorated, typename Group >
void sycl::_V1::ext::oneapi::detail::load_accumulator_layoutT (joint_matrix_hip< S, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &res, multi_ptr< T, Space, IsDecorated > src, size_t stride, Group &sg)
 
template<typename Group , typename S , typename T , size_t M, size_t N, access::address_space Space, access::decorated IsDecorated, typename = std::enable_if_t<std::is_same_v<S, std::remove_const_t<T>>>>
void sycl::_V1::ext::oneapi::detail::load_accumulator_hip (joint_matrix_hip< S, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &res, multi_ptr< T, Space, IsDecorated > src, size_t stride, sycl::ext::oneapi::experimental::matrix::layout layout, Group &sg)
 
template<typename Group , typename S , typename T , size_t M, size_t N, sycl::ext::oneapi::experimental::matrix::use Use, sycl::ext::oneapi::experimental::matrix::layout Layout, access::address_space Space, access::decorated IsDecorated, typename = typename std::enable_if_t< (Layout == sycl::ext::oneapi::experimental::matrix::layout::row_major || Layout == sycl::ext::oneapi::experimental::matrix::layout::col_major) && std::is_same_v<S, std::remove_const_t<T>>>>
void sycl::_V1::ext::oneapi::detail::load_multiplicand_hip (joint_matrix_hip< S, Use, M, N, Layout > &res, multi_ptr< T, Space, IsDecorated > src, size_t stride, Group &sg)
 
template<typename Group , sycl::ext::oneapi::experimental::matrix::layout Layout, typename T , size_t M, size_t N, access::address_space Space, access::decorated IsDecorated>
void sycl::_V1::ext::oneapi::detail::store_layoutT (const joint_matrix_hip< T, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &src, multi_ptr< T, Space, IsDecorated > dst, size_t stride, Group &sg)
 
template<typename Group , typename T , size_t M, size_t N, access::address_space Space, access::decorated IsDecorated>
void sycl::_V1::ext::oneapi::detail::joint_matrix_store_hip (const joint_matrix_hip< T, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &src, multi_ptr< T, Space, IsDecorated > dst, size_t stride, sycl::ext::oneapi::experimental::matrix::layout layout, Group &sg)
 
template<typename Tm , typename Tc , std::size_t M, std::size_t K, std::size_t N, sycl::ext::oneapi::experimental::matrix::layout LayoutA, sycl::ext::oneapi::experimental::matrix::layout LayoutB>
void sycl::_V1::ext::oneapi::detail::joint_matrix_mad_hip (joint_matrix_hip< Tc, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &D, const joint_matrix_hip< Tm, sycl::ext::oneapi::experimental::matrix::use::a, M, K, LayoutA > &A, const joint_matrix_hip< Tm, sycl::ext::oneapi::experimental::matrix::use::b, K, N, LayoutB > &B, const joint_matrix_hip< Tc, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic > &C)
 

Variables

constexpr int sycl::_V1::ext::oneapi::detail::WAVEFRONT_SIZE = 64
 

Macro Definition Documentation

◆ __HIP_PLATFORM_AMD_MFMA__

#define __HIP_PLATFORM_AMD_MFMA__

Definition at line 14 of file matrix-hip.hpp.

◆ __SYCL_JOINT_MATRIX_OVERLOAD_ARR

#define __SYCL_JOINT_MATRIX_OVERLOAD_ARR (   TYPE,
  USE,
  M,
  N,
  SIZE 
)
Value:
template <sycl::ext::oneapi::experimental::matrix::layout Layout> \
struct joint_matrix_hip< \
TYPE, sycl::ext::oneapi::experimental::matrix::use::USE, M, N, Layout, \
typename std::enable_if_t< \
Layout == \
sycl::ext::oneapi::experimental::matrix::layout::row_major || \
Layout == \
sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \
sycl::marray<TYPE, SIZE> wi_marray; \
};

Definition at line 57 of file matrix-hip.hpp.

◆ __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC

#define __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC (   TYPE,
  M,
 
)
Value:
template <> \
struct joint_matrix_hip< \
TYPE, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, \
sycl::ext::oneapi::experimental::matrix::layout::dynamic> { \
sycl::marray<TYPE, (M * N) / WAVEFRONT_SIZE> wi_marray; \
};

Definition at line 89 of file matrix-hip.hpp.