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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/tensor/tile_window_utils.hpp Source File
tile_window_utils.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #include "ck_tile/core/arch/arch.hpp"
5 #include "ck_tile/core/arch/utility.hpp"
6 #include "ck_tile/core/algorithm/space_filling_curve.hpp"
7 #include "ck_tile/core/config.hpp"
8 #include "ck_tile/core/container/array.hpp"
9 #include "ck_tile/core/container/sequence.hpp"
10 #include "ck_tile/core/container/tuple.hpp"
11 #include "ck_tile/core/container/container_helper.hpp"
12 #include "ck_tile/core/tensor/static_distributed_tensor.hpp"
13 #include "ck_tile/core/tensor/tensor_adaptor.hpp"
14 #include "ck_tile/core/tensor/tile_distribution.hpp"
15 #include "ck_tile/core/utility/functional.hpp"
16 #include "ck_tile/core/utility/type_traits.hpp"
17 
18 #pragma once
19 namespace ck_tile {
20 
21 template <typename TileWindow_>
22 CK_TILE_DEVICE void move_tile_window(TileWindow_& window,
23  const typename TileWindow_::BottomTensorIndex& step)
24 {
25  window.move(step);
26 }
27 
28 // input a lds store tile, extract some information from it
29 // used to set m0 value for gfx9 serious
30 template <typename LdsTileWindow_>
31 CK_TILE_DEVICE auto get_async_store_smem_info(LdsTileWindow_&& lds_tile)
32 {
33  using LdsTileWindow = remove_cvref_t<LdsTileWindow_>;
34  using LdsDataType = typename LdsTileWindow::DataType;
35 
36  // issues * warps * lanes
37  static_assert(LdsTileWindow::get_num_of_dimension() == 3); // TODO: hard coded
38 
39  const index_t size_per_buf =
40  lds_tile.get_bottom_tensor_view().get_tensor_descriptor().calculate_offset(
41  make_tuple(number<0>{}, number<0>{}, number<0>{})) *
42  sizeof(LdsDataType);
43 
44  const index_t size_per_wave =
45  lds_tile.get_bottom_tensor_view().get_tensor_descriptor().calculate_offset(
46  make_tuple(number<0>{}, number<1>{}, number<0>{})) *
47  sizeof(LdsDataType) -
48  size_per_buf;
49 
50  const index_t size_per_issue =
51  lds_tile.get_bottom_tensor_view().get_tensor_descriptor().calculate_offset(
52  make_tuple(number<1>{}, number<0>{}, number<0>{})) *
53  sizeof(LdsDataType) -
54  size_per_buf;
55 
56  const index_t m0_init_value = size_per_buf + size_per_wave * get_warp_id();
57 
58  return make_tuple(m0_init_value, size_per_issue);
59 }
60 
61 } // namespace ck_tile
CK_TILE_DEVICE
#define CK_TILE_DEVICE
Definition: config.hpp:41
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::index_t
int32_t index_t
Definition: integer.hpp:9
ck_tile::remove_cvref_t
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
ck_tile::get_async_store_smem_info
CK_TILE_DEVICE auto get_async_store_smem_info(LdsTileWindow_ &&lds_tile)
Definition: tile_window_utils.hpp:31
ck_tile::move_tile_window
CK_TILE_DEVICE void move_tile_window(null_tile_window< WindowLengths > &, const typename null_tile_window< WindowLengths >::BottomTensorIndex &)
Definition: null_tile_window.hpp:95
ck_tile::make_tuple
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:360
space_filling_curve.hpp
static_distributed_tensor.hpp
ck_tile::constant
Definition: integral_constant.hpp:13
container_helper.hpp
tensor_adaptor.hpp
tile_distribution.hpp
type_traits.hpp