This class represents a matrix that's stored on the GPU if we have one, and in memory if not. More...

#include <matrix-common.h>

Inheritance diagram for CuMatrix< Real >:

Collaboration diagram for CuMatrix< Real >:

Public Member Functions
	CuMatrix ()

	CuMatrix (MatrixIndexT rows, MatrixIndexT cols, MatrixResizeType resize_type=kSetZero, MatrixStrideType stride_type=kDefaultStride)
	Constructor with memory initialisation. More...

	CuMatrix (const CuMatrix< Real > &other, MatrixTransposeType trans=kNoTrans)

	CuMatrix (const CuBlockMatrix< Real > &other, MatrixTransposeType trans=kNoTrans)

	CuMatrix (const CuMatrixBase< Real > &other, MatrixTransposeType trans=kNoTrans)

template<typename OtherReal >
	CuMatrix (const MatrixBase< OtherReal > &other, MatrixTransposeType trans=kNoTrans)

	CuMatrix (const CuSpMatrix< Real > &M)
	Copy constructor taking SpMatrix... More...

template<typename OtherReal >
	CuMatrix (const CuTpMatrix< OtherReal > &M, MatrixTransposeType trans=kNoTrans)
	Copy constructor taking TpMatrix... More...

template<typename OtherReal >
	CuMatrix (const CuMatrixBase< OtherReal > &M, MatrixTransposeType trans=kNoTrans)
	Copy constructor: as above, but from another type. More...

CuMatrix< Real > &	operator= (const CuMatrixBase< Real > &other)

CuMatrix< Real > &	operator= (const CuMatrix< Real > &other)

CuMatrix< Real > &	operator= (const MatrixBase< Real > &other)

void	Transpose ()

void	Resize (MatrixIndexT rows, MatrixIndexT cols, MatrixResizeType resize_type=kSetZero, MatrixStrideType stride_type=kDefaultStride)
	Allocate the memory. More...

void	Swap (Matrix< Real > *mat)

void	Swap (CuMatrix< Real > *mat)

template<typename OtherReal >
void	Swap (CuMatrix< OtherReal > *mat)

void	Read (std::istream &is, bool binary)
	I/O functions. More...

	~CuMatrix ()
	Destructor. More...

const Matrix< Real > &	Mat () const

Matrix< Real > &	Mat ()

void	CompObjfAndDeriv (const std::vector< MatrixElement< Real > > &elements, const CuMatrix< Real > &A, Real tot_objf, Real tot_weight)
	Here, A is interpreted as a matrix of probabilities, and "elements" as a list of posteriors (possibly zero-one), and "*this" as a matrix of derivatives w.r.t. More...

Public Member Functions inherited from CuMatrixBase< Real >
void	CopyCols (const CuMatrixBase< Real > &src, const CuArrayBase< MatrixIndexT > &indexes)
	Copies column r from column indexes[r] of src. More...

void	AddCols (const CuMatrixBase< Real > &src, const CuArrayBase< MatrixIndexT > &indices)
	Add column indices[r] of src to column r. More...

void	CopyRows (const CuMatrixBase< Real > &src, const CuArrayBase< MatrixIndexT > &indexes)
	Copies row r from row indexes[r] of src. More...

