/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/tensor/tensor_view.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/tensor/tensor_view.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/tensor/tensor_view.hpp Source File
tensor_view.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/arch/arch.hpp"
7 #include "ck_tile/core/config.hpp"
8 #include "ck_tile/core/numeric/integer.hpp"
9 #include "ck_tile/core/numeric/integral_constant.hpp"
10 #include "ck_tile/core/algorithm/coordinate_transform.hpp"
11 #include "ck_tile/core/container/container_helper.hpp"
12 #include "ck_tile/core/numeric/math.hpp"
13 #include "ck_tile/core/tensor/tensor_descriptor.hpp"
14 #include "ck_tile/core/utility/functional.hpp"
15 #include "ck_tile/core/utility/type_traits.hpp"
16 
17 namespace ck_tile {
18 
19 /*
20  * tensor_view
21  * abstract the underneath memory buffer(global, LDS, etc...)
22  * and provide a unified get/set function for access
23  *
24  * For addressing into the buffer we use 2 variable to control:
25  * coord : ND tensor coordinate, will calculate the actual offset inside
26  * linear_offset : 1D offset, will be used in the immediate field of
27  * the buffer instruction to help reduce register usage
28  *
29  * User can use either of the field, or both to indexing into the tensor
30  *
31  * We usually provide 2 set of API for buffer get/set, e.g.
32  * get_vectorized_elements()/get_vectorized_elements_raw()
33  * the former usually will call intrinsic or normal C function, the later
34  * usually will call inline-asm function
35  *
36  */
37 template <typename BufferView_,
38  typename TensorDesc_,
39  memory_operation_enum DstInMemOp_ = memory_operation_enum::set>
40 struct tensor_view
41 {
42  using buffer_view = remove_reference_t<BufferView_>;
43  using DataType = typename buffer_view::type;
44  using TensorDesc = remove_cvref_t<TensorDesc_>;
45  using TensorIndex = array<index_t, TensorDesc::get_num_of_top_dimension()>;
46  using TensorCoord = decltype(make_tensor_coordinate(TensorDesc{}, TensorIndex{}));
47  static constexpr auto DstInMemOp = DstInMemOp_;
48  static constexpr index_t PackedSize =
49  ck_tile::numeric_traits<remove_cvref_t<DataType>>::PackedSize;
50 
51  CK_TILE_HOST_DEVICE constexpr tensor_view() = default;
52 
53  CK_TILE_HOST_DEVICE constexpr tensor_view(const buffer_view& buffer_view,
54  const TensorDesc& desc)
55  : buf_{buffer_view}, desc_{desc}
56  {
57  }
58 
59  CK_TILE_HOST_DEVICE void init_raw() { buf_.init_raw(); }
60 
61  CK_TILE_HOST_DEVICE constexpr auto& get_tensor_descriptor() const { return desc_; }
62 
63  CK_TILE_HOST_DEVICE static constexpr index_t get_num_of_dimension()
64  {
65  return TensorDesc::get_num_of_top_dimension();
66  }
67 
68  CK_TILE_HOST_DEVICE constexpr const auto& get_buffer_view() const { return buf_; }
69 
70  CK_TILE_HOST_DEVICE constexpr auto& get_buffer_view() { return buf_; }
71 
72  // X is vector of DataType.
73  // "coord" is coordinate of DataType, not X. "coord" should be aligned to X
74  template <typename X,
75  bool oob_conditional_check = true,
76  typename std::enable_if<
77  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
78  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
79  bool>::type = false>
80  CK_TILE_HOST_DEVICE constexpr remove_cvref_t<X>
81  get_vectorized_elements(const TensorCoord& coord,
82  index_t linear_offset,
83  bool_constant<oob_conditional_check> = {}) const
84  {
85  return buf_.template get<X>(
86  coord.get_offset() / PackedSize,
87  linear_offset / PackedSize,
88  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
89  bool_constant<oob_conditional_check>{});
90  }
91 
92  template <typename X,
93  bool oob_conditional_check = true,
94  typename std::enable_if<
95  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
96  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
97  bool>::type = false>
98  CK_TILE_HOST_DEVICE constexpr remove_cvref_t<X>
99  get_vectorized_elements(const TensorCoord& coord,
100  index_t linear_offset,
101  bool is_valid_element, // flag
102  bool_constant<oob_conditional_check> = {}) const
103  {
104  return buf_.template get<X>(coord.get_offset() / PackedSize,
105  linear_offset / PackedSize,
106  is_valid_element,
107  bool_constant<oob_conditional_check>{});
108  }
109 
110  // X is vector of DataType.
111  // "coord" is coordinate of DataType, not X. "coord" should be aligned to X
112  template <typename X,
113  bool oob_conditional_check = true,
114  bool pre_nop = false,
115  typename std::enable_if<
116  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
117  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
118  bool>::type = false>
119  CK_TILE_HOST_DEVICE void get_vectorized_elements_raw(remove_cvref_t<X>& dst,
120  const TensorCoord& coord,
121  index_t linear_offset,
122  bool_constant<oob_conditional_check> = {},
123  bool_constant<pre_nop> = {}) const
124  {
125  return buf_.template get_raw<X, oob_conditional_check, pre_nop>(
126  dst,
127  coord.get_offset() / PackedSize,
128  linear_offset / PackedSize,
129  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
130  bool_constant<pre_nop>{});
131  }
132 
133  template <typename X,
134  bool oob_conditional_check = true,
135  bool pre_nop = false,
136  typename std::enable_if<
137  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
138  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
139  bool>::type = false>
140  CK_TILE_HOST_DEVICE void get_vectorized_elements_raw(remove_cvref_t<X>& dst,
141  const TensorCoord& coord,
142  index_t linear_offset,
143  bool is_valid_element,
144  bool_constant<oob_conditional_check> = {},
145  bool_constant<pre_nop> = {}) const
146  {
147  return buf_.template get_raw<X, oob_conditional_check, pre_nop>(dst,
148  coord.get_offset() /
149  PackedSize,
150  linear_offset / PackedSize,
151  is_valid_element,
152  bool_constant<pre_nop>{});
153  }
154 
155  template <typename X,
156  bool oob_conditional_check = true,
157  typename std::enable_if<
158  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
159  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
160  bool>::type = false>
161  CK_TILE_HOST_DEVICE constexpr void
162  async_get_vectorized_elements(CK_TILE_LDS_ADDR remove_cvref_t<DataType>* smem,
163  const TensorCoord& coord,
164  index_t linear_offset,
165  bool_constant<oob_conditional_check> = {}) const
166  {
167  return buf_.template async_get<X>(
168  smem,
169  coord.get_offset() / PackedSize,
170  linear_offset / PackedSize,
171  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
172  bool_constant<oob_conditional_check>{});
173  }
174 
175  template <typename X,
176  bool oob_conditional_check = true,
177  typename std::enable_if<
178  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
179  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
180  bool>::type = false>
181  CK_TILE_HOST_DEVICE constexpr void
182  async_get_vectorized_elements(CK_TILE_LDS_ADDR remove_cvref_t<DataType>* smem,
183  const TensorCoord& coord,
184  index_t linear_offset,
185  bool is_valid_element,
186  bool_constant<oob_conditional_check> = {}) const
187  {
188  return buf_.template async_get<X>(smem,
189  coord.get_offset() / PackedSize,
190  linear_offset / PackedSize,
191  is_valid_element,
192  bool_constant<oob_conditional_check>{});
193  }
194 
195  template <typename X,
196  bool pre_nop = false,
197  typename std::enable_if<
198  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
199  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
200  bool>::type = false>
201  CK_TILE_HOST_DEVICE constexpr void
202  async_get_vectorized_elements_raw(remove_cvref_t<DataType>* smem,
203  const TensorCoord& coord,
204  index_t linear_offset,
205  bool_constant<pre_nop> = {}) const
206  {
207  return buf_.template async_get_raw<X>(
208  smem,
209  coord.get_offset() / PackedSize,
210  linear_offset / PackedSize,
211  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
212  bool_constant<pre_nop>{});
213  }
214 
215  template <typename X,
216  bool pre_nop = false,
217  typename std::enable_if<
218  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
219  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
220  bool>::type = false>
221  CK_TILE_HOST_DEVICE constexpr void
222  async_get_vectorized_elements_raw(remove_cvref_t<DataType>* smem,
223  const TensorCoord& coord,
224  index_t coord_extra_offset,
225  index_t linear_offset,
226  bool_constant<pre_nop> = {}) const
227  {
228  return buf_.template async_get_raw<X>(
229  smem,
230  (coord.get_offset() + coord_extra_offset) / PackedSize,
231  linear_offset / PackedSize,
232  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
233  bool_constant<pre_nop>{});
234  }
235 
236  template <typename X,
237  bool pre_nop = false,
238  typename std::enable_if<
239  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
240  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
241  bool>::type = false>
242  CK_TILE_HOST_DEVICE constexpr void
243  async_get_vectorized_elements_raw(remove_cvref_t<DataType>* smem,
244  const TensorCoord& coord,
245  index_t linear_offset,
246  bool is_valid_element,
247  bool_constant<pre_nop> = {}) const
248  {
249  return buf_.template async_get_raw<X>(smem,
250  coord.get_offset() / PackedSize,
251  linear_offset / PackedSize,
252  is_valid_element,
253  bool_constant<pre_nop>{});
254  }
255 
256  template <typename X,
257  typename std::enable_if<
258  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
259  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
260  bool>::type = false>
261  CK_TILE_HOST_DEVICE constexpr remove_cvref_t<X>
262  get_transpose_vectorized_elements(const TensorCoord& coord, index_t linear_offset) const
263  {
264  return buf_.template transpose_get<X>(
265  coord.get_offset(),
266  linear_offset,
267  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord));
268  }
269 
270  template <typename X,
271  typename std::enable_if<
272  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
273  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
274  bool>::type = false>
275  CK_TILE_HOST_DEVICE constexpr remove_cvref_t<X>
276  get_transpose_vectorized_elements(const TensorCoord& coord,
277  index_t linear_offset,
278  bool is_valid_element // flag
279  ) const
280  {
281  return buf_.template transpose_get<X>(coord.get_offset(), linear_offset, is_valid_element);
282  }
283  // X is vector of DataType.
284  // "coord" is coordinate of DataType, not X. "coord" should be aligned to X
285  template <typename X,
286  bool oob_conditional_check = true,
287  typename std::enable_if<
288  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
289  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
290  bool>::type = false>
291  CK_TILE_HOST_DEVICE constexpr void
292  set_vectorized_elements(const TensorCoord& coord,
293  index_t linear_offset,
294  const X& x,
295  bool_constant<oob_conditional_check> = {})
296  {
297  buf_.template set<X, oob_conditional_check>(
298  coord.get_offset() / PackedSize,
299  linear_offset / PackedSize,
300  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
301  x);
302  }
303 
304  template <typename X,
305  bool oob_conditional_check = true,
306  typename std::enable_if<
307  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
308  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
309  bool>::type = false>
310  CK_TILE_HOST_DEVICE constexpr void
311  set_vectorized_elements(const TensorCoord& coord,
312  index_t linear_offset,
313  bool is_valid_element,
314  const X& x,
315  bool_constant<oob_conditional_check> = {})
316  {
317  buf_.template set<X, oob_conditional_check>(
318  coord.get_offset(), linear_offset, is_valid_element, x);
319  }
320 
321  template <typename X,
322  bool oob_conditional_check = true,
323  typename std::enable_if<
324  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
325  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
326  bool>::type = false>
327  CK_TILE_HOST_DEVICE constexpr void
328  set_vectorized_elements_raw(const TensorCoord& coord,
329  index_t linear_offset,
330  const X& x,
331  bool_constant<oob_conditional_check> = {})
332  {
333  buf_.template set_raw<X, oob_conditional_check>(
334  coord.get_offset() / PackedSize,
335  linear_offset / PackedSize,
336  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
337  x);
338  }
339 
340  template <typename X,
341  bool oob_conditional_check = true,
342  typename std::enable_if<
343  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
344  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
345  bool>::type = false>
346  CK_TILE_HOST_DEVICE constexpr void
347  set_vectorized_elements_raw(const TensorCoord& coord,
348  index_t linear_offset,
349  bool is_valid_element,
350  const X& x,
351  bool_constant<oob_conditional_check> = {})
352  {
353  buf_.template set_raw<X, oob_conditional_check>(
354  coord.get_offset() / PackedSize, linear_offset / PackedSize, is_valid_element, x);
355  }
356 
357  // X is vector of DataType.
358  // "coord" is coordinate of DataType, not X. "coord" should be aligned to X
359  template <typename X,
360  bool oob_conditional_check = true,
361  typename std::enable_if<
362  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
363  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
364  bool>::type = false>
365  CK_TILE_HOST_DEVICE constexpr void
366  update_vectorized_elements(const TensorCoord& coord,
367  index_t linear_offset,
368  const X& x,
369  bool_constant<oob_conditional_check> = {})
370  {
371  buf_.template update<DstInMemOp, X, oob_conditional_check>(
372  coord.get_offset() / PackedSize,
373  linear_offset / PackedSize,
374  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
375  x);
376  }
377 
378  template <typename X,
379  bool oob_conditional_check = true,
380  typename std::enable_if<
381  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
382  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
383  bool>::type = false>
384  CK_TILE_HOST_DEVICE constexpr void
385  update_vectorized_elements(const TensorCoord& coord,
386  index_t linear_offset,
387  bool is_valid_element,
388  const X& x,
389  bool_constant<oob_conditional_check> = {})
390  {
391  buf_.template update<DstInMemOp, X, oob_conditional_check>(
392  coord.get_offset() / PackedSize, linear_offset / PackedSize, is_valid_element, x);
393  }
394 
395  // X is vector of DataType.
396  // "coord" is coordinate of DataType, not X. "coord" should be aligned to X
397  template <typename X,
398  bool oob_conditional_check = true,
399  bool pre_nop = false,
400  typename std::enable_if<
401  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
402  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
403  bool>::type = false>
404  CK_TILE_HOST_DEVICE constexpr void
405  update_vectorized_elements_raw(const TensorCoord& coord,
406  index_t linear_offset,
407  const X& x,
408  bool_constant<oob_conditional_check> = {},
409  bool_constant<pre_nop> = {})
410  {
411  buf_.template update_raw<DstInMemOp, X, oob_conditional_check, pre_nop>(
412  coord.get_offset() / PackedSize,
413  linear_offset / PackedSize,
414  coordinate_has_valid_offset_assuming_top_index_is_valid(desc_, coord),
415  x);
416  }
417 
418  template <typename X,
419  bool oob_conditional_check = true,
420  bool pre_nop = false,
421  typename std::enable_if<
422  std::is_same_v<typename vector_traits<remove_cvref_t<X>>::scalar_type,
423  typename vector_traits<remove_cvref_t<DataType>>::scalar_type>,
424  bool>::type = false>
425  CK_TILE_HOST_DEVICE constexpr void
426  update_vectorized_elements_raw(const TensorCoord& coord,
427  index_t linear_offset,
428  bool is_valid_element,
429  const X& x,
430  bool_constant<oob_conditional_check> = {},
431  bool_constant<pre_nop> = {})
432  {
433  buf_.template update_raw<DstInMemOp, X, oob_conditional_check, pre_nop>(
434  coord.get_offset() / PackedSize, linear_offset / PackedSize, is_valid_element, x);
435  }
436 
437  // member
438  buffer_view buf_;
439  TensorDesc desc_;
440 };
441 
442 // placeholder type if we want to opt-out a tile view parameter
443 struct null_tensor_view
444 {
445 };
446 
447 template <address_space_enum BufferAddressSpace = address_space_enum::generic,
448  memory_operation_enum DstInMemOp = memory_operation_enum::set,
449  amd_buffer_coherence_enum Coherence = amd_buffer_coherence_enum::coherence_default,
450  typename DataType,
451  typename... Ts>
452 CK_TILE_HOST_DEVICE constexpr auto make_tensor_view(DataType* __restrict__ p,
453  const tensor_descriptor<Ts...>& desc)
454 {
455  auto buffer_view =
456  make_buffer_view<BufferAddressSpace, Coherence>(p, desc.get_element_space_size());
457 
458  return tensor_view<decltype(buffer_view), decltype(desc)>{buffer_view, desc};
459 }
460 
461 template <address_space_enum BufferAddressSpace = address_space_enum::generic,
462  memory_operation_enum DstInMemOp = memory_operation_enum::set,
463  amd_buffer_coherence_enum Coherence = amd_buffer_coherence_enum::coherence_default,
464  typename DataType,
465  typename... Lengths,
466  typename... Strides,
467  index_t GuaranteedLastDimensionVectorLength = -1,
468  index_t GuaranteedLastDimensionVectorStride = -1,
469  typename std::enable_if<sizeof...(Lengths) == sizeof...(Strides), bool>::type = false>
470 CK_TILE_HOST_DEVICE constexpr auto
471 make_naive_tensor_view(DataType* __restrict__ p,
472  const tuple<Lengths...>& lengths,
473  const tuple<Strides...>& strides,
474  number<GuaranteedLastDimensionVectorLength> = number<-1>{},
475  number<GuaranteedLastDimensionVectorStride> = number<-1>{})
476 {
477  auto desc = make_naive_tensor_descriptor(lengths,
478  strides,
479  number<GuaranteedLastDimensionVectorLength>{},
480  number<GuaranteedLastDimensionVectorStride>{});
481 
482  auto buffer_view =
483  make_buffer_view<BufferAddressSpace, Coherence>(p, desc.get_element_space_size());
484 
485  return tensor_view<decltype(buffer_view), decltype(desc), DstInMemOp>{buffer_view, desc};
486 }
487 
488 template <address_space_enum BufferAddressSpace = address_space_enum::generic,
489  amd_buffer_coherence_enum Coherence = amd_buffer_coherence_enum::coherence_default,
490  typename DataType,
491  typename... Lengths,
492  index_t GuaranteedLastDimensionVectorLength = -1>
493 CK_TILE_HOST_DEVICE constexpr auto
494 make_naive_tensor_view_packed(DataType* __restrict__ p,
495  const tuple<Lengths...>& lengths,
496  number<GuaranteedLastDimensionVectorLength> = number<-1>{})
497 {
498  auto desc =
499  make_naive_tensor_descriptor_packed(lengths, number<GuaranteedLastDimensionVectorLength>{});
500 
501  auto buffer_view =
502  make_buffer_view<BufferAddressSpace, Coherence>(p, desc.get_element_space_size());
503 
504  return tensor_view<decltype(buffer_view), decltype(desc)>{buffer_view, desc};
505 }
506 
507 template <typename OldTensorView,
508  typename NewTransforms,
509  typename NewLowerDimensionOldVisibleIdss,
510  typename NewUpperDimensionNewVisibleIdss>
511 CK_TILE_HOST_DEVICE constexpr auto transform_tensor_view(const OldTensorView& old_tensor_view,
512  const NewTransforms& new_transforms,
513  NewLowerDimensionOldVisibleIdss,
514  NewUpperDimensionNewVisibleIdss)
515 {
516  auto new_desc = transform_tensor_descriptor(old_tensor_view.desc_,
517  new_transforms,
518  NewLowerDimensionOldVisibleIdss{},
519  NewUpperDimensionNewVisibleIdss{});
520 
521  return tensor_view<typename OldTensorView::buffer_view,
522  remove_cvref_t<decltype(new_desc)>,
523  remove_cvref_t<OldTensorView>::DstInMemOp>{old_tensor_view.buf_, new_desc};
524 }
525 
526 template <typename TensorView,
527  typename TileLengths, // tuple<...>
528  typename DoPads> // sequence<bool, bool, ...>
529 CK_TILE_HOST_DEVICE constexpr auto
530 pad_tensor_view(const TensorView& tensor_view, const TileLengths& tile_lengths, DoPads)
531 {
532 
533  constexpr index_t num_dim = DoPads::size();
534 
535  static_assert(num_dim == TileLengths::size() && num_dim == TensorView::get_num_of_dimension(),
536  "wrong! inconsistent # of dimensions");
537 
538  // transforms
539  const auto transforms = generate_tuple(
540  [&](auto idim) {
541  const auto old_length = tensor_view.get_tensor_descriptor().get_length(idim);
542 
543  const auto tile_length = tile_lengths[idim];
544 
545  const auto new_length = integer_divide_ceil(old_length, tile_length) * tile_length;
546 
547  const auto pad_length = new_length - old_length;
548 
549  constexpr bool DoPad = DoPads::at(idim);
550 
551  const auto transform =
552  conditional_expr<DoPad>(make_right_pad_transform(old_length, pad_length),
553  make_pass_through_transform(old_length));
554 
555  return transform;
556  },
557  number<num_dim>{});
558 
559  // lower dimension Id
560  const auto lower_dimss =
561  generate_tuple([&](auto idim) { return sequence<idim.value>{}; }, number<num_dim>{});
562 
563  // upper dimension Id
564  const auto upper_dimss = lower_dimss;
565 
566  return transform_tensor_view(tensor_view, transforms, lower_dimss, upper_dimss);
567 }
568 
569 } // namespace ck_tile
CK_TILE_LDS_ADDR
#define CK_TILE_LDS_ADDR
Definition: config.hpp:58
CK_TILE_HOST_DEVICE
#define CK_TILE_HOST_DEVICE
Definition: config.hpp:42
coordinate_transform.hpp
ck::ranges::transform
auto transform(InputRange &&range, OutputIterator iter, UnaryOperation unary_op) -> decltype(std::transform(std::begin(range), std::end(range), iter, unary_op))
Definition: algorithm.hpp:36
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::make_right_pad_transform
constexpr CK_TILE_HOST_DEVICE auto make_right_pad_transform(const LowLength &low_length, const RightPadLength &right_pad_, bool_constant< SkipIsValidCheck >=bool_constant< false >{})
Definition: coordinate_transform.hpp:1584
ck_tile::make_naive_tensor_descriptor
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
ck_tile::integer_divide_ceil
constexpr CK_TILE_HOST_DEVICE auto integer_divide_ceil(X x, Y y)
Definition: math.hpp:149
ck_tile::amd_buffer_coherence_enum
amd_buffer_coherence_enum
Definition: amd_buffer_addressing.hpp:1332
ck_tile::amd_buffer_coherence_enum::coherence_default
@ coherence_default
ck_tile::transform_tensor_view
constexpr CK_TILE_HOST_DEVICE auto transform_tensor_view(const OldTensorView &old_tensor_view, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition: tensor_view.hpp:511
ck_tile::make_naive_tensor_view_packed
constexpr CK_TILE_HOST_DEVICE auto make_naive_tensor_view_packed(DataType *__restrict__ p, const tuple< Lengths... > &lengths, number< GuaranteedLastDimensionVectorLength >=number<-1 >{})
Definition: tensor_view.hpp:494
ck_tile::index_t
int32_t index_t
Definition: integer.hpp:9
ck_tile::pad_tensor_view
constexpr CK_TILE_HOST_DEVICE auto pad_tensor_view(const TensorView &tensor_view, const TileLengths &tile_lengths, DoPads)
Definition: tensor_view.hpp:530
ck_tile::make_pass_through_transform
constexpr CK_TILE_HOST_DEVICE auto make_pass_through_transform(const LowLength &low_length)
Definition: coordinate_transform.hpp:1558
ck_tile::make_tensor_coordinate
constexpr CK_TILE_HOST_DEVICE auto make_tensor_coordinate(const TensorDesc &tensor_desc, const TopIndex &idx_top)
Definition: tensor_coordinate.hpp:60
ck_tile::remove_cvref_t
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
ck_tile::number
constant< v > number
Definition: integral_constant.hpp:37
ck_tile::make_tensor_view
constexpr CK_TILE_HOST_DEVICE auto make_tensor_view(DataType *__restrict__ p, const tensor_descriptor< Ts... > &desc)
Definition: tensor_view.hpp:452
ck_tile::coordinate_has_valid_offset_assuming_top_index_is_valid
constexpr CK_TILE_HOST_DEVICE bool coordinate_has_valid_offset_assuming_top_index_is_valid(const TensorDesc &tensor_desc, const TensorCoord &coord)
Definition: tensor_coordinate.hpp:79
ck_tile::transform_tensor_descriptor
constexpr CK_TILE_HOST_DEVICE auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldTopIdss, NewUpperDimensionNewTopIdss)
Definition: tensor_descriptor.hpp:197
ck_tile::make_naive_tensor_descriptor_packed
constexpr CK_TILE_HOST_DEVICE auto make_naive_tensor_descriptor_packed(const tuple< Lengths... > &lengths, number< GuaranteedLastDimensionVectorLength >=number<-1 >{})
Definition: tensor_descriptor.hpp:365
ck_tile::generate_tuple
constexpr CK_TILE_HOST_DEVICE auto generate_tuple(F &&f, number< N >)
Definition: tuple.hpp:429
ck_tile::remove_reference_t
typename std::remove_reference< T >::type remove_reference_t
Definition: type_traits.hpp:15
ck_tile::make_naive_tensor_view
constexpr CK_TILE_HOST_DEVICE auto make_naive_tensor_view(DataType *__restrict__ p, const tuple< Lengths... > &lengths, const tuple< Strides... > &strides, number< GuaranteedLastDimensionVectorLength >=number<-1 >{}, number< GuaranteedLastDimensionVectorStride >=number<-1 >{})
Definition: tensor_view.hpp:471
ck::enable_if
std::enable_if< B, T > enable_if
Definition: enable_if.hpp:24
ck_tile::array
A fixed-size array container similar to std::array with additional utilities.
Definition: array.hpp:43
ck_tile::buffer_view
Definition: buffer_view.hpp:38
ck_tile::constant
Definition: integral_constant.hpp:13
ck_tile::null_tensor_view
Definition: tensor_view.hpp:444
ck_tile::numeric_traits
Definition: numeric.hpp:81
ck_tile::sequence
Definition: sequence.hpp:49
ck_tile::tensor_descriptor
Definition: tensor_descriptor.hpp:34
ck_tile::tensor_descriptor::get_element_space_size
constexpr CK_TILE_HOST_DEVICE auto get_element_space_size() const
Definition: tensor_descriptor.hpp:96
ck_tile::tensor_view
Definition: tensor_view.hpp:41
ck_tile::tensor_view::set_vectorized_elements
constexpr CK_TILE_HOST_DEVICE void set_vectorized_elements(const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={})
Definition: tensor_view.hpp:292
ck_tile::tensor_view::desc_
TensorDesc desc_
Definition: tensor_view.hpp:439
ck_tile::tensor_view::tensor_view
constexpr CK_TILE_HOST_DEVICE tensor_view()=default
ck_tile::tensor_view::get_transpose_vectorized_elements
constexpr CK_TILE_HOST_DEVICE remove_cvref_t< X > get_transpose_vectorized_elements(const TensorCoord &coord, index_t linear_offset) const
Definition: tensor_view.hpp:262
ck_tile::tensor_view::update_vectorized_elements_raw
constexpr CK_TILE_HOST_DEVICE void update_vectorized_elements_raw(const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})
Definition: tensor_view.hpp:405
ck_tile::tensor_view::TensorCoord
decltype(make_tensor_coordinate(TensorDesc{}, TensorIndex{})) TensorCoord
Definition: tensor_view.hpp:46
ck_tile::tensor_view::PackedSize
static constexpr index_t PackedSize
Definition: tensor_view.hpp:48
ck_tile::tensor_view::set_vectorized_elements_raw
constexpr CK_TILE_HOST_DEVICE void set_vectorized_elements_raw(const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={})
Definition: tensor_view.hpp:347
ck_tile::tensor_view::async_get_vectorized_elements_raw
constexpr CK_TILE_HOST_DEVICE void async_get_vectorized_elements_raw(remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset, bool_constant< pre_nop >={}) const
Definition: tensor_view.hpp:202
ck_tile::tensor_view::get_transpose_vectorized_elements
constexpr CK_TILE_HOST_DEVICE remove_cvref_t< X > get_transpose_vectorized_elements(const TensorCoord &coord, index_t linear_offset, bool is_valid_element) const
Definition: tensor_view.hpp:276
ck_tile::tensor_view::DataType
typename buffer_view::type DataType
Definition: tensor_view.hpp:43
ck_tile::tensor_view::update_vectorized_elements
constexpr CK_TILE_HOST_DEVICE void update_vectorized_elements(const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={})
Definition: tensor_view.hpp:366
ck_tile::tensor_view::set_vectorized_elements
constexpr CK_TILE_HOST_DEVICE void set_vectorized_elements(const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={})
Definition: tensor_view.hpp:311
ck_tile::tensor_view::update_vectorized_elements_raw
constexpr CK_TILE_HOST_DEVICE void update_vectorized_elements_raw(const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})
Definition: tensor_view.hpp:426
ck_tile::tensor_view::async_get_vectorized_elements
constexpr CK_TILE_HOST_DEVICE void async_get_vectorized_elements(CK_TILE_LDS_ADDR remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset, bool is_valid_element, bool_constant< oob_conditional_check >={}) const
Definition: tensor_view.hpp:182
ck_tile::tensor_view::get_buffer_view
constexpr CK_TILE_HOST_DEVICE auto & get_buffer_view()
Definition: tensor_view.hpp:70
ck_tile::tensor_view::buf_
buffer_view buf_
Definition: tensor_view.hpp:438
ck_tile::tensor_view::TensorIndex
array< index_t, TensorDesc::get_num_of_top_dimension()> TensorIndex
Definition: tensor_view.hpp:45
ck_tile::tensor_view::get_vectorized_elements_raw
CK_TILE_HOST_DEVICE void get_vectorized_elements_raw(remove_cvref_t< X > &dst, const TensorCoord &coord, index_t linear_offset, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
Definition: tensor_view.hpp:119
ck_tile::tensor_view::async_get_vectorized_elements
constexpr CK_TILE_HOST_DEVICE void async_get_vectorized_elements(CK_TILE_LDS_ADDR remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset, bool_constant< oob_conditional_check >={}) const
Definition: tensor_view.hpp:162
ck_tile::tensor_view::DstInMemOp
static constexpr auto DstInMemOp
Definition: tensor_view.hpp:47
ck_tile::tensor_view::async_get_vectorized_elements_raw
constexpr CK_TILE_HOST_DEVICE void async_get_vectorized_elements_raw(remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t coord_extra_offset, index_t linear_offset, bool_constant< pre_nop >={}) const
Definition: tensor_view.hpp:222
ck_tile::tensor_view::update_vectorized_elements
constexpr CK_TILE_HOST_DEVICE void update_vectorized_elements(const TensorCoord &coord, index_t linear_offset, bool is_valid_element, const X &x, bool_constant< oob_conditional_check >={})
Definition: tensor_view.hpp:385
ck_tile::tensor_view::set_vectorized_elements_raw
constexpr CK_TILE_HOST_DEVICE void set_vectorized_elements_raw(const TensorCoord &coord, index_t linear_offset, const X &x, bool_constant< oob_conditional_check >={})
Definition: tensor_view.hpp:328
ck_tile::tensor_view::get_buffer_view
constexpr CK_TILE_HOST_DEVICE const auto & get_buffer_view() const
Definition: tensor_view.hpp:68
ck_tile::tensor_view::init_raw
CK_TILE_HOST_DEVICE void init_raw()
Definition: tensor_view.hpp:59
ck_tile::tensor_view::tensor_view
constexpr CK_TILE_HOST_DEVICE tensor_view(const buffer_view &buffer_view, const TensorDesc &desc)
Definition: tensor_view.hpp:53
ck_tile::tensor_view::async_get_vectorized_elements_raw
constexpr CK_TILE_HOST_DEVICE void async_get_vectorized_elements_raw(remove_cvref_t< DataType > *smem, const TensorCoord &coord, index_t linear_offset, bool is_valid_element, bool_constant< pre_nop >={}) const
Definition: tensor_view.hpp:243
ck_tile::tensor_view::get_vectorized_elements
constexpr CK_TILE_HOST_DEVICE remove_cvref_t< X > get_vectorized_elements(const TensorCoord &coord, index_t linear_offset, bool is_valid_element, bool_constant< oob_conditional_check >={}) const
Definition: tensor_view.hpp:99
ck_tile::tensor_view::get_vectorized_elements_raw
CK_TILE_HOST_DEVICE void get_vectorized_elements_raw(remove_cvref_t< X > &dst, const TensorCoord &coord, index_t linear_offset, bool is_valid_element, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
Definition: tensor_view.hpp:140
ck_tile::tensor_view::get_tensor_descriptor
constexpr CK_TILE_HOST_DEVICE auto & get_tensor_descriptor() const
Definition: tensor_view.hpp:61
ck_tile::tensor_view::get_vectorized_elements
constexpr CK_TILE_HOST_DEVICE remove_cvref_t< X > get_vectorized_elements(const TensorCoord &coord, index_t linear_offset, bool_constant< oob_conditional_check >={}) const
Definition: tensor_view.hpp:81
ck_tile::tensor_view::get_num_of_dimension
static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_dimension()
Definition: tensor_view.hpp:63
ck_tile::tensor_view::buffer_view
remove_reference_t< BufferView_ > buffer_view
Definition: tensor_view.hpp:42
ck_tile::tensor_view::TensorDesc
remove_cvref_t< TensorDesc_ > TensorDesc
Definition: tensor_view.hpp:44
ck_tile::tuple
Definition: tuple.hpp:192
ck_tile::vector_traits
Definition: vector_type.hpp:89
container_helper.hpp
integral_constant.hpp
tensor_descriptor.hpp
type_traits.hpp