tile_iterator_volta_tensor_op.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without modification, are permitted
* provided that the following conditions are met:
* * Redistributions of source code must retain the above copyright notice, this list of
* conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright notice, this list of
* conditions and the following disclaimer in the documentation and/or other materials
* provided with the distribution.
* * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
* to endorse or promote products derived from this software without specific prior written
* permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
* IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
* FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
* STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
**************************************************************************************************/
#pragma once
#include "cutlass/array.h"
#include "cutlass/layout/matrix.h"
#include "cutlass/layout/pitch_linear.h"
#include "cutlass/epilogue/warp/tensor_op_policy.h"
namespace cutlass {
namespace epilogue {
namespace warp {
template <
typename WarpShape,
typename InterleavedTileShape,
typename ElementC,
typename Layout
>
struct TileIteratorVoltaTensorOp;
template <
typename WarpShape_
>
class TileIteratorVoltaTensorOp<WarpShape_, gemm::GemmShape<32, 32, 4>, half_t, layout::RowMajor> {
public:
using WarpShape = WarpShape_;
using InterleavedTileShape = gemm::GemmShape<32, 32, 4>;
using Element = half_t;
using Layout = layout::RowMajor;
using TensorRef = TensorRef<Element, Layout>;
using TensorCoord = MatrixCoord;
using Index = typename TensorRef::Index;
using LongIndex = typename TensorRef::LongIndex;
using Policy = VoltaTensorOpPolicy<WarpShape, InterleavedTileShape, Element, Layout>;
using Shape = MatrixShape<
Policy::kRowsPerIteration,
WarpShape::kN
>;
using AccessType = typename Policy::AccessType;
using Fragment = typename Policy::Fragment;
using AccumulatorTile = typename Policy::AccumulatorTile;
static int const kIterations = Policy::kIterations;
static int const kElementsPerAccess = Policy::kElementsPerAccess;
// Internal constants
struct Detail {
static int const kLanesInQuad = 4;
static int const kRowsPerQuad = 4;
static int const kColumnsPerQuad = 8;
static int const kAccessesPerQuad = kColumnsPerQuad / Policy::kElementsPerAccess;
static int const kAccessQuadDelta = 16;
};
using Padding = MatrixShape<
0,
Policy::kElementsPerAccess>;
private:
//
// Data members
//
AccessType *pointer_;
Layout layout_;
public:
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp(): pointer_(nullptr) { }
CUTLASS_DEVICE
TileIteratorVoltaTensorOp(
TensorRef const &ref,
unsigned lane_id
):
pointer_(reinterpret_cast<AccessType *>(ref.data())),
layout_(ref.stride()[0] / Policy::kElementsPerAccess) {
int quad_id = lane_id / Detail::kLanesInQuad;
int lane_in_quad = (lane_id % Detail::kLanesInQuad);
int quad_row_idx = ((quad_id & 4) >> 1) + (quad_id & 1);
int quad_col_idx = ((quad_id & 2) >> 1);
int row = quad_row_idx * Detail::kRowsPerQuad + lane_in_quad;
int column = quad_col_idx * Detail::kColumnsPerQuad;
pointer_ += layout_({row, column / kElementsPerAccess});
}
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp & add_pointer_offset(Index pointer_offset) {
pointer_ += pointer_offset / Policy::kElementsPerAccess;
return *this;
}
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp & add_tile_offset(TensorCoord const &tile_offset) {
pointer_ += layout_({
tile_offset.row() * Shape::kRow,
tile_offset.column() * Shape::kColumn / Policy::kElementsPerAccess});
return *this;
}
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp & operator+=(TensorCoord const &tile_offset) {
add_tile_offset(tile_offset);
return *this;
}
CUTLASS_DEVICE
void store_with_pointer_offset(Fragment const &frag, Index pointer_offset) {
AccessType const *frag_ptr = reinterpret_cast< AccessType const *>(&frag);
CUTLASS_PRAGMA_UNROLL
for (int tile_idx = 0; tile_idx < Policy::TileIterations::kColumn; ++tile_idx) {
CUTLASS_PRAGMA_UNROLL
for (int access_idx = 0; access_idx < Policy::kAccessesPerInterleavedTile; ++access_idx) {
int access_quad = access_idx / 2;
int access = access_idx % 2;
int ptr_offset = tile_idx * InterleavedTileShape::kN / Policy::kElementsPerAccess +
access_quad * Detail::kAccessQuadDelta / Policy::kElementsPerAccess + access;
int frag_idx = tile_idx * Policy::kAccessesPerInterleavedTile + access_idx;
AccessType access_vector = frag_ptr[frag_idx];
pointer_[ptr_offset] = access_vector;
}
}
}
CUTLASS_HOST_DEVICE
void store(Fragment const &frag) {
store_with_pointer_offset(frag, 0);
}
CUTLASS_HOST_DEVICE
void load_with_pointer_offset(Fragment const &frag, Index pointer_offset) {
AccessType *frag_ptr = reinterpret_cast< AccessType *>(&frag);
CUTLASS_PRAGMA_UNROLL
for (int tile_idx = 0; tile_idx < Policy::TileIterations::kColumn; ++tile_idx) {
CUTLASS_PRAGMA_UNROLL
for (int access_idx = 0; access_idx < Policy::kAccessesPerInterleavedTile; ++access_idx) {
int access_quad = access_idx / 2;
int access = access_idx % 2;
int ptr_offset = tile_idx * Detail::kTileDelta + access_quad * Detail::kAccessQuadDelta + access;
int frag_idx = tile_idx * Policy::kAccessesPerInterleavedTile + access_idx;
frag_ptr[frag_idx] = pointer_[ptr_offset];
}
}
}
CUTLASS_HOST_DEVICE
void load(Fragment const &frag) {
load_with_pointer_offset(frag, 0);
}
};
template <
typename WarpShape_
>
class TileIteratorVoltaTensorOp<WarpShape_, gemm::GemmShape<32, 32, 4>, float, layout::RowMajor> {
public:
using WarpShape = WarpShape_;
using InterleavedTileShape = gemm::GemmShape<32, 32, 4>;
using Element = float;
using Layout = layout::RowMajor;
using TensorRef = TensorRef<Element, Layout>;
using TensorCoord = MatrixCoord;
using Index = typename TensorRef::Index;
using LongIndex = typename TensorRef::LongIndex;
using Policy = VoltaTensorOpPolicy<WarpShape, InterleavedTileShape, Element, Layout>;
using Shape = MatrixShape<
Policy::kRowsPerIteration,
WarpShape::kN
>;
using AccessType = typename Policy::AccessType;
using Fragment = typename Policy::Fragment;
using AccumulatorTile = typename Policy::AccumulatorTile;
static int const kIterations = Policy::kIterations;
static int const kElementsPerAccess = Policy::kElementsPerAccess;
// Internal constants
struct Detail {
static int const kLanesInQuad = 4;
static int const kRowsPerQuad = 4;
static int const kColumnsPerQuad = 8;
static int const kAccessesPerQuad = kColumnsPerQuad / Policy::kElementsPerAccess;
static int const kAccessQuadDelta = 16;
};
using Padding = MatrixShape<
0,
Policy::kElementsPerAccess>;
private:
//
// Data members
//
AccessType *pointer_;
Layout layout_;
public:
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp(): pointer_(nullptr) { }
CUTLASS_DEVICE
TileIteratorVoltaTensorOp(
TensorRef const &ref,
unsigned lane_id
):
pointer_(reinterpret_cast<AccessType *>(ref.data())),
layout_(ref.stride()[0] / Policy::kElementsPerAccess) {
int quad_id = lane_id / Detail::kLanesInQuad;
int lane_in_quad = (lane_id % Detail::kLanesInQuad);
int const kQuadRowDelta = 4;
int const kQuadColumnDelta = 2 * Policy::MmaIterations::kColumn;
int quad_row_offset = ((quad_id & 4) / 2 + (quad_id & 1)) * kQuadRowDelta;
int quad_column_offset = (quad_id & 2) / 2 * kQuadColumnDelta;
int thread_row_offset = (lane_in_quad & 1);
int thread_column_offset = (lane_in_quad & 2) / 2;
int row = quad_row_offset + thread_row_offset;
int column = quad_column_offset + thread_column_offset;
pointer_ += layout_({row, column});
}
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp & add_pointer_offset(Index pointer_offset) {
pointer_ += pointer_offset / Policy::kElementsPerAccess;
return *this;
}
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp & add_tile_offset(TensorCoord const &tile_offset) {
pointer_ += layout_({
tile_offset.row() * Shape::kRow,
tile_offset.column() * Shape::kColumn / Policy::kElementsPerAccess});
return *this;
}
CUTLASS_HOST_DEVICE
TileIteratorVoltaTensorOp & operator+=(TensorCoord const &tile_offset) {
add_tile_offset(tile_offset);
return *this;
}
CUTLASS_DEVICE
void store_with_pointer_offset(Fragment const &frag, Index pointer_offset) {
AccessType const *frag_ptr = reinterpret_cast< AccessType const *>(&frag);
int const kAccessesPerRow = Policy::TileIterations::kColumn * Policy::MmaIterations::kColumn * 2;
CUTLASS_PRAGMA_UNROLL
for (int row_idx = 0; row_idx < Policy::kRowsPerMmaTile; ++row_idx) {
CUTLASS_PRAGMA_UNROLL
for (int access_idx = 0; access_idx < kAccessesPerRow; ++access_idx) {
int frag_idx = row_idx * kAccessesPerRow + access_idx;
int ptr_column_offset = (access_idx & 1) * 2 +
(access_idx & 2) * Policy::MmaIterations::kColumn * 2 +
(access_idx & 4) * Policy::MmaIterations::kColumn * 2;
int ptr_row_offset = row_idx * 2;
int ptr_offset = layout_({ptr_row_offset, ptr_column_offset});
pointer_[ptr_offset] = frag_ptr[frag_idx];
}
}
}
CUTLASS_HOST_DEVICE
void store(Fragment const &frag) {
store_with_pointer_offset(frag, 0);
}
CUTLASS_HOST_DEVICE
void load_with_pointer_offset(Fragment const &frag, Index pointer_offset) {
AccessType *frag_ptr = reinterpret_cast< AccessType *>(&frag);
assert(0); // TODO
}
CUTLASS_HOST_DEVICE
void load(Fragment const &frag) {
load_with_pointer_offset(frag, 0);
}
};
} // namespace warp
} // namespace epilogue
} // namespace cutlass
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::load_with_pointer_offset
CUTLASS_HOST_DEVICE void load_with_pointer_offset(Fragment const &frag, Index pointer_offset)
Load.
Definition: tile_iterator_volta_tensor_op.h:400
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::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::Index
typename TensorRef::Index Index
Definition: tile_iterator_volta_tensor_op.h:70
cutlass
Definition: aligned_buffer.h:35
tensor_op_policy.h
Defines basic structures needed for implementing the warp-scoped phase of the epilogue. These quantities assume a &#39;column-major&#39; arrangement of TensorOp instructions, of which a row-oriented slice is visible per iteration.
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::load
CUTLASS_HOST_DEVICE void load(Fragment const &frag)
Load.
Definition: tile_iterator_volta_tensor_op.h:232
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::store_with_pointer_offset
CUTLASS_DEVICE void store_with_pointer_offset(Fragment const &frag, Index pointer_offset)
Store.
Definition: tile_iterator_volta_tensor_op.h:365
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::LongIndex
typename TensorRef::LongIndex LongIndex
Definition: tile_iterator_volta_tensor_op.h:71
cutlass::half_t
IEEE half-precision floating-point type.
Definition: half.h:126
cutlass::MatrixCoord::row
CUTLASS_HOST_DEVICE Index const & row() const
Returns the row of the coordinate.
Definition: matrix_coord.h:77
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::AccessType
typename Policy::AccessType AccessType
Array type for aligned memory accesses.
Definition: tile_iterator_volta_tensor_op.h:265
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::add_tile_offset
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp & add_tile_offset(TensorCoord const &tile_offset)
advances in units of whole tiles along the logical coordinate space of the tensor ...
Definition: tile_iterator_volta_tensor_op.h:158
array.h
Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::WarpShape
WarpShape_ WarpShape
Definition: tile_iterator_volta_tensor_op.h:63
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: cutlass.h:110
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::operator+=
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp & operator+=(TensorCoord const &tile_offset)
Definition: tile_iterator_volta_tensor_op.h:358
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::load
CUTLASS_HOST_DEVICE void load(Fragment const &frag)
Load.
Definition: tile_iterator_volta_tensor_op.h:409
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::operator+=
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp & operator+=(TensorCoord const &tile_offset)
Definition: tile_iterator_volta_tensor_op.h:169
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::Fragment
typename Policy::Fragment Fragment
This is the fragment size produced by one access of the iterator.
Definition: tile_iterator_volta_tensor_op.h:268
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::add_pointer_offset
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp & add_pointer_offset(Index pointer_offset)
Adds a pointer offset.
Definition: tile_iterator_volta_tensor_op.h:151
cutlass::epilogue::warp::TileIteratorVoltaTensorOp
Template for reading and writing tiles of accumulators to shared memory.
Definition: tile_iterator_volta_tensor_op.h:52
nullptr
#define nullptr
nullptr
Definition: platform.h:144
cutlass::TensorRef< Element, Layout >
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::TileIteratorVoltaTensorOp
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp()
Default constructor.
Definition: tile_iterator_volta_tensor_op.h:309
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::store
CUTLASS_HOST_DEVICE void store(Fragment const &frag)
Store.
Definition: tile_iterator_volta_tensor_op.h:203
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::add_tile_offset
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp & add_tile_offset(TensorCoord const &tile_offset)
advances in units of whole tiles along the logical coordinate space of the tensor ...
Definition: tile_iterator_volta_tensor_op.h:347
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::store_with_pointer_offset
CUTLASS_DEVICE void store_with_pointer_offset(Fragment const &frag, Index pointer_offset)
Store.
Definition: tile_iterator_volta_tensor_op.h:176
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::AccumulatorTile
typename Policy::AccumulatorTile AccumulatorTile
This is the complete warp-level accumulator tile.
Definition: tile_iterator_volta_tensor_op.h:271
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::gemm::GemmShape
Shape of a matrix multiply-add operation.
Definition: include/cutlass/gemm/gemm.h:57
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::Element
float Element
Definition: tile_iterator_volta_tensor_op.h:248
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::WarpShape
WarpShape_ WarpShape
Definition: tile_iterator_volta_tensor_op.h:246
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::AccessType
typename Policy::AccessType AccessType
Array type for aligned memory accesses.
Definition: tile_iterator_volta_tensor_op.h:82
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::LongIndex
typename TensorRef::LongIndex LongIndex
Definition: tile_iterator_volta_tensor_op.h:254
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::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::TileIteratorVoltaTensorOp
CUTLASS_DEVICE TileIteratorVoltaTensorOp(TensorRef const &ref, unsigned lane_id)
Constructor from TensorRef.
Definition: tile_iterator_volta_tensor_op.h:130
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::Fragment
typename Policy::Fragment Fragment
This is the fragment size produced by one access of the iterator.
Definition: tile_iterator_volta_tensor_op.h:85
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::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::Index
typename TensorRef::Index Index
Definition: tile_iterator_volta_tensor_op.h:253
cutlass::epilogue::warp::VoltaTensorOpPolicy
Policy details related to the epilogue.
Definition: volta_tensor_op_policy.h:52
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::store
CUTLASS_HOST_DEVICE void store(Fragment const &frag)
Store.
Definition: tile_iterator_volta_tensor_op.h:394
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::load_with_pointer_offset
CUTLASS_HOST_DEVICE void load_with_pointer_offset(Fragment const &frag, Index pointer_offset)
Load.
Definition: tile_iterator_volta_tensor_op.h:209
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::add_pointer_offset
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp & add_pointer_offset(Index pointer_offset)
Adds a pointer offset.
Definition: tile_iterator_volta_tensor_op.h:340
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, float, layout::RowMajor >::TileIteratorVoltaTensorOp
CUTLASS_DEVICE TileIteratorVoltaTensorOp(TensorRef const &ref, unsigned lane_id)
Constructor from TensorRef.
Definition: tile_iterator_volta_tensor_op.h:313
cutlass::MatrixCoord
Definition: matrix_coord.h:39
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::AccumulatorTile
typename Policy::AccumulatorTile AccumulatorTile
This is the complete warp-level accumulator tile.
Definition: tile_iterator_volta_tensor_op.h:88
cutlass::epilogue::warp::TileIteratorVoltaTensorOp< WarpShape_, gemm::GemmShape< 32, 32, 4 >, half_t, layout::RowMajor >::TileIteratorVoltaTensorOp
CUTLASS_HOST_DEVICE TileIteratorVoltaTensorOp()
Default constructor.
Definition: tile_iterator_volta_tensor_op.h:126
cutlass::TensorRef::LongIndex
typename Layout::LongIndex LongIndex
Long index used for pointer offsets.
Definition: tensor_ref.h:168