void	CopyRows (const CuArrayBase< const Real *> &src)
	Copies row r of this matrix from an array of floats at the location given by src[r], where src[r] is assumed to be obtained from the RowData() function of another CuMatrix, or from CuVector::Data() (the point is: the data it points to should be on the GPU if we're using a GPU, and on a CPU otherwise). More...

void	CopyToRows (const CuArrayBase< Real *> &dst) const
	For each row r of this matrix, copies it to the array of floats at the location given by dst[r], where dst[r] is assumed to be obtained from the RowData() function of another CuMatrix, or from CuVector::Data() (i.e. More...

void	AddRows (Real alpha, const CuMatrixBase< Real > &src, const CuArrayBase< MatrixIndexT > &indexes)
	Does for each row r, this.Row(r) += alpha * src.row(indexes[r]). More...

void	MulRows (const CuMatrixBase< Real > &src, const CuArrayBase< MatrixIndexT > &indexes)
	Does for each row r, this.Row(r) = alpha src.row(indexes[r]), where '*=' is elementwise multiplication. More...

void	AddRows (Real alpha, const CuArrayBase< const Real *> &src)
	Does for each row r, this.Row(r) += alpha * src[r], treating src[r] as the beginning of a region of memory representing a vector of floats, of the same length as this.NumCols(). More...

void	AddToRows (Real alpha, const CuArrayBase< MatrixIndexT > &indexes, CuMatrixBase< Real > *dst) const
	For each row i of *this, adds this->Row(i) to dst->Row(indexes(i)) if indexes(i) >= 0, else do nothing. More...

void	AddToRows (Real alpha, const CuArrayBase< Real *> &dst) const
	For each row r of this matrix, adds it (times alpha) to the array of floats at the location given by dst[r], where dst[r] is assumed to be obtained from the RowData() function of another CuMatrix, or from CuVector::Data() (i.e. More...

void	SumColumnRanges (const CuMatrixBase< Real > &src, const CuArrayBase< Int32Pair > &indexes)
	For each row r of this and for each column c, sets (*this)(r, c) to the sum src(r, j), where j ranges from indexes[c].first through indexes[c].second - 1. More...

void	AddRowRanges (const CuMatrixBase< Real > &src, const CuArrayBase< Int32Pair > &indexes)
	For each row r of this and for each column c, do (*this)(r, c) += src(j, c), where j ranges from indexes[r].first through indexes[r].second - 1. More...

void	AddToDiag (Real value)
	Adds "value" to the diagonal elements of the matrix. More...

MatrixIndexT	NumRows () const
	Dimensions. More...

MatrixIndexT	NumCols () const

MatrixIndexT	Stride () const

::MatrixDim	Dim () const

Real	FrobeniusNorm () const

bool	IsUnit (Real tol=0.001) const

bool	ApproxEqual (const CuMatrixBase< Real > &other, float tol=0.01) const
	True if ((this)-other).FrobeniusNorm() <= tol this->FrobeniusNorm() More...

MatrixIndexT	SizeInBytes () const
	Get size of matrix in bytes. More...

template<typename OtherReal >
void	CopyFromMat (const MatrixBase< OtherReal > &src, MatrixTransposeType trans=kNoTrans)

void	CopyFromGeneralMat (const GeneralMatrix &src, MatrixTransposeType trans=kNoTrans)

void	CopyFromMat (const MatrixBase< Real > &src, MatrixTransposeType trans=kNoTrans)

void	CopyFromSp (const CuSpMatrix< Real > &M)

template<typename OtherReal >
void	CopyFromTp (const CuTpMatrix< OtherReal > &M, MatrixTransposeType trans=kNoTrans)

void	CopyRangeFromMatClamped (const CuMatrixBase< Real > &src, int32_t start_range, int32_t end_range, int32_t clamp_low, int32_t clamp_high)

template<typename OtherReal >
void	CopyFromMat (const CuMatrixBase< OtherReal > &M, MatrixTransposeType trans=kNoTrans)

template<typename OtherReal >
void	CopyToMat (MatrixBase< OtherReal > *dst, MatrixTransposeType trans=kNoTrans) const

void	CopyRowsFromVec (const CuVectorBase< Real > &v)
	This function has two modes of operation. More...

void	CopyRowsFromVec (const VectorBase< Real > &v)
	Version of CopyRowsFromVec() that takes a CPU-based vector. More...

void	CopyColsFromVec (const CuVectorBase< Real > &v)
	Copies vector into matrix, column-by-column. More...

void	CopyColFromVec (const CuVectorBase< Real > &v, const MatrixIndexT col)
	Copy vector into specific column of matrix. More...

void	Sigmoid (const CuMatrixBase< Real > &src)
	Set each element to the sigmoid of the corresponding element of "src": element by element, x = 1 / (1 + exp(-x)) More...

void	Heaviside (const CuMatrixBase< Real > &src)
	Set each element to the Heaviside function of the corresponding element of "src", which we define as the function (x > 0 ? 1.0 : 0.0) [note: in general, there are different ways to deal with the situation when x==0. More...

void	Exp (const CuMatrixBase< Real > &src)

void	Log (const CuMatrixBase< Real > &src)

void	Pow (const CuMatrixBase< Real > &src, Real power)

void	PowAbs (const CuMatrixBase< Real > &src, Real power, bool include_sign=false)
	Apply power to the absolute value of each element. More...

void	Floor (const CuMatrixBase< Real > &src, Real floor_val)

void	Ceiling (const CuMatrixBase< Real > &src, Real ceiling_val)

void	ExpLimited (const CuMatrixBase< Real > &src, Real lower_limit, Real upper_limit)
	This is equivalent to running: Floor(src, lower_limit); Ceiling(src, upper_limit); Exp(src) More...

void	ExpSpecial (const CuMatrixBase< Real > &src)
	For each element x of the matrix, set it to (x < 0 ? exp(x) : x + 1). More...

void	SoftMaxPerRow (const CuMatrixBase< Real > &src)
	Softmax nonlinearity Y = Softmax(X) : Yij = e^Xij / sum_k(e^Xik), done to each row, with attention to avoiding overflow or underflow. More...

void	LogSoftMaxPerRow (const CuMatrixBase< Real > &src)
	LogSoftmax nonlinearity Y = LogSoftmax(X) : Yij = Xij - log(sum_k(e^Xik)), done to each row, with attention to avoiding overflow or underflow. More...

void	SoftHinge (const CuMatrixBase< Real > &src)
	Apply the function y = log(1 + exp(x)), to each element. More...

void	GroupPnorm (const CuMatrixBase< Real > &src, Real pow)
	Apply the function y(i) = (sum_{j = iG}^{(i+1)G-1} x_j ^ (power)) ^ (1 / p) where G = x.NumCols() / y.NumCols() must be an integer. More...

void	DiffGroupPnorm (const CuMatrixBase< Real > &in_value, const CuMatrixBase< Real > &out_value, const CuMatrixBase< Real > &out_deriv, Real power)
	Differentiate backward through the GroupPnorm function. More...

void	GroupMax (const CuMatrixBase< Real > &src)
	Apply the function y(i) = (max_{j = iG}^{(i+1)G-1} x_j where G = x.NumCols() / y.NumCols() must be an integer. More...

void	GroupMaxDeriv (const CuMatrixBase< Real > &input, const CuMatrixBase< Real > &output)
	Calculate derivatives for the GroupMax function above, where "input" is the input to the GroupMax function above (i.e. More...

void	ParametricRelu (const CuMatrixBase< Real > &src, const CuVectorBase< Real > &alpha, const CuVectorBase< Real > &beta)
	Compute the parametric rectified linear unit function; element by element, this = src (src > 0 ? alpha : beta) More...

void	DiffParametricRelu (const CuMatrixBase< Real > &value, const CuMatrixBase< Real > &diff, const CuVectorBase< Real > &alpha, const CuVectorBase< Real > &beta)
	Differentiate backward through the parametric relu function. More...

void	Tanh (const CuMatrixBase< Real > &src)
	Compute the hyperbolic tangent (tanh) function; element by element, *this = tanh(src). More...

void	DiffSigmoid (const CuMatrixBase< Real > &value, const CuMatrixBase< Real > &diff)
	Differentiate backward through the sigmoid function. More...

void	DiffTanh (const CuMatrixBase< Real > &value, const CuMatrixBase< Real > &diff)
	Differentiate backward through the tanh function. More...

void	DiffSoftmaxPerRow (const CuMatrixBase< Real > &value, const CuMatrixBase< Real > &diff)
	Differentiate backward through the softmax function. More...

void	DiffLogSoftmaxPerRow (const CuMatrixBase< Real > &out_value, const CuMatrixBase< Real > &out_deriv)
	Differentiate backward through the log softmax function. More...

void	DiffXent (const CuArrayBase< int32 > &tgt, CuVector< Real > *log_post_tgt)
	Differentiate the block [softmax+cross-entropy] : dE/da = posterior_mat - target_mat, 'E' is error function, 'a' is activation on softmax input. More...

void	Cholesky (CuMatrixBase< Real > *inv_cholesky=NULL)
	This function does sets this to the Cholesky factor of this (i.e. More...

void	SymInvertPosDef ()
	Inversion for positive definite symmetric matrices. More...

void	ApplyPow (Real power)

void	ApplyPowAbs (Real power, bool include_sign=false)

void	ApplyHeaviside ()

void	ApplyFloor (Real floor_val)

void	ApplyCeiling (Real ceiling_val)

void	ApplyExp ()

void	ApplyExpLimited (Real lower_limit, Real upper_limit)

void	ApplyExpSpecial ()

void	ApplySoftMaxPerRow ()

void	ApplyLogSoftMaxPerRow ()

void	ApplyLog ()

void	FindRowMaxId (CuArray< int32 > *id) const
	Find the id of the maximal element for each row (resizes the 'id' array to the appropriate size). More...

void	SetZero ()
	Math operations, some calling kernels. More...

void	Set (Real value)

void	Add (Real value)

void	SetZeroAboveDiag ()
	Zeroes all elements for which col > row. More...

void	Scale (Real value)

void	MulElements (const CuMatrixBase< Real > &A)
	Multiply two matrices elementwise: C = C .* A. More...

void	DivElements (const CuMatrixBase< Real > &A)
	Divide two matrices elementwise: C = A ./ A. More...

void	Max (const CuMatrixBase< Real > &A)
	Do, elementwise, this = max(this, A). More...

void	Min (const CuMatrixBase< Real > &A)
	Do, elementwise, this = min(this, A). More...

void	MulColsVec (const CuVectorBase< Real > &scale)
	scale i'th column by scale[i] More...

void	MulRowsVec (const CuVectorBase< Real > &scale)
	scale i'th row by scale[i] More...

void	MulRowsGroupMat (const CuMatrixBase< Real > &src)
	divide each row into src.NumCols() groups, and then scale i'th row's jth group of elements by src[i, j]. More...

void	DivRowsVec (const CuVectorBase< Real > &div)
	divide i'th row by scale[i] More...

void	InvertElements ()
	invert the matrix by elements. More...

void	AddMat (Real alpha, const CuMatrixBase< Real > &A, MatrixTransposeType trans=kNoTrans)
	this += alpha A More...

void	AddSmat (Real alpha, const CuSparseMatrix< Real > &A, MatrixTransposeType trans=kNoTrans)
	this += alpha A. More...

void	AddSmatMat (Real alpha, const CuSparseMatrix< Real > &A, MatrixTransposeType transA, const CuMatrixBase< Real > &B, Real beta)
	(this) = alpha op(A) * B + beta * (*this), where A is sparse. More...

void	AddMatSmat (Real alpha, const CuMatrixBase< Real > &A, const CuSparseMatrix< Real > &B, MatrixTransposeType transB, Real beta)
	(this) = alpha A * op(B) + beta * (*this), where B is sparse and op(B) is either B or trans(B) depending on the 'transB' argument. More...

void	AddToElements (Real alpha, const CuArrayBase< int32 > &elements)
	This is a rather special purpose function; we might generalize it later by adding a transpose-type option. More...

void	AddMatBlocks (Real alpha, const CuMatrixBase< Real > &A, MatrixTransposeType trans=kNoTrans)
	This function is like AddMat (it does this += alpha src), except that it supports cases where *this and src have different dimension. More...

void	AddVecToCols (Real alpha, const CuVectorBase< Real > &col, Real beta=1.0)
	(for each column c of this), c = alpha col + beta * c More...

void	AddVecToRows (Real alpha, const CuVectorBase< Real > &row, Real beta=1.0)
	(for each row r of this), r = alpha row + beta * r More...

void	AddMatMat (Real alpha, const CuMatrixBase< Real > &A, MatrixTransposeType transA, const CuMatrixBase< Real > &B, MatrixTransposeType transB, Real beta)
	C = alpha * A(^T)B(^T) + beta C. More...

void	AddVecVec (Real alpha, const CuVectorBase< Real > &x, const CuVectorBase< Real > &y)
	A = alpha * x * y^T + A . More...

void	SetMatMatDivMat (const CuMatrixBase< Real > &A, const CuMatrixBase< Real > &B, const CuMatrixBase< Real > &C)
	this = a b / c (by element; when c = 0, this = a) this can be an alias of a, b or c safely and get expected result. More...

void	SymAddMat2 (const Real alpha, const CuMatrixBase< Real > &M, MatrixTransposeType transA, Real beta)
	this = beta this + alpha M M^T, for symmetric matrices. More...

void	AddMatBlock (Real alpha, const CuMatrixBase< Real > &A, MatrixTransposeType transA, const CuBlockMatrix< Real > &B, MatrixTransposeType transB, Real beta)
	This function is like AddMatMat but for where the second argument is of type CuBlockMatrix (a block-diagonal matrix of blocks). More...

void	AddDiagVecMat (const Real alpha, const CuVectorBase< Real > &v, const CuMatrixBase< Real > &M, MatrixTransposeType transM, Real beta=1.0)
	this = beta this + alpha diag(v) * M [or M^T]. More...

void	AddMatDiagVec (const Real alpha, const CuMatrixBase< Real > &M, MatrixTransposeType transM, CuVectorBase< Real > &v, Real beta=1.0)

void	AddMatMatElements (const Real alpha, const CuMatrixBase< Real > &A, const CuMatrixBase< Real > &B, const Real beta)
	this = beta this + alpha A .* B (.* element by element multiplication) More...

void	AddMatSp (const Real alpha, const CuMatrixBase< Real > &A, MatrixTransposeType transA, const CuSpMatrix< Real > &B, const Real beta)
	this <– betathis + alphaA*B More...

void	AddSpMat (const Real alpha, const CuSpMatrix< Real > &A, const CuMatrixBase< Real > &B, MatrixTransposeType transB, const Real beta)
	this <– betathis + alphaSpA*B More...

void	AddTpMat (const Real alpha, const CuTpMatrix< Real > &A, MatrixTransposeType transA, const CuMatrixBase< Real > &B, MatrixTransposeType transB, const Real beta)
	this <– betathis + alphaA*B. More...

void	AddMatTp (const Real alpha, const CuMatrixBase< Real > &A, MatrixTransposeType transA, const CuTpMatrix< Real > &B, MatrixTransposeType transB, const Real beta)
	this <– betathis + alphaA*B. More...

void	CopyFromBlock (const CuBlockMatrix< Real > &B, MatrixTransposeType trans=kNoTrans)

void	CopyLowerToUpper ()

void	CopyUpperToLower ()

CuSubMatrix< Real >	Range (const MatrixIndexT row_offset, const MatrixIndexT num_rows, const MatrixIndexT col_offset, const MatrixIndexT num_cols) const

CuSubMatrix< Real >	RowRange (const MatrixIndexT row_offset, const MatrixIndexT num_rows) const

CuSubMatrix< Real >	ColRange (const MatrixIndexT col_offset, const MatrixIndexT num_cols) const

const CuSubVector< Real >	Row (MatrixIndexT i) const

CuSubVector< Real >	Row (MatrixIndexT i)

CuValue< Real >	operator() (MatrixIndexT r, MatrixIndexT c)

Real	operator() (MatrixIndexT r, MatrixIndexT c) const

Real	Sum () const

Real	Max () const

Real	Min () const

Real	Trace (bool check_square=true) const
	Return the trace. If check_square = true, will crash if matrix is not square. More...

void	SetRandn ()

void	SetRandUniform ()

void	Write (std::ostream &os, bool binary) const

void	AddElements (Real alpha, const std::vector< MatrixElement< Real > > &input)

void	AddElements (Real alpha, const CuArrayBase< Int32Pair > &indexes, const Real *input)

void	Lookup (const std::vector< Int32Pair > &indexes, Real *output) const

void	Lookup (const CuArrayBase< Int32Pair > &indexes, Real *output) const

void	EqualElementMask (const CuMatrixBase< Real > &mat, CuMatrix< Real > *mask) const

const Real *	RowData (MatrixIndexT r) const
	Get raw row pointer (const). More...

Real *	RowData (MatrixIndexT r)
	Get raw row pointer. More...

const Real *	Data () const
	Return data pointer (const). More...

Real *	Data ()
	Return data pointer. More...

const MatrixBase< Real > &	Mat () const

MatrixBase< Real > &	Mat ()

Private Member Functions
void	Destroy ()

Additional Inherited Members
Protected Member Functions inherited from CuMatrixBase< Real >
	CuMatrixBase ()

	CuMatrixBase (Real *data, MatrixIndexT num_rows, MatrixIndexT num_cols, MatrixIndexT stride)
	This constructor takes the #rows, #cols and stride; it's called from the constructor of CuSubMatrix. More...

Protected Attributes inherited from CuMatrixBase< Real >
Real *	data_
	GPU data pointer (or regular matrix data pointer,. More...

MatrixIndexT	num_cols_

MatrixIndexT	num_rows_

MatrixIndexT	stride_

Detailed Description

template<typename Real>
class kaldi::CuMatrix< Real >

This class represents a matrix that's stored on the GPU if we have one, and in memory if not.

Definition at line 71 of file matrix-common.h.

Constructor & Destructor Documentation

◆ CuMatrix() [1/9]

CuMatrix ( )

inline

Definition at line 799 of file cu-matrix.h.

799 { }

◆ CuMatrix() [2/9]

CuMatrix	(	MatrixIndexT	rows,
		MatrixIndexT	cols,
		MatrixResizeType	resize_type = `kSetZero`,
		MatrixStrideType	stride_type = `kDefaultStride`
	)

inline

Constructor with memory initialisation.

Definition at line 802 of file cu-matrix.h.

                                                           {
     Resize(rows, cols, resize_type, stride_type);
   }

◆ CuMatrix() [3/9]

CuMatrix	(	const CuMatrix< Real > &	other,
		MatrixTransposeType	trans = `kNoTrans`
	)

Definition at line 380 of file cu-matrix.cc.

                                                                                {
   if (trans == kNoTrans)
     this->Resize(other.NumRows(), other.NumCols(), kUndefined);
   else
     this->Resize(other.NumCols(), other.NumRows(), kUndefined);
   this->CopyFromMat(other, trans);
 }

◆ CuMatrix() [4/9]

CuMatrix	(	const CuBlockMatrix< Real > &	other,
		MatrixTransposeType	trans = `kNoTrans`
	)

explicit

Definition at line 195 of file cu-matrix.cc.

                                                     : CuMatrixBase<Real>() {
   if (trans == kNoTrans) {
     Resize(B.NumRows(), B.NumCols(), kUndefined);
     this->CopyFromBlock(B);
   } else {
     Resize(B.NumCols(), B.NumRows(), kUndefined);
     this->CopyFromBlock(B, kTrans);
   }
 }

◆ CuMatrix() [5/9]

CuMatrix	(	const CuMatrixBase< Real > &	other,
		MatrixTransposeType	trans = `kNoTrans`
	)

explicit

Definition at line 389 of file cu-matrix.cc.

                                                                                    {
   if (trans == kNoTrans)
     this->Resize(other.NumRows(), other.NumCols(), kUndefined);
   else
     this->Resize(other.NumCols(), other.NumRows(), kUndefined);
   this->CopyFromMat(other, trans);
 }

◆ CuMatrix() [6/9]

CuMatrix	(	const MatrixBase< OtherReal > &	other,
		MatrixTransposeType	trans = `kNoTrans`
	)

explicit

Definition at line 400 of file cu-matrix.cc.

                                                                                       {
   if (trans == kNoTrans)
     this->Resize(other.NumRows(), other.NumCols(), kUndefined);
   else
     this->Resize(other.NumCols(), other.NumRows(), kUndefined);
   this->CopyFromMat(other, trans);
 }

◆ CuMatrix() [7/9]

CuMatrix ( const CuSpMatrix< Real > & M )

inlineexplicit

Copy constructor taking SpMatrix...

Definition at line 824 of file cu-matrix.h.

                                                : CuMatrixBase<Real>() {
     Resize(M.NumRows(), M.NumRows(), kUndefined);
     this->CopyFromSp(M);
   }

◆ CuMatrix() [8/9]

CuMatrix	(	const CuTpMatrix< OtherReal > &	M,
		MatrixTransposeType	trans = `kNoTrans`
	)

inlineexplicit

Copy constructor taking TpMatrix...

Definition at line 831 of file cu-matrix.h.

                                                           : CuMatrixBase<Real>() {
     Resize(M.NumCols(), M.NumRows(), kUndefined);
     this->CopyFromTp(M, trans);
   }

◆ CuMatrix() [9/9]

CuMatrix	(	const CuMatrixBase< OtherReal > &	M,
		MatrixTransposeType	trans = `kNoTrans`
	)

explicit

Copy constructor: as above, but from another type.

Copy constructor from another type.

Definition at line 3168 of file cu-matrix.cc.

                                                     : CuMatrixBase<Real>() {
 
   if (trans == kNoTrans) {
     Resize(M.NumRows(), M.NumCols());
     this->CopyFromMat(M);
   } else {
     Resize(M.NumCols(), M.NumRows());
     this->CopyFromMat(M, kTrans);
   }
 }

◆ ~CuMatrix()

~CuMatrix ( )

inline

Destructor.

Definition at line 877 of file cu-matrix.h.

877 { Destroy(); }

kaldi::CuMatrix::Destroy

void Destroy()

Definition: cu-matrix.cc:94

Member Function Documentation

◆ CompObjfAndDeriv()

void CompObjfAndDeriv	(	const std::vector< MatrixElement< Real > > &	elements,
		const CuMatrix< Real > &	A,
		Real *	tot_objf,
		Real *	tot_weight
	)

Here, A is interpreted as a matrix of probabilities, and "elements" as a list of posteriors (possibly zero-one), and "*this" as a matrix of derivatives w.r.t.

the log-probs. This function does: for each element { row, column, weight } indexed i in the vector "elements", let x(i) = A(row(i), column(i)); then it does (*this)(row(i), column(i)) += weight(i) / x(i), and *tot_objf = weight(i) * log(x(i)), and *tot_weight = weight(i) Preconditions: A must be strictly positive, and no (row, column) pair may be repeated within "elements"

Definition at line 1661 of file cu-matrix.cc.

Referenced by NnetUpdater::ComputeObjfAndDeriv(), kaldi::TestCuMatrixCompObjfAndDeriv(), and kaldi::UnitTestCuMatrixObjfDeriv().

                                                                         {
   { // check the input.
     typedef typename std::vector<MatrixElement<Real> >::const_iterator Iter;
     MatrixIndexT num_rows = this->num_rows_, num_cols = this->num_cols_;
     for (Iter iter = sv_labels.begin(); iter != sv_labels.end(); ++iter) {
       KALDI_ASSERT(iter->row < num_rows && iter->row >= 0 &&
                    iter->column < num_cols && iter->column >= 0);
     }
   }
 
 #if HAVE_CUDA == 1
   if (CuDevice::Instantiate().Enabled()) {
     if (sv_labels.empty()) {
       KALDI_WARN << "Empty supervision labels";
       *tot_objf = 0.0;
       *tot_weight = 0.0;
       return;
     }
     void *addr = CuDevice::Instantiate().Malloc(sv_labels.size() * sizeof(MatrixElement<Real>));
     CU_SAFE_CALL(cudaMemcpyAsync(addr, sv_labels.data(), sv_labels.size() *
                                  sizeof(MatrixElement<Real>),
                                  cudaMemcpyHostToDevice,
                                  cudaStreamPerThread));
     CuTimer tim;
     CuVector<Real> tmp(2, kUndefined);
     int dimBlock(CU1DBLOCK);
     int dimGrid = 1; // only 1 block here. we have loops in each thread.
     cuda_comp_obj_deriv(dimGrid, dimBlock, (MatrixElement<Real>*)addr,
                         sv_labels.size(), output.Data(), output.Dim(),
                         this->Data(), this->Dim(), tmp.Data());
     Vector<Real> tmp_cpu(tmp);
     *tot_objf = tmp_cpu(0);
     *tot_weight = tmp_cpu(1);
     CuDevice::Instantiate().Free(addr);
     CuDevice::Instantiate().AccuProfile(__func__, tim);
   } else
 #endif
   {
     *tot_objf = 0.0;
     *tot_weight = 0.0;
     for(int32 i = 0; i<sv_labels.size(); i++) {
       int32 m = sv_labels[i].row, label = sv_labels[i].column;
       Real weight = sv_labels[i].weight;
       //KALDI_ASSERT(label >= 0 && label < nnet_.OutputDim());
       Real this_prob = output(m, label);
       KALDI_ASSERT(this_prob >= 0.99e-20); // we floored to 1.0e-20 in SoftmaxLayer.
       *tot_objf += weight * kaldi::Log(this_prob);
       *tot_weight += weight;
       (*this)(m, label) += weight / this_prob;
     }
   }
 }

◆ Destroy()

void Destroy ( )

private

Definition at line 94 of file cu-matrix.cc.

                              {
 #if HAVE_CUDA == 1
   if (CuDevice::Instantiate().Enabled()) {
     if (this->data_ != NULL) {
       CuTimer tim;
       CuDevice::Instantiate().Free(this->data_);
       CuDevice::Instantiate().AccuProfile(__func__, tim);
     }
   } else
 #endif
   {
     if (this->data_ != NULL) KALDI_MEMALIGN_FREE(this->data_);
   }
   this->data_ = NULL;
   this->num_rows_ = 0;
   this->num_cols_ = 0;
   this->stride_ = 0;
 }

◆ Mat() [1/2]

const Matrix<Real>& Mat ( ) const

inline

Definition at line 879 of file cu-matrix.h.

Referenced by CuVectorBase< float >::CopyDiagFromMat(), CuRand< float >::RandGaussian(), and CuRand< float >::RandUniform().

                                          {
     return *(reinterpret_cast<const Matrix<Real>* >(this));
   }

◆ Mat() [2/2]

Matrix<Real>& Mat ( )

inline

Definition at line 882 of file cu-matrix.h.

                              {
     return *(reinterpret_cast<Matrix<Real>* >(this));
   }

◆ operator=() [1/3]

CuMatrix<Real>& operator= ( const CuMatrixBase< Real > & other )

inline

Definition at line 842 of file cu-matrix.h.

                                                                {
     this->Resize(other.NumRows(), other.NumCols(), kUndefined);
     this->CopyFromMat(other);
     return *this;
   }

◆ operator=() [2/3]

CuMatrix<Real>& operator= ( const CuMatrix< Real > & other )

inline

Definition at line 848 of file cu-matrix.h.

                                                            {
     this->Resize(other.NumRows(), other.NumCols(), kUndefined);
     this->CopyFromMat(other);
     return *this;
   }

◆ operator=() [3/3]

CuMatrix<Real>& operator= ( const MatrixBase< Real > & other )

inline

Definition at line 854 of file cu-matrix.h.

                                                              {
     this->Resize(other.NumRows(), other.NumCols(), kUndefined);
     this->CopyFromMat(other);
     return *this;
   }

◆ Read()

void Read	(	std::istream &	is,
		bool	binary
	)

I/O functions.

Definition at line 494 of file cu-matrix.cc.

Referenced by LstmNonlinearityComponent::ConsolidateMemory(), FixedAffineComponent::InitFromConfig(), FixedLinearComponent::InitFromString(), FixedAffineComponent::InitFromString(), RestrictedAttentionComponent::Read(), kaldi::nnet1::ReadCuMatrixFromString(), and kaldi::UnitTestCuMatrixIO().

                                                      {
   Matrix<Real> temp;
   temp.Read(is, binary);
   Destroy();
   Swap(&temp);
 }

◆ Resize()

void Resize	(	MatrixIndexT	rows,
		MatrixIndexT	cols,
		MatrixResizeType	resize_type = `kSetZero`,
		MatrixStrideType	stride_type = `kDefaultStride`
	)

Allocate the memory.

Definition at line 50 of file cu-matrix.cc.

                                                           {
   // This code does not currently support the other resize_type options.
   KALDI_ASSERT(resize_type == kSetZero || resize_type == kUndefined);
   if (rows * cols == 0) KALDI_ASSERT(rows == 0 && cols == 0);
   if (this->num_rows_ == rows && this->num_cols_ == cols) {
     if (resize_type == kSetZero) this->SetZero();
     return;
   }
   if (this->num_rows_ != 0)
     this->Destroy();
   if (rows == 0) return;
 #if HAVE_CUDA == 1
   if (CuDevice::Instantiate().Enabled()) {
     CuTimer tim;
     MatrixIndexT row_bytes = cols * sizeof(Real);
     size_t pitch;
     if (stride_type == kDefaultStride) {
       this->data_ = static_cast<Real*>(CuDevice::Instantiate().MallocPitch(
           row_bytes, rows, &pitch));
       this->num_rows_ = rows;
       this->num_cols_ = cols;
       this->stride_ = pitch / sizeof(Real);
     } else {  // kStrideEqualNumCols
       size_t bytes = rows * cols * sizeof(Real);
       this->data_ = static_cast<Real*>(CuDevice::Instantiate().Malloc(bytes));
       this->num_rows_ = rows;
       this->num_cols_ = cols;
       this->stride_ = cols;
     }
     if (resize_type == kSetZero) this->SetZero();
     CuDevice::Instantiate().AccuProfile("CuMatrix::Resize", tim);
   } else
 #endif
   { // Let the initializer of Matrix<Real> handle the allocation,
     // and then just do Swap which will switch the pointers.
     // This wastes a few instructions but is simple to code.
     Matrix<Real> mat(rows, cols, resize_type, stride_type);
     this->Swap(&mat);
   }
 }

◆ Swap() [1/3]

void Swap ( Matrix< Real > * mat )

Definition at line 123 of file cu-matrix.cc.

Referenced by DecodableNnetLoopedOnlineBase::AdvanceChunk(), DecodableNnet2Online::ComputeForFrame(), kaldi::nnet3::ComputeObjectiveFunction(), GeneralMatrix::CopyToMat(), DecodableNnetSimple::DoNnetComputation(), Nnet::Feedforward(), main(), kaldi::nnet3::ReduceRankOfComponents(), NnetRescaler::Rescale(), SingleUtteranceNnet2DecoderThreaded::RunNnetEvaluationInternal(), Matrix< BaseFloat >::Swap(), kaldi::UnitTestSwapCu2Cu(), and kaldi::UnitTestSwapCu2M().

                                            {
 #if HAVE_CUDA == 1
   if (CuDevice::Instantiate().Enabled()) {
     if (this->num_rows_ == 0) {
       if (mat->num_rows_ != 0) {
         // *this is empty, but mat is nonempty.
         this->Resize(mat->num_rows_, mat->num_cols_, kUndefined);
         this->CopyFromMat(*mat);
         mat->Resize(0, 0);
       }
       // else both are empty.
     } else { // *this is nonempty.
       if (mat->num_rows_ != 0) {
         // Both *this and *mat are nonempty.  Recurse to simpler cases.
         // this could be done more efficiently in the case where
         // the size does not change.
         Matrix<Real> temp;
         this->Swap(&temp); // now temp is full, *this is empty.
         mat->Swap(&temp); // now mat has data from *this, temp has
         // data from mat.
         this->Swap(&temp); // copy data in mat to *this, which is now empty.
       } else { // *this is full but *mat is empty.
         mat->Resize(this->num_rows_, this->num_cols_, kUndefined);
         this->CopyToMat(mat);
         this->Destroy();
       }
     }
   } else
 #endif
   {
     std::swap(mat->data_, this->data_);
     std::swap(mat->num_cols_, this->num_cols_);
     std::swap(mat->num_rows_, this->num_rows_);
     std::swap(mat->stride_, this->stride_);
   }
 }

◆ Swap() [2/3]

void Swap ( CuMatrix< Real > * mat )

Definition at line 114 of file cu-matrix.cc.

                                              {
   std::swap(mat->data_, this->data_);
   std::swap(mat->num_cols_, this->num_cols_);
   std::swap(mat->num_rows_, this->num_rows_);
   std::swap(mat->stride_, this->stride_);
 }

◆ Swap() [3/3]

void Swap ( CuMatrix< OtherReal > * mat )

◆ Transpose()

void Transpose ( )

Definition at line 3190 of file cu-matrix.cc.

Referenced by MultiBasisComponent::BackpropagateFnc(), kaldi::TestCuMatrixTransposeNS(), kaldi::TestCuMatrixTransposeS(), kaldi::UnitTestCuMatrixAddDiagVecMat(), kaldi::UnitTestCuMatrixSymAddMat2(), and kaldi::UnitTestCuMatrixSymInvertPosDef().

                                {
   if (this->num_rows_ == 0)
     return;
   // Copy and swap for all cases.
   // No need for a separate kernel of squared matrix in-place transpose.
   // It has the same possible peak performance as copy_transpose,
   // if allocate/deallocate overhead can be ignored.
   CuMatrix<Real> tmp(*this, kTrans);
   this->Swap(&tmp);
 }

The documentation for this class was generated from the following files:

matrix/matrix-common.h
cudamatrix/cu-matrix.h
cudamatrix/cu-matrix.cc

Public Member Functions

Private Member Functions

Additional Inherited Members

Detailed Description

template<typename Real> class kaldi::CuMatrix< Real >

Constructor & Destructor Documentation

◆ CuMatrix() [1/9]

◆ CuMatrix() [2/9]

◆ CuMatrix() [3/9]

◆ CuMatrix() [4/9]

◆ CuMatrix() [5/9]

◆ CuMatrix() [6/9]

◆ CuMatrix() [7/9]

◆ CuMatrix() [8/9]

◆ CuMatrix() [9/9]

◆ ~CuMatrix()

Member Function Documentation

◆ CompObjfAndDeriv()

◆ Destroy()

◆ Mat() [1/2]

◆ Mat() [2/2]

◆ operator=() [1/3]

◆ operator=() [2/3]

◆ operator=() [3/3]

◆ Read()

◆ Resize()

◆ Swap() [1/3]

◆ Swap() [2/3]

◆ Swap() [3/3]

◆ Transpose()

template<typename Real>
class kaldi::CuMatrix< Real >