cudamatrix/cublas-wrappers.h
该头文件对cuBLAS的接口进行了简单的封装(函数名的简化和部分kaldi函数的封装)。
比如
cublasSgemm_v2封装为cublas_gemm
cublas_copy_kaldi_fd和cublas_copy_kaldi_df封装为cublas_copy
cudamatrix/cu-kernels.{h,cu}
以cuda_add_col_sum_mat函数为例
对Kaldi cuda kernel或cublas进行了简单的封装(针对不同精度浮点型)
cudamatrix/cu-kernels.h inline void cuda_add_col_sum_mat(int Gr, int Bl, double* result, const double* mat, const MatrixDim d, const double alpha, const double beta) { cudaD_add_col_sum_mat(Gr, Bl, result, mat, d, alpha, beta); } inline void cuda_add_col_sum_mat(int Gr, int Bl, float* result, const float* mat, const MatrixDim d, const float alpha, const float beta) { cudaF_add_col_sum_mat(Gr, Bl, result, mat, d, alpha, beta); } //... } |
kernel的定义
cudamatrix/cu-kernels.cu // Reduce a matrix 'mat' to a column vector 'result' template<EnumTransformReduce TransReduceType, typename Real> __global__ static void _transform_reduce_mat_cols( Real *result, const Real *mat, const MatrixDim d, const TransReduceOp<TransReduceType, Real> op) { __shared__ Real sdata[CU1DBLOCK]; const int tid = threadIdx.x; const int i = blockIdx.x; const int row_start = i * d.stride; Real tdata = op.InitValue(); for (int j = tid; j < d.cols; j += CU1DBLOCK) { tdata = op.Reduce(tdata, op.Transform(mat[row_start + j])); } sdata[tid] = tdata; __syncthreads(); // Tree reduce # pragma unroll for (int shift = CU1DBLOCK / 2; shift > warpSize; shift >>= 1) { if (tid < shift) sdata[tid] = op.Reduce(sdata[tid], sdata[tid + shift]); __syncthreads(); } // Reduce last warp. Threads implicitly synchronized within a warp. if (tid < warpSize) { for (int shift = warpSize; shift > 0; shift >>= 1) sdata[tid] = op.Reduce(sdata[tid], sdata[tid + shift]); } // Output to vector result. if (tid == 0) { result[i] = op.PostReduce(sdata[0], result[i]); } } void cudaD_add_col_sum_mat(int Gr, int Bl, double* result, const double* mat, const MatrixDim d, const double alpha, const double beta) { _transform_reduce_mat_cols<<<Gr, Bl>>>(result, mat, d, TransReduceOp<SUMAB, double>(alpha, beta)); } |
cudamatrix/cu-vector.h
与matrix/kaldi-vector.h类似的,该头文件声明了几个向量类。与之不同的是,但其运算的实现基于CUDA或CBLAS。
class CuVectorBase
Cuda向量抽象类。该类对基础运算与内存优化进行了封装,只提供向量运算。不涉及尺寸缩放和构造函数。
尺寸缩放和构造函数由派生类CuVector和CuSubVector负责。
向量初始化
void SetZero();
向量信息
MatrixIndexT Dim() const { return dim_; }
向量的读取与转换
inline Real* Data() { return data_; }
inline Real operator() (MatrixIndexT i) const
CuSubVector<Real> Range(const MatrixIndexT o, const MatrixIndexT l)
向量的拷贝函数
void CopyFromVec(const CuVectorBase<Real> &v);
向量的运算
void ApplyLog();
void AddVec(const Real alpha, const CuVectorBase<OtherReal> &v, Real beta = 1.0);
//*this += alpha * M [or M^T]
//linear_params_.AddMat(alpha, other->linear_params_);
//linear_params_ += alpha * other->linear_params_
void AddMat ( const Real alpha,
const MatrixBase< Real > & M,
MatrixTransposeType transA = kNoTrans
)
//*this = alpha * diag(M * M^T) + beta * *this
diag(M M^T)+beta ** M (1 2 3) (4 5 6) (7 8 9) (1 4 7) (2 5 8) (3 6 9) (1^2+2^2+3^2, *, *) (*, 4^2+5^2+6^2, *) (*, *, 7^2+8^2+9^2) diag=() |
void CuVectorBase<Real>::AddDiagMat2(Real alpha, const CuMatrixBase<Real> &M,
MatrixTransposeType trans, Real beta) {
//*this = alpha * diag(M * M^T) + beta * *this
this->AddDiagMatMat(alpha, M, trans, M, other_trans, beta);
}
//*this = alpha * diag(M * N^T) + beta * *this
void CuVectorBase<Real>::AddDiagMatMat(Real alpha, const CuMatrixBase<Real> &M,
MatrixTransposeType transM,
const CuMatrixBase<Real> &N,
MatrixTransposeType transN, Real beta) {
// v = alpha * diag(M * N^T) + beta * v
static void _add_diag_mat_mat_MNT(const Real alpha, const Real* M,
const MatrixDim dim_M, const Real* N,
const int stride_N, const Real beta,
Real* v)
//data_ = alpha * diag(M.Data() * N.Data()^T) + beta * data_
cuda_add_diag_mat_mat_MNT(dimGrid, dimBlock, alpha, M.Data(), M.Dim(),
N.Data(), N.Stride(), beta, data_);
class CuVector: public CuVectorBase<Real>
该类表示普通Cuda向量,并实现尺寸缩放和一般的构造函数。
多种构造函数
explicit CuVector(const CuVector<Real> &v) : CuVectorBase<Real>() {
Resize(v.Dim(), kUndefined);
this->CopyFromVec(v);
}
template<typename OtherReal>
explicit CuVector(const CuVectorBase<OtherReal> &v) : CuVectorBase<Real>() {
Resize(v.Dim(), kUndefined);
this->CopyFromVec(v);
}
template<typename OtherReal>
explicit CuVector(const VectorBase<OtherReal> &v) : CuVectorBase<Real>() {
Resize(v.Dim(), kUndefined);
this->CopyFromVec(Vector<Real>(v));
}
重载赋值运算符
CuVector<Real> &operator = (const CuVectorBase<Real> &other) {
Resize(other.Dim(), kUndefined);
this->CopyFromVec(other);
return *this;
}
CuVector<Real> &operator = (const CuVector<Real> &other) {
Resize(other.Dim(), kUndefined);
this->CopyFromVec(other);
return *this;
}
CuVector<Real> &operator = (const VectorBase<Real> &other) {
Resize(other.Dim());
this->CopyFromVec(other);
return *this;
}
Utils
void Swap(CuVector<Real> *vec);
void Swap(Vector<Real> *vec);
void Resize(MatrixIndexT length, MatrixResizeType resize_type = kSetZero);
class CuSubVector: public CuVectorBase<Real>
该类表示一个不占有实际数据的泛化向量或向量索引,可以表示高级向量的子向量或矩阵的行。实现多种用于索引的构造函数。
多种构造函数
CuSubVector(const CuVectorBase<Real> &t, const MatrixIndexT origin,
const MatrixIndexT length) : CuVectorBase<Real>() {
KALDI_ASSERT(static_cast<UnsignedMatrixIndexT>(origin)+
static_cast<UnsignedMatrixIndexT>(length) <=
static_cast<UnsignedMatrixIndexT>(t.Dim()));
CuVectorBase<Real>::data_ = const_cast<Real*>(t.Data()+origin);
CuVectorBase<Real>::dim_ = length;
}
/// Copy constructor
/// this constructor needed for Range() to work in base class.
CuSubVector(const CuSubVector &other) : CuVectorBase<Real> () {
CuVectorBase<Real>::data_ = other.data_;
CuVectorBase<Real>::dim_ = other.dim_;
}
CuSubVector(const Real* data, MatrixIndexT length) : CuVectorBase<Real> () {
// Yes, we're evading C's restrictions on const here, and yes, it can be used
// to do wrong stuff; unfortunately the workaround would be very difficult.
CuVectorBase<Real>::data_ = const_cast<Real*>(data);
CuVectorBase<Real>::dim_ = length;
}
cudamatrix/cu-matrix.h
与matrix/kaldi-matrixr.h类似的,该头文件声明了几个矩阵类。与之不同的是,但其运算的实现基于CUDA或CBLAS。当Kaldi基于CUDA环境编译且GPU可用(CuDevice::Instantiate().Enabled() == true)则使用CUDA卡进行计算,否则使用CPU进行计算(CBLAS)。
class CuMatrixBase
Cuda矩阵抽象类。该类对基础运算与内存优化进行了封装,只提供矩阵运算。不涉及尺寸缩放和构造函数。
尺寸缩放和构造函数由派生类CuMatrix和CuSubMatrix负责。
class CuMatrix
该类表示普通Cuda矩阵,并实现尺寸缩放和一般的构造函数。
class CuSubMatrix
该类表示一个不占有实际数据的泛化矩阵或矩阵索引,可以表示其他矩阵的矩阵。实现多种用于索引的构造函数。
继承于CuMatrixBase,用于对矩阵的子矩阵(块矩阵)进行运算。
kaldi通用底层矩阵运算库——CUDA的更多相关文章
-
kaldi通用底层矩阵运算库——CBLAS
matrix/cblas-wrappers.h 该头文件对CBLAS与CLAPACK的接口进行了简单的封装(将不同数据类型的多个接口封装为一个). 比如 cblas_scopy和cblas_dcopy ...
-
C++矩阵运算库推荐
最近在几个地方都看到有人问C++下用什么矩阵运算库比较好,顺便做了个调查,做一些相关的推荐吧.主要针对稠密矩阵,有时间会再写一个稀疏矩阵的推荐. Armadillo:C++下的Matlab替代品 地址 ...
-
C++通用框架和库
C++通用框架和库 来源 https://www.cnblogs.com/skyus/articles/8524408.html 关于 C++ 框架.库和资源的一些汇总列表,内容包括:标准库.Web应 ...
-
Python底层socket库
Python底层socket库将Unix关于网络通信的系统调用对象化处理,是底层函数的高级封装,socket()函数返回一个套接字,它的方法实现了各种套接字系统调用.read与write与Python ...
-
C++矩阵运算库armadillo配置笔记
前言 最近在用C++实现神经网络模型,优化算法需要用到矩阵操作,一开始我用的是boost的ublas库,但用着用着感觉很不习惯,接口不够友好.于是上网搜索矩阵运算哪家强,大神们都推荐armadillo ...
-
uTenux&mdash;&mdash;重新整理底层驱动库
重新整理底层驱动库 1. 整理chip.h 在chip.h文件中的07----13的宏定义设置位如下,这样我们就不用在工程配中定义sam3s4c这个宏了,为我们以后通用少了一件麻烦事. //#if d ...
-
.net通用底层搭建
.net通用底层搭建 之前写过几篇,有朋友说看不懂,有朋友说写的有点乱,自己看了下,的确是需要很认真的看才能看懂整套思路. 于是写下了这篇. 1.这个底层,使用的是ado.net,微软企业库 2.实体 ...
-
YARN底层基础库
YARN基础库是其他一切模块的基础,它的设计直接决定了YARN的稳定性和扩展性,YARN借用了MRV1的一些底层基础库,比如RPC库等,但因为引入了很多新的软件设计方式,所以它的基础库更多,包括直 ...
-
Duanxx的Design abroad: C++矩阵运算库Eigen 概要
一.概要 这两天想起来要做神经网络的作业了,要求用C++完毕神经网络的算法. 摆在面前的第一个问题就是,神经网络算法中大量用到了矩阵运算.可是C++不像matlab那样对矩阵运算有非常好的支持.本来准 ...
随机推荐
-
Uber优步宁波司机注册正式开始啦! UBER宁波司机注册指南!
自2012年Uber开始向全球进军以来,目前已进入全球56个国家和地区的市场,在全球超过270个城市提供服务, 而Uber公司的估值已高达412亿美元. [目前开通Uber优步叫车服务的中国城市] ...
-
mssql游标实现递归查询
表值函数 存储过程
-
C语课设心得分享(三)
调试. 以前咱们写课后习题,一般也不需要使用调试,如果程序编译error,根据错误信息就可以改好:如果是结果错误,那么在稿纸上过几遍基本也可以得出结果. 但咱们这个课设比较大,就需要很多调试的过程,尤 ...
-
嵌套结构使用:struc1-struc2-XXX
声明:原创作品,转载时请注明文章来自SAP师太技术博客( 博/客/园www.cnblogs.com):www.cnblogs.com/jiangzhengjun,并以超链接形式标明文章原始出处,否则将 ...
-
OpenGL开发时,fatal error C1083: 无法打开包括文件:“gl\glut.h”: No such file or directory
本人使用的是vs2012,编写一个简单的opengl程序,运行的时候总是提示: fatal error C1083: 无法打开包括文件:“gl/glut.h”: No such file or dir ...
-
代码分享:php判断数组是否有序
发布:脚本学堂/PHP编程 编辑:JB02 2013-12-17 14:59:02 [大 中 小] 转自:http://www.jbxue.com/article/14723.html如何判断 ...
-
windows系统下快捷命令
mstsc 远程计算机 regedit 注册表
-
Discuz经典函数注释之authcode
Discuz函数中最经典的函数是authcode函数,因为supesite,UCenterHome,UCenter,Discuz X都使用了这个函数进行加密啊传输串与cookie 今天为大家带来aut ...
-
盒子模型,定位技术,负边距,html5 新增标签
盒子模型 /*[margin 外边距] margin属性最多四个 1.只写一个值,四个方向的margin均为这个值 2.写两个值:上,右两个方向,下默认=上,右 默认=左 3.写三个值:上.右.下三个 ...
-
大数据 -->; Kafka集群搭建
Kafka集群搭建 下面是以三台机器搭建为例,(扩展到4台以上一样,修改下配置文件即可) 1.下载kafka http://apache.fayea.com/kafka/0.9.0.1/ ,拷贝到三台 ...