CuVectorBase< Real > Class Template Reference

Vector for CUDA computing. More...

#include <matrix-common.h>

Inheritance diagram for CuVectorBase< Real >:

Collaboration diagram for CuVectorBase< Real >:

Public Member Functions
MatrixIndexT	Dim () const
	Dimensions. More...
Real *	Data ()
	Returns a pointer to the start of the vector's data. More...
const Real *	Data () const
	Returns a pointer to the start of the vector's data (const). More...
void	CopyFromVec (const CuVectorBase< Real > &src)
	Copy functions; these will crash if the dimension do not match. More...
template<typename OtherReal >
void	CopyFromVec (const CuVectorBase< OtherReal > &M)
template<typename OtherReal >
void	CopyFromVec (const VectorBase< OtherReal > &src)
template<typename OtherReal >
void	CopyToVec (VectorBase< OtherReal > *dst) const
void	CopyRowsFromMat (const CuMatrixBase< Real > &M)
void	CopyRowsFromMat (const MatrixBase< Real > &M)
void	SetZero ()
	Math operations. More...
void	Set (Real value)
void	Add (Real value)
void	Scale (Real value)
void	AddVec (Real alpha, const CuVectorBase< Real > &vec, Real beta=1.0)
template<typename OtherReal >
void	AddVec (Real alpha, const CuVectorBase< OtherReal > &vec, Real beta=1.0)
void	AddRowSumMat (Real alpha, const CuMatrixBase< Real > &mat, Real beta=1.0)
	Sum the rows of the matrix, add to vector. More...
void	AddColSumMat (Real alpha, const CuMatrixBase< Real > &mat, Real beta=1.0)
	Sum the columns of the matrix, add to vector. More...
void	AddTpVec (const Real alpha, const CuTpMatrix< Real > &M, const MatrixTransposeType trans, const CuVectorBase< Real > &v, const Real beta)
	Add triangular matrix times vector: this <– beta*this + alpha*M*v. More...
void	MulTp (const CuTpMatrix< Real > &M, const MatrixTransposeType trans)
	Multiplies this vector by lower-triangular marix: *this <– *this *M. More...
bool	ApproxEqual (const CuVectorBase< Real > &other, float tol=0.01) const
void	InvertElements ()
void	CopyElements (const CuMatrixBase< Real > &mat, const MatrixTransposeType trans, const CuArrayBase< int32 > &elements)
	Copies selected elements from 'mat' to *this. More...
void	Floor (const CuVectorBase< Real > &src, Real floor_val, MatrixIndexT *floored_count=NULL)
void	Ceiling (const CuVectorBase< Real > &src, Real ceiling_val, MatrixIndexT *ceiled_count=NULL)
void	Pow (const CuVectorBase< Real > &src, Real power)
void	ApplyFloor (Real floor_val, MatrixIndexT *floored_count=NULL)
void	ApplyCeiling (Real ceiling_val, MatrixIndexT *ceiled_count=NULL)
void	ApplyPow (Real power)
void	ApplySoftMax ()
void	ApplyLogSoftMax ()
void	ApplyExp ()
void	ApplyLog ()
Real	Sum () const
void	SetRandn ()
void	SetRandUniform ()
CuSubVector< Real >	Range (const MatrixIndexT o, const MatrixIndexT l)
const CuSubVector< Real >	Range (const MatrixIndexT o, const MatrixIndexT l) const
void	CopyColFromMat (const CuMatrixBase< Real > &mat, MatrixIndexT col)
template<typename OtherReal >
void	CopyColFromMat (const CuMatrixBase< OtherReal > &mat, MatrixIndexT col)
void	AddMatVec (const Real alpha, const CuMatrixBase< Real > &M, MatrixTransposeType trans, const CuVectorBase< Real > &v, const Real beta)
void	AddVecVec (Real alpha, const CuVectorBase< Real > &v, const CuVectorBase< Real > &r, Real beta)
void	AddSpVec (const Real alpha, const CuSpMatrix< Real > &S, const CuVectorBase< Real > &v, const Real beta)
void	AddDiagMat2 (Real alpha, const CuMatrixBase< Real > &M, MatrixTransposeType trans, Real beta)
	Add the diagonal of a matrix times itself: *this = diag(M M^T) + beta * *this (if trans == kNoTrans), or *this = diag(M^T M) + beta * *this (if trans == kTrans). More...
void	AddDiagMatMat (Real alpha, const CuMatrixBase< Real > &M, MatrixTransposeType transM, const CuMatrixBase< Real > &N, MatrixTransposeType transN, Real beta=1.0)
	Add the diagonal of a matrix product: *this = diag(M N), assuming the "trans" arguments are both kNoTrans; for transpose arguments, it behaves as you would expect. More...
CuValue< Real >	operator() (MatrixIndexT i)
Real	Norm (Real p)
Real	operator() (MatrixIndexT i) const
void	CopyDiagFromPacked (const CuPackedMatrix< Real > &M)
	Extracts the diagonal of a packed matrix M; works for Sp or Tp. More...
void	CopyDiagFromMat (const CuMatrix< Real > &M)
	Extracts the diagonal of a matrix. More...
Real	Max () const
	Returns the maximum value of any element, or -infinity for the empty vector. More...
Real	Min () const
	Returns the minimum value of any element, or +infinity for the empty vector. More...
void	ReplaceValue (Real orig, Real changed)
void	MulElements (const CuVectorBase< Real > &v)
void	DivElements (const CuVectorBase< Real > &v)
const VectorBase< Real > &	Vec () const
VectorBase< Real > &	Vec ()
template<>
void	CopyColFromMat (const CuMatrixBase< float > &mat, MatrixIndexT col)
template<>
void	CopyColFromMat (const CuMatrixBase< double > &mat, MatrixIndexT col)
template<>
void	CopyFromVec (const CuVectorBase< float > &src)
template<>
void	CopyFromVec (const CuVectorBase< double > &src)
template<>
void	CopyFromVec (const CuVectorBase< float > &src)
template<>
void	CopyFromVec (const CuVectorBase< double > &src)

Protected Member Functions
	CuVectorBase ()
	Default constructor: make it protected so the user cannot instantiate this class. More...

Protected Attributes
Real *	data_
	GPU data pointer (or regular data pointer if CUDA is not compiled in or we have no GPU). More...
MatrixIndexT	dim_
	dimension of the vector More...

Private Member Functions
	KALDI_DISALLOW_COPY_AND_ASSIGN (CuVectorBase)

Friends
class	CuVectorBase< float >
class	CuVectorBase< double >
class	CuMatrixBase< Real >
class	MatrixBase< Real >
class	CuPackedMatrix< Real >
class	CuSpMatrix< Real >
class	CuTpMatrix< Real >
class	CuRand< Real >
template<typename OtherReal >
OtherReal	VecVec (const CuVectorBase< OtherReal > &v1, const CuVectorBase< OtherReal > &v2)
void	cu::Splice (const CuMatrixBase< Real > &src, const CuArray< int32 > &frame_offsets, CuMatrixBase< Real > *tgt)

Detailed Description

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

Vector for CUDA computing.

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

Constructor & Destructor Documentation

CuVectorBase()

CuVectorBase ( )

inlineprotected

Default constructor: make it protected so the user cannot instantiate this class.

Definition at line 246 of file cu-vector.h.

: data_(NULL), dim_(0) { }

Member Function Documentation

Add()

void Add

(

Real

value

)

Definition at line 1157 of file cu-vector.cc.

Referenced by BlockSoftmax::BackpropagateFnc(), LstmNonlinearityComponent::ConsolidateMemory(), CuVectorBase< float >::Data(), kaldi::MeanVariance(), BatchNormComponent::Propagate(), DropoutMaskComponent::Propagate(), SigmoidComponent::RepairGradients(), TanhComponent::RepairGradients(), RectifiedLinearComponent::RepairGradients(), and Xent::ReportPerClass().

                                       {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;

    dim3 dimBlock(CU1DBLOCK);
    dim3 dimGrid(n_blocks(Dim(), CU1DBLOCK));
    ::MatrixDim d = { 1, Dim(), Dim() };

    cuda_add(dimGrid, dimBlock, data_, value, d);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().Add(value);
  }
}

AddColSumMat()

void AddColSumMat	(	Real	alpha,
		const CuMatrixBase< Real > &	mat,
		Real	beta = 1.0
	)

Sum the columns of the matrix, add to vector.

Definition at line 1298 of file cu-vector.cc.

Referenced by LstmNonlinearityComponent::Backprop(), BlockSoftmax::BackpropagateFnc(), kaldi::CuVectorUnitTestAddColSumMat(), CuVectorBase< float >::Data(), CuMatrixBase< float >::DiffLogSoftmaxPerRow(), KlHmm::PropagateFnc(), MultiBasisComponent::PropagateFnc(), kaldi::TestCuVectorAddColSumMat(), kaldi::UnitTestCuVectorAddColSumMat(), kaldi::UnitTestCuVectorAddColSumMatLarge(), kaldi::nnet1::UnitTestLengthNorm(), AffineTransform::Update(), and ConvolutionalComponent::Update().

                                                 {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    KALDI_ASSERT(mat.NumRows() == Dim());

    cuda_add_col_sum_mat(mat.NumRows(), CU1DBLOCK, Data(), mat.Data(),
                         mat.Dim(), alpha, beta);
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().AddColSumMat(alpha, mat.Mat(), beta);
  }
}

