CuArrayBase< T > Class Template Reference

Class CuArrayBase, CuSubArray and CuArray are analogues of classes CuVectorBase, CuSubVector and CuVector, except that they are intended to store things other than float/double: they are intended to store integers or small structs. More...

#include <cu-array.h>

Inheritance diagram for CuArrayBase< T >:

Collaboration diagram for CuArrayBase< T >:

Public Member Functions
MatrixIndexT	Dim () const
	Return the vector dimension. More...
const T *	Data () const
	Get raw pointer. More...
T *	Data ()
void	SetZero ()
	Sets the memory for the object to zero, via memset. More...
void	CopyFromArray (const CuArrayBase< T > &src)
	The caller is responsible to ensure dim is equal between *this and src. More...
void	CopyFromVec (const std::vector< T > &src)
	The caller is responsible to ensure dim is equal between *this and src. More...
void	CopyToVec (std::vector< T > *dst) const
	This function resizes *dst if needed. More...
void	CopyToHost (T *dst) const
	Version of the above function that copies contents to a host array (i.e. More...
void	Set (const T &value)
	Set to a constant value. More...
void	Sequence (const T base)
	Fill with the sequence [base ... More...
void	Add (const T &value)
	Add a constant value. More...
T	Min () const
	Get minimum value (for now implemented on CPU, reimplement if slow). More...
T	Max () const
	Get minimum value (for now implemented on CPU, reimplement if slow). More...
template<>
void	Set (const int32 &value)
template<>
void	Sequence (const int32 base)
template<>
void	Add (const int32 &value)
template<>
void	Sequence (const int32 base)
template<>
void	Set (const int32 &value)
template<>
void	Add (const int32 &value)

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

Protected Attributes
T *	data_
	GPU data pointer (if GPU not available, will point to CPU memory). More...
MatrixIndexT	dim_
	dimension of the vector More...

Friends
class	CuArray< T >
class	CuSubArray< T >

Detailed Description

template<typename T>
class kaldi::CuArrayBase< T >

Class CuArrayBase, CuSubArray and CuArray are analogues of classes CuVectorBase, CuSubVector and CuVector, except that they are intended to store things other than float/double: they are intended to store integers or small structs.

Their CPU-based equivalents are std::vector, and we provide ways to copy to/from a std::vector of the same type.

Definition at line 44 of file cu-array.h.

Constructor & Destructor Documentation

CuArrayBase()

CuArrayBase ( )

inlineprotected

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

Definition at line 108 of file cu-array.h.

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

Member Function Documentation

Add() [1/3]

void Add

(

const int32 &

value

)

Definition at line 88 of file cu-array.cc.

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

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

    cuda_int32_add(dimGrid, dimBlock, data_, value, d);
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    for (int32 i = 0; i < dim_; i++) {
      data_[i] += value;
    }
  }
}

Add() [2/3]

void Add

(

const T &

value

)

Add a constant value.

This is NOT IMPLEMENTED YET except for T == int32 (the current implementation will just crash).

Definition at line 254 of file cu-array-inl.h.

Referenced by CuArrayBase< Int32Pair >::Add(), CuArrayBase< Int32Pair >::Data(), and CopyComponent::WriteData().

                                       {
  // This is not implemented yet, we'll do so if it's needed.
  KALDI_ERR << "CuArray<T>::Add not implemented yet for this type.";
}

Add() [3/3]

void Add

(

const int32 &

value

)

CopyFromArray()

void CopyFromArray

(

const CuArrayBase< T > &

src

)

The caller is responsible to ensure dim is equal between *this and src.

Note: copying to GPU is done via memcpy, and any constructors or assignment operators are not called.

Definition at line 157 of file cu-array-inl.h.

Referenced by CuSparseMatrix< Real >::CopyFromSmat(), CuArray< Int32Pair >::CuArray(), CuSparseMatrix< Real >::CuSparseMatrix(), CuArrayBase< Int32Pair >::Data(), CuArray< Int32Pair >::operator=(), and CuArray< Int32Pair >::~CuArray().

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

