diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index c8a9516b0309b..c3a4dac369c88 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -27,6 +27,7 @@ extern __DPCPP_SYCL_EXTERNAL float __spirv_RoundFToTF32INTEL(float a); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX template * __spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL *, Ts val, size_t i); +#else // __SPIRV_USE_COOPERATIVE_MATRIX +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CooperativeMatrixLoadKHR(T *Ptr, __spv::MatrixLayout Layout = L, + std::size_t Stride = 0, + int MemOperand = 0); + +template +extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixStoreKHR( + T *Ptr, __spv::__spirv_CooperativeMatrixKHR *Object, + __spv::MatrixLayout Layout = L, std::size_t Stride = 0, int MemOperand = 0); + +template +extern __DPCPP_SYCL_EXTERNAL size_t __spirv_CooperativeMatrixLengthKHR( + __spv::__spirv_CooperativeMatrixKHR *); + +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CooperativeMatrixConstructCheckedINTEL(const T Value, size_t Height, + size_t Stride, size_t Width, + size_t CoordX, + size_t CoordY); + +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CooperativeMatrixLoadCheckedINTEL(T *Ptr, std::size_t Stride, + size_t Height, size_t Width, + size_t CoordX, size_t CoordY, + __spv::MatrixLayout Layout = L, + int MemOperand = 0); + +template +extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixStoreCheckedINTEL( + T *Ptr, __spv::__spirv_CooperativeMatrixKHR *Object, + std::size_t Stride, size_t Height, size_t Width, size_t CoordX, + size_t CoordY, __spv::MatrixLayout Layout = L, int MemOperand = 0); + +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CooperativeMatrixMulAddKHR( + __spv::__spirv_CooperativeMatrixKHR *A, + __spv::__spirv_CooperativeMatrixKHR *B, + __spv::__spirv_CooperativeMatrixKHR *C, + size_t Operands = 0); + +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CompositeConstruct(const T v); + +// TODO: replace with __spirv_CooperativeMatrixGetElementCoordINTEL when ready +template +extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t +__spirv_JointMatrixGetElementCoordINTEL( + __spv::__spirv_CooperativeMatrixKHR *, size_t i); + +// AccessChain followed by load/store serves to extract/insert and element +// from/to the matrix +template +extern __DPCPP_SYCL_EXTERNAL Ts * +__spirv_AccessChain(__spv::__spirv_CooperativeMatrixKHR **, + size_t i); + +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CooperativeMatrixConstructCheckedINTEL(int32_t CoordX, + int32_t CoordY, + uint32_t Height, + uint32_t Width, + const T Value); + +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CooperativeMatrixLoadCheckedINTEL( + T *Ptr, int32_t CoordX, int32_t CoordY, __spv::MatrixLayout Layout = L, + uint32_t Height = 0, uint32_t Width = 0, std::size_t Stride = 0, + int MemOperand = 0); + +template +extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixStoreCheckedINTEL( + T *Ptr, int32_t CoordX, int32_t CoordY, + __spv::__spirv_CooperativeMatrixKHR *Object, + __spv::MatrixLayout Layout = L, uint32_t Height = 0, uint32_t Width = 0, + std::size_t Stride = 0, int MemOperand = 0); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX template extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL( diff --git a/sycl/include/CL/__spirv/spirv_types.hpp b/sycl/include/CL/__spirv/spirv_types.hpp index 013c2e1acc564..21563c3b8a67a 100644 --- a/sycl/include/CL/__spirv/spirv_types.hpp +++ b/sycl/include/CL/__spirv/spirv_types.hpp @@ -118,10 +118,34 @@ enum class MatrixLayout : uint32_t { enum class MatrixUse : uint32_t { MatrixA = 0, MatrixB = 1, Accumulator = 2 }; +#ifdef __SPIRV_USE_COOPERATIVE_MATRIX +enum class MatrixOperands : uint32_t { + // SPV_KHR_cooperative_matrix operands + NoneKHR = 0, + MatrixASignedComponentsKHR = 0x1, + MatrixBSignedComponentsKHR = 0x2, + MatrixCSignedComponentsKHR = 0x4, + MatrixResultSignedComponentsKHR = 0x8, + SaturatingAccumulationKHR = 0x10, + // SPV_INTEL_joint_matrix operands + MatrixAAndBTF32ComponentsINTEL = 0x20, + MatrixAAndBBFloat16ComponentsINTEL = 0x40, + MatrixCBFloat16ComponentsINTEL = 0x80, + MatrixResultBFloat16ComponentsINTEL = 0x100 +}; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX + +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX + template struct __spirv_JointMatrixINTEL; +#else +template +struct __spirv_CooperativeMatrixKHR; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX struct __spirv_TaskSequenceINTEL; diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp index 9fb099e7eb8e4..a3749a0137e78 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp @@ -115,8 +115,13 @@ class wi_element { inline __SYCL_ALWAYS_INLINE std::tuple get_coord() { #if defined(__SYCL_DEVICE_ONLY__) +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __ocl_vec_t coord = __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); +#else + __ocl_vec_t coord = + __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX const size_t row = coord[0]; const size_t col = coord[1]; return std::make_tuple(row, col); @@ -128,12 +133,20 @@ class wi_element { operator storage_element_type() { #ifdef __SYCL_DEVICE_ONLY__ +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX storage_element_type elem = __spirv_VectorExtractDynamic::value, spv_matrix_layout_traits::value, spv_scope_traits::value>(M.spvm, idx); +#else + storage_element_type *ExtractP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&M.spvm, idx); + storage_element_type elem = *ExtractP; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX return elem; #else throw exception(make_error_code(errc::runtime), @@ -143,12 +156,20 @@ class wi_element { explicit operator bool() { #ifdef __SYCL_DEVICE_ONLY__ +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX return __spirv_VectorExtractDynamic::value, spv_matrix_layout_traits::value, spv_scope_traits::value>( M.spvm, idx) != static_cast(0); +#else + storage_element_type *ExtractP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&M.spvm, idx); + return *ExtractP != static_cast(0); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -157,8 +178,16 @@ class wi_element { template wi_element &operator=(const T2 &rhs) { #ifdef __SYCL_DEVICE_ONLY__ +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX M.spvm = __spirv_VectorInsertDynamic( M.spvm, static_cast(rhs), idx); +#else + storage_element_type *InsertP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&M.spvm, idx); + *InsertP = static_cast(rhs); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX return *this; #else (void)rhs; @@ -170,6 +199,7 @@ class wi_element { wi_element & operator=(const wi_element &rhs) { #ifdef __SYCL_DEVICE_ONLY__ +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX M.spvm = __spirv_VectorInsertDynamic( M.spvm, __spirv_VectorExtractDynamic::value>(rhs.M.spvm, rhs.idx), idx); +#else + storage_element_type *ExtractP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&rhs.M.spvm, + rhs.idx); + storage_element_type *InsertP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&M.spvm, idx); + *InsertP = *ExtractP; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX return *this; #else (void)rhs; @@ -187,6 +229,7 @@ class wi_element { } #if __SYCL_DEVICE_ONLY__ +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX #define OP(op) \ template wi_element &operator op##=(const T2 & rhs) { \ M.spvm = __spirv_VectorInsertDynamic( \ @@ -201,6 +244,22 @@ class wi_element { idx); \ return *this; \ } +#else // __SPIRV_USE_COOPERATIVE_MATRIX +#define OP(op) \ + template wi_element &operator op##=(const T2 & rhs) { \ + storage_element_type *ExtractP = \ + __spirv_AccessChain::value, \ + spv_scope_traits::value>(&rhs.M.spvm, \ + rhs.idx); \ + storage_element_type *InsertP = \ + __spirv_AccessChain::value, \ + spv_scope_traits::value>(&M.spvm, idx); \ + *InsertP = *ExtractP op static_cast(rhs); \ + return *this; \ + } +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else // __SYCL_DEVICE_ONLY__ #define OP(op) \ template wi_element &operator op##=(const T2 & rhs) { \ @@ -235,8 +294,13 @@ class wi_element get_coord() { #if defined(__SYCL_DEVICE_ONLY__) +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __ocl_vec_t coord = __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); +#else + __ocl_vec_t coord = + __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX const uint32_t row = coord[0]; const uint32_t col = coord[1]; return std::make_tuple(row, col); @@ -248,11 +312,20 @@ class wi_element::value, spv_matrix_layout_traits::value, spv_scope_traits::value>(M.spvm, idx); +#else + sycl::ext::oneapi::bfloat16 *ExtractP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&M.spvm, idx); + return *ExtractP; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -261,6 +334,7 @@ class wi_element( __spirv_VectorExtractDynamic< sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, @@ -268,6 +342,16 @@ class wi_element::value, spv_scope_traits::value>(M.spvm, idx))) >= std::numeric_limits::epsilon(); +#else + sycl::ext::oneapi::bfloat16 *ExtractP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&M.spvm, idx); + sycl::ext::oneapi::bfloat16 Elem = *ExtractP; + return sycl::fabs(static_cast(Elem)) >= + std::numeric_limits::epsilon(); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -276,7 +360,16 @@ class wi_element::value, + spv_scope_traits::value>(&M.spvm, idx); + *InsertP = rhs; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX return *this; #else (void)rhs; @@ -288,6 +381,7 @@ class wi_element &rhs) { #ifdef __SYCL_DEVICE_ONLY__ +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX M.spvm = __spirv_VectorInsertDynamic( M.spvm, __spirv_VectorExtractDynamic::value, + spv_scope_traits::value>(&rhs.M.spvm, + rhs.idx); + sycl::ext::oneapi::bfloat16 *InsertP = + __spirv_AccessChain::value, + spv_scope_traits::value>(&M.spvm, idx); + *InsertP = *ExtractP; + return *this; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else (void)rhs; throw exception(make_error_code(errc::runtime), @@ -306,6 +415,7 @@ class wi_element::value, \ + spv_scope_traits::value>(&M.spvm, idx); \ + sycl::ext::oneapi::bfloat16 *InsertP = \ + __spirv_AccessChain::value, \ + spv_scope_traits::value>(&M.spvm, idx); \ + *InsertP = *ExtractP op rhs; \ + return *this; \ + } +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else // __SYCL_DEVICE_ONLY__ #define OP(opassign, op) \ wi_element &operator opassign(const sycl::ext::oneapi::bfloat16 & rhs) { \ @@ -333,6 +460,7 @@ class wi_element::value, \ spv_scope_traits::value>(rhs.M.spvm, rhs.idx) op lhs; \ } +#else +#define OP(type, op) \ + friend type operator op( \ + const wi_element &lhs, \ + const sycl::ext::oneapi::bfloat16 &rhs) { \ + sycl::ext::oneapi::bfloat16 *ExtractP = \ + __spirv_AccessChain::value, \ + spv_scope_traits::value>(&lhs.M.spvm, \ + lhs.idx); \ + return *ExtractP op rhs; \ + } \ + friend type operator op( \ + const sycl::ext::oneapi::bfloat16 &lhs, \ + const wi_element &rhs) { \ + sycl::ext::oneapi::bfloat16 *ExtractP = \ + __spirv_AccessChain::value, \ + spv_scope_traits::value>(&rhs.M.spvm, \ + rhs.idx); \ + return *ExtractP op lhs; \ + } +#endif // __SPIRV_USE_COOPERATIVE_MATRIX OP(sycl::ext::oneapi::bfloat16, +) OP(sycl::ext::oneapi::bfloat16, -) OP(sycl::ext::oneapi::bfloat16, *) OP(sycl::ext::oneapi::bfloat16, /) #undef OP +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX #define OP(type, op) \ friend type operator op( \ const wi_element::value>(rhs.M.spvm, rhs.idx)) \ op static_cast(lhs)}; \ } +#else +#define OP(type, op) \ + friend type operator op( \ + const wi_element &lhs, \ + const sycl::ext::oneapi::bfloat16 &rhs) { \ + sycl::ext::oneapi::bfloat16 *ExtractP = \ + __spirv_AccessChain::value, \ + spv_scope_traits::value>(&lhs.M.spvm, \ + lhs.idx); \ + return type{static_cast(*ExtractP) op static_cast(rhs)}; \ + } \ + friend type operator op( \ + const sycl::ext::oneapi::bfloat16 &lhs, \ + const wi_element &rhs) { \ + sycl::ext::oneapi::bfloat16 *ExtractP = \ + __spirv_AccessChain::value, \ + spv_scope_traits::value>(&rhs.M.spvm, \ + rhs.idx); \ + return type{static_cast(*ExtractP) op static_cast(lhs)}; \ + } +#endif // __SPIRV_USE_COOPERATIVE_MATRIX OP(bool, ==) OP(bool, !=) OP(bool, <) @@ -448,7 +631,11 @@ class wi_data { public: size_t length() { #if __SYCL_DEVICE_ONLY__ +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX return __spirv_JointMatrixWorkItemLengthINTEL(jm.spvm); +#else + return __spirv_CooperativeMatrixLengthKHR(jm.spvm); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -503,6 +690,7 @@ joint_matrix_store(Group, // intel's impl using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(dst); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __spirv_JointMatrixStoreINTEL::value, @@ -512,6 +700,18 @@ joint_matrix_store(Group, sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< Layout>::value, sycl::ext::oneapi::experimental::matrix::spv_scope_traits::value); +#else + __spirv_CooperativeMatrixStoreKHR< + DecorT, Tp, NumRows, NumCols, + sycl::ext::oneapi::experimental::matrix::spv_matrix_use_traits< + Use>::value, + sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< + Layout>::value>( + Ptr, src.spvm, + sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< + Layout>::value, + stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = src; @@ -548,6 +748,7 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( #else // intel's impl T *Ptr = dst.get(); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __spirv_JointMatrixStoreINTEL::value, @@ -557,6 +758,18 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< Layout>::value, sycl::ext::oneapi::experimental::matrix::spv_scope_traits::value); +#else + __spirv_CooperativeMatrixStoreKHR< + T, Tp, NumRows, NumCols, + sycl::ext::oneapi::experimental::matrix::spv_matrix_use_traits< + Use>::value, + sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< + Layout>::value>( + Ptr, src.spvm, + sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< + Layout>::value, + stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = src; diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index cba25494ca65d..d3d57f24c56e6 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -50,9 +50,15 @@ struct joint_matrix { sycl::ext::oneapi::detail::joint_matrix_hip matrix_impl; #elif defined(__SPIR__) || defined(__SPIRV__) +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __spv::__spirv_JointMatrixINTEL< T, Rows, Cols, spv_matrix_layout_traits::value, spv_scope_traits::value, spv_matrix_use_traits::value> *spvm; +#else + __spv::__spirv_CooperativeMatrixKHR::value, Rows, + Cols, spv_matrix_use_traits::value> + *spvm; +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else static_assert(false, "The joint_matrix API is only supported by the Intel, " "CUDA and HIP (GFX90A) backends"); @@ -200,12 +206,20 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load( std::ignore = sg; using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(src); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX res.spvm = __spirv_JointMatrixLoadINTEL< DecorT, S, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), spv_scope_traits::value); +#else + res.spvm = __spirv_CooperativeMatrixLoadKHR< + DecorT, S, NumRows, NumCols, + spv_matrix_use_traits::value, + spv_matrix_layout_traits::value>( + Ptr, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -246,12 +260,20 @@ joint_matrix_load(Group sg, std::ignore = sg; using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(src); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX res.spvm = __spirv_JointMatrixLoadINTEL::value, spv_matrix_layout_traits::value>( Ptr, stride, spv_matrix_layout_traits::value, spv_scope_traits::value); +#else + res.spvm = + __spirv_CooperativeMatrixLoadKHR::value, + spv_matrix_layout_traits::value>( + Ptr, spv_matrix_layout_traits::value, stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -284,11 +306,18 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load( #else std::ignore = sg; T *Ptr = src.get(); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX res.spvm = __spirv_JointMatrixLoadINTEL< T, S, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), spv_scope_traits::value); +#else + res.spvm = __spirv_CooperativeMatrixLoadKHR< + T, S, NumRows, NumCols, spv_matrix_use_traits::value, + spv_matrix_layout_traits::value>( + Ptr, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -323,12 +352,20 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load( #else std::ignore = sg; T *Ptr = src.get(); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX res.spvm = __spirv_JointMatrixLoadINTEL::value, spv_matrix_layout_traits::value>( Ptr, stride, spv_matrix_layout_traits::value, spv_scope_traits::value); +#else + res.spvm = + __spirv_CooperativeMatrixLoadKHR::value, + spv_matrix_layout_traits::value>( + Ptr, spv_matrix_layout_traits::value, stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -365,12 +402,20 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( std::ignore = sg; using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(dst); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __spirv_JointMatrixStoreINTEL< DecorT, T, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, src.spvm, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), spv_scope_traits::value); +#else + __spirv_CooperativeMatrixStoreKHR< + DecorT, T, NumRows, NumCols, + spv_matrix_use_traits::value, + spv_matrix_layout_traits::value>( + Ptr, src.spvm, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -403,11 +448,18 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( #else std::ignore = sg; T *Ptr = dst.get(); +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __spirv_JointMatrixStoreINTEL< T, T, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, src.spvm, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), spv_scope_traits::value); +#else + __spirv_CooperativeMatrixStoreKHR< + T, T, NumRows, NumCols, spv_matrix_use_traits::value, + spv_matrix_layout_traits::value>( + Ptr, src.spvm, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -464,8 +516,9 @@ joint_matrix_mad( "requires that joint_matrix data types Ta and Tb match"); } #else - if constexpr (std::is_same::value && - std::is_same::value && +#ifndef __SPIRV_USE_COOPERATIVE_MATRIX + if constexpr (std::is_same::value && + std::is_same::value && std::is_same::value) D.spvm = __spirv_JointMatrixMadINTEL(A.spvm, B.spvm, C.spvm); else if constexpr (std::is_unsigned::value && std::is_unsigned::value) @@ -476,6 +529,38 @@ joint_matrix_mad( D.spvm = __spirv_JointMatrixUSMadINTEL(A.spvm, B.spvm, C.spvm); else D.spvm = __spirv_JointMatrixMadINTEL(A.spvm, B.spvm, C.spvm); +#else + if constexpr (std::is_same::value && + std::is_same::value && + std::is_same::value) { + constexpr uint32_t MatrixOperand = static_cast( + __spv::MatrixOperands::MatrixAAndBBFloat16ComponentsINTEL); + D.spvm = __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm, + MatrixOperand); + } else if constexpr (std::is_signed::value && + std::is_unsigned::value) { + constexpr uint32_t MatrixOperand = static_cast( + __spv::MatrixOperands::MatrixASignedComponentsKHR); + D.spvm = __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm, + MatrixOperand); + } else if constexpr (std::is_unsigned::value && + std::is_signed::value) { + constexpr uint32_t MatrixOperand = static_cast( + __spv::MatrixOperands::MatrixBSignedComponentsKHR); + D.spvm = __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm, + MatrixOperand); + } else if constexpr (std::is_signed::value && std::is_signed::value) { + constexpr uint32_t MatrixOperand = + static_cast( + __spv::MatrixOperands::MatrixASignedComponentsKHR) + + static_cast( + __spv::MatrixOperands::MatrixBSignedComponentsKHR); + D.spvm = __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm, + MatrixOperand); + } else { + D.spvm = __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm); + } +#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = A; diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp new file mode 100644 index 0000000000000..16e1951eab262 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp @@ -0,0 +1,22 @@ +//==----------- element_wise_abc.cpp - DPC++ joint_matrix------------- ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../element_wise_abc_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp new file mode 100644 index 0000000000000..6c80692109ca8 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp @@ -0,0 +1,22 @@ +//==------------ element_wise_all_ops.cpp - DPC++ joint_matrix-------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp new file mode 100644 index 0000000000000..c90aa1b824a6d --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp @@ -0,0 +1,23 @@ +//==----------- element_wise_all_ops_half.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-fp16 +// REQUIRES: aspect-ext_intel_matrix, gpu +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../element_wise_all_ops_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp new file mode 100644 index 0000000000000..4d7fdb9285023 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp @@ -0,0 +1,22 @@ +//==----------- element_wise_all_ops_int8.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../element_wise_all_ops_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp new file mode 100644 index 0000000000000..87ede89ab00c8 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp @@ -0,0 +1,24 @@ +//==------ element_wise_all_ops_int8_packed.cpp - DPC++ joint_matrix-------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +// This test stores the matrix B that is VNNIed (packed). + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../element_wise_all_ops_int8_packed_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp new file mode 100644 index 0000000000000..613bcd1f7650e --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp @@ -0,0 +1,21 @@ +//==----------- element_wise_all_ops_tf32.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix-tf32 +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 +constexpr size_t TN = 16; + +#include "../../element_wise_all_ops_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp new file mode 100644 index 0000000000000..b473dc00a1844 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp @@ -0,0 +1,20 @@ +//==----------- element_wise_all_sizes.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../element_wise_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp new file mode 100644 index 0000000000000..611c369b99011 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp @@ -0,0 +1,22 @@ +//==----------- element_wise_ops.cpp - DPC++ joint_matrix------------- ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../element_wise_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp new file mode 100644 index 0000000000000..8ef78f76b3509 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp @@ -0,0 +1,22 @@ +//==----------- get_coord_float_matC.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../get_coord_float_matC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp new file mode 100644 index 0000000000000..9d3e62726720c --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp @@ -0,0 +1,22 @@ +//==----------- get_coord_int8_matA.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../get_coord_int8_matA_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp new file mode 100644 index 0000000000000..0b6dac6047681 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp @@ -0,0 +1,21 @@ +//==----------- get_coord_int8_matB.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../get_coord_int8_matB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp new file mode 100644 index 0000000000000..cddc7659167a2 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp @@ -0,0 +1,22 @@ +//==-------- joint_matrix_all_sizes.cpp - DPC++ joint_matrix---------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp new file mode 100644 index 0000000000000..21b9014e7bff9 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp @@ -0,0 +1,24 @@ +//==-------- joint_matrix_annotated_ptr.cpp - DPC++ joint_matrix-----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +// Currently row major B fails when annotated_ptr is used +// XFAIL: gpu + +#include "../../common.hpp" + +#define SG_SZ 32 +constexpr size_t TN = 16; + +#include "../../joint_matrix_annotated_ptr_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp new file mode 100644 index 0000000000000..b5fb8bf2c6dfe --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp @@ -0,0 +1,22 @@ +//==----------- joint_matrix_apply_bf16.cpp - DPC++ joint_matrix-----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_apply_bf16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp new file mode 100644 index 0000000000000..5fa27fc66df1c --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp @@ -0,0 +1,23 @@ +//==------ joint_matrix_apply_two_matrices.cpp - DPC++ joint_matrix--------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -ffp-model=precise -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_apply_two_matrices_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp new file mode 100644 index 0000000000000..91175103be6cc --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp @@ -0,0 +1,27 @@ +//==--- joint_matrix_bf16_fill_k_cache.cpp - DPC++ joint_matrix----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_vnni.out -DVNNI -ffp-model=precise +// RUN: %{run} %t_vnni.out + +// TODO: add row major compilation and run once Sub-group size 32 +// support becomes available in IGC for row major + +// XFAIL: cpu + +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp new file mode 100644 index 0000000000000..eda198c7b2a41 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp @@ -0,0 +1,22 @@ +//==---joint_matrix_bf16_fill_k_cache_init.cpp - DPC++ joint_matrix--------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix, gpu +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -DINIT_LIST -DVNNI -ffp-model=precise +// RUN: %{run} %t.out + +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp new file mode 100644 index 0000000000000..df0b207c0605b --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp @@ -0,0 +1,26 @@ +//==---joint_matrix_bf16_fill_k_cache_unroll.cpp - DPC++ joint_matrix------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -mllvm -inline-threshold=5000 -ffp-model=precise -o %t.out -DMANUAL_UNROLL -DVNNI +// RUN: %{run} %t.out + +// XFAIL: cpu + +// -mllvm -inline-threshold added as a workaround, +// since IGC doesn't support some variants of IR for Joint Matrix currently +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp new file mode 100644 index 0000000000000..7fdd5e62736bc --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp @@ -0,0 +1,24 @@ +//==--joint_matrix_bf16_fill_k_cache_unroll_init.cpp - DPC++ joint_matrix--==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix, gpu +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -mllvm -inline-threshold=5000 -ffp-model=precise -o %t_gpu.out -DINIT_LIST -DMANUAL_UNROLL -DVNNI +// RUN: %{run} %t_gpu.out + +// -mllvm -inline-threshold added as a workaround, +// since IGC doesn't support some variants of IR for Joint Matrix currently +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp new file mode 100644 index 0000000000000..637e9ebcd858c --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp @@ -0,0 +1,22 @@ +//==-------- joint_matrix_bfloat16.cpp - DPC++ joint_matrix----------- ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_bfloat16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp new file mode 100644 index 0000000000000..eee85175d678d --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp @@ -0,0 +1,22 @@ +//==-------- joint_matrix_bfloat16_array.cpp - DPC++ joint_matrix----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_bfloat16_array_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp new file mode 100644 index 0000000000000..80de9af24bb85 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp @@ -0,0 +1,31 @@ +//==-- joint_matrix_bfloat16_colmajorA_colmajorB.cpp - DPC++ joint_matrix--==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// This tests support of col major layout for matrix B which does transpose and +// then VNNI transform. This is currently only available on AMX + +// XFAIL: gpu + +#include "../../common.hpp" +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; +using bfloat16 = sycl::ext::oneapi::bfloat16; + +#define SG_SZ 32 +constexpr size_t TN = 16; + +#include "../../joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp new file mode 100644 index 0000000000000..aa8e4bf7758ca --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp @@ -0,0 +1,20 @@ +//==----- joint_matrix_bfloat16_packedB.cpp - DPC++ joint_matrix----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 +#include "../../joint_matrix_bfloat16_packedB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp new file mode 100644 index 0000000000000..6a51e03bebd23 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp @@ -0,0 +1,21 @@ +//==---------- joint_matrix_colA_rowB_colC.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL:* + +#include "../../common.hpp" + +#define SG_SZ 32 +constexpr size_t TN = 16; + +#include "../../joint_matrix_colA_rowB_colC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp new file mode 100644 index 0000000000000..efe3c2b456802 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp @@ -0,0 +1,20 @@ +//==-------- joint_matrix_down_convert.cpp - DPC++ joint_matrix------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_down_convert_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp new file mode 100644 index 0000000000000..16953caff99e0 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp @@ -0,0 +1,24 @@ +//==-------- joint_matrix_half.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-fp16 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES: matrix-fp16 +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp new file mode 100644 index 0000000000000..375e3bb958733 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp @@ -0,0 +1,27 @@ +//==----- joint_matrix_int8_colmajorA_colmajorB.cpp - DPC++ joint_matrix---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// This tests support of col major layout for matrix B which does transpose and +// then VNNI transform. This is currently only available on AMX + +// XFAIL: gpu + +#include "../../common.hpp" + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; + +#define SG_SZ 32 +constexpr size_t TN = 16; + +#include "../../joint_matrix_int8_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp new file mode 100644 index 0000000000000..654bdfe695116 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp @@ -0,0 +1,22 @@ +//==--- joint_matrix_int8_rowmajorA_rowmajorB.cpp - DPC++ joint_matrix-----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp new file mode 100644 index 0000000000000..6926e7951b0db --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp @@ -0,0 +1,24 @@ +//==-------- joint_matrix_out_bounds.cpp - DPC++ joint_matrix--------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// UNSUPPORTED: gpu-intel-dg2, cpu + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL:gpu + +#include "../../common.hpp" + +#define SG_SZ 32 +constexpr size_t TN = 16; +constexpr size_t MATRIX_K = 1024 + 24; + +#include "../../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp new file mode 100644 index 0000000000000..7073feae64ac4 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp @@ -0,0 +1,23 @@ +//==-------- joint_matrix_prefetch.cpp - DPC++ joint_matrix----------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// REQUIRES: aspect-ext_intel_matrix +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + +#include "../../common.hpp" + +#define SG_SZ 32 +constexpr size_t TN = 16; +#include "../../joint_matrix_prefetch_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp new file mode 100644 index 0000000000000..610b2b5bf6e5c --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp @@ -0,0 +1,29 @@ +//==--------joint_matrix_rowmajorA_rowmajorB.cpp - DPC++ joint_matrix------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// This tests support of row major layout for matrix B which does automatic VNNI +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// VNNI transform and sub-group size 32 are not supported yet on DG2 by IGC +// UNSUPPORTED: gpu-intel-dg2 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +// Sub-group size 32 support for this test is not currently available in IGC +// XFAIL: gpu + +#include "../../common.hpp" + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; + +#define SG_SZ 32 + +#include "../../joint_matrix_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp new file mode 100644 index 0000000000000..fce90d5a42fa3 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp @@ -0,0 +1,22 @@ +//==-------- joint_matrix_ss_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_ss_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp new file mode 100644 index 0000000000000..c3d0302a3f187 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp @@ -0,0 +1,22 @@ +//==-------- joint_matrix_su_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_su_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp new file mode 100644 index 0000000000000..e2ae342ed4598 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp @@ -0,0 +1,24 @@ +//==---------------- joint_matrix_tf32.cpp - DPC++ joint_matrix------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix-tf32 +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; + +#define SG_SZ 32 +constexpr size_t TN = 16; + +#include "../../joint_matrix_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp new file mode 100644 index 0000000000000..d0fd090b2e371 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp @@ -0,0 +1,20 @@ +//==----------- joint_matrix_transposeC.cpp - DPC++ joint_matrix-----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_transposeC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp new file mode 100644 index 0000000000000..d0ac32d7661e6 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp @@ -0,0 +1,24 @@ +//==-------- joint_matrix_unaligned_k.cpp - DPC++ joint_matrix-------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// UNSUPPORTED: gpu-intel-dg2, cpu + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu + +#include "../../common.hpp" + +#define SG_SZ 32 +constexpr size_t TN = 16; +static constexpr size_t MATRIX_K = 1024 + 14; + +#include "../../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp new file mode 100644 index 0000000000000..56ae3f112bb85 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp @@ -0,0 +1,22 @@ +//==-------- joint_matrix_us_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_us_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp new file mode 100644 index 0000000000000..daf87d386d3c8 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp @@ -0,0 +1,22 @@ +//==-------- joint_matrix_uu_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../../common.hpp" + +#define SG_SZ 32 + +#include "../../joint_matrix_uu_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_abc.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_abc.cpp new file mode 100644 index 0000000000000..9a27ccf807076 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_abc.cpp @@ -0,0 +1,17 @@ +//==----------- element_wise_abc.cpp - DPC++ joint_matrix------------- ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_abc_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp new file mode 100644 index 0000000000000..5cce8d34344c7 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp @@ -0,0 +1,17 @@ +//==------------ element_wise_all_ops.cpp - DPC++ joint_matrix-------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d.cpp new file mode 100644 index 0000000000000..73afef5afa996 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d.cpp @@ -0,0 +1,16 @@ +//==-------- element_wise_all_ops_1d.cpp - DPC++ joint_matrix ---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix, gpu + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out + +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d_cont.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d_cont.cpp new file mode 100644 index 0000000000000..076e1c9114523 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d_cont.cpp @@ -0,0 +1,16 @@ +//==-------- element_wise_all_ops_1d_cont.cpp - DPC++ joint_matrix ---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix, gpu + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out + +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp new file mode 100644 index 0000000000000..8853c0306edb1 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp @@ -0,0 +1,19 @@ +//==----------- element_wise_all_ops_half.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-fp16 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES: matrix-fp16 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_all_ops_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp new file mode 100644 index 0000000000000..ba10657af03ae --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp @@ -0,0 +1,17 @@ +//==----------- element_wise_all_ops_int8.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_all_ops_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp new file mode 100644 index 0000000000000..105c1ed763154 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp @@ -0,0 +1,19 @@ +//==------ element_wise_all_ops_int8_packed.cpp - DPC++ joint_matrix-------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +// This test stores the matrix B that is VNNIed (packed). + +#include "../common.hpp" +#include "../element_wise_all_ops_int8_packed_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp new file mode 100644 index 0000000000000..d86cd6f469a66 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp @@ -0,0 +1,17 @@ +//==---- element_wise_all_ops_scalar.cpp - DPC++ joint_matrix ------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix, gpu + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp new file mode 100644 index 0000000000000..cc26efe30d8b8 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp @@ -0,0 +1,20 @@ +//==----------- element_wise_all_ops_tf32.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix-tf32 +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" + +constexpr size_t TN = 16; + +#include "../element_wise_all_ops_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp new file mode 100644 index 0000000000000..27b2f4b21f0dc --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp @@ -0,0 +1,23 @@ +//==----------- element_wise_all_sizes.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// This is a version of the test with disabled device code +// split to test against fixed bug in IGC +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -fsycl-device-code-split=off -o %t_split.out +// RUN: %if gpu-intel-dg2 %{ %{run} %t_split.out %} + +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../element_wise_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_ops.cpp new file mode 100644 index 0000000000000..e45b7507ed828 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_ops.cpp @@ -0,0 +1,17 @@ +//==----------- element_wise_ops.cpp - DPC++ joint_matrix------------- ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu +// XFAIL: cpu + +#include "../common.hpp" +#include "../element_wise_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp new file mode 100644 index 0000000000000..41d9971c6f003 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp @@ -0,0 +1,17 @@ +//==----------- get_coord_float_matC.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu +// XFAIL: cpu + +#include "../common.hpp" +#include "../get_coord_float_matC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp new file mode 100644 index 0000000000000..e9f19ccc5b04b --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp @@ -0,0 +1,17 @@ +//==----------- get_coord_int8_matA.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu +// XFAIL: cpu + +#include "../common.hpp" +#include "../get_coord_int8_matA_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp new file mode 100644 index 0000000000000..0f370e1be8898 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp @@ -0,0 +1,16 @@ +//==----------- get_coord_int8_matB.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../get_coord_int8_matB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp new file mode 100644 index 0000000000000..10d6d9ee62d56 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp @@ -0,0 +1,16 @@ +//==-------- joint_matrix_all_sizes.cpp - DPC++ joint_matrix---------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp new file mode 100644 index 0000000000000..52e680fb25070 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp @@ -0,0 +1,21 @@ +//==-------- joint_matrix_annotated_ptr.cpp - DPC++ joint_matrix-----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} + +// XFAIL: cpu + +#include "../common.hpp" + +constexpr size_t TN = 16; + +#include "../joint_matrix_annotated_ptr_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp new file mode 100644 index 0000000000000..24bf88f9d3f34 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp @@ -0,0 +1,17 @@ +//==----------- joint_matrix_apply_bf16.cpp - DPC++ joint_matrix-----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_apply_bf16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp new file mode 100644 index 0000000000000..1d1927c3a1d5d --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp @@ -0,0 +1,17 @@ +//==------ joint_matrix_apply_two_matrices.cpp - DPC++ joint_matrix--------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -ffp-model=precise -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: gpu +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_apply_two_matrices_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp new file mode 100644 index 0000000000000..eb68e332eaa7c --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp @@ -0,0 +1,21 @@ +//==--- joint_matrix_bf16_fill_k_cache.cpp - DPC++ joint_matrix----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_vnni.out -DVNNI -ffp-model=precise +// RUN: %{run} %t_vnni.out + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -ffp-model=precise +// RUN: %{run} %t.out + +// XFAIL: cpu + +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../common.hpp" +#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp new file mode 100644 index 0000000000000..af4b28e090b4c --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp @@ -0,0 +1,20 @@ +//==--- joint_matrix_bf16_fill_k_cache_OOB.cpp - DPC++ joint_matrix--------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix, gpu +// UNSUPPORTED: gpu-intel-dg2 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu_vnni.out -ffp-model=precise -DOOB -DVNNI +// RUN: %{run} %t_gpu_vnni.out + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu.out -ffp-model=precise -DOOB +// RUN: %{run} %t_gpu.out + +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../common.hpp" +#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_init.cpp new file mode 100644 index 0000000000000..f4c77bb50c5da --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_init.cpp @@ -0,0 +1,16 @@ +//==---joint_matrix_bf16_fill_k_cache_init.cpp - DPC++ joint_matrix--------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix, gpu + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -DINIT_LIST -DVNNI -ffp-model=precise +// RUN: %{run} %t.out + +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../common.hpp" +#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp new file mode 100644 index 0000000000000..89e9048170837 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp @@ -0,0 +1,20 @@ +//==---joint_matrix_bf16_fill_k_cache_unroll.cpp - DPC++ joint_matrix------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -mllvm -inline-threshold=2000 -ffp-model=precise -o %t.out -DMANUAL_UNROLL -DVNNI +// RUN: %{run} %t.out + +// XFAIL: cpu + +// -mllvm -inline-threshold=2000 added as a workaround, +// since IGC doesn't support some variants of IR for Joint Matrix currently +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../common.hpp" +#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp new file mode 100644 index 0000000000000..d653268345dcc --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp @@ -0,0 +1,18 @@ +//==--joint_matrix_bf16_fill_k_cache_unroll_init.cpp - DPC++ joint_matrix--==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix, gpu + +// RUN: %{build} -mllvm -inline-threshold=2000 -ffp-model=precise -o %t_gpu.out -DINIT_LIST -DMANUAL_UNROLL -DVNNI +// RUN: %{run} %t_gpu.out + +// -mllvm -inline-threshold=2000 added as a workaround, +// since IGC doesn't support some variants of IR for Joint Matrix currently +// -ffp-model=precise is added to not depend on compiler defaults. + +#include "../common.hpp" +#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp new file mode 100644 index 0000000000000..c6f1db06e625f --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp @@ -0,0 +1,16 @@ +//==-------- joint_matrix_bfloat16.cpp - DPC++ joint_matrix----------- ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_bfloat16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp new file mode 100644 index 0000000000000..69f9aa8553bd3 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp @@ -0,0 +1,16 @@ +//==-------- joint_matrix_bfloat16_array.cpp - DPC++ joint_matrix----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_bfloat16_array_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp new file mode 100644 index 0000000000000..ff7d2f1e05882 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp @@ -0,0 +1,22 @@ +//==-- joint_matrix_bfloat16_colmajorA_colmajorB.cpp - DPC++ joint_matrix--==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// This tests support of col major layout for matrix B which does transpose and +// then VNNI transform. This is currently only available on AMX + +// XFAIL: gpu + +#include "../common.hpp" + +constexpr size_t TN = 16; + +#include "../joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp new file mode 100644 index 0000000000000..1b7de0a673641 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp @@ -0,0 +1,20 @@ +//==----- joint_matrix_bfloat16_packedB.cpp - DPC++ joint_matrix----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix +// REQUIRES-INTEL-DRIVER: lin: 27868, win: 101.5181 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out %} +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_bfloat16_packedB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp new file mode 100644 index 0000000000000..2c6323364ead9 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp @@ -0,0 +1,19 @@ +//==---------- joint_matrix_colA_rowB_colC.cpp - DPC++ joint_matrix---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL:* + +#include "../common.hpp" + +constexpr size_t TN = 16; + +#include "../joint_matrix_colA_rowB_colC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp new file mode 100644 index 0000000000000..760402d0a35cc --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp @@ -0,0 +1,17 @@ +//==-------- joint_matrix_down_convert.cpp - DPC++ joint_matrix------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu +// XFAIL: gpu + +#include "../common.hpp" +#include "../joint_matrix_down_convert_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp new file mode 100644 index 0000000000000..43b70b990b7d1 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp @@ -0,0 +1,18 @@ +//==-------- joint_matrix_half.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-fp16 +// REQUIRES: aspect-ext_intel_matrix +// REQUIRES: matrix-fp16 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp new file mode 100644 index 0000000000000..e0d2c5bd81769 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp @@ -0,0 +1,22 @@ +//==----- joint_matrix_int8_colmajorA_colmajorB.cpp - DPC++ joint_matrix---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// This tests support of col major layout for matrix B which does transpose and +// then VNNI transform. This is currently only available on AMX + +// XFAIL: gpu + +#include "../common.hpp" + +constexpr size_t TN = 16; + +#include "../joint_matrix_int8_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_rowmajorA_rowmajorB.cpp new file mode 100644 index 0000000000000..31bc890fdd197 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_rowmajorA_rowmajorB.cpp @@ -0,0 +1,19 @@ +//==----- joint_matrix_int8_rowmajorA_rowmajorB.cpp - DPC++ joint_matrix---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// Run these 2 tests on PVC only for now. Check can be updated to "gpu", +// when newer IGC is used in intel/llvm pre-checkin testing on Intel Arc +// RUN: %if arch-intel_gpu_pvc %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} +// RUN: %if arch-intel_gpu_pvc %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} + +#include "../common.hpp" +#include "../joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_opt_kernel_feature.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_opt_kernel_feature.cpp new file mode 100644 index 0000000000000..2ef42e3b499e9 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_opt_kernel_feature.cpp @@ -0,0 +1,18 @@ +//===---joint_matrix_opt_kernel_feature.cpp - DPC++ joint_matrix-----------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// Test checks that exception will be thrown in case matrix parameters are +// incompatible on the current device + +#include "../common.hpp" +#include "../joint_matrix_opt_kernel_feature_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp new file mode 100644 index 0000000000000..138f6738155f9 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp @@ -0,0 +1,20 @@ +//==-------- joint_matrix_out_bounds.cpp - DPC++ joint_matrix--------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// UNSUPPORTED: gpu-intel-dg2, cpu +// +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +#include "../common.hpp" + +constexpr size_t TN = 16; +constexpr size_t MATRIX_K = 1024 + 24; + +#include "../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp new file mode 100644 index 0000000000000..0c99e7a09063e --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp @@ -0,0 +1,17 @@ +//==-------- joint_matrix_prefetch.cpp - DPC++ joint_matrix----------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" + +constexpr size_t TN = 16; +#include "../joint_matrix_prefetch_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp new file mode 100644 index 0000000000000..6d213ba4ed870 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp @@ -0,0 +1,24 @@ +//==-------joint_matrix_rowmajorA_rowmajorB.cpp - DPC++ joint_matrix-------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix +// VNNI transform is not supported yet by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out %} +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} +// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} + +// XFAIL: cpu + +// This tests support of row major layout for matrix B which does automatic VNNI +// transform. This is currently only available on AMX and XMX of PVC + +#include "../common.hpp" +#include "../joint_matrix_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp new file mode 100644 index 0000000000000..d43e10ffed568 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp @@ -0,0 +1,16 @@ +//==-------- joint_matrix_ss_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_ss_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp new file mode 100644 index 0000000000000..1432abf115508 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp @@ -0,0 +1,16 @@ +//==-------- joint_matrix_su_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_su_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp new file mode 100644 index 0000000000000..3033631282058 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp @@ -0,0 +1,20 @@ +//==---------------- joint_matrix_tf32.cpp - DPC++ joint_matrix------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix-tf32 +// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" + +constexpr size_t TN = 16; + +#include "../joint_matrix_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_transposeC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_transposeC.cpp new file mode 100644 index 0000000000000..2a2ee8eccaf66 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_transposeC.cpp @@ -0,0 +1,14 @@ +//==----------- joint_matrix_transposeC.cpp - DPC++ joint_matrix-----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +#include "../common.hpp" +#include "../joint_matrix_transposeC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp new file mode 100644 index 0000000000000..43f3c12b70079 --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp @@ -0,0 +1,20 @@ +//==-------- joint_matrix_unaligned_k.cpp - DPC++ joint_matrix-------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: matrix + +// UNSUPPORTED: gpu-intel-dg2, cpu + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +#include "../common.hpp" + +constexpr size_t TN = 16; +static constexpr size_t MATRIX_K = 1024 + 14; + +#include "../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp new file mode 100644 index 0000000000000..93050559c1dfe --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp @@ -0,0 +1,16 @@ +//==-------- joint_matrix_us_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_us_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp new file mode 100644 index 0000000000000..2b988bdf92bcc --- /dev/null +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp @@ -0,0 +1,16 @@ +//==-------- joint_matrix_uu_int8.cpp - DPC++ joint_matrix------------ ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: aspect-ext_intel_matrix + +// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out +// RUN: %{run} %t.out + +// XFAIL: cpu + +#include "../common.hpp" +#include "../joint_matrix_uu_int8_impl.hpp"