AddDiagMat2()

void AddDiagMat2	(	Real	alpha,
		const CuMatrixBase< Real > &	M,
		MatrixTransposeType	trans,
		Real	beta
	)

Add the diagonal of a matrix times itself: *this = diag(M M^T) + beta * *this (if trans == kNoTrans), or *this = diag(M^T M) + beta * *this (if trans == kTrans).

Definition at line 595 of file cu-vector.cc.

Referenced by BackpropTruncationComponent::Backprop(), ClipGradientComponent::Backprop(), DiscriminativeComputation::Compute(), kaldi::CuVectorUnitTestAddDiagMat2(), kaldi::cu::DiffNormalizePerRow(), AffineComponentPreconditioned::GetScalingFactor(), OnlinePreconditioner::PreconditionDirections(), OnlinePreconditioner::PreconditionDirectionsInternal(), kaldi::nnet3::PrintParameterStats(), CuVectorBase< float >::Range(), ClipGradientComponent::RepairGradients(), NonlinearComponent::StoreBackpropStats(), kaldi::TestCuVectorAddDiagMat2(), and kaldi::TestCuVectorAddDiagMat2OnVariousShapes().

                                                                           {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    MatrixTransposeType other_trans = (trans == kTrans ? kNoTrans : kTrans);
    KALDI_ASSERT(dim_ == (trans == kNoTrans ? M.NumRows() : M.NumCols()));
    this->AddDiagMatMat(alpha, M, trans, M, other_trans, beta);
  } else
#endif
  {
    Vec().AddDiagMat2(alpha, M.Mat(), trans, beta);
  }
}

AddDiagMatMat()

void AddDiagMatMat	(	Real	alpha,
		const CuMatrixBase< Real > &	M,
		MatrixTransposeType	transM,
		const CuMatrixBase< Real > &	N,
		MatrixTransposeType	transN,
		Real	beta = 1.0
	)

Add the diagonal of a matrix product: *this = diag(M N), assuming the "trans" arguments are both kNoTrans; for transpose arguments, it behaves as you would expect.

Definition at line 611 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestAddDiagMatMat(), kaldi::cu::DiffNormalizePerRow(), CuMatrixBase< float >::DiffSoftmaxPerRow(), kaldi::nnet3::attention::GetAttentionDotProducts(), kaldi::nnet2::PreconditionDirections(), CuVectorBase< float >::Range(), RestrictedAttentionComponent::StoreStats(), kaldi::TestCuVectorAddDiagMatMat(), kaldi::TestCuVectorAddDiagMatMatShape(), kaldi::nnet3::UnitTestPreconditionDirectionsOnline(), and kaldi::nnet2::UnitTestPreconditionDirectionsOnline().

                                                                              {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;

    if (transM != transN) {
      KALDI_ASSERT(M.NumCols() == N.NumCols());
      KALDI_ASSERT(M.NumRows() == N.NumRows());
      if (transM == kNoTrans) {
        // Case 1: diag(M*N') == sum(M.*N, 2)
        // 1D grid and 1D block. One block per row of N.
        // 1D grid expands along the column of N.
        int dimBlock(CU1DBLOCK);
        int dimGrid(M.NumRows());
        cuda_add_diag_mat_mat_MNT(dimGrid, dimBlock, alpha, M.Data(), M.Dim(),
                                  N.Data(), N.Stride(), beta, data_);
      } else {
        // Case 2: diag(M'*N) == sum(M.*N, 1)
        // 16x16 or 8x32 2D block for coalesced memory access.
        // Grid shape is designed as follows,
        // 1. for small matrices, use 1D grid with only 1 row of 16x16 block,
        //    to avoid multiple kernel launch;
        // 2. for large enough matrices (no matter thin or fat),
        //    use 1- or 2-D grid so that the grid contains
        //    at least and not much larger than 'kOptNumBlocks' blocks
        //    to fully utilize the GPU;
        const int32 warpSize = 32;
        const int32 kOptNumBlocks = 512;
        const int32 tile_dim =
            (N.NumRows() < 4096 && N.NumCols() < kOptNumBlocks * warpSize) ?
                16 : 32;
        dim3 dimBlock(tile_dim, CU1DBLOCK / tile_dim);
        dim3 dimGrid(n_blocks(N.NumCols(), dimBlock.x),
                     n_blocks(N.NumRows(), dimBlock.y));
        dimGrid.y = std::min(dimGrid.y, (kOptNumBlocks - 1) / dimGrid.x + 1);
        dimGrid.y = tile_dim == 16 ? 1 : dimGrid.y;
        if (dimGrid.y > 1) {
          CuMatrix<Real> buf(dimGrid.y, N.NumCols());
          cuda_add_diag_mat_mat_MTN(dimGrid, dimBlock, Real(1), M.Data(),
                                    M.Stride(), N.Data(), N.Dim(), Real(0),
                                    buf.Data(), buf.Stride());
          this->AddRowSumMat(alpha, buf, beta);
        } else {
          cuda_add_diag_mat_mat_MTN(dimGrid, dimBlock, alpha, M.Data(),
                                    M.Stride(), N.Data(), N.Dim(), beta, data_,
                                    dim_);
        }
      }
    } else {
      KALDI_ASSERT(M.NumCols() == N.NumRows());
      KALDI_ASSERT(N.NumCols() == M.NumRows());
      if (transM == kNoTrans) {
        // Case 3: diag(M*N) == sum(M'.*N, 1)
        // 16x16 or 8x32 2D block for matrix transpose and coalesced memory access.
        // One block per 'tile_dim' columns of N.
        // 1D grid expands along the row of N.
        int tile_dim =
            sizeof(Real) == sizeof(float) && N.NumCols() >= 2048 ? 32 : 16;
        dim3 dimBlock(tile_dim, CU1DBLOCK / tile_dim);
        dim3 dimGrid(n_blocks(N.NumCols(), tile_dim));
        cuda_add_diag_mat_mat_MN(dimGrid, dimBlock, alpha, M.Data(), M.Stride(),
                                 N.Data(), N.Dim(), beta, data_);
      } else {
        // Case 4: diag(M'*N') == sum(N'.*M, 1)
        // Same kernel and config as case 3 except M and N are swapped.
        int tile_dim =
            sizeof(Real) == sizeof(float) && N.NumCols() >= 2048 ? 32 : 16;
        dim3 dimBlock(tile_dim, CU1DBLOCK / tile_dim);
        dim3 dimGrid(n_blocks(M.NumCols(), tile_dim));
        cuda_add_diag_mat_mat_MN(dimGrid, dimBlock, alpha, N.Data(), N.Stride(),
                                 M.Data(), M.Dim(), beta, data_);
      }
    }
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().AddDiagMatMat(alpha, M.Mat(), transM, N.Mat(), transN, beta);
  }
}

AddMatVec()

void AddMatVec	(	const Real	alpha,
		const CuMatrixBase< Real > &	M,
		MatrixTransposeType	trans,
		const CuVectorBase< Real > &	v,
		const Real	beta
	)

Definition at line 506 of file cu-vector.cc.

Referenced by ModelCollapser::CollapseComponentsAffine(), kaldi::CuVectorUnitTestAddMatVec(), ModelCollapser::PreMultiplyAffineParameters(), CuVectorBase< float >::Range(), kaldi::nnet2::UnitTestGenericComponentInternal(), and kaldi::VecMatVec().

                                                    {
  KALDI_ASSERT((trans == kNoTrans && M.NumCols() == v.dim_ && M.NumRows() == dim_) ||
               (trans == kTrans && M.NumRows() == v.dim_ && M.NumCols() == dim_));
  KALDI_ASSERT(&v != this);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;

    // Everything is backwards in CuBlas.  We need to reverse rows, columns,
    // transpose-ness.
    CUBLAS_SAFE_CALL(cublas_gemv(GetCublasHandle(),
                                 (trans==kTrans? CUBLAS_OP_N:CUBLAS_OP_T),
                                 M.NumCols(), M.NumRows(), alpha, M.Data(),
                                 M.Stride(), v.Data(), 1, beta, data_, 1));

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().AddMatVec(alpha,M.Mat(),trans,v.Vec(),beta);
  }
}

AddRowSumMat()

void AddRowSumMat	(	Real	alpha,
		const CuMatrixBase< Real > &	mat,
		Real	beta = 1.0
	)

Sum the rows of the matrix, add to vector.

Definition at line 1277 of file cu-vector.cc.

Referenced by LstmNonlinearityComponent::ConsolidateMemory(), kaldi::CuVectorUnitTestAddRowSumMat(), CuVectorBase< float >::Data(), Xent::Eval(), SentenceAveragingComponent::PropagateFnc(), RestrictedAttentionComponent::StoreStats(), NonlinearComponent::StoreStatsInternal(), kaldi::TestCuVectorAddRowSumMat(), kaldi::UnitTestCuVectorAddRowSumMat(), kaldi::UnitTestCuVectorAddRowSumMatLarge(), ConvolutionComponent::Update(), SentenceAveragingComponent::Update(), Convolutional1dComponent::Update(), NaturalGradientPerElementScaleComponent::Update(), and NonlinearComponent::UpdateStats().

                                                 {
  KALDI_ASSERT(mat.NumCols() == Dim());
  if (Dim() == 0)
    return;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    cuda_add_row_sum_mat(mat.NumCols(), CU1DBLOCK, Data(), mat.Data(),
                         mat.Dim(), alpha, beta);
    CU_SAFE_CALL(cudaGetLastError());
    
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else 
#endif
  {
    Vec().AddRowSumMat(alpha, mat.Mat(), beta);
  }
}

