layout/matrix.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/cutlass.h"
#include "cutlass/matrix_coord.h"
#include "cutlass/matrix_traits.h"
namespace cutlass {
namespace layout {
//
// Defines data layouts of various matrix formats usable by TensorRef and other classes.
//
class RowMajor {
public:
static int const kRank = 2;
static int const kStrideRank = 1;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
private:
//
// Data members
//
Stride stride_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
RowMajor(Index ldm = 0): stride_(ldm) { }
CUTLASS_HOST_DEVICE
RowMajor(Stride stride): stride_(stride) { }
CUTLASS_HOST_DEVICE
static RowMajor packed(MatrixCoord const &extent) {
return RowMajor(extent.column());
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
return coord.row() * stride_[0] + coord.column();
}
CUTLASS_HOST_DEVICE
MatrixCoord inverse(LongIndex offset) const {
return MatrixCoord(Index(offset / stride_[0]), Index(offset % stride_[0]));
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
return extent.row() * stride_[0];
}
};
class ColumnMajor {
public:
static int const kRank = 2;
static int const kStrideRank = 1;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
private:
//
// Data members
//
Stride stride_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
ColumnMajor(Index ldm = 0): stride_(ldm) { }
CUTLASS_HOST_DEVICE
ColumnMajor(Stride stride): stride_(stride) { }
CUTLASS_HOST_DEVICE
static ColumnMajor packed(MatrixCoord const &extent) {
return ColumnMajor(extent.row());
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
return coord.row() + coord.column() * stride_[0];
}
CUTLASS_HOST_DEVICE
MatrixCoord inverse(LongIndex offset) const {
return MatrixCoord(Index(offset % stride_[0]), Index(offset / stride_[0]));
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
return extent.column() * stride_[0];
}
};
template <int Interleave>
struct RowMajorInterleaved {
static int const kRank = 2;
static int const kStrideRank = 1;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
static int const kInterleave = Interleave;
private:
//
// Data members
//
Stride stride_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
RowMajorInterleaved(Index ldm = 0): stride_(ldm) { }
CUTLASS_HOST_DEVICE
RowMajorInterleaved(Stride stride): stride_(stride) { }
CUTLASS_HOST_DEVICE
static RowMajorInterleaved packed(MatrixCoord const &extent) {
return RowMajorInterleaved(extent.column() * kInterleave);
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
Index row_major = coord.row() / kInterleave;
Index row_minor = coord.row() % kInterleave;
return row_major * stride_[0] + coord.column() * kInterleave + row_minor;
}
CUTLASS_HOST_DEVICE
MatrixCoord inverse(LongIndex offset) const {
Index row_major = Index(offset / stride_[0]);
Index residual = Index(offset % stride_[0]);
Index column = residual / kInterleave;
Index row_minor = residual % kInterleave;
return MatrixCoord(row_major * kInterleave + row_minor, column);
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
return (extent.row() + kInterleave - 1) / kInterleave * stride_[0];
}
};
template <int Interleave>
struct ColumnMajorInterleaved {
static int const kRank = 2;
static int const kStrideRank = 1;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
static int const kInterleave = Interleave;
private:
//
// Data members
//
Stride stride_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
ColumnMajorInterleaved(Index ldm = 0): stride_(ldm) { }
CUTLASS_HOST_DEVICE
ColumnMajorInterleaved(Stride stride): stride_(stride) { }
CUTLASS_HOST_DEVICE
static ColumnMajorInterleaved packed(MatrixCoord const &extent) {
return ColumnMajorInterleaved(extent.row() * kInterleave);
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
Index column_major = coord.column() / kInterleave;
Index column_minor = coord.column() % kInterleave;
return column_major * stride_[0] + coord.row() * kInterleave + column_minor;
}
CUTLASS_HOST_DEVICE
MatrixCoord inverse(LongIndex offset) const {
Index column_major = Index(offset / stride_[0]);
Index residual = Index(offset % stride_[0]);
Index row = residual / kInterleave;
Index column_minor = residual % kInterleave;
return MatrixCoord(row, column_major * kInterleave + column_minor);
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
return (extent.column() + kInterleave - 1) / kInterleave * stride_[0];
}
};
enum class Matrix {
kColumnMajor,
kRowMajor
};
struct ContiguousMatrix {
static int const kRank = 2;
static int const kStrideRank = 1;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
private:
//
// Data members
//
Stride stride_;
Matrix layout_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
ContiguousMatrix(
Index ldm = 0,
Matrix layout = Matrix::kColumnMajor
):
stride_(ldm), layout_(layout) { }
CUTLASS_HOST_DEVICE
static ContiguousMatrix packed(
MatrixCoord const &extent,
Matrix layout = Matrix::kColumnMajor) {
Index ldm = 0;
if (layout == Matrix::kColumnMajor) {
ldm = extent.row();
}
else if (layout == Matrix::kRowMajor) {
ldm = extent.column();
}
return ContiguousMatrix(ldm, layout);
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
if (layout_ == Matrix::kColumnMajor) {
return coord.row() + coord.column() * stride_[0];
}
else if (layout_ == Matrix::kRowMajor) {
return coord.row() * stride_[0] + coord.column();
}
else {
// degenerate case
return 0;
}
}
CUTLASS_HOST_DEVICE
MatrixCoord inverse(LongIndex offset) const {
// TODO
return MatrixCoord(0, 0);
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
if (layout_ == Matrix::kColumnMajor) {
return stride_[0] * extent.column();
}
else if (layout_ == Matrix::kRowMajor) {
return stride_[0] * extent.row();
}
else {
// degenerate case
return 0;
}
}
};
template <int BlockRows, int BlockColumns>
struct ColumnMajorBlockLinear {
static int const kRank = 2;
static int const kStrideRank = 1;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
static int const kBlockRows = BlockRows;
static int const kBlockColumns = BlockColumns;
private:
//
// Data members
//
Stride stride_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
ColumnMajorBlockLinear(Index ldm = 0): stride_(ldm) { }
CUTLASS_HOST_DEVICE
static ColumnMajorBlockLinear packed(MatrixCoord const &extent) {
return ColumnMajorBlockLinear(extent.row() * kBlockRows * kBlockColumns);
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
return
(coord.row() % kBlockRows) +
(coord.column() % kBlockColumns) * kBlockRows +
(coord.row() / kBlockRows) * kBlockRows * kBlockColumns +
(coord.column() / kBlockColumns) * stride_[0];
}
CUTLASS_HOST_DEVICE
MatrixCoord inverse(LongIndex offset) const {
// TODO
return MatrixCoord(0, 0);
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
return (extent.column() + kBlockColumns - 1) / kBlockColumns * stride_[0];
}
};
template <int BlockRows, int BlockColumns>
struct RowMajorBlockLinear {
static int const kRank = 2;
static int const kStrideRank = 1;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
static int const kBlockRows = BlockRows;
static int const kBlockColumns = BlockColumns;
private:
//
// Data members
//
Stride stride_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
RowMajorBlockLinear(Index ldm = 0): stride_(ldm) { }
CUTLASS_HOST_DEVICE
static RowMajorBlockLinear packed(MatrixCoord const &extent) {
return RowMajorBlockLinear(extent.column() * kBlockRows * kBlockColumns);
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
return
(coord.column() % kBlockColumns) +
(coord.row() % kBlockRows) * kBlockColumns +
(coord.column() / kBlockColumns) * kBlockRows * kBlockColumns +
(coord.row() / kBlockRows) * stride_[0];
}
CUTLASS_HOST_DEVICE
MatrixCoord inverse(LongIndex offset) const {
// TODO
return MatrixCoord(0, 0);
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
return (extent.row() + kBlockRows - 1) / kBlockRows * stride_[0];
}
};
struct GeneralMatrix {
static int const kRank = 2;
static int const kStrideRank = 2;
using Index = int32_t;
using LongIndex = int64_t;
using TensorCoord = MatrixCoord;
using Stride = Coord<kStrideRank, Index>;
private:
//
// Data members
//
MatrixLayout layout_id_;
Stride stride_;
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
GeneralMatrix(): layout_id_(MatrixLayout::kColumnMajor), stride_(make_Coord(0, 1)) { }
CUTLASS_HOST_DEVICE
GeneralMatrix(
MatrixLayout layout_id,
Index ldm,
Index interleave): layout_id_(layout_id), stride_(make_Coord(ldm, interleave)) { }
CUTLASS_HOST_DEVICE
static GeneralMatrix packed(
MatrixCoord const &extent,
MatrixLayout layout_id = MatrixLayout::kColumnMajor,
Index interleave = 1) {
Index c;
if (layout_id == MatrixLayout::kRowMajor) {
c = extent.column();
}
else {
c = extent.row();
}
Index ldm = c * interleave;
return GeneralMatrix(layout_id, ldm, interleave);
}
CUTLASS_HOST_DEVICE
LongIndex operator()(MatrixCoord const &coord) const {
Index c, s;
if (layout_id_ == MatrixLayout::kRowMajor) {
c = coord.column();
s = coord.row();
}
else {
s = coord.column();
c = coord.row();
}
Index v = s / stride_[1];
Index residual = (s % stride_[1]);
return LongIndex(c) * LongIndex(stride_[1]) + LongIndex(v) * LongIndex(stride_[0]) + residual;
}
CUTLASS_HOST_DEVICE
Stride stride() const {
return stride_;
}
CUTLASS_HOST_DEVICE
MatrixLayout layout_id() const {
return layout_id_;
}
CUTLASS_HOST_DEVICE
Stride & stride() {
return stride_;
}
CUTLASS_HOST_DEVICE
MatrixLayout & layout_id() {
return layout_id_;
}
CUTLASS_HOST_DEVICE
Index stride(int idx) const {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
Index & stride(int idx) {
return stride_[idx];
}
CUTLASS_HOST_DEVICE
LongIndex capacity(MatrixCoord const &extent) const {
Index s;
if (layout_id_ == MatrixLayout::kRowMajor) {
s = extent.row();
}
else {
s = extent.column();
}
Index v = Index((s + stride_[1] - 1) / stride_[1]);
return LongIndex(v) * LongIndex(stride_[0]);
}
};
template <typename Layout>
struct LayoutTranspose;
template <>
struct LayoutTranspose<layout::RowMajor> {
using type = layout::ColumnMajor;
};
template <>
struct LayoutTranspose<layout::ColumnMajor> {
using type = layout::RowMajor;
};
} // namespace layout
} // namespace cutlass
cutlass::layout::ContiguousMatrix::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:464
cutlass::layout::RowMajorInterleaved::packed
static CUTLASS_HOST_DEVICE RowMajorInterleaved packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:283
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:323
cutlass::layout::RowMajor::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:62
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:668
cutlass::layout::ColumnMajorInterleaved< 4 >::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:355
cutlass::MatrixCoord::column
CUTLASS_HOST_DEVICE Index const & column() const
Returns the column of the coordinate.
Definition: matrix_coord.h:85
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:540
cutlass::layout::RowMajorBlockLinear::packed
static CUTLASS_HOST_DEVICE RowMajorBlockLinear packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:726
cutlass
Definition: aligned_buffer.h:35
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:750
cutlass::layout::Matrix::kColumnMajor
leading dimension refers to stride between columns; stride along rows is 1
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:311
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:768
cutlass::layout::ColumnMajorInterleaved::ColumnMajorInterleaved
CUTLASS_HOST_DEVICE ColumnMajorInterleaved(Index ldm=0)
Ctor.
Definition: layout/matrix.h:381
cutlass::layout::GeneralMatrix::GeneralMatrix
CUTLASS_HOST_DEVICE GeneralMatrix(MatrixLayout layout_id, Index ldm, Index interleave)
Ctor.
Definition: layout/matrix.h:822
cutlass::layout::ColumnMajorBlockLinear::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:592
cutlass::layout::ContiguousMatrix
Definition: layout/matrix.h:455
cutlass::layout::GeneralMatrix
Definition: layout/matrix.h:781
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:662
cutlass::layout::ColumnMajorInterleaved::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:442
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:130
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:124
cutlass::layout::ContiguousMatrix::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:467
cutlass::layout::RowMajorInterleaved< 4 >::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:249
cutlass::layout::ColumnMajorInterleaved::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:397
cutlass::layout::GeneralMatrix::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:793
cutlass::layout::RowMajorBlockLinear::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:693
cutlass::layout::RowMajor::kRank
static int const kRank
Logical rank of tensor.
Definition: layout/matrix.h:53
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:211
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:329
cutlass::layout::ColumnMajorInterleaved< 4 >::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:352
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:558
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:387
cutlass::layout::ColumnMajor::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: layout/matrix.h:154
cutlass::layout::ColumnMajor::ColumnMajor
CUTLASS_HOST_DEVICE ColumnMajor(Index ldm=0)
Ctor.
Definition: layout/matrix.h:177
cutlass::layout::RowMajorInterleaved::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:298
cutlass::layout::RowMajorBlockLinear::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:733
cutlass::layout::RowMajorInterleaved< 4 >::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:246
cutlass::MatrixCoord::row
CUTLASS_HOST_DEVICE Index const & row() const
Returns the row of the coordinate.
Definition: matrix_coord.h:77
cutlass::layout::RowMajor::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:136
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:418
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:112
cutlass::layout::RowMajorInterleaved::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:335
cutlass::layout::RowMajor::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:106
cutlass::layout::ColumnMajorBlockLinear
Definition: layout/matrix.h:581
cutlass::layout::ColumnMajor
Mapping function for column-major matrices.
Definition: layout/matrix.h:142
cutlass::layout::GeneralMatrix::layout_id
CUTLASS_HOST_DEVICE MatrixLayout & layout_id()
Definition: layout/matrix.h:885
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:762
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:891
cutlass::layout::GeneralMatrix::GeneralMatrix
CUTLASS_HOST_DEVICE GeneralMatrix()
Ctor.
Definition: layout/matrix.h:818
cutlass::layout::RowMajor::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:59
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:650
cutlass::layout::ColumnMajorBlockLinear::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:642
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:897
cutlass::layout::RowMajor::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:100
cutlass::layout::GeneralMatrix::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:790
cutlass::layout::LayoutTranspose
Defines transposes of matrix layouts.
Definition: layout/matrix.h:921
cutlass::layout::ColumnMajorBlockLinear::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:632
cutlass::layout::ContiguousMatrix::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:564
cutlass::layout::ColumnMajorBlockLinear::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:589
cutlass::layout::Matrix
Matrix
Enumerated type for canonical pitch-linear matrix layouts.
Definition: layout/matrix.h:448
cutlass::layout::ContiguousMatrix::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:518
cutlass::layout::ColumnMajor::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:199
cutlass::layout::RowMajorBlockLinear
Definition: layout/matrix.h:682
cutlass::MatrixLayout
MatrixLayout
Definition: matrix_traits.h:36
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:436
cutlass::layout::RowMajorBlockLinear::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:690
cutlass::layout::RowMajorInterleaved::RowMajorInterleaved
CUTLASS_HOST_DEVICE RowMajorInterleaved(Stride stride)
Ctor.
Definition: layout/matrix.h:279
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::layout::ContiguousMatrix::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:533
cutlass::layout::RowMajorBlockLinear::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:743
cutlass::layout::ColumnMajor::packed
static CUTLASS_HOST_DEVICE ColumnMajor packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:186
cutlass::layout::RowMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:756
cutlass::layout::RowMajorBlockLinear::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:774
cutlass::layout::GeneralMatrix::packed
static CUTLASS_HOST_DEVICE GeneralMatrix packed(MatrixCoord const &extent, MatrixLayout layout_id=MatrixLayout::kColumnMajor, Index interleave=1)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:829
cutlass::layout::ColumnMajor::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:229
cutlass::layout::RowMajor::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:118
cutlass::layout::ContiguousMatrix::packed
static CUTLASS_HOST_DEVICE ContiguousMatrix packed(MatrixCoord const &extent, Matrix layout=Matrix::kColumnMajor)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:501
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:205
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:424
cutlass::layout::Matrix::kRowMajor
leading dimension refers to stride between rows; stride along columns is 1
cutlass::Coord< kStrideRank, Index >
cutlass::layout::RowMajor
Mapping function for row-major matrices.
Definition: layout/matrix.h:50
cutlass::layout::ColumnMajor::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:193
matrix_coord.h
Defines a canonical coordinate for rank=2 matrices offering named indices.
cutlass::layout::RowMajor::RowMajor
CUTLASS_HOST_DEVICE RowMajor(Stride stride)
Ctor.
Definition: layout/matrix.h:89
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:880
cutlass::layout::RowMajorInterleaved::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:290
cutlass::layout::ContiguousMatrix::ContiguousMatrix
CUTLASS_HOST_DEVICE ContiguousMatrix(Index ldm=0, Matrix layout=Matrix::kColumnMajor)
Ctor.
Definition: layout/matrix.h:493
cutlass::layout::ColumnMajorInterleaved::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:430
cutlass::layout::RowMajorInterleaved::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:317
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:552
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Index stride(int idx) const
Returns the stride of the layout.
Definition: layout/matrix.h:217
cutlass::layout::ColumnMajorBlockLinear::packed
static CUTLASS_HOST_DEVICE ColumnMajorBlockLinear packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:625
cutlass::layout::ColumnMajorInterleaved::ColumnMajorInterleaved
CUTLASS_HOST_DEVICE ColumnMajorInterleaved(Stride stride)
Ctor.
Definition: layout/matrix.h:385
cutlass::layout::GeneralMatrix::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(MatrixCoord const &coord) const
Definition: layout/matrix.h:850
cutlass::layout::ColumnMajorInterleaved
Definition: layout/matrix.h:343
cutlass::layout::ColumnMajor::Index
int32_t Index
Index type used for coordinates.
Definition: layout/matrix.h:151
cutlass::layout::GeneralMatrix::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: layout/matrix.h:869
cutlass::layout::ColumnMajorBlockLinear::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:656
cutlass::layout::RowMajor::packed
static CUTLASS_HOST_DEVICE RowMajor packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:93
cutlass::layout::RowMajorInterleaved::RowMajorInterleaved
CUTLASS_HOST_DEVICE RowMajorInterleaved(Index ldm=0)
Ctor.
Definition: layout/matrix.h:275
cutlass::layout::RowMajorBlockLinear::RowMajorBlockLinear
CUTLASS_HOST_DEVICE RowMajorBlockLinear(Index ldm=0)
Ctor.
Definition: layout/matrix.h:722
cutlass::layout::GeneralMatrix::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:903
cutlass::layout::ColumnMajorInterleaved::inverse
CUTLASS_HOST_DEVICE MatrixCoord inverse(LongIndex offset) const
Inverse of layout function, mapping linear offset to logical coordinate.
Definition: layout/matrix.h:405
matrix_traits.h
Defines properties of matrices used to denote layout and operands to GEMM kernels.
cutlass::layout::ColumnMajorInterleaved::packed
static CUTLASS_HOST_DEVICE ColumnMajorInterleaved packed(MatrixCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: layout/matrix.h:390
cutlass::layout::ColumnMajor::stride
CUTLASS_HOST_DEVICE Index & stride(int idx)
Returns the stride of the layout.
Definition: layout/matrix.h:223
cutlass::layout::ContiguousMatrix::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: layout/matrix.h:546
cutlass.h
Basic include for CUTLASS.
cutlass::MatrixCoord
Definition: matrix_coord.h:39
cutlass::layout::ColumnMajorBlockLinear::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(MatrixCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: layout/matrix.h:674
cutlass::layout::ColumnMajor::ColumnMajor
CUTLASS_HOST_DEVICE ColumnMajor(Stride stride)
Ctor.
Definition: layout/matrix.h:181
cutlass::layout::RowMajor::kStrideRank
static int const kStrideRank
Rank of stride vector.
Definition: layout/matrix.h:56
cutlass::layout::ColumnMajorBlockLinear::ColumnMajorBlockLinear
CUTLASS_HOST_DEVICE ColumnMajorBlockLinear(Index ldm=0)
Ctor.
Definition: layout/matrix.h:621
cutlass::layout::RowMajor::RowMajor
CUTLASS_HOST_DEVICE RowMajor(Index ldm=0)
Constructor.
Definition: layout/matrix.h:85
cutlass::layout::GeneralMatrix::layout_id
CUTLASS_HOST_DEVICE MatrixLayout layout_id() const
Definition: layout/matrix.h:874
cutlass::layout::RowMajorInterleaved
Definition: layout/matrix.h:237