include/ck_tile/core/tensor/load_tile.hpp Source File

• 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 BottomTensorView_,

22  typename WindowLengths_,

23  typename TileDistribution_,

24  index_t NumCoord,

25  index_t i_access = -1,

26  bool oob_conditional_check = true>

27 CK_TILE_DEVICE auto load_tile(const tile_window_with_static_distribution<BottomTensorView_,

28  WindowLengths_,

29  TileDistribution_,

30  NumCoord>& tile_window,

31  number<i_access> = {},

32  bool_constant<oob_conditional_check> = {})

33 {

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

35 }

36 

37 template <typename BottomTensorView_,

38  typename WindowLengths_,

39  typename TileDistribution_,

40  typename LinearBottomDims_,

41  index_t i_access = -1,

42  bool oob_conditional_check = true>

43 CK_TILE_DEVICE auto load_tile(const tile_window_linear<BottomTensorView_,

44  WindowLengths_,

45  TileDistribution_,

46  LinearBottomDims_>& tile_window,

47  number<i_access> = {},

48  bool_constant<oob_conditional_check> = {})

49 {

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

51 }

52 

53 template <typename DistributedTensor_,

54  typename BottomTensorView_,

55  typename WindowLengths_,

56  typename TileDistribution_,

57  index_t NumCoord,

58  index_t i_access = -1,

59  bool oob_conditional_check = true>

60 CK_TILE_DEVICE auto load_tile(DistributedTensor_& dst_tile,

61  const tile_window_with_static_distribution<BottomTensorView_,

62  WindowLengths_,

63  TileDistribution_,

64  NumCoord>& tile_window,

65  number<i_access> = {},

66  bool_constant<oob_conditional_check> = {})

67 {

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

69 }

70 

71 template <typename DistributedTensor_,

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 CK_TILE_DEVICE auto load_tile(DistributedTensor_& dst_tile,

79  const tile_window_linear<BottomTensorView_,

80  WindowLengths_,

81  TileDistribution_,

82  LinearBottomDims_>& tile_window,

83  number<i_access> = {},

84  bool_constant<oob_conditional_check> = {})

85 {

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

87 }

88 

98 template <typename T,

99  typename BottomTensorView_,

100  typename WindowLengths_,

101  typename TileDistribution_,

102  index_t NumCoord,

103  index_t i_access = -1,

104  bool oob_conditional_check = true,

105  bool pre_nop = false>

106 CK_TILE_DEVICE auto load_tile_raw(T& tile,

107  const tile_window_with_static_distribution<BottomTensorView_,

108  WindowLengths_,

109  TileDistribution_,

110  NumCoord>& tile_window,

111  number<i_access> = {},

112  bool_constant<oob_conditional_check> = {},

113  bool_constant<pre_nop> = {})

114 {

115  tile_window.load_raw(

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

117 }

118 

119 template <typename T,

120  typename BottomTensorView_,

121  typename WindowLengths_,

122  typename TileDistribution_,

123  typename LinearBottomDims_,

124  index_t i_access = -1,

125  bool oob_conditional_check = true,

126  bool pre_nop = false>

127 CK_TILE_DEVICE auto load_tile_raw(T& tile,

128  const tile_window_linear<BottomTensorView_,

129  WindowLengths_,

130  TileDistribution_,

131  LinearBottomDims_>& tile_window,

132  number<i_access> = {},

133  bool_constant<oob_conditional_check> = {},

134  bool_constant<pre_nop> = {})

135 {

136  tile_window.load_raw(

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

138 }

139 

140 template <typename LdsTileWindow_,

141  typename BottomTensorView_,

142  typename WindowLengths_,

143  typename TileDistribution_,

144  index_t NumCoord,

145  index_t i_access = -1,

146  bool oob_conditional_check = true,

147  bool pre_nop = false>

148 CK_TILE_DEVICE auto

149 async_load_tile_raw(LdsTileWindow_&& lds_tile,

150  const tile_window_with_static_distribution<BottomTensorView_,

151  WindowLengths_,

152  TileDistribution_,

153  NumCoord>& tile_window,

154  number<i_access> = {},

155  bool_constant<oob_conditional_check> = {},

156  bool_constant<pre_nop> = {})

157 {

158  return tile_window.async_load_raw(lds_tile,

159  number<i_access>{},

160  bool_constant<oob_conditional_check>{},

161  bool_constant<pre_nop>{});

162 }

163 

164 template <typename LdsTileWindow_,

165  typename BottomTensorView_,

166  typename WindowLengths_,

167  typename TileDistribution_,

168  typename LinearBottomDims_,

169  index_t i_access = -1,

170  bool oob_conditional_check = true,

171  bool pre_nop = false>

172 CK_TILE_DEVICE auto async_load_tile_raw(LdsTileWindow_&& lds_tile,

173  const tile_window_linear<BottomTensorView_,

174  WindowLengths_,

175  TileDistribution_,

176  LinearBottomDims_>& tile_window,

177  number<i_access> = {},

178  bool_constant<oob_conditional_check> = {},

179  bool_constant<pre_nop> = {})

180 {

181  return tile_window.async_load_raw(lds_tile,

182  number<i_access>{},

183  bool_constant<oob_conditional_check>{},

184  bool_constant<pre_nop>{});

185 }

186 

187 CK_TILE_DEVICE auto async_load_fence(index_t cnt = 0)

188 {

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

190 }

191 

192 template <typename WindowLengths>

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

194 {

195  return null_tensor{};

196 }

197 

198 template <typename T, typename WindowLengths>

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

200 {

201 }

202 

203 } // namespace ck_tile

config.hpp

CK_TILE_DEVICE

#define CK_TILE_DEVICE

Definition: config.hpp:40

coordinate_transform.hpp

integer.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::async_load_fence

CK_TILE_DEVICE auto async_load_fence(index_t cnt=0)

Definition: load_tile.hpp:187

ck_tile::async_load_tile_raw

CK_TILE_DEVICE auto async_load_tile_raw(LdsTileWindow_ &&lds_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 >={})

Definition: load_tile.hpp:149

ck_tile::index_t

int32_t index_t

Definition: integer.hpp:9

ck_tile::load_tile

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

Definition: load_tile.hpp:27

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

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

ck_tile::tile_window_with_static_distribution

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

Definition: tile_window.hpp:37

container_helper.hpp

integral_constant.hpp

math.hpp

tile_window.hpp

tile_window_linear.hpp

type_traits.hpp