CopyFromVec()

void CopyFromVec

(

const std::vector< T > &

src

)

The caller is responsible to ensure dim is equal between *this and src.

Note: copying to GPU is done via memcpy, and any constructors or assignment operators are not called.

Definition at line 100 of file cu-array-inl.h.

Referenced by CuSparseMatrix< Real >::CopyFromSmat(), CuArray< Int32Pair >::CuArray(), CuArrayBase< Int32Pair >::Data(), CuArray< Int32Pair >::operator=(), and CuArray< Int32Pair >::~CuArray().

                                                        {
  KALDI_ASSERT(dim_ == src.size());
  if (src.empty())
    return;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    CU_SAFE_CALL(
        cudaMemcpyAsync(data_, &src.front(), src.size() * sizeof(T),
                   cudaMemcpyHostToDevice, cudaStreamPerThread));
    CU_SAFE_CALL(cudaStreamSynchronize(cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    memcpy(data_, &src.front(), src.size() * sizeof(T));
  }
}

CopyToHost()

void CopyToHost

(

T *

dst

)

const

Version of the above function that copies contents to a host array (i.e.

to regular memory, not GPU memory, assuming we're using a GPU). This function requires *dst to be allocated before calling. The allocated size should be dim_ * sizeof(T)

Definition at line 198 of file cu-array-inl.h.

Referenced by CuArrayBase< Int32Pair >::Data(), CuMatrixBase< float >::Lookup(), and kaldi::UnitTestCuArray().

                                            {
  if (this->dim_ == 0) return;
  KALDI_ASSERT(dst != NULL);
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    CU_SAFE_CALL(cudaMemcpyAsync(dst, Data(), this->dim_ * sizeof(T),
          cudaMemcpyDeviceToHost, cudaStreamPerThread));
    CU_SAFE_CALL(cudaStreamSynchronize(cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile("CuArray::CopyToVecD2H", tim);
  } else
#endif
  {
    memcpy(dst, this->data_, this->dim_ * sizeof(T));
  }
}

CopyToVec()

void CopyToVec

(

std::vector< T > *

dst

)

const

This function resizes *dst if needed.

On resize of "dst", the STL vector may call copy-constructors, initializers, and assignment operators for existing objects (which will be overwritten), but the copy from GPU to CPU is done via memcpy. So be very careful calling this function if your objects are more than plain structs.

Definition at line 177 of file cu-array-inl.h.

Referenced by Splice::BackpropagateFnc(), ConvolutionComputation::Check(), kaldi::nnet3::ComputeAccuracy(), kaldi::nnet3::CopyPairVector(), CuSparseMatrix< Real >::CopyToSmat(), kaldi::nnet1::CountCorrectFramesWeighted(), CuSparseMatrix< Real >::CuSparseMatrix(), CuArrayBase< Int32Pair >::Data(), CuSparseMatrix< Real >::SelectRows(), kaldi::TestCuFindRowMaxId(), kaldi::UnitTestCuArray(), and kaldi::UnitTestCuFindRowMaxId().

                                                      {
  if (static_cast<MatrixIndexT>(dst->size()) != this->dim_) {
    dst->resize(this->dim_);
  }
  if (this->dim_ == 0) return;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    CU_SAFE_CALL(cudaMemcpyAsync(&dst->front(), Data(), this->dim_ * sizeof(T),
          cudaMemcpyDeviceToHost, cudaStreamPerThread));
    CU_SAFE_CALL(cudaStreamSynchronize(cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile("CuArray::CopyToVecD2H", tim);
  } else
#endif
  {
    memcpy(&dst->front(), this->data_, this->dim_ * sizeof(T));
  }
}

Data() [1/2]

const T* Data ( ) const

inline

Get raw pointer.

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

Referenced by CuMatrixBase< float >::AddCols(), CuMatrixBase< float >::AddElements(), CuMatrixBase< float >::AddRowRanges(), CuMatrixBase< float >::AddRows(), CuMatrixBase< float >::AddToElements(), CuMatrixBase< float >::AddToRows(), kaldi::cu::Copy(), CuMatrixBase< float >::CopyCols(), CuVectorBase< float >::CopyElements(), CuMatrixBase< float >::CopyRows(), CuMatrixBase< float >::CopyToRows(), CuMatrixBase< float >::DiffXent(), CuMatrixBase< float >::FindRowMaxId(), CuMatrixBase< float >::Lookup(), CuMatrixBase< float >::MulRows(), kaldi::cu::Randomize(), CuSparseMatrix< Real >::SelectRows(), kaldi::cu::Splice(), and CuMatrixBase< float >::SumColumnRanges().

{ return data_; }

Data() [2/2]

T* Data ( )

inline

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

{ return data_; }

Dim()

MatrixIndexT Dim ( ) const

inline

Return the vector dimension.

Definition at line 49 of file cu-array.h.

Referenced by CuMatrixBase< float >::AddCols(), CuMatrixBase< float >::AddElements(), CuMatrixBase< float >::AddRowRanges(), CuMatrixBase< float >::AddRows(), CuMatrixBase< float >::AddToElements(), CuMatrixBase< float >::AddToRows(), Splice::BackpropagateFnc(), ConvolutionComputation::Check(), kaldi::nnet3::time_height_convolution::ConvolveBackwardDataInternal(), kaldi::nnet3::time_height_convolution::ConvolveBackwardParamsInternal(), kaldi::nnet3::time_height_convolution::ConvolveForwardInternal(), kaldi::cu::Copy(), CuMatrixBase< float >::CopyCols(), CuVectorBase< float >::CopyElements(), CuArrayBase< Int32Pair >::CopyFromArray(), CuArray< Int32Pair >::CopyFromArray(), CuMatrixBase< float >::CopyRows(), CuMatrixBase< float >::CopyToRows(), kaldi::nnet1::CountCorrectFramesWeighted(), CuSparseMatrix< Real >::CuSparseMatrix(), CuSubArray< T >::CuSubArray(), CuMatrixBase< float >::DiffXent(), Splice::InitData(), CuMatrixBase< float >::Lookup(), CuMatrixBase< float >::MulRows(), StatisticsExtractionComponent::Propagate(), StatisticsPoolingComponent::Propagate(), kaldi::cu::Randomize(), Splice::ReadData(), CuSparseMatrix< Real >::SelectRows(), kaldi::cu::Splice(), CuMatrixBase< float >::SumColumnRanges(), and kaldi::UnitTestCuArray().

{ return dim_;  }

Max()

T Max ( ) const

inline

Get minimum value (for now implemented on CPU, reimplement if slow).

Asserts the vector is non-empty, otherwise crashes.

Definition at line 282 of file cu-array-inl.h.

Referenced by CuArrayBase< Int32Pair >::Data().

                                   {
  KALDI_ASSERT(this->Dim() > 0);
#if HAVE_CUDA == 1
  CuTimer tim;
#endif
  std::vector<T> tmp(Dim());
  CopyToVec(&tmp);
  T ans = *std::max_element(tmp.begin(), tmp.end());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  }
#endif
  return ans;
}

Min()

T Min ( ) const

inline

Get minimum value (for now implemented on CPU, reimplement if slow).

Asserts the vector is non-empty, otherwise crashes.

Definition at line 264 of file cu-array-inl.h.

Referenced by CuArrayBase< Int32Pair >::Data().

                                   {
  KALDI_ASSERT(this->Dim() > 0);
#if HAVE_CUDA == 1
  CuTimer tim;
#endif
  std::vector<T> tmp(Dim());
  CopyToVec(&tmp);
  T ans = *std::min_element(tmp.begin(), tmp.end());
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuDevice::Instantiate().AccuProfile(__func__, tim);
  }
#endif
  return ans;
}

Sequence() [1/3]

void Sequence

(

const int32

base

)

Definition at line 39 of file cu-array.cc.

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

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

    cuda_sequence(dimGrid, dimBlock, Data(), Dim(), base);
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    for (int32 i = 0; i < dim_; i++) {
      data_[i] = base + i;
    }
  }
}

