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

• develop
• include
• ck_tile
• core
• tensor

update_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 

6 #include "ck_tile/core/config.hpp"

7 #include "ck_tile/core/numeric/integer.hpp"

8 #include "ck_tile/core/numeric/integral_constant.hpp"

9 #include "ck_tile/core/algorithm/coordinate_transform.hpp"

10 #include "ck_tile/core/container/container_helper.hpp"

11 #include "ck_tile/core/numeric/math.hpp"

12 #include "ck_tile/core/tensor/tile_window.hpp"

13 #include "ck_tile/core/utility/type_traits.hpp"

14 

15 namespace ck_tile {

16 

17 template <typename BottomTensorView_,

18  typename WindowLengths_,

19  typename TileDistribution_,

20  typename DataType_>

21 CK_TILE_DEVICE void

22 update_tile(tile_window_with_static_lengths<BottomTensorView_, WindowLengths_>& tile_window_tmp,

23  const static_distributed_tensor<DataType_, TileDistribution_>& dstr_tensor)

24 {

25  using DataType = remove_cvref_t<typename BottomTensorView_::DataType>;

26  using TileDstr = remove_cvref_t<TileDistribution_>;

27 

28  static_assert(std::is_same_v<remove_cvref_t<DataType_>, DataType>, "wrong!");

29 

30  constexpr auto tile_dstr = TileDstr{};

31 

32  auto tile_window = make_tile_window(tile_window_tmp.get_bottom_tensor_view(),

33  tile_window_tmp.get_window_lengths(),

34  tile_window_tmp.get_window_origin(),

35  tile_dstr);

36 

37  tile_window.update(dstr_tensor);

38 }

39 

40 template <typename BottomTensorView_,

41  typename WindowLengths_,

42  typename TileDistribution_,

43  index_t NumCoord,

44  typename DataType_,

45  index_t i_access = -1,

46  bool oob_conditional_check = true>

47 CK_TILE_DEVICE void

48 update_tile(tile_window_with_static_distribution<BottomTensorView_,

49  WindowLengths_,

50  TileDistribution_,

51  NumCoord>& tile_window,

52  const static_distributed_tensor<DataType_, TileDistribution_>& dstr_tensor,

53  number<i_access> = {},

54  bool_constant<oob_conditional_check> = {})

55 {

56  tile_window.update(dstr_tensor, number<i_access>{}, bool_constant<oob_conditional_check>{});

57 }

58 

59 template <typename BottomTensorView_,

60  typename WindowLengths_,

61  typename TileDistribution_,

62  index_t NumCoord,

63  typename DataType_,

64  index_t i_access = -1,

65  bool oob_conditional_check = true,

66  bool pre_nop = false>

67 CK_TILE_DEVICE void

68 update_tile_raw(tile_window_with_static_distribution<BottomTensorView_,

69  WindowLengths_,

70  TileDistribution_,

71  NumCoord>& tile_window,

72  const static_distributed_tensor<DataType_, TileDistribution_>& dstr_tensor,

73  number<i_access> = {},

74  bool_constant<oob_conditional_check> = {},

75  bool_constant<pre_nop> = {})

76 {

77  tile_window.update_raw(dstr_tensor,

78  number<i_access>{},

79  bool_constant<oob_conditional_check>{},

80  bool_constant<pre_nop>{});

81 }

82 

83 template <typename BottomTensorView_,

84  typename WindowLengths_,

85  typename TileDistribution_,

86  typename LinearBottomDims_,

87  typename DataType_,

88  index_t i_access = -1,

89  bool oob_conditional_check = true,

90  bool pre_nop = false>

91 CK_TILE_DEVICE auto update_tile_raw(

92  tile_window_linear<BottomTensorView_, WindowLengths_, TileDistribution_, LinearBottomDims_>&

93  tile_window,

94  const static_distributed_tensor<DataType_, TileDistribution_>& dstr_tensor,

95  number<i_access> = {},

96  bool_constant<oob_conditional_check> = {},

97  bool_constant<pre_nop> = {})

98 {

99  tile_window.update_raw(dstr_tensor,

100  number<i_access>{},

101  bool_constant<oob_conditional_check>{},

102  bool_constant<pre_nop>{});

103 }

104 

105 } // namespace ck_tile

config.hpp

CK_TILE_DEVICE

#define CK_TILE_DEVICE

Definition: config.hpp:41

coordinate_transform.hpp

integer.hpp

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::make_tile_window

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::update_tile

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::update_tile_raw

CK_TILE_DEVICE void update_tile_raw(tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, TileDistribution_, NumCoord > &tile_window, const static_distributed_tensor< DataType_, TileDistribution_ > &dstr_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})

Definition: update_tile.hpp:68

ck::is_same_v

constexpr bool is_same_v

Definition: type.hpp:283

ck_tile::constant

Definition: integral_constant.hpp:13

ck_tile::static_distributed_tensor

Definition: static_distributed_tensor.hpp:21

ck_tile::tile_window_base::get_window_origin

constexpr CK_TILE_DEVICE auto get_window_origin() const

Definition: tile_window_base.hpp:45

ck_tile::tile_window_base::get_bottom_tensor_view

constexpr CK_TILE_DEVICE auto get_bottom_tensor_view() const

Definition: tile_window_base.hpp:47

ck_tile::tile_window_base::get_window_lengths

constexpr CK_TILE_DEVICE auto get_window_lengths() const

Definition: tile_window_base.hpp:46

ck_tile::tile_window_linear

Definition: tile_window_linear.hpp:55

ck_tile::tile_window_linear::update_raw

CK_TILE_DEVICE void update_raw(const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const

Definition: tile_window_linear.hpp:811

ck_tile::tile_window_with_static_distribution

This class provides tile (windowed) view and access to the device memory.

Definition: tile_window.hpp:46

ck_tile::tile_window_with_static_lengths

This class provides description of tile windowed view on the device memory.

Definition: tile_window.hpp:873

container_helper.hpp

integral_constant.hpp

math.hpp

tile_window.hpp

type_traits.hpp