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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/tensor/store_tile.hpp Source File
store_tile.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
15 
16 namespace ck_tile {
17 
18 template <typename BottomTensorView_,
19  typename WindowLengths_,
20  typename TileDistribution_,
21  typename DataType_>
22 CK_TILE_DEVICE void
25 {
27  using TileDstr = remove_cvref_t<TileDistribution_>;
28 
29  static_assert(std::is_same_v<remove_cvref_t<DataType_>, DataType>, "wrong!");
30 
31  constexpr auto tile_dstr = TileDstr{};
32 
33  auto tile_window = make_tile_window(tile_window_tmp.get_bottom_tensor_view(),
34  tile_window_tmp.get_window_lengths(),
35  tile_window_tmp.get_window_origin(),
36  tile_dstr);
37 
38  tile_window.store(dstr_tensor);
39 }
40 
41 template <typename BottomTensorView_,
42  typename WindowLengths_,
43  typename TileDistribution_,
44  typename DataType_>
45 CK_TILE_DEVICE void
48 {
50  using TileDstr = remove_cvref_t<TileDistribution_>;
51 
52  static_assert(std::is_same_v<remove_cvref_t<DataType_>, DataType>, "wrong!");
53 
54  constexpr auto tile_dstr = TileDstr{};
55 
56  auto tile_window = make_tile_window(tile_window_tmp.get_bottom_tensor_view(),
57  tile_window_tmp.get_window_lengths(),
58  tile_window_tmp.get_window_origin(),
59  tile_dstr);
60 
61  tile_window.store_raw(dstr_tensor);
62 }
63 
64 template <typename BottomTensorView_,
65  typename WindowLengths_,
66  typename TileDistribution_,
67  index_t NumCoord,
68  typename DataType_>
69 CK_TILE_DEVICE void
71  WindowLengths_,
72  TileDistribution_,
73  NumCoord>& tile_window,
75 {
76  tile_window.store(dstr_tensor, number<-1>{});
77 }
78 
79 template <typename BottomTensorView_,
80  typename WindowLengths_,
81  typename TileDistribution_,
82  index_t NumCoord,
83  typename DataType_>
84 CK_TILE_DEVICE void
86  WindowLengths_,
87  TileDistribution_,
88  NumCoord>& tile_window,
90 {
91  tile_window.store_raw(dstr_tensor, number<-1>{});
92 }
93 
94 template <typename BottomTensorView_,
95  typename WindowLengths_,
96  typename TileDistribution_,
97  typename LinearBottomDims_,
98  typename DataType_>
101  tile_window,
103 {
104  tile_window.store(dstr_tensor, number<-1>{});
105 }
106 
107 template <typename BottomTensorView_,
108  typename WindowLengths_,
109  typename TileDistribution_,
110  typename LinearBottomDims_,
111  typename DataType_>
114  tile_window,
116 {
117  tile_window.store_raw(dstr_tensor, number<-1>{});
118 }
119 
120 } // namespace ck_tile
#define CK_TILE_DEVICE
Definition: config.hpp:41
Definition: cluster_descriptor.hpp:13
int32_t index_t
Definition: integer.hpp:9
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
CK_TILE_DEVICE void store_tile_raw(tile_window_with_static_lengths< BottomTensorView_, WindowLengths_ > &tile_window_tmp, const static_distributed_tensor< DataType_, TileDistribution_ > &dstr_tensor)
Definition: store_tile.hpp:46
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 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 bool is_same_v
Definition: type.hpp:283
Definition: integral_constant.hpp:13
Definition: static_distributed_tensor.hpp:21
constexpr CK_TILE_DEVICE auto get_window_origin() const
Definition: tile_window_base.hpp:45
constexpr CK_TILE_DEVICE auto get_bottom_tensor_view() const
Definition: tile_window_base.hpp:47
constexpr CK_TILE_DEVICE auto get_window_lengths() const
Definition: tile_window_base.hpp:46
Definition: tile_window_linear.hpp:55
CK_TILE_DEVICE void store_raw(const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access >={}) const
Definition: tile_window_linear.hpp:710
CK_TILE_DEVICE void store(const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}) const
Definition: tile_window_linear.hpp:657
This class provides tile (windowed) view and access to the device memory.
Definition: tile_window.hpp:46
This class provides description of tile windowed view on the device memory.
Definition: tile_window.hpp:873