ContextQMD
Back to Cutlass

CUTLASS: tile_iterator_simt.h Source File

docs/tile__iterator__simt_8h_source.html

4.4.226.5 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

tile_iterator_simt.h

[Go to the documentation of this file.](tile iterator simt_8h.html)

1 /***************************************************************************************************

2 * Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.

3 *

4 * Redistribution and use in source and binary forms, with or without modification, are permitted

5 * provided that the following conditions are met:

6 * * Redistributions of source code must retain the above copyright notice, this list of

7 * conditions and the following disclaimer.

8 * * Redistributions in binary form must reproduce the above copyright notice, this list of

9 * conditions and the following disclaimer in the documentation and/or other materials

10 * provided with the distribution.

11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used

12 * to endorse or promote products derived from this software without specific prior written

13 * permission.

14 *

15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR

16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND

17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE

18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,

19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;

20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,

21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE

22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

23 *

24 **************************************************************************************************/

29 #pragma once

30

31 #include "cutlass/array.h"

32 #include "cutlass/layout/matrix.h"

33 #include "cutlass/layout/pitch_linear.h"

34

35 #include "cutlass/epilogue/warp/simt_policy.h"

36

[37](tile iterator simt_8h.html#af5392199c6e71a31335abd0bb1d9ba36) #define CUTLASS_SIMT_EPILOGUE_USE_SCALAR_STORES 1

38

40

41 namespace cutlass {

42 namespace epilogue {

43 namespace warp {

44

46

48 template <

49typename WarpShape,

50typename Operator,

51typename Element,

52typename Layout,

53typename MmaSimtPolicy

54 >

55 class TileIteratorSimt;

56

58

60 template <

61typename WarpShape_,

62typename Operator_,

63typename Element_,

64typename MmaSimtPolicy_

65 >

[66](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html) class TileIteratorSimt<WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_> {

67 public:

68

[69](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aff99d6853af936cc23506b16b780ce09)using [WarpShape](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aff99d6853af936cc23506b16b780ce09) = WarpShape_;

[70](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa712e96097d9f487858208c04c83d26b)using Operator = Operator_;

[71](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a7538a71b8c5c38ffbb2056e2c425e47b)using Element = Element_;

[72](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a9928349624ec760485d8921255edf9e4)using Layout = layout::RowMajor;

73

[74](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a9124f663869b4a8d67664c0b5e41e912)using TensorRef = TensorRef<Element, Layout>;

[75](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a126b9160dc4a4c6156be464673202e1d)using TensorCoord = MatrixCoord;

[76](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a05ff07efb8d79a930472de75ea30938a)using [Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a05ff07efb8d79a930472de75ea30938a) = typename TensorRef::Index;

[77](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a5c553a18b6215ed1158e863851be631a)using [LongIndex](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a5c553a18b6215ed1158e863851be631a) = typename TensorRef::LongIndex;

78

[79](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ad70c4e067ea653db52e72cd07a918cc7)using Policy = SimtPolicy<WarpShape, Operator, Layout, MmaSimtPolicy_>;

80

82using Shape = MatrixShape<

83 Policy::kRowsPerIteration,

84 WarpShape::kN

[85](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ad45557644f48e06a4d68c4eb6f8515e8) >;

86

88using [Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa1caeb6928f62ceabba7817c754b08b7) = Array<

89typename Operator::ElementC,

[90](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa1caeb6928f62ceabba7817c754b08b7) Policy::kElementsPerIteration>;

91

93using [AccumulatorTile](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a6c100257da1ab23db35da3c4818a32f3) = Array<

94typename Operator::ElementC,

[95](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a6c100257da1ab23db35da3c4818a32f3) Policy::kAccumulatorElementCount>;

96

[98](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ad22571b7a352a6c9834341ff15614759)static int const kIterations = Policy::kIterations;

99

101using Padding = MatrixShape<

102 0,

[103](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa56eb3d1c6b3aea627b8ee024be0e451) 4 * Policy::kElementsPerAccess>;

104

105 private:

106

108using AccessType = AlignedArray<Element, Policy::kElementsPerAccess>;

109

110//

111// Data members

112//

113

115AccessType *pointer_;

116

118 Layout layout_;

119

120 public:

121

123CUTLASS_HOST_DEVICE

[124](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a0e0346d2b4a2e5b111a2d1e6d3ac775e)[TileIteratorSimt](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a0e0346d2b4a2e5b111a2d1e6d3ac775e)(): pointer_(nullptr) { }

125

127CUTLASS_HOST_DEVICE

[128](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a25c74737e253375473171917c4f3df6f)[TileIteratorSimt](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a25c74737e253375473171917c4f3df6f)(

129TensorRef const &ref,

130unsigned lane_id

131 ):

132 pointer_(reinterpret_cast<AccessType *>(ref.data())),

133 layout_(ref.stride()[0] / Policy::kElementsPerAccess) {

134

135auto lane_layout = Policy::MmaSimtPolicy::get_lane_layout();

136MatrixCoord lane_offset = lane_layout.inverse(lane_id);

137

138 pointer_ += layout_(lane_offset);

139 }

140

142CUTLASS_HOST_DEVICE

[143](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a59a4695befd919265485918ee935aab1)TileIteratorSimt & [add_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a59a4695befd919265485918ee935aab1)([Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a05ff07efb8d79a930472de75ea30938a) pointer_offset) {

144 pointer_ += pointer_offset / Policy::kElementsPerAccess;

145return *this;

146 }

147

149CUTLASS_HOST_DEVICE

[150](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ae9f188af3f314fb898e7ad842cba98df)TileIteratorSimt & [add_tile_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ae9f188af3f314fb898e7ad842cba98df)(TensorCoord const &tile_offset) {

151

152 pointer_ += layout_({

153 tile_offset.row() * Shape::kRow,

154 (tile_offset.column() * Shape::kColumn / Policy::kElementsPerAccess)

155 });

156

157return *this;

158 }

159

161CUTLASS_HOST_DEVICE

[162](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a1a76d3b0fd2419e78d66640650ac511c)TileIteratorSimt & [operator+=](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a1a76d3b0fd2419e78d66640650ac511c)(TensorCoord const &tile_offset) {

163

164 add_tile_offset(tile_offset);

165

166return *this;

167 }

168

170CUTLASS_HOST_DEVICE

[171](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a90b52f1411169d8f31a9b336cbb7390b)void [store_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a90b52f1411169d8f31a9b336cbb7390b)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa1caeb6928f62ceabba7817c754b08b7) const &frag, [Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a05ff07efb8d79a930472de75ea30938a) pointer_offset) {

172 #if CUTLASS_SIMT_EPILOGUE_USE_SCALAR_STORES

173// de-vectorized stores

174using ScalarAccessType = AlignedArray<Element, 1>;

175 ScalarAccessType const *scalarFragPtr = reinterpret_cast<ScalarAccessType const *>(&frag);

176 ScalarAccessType *scalarPointer = reinterpret_cast<ScalarAccessType *>(pointer_);

177

178CUTLASS_PRAGMA_UNROLL

179for (int n = 0; n < Policy::kAccessesPerIteration; ++n) {

180CUTLASS_PRAGMA_UNROLL

181for (int s = 0; s < Policy::kElementsPerAccess; s++) {

182 scalarPointer[n * Policy::MmaSimtPolicy::WarpShape::kColumn * Policy::kElementsPerAccess + s] = scalarFragPtr[n * Policy::kElementsPerAccess + s];

183 }

184 }

185 #else

186// original vector stores

187AccessType const *frag_ptr = reinterpret_cast<AccessType const *>(&frag);

188CUTLASS_PRAGMA_UNROLL

189for (int n = 0; n < Policy::kAccessesPerIteration; ++n) {

190 pointer_[n * Policy::MmaSimtPolicy::WarpShape::kColumn] = frag_ptr[n];

191 }

192 #endif

193 }

194

196CUTLASS_HOST_DEVICE

[197](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ad3979de3f4e9abec2cbc7e8d8f41641c)void [store](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ad3979de3f4e9abec2cbc7e8d8f41641c)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa1caeb6928f62ceabba7817c754b08b7) const &frag) {

198 store_with_pointer_offset(frag, 0);

199 }

200

202CUTLASS_HOST_DEVICE

[203](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a91fa2b56deb9d9e4a5d58c0103198559)void [load_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a91fa2b56deb9d9e4a5d58c0103198559)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa1caeb6928f62ceabba7817c754b08b7) &frag, [Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a05ff07efb8d79a930472de75ea30938a) pointer_offset) const {

204

205AccessType *frag_ptr = reinterpret_cast<AccessType *>(&frag);

206

207CUTLASS_PRAGMA_UNROLL

208for (int n = 0; n < Policy::kAccessesPerIteration; ++n) {

209 frag_ptr[n] = pointer_[n * Policy::MmaSimtPolicy::WarpShape::kColumn];

210 }

211 }

