/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp Source File
cshuffle_epilogue.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include "ck_tile/core.hpp"
9 
10 namespace ck_tile {
11 
12 template <typename ADataType_,
13  typename BDataType_,
14  typename DsDataType_,
15  typename AccDataType_,
16  typename ODataType_,
17  typename DsLayout_,
18  typename ELayout_,
19  typename CDElementwise_,
20  index_t kM_,
21  index_t kN_,
22  index_t MWave_,
23  index_t NWave_,
24  index_t MPerXdl_,
25  index_t NPerXdl_,
26  index_t KPerXdl_,
27  bool isCTransposed_,
28  memory_operation_enum MemoryOperation_,
29  index_t kNumWaveGroups_ = 1,
30  bool FixedVectorSize_ = false,
31  index_t VectorSizeC_ = 1>
33 {
42  static constexpr index_t kBlockSize = MWave_ * NWave_ * get_warp_size();
43  static constexpr index_t kMPerBlock = kM_;
44  static constexpr index_t kNPerBlock = kN_;
45  static constexpr index_t MWave = MWave_;
46  static constexpr index_t NWave = NWave_;
47  static constexpr index_t MPerXdl = MPerXdl_;
48  static constexpr index_t NPerXdl = NPerXdl_;
49  static constexpr index_t KPerXdl = KPerXdl_;
50  static constexpr index_t isCTransposed = isCTransposed_;
51  static constexpr memory_operation_enum MemoryOperation = MemoryOperation_;
52  static constexpr bool FixedVectorSize = FixedVectorSize_;
53  static constexpr index_t VectorSizeC = VectorSizeC_;
54  static constexpr index_t kNumWaveGroups = kNumWaveGroups_;
55  static constexpr index_t NumDTensor = DsDataType::size();
56 
57  static_assert(NumDTensor == DsLayout::size(),
58  "The size of DsDataType and DsLayout should be the same");
59 };
60 
61 template <typename Problem_, typename Policy_ = void>
63 {
71  using ATypeToUse =
72  std::conditional_t<std::is_same_v<ADataType, pk_int4_t>, BDataType, ADataType>;
73  // Used for weight-only quantization kernel, B would be dequantized to the same data type as A
74  using BTypeToUse =
75  std::conditional_t<std::is_same_v<BDataType, pk_int4_t>, ADataType, BDataType>;
78  static constexpr memory_operation_enum MemoryOperation = Problem::MemoryOperation;
79  static constexpr index_t kBlockSize = Problem::kBlockSize;
80  static constexpr index_t kMPerBlock = Problem::kMPerBlock;
81  static constexpr index_t kNPerBlock = Problem::kNPerBlock;
82  static constexpr index_t MWave = Problem::MWave;
83  static constexpr index_t NWave = Problem::NWave;
84  static constexpr index_t MPerXdl = Problem::MPerXdl;
85  static constexpr index_t NPerXdl = Problem::NPerXdl;
86  static constexpr index_t KPerXdl = Problem::KPerXdl;
87  static constexpr index_t isCTransposed = Problem::isCTransposed;
88  static constexpr bool FixedVectorSize = Problem::FixedVectorSize;
89  static constexpr index_t VectorSizeC = Problem::VectorSizeC;
90  static constexpr index_t MPerIteration = MPerXdl * MWave;
91  static constexpr index_t NPerIteration = NPerXdl * NWave;
92  static constexpr index_t NumDTensor = Problem::NumDTensor;
93 
94  static_assert(NumDTensor == DsLayout::size(),
95  "The size of DsDataType and DsLayout should be the same");
107  {
108  if constexpr(FixedVectorSize)
109  {
110  return VectorSizeC;
111  }
112  constexpr index_t max_vector_size = 16;
113  if constexpr(std::is_same_v<ELayout, tensor_layout::gemm::RowMajor>)
114  {
115  return std::min(static_cast<int>(NPerIteration),
116  static_cast<int>(max_vector_size / sizeof(ODataType)));
117  }
118  else if constexpr(std::is_same_v<ELayout, tensor_layout::gemm::ColumnMajor>)
119  {
120  return std::min(static_cast<int>(MPerIteration),
121  static_cast<int>(max_vector_size / sizeof(ODataType)));
122  }
123  else
124  {
125  static_assert(false, "Unsupported ELayout!");
126  }
127  }
128 
134  template <index_t I>
136  {
137  constexpr index_t max_vector_size = 16;
138  using DiDataType = remove_cvref_t<std::tuple_element_t<index.value, DsDataType>>;
139  using DiLayout = remove_cvref_t<std::tuple_element_t<index.value, DsLayout>>;
140  if constexpr(std::is_same_v<DiLayout, tensor_layout::gemm::RowMajor>)
141  {
142  return std::min(static_cast<int>(NPerIteration),
143  static_cast<int>(max_vector_size / sizeof(DiDataType)));
144  }
145  else if constexpr(std::is_same_v<DiLayout, tensor_layout::gemm::ColumnMajor>)
146  {
147  return std::min(static_cast<int>(MPerIteration),
148  static_cast<int>(max_vector_size / sizeof(DiDataType)));
149  }
150  else
151  {
152  static_assert(false, "Unsupported DLayout!");
153  }
154  return max_vector_size / sizeof(DiDataType);
155  }
164  static constexpr auto shuffle_tile_tuple = [] {
165  constexpr index_t elem_per_thread = MPerXdl * NPerXdl / get_warp_size();
166  if constexpr(elem_per_thread >= GetVectorSizeC())
167  {
168  return std::make_tuple(1, 1);
169  }
170  else
171  {
172  constexpr index_t num_xdl_shuffles = GetVectorSizeC() / elem_per_thread;
173  if constexpr(std::is_same_v<ELayout, tensor_layout::gemm::RowMajor>)
174  {
175  static_assert((kMPerBlock % (MPerXdl * MWave) == 0) &&
176  (kMPerBlock % num_xdl_shuffles == 0),
177  "kMPerBlock must be divisible by MPerXdl*MWave and "
178  "num_xdl_shuffles for CShuffleEpilogue");
179  return std::make_tuple(min(num_xdl_shuffles, kMPerBlock / (MPerXdl * MWave)), 1);
180  }
181  else
182  {
183  static_assert((kNPerBlock % (NPerXdl * NWave) == 0) &&
184  (kNPerBlock % num_xdl_shuffles == 0),
185  "kNPerBlock must be divisible by NPerXdl*NWave and "
186  "num_xdl_shuffles for CShuffleEpilogue");
187  return std::make_tuple(1, min(num_xdl_shuffles, kNPerBlock / (NPerXdl * NWave)));
188  }
189  }
190  }();
191  static constexpr index_t NumMXdlPerWavePerShuffle = std::get<0>(shuffle_tile_tuple);
192  static constexpr index_t NumNXdlPerWavePerShuffle = std::get<1>(shuffle_tile_tuple);
193 
194  static constexpr auto MNPerIterationShuffle = [] {
195  constexpr index_t m_val = MPerXdl * MWave * NumMXdlPerWavePerShuffle;
196  constexpr index_t n_val = NPerXdl * NWave * NumNXdlPerWavePerShuffle;
197  if constexpr(kMPerBlock % m_val != 0 || kNPerBlock % n_val != 0)
199  else
200  return std::make_tuple(m_val, n_val);
201  }();
202  static constexpr index_t MPerIterationShuffle = std::get<0>(MNPerIterationShuffle);
203  static constexpr index_t NPerIterationShuffle = std::get<1>(MNPerIterationShuffle);
204 
206  BTypeToUse,
207  AccDataType,
208  MPerXdl,
209  NPerXdl,
210  KPerXdl,
211  isCTransposed>;
212 
213  using CWarpDstr = typename WG::CWarpDstr;
214  using CWarpTensor = typename WG::CWarpTensor;
215 
216  template <typename Problem>
218  {
219  // N is contiguous dimension
220  if constexpr(std::is_same_v<ELayout, tensor_layout::gemm::RowMajor>)
221  {
225  }
226  // M is contiguous dimension
227  else if constexpr(std::is_same_v<ELayout, tensor_layout::gemm::ColumnMajor>)
228  {
232  }
233  else
234  {
235  static_assert(false, "Unsupported ELayout!");
236  }
237  }
238 
240  {
241  constexpr auto block_outer_dstr_encoding =
248  sequence<0, 0>>{};
249  constexpr auto block_dstr_encoding = detail::make_embed_tile_distribution_encoding(
250  block_outer_dstr_encoding, typename CWarpDstr::DstrEncode{});
251 
252  return block_dstr_encoding;
253  }
254 
256  {
258  }
259 
260  template <typename ODramWindow, typename OAccTile, typename DsDramWindows>
261  CK_TILE_DEVICE auto operator()(ODramWindow& out_dram_window,
262  const OAccTile& o_acc_tile,
263  const DsDramWindows& ds_dram_windows,
264  void* p_smem)
265  {
266  constexpr auto LdsTileDistr = make_static_tile_distribution(MakeLdsDistributionEncode());
267 
268  auto lds_tile = make_static_distributed_tensor<AccDataType>(LdsTileDistr);
269 
270  constexpr auto lds_block_desc = MakeLdsBlockDescriptor<Problem>();
271  auto o_lds_block = make_tensor_view<address_space_enum::lds>(
272  static_cast<ODataType*>(p_smem), lds_block_desc);
273 
274  auto in_lds_window = make_tile_window(
275  o_lds_block,
277  {0, 0},
278  LdsTileDistr);
279 
280  auto out_lds_window = make_tile_window(
281  o_lds_block,
283  {0, 0});
284 
288  constexpr index_t num_access = SFC::get_num_of_access();
289 
290  static_assert(std::is_same_v<ELayout, tensor_layout::gemm::RowMajor>,
291  "Currently, the CShuffle Epilogue only supports the Row Major Output layout");
292 
293  using TileEncodingPattern =
297  GetVectorSizeC(),
299  Problem::kNumWaveGroups>;
300  constexpr auto dram_tile_distribution =
301  TileEncodingPattern::make_2d_static_tile_distribution();
302 
303  auto d_dram_windows = generate_tuple(
304  [&](auto idx) {
305  return make_tile_window(ds_dram_windows[idx], dram_tile_distribution);
306  },
308 
309  constexpr auto c_warp_y_lengths =
310  to_sequence(CWarpDstr{}.get_ys_to_d_descriptor().get_lengths());
311  constexpr auto c_warp_y_index_zeros = uniform_sequence_gen_t<CWarpDstr::NDimY, 0>{};
312 
313  static_for<0, num_access, 1>{}([&](auto iAccess) {
314  block_sync_lds();
315  constexpr auto idx_y_start = SFC::get_index(iAccess);
316 
317  constexpr auto mIter = number<idx_y_start.at(number<0>{}) / (MPerIterationShuffle)>{};
318  constexpr auto nIter = number<idx_y_start.at(number<1>{}) / (NPerIterationShuffle)>{};
319 
320  lds_tile.get_thread_buffer() = o_acc_tile.get_y_sliced_thread_data(
323  c_warp_y_index_zeros),
325  c_warp_y_lengths));
326 
327  const auto c_warptile_in_tensor_casted = cast_tile<ODataType>(lds_tile);
328 
329  store_tile(in_lds_window, c_warptile_in_tensor_casted);
330  block_sync_lds();
331 
332  auto c_out_tensor = load_tile(make_tile_window(out_lds_window, dram_tile_distribution));
333 
334  const auto ds_tensor = generate_tuple(
335  [&](auto idx) { return load_tile(d_dram_windows[idx]); }, number<NumDTensor>{});
336 
337  const auto c_ds_tiles = concat_tuple_of_reference(
338  tie(c_out_tensor, c_out_tensor),
339  generate_tie([&](auto idx) -> const auto& { return ds_tensor[idx]; },
340  number<NumDTensor>{}));
341 
342  tile_elementwise_inout_unpack(typename Problem::CDElementwise{}, c_ds_tiles);
343 
344  if constexpr(MemoryOperation == memory_operation_enum::set)
345  {
346  store_tile(out_dram_window, c_out_tensor);
347  }
348  else
349  {
350  update_tile(out_dram_window, c_out_tensor);
351  }
352  if constexpr(iAccess != num_access - 1)
353  {
354  constexpr auto step = SFC::get_forward_step(iAccess);
355 
356  move_tile_window(out_dram_window, {step.at(number<0>{}), step.at(number<1>{})});
357 
358  static_for<0, NumDTensor, 1>{}([&](auto idx) {
359  move_tile_window(d_dram_windows[idx],
360  {step.at(number<0>{}), step.at(number<1>{})});
361  });
362  }
363  });
364  }
365 };
366 } // namespace ck_tile
CK_TILE_DEVICE void block_sync_lds()
Definition: arch.hpp:190
#define CK_TILE_DEVICE
Definition: config.hpp:41
#define CK_TILE_HOST_DEVICE
Definition: config.hpp:42
__host__ constexpr __device__ T min(T x)
Definition: math.hpp:116
constexpr CK_TILE_HOST_DEVICE auto make_embed_tile_distribution_encoding(OuterDstr, InnerDstr)
Definition: tile_distribution_encoding.hpp:457
Definition: cluster_descriptor.hpp:13
constexpr CK_TILE_HOST_DEVICE auto make_naive_tensor_descriptor(const tuple< Lengths... > &lengths, const tuple< Strides... > &strides, number< GuaranteedLastDimensionVectorLength >=number<-1 >{}, number< GuaranteedLastDimensionVectorStride >=number<-1 >{})
Definition: tensor_descriptor.hpp:268
constexpr tuple< Args &... > tie(Args &... args) noexcept
Definition: tuple.hpp:376
int32_t index_t
Definition: integer.hpp:9
constexpr CK_TILE_HOST_DEVICE auto generate_tie(F &&f, number< N >)
Definition: tuple.hpp:435
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
CK_TILE_DEVICE auto tile_elementwise_inout_unpack(const InElementFunc &in_element_func, const Tuple &t, std::index_sequence< I... >)
Template function that "unpacks" a tuple and applies an element-wise operation.
Definition: tile_elementwise.hpp:71
@ thread_raked
Thread raked pattern.
constexpr CK_TILE_HOST_DEVICE auto to_sequence(tuple< number< Is >... >)
Definition: sequence.hpp:1052
constexpr CK_TILE_HOST_DEVICE auto merge_sequences(Seqs...)
Definition: sequence.hpp:823
constexpr CK_TILE_DEVICE auto make_tile_window(null_tensor_view, const WindowLengths &window_lengths, const multi_index< WindowLengths::size()> &, Ts &&...)
Definition: null_tile_window.hpp:75
CK_TILE_DEVICE void move_tile_window(null_tile_window< WindowLengths > &, const typename null_tile_window< WindowLengths >::BottomTensorIndex &)
Definition: null_tile_window.hpp:95
constexpr CK_TILE_HOST_DEVICE auto generate_tuple(F &&f, number< N >)
Definition: tuple.hpp:429
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:360
constexpr CK_TILE_HOST_DEVICE auto concat_tuple_of_reference(const tuple< X &... > &tx, const tuple< Y &... > &ty)
Definition: tuple.hpp:443
CK_TILE_DEVICE void update_tile(tile_window_with_static_lengths< BottomTensorView_, WindowLengths_ > &tile_window_tmp, const static_distributed_tensor< DataType_, TileDistribution_ > &dstr_tensor)
Definition: update_tile.hpp:22
CK_TILE_DEVICE void store_tile(tile_window_with_static_lengths< BottomTensorView_, WindowLengths_ > &tile_window_tmp, const static_distributed_tensor< DataType_, TileDistribution_ > &dstr_tensor)
Definition: store_tile.hpp:23
constexpr CK_TILE_HOST_DEVICE T min(T x)
Definition: math.hpp:210
CK_TILE_DEVICE auto load_tile(const TileWindow_ &tile_window, number< i_access >={}, bool_constant< oob_conditional_check >={})
Definition: load_tile.hpp:22
constexpr CK_TILE_HOST_DEVICE auto make_static_tile_distribution(StaticTileDistributionEncoding_)
Definition: tile_distribution.hpp:480
typename impl::WarpGemmDispatcher< AType, BType, AccType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity, AttrNumAccess >::Type WarpGemmDispatcher
Definition: warp_gemm_dispatcher.hpp:178
typename uniform_sequence_gen< NSize, I >::type uniform_sequence_gen_t
Definition: sequence.hpp:1023
typename tuple_element< I, TTuple >::type tuple_element_t
Definition: tuple.hpp:208
__host__ constexpr __device__ index_t get_warp_size()
Definition: get_id.hpp:42
__host__ constexpr __device__ auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:211
Definition: cshuffle_epilogue.hpp:63
static constexpr index_t kBlockSize
Definition: cshuffle_epilogue.hpp:79
static constexpr CK_TILE_HOST_DEVICE auto MakeLdsBlockDescriptor()
Definition: cshuffle_epilogue.hpp:217
typename WG::CWarpTensor CWarpTensor
Definition: cshuffle_epilogue.hpp:214
remove_cvref_t< Problem_ > Problem
Definition: cshuffle_epilogue.hpp:64
static constexpr index_t MPerXdl
Definition: cshuffle_epilogue.hpp:84
static constexpr bool FixedVectorSize
Definition: cshuffle_epilogue.hpp:88
static constexpr CK_TILE_HOST_DEVICE index_t GetVectorSizeD(number< I > index)
Get the vector store size for Di tensor.
Definition: cshuffle_epilogue.hpp:135
remove_cvref_t< typename Problem::ODataType > ODataType
Definition: cshuffle_epilogue.hpp:68
static constexpr index_t kNPerBlock
Definition: cshuffle_epilogue.hpp:81
CK_TILE_DEVICE auto operator()(ODramWindow &out_dram_window, const OAccTile &o_acc_tile, const DsDramWindows &ds_dram_windows, void *p_smem)
Definition: cshuffle_epilogue.hpp:261
remove_cvref_t< typename Problem::ELayout > ELayout
Definition: cshuffle_epilogue.hpp:76
static constexpr memory_operation_enum MemoryOperation
Definition: cshuffle_epilogue.hpp:78
remove_cvref_t< typename Problem::DsLayout > DsLayout
Definition: cshuffle_epilogue.hpp:70
static constexpr CK_TILE_DEVICE auto MakeLdsDistributionEncode()
Definition: cshuffle_epilogue.hpp:239
static constexpr index_t MPerIteration
Definition: cshuffle_epilogue.hpp:90
static constexpr auto MNPerIterationShuffle
Definition: cshuffle_epilogue.hpp:194
static constexpr index_t isCTransposed
Definition: cshuffle_epilogue.hpp:87
static constexpr CK_TILE_HOST_DEVICE index_t GetSmemSize()
Definition: cshuffle_epilogue.hpp:255
static constexpr index_t MWave
Definition: cshuffle_epilogue.hpp:82
static constexpr CK_TILE_HOST_DEVICE index_t GetVectorSizeC()
Get the vector store size for C tensor.
Definition: cshuffle_epilogue.hpp:106
static constexpr index_t VectorSizeC
Definition: cshuffle_epilogue.hpp:89
remove_cvref_t< typename Problem::DsDataType > DsDataType
Definition: cshuffle_epilogue.hpp:69
std::conditional_t< std::is_same_v< BDataType, pk_int4_t >, ADataType, BDataType > BTypeToUse
Definition: cshuffle_epilogue.hpp:75
remove_cvref_t< typename Problem::CDElementwise > CDElementwise
Definition: cshuffle_epilogue.hpp:77
static constexpr index_t NPerIterationShuffle
Definition: cshuffle_epilogue.hpp:203
remove_cvref_t< typename Problem::AccDataType > AccDataType
Definition: cshuffle_epilogue.hpp:67
static constexpr index_t NumDTensor
Definition: cshuffle_epilogue.hpp:92
static constexpr index_t KPerXdl
Definition: cshuffle_epilogue.hpp:86
static constexpr index_t NumMXdlPerWavePerShuffle
Definition: cshuffle_epilogue.hpp:191
static constexpr index_t NumNXdlPerWavePerShuffle
Definition: cshuffle_epilogue.hpp:192
WarpGemmDispatcher< ATypeToUse, BTypeToUse, AccDataType, MPerXdl, NPerXdl, KPerXdl, isCTransposed > WG
Definition: cshuffle_epilogue.hpp:211
static constexpr index_t MPerIterationShuffle
Definition: cshuffle_epilogue.hpp:202
remove_cvref_t< typename Problem::BDataType > BDataType
Definition: cshuffle_epilogue.hpp:66
static constexpr auto shuffle_tile_tuple
Shuffle tile configuration parameters.
Definition: cshuffle_epilogue.hpp:164
static constexpr index_t NWave
Definition: cshuffle_epilogue.hpp:83
remove_cvref_t< typename Problem::ADataType > ADataType
Definition: cshuffle_epilogue.hpp:65
static constexpr index_t kMPerBlock
Definition: cshuffle_epilogue.hpp:80
static constexpr index_t NPerIteration
Definition: cshuffle_epilogue.hpp:91
static constexpr index_t NPerXdl
Definition: cshuffle_epilogue.hpp:85
typename WG::CWarpDstr CWarpDstr
Definition: cshuffle_epilogue.hpp:213
std::conditional_t< std::is_same_v< ADataType, pk_int4_t >, BDataType, ADataType > ATypeToUse
Definition: cshuffle_epilogue.hpp:72
Definition: cshuffle_epilogue.hpp:33
remove_cvref_t< DsDataType_ > DsDataType
Definition: cshuffle_epilogue.hpp:38
remove_cvref_t< BDataType_ > BDataType
Definition: cshuffle_epilogue.hpp:35
remove_cvref_t< AccDataType_ > AccDataType
Definition: cshuffle_epilogue.hpp:36
static constexpr index_t NPerXdl
Definition: cshuffle_epilogue.hpp:48
static constexpr index_t kBlockSize
Definition: cshuffle_epilogue.hpp:42
static constexpr index_t kMPerBlock
Definition: cshuffle_epilogue.hpp:43
static constexpr index_t KPerXdl
Definition: cshuffle_epilogue.hpp:49
remove_cvref_t< ELayout_ > ELayout
Definition: cshuffle_epilogue.hpp:40
static constexpr index_t MPerXdl
Definition: cshuffle_epilogue.hpp:47
remove_cvref_t< ADataType_ > ADataType
Definition: cshuffle_epilogue.hpp:34
static constexpr index_t MWave
Definition: cshuffle_epilogue.hpp:45
static constexpr index_t kNPerBlock
Definition: cshuffle_epilogue.hpp:44
remove_cvref_t< CDElementwise_ > CDElementwise
Definition: cshuffle_epilogue.hpp:41
static constexpr bool FixedVectorSize
Definition: cshuffle_epilogue.hpp:52
static constexpr memory_operation_enum MemoryOperation
Definition: cshuffle_epilogue.hpp:51
static constexpr index_t isCTransposed
Definition: cshuffle_epilogue.hpp:50
static constexpr index_t NumDTensor
Definition: cshuffle_epilogue.hpp:55
remove_cvref_t< DsLayout_ > DsLayout
Definition: cshuffle_epilogue.hpp:39
remove_cvref_t< ODataType_ > ODataType
Definition: cshuffle_epilogue.hpp:37
static constexpr index_t kNumWaveGroups
Definition: cshuffle_epilogue.hpp:54
static constexpr index_t NWave
Definition: cshuffle_epilogue.hpp:46
static constexpr index_t VectorSizeC
Definition: cshuffle_epilogue.hpp:53
Definition: integral_constant.hpp:13
static constexpr value_type value
Definition: integral_constant.hpp:16
Definition: sequence.hpp:49
Definition: space_filling_curve.hpp:20
Definition: functional.hpp:43
Class creating 2D static tile distribution with different load/store patterns.
Definition: static_encoding_pattern.hpp:130
Definition: tile_distribution_encoding.hpp:26
Definition: tuple.hpp:192