Sequence() [2/3]

void Sequence

(

const T

base

)

Fill with the sequence [base ...

base + Dim()) This is not implemented except for T=int32

Definition at line 244 of file cu-array-inl.h.

Referenced by CuSparseMatrix< Real >::CuSparseMatrix(), CuArrayBase< Int32Pair >::Data(), and CuArrayBase< Int32Pair >::Sequence().

                                          {
  // This is not implemented yet, we'll do so if it's needed.
  KALDI_ERR << "CuArray<T>::Sequence not implemented yet for this type.";
}

Sequence() [3/3]

void Sequence

(

const int32

base

)

Set() [1/3]

void Set

(

const int32 &

value

)

Definition at line 63 of file cu-array.cc.

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

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

    cuda_int32_set_const(dimGrid, dimBlock, data_, value, d);
    CU_SAFE_CALL(cudaGetLastError());

    CuDevice::Instantiate().AccuProfile(__func__, tim);
  } else
#endif
  {
    for (int32 i = 0; i < dim_; i++) {
      data_[i] = value;
    }
  }
}

Set() [2/3]

void Set

(

const T &

value

)

Set to a constant value.

Note: any copying is done as if using memcpy, and assignment operators or destructors are not called. This is NOT IMPLEMENTED YET except for T == int32 (the current implementation will just crash).

