CuBlockMatrix< Real > Class Template Reference

The class CuBlockMatrix holds a vector of objects of type CuMatrix, say, M_1, M_2, . More...

#include <cu-block-matrix.h>

Collaboration diagram for CuBlockMatrix< Real >:

Classes
struct	BlockMatrixData

Public Member Functions
	CuBlockMatrix ()
	CuBlockMatrix (const std::vector< CuMatrix< Real > > &data)
	~CuBlockMatrix ()
	CuBlockMatrix (const CuBlockMatrix &other)
	Copy constructor. More...
CuBlockMatrix &	operator= (const CuBlockMatrix &other)
	Assignment operator. More...
void	Write (std::ostream &os, bool binary) const
void	Read (std::istream &is, bool binary)
MatrixIndexT	NumRows () const
MatrixIndexT	NumCols () const
MatrixIndexT	NumBlocks () const
MatrixIndexT	MaxBlockCols () const
MatrixIndexT	MaxBlockRows () const
const CuSubMatrix< Real >	Block (MatrixIndexT b) const
CuSubMatrix< Real >	Block (MatrixIndexT b)
void	AddMatMat (BaseFloat alpha, const CuMatrix< Real > &A, MatrixTransposeType transA, const CuMatrix< Real > &B, MatrixTransposeType transB, BaseFloat beta)
	Does *this = alpha A B + beta * *this, discarding elements of the product outside the block structure of the *this matrix. More...
void	CopyFromMat (const CuMatrix< Real > &M)
	Copies elements within the block structure from matrix M, discarding others. More...
void	NormalizeColumns ()
	Normalizes the columns of *this so that each one sums to one. More...
void	Swap (CuBlockMatrix *other)

Protected Attributes
CuMatrix< Real >	data_

Private Member Functions
void	FreeCudaData ()
	If using GPU and cu_data_ != NULL, free cu_data_ and set it to NULL. More...
void	SetCudaData ()
	If using GPU, allocate and set cu_data_ on the GPU to reflect "data_". More...
void	Destroy ()
	Frees and deinitializes everything. More...

Private Attributes
std::vector< BlockMatrixData >	block_data_
MatrixIndexT	num_rows_

Friends
class	CuMatrixBase< Real >

Detailed Description

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

The class CuBlockMatrix holds a vector of objects of type CuMatrix, say, M_1, M_2, .

. M_N and it represents the matrix diag(M_1, M_2, ... M_N). Note: the individual matrices do not have to be square. The reason the class is needed is mostly so that we can efficiently multiply by this block-diagonal structure in a parallel way.

If we have a GPU available, CuBlockMatrix will store a copy of the individual CuMatrix quantities M_1 .. M_N on the GPU, but their 'primary' home remains on the CPU.. what we mean by this is that while the data remains on the GPU, the "primary" version of the Matrix object that holds the pointers will remain on the CPU. We just copy it over to the GPU whenever it is changed.

Definition at line 51 of file cu-block-matrix.h.

Constructor & Destructor Documentation

CuBlockMatrix() [1/3]

CuBlockMatrix

(

)

Definition at line 35 of file cu-block-matrix.cc.

Referenced by CuBlockMatrix< Real >::Block(), and CuBlockMatrix< Real >::~CuBlockMatrix().

                                   {
#if HAVE_CUDA == 1
  cu_data_ = NULL;
#endif
}

CuBlockMatrix() [2/3]

CuBlockMatrix

(

const std::vector< CuMatrix< Real > > &

data

)

Definition at line 42 of file cu-block-matrix.cc.

References CuBlockMatrix< Real >::BlockMatrixData::col_offset, data_, KALDI_ASSERT, CuBlockMatrix< Real >::BlockMatrixData::num_cols, CuBlockMatrix< Real >::BlockMatrixData::num_rows, and CuBlockMatrix< Real >::BlockMatrixData::row_offset.

                                                                        {
#if HAVE_CUDA == 1
  cu_data_ = NULL;
#endif
  block_data_.resize(data.size());
  MatrixIndexT row_offset = 0, col_offset = 0, max_num_rows = 0;
  for (size_t b = 0; b < data.size(); b++) {
    MatrixIndexT num_rows = data[b].NumRows(), num_cols = data[b].NumCols();
    KALDI_ASSERT(num_rows > 0 && num_cols > 0);
    BlockMatrixData block_data;
    block_data.num_rows = num_rows;
    block_data.num_cols = num_cols;
    block_data.row_offset = row_offset;
    block_data.col_offset = col_offset;
    row_offset += num_rows;
    col_offset += num_cols;
    max_num_rows = std::max(max_num_rows, num_rows);
    block_data_[b] = block_data;
  }
  num_rows_ = row_offset;
  data_.Resize(max_num_rows, col_offset);
  for (int32 b = 0; b < NumBlocks(); b++)
    Block(b).CopyFromMat(data[b]);
  SetCudaData();
}

