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

• develop
• include
• ck_tile
• core
• tensor

load_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 #include "ck_tile/core/tensor/tile_window.hpp"

15 #include "ck_tile/core/tensor/tile_window_linear.hpp"

16 #include "ck_tile/core/tensor/null_tile_window.hpp"

17 #include "ck_tile/core/tensor/null_tensor.hpp"

18 

19 namespace ck_tile {

20 

21 template <typename TileWindow_, index_t i_access = -1, bool oob_conditional_check = true>

22 CK_TILE_DEVICE auto load_tile(const TileWindow_& tile_window,

23  number<i_access> = {},

24  bool_constant<oob_conditional_check> = {})

25 {

26  return tile_window.load(number<i_access>{}, bool_constant<oob_conditional_check>{});

27 }

28 

29 template <typename DistributedTensor_,

30  typename TileWindow_,

31  index_t i_access = -1,

32  bool oob_conditional_check = true>

33 CK_TILE_DEVICE auto load_tile(DistributedTensor_& dst_tile,

34  const TileWindow_& tile_window,

35  number<i_access> = {},

36  bool_constant<oob_conditional_check> = {})

37 {

38  return tile_window.load(dst_tile, number<i_access>{}, bool_constant<oob_conditional_check>{});

39 }

40 

50 template <typename T,

51  typename BottomTensorView_,

52  typename WindowLengths_,

53  typename TileDistribution_,

54  index_t NumCoord,

55  index_t i_access = -1,

56  bool oob_conditional_check = true,

57  bool pre_nop = false>

58 CK_TILE_DEVICE auto load_tile_raw(T& tile,

59  const tile_window_with_static_distribution<BottomTensorView_,

60  WindowLengths_,

61  TileDistribution_,

62  NumCoord>& tile_window,

63  number<i_access> = {},

64  bool_constant<oob_conditional_check> = {},

65  bool_constant<pre_nop> = {})

66 {

67  tile_window.load_raw(

68  tile, number<i_access>{}, bool_constant<oob_conditional_check>{}, bool_constant<pre_nop>{});

69 }

70 

71 template <typename T,

72  typename BottomTensorView_,

73  typename WindowLengths_,

74  typename TileDistribution_,

75  typename LinearBottomDims_,

76  index_t i_access = -1,

77  bool oob_conditional_check = true,

78  bool pre_nop = false>

79 CK_TILE_DEVICE auto load_tile_raw(T& tile,

80  const tile_window_linear<BottomTensorView_,

81  WindowLengths_,

82  TileDistribution_,

83  LinearBottomDims_>& tile_window,

84  number<i_access> = {},

85  bool_constant<oob_conditional_check> = {},

86  bool_constant<pre_nop> = {})

87 {

88  tile_window.load_raw(

89  tile, number<i_access>{}, bool_constant<oob_conditional_check>{}, bool_constant<pre_nop>{});

90 }

91 

92 template <typename LdsTileWindow_,

93  typename TileWindow_,

94  index_t i_access = -1,

95  bool oob_conditional_check = true>

96 CK_TILE_DEVICE auto async_load_tile(LdsTileWindow_&& lds_tile,

97  const TileWindow_& tile_window,

98  number<i_access> = {},

99  bool_constant<oob_conditional_check> = {})

100 {

101  return tile_window.async_load(

102  lds_tile, number<i_access>{}, bool_constant<oob_conditional_check>{});

103 }

104 

105 template <typename LdsTileWindow_,

106  typename TileWindow_,

107  index_t i_access = -1,

108  bool oob_conditional_check = true,

109  bool pre_nop = false>

110 CK_TILE_DEVICE auto async_load_tile_raw(LdsTileWindow_&& lds_tile,

111  const TileWindow_& tile_window,

112  number<i_access> = {},

113  bool_constant<oob_conditional_check> = {},

114  bool_constant<pre_nop> = {})

115 {

116  return tile_window.async_load_raw(lds_tile,

117  number<i_access>{},

118  bool_constant<oob_conditional_check>{},

119  bool_constant<pre_nop>{});

120 }

121 

122 CK_TILE_DEVICE auto async_load_fence(index_t cnt = 0)

123 {

124  asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory");

125 }

126 

127 template <typename WindowLengths>

128 CK_TILE_DEVICE auto load_tile(const null_tile_window<WindowLengths>&)

129 {

130  return null_tensor{};

131 }

132 

133 template <typename T, typename WindowLengths>

134 CK_TILE_DEVICE auto load_tile_raw(T& /*null_tile*/, const null_tile_window<WindowLengths>&)

135 {

136 }

137 

138 } // 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::async_load_tile

CK_TILE_DEVICE auto async_load_tile(LdsTileWindow_ &&lds_tile, const TileWindow_ &tile_window, number< i_access >={}, bool_constant< oob_conditional_check >={})

Definition: load_tile.hpp:96

ck_tile::async_load_fence

CK_TILE_DEVICE auto async_load_fence(index_t cnt=0)

Definition: load_tile.hpp:122

ck_tile::index_t

int32_t index_t

Definition: integer.hpp:9

ck_tile::load_tile_raw

CK_TILE_DEVICE auto load_tile_raw(T &tile, const tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, TileDistribution_, NumCoord > &tile_window, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})

Loads a tile of data using inline assembly.

Definition: load_tile.hpp:58

ck_tile::async_load_tile_raw

CK_TILE_DEVICE auto async_load_tile_raw(LdsTileWindow_ &&lds_tile, const TileWindow_ &tile_window, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={})

Definition: load_tile.hpp:110

ck_tile::load_tile

CK_TILE_DEVICE auto load_tile(const TileWindow_ &tile_window, number< i_access >={}, bool_constant< oob_conditional_check >={})

Definition: load_tile.hpp:22

null_tensor.hpp

null_tile_window.hpp

ck_tile::constant

Definition: integral_constant.hpp:13

ck_tile::null_tensor

Definition: null_tensor.hpp:9

ck_tile::null_tile_window

Definition: null_tile_window.hpp:19

ck_tile::tile_window_linear

Definition: tile_window_linear.hpp:55

ck_tile::tile_window_with_static_distribution

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

Definition: tile_window.hpp:46

container_helper.hpp

integral_constant.hpp

math.hpp

tile_window.hpp

tile_window_linear.hpp

type_traits.hpp