DPC++ Runtime
Runtime libraries for oneAPI DPC++
static-query.hpp
Go to the documentation of this file.
1 //===-------------- static-query.hpp - SYCL matrix ------------*- C++ -*---===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 // ===--------------------------------------------------------------------=== //
8 // This file implements the static query interface for the joint_matrix
9 // experimental extension. AMX, DPAS and different other TPUs support different
10 // logical sizes and types. The query interface is used to validate user code
11 // and inform them about supported types, sizes, scope, and layouts by the
12 // current implementation. Note that this query interface is a compile-time
13 // query, so there will be no runtime errors. The query interface provides
14 // three functionalities:
15 // 1- At compile time, inform the user whether a specific
16 // combination is valid or not.
17 // 2- Construct the matrices using a default shape
18 // if user does not provide a combination
19 // 3- General query interface for sizes, types,
20 // static/dynamic, scope. This is needed to void padding by the user,
21 // for tuning, and efficient code generation if used by a library.
22 
23 #pragma once
24 
26 namespace sycl {
27 namespace ext {
28 namespace oneapi {
29 namespace experimental::matrix {
30 
31 enum class tpu {
32  dpas,
33  amx,
34 };
35 enum class matrix_type {
36  bf8,
37  bf16,
38  fp16,
39  fp19, // tfloat32
40  fp32,
41  fp64,
42  sint2,
43  sint4,
44  sint8,
45  sint16,
46  sint32,
47  sint64,
48  uint2,
49  uint4,
50  uint8,
51  uint16,
52  uint32,
53  uint64
54 };
55 
56 enum class scope_t { sub_group, work_group };
57 
58 template <tpu u, typename Ta = void, typename Tb = void, typename Tc = void,
59  int M = 0, int N = 0, int K = 0, typename Enabled = void>
60 struct tpu_params;
61 
62 #if __cplusplus >= 201703L
63 template <typename Ta, typename Tb, typename Tc>
64 constexpr bool is_combination_valid_amx(int M, int N, int K) {
65  // is_same_v is a C++17 feature
66  if ((std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, int8_t> &&
67  std::is_same_v<Tc, int> && M <= 16 && N <= 16 && K <= 64) ||
68  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, uint8_t> &&
69  std::is_same_v<Tc, int> && M <= 16 && N <= 16 && K <= 64) ||
70  (std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, uint8_t> &&
71  std::is_same_v<Tc, int> && M <= 16 && N <= 16 && K <= 64) ||
72  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, int8_t> &&
73  std::is_same_v<Tc, int> && M <= 16 && N <= 16 && K <= 64) ||
74  // bf16
75  (std::is_same_v<Ta, unsigned short> &&
76  std::is_same_v<Tb, unsigned short> && std::is_same_v<Tc, float> &&
77  M <= 16 && N <= 16 && K <= 32))
78  return true;
79  else
80  return false;
81 }
82 
83 template <typename Ta, typename Tb, typename Tc>
84 constexpr bool are_types_valid_amx() {
85  if ((std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, int8_t> &&
86  std::is_same_v<Tc, int>) ||
87  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, uint8_t> &&
88  std::is_same_v<Tc, int>) ||
89  (std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, uint8_t> &&
90  std::is_same_v<Tc, int>) ||
91  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, int8_t> &&
92  std::is_same_v<Tc, int>) ||
93  (std::is_same_v<Ta, unsigned short> &&
94  std::is_same_v<Tb, unsigned short> && std::is_same_v<Tc, float>))
95  return true;
96  else
97  return false;
98 }
99 #endif
100 
101 // General query:
102 // types are not given, no default sizes and no implicit matrix construction
103 template <int M, int N, int K>
104 struct tpu_params<tpu::amx, void, void, void, M, N, K> {
105  static constexpr std::size_t defaultM = -1; // depends on the type
106  static constexpr std::size_t defaultN = -1;
107  static constexpr std::size_t defaultK = -1;
108 
109  bool dynamic_p = false; // should be true in future implementations because
110  // AMX hardware supports dynamic sizes
111  uint32_t numtiles = 8;
112  scope_t scope = scope_t::sub_group;
113  struct combination {
114  uint32_t max_msize;
115  uint32_t max_nsize;
116  uint32_t max_ksize;
120  uint32_t msize;
121  uint32_t nsize;
122  uint32_t ksize;
123  };
124  using mt = matrix_type;
125  static constexpr combination combinations[] = {
126  {16, 16, 64, mt::sint8, mt::sint8, mt::sint32},
127  {16, 16, 64, mt::sint8, mt::uint8, mt::sint32},
128  {16, 16, 64, mt::uint8, mt::sint8, mt::sint32},
129  {16, 16, 64, mt::uint8, mt::uint8, mt::sint32},
130  {16, 16, 32, mt::bf16, mt::bf16, mt::fp32}};
131  static constexpr int num_combinations =
132  sizeof(combinations) / sizeof(combination);
133 };
134 
135 #if __cplusplus >= 201703L
136 // Sizes-only query
137 // Specialization for when only types are given, need to query only sizes
138 template <typename Ta, typename Tb, typename Tc>
139 struct tpu_params<tpu::amx, Ta, Tb, Tc, 0, 0, 0,
140  typename std::enable_if<(!std::is_same_v<Ta, void> &&
141  !std::is_same_v<Tb, void> &&
142  !std::is_same_v<Tc, void>)>::type> {
143  static_assert((are_types_valid_amx<Ta, Tb, Tc>()),
144  "Invalid types for AMX, supported types are int8_t, uint8_t, "
145  "and bf16 (Note that unsigned short should be used in the"
146  "DPC++ code to implement bf16) ");
147 
148  // construct the matrices using the default sizes
149  static constexpr std::size_t defaultM = 16;
150  static constexpr std::size_t defaultN = 16;
151  static constexpr std::size_t defaultK = ((sizeof(Ta) == 1) ? 64 : 32);
152 
153  template <typename Group>
154  using joint_matrix_a =
156  template <typename Group>
157  using joint_matrix_b =
159  template <typename Group>
160  using joint_matrix_c =
162 
163  bool dynamic_p = false; // should be true in future implementations because
164  // AMX hardware supports dynamic sizes
165  uint32_t numtiles = 8;
166  scope_t scope = scope_t::sub_group;
167  struct combination {
168  uint32_t max_msize;
169  uint32_t max_nsize;
170  uint32_t max_ksize;
171  matrix_type atype;
172  matrix_type btype;
173  matrix_type ctype;
174  uint32_t msize;
175  uint32_t nsize;
176  uint32_t ksize;
177  };
178  static constexpr combination combinations[] = {
179  {16, 16, (sizeof(Ta) == 1) ? 64 : 32}};
180  static constexpr int num_combinations =
181  sizeof(combinations) / sizeof(combination);
182 };
183 
184 // Valid or not:
185 // Specialization when both types and sizes are given
186 template <typename Ta, typename Tb, typename Tc, int M, int N, int K>
187 struct tpu_params<
188  tpu::amx, Ta, Tb, Tc, M, N, K,
189  typename std::enable_if<(
190  !std::is_same_v<Ta, void> && !std::is_same_v<Tb, void> &&
191  !std::is_same_v<Tc, void> && M != 0 && N != 0 && K != 0)>::type> {
192  // Validate that parameters are supported
193  static_assert(
194  (M == 0 && N == 0 && K == 0) ||
195  (is_combination_valid_amx<Ta, Tb, Tc>(M, N, K)),
196  "Invalid parameters for AMX, query valid types and maximum sizes "
197  "using: tpu_params<tpu::amx> myparams; and then check out "
198  "myparams.combinations array");
199 
200  // if combination is valid, construct the matrices
201 
202  static constexpr std::size_t defaultM = (M != 0) ? M : 16;
203  static constexpr std::size_t defaultN = (N != 0) ? N : 16;
204  static constexpr std::size_t defaultK =
205  (K != 0) ? K : ((sizeof(Ta) == 1) ? 64 : 32);
206 
207  template <typename Group>
208  using joint_matrix_a =
209  joint_matrix<Ta, defaultM, defaultK, matrix_layout::row_major, Group>;
210  template <typename Group>
211  using joint_matrix_b =
212  joint_matrix<Tb, defaultK, defaultN, matrix_layout::packed_b, Group>;
213  template <typename Group>
214  using joint_matrix_c =
215  joint_matrix<Tc, defaultM, defaultN, matrix_layout::row_major, Group>;
216 
217  bool dynamic_p = false; // should be true in future implementations
218  // because AMX hardware supports dynamic sizes
219  uint32_t numtiles = 8;
220  scope_t scope = scope_t::sub_group;
221 };
222 
223 // DPAS case
224 // The DPAS implementation supports the logical capability support of the HW
225 // So in this case, M, N, K sizes returned by the query represent the logical
226 // capabilities of the DPAS hardware.
227 
228 template <typename Ta, typename Tb, typename Tc>
229 constexpr bool is_combination_valid_dpas(int M, int N, int K) {
230  if ((std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, int8_t> &&
231  std::is_same_v<Tc, int> && (M == 1 || M == 2 || M == 4 || M == 8) &&
232  N == 8 && K == 32) ||
233  (std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, uint8_t> &&
234  std::is_same_v<Tc, int> && (M == 1 || M == 2 || M == 4 || M == 8) &&
235  N == 8 && K == 32) ||
236  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, int8_t> &&
237  std::is_same_v<Tc, int> && (M == 1 || M == 2 || M == 4 || M == 8) &&
238  N == 8 && K == 32) ||
239  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, uint8_t> &&
240  std::is_same_v<Tc, int> && (M == 1 || M == 2 || M == 4 || M == 8) &&
241  N == 8 && K == 32) ||
242  (std::is_same_v<Ta, half> && std::is_same_v<Tb, half> &&
243  std::is_same_v<Tc, float> && (M == 1 || M == 2 || M == 4 || M == 8) &&
244  N == 8 && K == 16) ||
245  (std::is_same_v<Ta, unsigned short> &&
246  std::is_same_v<Tb, unsigned short> && std::is_same_v<Tc, float> &&
247  (M == 1 || M == 2 || M == 4 || M == 8) && N == 8 && K == 16))
248  return true;
249  else
250  return false;
251 }
252 
253 template <typename Ta, typename Tb, typename Tc>
254 constexpr bool are_types_valid_dpas() {
255  if ((std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, int8_t> &&
256  std::is_same_v<Tc, int>) ||
257  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, int8_t> &&
258  std::is_same_v<Tc, int>) ||
259  (std::is_same_v<Ta, int8_t> && std::is_same_v<Tb, uint8_t> &&
260  std::is_same_v<Tc, int>) ||
261  (std::is_same_v<Ta, uint8_t> && std::is_same_v<Tb, uint8_t> &&
262  std::is_same_v<Tc, int>) ||
263  (std::is_same_v<Ta, half> && std::is_same_v<Tb, half> &&
264  std::is_same_v<Tc, float>) ||
265  (std::is_same_v<Ta, unsigned short> &&
266  std::is_same_v<Tb, unsigned short> && std::is_same_v<Tc, float>))
267  return true;
268  else
269  return false;
270 }
271 #endif
272 
273 // General Query
274 // specialization for when types are not given --> no default values
275 template <int M, int N, int K>
276 struct tpu_params<tpu::dpas, void, void, void, M, N, K> {
277  static constexpr std::size_t defaultM = -1; // depends on the type
278  static constexpr std::size_t defaultN = -1;
279  static constexpr std::size_t defaultK = -1;
280 
281  bool dynamic_p = false; // no dynamic allocation on the GPU
282  uint32_t numtiles = -1; // does not apply for DPAS
283  scope_t scope = scope_t::sub_group;
284 
285  struct combination {
286  uint32_t max_msize;
287  uint32_t max_nsize;
288  uint32_t max_ksize;
292  uint32_t msize;
293  uint32_t nsize;
294  uint32_t ksize;
295  };
296  using mt = matrix_type;
297  static constexpr combination combinations[] = {
298  {0, 0, 0, mt::sint8, mt::sint8, mt::sint32, 1, 8, 32},
299  {0, 0, 0, mt::sint8, mt::sint8, mt::sint32, 2, 8, 32},
300  {0, 0, 0, mt::sint8, mt::sint8, mt::sint32, 4, 8, 32},
301  {0, 0, 0, mt::sint8, mt::sint8, mt::sint32, 8, 8, 32},
302  {0, 0, 0, mt::sint8, mt::uint8, mt::sint32, 1, 8, 32},
303  {0, 0, 0, mt::sint8, mt::uint8, mt::sint32, 2, 8, 32},
304  {0, 0, 0, mt::sint8, mt::uint8, mt::sint32, 4, 8, 32},
305  {0, 0, 0, mt::sint8, mt::uint8, mt::sint32, 8, 8, 32},
306  {0, 0, 0, mt::uint8, mt::sint8, mt::sint32, 1, 8, 32},
307  {0, 0, 0, mt::uint8, mt::sint8, mt::sint32, 2, 8, 32},
308  {0, 0, 0, mt::uint8, mt::sint8, mt::sint32, 4, 8, 32},
309  {0, 0, 0, mt::uint8, mt::sint8, mt::sint32, 8, 8, 32},
310  {0, 0, 0, mt::uint8, mt::uint8, mt::sint32, 1, 8, 32},
311  {0, 0, 0, mt::uint8, mt::uint8, mt::sint32, 2, 8, 32},
312  {0, 0, 0, mt::uint8, mt::uint8, mt::sint32, 4, 8, 32},
313  {0, 0, 0, mt::uint8, mt::uint8, mt::sint32, 8, 8, 32},
314  {0, 0, 0, mt::fp16, mt::fp16, mt::fp32, 1, 8, 16},
315  {0, 0, 0, mt::fp16, mt::fp16, mt::fp32, 2, 8, 16},
316  {0, 0, 0, mt::fp16, mt::fp16, mt::fp32, 4, 8, 16},
317  {0, 0, 0, mt::fp16, mt::fp16, mt::fp32, 8, 8, 16},
318  {0, 0, 0, mt::bf16, mt::bf16, mt::fp32, 1, 8, 16},
319  {0, 0, 0, mt::bf16, mt::bf16, mt::fp32, 2, 8, 16},
320  {0, 0, 0, mt::bf16, mt::bf16, mt::fp32, 4, 8, 16},
321  {0, 0, 0, mt::bf16, mt::bf16, mt::fp32, 8, 8, 16},
322  };
323  static constexpr int num_combinations =
324  sizeof(combinations) / sizeof(combination);
325 };
326 
327 // Sizes-only query:
328 // Specialization for when only types are given, need to query only sizes
329 
330 #if __cplusplus >= 201703L
331 template <typename Ta, typename Tb, typename Tc>
332 struct tpu_params<tpu::dpas, Ta, Tb, Tc, 0, 0, 0,
333  typename std::enable_if<(!std::is_same_v<Ta, void> &&
334  !std::is_same_v<Tb, void> &&
335  !std::is_same_v<Tc, void>)>::type> {
336  static_assert((are_types_valid_dpas<Ta, Tb, Tc>()),
337  "Invalid types for DPAS, supported types are int8_t, uint8_t, "
338  "half, and bf16 (Note that unsigned short should be used in the"
339  "DPC++ code to implement bf16)");
340 
341  // construct the matrices using the default sizes
342 
343  static constexpr std::size_t defaultM = 8;
344  static constexpr std::size_t defaultN = 8;
345  static constexpr std::size_t defaultK = ((sizeof(Ta) == 1) ? 32 : 16);
346 
347  template <typename Group>
348  using joint_matrix_a =
350  template <typename Group>
351  using joint_matrix_b =
353  template <typename Group>
354  using joint_matrix_c =
356 
357  bool dynamic_p = false; // no dynamic allocation on the GPU
358  uint32_t numtiles = -1; // does not apply for DPAS
359  scope_t scope = scope_t::sub_group;
360  struct combination {
361  uint32_t max_msize;
362  uint32_t max_nsize;
363  uint32_t max_ksize;
364  matrix_type atype;
365  matrix_type btype;
366  matrix_type ctype;
367  uint32_t msize;
368  uint32_t nsize;
369  uint32_t ksize;
370  };
371  using mt = matrix_type;
372  static constexpr combination combinations[] = {
373  // The types used in the initialization below are fake and not used. In
374  // this case, users already chose the types, they are only looking for the
375  // sizes
376  {0, 0, 0, mt::bf8, mt::bf8, mt::bf8, 1, 8, (sizeof(Ta) == 1) ? 32 : 16},
377  {0, 0, 0, mt::bf8, mt::bf8, mt::bf8, 2, 8, (sizeof(Ta) == 1) ? 32 : 16},
378  {0, 0, 0, mt::bf8, mt::bf8, mt::bf8, 4, 8, (sizeof(Ta) == 1) ? 32 : 16},
379  {0, 0, 0, mt::bf8, mt::bf8, mt::bf8, 8, 8, (sizeof(Ta) == 1) ? 32 : 16},
380  };
381  static constexpr int num_combinations =
382  sizeof(combinations) / sizeof(combination);
383 };
384 
385 // Valid or not:
386 // Specialization when both types and sizes are given
387 template <typename Ta, typename Tb, typename Tc, int M, int N, int K>
388 struct tpu_params<
389  tpu::dpas, Ta, Tb, Tc, M, N, K,
390  typename std::enable_if<((!std::is_same_v<Ta, void> && M != 0))>::type> {
391  // Validate that parameters are supported
392  static_assert((M == 0 && N == 0 && K == 0) ||
393  (is_combination_valid_dpas<Ta, Tb, Tc>(M, N, K)),
394  "Invalid parameters for DPAS, query valid combinations "
395  "using: tpu_params<tpu::dpas> myparams; and then check out "
396  "myparams.combinations array");
397 
398  // if combination is valid, construct the matrices
399  static constexpr std::size_t defaultM = (M != 0) ? M : 8;
400  static constexpr std::size_t defaultN = (N != 0) ? N : 8;
401  static constexpr std::size_t defaultK =
402  (K != 0) ? K : ((sizeof(Ta) == 1) ? 32 : 16);
403 
404  template <typename Group>
405  using joint_matrix_a =
406  joint_matrix<Ta, defaultM, defaultK, matrix_layout::row_major, Group>;
407  template <typename Group>
408  using joint_matrix_b =
409  joint_matrix<Tb, defaultK, defaultN, matrix_layout::packed_b, Group>;
410  template <typename Group>
411  using joint_matrix_c =
412  joint_matrix<Tc, defaultM, defaultN, matrix_layout::row_major, Group>;
413 
414  bool dynamic_p = false; // no dynamic allocation on the GPU
415  uint32_t numtiles = -1; // does not apply for DPAS
416  scope_t scope = scope_t::sub_group;
417 };
418 #endif
419 } // namespace experimental::matrix
420 } // namespace oneapi
421 } // namespace ext
422 } // namespace sycl
423 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::nsize
uint32_t nsize
Definition: static-query.hpp:121
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::max_msize
uint32_t max_msize
Definition: static-query.hpp:286
cl::sycl::ext::oneapi::experimental::matrix::tpu_params
Definition: static-query.hpp:60
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::sint4
@ sint4
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::ctype
matrix_type ctype
Definition: static-query.hpp:291
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::max_ksize
uint32_t max_ksize
Definition: static-query.hpp:116
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::sint8
@ sint8
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::ctype
matrix_type ctype
Definition: static-query.hpp:119
sycl
Definition: invoke_simd.hpp:68
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::sint64
@ sint64
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::sint32
@ sint32
cl::sycl::ext::oneapi::experimental::matrix::tpu
tpu
Definition: static-query.hpp:31
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::sint2
@ sint2
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::uint4
@ uint4
cl::sycl::image_channel_type::fp32
@ fp32
cl::sycl::aspect::fp64
@ fp64
cl::sycl::ext::intel::experimental::esimd::dpas
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpas(sycl::ext::intel::esimd::simd< T0, N > src0, sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, Sat sat={})
Definition: math.hpp:1742
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::bf16
@ bf16
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::btype
matrix_type btype
Definition: static-query.hpp:290
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::max_nsize
uint32_t max_nsize
Definition: static-query.hpp:287
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::btype
matrix_type btype
Definition: static-query.hpp:118
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::msize
uint32_t msize
Definition: static-query.hpp:292
cl::sycl::aspect::fp16
@ fp16
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::ksize
uint32_t ksize
Definition: static-query.hpp:122
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::max_ksize
uint32_t max_ksize
Definition: static-query.hpp:288
cl::sycl::ext::oneapi::sub_group
Definition: sub_group.hpp:108
cl::sycl::ext::oneapi::experimental::matrix::matrix_type
matrix_type
Definition: static-query.hpp:35
cl::sycl::ext::oneapi::experimental::matrix::tpu::amx
@ amx
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::msize
uint32_t msize
Definition: static-query.hpp:120
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::uint32
@ uint32
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::fp19
@ fp19
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::atype
matrix_type atype
Definition: static-query.hpp:289
cl::sycl::ext::oneapi::experimental::matrix::joint_matrix
Definition: matrix-jit.hpp:56
cl::sycl::memory_scope::work_group
@ work_group
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::ksize
uint32_t ksize
Definition: static-query.hpp:294
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::uint64
@ uint64
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::uint16
@ uint16
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::max_msize
uint32_t max_msize
Definition: static-query.hpp:114
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::uint8
@ uint8
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::bf8
@ bf8
cl::sycl::ext::oneapi::experimental::matrix::scope_t
scope_t
Definition: static-query.hpp:56
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::uint2
@ uint2
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::dpas, void, void, void, M, N, K >::combination::nsize
uint32_t nsize
Definition: static-query.hpp:293
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::atype
matrix_type atype
Definition: static-query.hpp:117
cl::sycl::ext::oneapi::experimental::matrix::tpu_params< tpu::amx, void, void, void, M, N, K >::combination::max_nsize
uint32_t max_nsize
Definition: static-query.hpp:115
cl::sycl::ext::oneapi::experimental::matrix::matrix_type::sint16
@ sint16
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12