~CuBlockMatrix()

~CuBlockMatrix ( )

inline

Definition at line 59 of file cu-block-matrix.h.

References CuBlockMatrix< Real >::CuBlockMatrix(), CuBlockMatrix< Real >::Destroy(), CuBlockMatrix< Real >::operator=(), CuBlockMatrix< Real >::Read(), and CuBlockMatrix< Real >::Write().

{ Destroy(); }

CuBlockMatrix() [3/3]

CuBlockMatrix

(

const CuBlockMatrix< Real > &

other

)

Copy constructor.

Member Function Documentation

AddMatMat()

void AddMatMat	(	BaseFloat	alpha,
		const CuMatrix< Real > &	A,
		MatrixTransposeType	transA,
		const CuMatrix< Real > &	B,
		MatrixTransposeType	transB,
		BaseFloat	beta
	)

Does *this = alpha A B + beta * *this, discarding elements of the product outside the block structure of the *this matrix.

The transA and transB parameters can be used to substitute A^T for A and B^T for B, respectively.

Definition at line 213 of file cu-block-matrix.cc.

References CuMatrixBase< Real >::AddMatMat(), CU1DBLOCK, CU2DBLOCK, CuMatrixBase< Real >::Data(), KALDI_ASSERT, kaldi::kNoTrans, kaldi::kTrans, CuMatrixBase< Real >::NumCols(), CuMatrixBase< Real >::NumRows(), CuMatrixBase< Real >::Range(), CuMatrixBase< Real >::Stride(), and kaldi::swap().

Referenced by CuBlockMatrix< Real >::NumBlocks(), and kaldi::UnitTestCuBlockMatrixAddMatMat().

                    {
  MatrixIndexT A_num_rows = A.NumRows(), A_num_cols = A.NumCols(),
      A_row_stride = A.Stride(), A_col_stride = 1,
      B_num_rows = B.NumRows(), B_num_cols = B.NumCols(),
      B_row_stride = B.Stride(), B_col_stride = 1;
  if (transA == kTrans) {
    std::swap(A_num_rows, A_num_cols);
    std::swap(A_row_stride, A_col_stride);
  }
  if (transB == kTrans) {
    std::swap(B_num_rows, B_num_cols);
    std::swap(B_row_stride, B_col_stride);
  }
  KALDI_ASSERT(A_num_rows == NumRows() && B_num_cols == NumCols()
               && A_num_cols == B_num_rows);
  if (NumBlocks() == 0) return; // empty matrix.
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;

    // (x,y,z) dimensions are (block-id, row-of-block, col-of-block)
    // First some logic to choose block dims...
    // we assume (which we can, safely) that CU1DBLOCK is <= the max threads per block.
    int32 x_blocksize = std::min(CU1DBLOCK, NumBlocks()); // x dim corresponds to block-idx.
    int32 max_block_rows = MaxBlockRows(), max_block_cols = MaxBlockCols();
    int32 y_blocksize = max_block_rows;
    while (y_blocksize * x_blocksize > CU1DBLOCK || y_blocksize > CU2DBLOCK)
      y_blocksize--;
    int32 z_blocksize = max_block_cols;
    while (z_blocksize * x_blocksize * y_blocksize > CU1DBLOCK || z_blocksize > CU2DBLOCK)
      z_blocksize--;
    
    dim3 dimBlock(x_blocksize, y_blocksize, z_blocksize);
    dim3 dimGrid(n_blocks(NumBlocks(), x_blocksize),
                 n_blocks(max_block_rows, y_blocksize),
                 n_blocks(max_block_cols, z_blocksize));
    cuda_block_add_mat_mat(dimGrid, dimBlock, cu_data_, NumBlocks(),
                           A.Data(), A_num_cols, A_row_stride, A_col_stride,
                           B.Data(), B_row_stride, B_col_stride, alpha, beta);
    CU_SAFE_CALL(cudaGetLastError());    
    CuDevice::Instantiate().AccuProfile(__func__, tim);    
  } else
