ContextQMD
Back to Cutlass

CUTLASS: tile_iterator_tensor_op.h Source File

docs/tile__iterator__tensor__op_8h_source.html

4.4.225.2 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

tile_iterator_tensor_op.h

[Go to the documentation of this file.](tile iterator tensor__op_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/tensor_op_policy.h](tensor op policy_8h.html)"

36

38

39 namespace cutlass {

40 namespace epilogue {

41 namespace warp {

42

44

46 template <

47typename WarpShape,

48typename OperatorShape,

49typename Element,

50typename Layout

51 >

52 class TileIteratorTensorOp;

53

55

57 template <

58typename WarpShape_,

59typename OperatorShape_,

60typename Element_

61 >

[62](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html) class TileIteratorTensorOp<WarpShape_, OperatorShape_, Element_, layout::RowMajor> {

63 public:

64

[65](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a587ef97d446f1857f1691189fac8374f)using [WarpShape](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a587ef97d446f1857f1691189fac8374f) = WarpShape_;

[66](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a90059f91fd1200e119441c83787b0fa2)using OperatorShape = OperatorShape_;

[67](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a98c66293f80489be1140fc19e15eaecb)using Element = Element_;

[68](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aeaef96491169d1e8f2e831a3d858382f)using Layout = layout::RowMajor;

69

[70](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a7151807a9143cfed1359ff6b186df6f6)using TensorRef = TensorRef<Element, Layout>;

[71](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a50ac8a7aa13d124f37d89b11f4d10e95)using TensorCoord = MatrixCoord;

[72](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a551137bff490b3b6acfd2f0685a81da5)using [Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a551137bff490b3b6acfd2f0685a81da5) = typename TensorRef::Index;

[73](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a38473d6ddc3a0eab3fde84840611e2d4)using [LongIndex](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a38473d6ddc3a0eab3fde84840611e2d4) = typename TensorRef::LongIndex;

74

[75](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aa43845436513f3eec39906a41a275953)using Policy = TensorOpPolicy<WarpShape, OperatorShape, Layout>;

76

78using Shape = MatrixShape<

79 Policy::kRowsPerIteration,

80 WarpShape::kN

[81](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a92e42402b1ee46fda6f6ded6825b7aed) >;

82

84using [Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a9156203bccdaf3b36b87286153c63147) = Array<

85 Element,

[86](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a9156203bccdaf3b36b87286153c63147) Policy::OperatorCount::kColumn * Policy::kElementsPerAccess>;

87

89//using AccumulatorTile = typename Operator::FragmentC;

90

[92](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a5f4cbb21e1c17bc6e6f7938415b53ee8)static int const kIterations = Policy::kIterations;

93

94// Internal constants

[95](structcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 05f11e023c9e6ee5f7a888fa4c5bbf6d1.html)struct Detail {

[96](structcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 05f11e023c9e6ee5f7a888fa4c5bbf6d1.html#a7848a076b84109151f08303706503063)static int const kLanesInQuad = 4;

97 };

98

100using Padding = MatrixShape<

101 0,

[102](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a96dd094804882c2103e1a457632cf182) Detail::kLanesInQuad * Policy::kElementsPerAccess>;

103

104 private:

105

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

108

109//

110// Data members

111//

112

114AccessType *pointer_;

115

117Layout layout_;

118

119 public:

120

122CUTLASS_HOST_DEVICE

[123](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a40db089c95d2e6aed4a652862bc09f32)[TileIteratorTensorOp](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a40db089c95d2e6aed4a652862bc09f32)(): pointer_(nullptr) { }

124

126CUTLASS_HOST_DEVICE

[127](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#ac3bd8284bb551d89dbb9c639654a06ee)[TileIteratorTensorOp](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#ac3bd8284bb551d89dbb9c639654a06ee)(

128TensorRef const &ref,

129unsigned lane_id

130 ):

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

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

133

134int quad_id = (lane_id / Detail::kLanesInQuad);

135int lane_in_quad = (lane_id % Detail::kLanesInQuad);

136

137 pointer_ += layout_({quad_id, lane_in_quad});

138 }

139

141CUTLASS_HOST_DEVICE

[142](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aeda0dac956a396df42a2d6dedf5da3f5)TileIteratorTensorOp & [add_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aeda0dac956a396df42a2d6dedf5da3f5)([Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a551137bff490b3b6acfd2f0685a81da5) pointer_offset) {

143 pointer_ += pointer_offset / Policy::kElementsPerAccess;

144return *this;

145 }

146

148CUTLASS_HOST_DEVICE

[149](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aa54bfe6b9c53d0e79cfbae74c0e52fe4)TileIteratorTensorOp & [add_tile_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aa54bfe6b9c53d0e79cfbae74c0e52fe4)(TensorCoord const &tile_offset) {

150

151 pointer_ += layout_({

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

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

154 });

155

156return *this;

157 }

158

160CUTLASS_HOST_DEVICE

[161](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a46b0395432768c15516edf5d4ce5af73)TileIteratorTensorOp & [operator+=](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a46b0395432768c15516edf5d4ce5af73)(TensorCoord const &tile_offset) {

162 add_tile_offset(tile_offset);

163return *this;

164 }

165

167CUTLASS_HOST_DEVICE

[168](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#ac47720c42f8242c6350c0c645a598c08)void [store_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#ac47720c42f8242c6350c0c645a598c08)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a9156203bccdaf3b36b87286153c63147) const &frag, [Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a551137bff490b3b6acfd2f0685a81da5) pointer_offset) {

169

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

171

172CUTLASS_PRAGMA_UNROLL

173for (int n = 0; n < Policy::OperatorCount::kColumn; ++n) {

174 pointer_[n * Detail::kLanesInQuad + pointer_offset / Policy::kElementsPerAccess] = frag_ptr[n];

175 }

176 }

177

179CUTLASS_HOST_DEVICE

[180](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#acf1c8f751d72ce97b2e6f94633c8fdd6)void [store](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#acf1c8f751d72ce97b2e6f94633c8fdd6)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a9156203bccdaf3b36b87286153c63147) const &frag) {

181 store_with_pointer_offset(frag, 0);

182 }

183

185CUTLASS_HOST_DEVICE

[186](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a3b35869528032ca52eb9e37d61265209)void [load_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a3b35869528032ca52eb9e37d61265209)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a9156203bccdaf3b36b87286153c63147) &frag, [Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a551137bff490b3b6acfd2f0685a81da5) pointer_offset) const {

187

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

189

190CUTLASS_PRAGMA_UNROLL

191for (int n = 0; n < Policy::OperatorCount::kColumn; ++n) {

192 frag_ptr[n] = pointer_[n * Detail::kLanesInQuad + pointer_offset / Policy::kElementsPerAccess];

193 }

194 }

195

197CUTLASS_HOST_DEVICE

[198](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a4e473670fd9c4fc7004db274ba89c9c5)void [load](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a4e473670fd9c4fc7004db274ba89c9c5)([Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a9156203bccdaf3b36b87286153c63147) &frag) const {

199 load_with_pointer_offset(frag, 0);

200 }

201 };

202

204

205 } // namespace warp

