Skip to content

Commit 5a188c7

Browse files
MrSidimsMrSidims
authored
[NFC] Remove MatrixUse::Unnecessary (#7335)
It was previously introduced for backwards compatibility with legacy (without Use matrix parameter) API. It is actually not needed. Signed-off-by: Sidorov, Dmitry <[email protected]>
1 parent 4f7787c commit 5a188c7

File tree

4 files changed

+167
-129
lines changed

4 files changed

+167
-129
lines changed

‎sycl/include/CL/__spirv/spirv_ops.hpp‎

Lines changed: 122 additions & 51 deletions
Original file line numberDiff line numberDiff line change
@@ -24,120 +24,191 @@
2424
#ifdef __SYCL_DEVICE_ONLY__
2525

2626
#if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
27-
#define JOINT_MATRIX_INTEL(T, R, C, L, S, U) \
28-
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U>
29-
#else
30-
#define JOINT_MATRIX_INTEL(T, R, C, L, S, U) \
31-
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S>
32-
#endif // SYCL_EXT_ONEAPI_MATRIX_VERSION
33-
34-
template <typename T, std::size_t R, std::size_t C,
35-
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
27+
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
3628
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
3729
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
38-
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T, R, C, L, S, U) *
30+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
3931
__spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
4032
__spv::MatrixLayout Layout = L,
4133
__spv::Scope::Flag Sc = S, int MemOperand = 0);
4234

43-
template <typename T, std::size_t R, std::size_t C,
44-
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
35+
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
4536
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
4637
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
4738
extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
48-
T *Ptr, JOINT_MATRIX_INTEL(T, R, C, L, S, U) *Object,
39+
T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *Object,
4940
std::size_t Stride, __spv::MatrixLayout Layout = L,
5041
__spv::Scope::Flag Sc = S, int MemOperand = 0);
5142