AddSpVec()

void AddSpVec	(	const Real	alpha,
		const CuSpMatrix< Real > &	S,
		const CuVectorBase< Real > &	v,
		const Real	beta
	)

Definition at line 535 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestAddSpVec(), CuVectorBase< float >::Range(), and kaldi::nnet2::UnitTestPreconditionDirections().

                                                   {
  KALDI_ASSERT(M.NumCols() == v.dim_ && M.NumRows() == dim_);
  KALDI_ASSERT(&v != this);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;

    // Note: in our opinion the CuSpMatrix represents a lower-triangular matrix, but
    // in CUBLAS, for some stupid reason, everything is reversed.
    CUBLAS_SAFE_CALL(cublas_spmv(GetCublasHandle(), CUBLAS_FILL_MODE_UPPER, Dim(),
                                 alpha, M.Data(), v.Data(), 1, beta, data_, 1));

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().AddSpVec(alpha,M.Mat(),v.Vec(),beta);
  }
}

AddTpVec()

void AddTpVec	(	const Real	alpha,
		const CuTpMatrix< Real > &	M,
		const MatrixTransposeType	trans,
		const CuVectorBase< Real > &	v,
		const Real	beta
	)

Add triangular matrix times vector: this <– beta*this + alpha*M*v.

Works even if rv == *this.

Definition at line 698 of file cu-vector.cc.

Referenced by CuVectorBase< float >::Data(), kaldi::UnitTestCuVectorAddTp(), and kaldi::UnitTestCuVectorAddTpVec().

                                                   {
  KALDI_ASSERT(dim_ == v.dim_ && dim_ == M.NumRows());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    if (beta == 0.0) {
      if (&v != this) CopyFromVec(v);
      MulTp(M, trans);
      if (alpha != 1.0) Scale(alpha);
    } else {
      CuVector<Real> tmp(v);
      tmp.MulTp(M, trans);
      if (beta != 1.0) Scale(beta);  // *this <-- beta * *this
      AddVec(alpha, tmp, 1.0);          // *this += alpha * M * v
    }
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().AddTpVec(alpha, M.Mat(), trans, v.Vec(), beta);
  }
}

AddVec() [1/2]

void AddVec	(	Real	alpha,
		const CuVectorBase< Real > &	vec,
		Real	beta = 1.0
	)

Definition at line 1237 of file cu-vector.cc.

Referenced by BatchNormComponent::Add(), CuVectorBase< float >::ApproxEqual(), LstmNonlinearityComponent::Backprop(), LstmNonlinearityComponent::ConsolidateMemory(), kaldi::CuVectorUnitTestAddDiagMatMat(), kaldi::CuVectorUnitTestAddVec(), kaldi::CuVectorUnitTestAddVecCross(), CuVectorBase< float >::Data(), SigmoidComponent::RepairGradients(), TanhComponent::RepairGradients(), RectifiedLinearComponent::RepairGradients(), Xent::ReportPerClass(), BatchNormComponent::StoreStats(), kaldi::UnitTestCuMatrixAddDiagVecMat(), and kaldi::UnitTestCuVectorAddVec().

                                           {
  KALDI_ASSERT(vec.Dim() == Dim());

#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    int32 dim = this->dim_;
    Real *data = this->data_;
    const Real *vec_data = vec.data_;
    if (beta != 1.0) CU_SAFE_CALL(cuda_scal(GetCublasHandle(), dim, beta, data, 1));
    if (alpha != 0.0) CU_SAFE_CALL(cuda_axpy(GetCublasHandle(), dim, alpha, vec_data, 1, data, 1));
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
  #endif
  {
    if (beta != 1.0) Vec().Scale(beta);
    Vec().AddVec(alpha, vec.Vec());
  }
}

AddVec() [2/2]

void AddVec	(	Real	alpha,
		const CuVectorBase< OtherReal > &	vec,
		Real	beta = 1.0
	)

Definition at line 1261 of file cu-vector.cc.

                                           {
  // We could implement this directly, without using a temporary-- this can
  // be done later, when we have time.
  CuVector<Real> temp(vec);
  this->AddVec(alpha, temp, beta);
}

AddVecVec()

void AddVecVec	(	Real	alpha,
		const CuVectorBase< Real > &	v,
		const CuVectorBase< Real > &	r,
		Real	beta
	)

Definition at line 560 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestAddVecVec(), BatchNormComponent::Propagate(), CuVectorBase< float >::Range(), BatchNormComponent::Read(), and BatchNormComponent::Write().

                                                                           {
  KALDI_ASSERT((dim_ == v.dim_ && dim_ == r.dim_));
  KALDI_ASSERT(this != &v && this != &r);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_,CU1DBLOCK));

    cuda_add_vec_vec(dimGrid, dimBlock, alpha, data_, v.Data(), r.Data(), beta, dim_);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile("CuVectorBase::AddVecVec", tim);
  } else
#endif
  {
    Vec().AddVecVec(alpha, v.Vec(), r.Vec(), beta);
  }
}

ApplyCeiling()

void ApplyCeiling ( Real  ceiling_val, MatrixIndexT *  ceiled_count = NULL  )

inline

Definition at line 143 of file cu-vector.h.

Referenced by kaldi::CuVectorUnitTestApplyCeiling(), kaldi::CuVectorUnitTestApplyCeilingNoCount(), kaldi::TestCuVectorApplyCeiling(), and kaldi::TestCuVectorApplyCeilingNoCount().

                                                                                {
    this -> Ceiling(*this, ceiling_val, ceiled_count);
  };

ApplyExp()

void ApplyExp

(

)

Definition at line 442 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow(), StatisticsPoolingComponent::Backprop(), and kaldi::CuVectorUnitTestApplyExp().

                                  {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_,CU1DBLOCK));

    cuda_vec_apply_exp(dimGrid, dimBlock, data_, dim_);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile("CuVectorBase::ApplyExp", tim);
  } else
#endif
  {
    Vec().ApplyExp();
  }
}

ApplyFloor()

void ApplyFloor ( Real  floor_val, MatrixIndexT *  floored_count = NULL  )

inline

Definition at line 139 of file cu-vector.h.

Referenced by kaldi::CuVectorUnitTestApplyFloor(), kaldi::CuVectorUnitTestApplyFloorNoCount(), kaldi::cu::DiffNormalizePerRow(), BatchNormComponent::Propagate(), SingleUtteranceNnet2DecoderThreaded::RunNnetEvaluationInternal(), kaldi::TestCuVectorApplyFloor(), kaldi::TestCuVectorApplyFloorNoCount(), AffineTransform::Update(), and ConvolutionalComponent::Update().

                                                                             {
    this -> Floor(*this, floor_val, floored_count);
  };

ApplyLog()

void ApplyLog

(

)

Definition at line 462 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow(), and kaldi::CuVectorUnitTestApplyLog().

                                  {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_,CU1DBLOCK));

    CuVector<Real> flag(1);
    cuda_vec_apply_log(dimGrid, dimBlock, data_, flag.Data(), dim_);
    CU_SAFE_CALL(cudaGetLastError());
    if (flag(0) > 0)
      KALDI_ERR << "Trying to take log of a negative number.";
    CuDevice::Instantiate().AccuProfile("CuVectorBase::ApplyLog", tim);
  } else
#endif
  {
    Vec().ApplyLog();
  }
}

ApplyLogSoftMax()

void ApplyLogSoftMax

(

)

Definition at line 484 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow().

                                         {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    size_t dimBlock = CU1DBLOCK;
    size_t dimGrid = 1;       // dimGrid value represent the number of rows
    ::MatrixDim dim = { 1, this->dim_, this->dim_};
    
    cuda_log_softmax_reduce(dimGrid, dimBlock, data_, data_, dim, this->dim_);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().ApplyLogSoftMax();
  }
}

ApplyPow()

void ApplyPow ( Real  power )

inline

Definition at line 147 of file cu-vector.h.

Referenced by kaldi::CuVectorUnitTestApplyPow(), AffineComponentPreconditionedOnline::GetScalingFactor(), kaldi::MeanVariance(), BatchNormComponent::Propagate(), RectifiedLinearComponent::RepairGradients(), Xent::ReportPerClass(), kaldi::nnet1::UnitTestLengthNorm(), and AffineTransform::Update().

                                   {
    this -> Pow(*this, power);
  };

ApplySoftMax()

void ApplySoftMax

(

)

Definition at line 334 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow(), kaldi::CuVectorUnitTestApplySoftMax(), and kaldi::TestCuVectorSoftmax().

                                      {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    size_t dimBlock = CU1DBLOCK;
    size_t dimGrid = 1;       // dimGrid value represent the number of rows
    ::MatrixDim dim = { 1, this->dim_, this->dim_};
    cuda_softmax_reduce(dimGrid, dimBlock, data_, data_, dim, this->dim_);//actually dim is not stride...
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().ApplySoftMax();
  }
}

