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

19namespace ck_tile {

20

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

22CK_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

37template <typename TileWindow_,

38 typename ElementWise_,

39 index_t i_access = -1,

40 bool oob_conditional_check = true>

41CK_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

52template <typename DistributedTensor_,

53 typename TileWindow_,

54 index_t i_access = -1,

55 bool oob_conditional_check = true>

56CK_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

73template <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>

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

94template <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>

102CK_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

115template <typename LdsTileWindow_,

116 typename TileWindow_,

117 index_t i_access = -1,

118 bool oob_conditional_check = true>

119CK_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

128template <typename LdsTileWindow_,

129 typename TileWindow_,

130 index_t i_access = -1,

131 bool oob_conditional_check = true,

132 bool pre_nop = false>

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

145CK_TILE_DEVICE auto async_load_fence(index_t cnt = 0)

146{

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

148}

149

150template <typename WindowLengths>

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

152{

153 return null_tensor{};

154}

155

156template <typename T, typename WindowLengths>

157CK_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 tile/core/algorithm/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::bool_constant

constant< b > bool_constant

Definition tile/core/numeric/integral_constant.hpp:43

ck_tile::number

constant< v > number

Definition tile/core/numeric/integral_constant.hpp:37

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

int32_t index_t

Definition integer.hpp:9

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