5243
template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
53-
__spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
54-
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
55-
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
44+
__spv::MatrixUse UA, __spv::MatrixUse UB, __spv::MatrixUse UC,
5645
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
5746
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
5847
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
5948
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
60-
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T2, M, N, LC, S, UC) *
49+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *
6150
__spirv_JointMatrixMadINTEL(
62-
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
63-
JOINT_MATRIX_INTEL(T1, K, N, LB, S, UB) *B,
64-
JOINT_MATRIX_INTEL(T2, M, N, LC, S, UC) *C,
51+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
52+
__spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S, UB> *B,
53+
__spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S, UC> *C,
6554
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
6655

6756
template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
68-
std::size_t N, __spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
69-
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
70-
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
57+
std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
58+
__spv::MatrixUse UC,
7159
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
7260
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
7361
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
7462
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
75-
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *
63+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
7664
__spirv_JointMatrixUUMadINTEL(
77-
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
78-
JOINT_MATRIX_INTEL(T2, K, N, LB, S, UB) *B,
79-
JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *C,
65+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
66+
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
67+
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
8068
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
8169

8270
template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
83-
std::size_t N, __spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
84-
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
85-
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
71+
std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
72+
__spv::MatrixUse UC,
8673
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
8774
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
8875
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
8976
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
90-
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *
77+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
9178
__spirv_JointMatrixUSMadINTEL(
92-
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
93-
JOINT_MATRIX_INTEL(T2, K, N, LB, S, UB) *B,
94-
JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *C,
79+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
80+
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
81+
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
9582
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
9683

9784
template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
98-
std::size_t N, __spv::MatrixUse UA = __spv::MatrixUse::Unnecessary,
99-
__spv::MatrixUse UB = __spv::MatrixUse::Unnecessary,
100-
__spv::MatrixUse UC = __spv::MatrixUse::Unnecessary,
85+
std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
86+
__spv::MatrixUse UC,
10187
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
10288
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
10389
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
10490
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
105-
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *
91+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
10692
__spirv_JointMatrixSUMadINTEL(
107-
JOINT_MATRIX_INTEL(T1, M, K, LA, S, UA) *A,
108-
JOINT_MATRIX_INTEL(T2, K, N, LB, S, UB) *B,
109-
JOINT_MATRIX_INTEL(T3, M, N, LC, S, UC) *C,
93+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
94+
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
95+
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
11096
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
11197

98+
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
99+
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
100+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
101+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
102+
__spirv_CompositeConstruct(const T v);
103+
104+
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
105+
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
106+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
107+
extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
108+
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *);
109+
110+
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
111+
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
112+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
113+
extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
114+
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *, size_t i);
115+
116+
template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
117+
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
118+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
119+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
120+
__spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *,
121+
T val, size_t i);
122+
#else
123+
template <typename T, std::size_t R, std::size_t C,
124+
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
125+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
126+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
127+
__spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
128+
__spv::MatrixLayout Layout = L,
129+
__spv::Scope::Flag Sc = S, int MemOperand = 0);
130+
112131
template <typename T, std::size_t R, std::size_t C,
113-
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
114132
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
115133
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
116-
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T, R, C, L, S, U) *
134+
extern SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
135+
T *Ptr, __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *Object,
136+
std::size_t Stride, __spv::MatrixLayout Layout = L,
137+
__spv::Scope::Flag Sc = S, int MemOperand = 0);
138+
139+
template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
140+
__spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
141+
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
142+
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
143+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
144+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *
145+
__spirv_JointMatrixMadINTEL(
146+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
147+
__spv::__spirv_JointMatrixINTEL<T1, K, N, LB, S> *B,
148+
__spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *C,
149+
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
150+
151+
template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
152+
std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
153+
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
154+
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
155+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
156+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T2, M, N, LC, S> *
157+
__spirv_JointMatrixUUMadINTEL(
158+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
159+
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
160+
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
161+
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
162+
163+
template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
164+
std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
165+
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
166+
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
167+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
168+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
169+
__spirv_JointMatrixUSMadINTEL(
170+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
171+
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
172+
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
173+
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
174+
175+
template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
176+
std::size_t N, __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
177+
__spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
178+
__spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
179+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
180+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *
181+
__spirv_JointMatrixSUMadINTEL(
182+
__spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S> *A,
183+
__spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S> *B,
184+
__spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S> *C,
185+
__spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
186+
187+
template <typename T, std::size_t R, std::size_t C,
188+
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
189+
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
190+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
117191
__spirv_CompositeConstruct(const T v);
118192

119193
template <typename T, std::size_t R, std::size_t C,
120-
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
121194
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
122195
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
123196
extern SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
124-
JOINT_MATRIX_INTEL(T, R, C, L, S, U) *);
197+
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *);
125198

126199
template <typename T, std::size_t R, std::size_t C,
127-
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
128200
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
129201
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
130202
extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(
131-
JOINT_MATRIX_INTEL(T, R, C, L, S, U) *, size_t i);
203+
__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *, size_t i);
132204

133205
template <typename T, std::size_t R, std::size_t C,
134-
__spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
135206
__spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
136207
__spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
137-
extern SYCL_EXTERNAL JOINT_MATRIX_INTEL(T, R, C, L, S, U) *
138-
__spirv_VectorInsertDynamic(JOINT_MATRIX_INTEL(T, R, C, L, S, U) *,
208+
extern SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *
209+
__spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S> *,
139210
T val, size_t i);
140-
#undef JOINT_MATRIX_INTEL
211+
#endif // SYCL_EXT_ONEAPI_MATRIX_VERSION
141212

142213
#ifndef __SPIRV_BUILTIN_DECLARATIONS__
143214
#error \

‎sycl/include/CL/__spirv/spirv_types.hpp‎

Lines changed: 2 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -116,17 +116,12 @@ enum class MatrixLayout : uint32_t {
116116
Unused = 4
117117
};
118118

119-
enum class MatrixUse : uint32_t {
120-
MatrixA = 0,
121-
MatrixB = 1,
122-
Accumulator = 2,
123-
Unnecessary = 3
124-
};
119+
enum class MatrixUse : uint32_t { MatrixA = 0, MatrixB = 1, Accumulator = 2 };
125120

126121
#if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
127122
template <typename T, std::size_t R, std::size_t C, MatrixLayout L,
128123
Scope::Flag S = Scope::Flag::Subgroup,
129-
MatrixUse U = MatrixUse::Unnecessary>
124+
MatrixUse U = MatrixUse::MatrixA>
130125
struct __spirv_JointMatrixINTEL;
131126
#else
132127
template <typename T, std::size_t R, std::size_t C, MatrixLayout L,

‎sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -39,9 +39,7 @@ SPV_MATRIX_LAYOUT_TRAITS(layout::packed_a, __spv::MatrixLayout::PackedA)
3939
SPV_MATRIX_LAYOUT_TRAITS(layout::packed_b, __spv::MatrixLayout::PackedB)
4040
SPV_MATRIX_LAYOUT_TRAITS(layout::unused, __spv::MatrixLayout::Unused)
4141

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

4644
template <use Use> struct spv_matrix_use_traits {
4745
static constexpr __spv::MatrixUse value = __spv::MatrixUse::MatrixA;
@@ -55,7 +53,6 @@ template <use Use> struct spv_matrix_use_traits {
5553
SPV_MATRIX_USE_TRAITS(use::a, __spv::MatrixUse::MatrixA)
5654
SPV_MATRIX_USE_TRAITS(use::b, __spv::MatrixUse::MatrixB)
5755
SPV_MATRIX_USE_TRAITS(use::accumulator, __spv::MatrixUse::Accumulator)
58-
SPV_MATRIX_USE_TRAITS(use::unnecessary, __spv::MatrixUse::Unnecessary)
5956

6057
template <typename G> struct spv_scope_traits {};
6158
template <> struct spv_scope_traits<sycl::sub_group> {

0 commit comments

Comments
 (0)