ApproxEqual()

bool ApproxEqual	(	const CuVectorBase< Real > &	other,
		float	tol = 0.01
	)		const

Definition at line 583 of file cu-vector.cc.

Referenced by CuVectorBase< float >::Data(), and kaldi::nnet2::UnitTestPreconditionDirections().

                                                                                     {
  if (dim_ != other.dim_) KALDI_ERR << "ApproxEqual: size mismatch "
                                    << dim_ << " vs. " << other.dim_;
  KALDI_ASSERT(tol >= 0.0);
  CuVector<Real> tmp(*this);
  tmp.AddVec(-1.0, other);
  BaseFloat tmp_norm = sqrt(VecVec(tmp, tmp)), this_norm = sqrt(VecVec(*this, *this));
  return tmp_norm <= static_cast<Real>(tol) * this_norm;
}

Ceiling()

void Ceiling	(	const CuVectorBase< Real > &	src,
		Real	ceiling_val,
		MatrixIndexT *	ceiled_count = NULL
	)

Definition at line 385 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyCeiling(), and CuVectorBase< float >::Data().

                                                             {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_,CU1DBLOCK));
    if (ceiled_count == nullptr) {
      if (dim_ == 0) return;
      CuTimer tim;
      // We are calling a function meant for matrices, by viewing the
      // vector as a matrix with a single row.
      ::MatrixDim dim = {1, Dim(), 1};
      cuda_ceiling(dimGrid, dimBlock, this->data_, src.Data(), ceiling_val, dim, 1);

      CuDevice::Instantiate().AccuProfile("CuVectorBase::CeilingNoCount", tim);
    } else {
      if (dim_ == 0) { *ceiled_count = 0; return; }
      CuTimer tim;

      CuVector<float> count_vec(dim_, kUndefined);

      cuda_vec_apply_ceiling(dimGrid, dimBlock, data_, ceiling_val, count_vec.Data(), dim_);
      CU_SAFE_CALL(cudaGetLastError());
      *ceiled_count = count_vec.Sum();
      CuDevice::Instantiate().AccuProfile("CuVectorBase::Ceiling", tim);
    }
  } else
#endif
  {
    Vec().Ceiling(src.Vec(), ceiling_val, ceiled_count);
  }
}

CopyColFromMat() [1/4]

void CopyColFromMat	(	const CuMatrixBase< float > &	mat,
		MatrixIndexT	col
	)

Definition at line 122 of file cu-vector.cc.

                                                                                          {
  KALDI_ASSERT(col < mat.NumCols());
  KALDI_ASSERT(dim_ == mat.NumRows());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_,CU1DBLOCK));

    cuda_copy_col_from_mat_df(dimGrid, dimBlock, data_, col, mat.Data(), mat.Dim(), dim_);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile("CuVectorBase::CopyColFromMat", tim);
  } else
#endif
  {
    Vec().CopyColFromMat(mat.Mat(), col);
  }
}

CopyColFromMat() [2/4]

void CopyColFromMat	(	const CuMatrixBase< double > &	mat,
		MatrixIndexT	col
	)

Definition at line 144 of file cu-vector.cc.

                                                                                          {
  KALDI_ASSERT(col < mat.NumCols());
  KALDI_ASSERT(dim_ == mat.NumRows());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_,CU1DBLOCK));

    cuda_copy_col_from_mat_fd(dimGrid, dimBlock, data_, col, mat.Data(), mat.Dim(), dim_);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile("CuVectorBase::CopyColFromMat", tim);
  } else
#endif
  {
    Vec().CopyColFromMat(mat.Mat(), col);
  }
}

CopyColFromMat() [3/4]

void CopyColFromMat	(	const CuMatrixBase< Real > &	mat,
		MatrixIndexT	col
	)

Definition at line 103 of file cu-vector.cc.

Referenced by StatisticsPoolingComponent::Backprop(), kaldi::CuVectorUnitTestCopyFromMat(), CuVectorBase< float >::Range(), kaldi::UnitTestCuMatrixAddDiagVecMat(), NaturalGradientAffineComponent::Update(), AffineComponentPreconditioned::Update(), AffineComponentPreconditionedOnline::Update(), BlockAffineComponentPreconditioned::Update(), and TdnnComponent::UpdateNaturalGradient().

                                                                                       {
  KALDI_ASSERT(col < mat.NumCols());
  KALDI_ASSERT(dim_ == mat.NumRows());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    cublas_copy(GetCublasHandle(),
                this->dim_, mat.Data() + col, mat.Stride(), this->data_, 1);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile("CuVectorBase::CopyColFromMat", tim);
  } else
#endif
  {
    Vec().CopyColFromMat(mat.Mat(),col);
  }
}

CopyColFromMat() [4/4]

void CopyColFromMat	(	const CuMatrixBase< OtherReal > &	mat,
		MatrixIndexT	col
	)

CopyDiagFromMat()

void CopyDiagFromMat

(

const CuMatrix< Real > &

M

)

Extracts the diagonal of a matrix.

Definition at line 1198 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestAddDiagMatMat(), kaldi::CuVectorUnitTestCopyDiagFromMat(), and CuVectorBase< float >::operator()().

                                                                {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    KALDI_ASSERT(dim_ == std::min(M.NumRows(), M.NumCols()));
    CuTimer tim;
    CUBLAS_SAFE_CALL(cublas_copy(GetCublasHandle(), dim_, M.Data(), M.Stride() + 1,
                                 data_, 1));

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().CopyDiagFromMat(M.Mat());
  }
}

CopyDiagFromPacked()

void CopyDiagFromPacked

(

const CuPackedMatrix< Real > &

M

)

Extracts the diagonal of a packed matrix M; works for Sp or Tp.

Definition at line 1177 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestCopyDiagFromPacked(), CuVectorBase< float >::operator()(), CuPackedMatrix< Real >::Trace(), and kaldi::TraceSpSp().

                                                                         {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    KALDI_ASSERT(dim_ == M.NumRows());
    if (dim_ == 0) return;
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(Dim(), CU1DBLOCK));
    cuda_vec_copy_diag_from_packed(dimGrid, dimBlock, data_, M.Data(), dim_);
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().CopyDiagFromPacked(M.Mat());
  }
}

CopyElements()

void CopyElements	(	const CuMatrixBase< Real > &	mat,
		const MatrixTransposeType	trans,
		const CuArrayBase< int32 > &	elements
	)

Copies selected elements from 'mat' to *this.

Expects this->Dim() to equal elements.Dim(). If trans == kNoTrans, expects mat.NumRows() to equal this.Dim(), and for each i, copies mat(i, elements[i]) to (*this)(i). If trans == kTrans, expects mat.NumCols() to equal this.Dim(), and for each i, copies mat(elements[i], i) to (*this)(i).

Definition at line 1340 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestCopyElements(), and CuVectorBase< float >::Data().

                                                                          {
  KALDI_ASSERT(elements.Dim() == Dim());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;

    dim3 dimBlock(CU1DBLOCK);
    dim3 dimGrid(n_blocks(Dim(), CU1DBLOCK));

    cuda_vector_copy_elements(dimGrid, dimBlock, this->data_, Dim(),
                                mat.Data(), mat.Stride(), trans == kTrans,
                                elements.Data());
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    VectorBase<Real> &this_vec = this->Vec();
    const MatrixBase<Real> &src_mat = mat.Mat();
    const int32* index_map = elements.Data();
    KALDI_ASSERT((Dim() == mat.NumRows() && trans == kNoTrans)
                 || (Dim() == mat.NumCols() && trans == kTrans));
    for (int32 i = 0; i < Dim(); i++) {
      int32 j = index_map[i];
      KALDI_ASSERT(j >= 0);
      if (trans == kNoTrans) {
        KALDI_ASSERT(j < mat.NumCols());
        this_vec(i) = src_mat(i, j);
      } else {
        KALDI_ASSERT(j < mat.NumRows());
        this_vec(i) = src_mat(j, i);
      }
    }
  }
}

CopyFromVec() [1/7]

void CopyFromVec

(

const CuVectorBase< Real > &

src

)

Copy functions; these will crash if the dimension do not match.

The operator = in class CuVector will also change the sizes for you.

Definition at line 1078 of file cu-vector.cc.

