/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-2025, 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 

37 template <typename TileWindow_,

38  typename ElementWise_,

39  index_t i_access = -1,

40  bool oob_conditional_check = true>

41 CK_TILE_DEVICE auto load_tile_with_elementwise(const TileWindow_& tile_window,

42  ElementWise_ elementwise,

43  number<i_access> = {},

44  bool_constant<oob_conditional_check> = {})

45 {

46  // TODO: Tile windows should works with unknow number of params

47  // Load element_wise API works only when the input typle is a tuple-tyupe

48  return tile_window[number<0>{}].load(

49  tile_window, elementwise, number<i_access>{}, bool_constant<oob_conditional_check>{});

50 }

51 

52 template <typename DistributedTensor_,

53  typename TileWindow_,

54  index_t i_access = -1,

55  bool oob_conditional_check = true>

56 CK_TILE_DEVICE auto load_tile(DistributedTensor_& dst_tile,

57  const TileWindow_& tile_window,

58  number<i_access> = {},

59  bool_constant<oob_conditional_check> = {})

60 {

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

62 }

63 

73 template <typename T,

74  typename BottomTensorView_,

75  typename WindowLengths_,

76  typename TileDistribution_,

77  index_t NumCoord,

78  index_t i_access = -1,

79  bool oob_conditional_check = true,

80  bool pre_nop = false>

81 CK_TILE_DEVICE auto load_tile_raw(T& tile,

82  const tile_window_with_static_distribution<BottomTensorView_,

83  WindowLengths_,

84  TileDistribution_,

85  NumCoord>& tile_window,

86  number<i_access> = {},

87  bool_constant<oob_conditional_check> = {},

88  bool_constant<pre_nop> = {})

89 {

90  tile_window.load_raw(

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

92 }

93 

94 template <typename T,

95  typename BottomTensorView_,

96  typename WindowLengths_,

97  typename TileDistribution_,

98  typename LinearBottomDims_,

99  index_t i_access = -1,

100  bool oob_conditional_check = true,

101  bool pre_nop = false>

102 CK_TILE_DEVICE auto load_tile_raw(T& tile,

103  const tile_window_linear<BottomTensorView_,

104  WindowLengths_,

105  TileDistribution_,

106  LinearBottomDims_>& tile_window,

107  number<i_access> = {},

108  bool_constant<oob_conditional_check> = {},

109  bool_constant<pre_nop> = {})

110 {

111  tile_window.load_raw(

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

113 }

114 

115 template <typename LdsTileWindow_,

116  typename TileWindow_,

117  index_t i_access = -1,

118  bool oob_conditional_check = true>

119 CK_TILE_DEVICE auto async_load_tile(LdsTileWindow_&& lds_tile,

120  const TileWindow_& tile_window,

121  number<i_access> = {},

122  bool_constant<oob_conditional_check> = {})

123 {

124  return tile_window.async_load(

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

126 }

127 

128 template <typename LdsTileWindow_,

129  typename TileWindow_,

130  index_t i_access = -1,

131  bool oob_conditional_check = true,

132  bool pre_nop = false>

133 CK_TILE_DEVICE auto async_load_tile_raw(LdsTileWindow_&& lds_tile,

134  const TileWindow_& tile_window,

135  number<i_access> = {},

136  bool_constant<oob_conditional_check> = {},

137  bool_constant<pre_nop> = {})

138 {

139  return tile_window.async_load_raw(lds_tile,

140  number<i_access>{},

141  bool_constant<oob_conditional_check>{},

142  bool_constant<pre_nop>{});

143 }

144 

145 CK_TILE_DEVICE auto async_load_fence(index_t cnt = 0)

146 {

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

148 }

149 

150 template <typename WindowLengths>

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

152 {

153  return null_tensor{};

154 }

155 

156 template <typename T, typename WindowLengths>

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

158 {

159 }

160 

161 } // 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:119

ck_tile::async_load_fence

CK_TILE_DEVICE auto async_load_fence(index_t cnt=0)

Definition: load_tile.hpp:145

ck_tile::load_tile_with_elementwise

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

Load tile with elementwise function.

Definition: load_tile.hpp:41

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

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

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