Definition at line 234 of file cu-array-inl.h.

Referenced by CuArrayBase< Int32Pair >::Data(), CuSparseMatrix< Real >::Resize(), CuArrayBase< Int32Pair >::Set(), and kaldi::UnitTestCuArray().

                                       {
  // This is not implemented yet, we'll do so if it's needed.
  KALDI_ERR << "CuArray<T>::Set not implemented yet for this type.";
}

Set() [3/3]

void Set

(

const int32 &

value

)

SetZero()

void SetZero

(

)

Sets the memory for the object to zero, via memset.

You should verify that this makes sense for type T.

Definition at line 217 of file cu-array-inl.h.

Referenced by CuArrayBase< Int32Pair >::Data().

                             {
  if (this->dim_ == 0) return;
#if HAVE_CUDA == 1
  if (CuDevice::Instantiate().Enabled()) {
    CuTimer tim;
    CU_SAFE_CALL(cudaMemsetAsync(this->data_, 0, this->dim_ * sizeof(T),
          cudaStreamPerThread));
    CuDevice::Instantiate().AccuProfile("CuArray::SetZero", tim);
  } else
#endif
  {
    memset(static_cast<void*>(this->data_), 0, this->dim_ * sizeof(T));
  }
}

Friends And Related Function Documentation

CuArray< T >

friend class CuArray< T >

friend

Definition at line 45 of file cu-array.h.

CuSubArray< T >

friend class CuSubArray< T >

friend

Definition at line 46 of file cu-array.h.

Member Data Documentation

data_

T* data_

protected

GPU data pointer (if GPU not available, will point to CPU memory).

Definition at line 111 of file cu-array.h.

Referenced by CuArrayBase< Int32Pair >::CopyFromArray(), CuArray< Int32Pair >::CopyFromArray(), CuSubArray< T >::CuSubArray(), CuArrayBase< Int32Pair >::Data(), and CuArray< Int32Pair >::Swap().

dim_

MatrixIndexT dim_

protected

dimension of the vector

Definition at line 113 of file cu-array.h.

Referenced by CuSubArray< T >::CuSubArray(), CuArrayBase< Int32Pair >::Dim(), and CuArray< Int32Pair >::Swap().

---

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

• cudamatrix/cu-array.h
• cudamatrix/cu-array-inl.h