Referenced by BlockAffineComponent::BlockAffineComponent(), CuSparseMatrix< Real >::CopyElementsToVec(), CuSparseMatrix< Real >::CopyFromSmat(), CuVectorBase< float >::CopyFromVec(), CuMatrixBase< float >::CopyRangeFromMatClamped(), CuSparseMatrix< Real >::CuSparseMatrix(), CuVector< float >::CuVector(), kaldi::CuVectorUnitTestCopyCross(), kaldi::CuVectorUnitTestCopyCross2(), CuVectorBase< float >::Data(), SoftmaxComponent::MixUp(), CuVector< float >::operator=(), OnlinePreconditionerSimple::PreconditionDirections(), OnlineNaturalGradientSimple::PreconditionDirections(), BatchNormComponent::Propagate(), CuRand< float >::RandGaussian(), kaldi::nnet1::RandUniform(), DctComponent::Reorder(), RectifiedLinearComponent::RepairGradients(), NnetBatchComputer::SplitUtteranceIntoTasks(), kaldi::TestCuVectorCopyFromVec(), kaldi::UnitTestCuMatrixAddDiagVecMat(), kaldi::UnitTestCuMatrixAddVecToCols(), kaldi::UnitTestCuMatrixAddVecToRows(), kaldi::UnitTestCuMatrixAddVecVec(), kaldi::UnitTestCuMatrixDivRowsVec(), kaldi::UnitTestCuMatrixMulColsVec(), kaldi::UnitTestCuMatrixMulRowsVec(), kaldi::UnitTestCuVectorAddColSumMat(), kaldi::UnitTestCuVectorAddColSumMatLarge(), kaldi::UnitTestCuVectorAddRowSumMat(), kaldi::UnitTestCuVectorAddRowSumMatLarge(), kaldi::UnitTestCuVectorAddTpVec(), kaldi::UnitTestCuVectorAddVec(), kaldi::UnitTestCuVectorCopyFromVec(), kaldi::UnitTestCuVectorInvertElements(), and kaldi::UnitTestCuVectorMulTp().

                                                                  {
  KALDI_ASSERT(src.Dim() == dim_);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    CU_SAFE_CALL(
      cudaMemcpyAsync(data_, src.data_, src.dim_ * sizeof(Real), 
                      cudaMemcpyDeviceToDevice, cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
  #endif
  {
    memcpy(static_cast<void*>(data_), static_cast<void*>(src.data_),
           dim_ * sizeof(Real));
  }
}

CopyFromVec() [2/7]

void CopyFromVec

(

const CuVectorBase< OtherReal > &

M

)

Definition at line 377 of file cu-vector.h.

                                                                     {
  v.CopyToVec(&this);
}

CopyFromVec() [3/7]

void CopyFromVec

(

const VectorBase< OtherReal > &

src

)

Definition at line 904 of file cu-vector.cc.

                                                                     {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (sizeof(Real) != sizeof(OtherReal)) {
      CuVector<OtherReal> temp(dim_, kUndefined);
      temp.CopyFromVec(src);
      this->CopyFromVec(temp);
    } else {
      KALDI_ASSERT(src.Dim() == dim_);
      if (dim_ == 0) return;
      CuTimer tim;
      CU_SAFE_CALL(cudaMemcpyAsync(data_, src.Data(), src.Dim()*sizeof(Real), 
                                   cudaMemcpyHostToDevice, cudaStreamPerThread));
      CU_SAFE_CALL(cudaStreamSynchronize(cudaStreamPerThread));
      CuDevice::Instantiate().AccuProfile("CuVector::CopyFromVecH2D", tim);
    }
  } else
  #endif
  {
    Vec().CopyFromVec(src);
  }
}

CopyFromVec() [4/7]

void CopyFromVec< float >

(

const CuVectorBase< float > &

src

)

CopyFromVec() [5/7]

void CopyFromVec< double >

(

const CuVectorBase< double > &

src

)

CopyFromVec() [6/7]

void CopyFromVec

(

const CuVectorBase< float > &

src

)

Definition at line 869 of file cu-vector.cc.

                                                                     {
  KALDI_ASSERT(src.Dim() == dim_);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    CUBLAS_SAFE_CALL(cublas_copy(GetCublasHandle(), dim_, src.Data(), 1, data_, 1));
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().CopyFromVec(src.Vec());
  }
}

CopyFromVec() [7/7]

void CopyFromVec

(

const CuVectorBase< double > &

src

)

Definition at line 886 of file cu-vector.cc.

                                                                     {
  KALDI_ASSERT(src.Dim() == dim_);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    CUBLAS_SAFE_CALL(cublas_copy(GetCublasHandle(), dim_, src.Data(), 1, data_, 1));
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().CopyFromVec(src.Vec());
  }
}

CopyRowsFromMat() [1/2]

void CopyRowsFromMat

(

const CuMatrixBase< Real > &

M

)

Definition at line 164 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestCopyFromMat(), and CuVectorBase< float >::Data().

                                                                      {
  KALDI_ASSERT(dim_ == mat.NumCols() * mat.NumRows());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    if (mat.Stride() == mat.NumCols() && mat.NumRows() != 0) {
      CU_SAFE_CALL(
        cudaMemcpyAsync(data_, mat.Data(), sizeof(Real)*dim_,
                        cudaMemcpyDeviceToDevice, cudaStreamPerThread));
    } else {
      Real* vec_data = data_;
      for (MatrixIndexT r = 0; r < mat.NumRows(); r++) {
        CU_SAFE_CALL(cudaMemcpyAsync(vec_data, mat.RowData(r),
                                     sizeof(Real) * mat.NumCols(),
                                     cudaMemcpyDeviceToDevice,
                                     cudaStreamPerThread));
        vec_data += mat.NumCols();
      }
    }
    CuDevice::Instantiate().AccuProfile("CuVectorBase::CopyRowsFromMat", tim);
  } else
#endif
  {
    Vec().CopyRowsFromMat(mat.Mat());
  }
}

CopyRowsFromMat() [2/2]

void CopyRowsFromMat

(

const MatrixBase< Real > &

M

)

Definition at line 218 of file cu-vector.cc.

                                                                    {
  KALDI_ASSERT(dim_ == mat.NumCols() * mat.NumRows());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    if (mat.Stride() == mat.NumCols()) {
      CU_SAFE_CALL(cudaMemcpyAsync(data_, mat.Data(), sizeof(Real)*dim_,
                              cudaMemcpyHostToDevice, cudaStreamPerThread));
    } else {
      Real* vec_data = data_;
      for (MatrixIndexT r = 0; r < mat.NumRows(); r++) {
        CU_SAFE_CALL(cudaMemcpyAsync(vec_data, mat.RowData(r),
                                sizeof(Real) * mat.NumCols(),
                                cudaMemcpyHostToDevice, cudaStreamPerThread));
        vec_data += mat.NumCols();
      }
    }
    CU_SAFE_CALL(cudaStreamSynchronize(cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().CopyRowsFromMat(mat);
  }
}

CopyToVec()

template void CopyToVec

(

VectorBase< OtherReal > *

dst

)

const

Definition at line 938 of file cu-vector.cc.

Referenced by CuSparseMatrix< Real >::CopyToSmat(), CuVectorBase< float >::CopyToVec(), kaldi::nnet1::CountCorrectFramesWeighted(), CuVectorBase< float >::Data(), kaldi::nnet1::MomentStatistics(), kaldi::UnitTestCuDiffXent(), kaldi::UnitTestCuVectorAddColSumMat(), kaldi::UnitTestCuVectorAddColSumMatLarge(), kaldi::UnitTestCuVectorAddRowSumMat(), kaldi::UnitTestCuVectorAddRowSumMatLarge(), kaldi::UnitTestCuVectorAddTpVec(), kaldi::UnitTestCuVectorAddVec(), kaldi::UnitTestCuVectorInvertElements(), kaldi::UnitTestCuVectorMulTp(), kaldi::UnitTestMulTp(), and kaldi::UnitTestVector().

                                                                   {
  KALDI_ASSERT(dim_ == dst->Dim());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (sizeof(Real) != sizeof(OtherReal)) {
      CuVector<OtherReal> temp(*this);
      temp.CopyToVec(dst);
    } else {
      if (dim_ == 0) return;
      CuTimer tim;
      CU_SAFE_CALL(cudaMemcpyAsync(dst->Data(), this->data_,
                              sizeof(Real) * dim_, cudaMemcpyDeviceToHost,
                              cudaStreamPerThread));
      CU_SAFE_CALL(cudaStreamSynchronize(cudaStreamPerThread));
      CuDevice::Instantiate().AccuProfile(__func__, tim);
    }
  } else
#endif
  {
    dst->CopyFromVec(this->Vec());
  }
}

Data() [1/2]

Real* Data ( )

inline

Returns a pointer to the start of the vector's data.

Definition at line 72 of file cu-vector.h.

Referenced by CuMatrixBase< float >::AddDiagVecMat(), CuMatrixBase< float >::AddMatDiagVec(), CuVectorBase< float >::AddMatVec(), CuVectorBase< float >::AddSpVec(), CuSpMatrix< Real >::AddVec2(), CuVectorBase< float >::AddVecVec(), CuMatrixBase< float >::AddVecVec(), CuVectorBase< float >::ApplyLog(), StatisticsPoolingComponent::Backprop(), kaldi::cu::BackpropLstmNonlinearity(), CuVectorBase< float >::Ceiling(), CuMatrix< float >::CompObjfAndDeriv(), CuMatrixBase< float >::CopyColFromVec(), CuMatrixBase< float >::CopyColsFromVec(), CuVectorBase< float >::CopyFromVec(), MatrixBase< float >::CopyRowsFromVec(), CuMatrixBase< float >::CopyRowsFromVec(), CuVectorBase< float >::DivElements(), kaldi::cu::EnsureNonzero(), CuVectorBase< float >::Floor(), NnetBatchComputer::FormatInputs(), CuVectorBase< float >::Max(), CuMatrixBase< float >::Max(), CuVectorBase< float >::Min(), CuMatrixBase< float >::Min(), CuVectorBase< float >::MulElements(), CuVectorBase< float >::Pow(), CuRand< float >::RandGaussian(), CuRand< float >::RandUniform(), NnetBatchComputer::SplitUtteranceIntoTasks(), RestrictedAttentionComponent::StoreStats(), CuVectorBase< float >::Sum(), CuMatrixBase< float >::Sum(), kaldi::TestCuVectorCopyFromVec(), CuMatrixBase< float >::Trace(), kaldi::TraceMatMat(), and kaldi::VecVec().

