Skip to content

[NFC] Remove MatrixUse::Unnecessary #7335

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
steffenlarsen merged 3 commits into from
Nov 17, 2022
Merged

[NFC] Remove MatrixUse::Unnecessary #7335

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
0a44077
[NFC] Remove MatrixUse::Unnecessary
Nov 9, 2022
b54cbf9
Fix clang-format
Nov 9, 2022
fa05293
Apply comments
Nov 14, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
173 changes: 122 additions & 51 deletions sycl/include/CL/__spirv/spirv_ops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,120 +24,191 @@
#ifdef __SYCL_DEVICE_ONLY__

#if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
#define JOINT_MATRIX_INTEL(T, R, C, L, S, U) \
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U>
#else
#define JOINT_MATRIX_INTEL(T, R, C, L, S, U) \
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S>
#endif // SYCL_EXT_ONEAPI_MATRIX_VERSION

template <typename T, std::size_t R, std::size_t C,
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
Copy link
Contributor

@dkhaldi dkhaldi Nov 9, 2022

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

we are making the change to dynamic in this PR right?
If yes, the default for layout should be dynamic

Copy link
Contributor Author

@MrSidims MrSidims Nov 10, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This PR was meant by me to just remove MatrixUse::Unnecessary. Adding dynamic requires a change in layout enum guarding it with if MATRIX_VERSION - and AFAIK Bing is doing this work.

__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T, R, C, L, S, U) *
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
__spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
__spv::MatrixLayout Layout = L,
__spv::Scope::Flag Sc = S, int MemOperand = 0);

template <typename T, std::size_t R, std::size_t C,
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
T *Ptr, JOINT_MATRIX_INTEL(T, R, C, L, S, U) *Object,
T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *Object,
std::size_t Stride, __spv::MatrixLayout Layout = L,
__spv::Scope::Flag Sc = S, int MemOperand = 0);

template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
__spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UA, __spv::MatrixUse UB, __spv::MatrixUse UC,
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T2, M, N, LC, S, UC) *
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *
__spirv_JointMatrixMadINTEL(
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
JOINT_MATRIX_INTEL(T1, K, N, LB, S, UB) *B,
JOINT_MATRIX_INTEL(T2, M, N, LC, S, UC) *C,
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
__spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S, UB> *B,
__spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
std::size_t N, __spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
__spv::MatrixUse UC,
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
__spirv_JointMatrixUUMadINTEL(
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
JOINT_MATRIX_INTEL(T2, K, N, LB, S, UB) *B,
JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *C,
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
std::size_t N, __spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
__spv::MatrixUse UC,
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
__spirv_JointMatrixUSMadINTEL(
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
JOINT_MATRIX_INTEL(T2, K, N, LB, S, UB) *B,
JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *C,
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
std::size_t N, __spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
__spv::MatrixUse UC,
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
__spirv_JointMatrixSUMadINTEL(
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
JOINT_MATRIX_INTEL(T2, K, N, LB, S, UB) *B,
JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *C,
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
__spirv_CompositeConstruct(const T v);

template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *);

template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *, size_t i);

template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
__spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *,
T val, size_t i);
#else
template <typename T, std::size_t R, std::size_t C,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
__spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
__spv::MatrixLayout Layout = L,
__spv::Scope::Flag Sc = S, int MemOperand = 0);

template <typename T, std::size_t R, std::size_t C,
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T, R, C, L, S, U) *
extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *Object,
std::size_t Stride, __spv::MatrixLayout Layout = L,
__spv::Scope::Flag Sc = S, int MemOperand = 0);

template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *
__spirv_JointMatrixMadINTEL(
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
__spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S> *B,
__spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *
__spirv_JointMatrixUUMadINTEL(
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
__spirv_JointMatrixUSMadINTEL(
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
__spirv_JointMatrixSUMadINTEL(
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);

template <typename T, std::size_t R, std::size_t C,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
__spirv_CompositeConstruct(const T v);

template <typename T, std::size_t R, std::size_t C,
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
JOINT_MATRIX_INTEL(T, R, C, L, S, U) *);
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *);

template <typename T, std::size_t R, std::size_t C,
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
JOINT_MATRIX_INTEL(T, R, C, L, S, U) *, size_t i);
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *, size_t i);

template <typename T, std::size_t R, std::size_t C,
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T, R, C, L, S, U) *
__spirv_VectorInsertDynamic(JOINT_MATRIX_INTEL(T, R, C, L, S, U) *,
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
__spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *,
T val, size_t i);
#undef JOINT_MATRIX_INTEL
#endif // SYCL_EXT_ONEAPI_MATRIX_VERSION

#ifndef __SPIRV_BUILTIN_DECLARATIONS__
#error \
Expand Down
9 changes: 2 additions & 7 deletions sycl/include/CL/__spirv/spirv_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -116,17 +116,12 @@ enum class MatrixLayout : uint32_t {
Unused = 4
};

enum class MatrixUse : uint32_t {
MatrixA = 0,
MatrixB = 1,
Accumulator = 2,
Unnecessary = 3
};
enum class MatrixUse : uint32_t { MatrixA = 0, MatrixB = 1, Accumulator = 2 };

#if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
template <typename T, std::size_t R, std::size_t C, MatrixLayout L,
Scope::Flag S = Scope::Flag::Subgroup,
MatrixUse U = MatrixUse::Unnecessary>
MatrixUse U = MatrixUse::MatrixA>
struct __spirv_JointMatrixINTEL;
#else
template <typename T, std::size_t R, std::size_t C, MatrixLayout L,
Expand Down
5 changes: 1 addition & 4 deletions sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,9 +39,7 @@ SPV_MATRIX_LAYOUT_TRAITS(layout::packed_a, __spv::MatrixLayout::PackedA)
SPV_MATRIX_LAYOUT_TRAITS(layout::packed_b, __spv::MatrixLayout::PackedB)
Copy link
Contributor

@dkhaldi dkhaldi Nov 9, 2022

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

when we should make the change: unused --> dynamic
packed_a and packed_b --> packed

SPV_MATRIX_LAYOUT_TRAITS(layout::unused, __spv::MatrixLayout::Unused)

// unnecessary was introduced for backward compatibility.
// Once the use implementation is stable, "unnecessary" value will be omitted
enum class use { a, b, accumulator, unnecessary };
enum class use { a, b, accumulator };

template <use Use> struct spv_matrix_use_traits {
static constexpr __spv::MatrixUse value = __spv::MatrixUse::MatrixA;
Expand All @@ -55,7 +53,6 @@ template <use Use> struct spv_matrix_use_traits {
SPV_MATRIX_USE_TRAITS(use::a, __spv::MatrixUse::MatrixA)
SPV_MATRIX_USE_TRAITS(use::b, __spv::MatrixUse::MatrixB)
SPV_MATRIX_USE_TRAITS(use::accumulator, __spv::MatrixUse::Accumulator)
SPV_MATRIX_USE_TRAITS(use::unnecessary, __spv::MatrixUse::Unnecessary)

template <typename G> struct spv_scope_traits {};
template <> struct spv_scope_traits<sycl::sub_group> {
Expand Down
Loading