#endif
  {
    int32 row_offset = 0, col_offset = 0;    
    for (MatrixIndexT b = 0; b < NumBlocks(); b++) {
      CuSubMatrix<Real> this_block = Block(b);
      MatrixIndexT this_num_rows = this_block.NumRows(),
          this_num_cols = this_block.NumCols();
      CuSubMatrix<Real> A_part = (transA == kNoTrans ?
                                  A.Range(row_offset, this_num_rows,
                                          0, A.NumCols()) :
                                  A.Range(0, A.NumRows(),
                                          row_offset, this_num_rows)),
          B_part = (transB == kNoTrans ?
                    B.Range(0, B.NumRows(),
                            col_offset, this_num_cols) :
                    B.Range(col_offset, this_num_cols,
                            0, B.NumCols()));
      this_block.AddMatMat(alpha, A_part, transA, B_part, transB, beta);
      row_offset += this_num_rows;
      col_offset += this_num_cols;
    }
    KALDI_ASSERT(row_offset == NumRows() && col_offset == NumCols());
  }
}

Block() [1/2]

const CuSubMatrix< Real > Block

(

MatrixIndexT

b

)

const

Definition at line 70 of file cu-block-matrix.cc.

References data_, and KALDI_ASSERT.

Referenced by CuMatrixBase< float >::AddMatBlock(), CuMatrixBase< float >::CopyFromBlock(), and CuBlockMatrix< Real >::NumBlocks().

                                                                {
  KALDI_ASSERT(static_cast<size_t>(b) < block_data_.size());
  const BlockMatrixData &block_data = block_data_[b];
  return CuSubMatrix<Real>(data_, 0, block_data.num_rows,
                           block_data.col_offset, block_data.num_cols);
}

Block() [2/2]

CuSubMatrix< Real > Block

(

MatrixIndexT

b

)

Definition at line 78 of file cu-block-matrix.cc.

References CuBlockMatrix< Real >::block_data_, CuBlockMatrix< Real >::CuBlockMatrix(), data_, CuBlockMatrix< Real >::data_, KALDI_ASSERT, and CuBlockMatrix< Real >::num_rows_.

                                                    {
  KALDI_ASSERT(static_cast<size_t>(b) < block_data_.size());
  BlockMatrixData &block_data = block_data_[b];
  return CuSubMatrix<Real>(data_, 0, block_data.num_rows,
                           block_data.col_offset, block_data.num_cols);
}

CopyFromMat()

void CopyFromMat

(

const CuMatrix< Real > &

M

)

Copies elements within the block structure from matrix M, discarding others.

Note: this has not been implemented in a very efficient way, it's used only for testing.

Definition at line 298 of file cu-block-matrix.cc.

References CuMatrixBase< Real >::CopyFromMat(), KALDI_ASSERT, CuMatrixBase< Real >::NumCols(), and CuMatrixBase< Real >::NumRows().

Referenced by CuBlockMatrix< Real >::NumBlocks(), and kaldi::UnitTestCuBlockMatrixAddMatMat().

                                                             {
  KALDI_ASSERT(NumRows() == M.NumRows() && NumCols() == M.NumCols());
  MatrixIndexT row_offset = 0, col_offset = 0;
  for (MatrixIndexT b = 0; b < NumBlocks(); b++) {
    CuSubMatrix<Real> this_block = Block(b);
    MatrixIndexT this_num_rows = this_block.NumRows(),
        this_num_cols = this_block.NumCols();
    const CuSubMatrix<Real> src(M, row_offset, this_num_rows,
                                col_offset, this_num_cols);
    this_block.CopyFromMat(src);
    row_offset += this_num_rows;
    col_offset += this_num_cols;
  }
  KALDI_ASSERT(row_offset == NumRows() && col_offset == NumCols());
}

Destroy()

void Destroy ( )

private

Frees and deinitializes everything.

Definition at line 203 of file cu-block-matrix.cc.

References data_.

Referenced by CuBlockMatrix< Real >::~CuBlockMatrix().

                                  {
  data_.Resize(0, 0);
  block_data_.clear();
  num_rows_ = 0;  
  FreeCudaData();
}

FreeCudaData()

void FreeCudaData ( )

private

If using GPU and cu_data_ != NULL, free cu_data_ and set it to NULL.

Definition at line 106 of file cu-block-matrix.cc.

References KALDI_ERR.

                                       {
#if HAVE_CUDA == 1
  if (cu_data_ != NULL) {
    if (CuDevice::Instantiate().Enabled()) {
      CuDevice::Instantiate().Free(cu_data_);
      cu_data_ = NULL;
    } else {
      KALDI_ERR << "CuBlockMatrix: you have CUDA data pointer but "
                << "no GPU is enabled: likely code error.";
    }
  }
#endif
}

MaxBlockCols()

MatrixIndexT MaxBlockCols