206 } // namespace epilogue

207 } // namespace cutlass

208

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

[tensor_op_policy.h](tensor op policy_8h.html)

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

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::load_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a3b35869528032ca52eb9e37d61265209)

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

Load.

Definition: tile_iterator_tensor_op.h:186

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::TileIteratorTensorOp](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#ac3bd8284bb551d89dbb9c639654a06ee)

CUTLASS_HOST_DEVICE TileIteratorTensorOp(TensorRef const &ref, unsigned lane_id)

Constructor from TensorRef.

Definition: tile_iterator_tensor_op.h:127

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::WarpShape](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a587ef97d446f1857f1691189fac8374f)

WarpShape_ WarpShape

Definition: tile_iterator_tensor_op.h:65

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::add_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aeda0dac956a396df42a2d6dedf5da3f5)

CUTLASS_HOST_DEVICE TileIteratorTensorOp & add_pointer_offset(Index pointer_offset)

Adds a pointer offset.

Definition: tile_iterator_tensor_op.h:142

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::add_tile_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#aa54bfe6b9c53d0e79cfbae74c0e52fe4)

CUTLASS_HOST_DEVICE TileIteratorTensorOp & add_tile_offset(TensorCoord const &tile_offset)

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

Definition: tile_iterator_tensor_op.h:149

cutlass::AlignedArray

Aligned array type.

Definition: array.h:511

cutlass::MatrixCoord::row

CUTLASS_HOST_DEVICE Index const & row() const

Returns the row of the coordinate.

Definition: matrix_coord.h:77

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::store_with_pointer_offset](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#ac47720c42f8242c6350c0c645a598c08)

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

Store.

Definition: tile_iterator_tensor_op.h:168

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::operator+=](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a46b0395432768c15516edf5d4ce5af73)

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

Definition: tile_iterator_tensor_op.h:161

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::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::store](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#acf1c8f751d72ce97b2e6f94633c8fdd6)

CUTLASS_HOST_DEVICE void store(Fragment const &frag)

Store.

Definition: tile_iterator_tensor_op.h:180

nullptr

#define nullptr

nullptr

Definition: platform.h:144

cutlass::epilogue::warp::TensorOpPolicy

Policy details related to the epilogue.

Definition: tensor_op_policy.h:50

cutlass::TensorRef< Element, Layout >

CUTLASS_HOST_DEVICE

#define CUTLASS_HOST_DEVICE

Definition: cutlass.h:89

cutlass::epilogue::warp::TileIteratorTensorOp

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

Definition: tile_iterator_tensor_op.h:52

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::LongIndex](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a38473d6ddc3a0eab3fde84840611e2d4)

typename TensorRef::LongIndex LongIndex

Definition: tile_iterator_tensor_op.h:73

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Index](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a551137bff490b3b6acfd2f0685a81da5)

typename TensorRef::Index Index

Definition: tile_iterator_tensor_op.h:72

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

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::TileIteratorTensorOp](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a40db089c95d2e6aed4a652862bc09f32)

CUTLASS_HOST_DEVICE TileIteratorTensorOp()

Default constructor.

Definition: tile_iterator_tensor_op.h:123

[cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Fragment](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a9156203bccdaf3b36b87286153c63147)

Array< Element, Policy::OperatorCount::kColumn *Policy::kElementsPerAccess > Fragment

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

Definition: tile_iterator_tensor_op.h:86

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::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::load](classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape 00_01OperatorShape 003cbb32beb84b4984cb7853662096d289.html#a4e473670fd9c4fc7004db274ba89c9c5)

CUTLASS_HOST_DEVICE void load(Fragment &frag) const

Load.

Definition: tile_iterator_tensor_op.h:198

cutlass::MatrixCoord

Definition: matrix_coord.h:39

cutlass::TensorRef::LongIndex

typename Layout::LongIndex LongIndex

Long index used for pointer offsets.

Definition: tensor_ref.h:168

Generated by 1.8.11