matrix-hip.hpp File Reference

#include "matrix-unified-utils.hpp"
#include <sycl/access/access.hpp>
#include <sycl/ext/oneapi/bfloat16.hpp>
#include <sycl/marray.hpp>
#include <sycl/multi_ptr.hpp>
#include <cstring>

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
	Objects of the accessor_property_list class are containers for the SYCL properties.

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 21 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 64 of file matrix-hip.hpp.

__SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC

#define __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC	(		TYPE,
			M,
			N
	)

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 96 of file matrix-hip.hpp.