212

214CUTLASS_HOST_DEVICE

[215](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#adc0257e19f8d1c6d3cfd8ef52d952f71)void [load](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#adc0257e19f8d1c6d3cfd8ef52d952f71)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa1caeb6928f62ceabba7817c754b08b7) &frag) const {

216 load_with_pointer_offset(frag, 0);

217 }

218 };

219

221

222 } // namespace warp

223 } // namespace epilogue

224 } // namespace cutlass

225

cutlass::MatrixShape

Describes the size of a matrix tile.

Definition: matrix_shape.h:42

cutlass::MatrixCoord::column

CUTLASS_HOST_DEVICE Index const & column() const

Returns the column of the coordinate.

Definition: matrix_coord.h:85

cutlass

Definition: aligned_buffer.h:35

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::AccumulatorTile](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a6c100257da1ab23db35da3c4818a32f3)

Array< typename Operator::ElementC, Policy::kAccumulatorElementCount > AccumulatorTile

This is the complete warp-level accumulator tile.

Definition: tile_iterator_simt.h:95

cutlass::AlignedArray

Aligned array type.

Definition: array.h:511

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aa1caeb6928f62ceabba7817c754b08b7)

Array< typename Operator::ElementC, Policy::kElementsPerIteration > Fragment

This is the fragment size produced by one access of the iterator.

