kaldi通用底层矩阵运算库——CUDA

时间:2022-12-18 16:59:03

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
namespace kaldi {

 
 

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

)

kaldi通用底层矩阵运算库——CUDA

 
 

//*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的更多相关文章

  1. kaldi通用底层矩阵运算库——CBLAS

    matrix/cblas-wrappers.h 该头文件对CBLAS与CLAPACK的接口进行了简单的封装(将不同数据类型的多个接口封装为一个). 比如 cblas_scopy和cblas_dcopy ...

  2. C&plus;&plus;矩阵运算库推荐

    最近在几个地方都看到有人问C++下用什么矩阵运算库比较好,顺便做了个调查,做一些相关的推荐吧.主要针对稠密矩阵,有时间会再写一个稀疏矩阵的推荐. Armadillo:C++下的Matlab替代品 地址 ...

  3. C&plus;&plus;通用框架和库

    C++通用框架和库 来源 https://www.cnblogs.com/skyus/articles/8524408.html 关于 C++ 框架.库和资源的一些汇总列表,内容包括:标准库.Web应 ...

  4. Python底层socket库

    Python底层socket库将Unix关于网络通信的系统调用对象化处理,是底层函数的高级封装,socket()函数返回一个套接字,它的方法实现了各种套接字系统调用.read与write与Python ...

  5. C&plus;&plus;矩阵运算库armadillo配置笔记

    前言 最近在用C++实现神经网络模型,优化算法需要用到矩阵操作,一开始我用的是boost的ublas库,但用着用着感觉很不习惯,接口不够友好.于是上网搜索矩阵运算哪家强,大神们都推荐armadillo ...

  6. uTenux&mdash&semi;&mdash&semi;重新整理底层驱动库

    重新整理底层驱动库 1. 整理chip.h 在chip.h文件中的07----13的宏定义设置位如下,这样我们就不用在工程配中定义sam3s4c这个宏了,为我们以后通用少了一件麻烦事. //#if d ...

  7. &period;net通用底层搭建

    .net通用底层搭建 之前写过几篇,有朋友说看不懂,有朋友说写的有点乱,自己看了下,的确是需要很认真的看才能看懂整套思路. 于是写下了这篇. 1.这个底层,使用的是ado.net,微软企业库 2.实体 ...

  8. YARN底层基础库

      YARN基础库是其他一切模块的基础,它的设计直接决定了YARN的稳定性和扩展性,YARN借用了MRV1的一些底层基础库,比如RPC库等,但因为引入了很多新的软件设计方式,所以它的基础库更多,包括直 ...

  9. Duanxx的Design abroad: C&plus;&plus;矩阵运算库Eigen 概要

    一.概要 这两天想起来要做神经网络的作业了,要求用C++完毕神经网络的算法. 摆在面前的第一个问题就是,神经网络算法中大量用到了矩阵运算.可是C++不像matlab那样对矩阵运算有非常好的支持.本来准 ...

随机推荐

  1. Uber优步宁波司机注册正式开始啦! UBER宁波司机注册指南!

      自2012年Uber开始向全球进军以来,目前已进入全球56个国家和地区的市场,在全球超过270个城市提供服务, 而Uber公司的估值已高达412亿美元. [目前开通Uber优步叫车服务的中国城市] ...

  2. mssql游标实现递归查询

    表值函数 存储过程

  3. C语课设心得分享&lpar;三&rpar;

    调试. 以前咱们写课后习题,一般也不需要使用调试,如果程序编译error,根据错误信息就可以改好:如果是结果错误,那么在稿纸上过几遍基本也可以得出结果. 但咱们这个课设比较大,就需要很多调试的过程,尤 ...

  4. 嵌套结构使用:struc1-struc2-XXX

    声明:原创作品,转载时请注明文章来自SAP师太技术博客( 博/客/园www.cnblogs.com):www.cnblogs.com/jiangzhengjun,并以超链接形式标明文章原始出处,否则将 ...

  5. OpenGL开发时,fatal error C1083&colon; 无法打开包括文件&colon;&OpenCurlyDoubleQuote;gl&bsol;glut&period;h”&colon; No such file or directory

    本人使用的是vs2012,编写一个简单的opengl程序,运行的时候总是提示: fatal error C1083: 无法打开包括文件:“gl/glut.h”: No such file or dir ...

  6. 代码分享:php判断数组是否有序

    发布:脚本学堂/PHP编程  编辑:JB02   2013-12-17 14:59:02  [大 中 小] 转自:http://www.jbxue.com/article/14723.html如何判断 ...

  7. windows系统下快捷命令

    mstsc 远程计算机 regedit 注册表

  8. Discuz经典函数注释之authcode

    Discuz函数中最经典的函数是authcode函数,因为supesite,UCenterHome,UCenter,Discuz X都使用了这个函数进行加密啊传输串与cookie 今天为大家带来aut ...

  9. 盒子模型,定位技术,负边距,html5 新增标签

    盒子模型 /*[margin 外边距] margin属性最多四个 1.只写一个值,四个方向的margin均为这个值 2.写两个值:上,右两个方向,下默认=上,右 默认=左 3.写三个值:上.右.下三个 ...

  10. 大数据 --&gt&semi; Kafka集群搭建

    Kafka集群搭建 下面是以三台机器搭建为例,(扩展到4台以上一样,修改下配置文件即可) 1.下载kafka http://apache.fayea.com/kafka/0.9.0.1/ ,拷贝到三台 ...