{ return data_; }

Data() [2/2]

const Real* Data ( ) const

inline

Returns a pointer to the start of the vector's data (const).

Definition at line 74 of file cu-vector.h.

{ return data_; }

Dim()

MatrixIndexT Dim ( ) const

inline

Dimensions.

Definition at line 69 of file cu-vector.h.

Referenced by DiscriminativeObjectiveInfo::AccumulateGradients(), DiscriminativeObjectiveInfo::AccumulateOutput(), NonlinearComponent::Add(), CuMatrixBase< float >::AddDiagVecMat(), CuMatrixBase< float >::AddMatDiagVec(), NnetStats::AddStatsFromNnet(), CuVectorBase< float >::AddVec(), CuSpMatrix< Real >::AddVec2(), CuMatrixBase< float >::AddVecToCols(), CuMatrixBase< float >::AddVecToRows(), CuMatrixBase< float >::AddVecVec(), AffineComponent::AffineComponent(), kaldi::cu::BackpropLstmNonlinearity(), Convolutional1dComponent::Convolutional1dComponent(), ConvolutionComponent::ConvolutionComponent(), CuMatrixBase< float >::CopyColFromVec(), CuMatrixBase< float >::CopyColsFromVec(), CuSparseMatrix< Real >::CopyElementsToVec(), CuVectorBase< float >::CopyFromVec(), MatrixBase< float >::CopyRowsFromVec(), CuMatrixBase< float >::CopyRowsFromVec(), kaldi::nnet1::CountCorrectFramesWeighted(), kaldi::CuRandGaussianVectorSpeedTest(), kaldi::CuRandUniformVectorSpeedTest(), CuVector< float >::CuVector(), CuMatrixBase< float >::DivRowsVec(), kaldi::cu::EnsureNonzero(), Xent::Eval(), ModelCollapser::GetDiagonallyPreModifiedComponentIndex(), AffineComponentPreconditionedOnline::GetScalingFactor(), FixedScaleComponent::Init(), FixedBiasComponent::Init(), PerElementScaleComponent::Init(), kaldi::MeanVariance(), SoftmaxComponent::MixUp(), kaldi::nnet1::MomentStatistics(), CuMatrixBase< float >::MulColsVec(), CuMatrixBase< float >::MulRowsVec(), NaturalGradientAffineComponent::NaturalGradientAffineComponent(), CuVector< float >::operator=(), CuMatrixBase< float >::ParametricRelu(), ModelCollapser::PreMultiplyAffineParameters(), kaldi::nnet3::PrintParameterStats(), CuRand< float >::RandGaussian(), CuRand< float >::RandUniform(), kaldi::nnet1::RandUniform(), AffineTransform::SetBias(), NnetBatchComputer::SplitUtteranceIntoTasks(), BatchNormComponent::StoreStats(), kaldi::TestCuVectorCopyFromVec(), kaldi::VecMatVec(), and kaldi::VecVec().

{ return dim_;  }

DivElements()

void DivElements

(

const CuVectorBase< Real > &

v

)

Definition at line 857 of file cu-vector.cc.

Referenced by CuVectorBase< float >::operator()().

                                                                {
  // this just creates a matrix and calls the matrix version.
  KALDI_ASSERT(dim_ == v.dim_);
  CuSubMatrix<Real> this_mat(this->Data(), 1, dim_, dim_),
      v_mat(v.Data(), 1, dim_, dim_);
  this_mat.DivElements(v_mat);
}

Floor()

void Floor	(	const CuVectorBase< Real > &	src,
		Real	floor_val,
		MatrixIndexT *	floored_count = NULL
	)

Definition at line 353 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyFloor(), and CuVectorBase< float >::Data().

                                                                                                         {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_,CU1DBLOCK));
    if (floored_count == nullptr) {
      if (dim_ == 0) return;
      CuTimer tim;
      // We are calling a function meant for matrices, by viewing the
      // vector as a matrix with a single row.
      ::MatrixDim dim = {1, Dim(), 1};
      cuda_floor(dimGrid, dimBlock, this->data_, src.Data(), floor_val, dim, 1);
      CuDevice::Instantiate().AccuProfile("CuVectorBase::FloorNoCount", tim);
    } else {
      if (dim_ == 0) { *floored_count = 0; return; }
      CuTimer tim;

      CuVector<float> count_vec(dim_, kUndefined);

      cuda_vec_apply_floor(dimGrid, dimBlock, data_, floor_val, count_vec.Data(), dim_);
      CU_SAFE_CALL(cudaGetLastError());
      *floored_count = count_vec.Sum();
      CuDevice::Instantiate().AccuProfile("CuVectorBase::Floor", tim);
    }
  } else
#endif
  {
    Vec().Floor(src.Vec(), floor_val, floored_count);
  }
}

InvertElements()

void InvertElements

(

)

Definition at line 1318 of file cu-vector.cc.

Referenced by ScaleAndOffsetComponent::BackpropInternal(), kaldi::CuVectorUnitTestInvertElements(), CuVectorBase< float >::Data(), kaldi::UnitTestCuVectorInvertElements(), AffineTransform::Update(), and ConvolutionalComponent::Update().

                                        {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;

    dim3 dimBlock(CU1DBLOCK, 1);
    dim3 dimGrid(n_blocks(dim_, CU1DBLOCK));
    MatrixDim d = {1, dim_, dim_};

    cuda_invert_elements(dimGrid, dimBlock, data_, d);
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().InvertElements();
  }
}

KALDI_DISALLOW_COPY_AND_ASSIGN()

KALDI_DISALLOW_COPY_AND_ASSIGN ( CuVectorBase< Real >  )

private

Max()

Real Max

(

)

const

Returns the maximum value of any element, or -infinity for the empty vector.

Definition at line 782 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestMax(), ModelCollapser::GetDiagonallyPreModifiedComponentIndex(), CuMatrixBase< float >::Max(), and CuVectorBase< float >::operator()().

                                   {
  Real result = 0.0;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) {  // max of an empty set is -infinity.
      return -std::numeric_limits<Real>::infinity();
    }
    CuTimer tim;

    // Small vectors are copied to RAM and reduced on CPU.
    // The length is chosen by cu-vector-speed-test
    if (dim_ < 4096) {
      Vector<Real> ans_cpu(*this);
      result = ans_cpu.Max();
    } else {
      // Use no more than 256 blocks (still too many?)
      int dimBlock = CU1DBLOCK;
      int dimGrid = n_blocks(dim_, dimBlock);
      if (dimGrid > 256) {
        dimGrid = 256;
      }
      CuVector<Real> ans(dimGrid, kUndefined);
      cuda_vec_max(dimGrid, dimBlock, data_, ans.Data(), dim_, 1);
      CU_SAFE_CALL(cudaGetLastError());
      Vector<Real> ans_cpu(ans);
      result = ans_cpu.Max();
    }

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    result = (this->Vec()).Max();
  }
  return result;
}

Min()

Real Min

(

)

const

Returns the minimum value of any element, or +infinity for the empty vector.

Definition at line 744 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestMin(), ModelCollapser::GetDiagonallyPreModifiedComponentIndex(), CuMatrixBase< float >::Min(), and CuVectorBase< float >::operator()().

                                   {
  Real result = 0.0;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) {  // min of an empty set is infinity.
      return std::numeric_limits<Real>::infinity();
    }
    CuTimer tim;

    // Small vectors are copied to RAM and reduced on CPU.
    // The length is chosen by cu-vector-speed-test
    if (dim_ < 4096) {
      Vector<Real> ans_cpu(*this);
      result = ans_cpu.Min();
    } else {
      // Use no more than 256 blocks (still too many?)
      int dimBlock = CU1DBLOCK;
      int dimGrid = n_blocks(dim_, dimBlock);
      if (dimGrid > 256) {
        dimGrid = 256;
      }
      CuVector<Real> ans(dimGrid, kUndefined);
      cuda_vec_min(dimGrid, dimBlock, data_, ans.Data(), dim_, 1);
      CU_SAFE_CALL(cudaGetLastError());
      Vector<Real> ans_cpu(ans);
      result = ans_cpu.Min();
    }

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    result = (this->Vec()).Min();
  }
  return result;
}

MulElements()

void MulElements

(

const CuVectorBase< Real > &

v

)

Definition at line 838 of file cu-vector.cc.

Referenced by BackpropTruncationComponent::Backprop(), kaldi::CuVectorUnitTestInvertElements(), AffineComponentPreconditionedOnline::GetScalingFactor(), CuVectorBase< float >::operator()(), Xent::ReportPerClass(), and kaldi::UnitTestVector().

                                                                {
  KALDI_ASSERT(dim_ == v.dim_);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_, CU1DBLOCK));
    cuda_vec_mul_elements(dimGrid, dimBlock, data_, v.Data(), dim_);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile("CuVectorBase::MulElements", tim);
  } else