Definition: tile_iterator_simt.h:90

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a05ff07efb8d79a930472de75ea30938a)

typename TensorRef::Index Index

Definition: tile_iterator_simt.h:76

cutlass::MatrixCoord::row

CUTLASS_HOST_DEVICE Index const & row() const

Returns the row of the coordinate.

Definition: matrix_coord.h:77

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::LongIndex](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a5c553a18b6215ed1158e863851be631a)

typename TensorRef::LongIndex LongIndex

Definition: tile_iterator_simt.h:77

cutlass::epilogue::warp::SimtPolicy

Definition: simt_policy.h:50

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::load](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#adc0257e19f8d1c6d3cfd8ef52d952f71)

CUTLASS_HOST_DEVICE void load(Fragment &frag) const

Load.

Definition: tile_iterator_simt.h:215

array.h

Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...

CUTLASS_PRAGMA_UNROLL

#define CUTLASS_PRAGMA_UNROLL

Definition: cutlass.h:110

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::store](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ad3979de3f4e9abec2cbc7e8d8f41641c)

CUTLASS_HOST_DEVICE void store(Fragment const &frag)

Store.

Definition: tile_iterator_simt.h:197

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::add_tile_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#ae9f188af3f314fb898e7ad842cba98df)

CUTLASS_HOST_DEVICE TileIteratorSimt & add_tile_offset(TensorCoord const &tile_offset)

advances in units of whole tiles along the logical coordinate space of the tensor ...

Definition: tile_iterator_simt.h:150

nullptr

#define nullptr

nullptr

Definition: platform.h:144

cutlass::TensorRef< Element, Layout >

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::store_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a90b52f1411169d8f31a9b336cbb7390b)

CUTLASS_HOST_DEVICE void store_with_pointer_offset(Fragment const &frag, Index pointer_offset)

Store.

Definition: tile_iterator_simt.h:171

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::add_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a59a4695befd919265485918ee935aab1)

CUTLASS_HOST_DEVICE TileIteratorSimt & add_pointer_offset(Index pointer_offset)

Adds a pointer offset.

Definition: tile_iterator_simt.h:143

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::TileIteratorSimt](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a25c74737e253375473171917c4f3df6f)

CUTLASS_HOST_DEVICE TileIteratorSimt(TensorRef const &ref, unsigned lane_id)

Constructor from TensorRef.

Definition: tile_iterator_simt.h:128

CUTLASS_HOST_DEVICE

#define CUTLASS_HOST_DEVICE

Definition: cutlass.h:89

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::TileIteratorSimt](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a0e0346d2b4a2e5b111a2d1e6d3ac775e)

CUTLASS_HOST_DEVICE TileIteratorSimt()

Default constructor.

Definition: tile_iterator_simt.h:124

cutlass::TensorRef::Index

typename Layout::Index Index

Index type.

Definition: tensor_ref.h:165

cutlass::layout::RowMajor

Mapping function for row-major matrices.

Definition: layout/matrix.h:50

matrix.h

Defines layout functions used by TensorRef and derived classes.

pitch_linear.h

Defines layout functions used by TensorRef and derived classes for pitch-linear memory.

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::operator+=](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a1a76d3b0fd2419e78d66640650ac511c)

CUTLASS_HOST_DEVICE TileIteratorSimt & operator+=(TensorCoord const &tile_offset)

Definition: tile_iterator_simt.h:162

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::load_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#a91fa2b56deb9d9e4a5d58c0103198559)

CUTLASS_HOST_DEVICE void load_with_pointer_offset(Fragment &frag, Index pointer_offset) const

Load.

Definition: tile_iterator_simt.h:203

simt_policy.h

Defines basic structures needed for implementing the warp-scoped phase of the epilogue. These quantities assume a 'column-major' arrangement of SimtOp instructions, of which a row-oriented slice is visible per iteration.

[cutlass::epilogue::warp::TileIteratorSimt< WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ >::WarpShape](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape 00_01Operator 00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html#aff99d6853af936cc23506b16b780ce09)

WarpShape_ WarpShape

Definition: tile_iterator_simt.h:69

cutlass::MatrixCoord

Definition: matrix_coord.h:39

cutlass::epilogue::warp::TileIteratorSimt

Template for reading and writing tiles of accumulators to shared memory.

Definition: tile_iterator_simt.h:55

cutlass::TensorRef::LongIndex

typename Layout::LongIndex LongIndex

Long index used for pointer offsets.

Definition: tensor_ref.h:168

Generated by 1.8.11