thread/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/array.h"
#include "cutlass/matrix_coord.h"
namespace cutlass {
namespace thread {
template <
typename Element,
int Rows,
int Columns,
typename Layout = layout::RowMajor
>
class Matrix : public Array<Element, Rows * Columns> {
public:
// Verify layout refers to a rank=2 matrix.
static_assert(
Layout::kRank == 2,
"Layout type must refer to a rank=2 matrix");
using Base = Array<Element, Rows * Columns>;
using Element = Element_;
static int const kRows = Rows;
static int const kColumns = Columns;
using Layout = Layout_;
using Reference = Element &;
static int const kRank = 2;
using Index = typename Layout::Index;
using LongIndex = typename Layout::LongIndex;
using TensorCoord = typename Layout::TensorCoord;
using Stride = typename Layout::Stride;
using TensorRef = TensorRef<Element, kRank, Layout>;
using ConstTensorRef = typename TensorRef::ConstTensorRef;
using TensorView = TensorView<Element, kRank, Layout>;
using ConstTensorView = typename TensorView::ConstTensorView;
using Diagonal = Vector<Element, __NV_STD_MIN(kRows, kColumns)>;
private:
public:
//
// Methods
//
CUTLASS_HOST_DEVICE
static MatrixCoord extent() {
return make_Coord(kRows, kColumns);
}
CUTLASS_HOST_DEVICE
static Layout layout() {
return Layout::packed(extent());
}
CUTLASS_HOST_DEVICE
Matrix() { }
CUTLASS_HOST_DEVICE
Matrix(Diagonal const &diag) {
// Todo - construct from diagonal
}
CUTLASS_HOST_DEVICE
TensorRef ref() {
return TensorRef(this->data(), layout());
}
CUTLASS_HOST_DEVICE
ConstTensorRef const_ref() const {
return ConstTensorRef(this->data(), layout());
}
CUTLASS_HOST_DEVICE
TensorView view() {
return TensorView(ref(), extent());
}
CUTLASS_HOST_DEVICE
ConstTensorView const_view() const {
return ConstTensorView(const_ref(), extent());
}
CUTLASS_HOST_DEVICE
Reference at(MatrixCoord const& coord) const {
typename Base::size_type offset_(layout().offset(coord));
return Base::at(offset_);
}
CUTLASS_HOST_DEVICE
LongIndex capacity() const {
return LongIndex(Base::size());
}
};
template <
typename Element,
int Rows,
typename Layout = layout::ColumnMajor
>
using ColumnVector = Matrix<Element, Rows, 1, Layout>;
template <
typename Element,
int Columns,
typename Layout = layout::RowMajor
>
using RowVector = Matrix<Element, 1, Columns, Layout>;
} // namespace thread
} // namespace cutlass
cutlass::thread::Matrix::TensorCoord
typename Layout::TensorCoord TensorCoord
Coordinate in logical tensor space.
Definition: thread/matrix.h:83
cutlass::thread::Matrix
Per-thread matrix object storing a packed matrix.
Definition: thread/matrix.h:47
cutlass
Definition: aligned_buffer.h:35
cutlass::thread::Matrix::Stride
typename Layout::Stride Stride
Stride type.
Definition: thread/matrix.h:86
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:387
cutlass::thread::Matrix::Base
Array< Element, Rows *Columns > Base
Base type.
Definition: thread/matrix.h:56
cutlass::thread::Matrix::ref
CUTLASS_HOST_DEVICE TensorRef ref()
Returns a TensorRef pointing to the first element of the tensor.
Definition: thread/matrix.h:136
cutlass::thread::Matrix::Diagonal
Vector< Element, __NV_STD_MIN(kRows, kColumns)> Diagonal
Diagonal vector.
Definition: thread/matrix.h:101
cutlass::thread::Matrix::extent
static CUTLASS_HOST_DEVICE MatrixCoord extent()
Returns the size of the object.
Definition: thread/matrix.h:114
cutlass::thread::Matrix::kRows
static int const kRows
Number of rows.
Definition: thread/matrix.h:62
cutlass::thread::Matrix::capacity
CUTLASS_HOST_DEVICE LongIndex capacity() const
Returns the number of scalar elements needed to store tensor.
Definition: thread/matrix.h:167
cutlass::TensorRef::ConstTensorRef
TensorRef< typename platform::remove_const< Element >::type const, Layout > ConstTensorRef
TensorRef to constant data.
Definition: tensor_ref.h:179
cutlass::thread::Matrix::ConstTensorRef
typename TensorRef::ConstTensorRef ConstTensorRef
TensorRef to constant matrix object.
Definition: thread/matrix.h:92
cutlass::layout::ColumnMajor
Mapping function for column-major matrices.
Definition: layout/matrix.h:142
cutlass::thread::Matrix::TensorView
TensorView< Element, kRank, Layout > TensorView
TensorRef to matrix object.
Definition: thread/matrix.h:95
cutlass::thread::Matrix::view
CUTLASS_HOST_DEVICE TensorView view()
Returns a TensorRef pointing to the first element of the tensor.
Definition: thread/matrix.h:148
array.h
Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...
cutlass::TensorView::ConstTensorView
TensorView< typename platform::remove_const< Element >::type const, Layout > ConstTensorView
TensorView pointing to constant memory.
Definition: tensor_view.h:95
cutlass::thread::Matrix::at
CUTLASS_HOST_DEVICE Reference at(MatrixCoord const &coord) const
Returns a reference to the element at a given Coord.
Definition: thread/matrix.h:160
cutlass::thread::Matrix::TensorRef
TensorRef< Element, kRank, Layout > TensorRef
TensorRef to matrix object.
Definition: thread/matrix.h:89
cutlass::thread::Matrix::Index
typename Layout::Index Index
Index type.
Definition: thread/matrix.h:77
cutlass::thread::Matrix::ConstTensorView
typename TensorView::ConstTensorView ConstTensorView
TensorRef to constant matrix object.
Definition: thread/matrix.h:98
cutlass::thread::Matrix::LongIndex
typename Layout::LongIndex LongIndex
Long index used for pointer offsets.
Definition: thread/matrix.h:80
cutlass::thread::Matrix::Element
Element_ Element
Element type.
Definition: thread/matrix.h:59
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
cutlass::layout::RowMajor
Mapping function for row-major matrices.
Definition: layout/matrix.h:50
cutlass::thread::Matrix::Matrix
CUTLASS_HOST_DEVICE Matrix()
Ctor.
Definition: thread/matrix.h:126
matrix_coord.h
Defines a canonical coordinate for rank=2 matrices offering named indices.
cutlass::thread::Matrix::Matrix
CUTLASS_HOST_DEVICE Matrix(Diagonal const &diag)
Ctor.
Definition: thread/matrix.h:130
cutlass::thread::Matrix::Reference
Element & Reference
Reference type to an element.
Definition: thread/matrix.h:71
cutlass::thread::Matrix::const_ref
CUTLASS_HOST_DEVICE ConstTensorRef const_ref() const
Returns a TensorRef pointing to the first element of the tensor.
Definition: thread/matrix.h:142
cutlass::thread::Matrix::layout
static CUTLASS_HOST_DEVICE Layout layout()
Returns the layout object.
Definition: thread/matrix.h:120
cutlass.h
Basic include for CUTLASS.
cutlass::MatrixCoord
Definition: matrix_coord.h:39
cutlass::thread::Matrix::const_view
CUTLASS_HOST_DEVICE ConstTensorView const_view() const
Returns a TensorView to const data.
Definition: thread/matrix.h:154
cutlass::thread::Matrix::kRank
static int const kRank
Logical rank of tensor index space.
Definition: thread/matrix.h:74
cutlass::thread::Matrix::kColumns
static int const kColumns
Number of columns.
Definition: thread/matrix.h:65