#endif
  {
    Vec().MulElements(v.Vec());
  }
}

MulTp()

void MulTp	(	const CuTpMatrix< Real > &	M,
		const MatrixTransposeType	trans
	)

Multiplies this vector by lower-triangular marix: *this <– *this *M.

Definition at line 727 of file cu-vector.cc.

Referenced by CuVectorBase< float >::AddTpVec(), CuVectorBase< float >::Data(), kaldi::UnitTestCuVectorMulTp(), and kaldi::UnitTestMulTp().

                                                                                         {
  KALDI_ASSERT(M.NumRows() == dim_);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    cublas_tpmv(GetCublasHandle(), (trans==kTrans? CUBLAS_OP_N:CUBLAS_OP_T),
                M.NumRows(), M.Data(), data_, 1);
    CuDevice::Instantiate().AccuProfile("CuVectorBase::MulTp", tim);
  } else
#endif
  {
    Vec().MulTp(M.Mat(), trans);
  }
}

Norm()

Real Norm

(

Real

p

)

Definition at line 193 of file cu-vector.cc.

Referenced by kaldi::CuVectorUnitTestNorm(), CuSparseMatrix< Real >::FrobeniusNorm(), and CuVectorBase< float >::operator()().

                                    {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    Real ans;
    KALDI_ASSERT(p == 1.0 || p == 2.0);
    if (dim_ == 0) return 0.0;
    if (p == 1.0) {
      cublas_asum(GetCublasHandle(), dim_, data_, 1, &ans);
    } else {
      cublas_nrm2(GetCublasHandle(), dim_, data_, 1, &ans);
    }
    CuDevice::Instantiate().AccuProfile(__func__, tim);
    if (ans != ans) {
      KALDI_ERR << "NaN in norm " << *this;
    }
    return ans;
  } else
#endif
  {
    return Vec().Norm(p);
  }
}

operator()() [1/2]

CuValue<Real> operator() ( MatrixIndexT  i )

inline

Definition at line 196 of file cu-vector.h.

                                                   {
    KALDI_PARANOID_ASSERT(static_cast<UnsignedMatrixIndexT>(i) <
                          static_cast<UnsignedMatrixIndexT>(dim_));
    return CuValue<Real>(data_ + i);
  }

operator()() [2/2]

Real operator() ( MatrixIndexT  i ) const

inline

Definition at line 204 of file cu-vector.h.

                                                {
    KALDI_PARANOID_ASSERT(static_cast<UnsignedMatrixIndexT>(i) <
                          static_cast<UnsignedMatrixIndexT>(dim_));
    return CuValue<Real>(data_ + i); // will be casted to Real.
  }

Pow()

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

Definition at line 419 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow(), and CuVectorBase< float >::Data().

                                                                      {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    // for this particular kernel, x is #rows, y is #cols.  so
    // fake matrix with 1 row, Dim() cols.
    dim3 dimBlock(CU1DBLOCK, 1);
    dim3 dimGrid(n_blocks(Dim(), CU1DBLOCK), 1);
    ::MatrixDim fake_matrix_dim = { 1, Dim(), 1 };
    // num_cols is Dim(), num_rows is 1, stride is 1 (it's a don't-care).
    cuda_pow(dimGrid, dimBlock, this->data_, src.Data(), power, fake_matrix_dim, 1);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile("CuVectorBase::ApplyPow", tim);
  } else
#endif
  {
    Vec().Pow(src.Vec(), power);
  }
}

Range() [1/2]

CuSubVector<Real> Range ( const MatrixIndexT  o, const MatrixIndexT  l  )

inline

Definition at line 160 of file cu-vector.h.

Referenced by ModelCollapser::CollapseComponentsAffine(), CuRand< float >::RandGaussian(), and kaldi::UnitTestCuSubVector().

                                                                      {
    return CuSubVector<Real>(*this, o, l);
  }

Range() [2/2]

const CuSubVector<Real> Range ( const MatrixIndexT  o, const MatrixIndexT  l  ) const

inline

Definition at line 164 of file cu-vector.h.

                                                            {
    return CuSubVector<Real>(*this, o, l);
  }

ReplaceValue()

void ReplaceValue	(	Real	orig,
		Real	changed
	)

Definition at line 820 of file cu-vector.cc.

Referenced by CuVectorBase< float >::operator()(), and kaldi::UnitTestCuVectorReplaceValue().

                                                             {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (dim_ == 0) return;
    CuTimer tim;
    int dimBlock(CU1DBLOCK);
    int dimGrid(n_blocks(dim_, CU1DBLOCK));
    cuda_replace_value(dimGrid, dimBlock, data_, dim_, orig, changed);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().ReplaceValue(orig, changed);
  }
}

Scale()

void Scale

(

Real

value

)

Definition at line 1216 of file cu-vector.cc.

Referenced by BlockSoftmax::BackpropagateFnc(), kaldi::CuVectorUnitTestAddDiagMatMat(), kaldi::CuVectorUnitTestScale(), CuVectorBase< float >::Data(), kaldi::nnet2::PreconditionDirections(), DropoutMaskComponent::Propagate(), BatchNormComponent::Read(), RectifiedLinearComponent::RepairGradients(), Xent::ReportPerClass(), BatchNormComponent::Scale(), kaldi::UnitTestCuMatrixAddDiagVecMat(), kaldi::UnitTestVector(), AffineTransform::Update(), ConvolutionalComponent::Update(), and BatchNormComponent::Write().

                                         {
  #if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    if (Dim() == 0 ) return;

    CuTimer tim;
    dim3 dimBlock(CU1DBLOCK);
    dim3 dimGrid(n_blocks(Dim(), CU1DBLOCK));
    ::MatrixDim d = { 1, Dim(), Dim() };
    cuda_scale(dimGrid, dimBlock, data_, value, d);
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
  #endif
  {
    Vec().Scale(value);
  }
}

Set()

void Set

(

Real

value

)

Definition at line 1135 of file cu-vector.cc.

Referenced by CuSparseMatrix< Real >::CuSparseMatrix(), kaldi::CuVectorUnitTestSum(), CuVectorBase< float >::Data(), FixedScaleComponent::InitFromConfig(), StatisticsPoolingComponent::Propagate(), CuSparseMatrix< Real >::Resize(), kaldi::TestCuVectorVecVecOne(), kaldi::nnet1::UnitTestLengthNorm(), and kaldi::nnet1::UnitTestSimpleSentenceAveragingComponent().

                                       {
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;

    dim3 dimBlock(CU1DBLOCK);
    dim3 dimGrid(n_blocks(Dim(), CU1DBLOCK));
    ::MatrixDim d = { 1, Dim(), Dim() };

    cuda_set_const(dimGrid, dimBlock, data_, value, d);
    CU_SAFE_CALL(cudaGetLastError());
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    Vec().Set(value);
  }
}

SetRandn()

void SetRandn

(

)

Definition at line 281 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow(), LstmNonlinearityComponent::ConsolidateMemory(), kaldi::CuVectorUnitTestAddDiagMat2(), kaldi::CuVectorUnitTestAddDiagMatMat(), kaldi::CuVectorUnitTestAddMatVec(), kaldi::CuVectorUnitTestAddSpVec(), kaldi::CuVectorUnitTestAddVec(), kaldi::CuVectorUnitTestAddVecCross(), kaldi::CuVectorUnitTestAddVecExtra(), kaldi::CuVectorUnitTestAddVecVec(), kaldi::CuVectorUnitTestApplyCeiling(), kaldi::CuVectorUnitTestApplyCeilingNoCount(), kaldi::CuVectorUnitTestApplyExp(), kaldi::CuVectorUnitTestApplyFloor(), kaldi::CuVectorUnitTestApplyFloorNoCount(), kaldi::CuVectorUnitTestApplyLog(), kaldi::CuVectorUnitTestApplyPow(), kaldi::CuVectorUnitTestApplySoftMax(), kaldi::CuVectorUnitTestApproxEqual(), kaldi::CuVectorUnitTestCopyCross(), kaldi::CuVectorUnitTestCopyCross2(), kaldi::CuVectorUnitTestCopyElements(), kaldi::CuVectorUnitTestInvertElements(), kaldi::CuVectorUnitTestMax(), kaldi::CuVectorUnitTestMin(), kaldi::CuVectorUnitTestScale(), kaldi::CuVectorUnitTestVecVec(), FixedScaleComponent::InitFromConfig(), FixedBiasComponent::InitFromConfig(), ConvolutionComponent::PerturbParams(), TimeHeightConvolutionComponent::PerturbParams(), TdnnComponent::PerturbParams(), RepeatedAffineComponent::PerturbParams(), ConstantComponent::PerturbParams(), AffineComponent::PerturbParams(), BlockAffineComponent::PerturbParams(), PerElementScaleComponent::PerturbParams(), PerElementOffsetComponent::PerturbParams(), ConstantFunctionComponent::PerturbParams(), Convolutional1dComponent::PerturbParams(), ScaleAndOffsetComponent::PerturbParams(), kaldi::TestCuMatrixAddDiagVecMat(), kaldi::TestCuMatrixDivRowsVec(), kaldi::TestCuVectorAddColSumMat(), kaldi::TestCuVectorAddDiagMat2(), kaldi::TestCuVectorAddDiagMat2OnVariousShapes(), kaldi::TestCuVectorAddDiagMatMat(), kaldi::TestCuVectorAddDiagMatMatShape(), kaldi::TestCuVectorAddRowSumMat(), kaldi::TestCuVectorApplyCeiling(), kaldi::TestCuVectorApplyCeilingNoCount(), kaldi::TestCuVectorApplyFloor(), kaldi::TestCuVectorApplyFloorNoCount(), kaldi::TestCuVectorCopyFromVec(), kaldi::TestCuVectorSoftmax(), kaldi::TestCuVectorSum(), kaldi::TestCuVectorVecVecOne(), kaldi::UnitTestBackpropLstmNonlinearity(), kaldi::UnitTestCuMatrixAddDiagVecMat(), kaldi::UnitTestCuMatrixAddMatDiagVec(), kaldi::UnitTestCuMatrixCopyColsFromVec(), kaldi::UnitTestCuMatrixCopyRowsFromVec(), kaldi::UnitTestCuSubVector(), kaldi::UnitTestCuVectorIO(), kaldi::nnet2::UnitTestFixedBiasComponent(), kaldi::nnet2::UnitTestFixedScaleComponent(), kaldi::nnet2::UnitTestGenericComponentInternal(), kaldi::UnitTestVecMatVec(), and kaldi::UnitTestVector().

                                  {
  if (dim_ == 0) return;
  CuRand<Real> tmp;
  tmp.RandGaussian(this);
}