(

)

const

Definition at line 285 of file cu-block-matrix.cc.

References rnnlm::i.

Referenced by CuBlockMatrix< Real >::NumBlocks().

                                                     {
  MatrixIndexT max_cols = 0;
  for (size_t i = 0; i < block_data_.size(); i++)
    max_cols = std::max(max_cols, block_data_[i].num_cols);
  return max_cols;
}

MaxBlockRows()

MatrixIndexT MaxBlockRows

(

)

const

Definition at line 293 of file cu-block-matrix.cc.

References data_.

Referenced by CuBlockMatrix< Real >::NumBlocks().

                                                     {
  return data_.NumRows();
}

NormalizeColumns()

void NormalizeColumns

(

)

Normalizes the columns of *this so that each one sums to one.

On error (e.g. inf's), will set the column to a constant value that sums to one.

Referenced by CuBlockMatrix< Real >::NumBlocks().

NumBlocks()

MatrixIndexT NumBlocks ( ) const

inline

Definition at line 75 of file cu-block-matrix.h.

References CuBlockMatrix< Real >::AddMatMat(), CuBlockMatrix< Real >::Block(), CuBlockMatrix< Real >::block_data_, CuBlockMatrix< Real >::CopyFromMat(), CuBlockMatrix< Real >::MaxBlockCols(), CuBlockMatrix< Real >::MaxBlockRows(), CuBlockMatrix< Real >::NormalizeColumns(), and CuBlockMatrix< Real >::Swap().

Referenced by CuMatrixBase< float >::AddMatBlock(), and CuMatrixBase< float >::CopyFromBlock().

{ return block_data_.size(); }

NumCols()

MatrixIndexT NumCols ( ) const

inline

Definition at line 73 of file cu-block-matrix.h.

References CuBlockMatrix< Real >::data_.

Referenced by CuMatrixBase< float >::AddMatBlock(), CuMatrixBase< float >::CopyFromBlock(), CuMatrix< float >::CuMatrix(), kaldi::UnitTestCuBlockMatrixAddMatBlock(), and kaldi::UnitTestCuBlockMatrixAddMatMat().

{ return data_.num_cols_; }

NumRows()

MatrixIndexT NumRows ( ) const

inline

Definition at line 71 of file cu-block-matrix.h.

References CuBlockMatrix< Real >::num_rows_.

Referenced by CuMatrixBase< float >::AddMatBlock(), CuMatrixBase< float >::CopyFromBlock(), CuMatrix< float >::CuMatrix(), kaldi::UnitTestCuBlockMatrixAddMatBlock(), and kaldi::UnitTestCuBlockMatrixAddMatMat().

{ return num_rows_; }

operator=()

CuBlockMatrix< Real > & operator=

(

const CuBlockMatrix< Real > &

other

)

Assignment operator.

Definition at line 96 of file cu-block-matrix.cc.

References CuBlockMatrix< Real >::block_data_, data_, CuBlockMatrix< Real >::data_, and CuBlockMatrix< Real >::num_rows_.

Referenced by CuBlockMatrix< Real >::~CuBlockMatrix().

                                                                                     {
  FreeCudaData();
  data_ = other.data_;
  block_data_ = other.block_data_;
  num_rows_ = other.num_rows_;
  SetCudaData();
  return *this;
}

Read()

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

Definition at line 173 of file cu-block-matrix.cc.

References kaldi::ExpectToken(), rnnlm::i, KALDI_ASSERT, kaldi::Peek(), and kaldi::ReadBasicType().

Referenced by kaldi::UnitTestCuBlockMatrixIO(), and CuBlockMatrix< Real >::~CuBlockMatrix().

                                                          {
  Destroy();
  int i = Peek(is, binary);
  std::vector<CuMatrix<Real> > data;
  if (i != static_cast<int>('<')) {
    // back-compatibility code so we can read the older format of
    // MixtureProbComponent.  This code should be deleted eventually.
    int32 size;
    ReadBasicType(is, binary, &size);
    KALDI_ASSERT(size >= 0);
    data.resize(size);
    for (int32 i = 0; i < size; i++)
      data[i].Read(is, binary);
  } else {
    ExpectToken(is, binary, "<CuBlockMatrix>");
    int32 size;
    ReadBasicType(is, binary, &size);
    KALDI_ASSERT(size >= 0);
    data.resize(size);
    for (int32 i = 0; i < size; i++)
      data[i].Read(is, binary);
    ExpectToken(is, binary, "</CuBlockMatrix>");    
  }

  CuBlockMatrix<Real> block_mat(data); // initializer from std::vector<CuMatrix<Real> > does
  // the main job of initialization.
  this->Swap(&block_mat);
}

SetCudaData()

void SetCudaData ( )

private

If using GPU, allocate and set cu_data_ on the GPU to reflect "data_".

Definition at line 122 of file cu-block-matrix.cc.

References CuBlockMatrixData_::col_offset, CuMatrixBase< Real >::Data(), CuMatrixBase< Real >::Dim(), KALDI_ASSERT, CuBlockMatrixData_::matrix_data, CuBlockMatrixData_::matrix_dim, CuMatrixBase< Real >::NumCols(), CuMatrixBase< Real >::NumRows(), and CuBlockMatrixData_::row_offset.

                                      {
#if HAVE_CUDA == 1
  KALDI_ASSERT(cu_data_ == NULL);
  if (block_data_.size() == 0) return; // Nothing to do.
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    std::vector<CuBlockMatrixData> tmp_cu_data(NumBlocks());
    int32 row_offset = 0, col_offset = 0;
    for (size_t b = 0; b < NumBlocks(); b++) {
      CuSubMatrix<Real> this_mat = Block(b);
      CuBlockMatrixData &this_cu_data = tmp_cu_data[b];
      this_cu_data.row_offset = row_offset;
      this_cu_data.col_offset = col_offset;
      this_cu_data.matrix_dim = this_mat.Dim();
      this_cu_data.matrix_data = static_cast<void*>(this_mat.Data());
      row_offset += this_mat.NumRows();
      col_offset += this_mat.NumCols();
    }
    size_t size = NumBlocks() * sizeof(CuBlockMatrixData);
    cu_data_ = static_cast<CuBlockMatrixData*>(
        CuDevice::Instantiate().Malloc(size));
    CU_SAFE_CALL(cudaMemcpyAsync(cu_data_, &(tmp_cu_data[0]), size, 
                                 cudaMemcpyHostToDevice, cudaStreamPerThread));
    CU_SAFE_CALL(cudaStreamSynchronize(cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile(__func__, tim);    
  }
#endif
}

Swap()

void Swap

(

CuBlockMatrix< Real > *

other

)

Definition at line 152 of file cu-block-matrix.cc.

References CuBlockMatrix< Real >::block_data_, data_, CuBlockMatrix< Real >::data_, CuBlockMatrix< Real >::num_rows_, and kaldi::swap().

Referenced by CuBlockMatrix< Real >::NumBlocks().

                                                         {
  data_.Swap(&other->data_);
  block_data_.swap(other->block_data_);
  std::swap(num_rows_, other->num_rows_);
#if HAVE_CUDA == 1
  std::swap(cu_data_, other->cu_data_);
#endif
}

Write()

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

Definition at line 162 of file cu-block-matrix.cc.

References kaldi::WriteBasicType(), and kaldi::WriteToken().

Referenced by kaldi::UnitTestCuBlockMatrixIO(), and CuBlockMatrix< Real >::~CuBlockMatrix().

                                                                 {
  WriteToken(os, binary, "<CuBlockMatrix>");
  int32 num_blocks = NumBlocks();
  WriteBasicType(os, binary, num_blocks);
  for (int32 b = 0; b < num_blocks; b++)
    this->Block(b).Write(os, binary);
  WriteToken(os, binary, "</CuBlockMatrix>");  
}

Friends And Related Function Documentation

CuMatrixBase< Real >

friend class CuMatrixBase< Real >

friend

Definition at line 53 of file cu-block-matrix.h.

Member Data Documentation

block_data_

std::vector<BlockMatrixData> block_data_

private

Definition at line 135 of file cu-block-matrix.h.

Referenced by CuBlockMatrix< Real >::Block(), CuBlockMatrix< Real >::NumBlocks(), CuBlockMatrix< Real >::operator=(), and CuBlockMatrix< Real >::Swap().

data_

CuMatrix<Real> data_

protected

Definition at line 110 of file cu-block-matrix.h.

Referenced by CuBlockMatrix< Real >::Block(), CuBlockMatrix< Real >::NumCols(), CuBlockMatrix< Real >::operator=(), and CuBlockMatrix< Real >::Swap().

num_rows_

MatrixIndexT num_rows_

private

Definition at line 137 of file cu-block-matrix.h.

Referenced by CuBlockMatrix< Real >::Block(), CuBlockMatrix< Real >::NumRows(), CuBlockMatrix< Real >::operator=(), and CuBlockMatrix< Real >::Swap().

---

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

• cudamatrix/cu-block-matrix.h
• cudamatrix/cu-block-matrix.cc