SetRandUniform()

void SetRandUniform

(

)

Definition at line 288 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow().

                                        {
  if (dim_ == 0) return;
  CuRand<Real> tmp;
  tmp.RandUniform(this);
}

SetZero()

void SetZero

(

)

Math operations.

Definition at line 1098 of file cu-vector.cc.

Referenced by StatisticsPoolingComponent::Backprop(), CuVectorBase< float >::Data(), DiscriminativeObjectiveInfo::Reset(), BatchNormComponent::Scale(), and BatchNormComponent::ZeroStats().

                                 {
  if (dim_==0 || data_==NULL) return;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    KALDI_ASSERT(dim_>=0);
    KALDI_ASSERT(data_!=NULL);
    CuTimer tim;
    CU_SAFE_CALL(cudaMemsetAsync(data_, 0, dim_*sizeof(Real),
          cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile("CuVector::SetZero", tim);
  } else
#endif
  {
    Vec().SetZero();
  }
}

Sum()

Real Sum

(

)

const

Definition at line 297 of file cu-vector.cc.

Referenced by CuVectorBase< float >::ApplyPow(), CuVectorBase< float >::Ceiling(), LstmNonlinearityComponent::ConsolidateMemory(), kaldi::CuVectorUnitTestCopyDiagFromMat(), Xent::Eval(), CuVectorBase< float >::Floor(), AffineComponentPreconditionedOnline::GetScalingFactor(), kaldi::MeanVariance(), SoftmaxComponent::MixUp(), OnlinePreconditioner::PreconditionDirectionsInternal(), kaldi::nnet3::PrintParameterStats(), SigmoidComponent::RepairGradients(), TanhComponent::RepairGradients(), RectifiedLinearComponent::RepairGradients(), Xent::Report(), CuSparseMatrix< Real >::Sum(), CuMatrixBase< float >::Sum(), kaldi::TestCuVectorSum(), CuPackedMatrix< Real >::Trace(), and kaldi::UnitTestCuMatrixAddDiagVecMat().

                                   {
  if (dim_ == 0)
    return 0.0;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    Real result;
    CuTimer tim;

    // Small vectors are copied to RAM and reduced on CPU.
    // The length is chosen by cu-vector-speed-test
    if (dim_ < 4096) {
      Vector<Real> ans_cpu(*this);
      result = ans_cpu.Sum();
    } else {
      // Use no more than 256 blocks (still too many?)
      int dimBlock = CU1DBLOCK;
      int dimGrid = n_blocks(dim_, dimBlock);
      if (dimGrid > 256) {
        dimGrid = 256;
      }
      CuVector<Real> ans(dimGrid, kUndefined);
      cuda_vec_sum(dimGrid, dimBlock, data_, ans.Data(), dim_, 1);
      CU_SAFE_CALL(cudaGetLastError());
      Vector<Real> ans_cpu(ans);
      result = ans_cpu.Sum();
    }

    CuDevice::Instantiate().AccuProfile(__func__, tim);
    return result;
  } else
#endif
  {
    return Vec().Sum();
  }
}

Vec() [1/2]

const VectorBase<Real>& Vec ( ) const

inline

Definition at line 235 of file cu-vector.h.

Referenced by CuMatrixBase< float >::AddDiagVecMat(), CuMatrixBase< float >::AddMatDiagVec(), CuVectorBase< float >::AddMatVec(), CuVectorBase< float >::AddSpVec(), CuVectorBase< float >::AddTpVec(), CuVectorBase< float >::AddVec(), CuSpMatrix< Real >::AddVec2(), CuMatrixBase< float >::AddVecToCols(), CuMatrixBase< float >::AddVecToRows(), CuVectorBase< float >::AddVecVec(), CuMatrixBase< float >::AddVecVec(), kaldi::cu::BackpropLstmNonlinearity(), CuVectorBase< float >::Ceiling(), CuMatrixBase< float >::CopyColFromVec(), CuMatrixBase< float >::CopyColsFromVec(), CuSparseMatrix< Real >::CopyElementsToVec(), CuVectorBase< float >::CopyFromVec(), MatrixBase< float >::CopyRowsFromVec(), CuMatrixBase< float >::CopyRowsFromVec(), CuSparseMatrix< Real >::CuSparseMatrix(), CuMatrixBase< float >::DiffParametricRelu(), CuMatrixBase< float >::DiffXent(), CuMatrixBase< float >::DivRowsVec(), CuVectorBase< float >::Floor(), CuMatrixBase< float >::MulColsVec(), CuVectorBase< float >::MulElements(), CuMatrixBase< float >::MulRowsVec(), CuMatrixBase< float >::ParametricRelu(), CuVectorBase< float >::Pow(), CuRand< float >::RandGaussian(), CuRand< float >::RandUniform(), and kaldi::VecVec().

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

Vec() [2/2]

VectorBase<Real>& Vec ( )

inline

Definition at line 238 of file cu-vector.h.

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

Friends And Related Function Documentation

cu::Splice

void cu::Splice ( const CuMatrixBase< Real > &  src, const CuArray< int32 > &  frame_offsets, CuMatrixBase< Real > *  tgt  )

friend

CuMatrixBase< Real >

friend class CuMatrixBase< Real >

friend

Definition at line 54 of file cu-vector.h.

CuPackedMatrix< Real >

friend class CuPackedMatrix< Real >

friend

Definition at line 56 of file cu-vector.h.

CuRand< Real >

friend class CuRand< Real >

friend

Definition at line 66 of file cu-vector.h.

CuSpMatrix< Real >

friend class CuSpMatrix< Real >

friend

Definition at line 57 of file cu-vector.h.

CuTpMatrix< Real >

friend class CuTpMatrix< Real >

friend

Definition at line 58 of file cu-vector.h.

CuVectorBase< double >

friend class CuVectorBase< double >

friend

Definition at line 53 of file cu-vector.h.

CuVectorBase< float >

friend class CuVectorBase< float >

friend

Definition at line 52 of file cu-vector.h.

MatrixBase< Real >

friend class MatrixBase< Real >

friend

Definition at line 55 of file cu-vector.h.

VecVec

OtherReal VecVec ( const CuVectorBase< OtherReal > &  v1, const CuVectorBase< OtherReal > &  v2  )

friend

Member Data Documentation

data_

Real* data_

protected

GPU data pointer (or regular data pointer if CUDA is not compiled in or we have no GPU).

Definition at line 248 of file cu-vector.h.

Referenced by CuVectorBase< float >::AddVec(), CuMatrixBase< float >::AddVecToCols(), CuMatrixBase< float >::AddVecToRows(), CuVectorBase< float >::CopyFromVec(), CuSubVector< Real >::CuSubVector(), CuVectorBase< float >::Data(), CuMatrixBase< float >::DiffParametricRelu(), CuMatrixBase< float >::DiffXent(), CuMatrixBase< float >::DivRowsVec(), CuMatrixBase< float >::MulColsVec(), CuMatrixBase< float >::MulRowsVec(), CuVectorBase< float >::operator()(), CuMatrixBase< float >::ParametricRelu(), and CuVector< float >::Swap().

dim_

MatrixIndexT dim_

protected

dimension of the vector

Definition at line 250 of file cu-vector.h.

Referenced by CuVectorBase< float >::AddMatVec(), CuVectorBase< float >::AddSpVec(), CuVectorBase< float >::AddTpVec(), CuVectorBase< float >::AddVecVec(), CuVectorBase< float >::ApproxEqual(), CuVectorBase< float >::CopyFromVec(), CuSubVector< Real >::CuSubVector(), CuVectorBase< float >::Dim(), CuVectorBase< float >::DivElements(), CuVectorBase< float >::MulElements(), CuVectorBase< float >::operator()(), and CuVector< float >::Swap().

---

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

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