Kaldi使用cudaD_*和cudaF_*对kernel进行封装的意义

以AddMat函数为例

Kaldi的CUDA Kernel函数的声明位于cudamatrix/cu-kernels-ansi.h；定义位于cudamatrix/cu-kernels.cu中：

src/cudamatrix/cu-kernels-ansi.h

extern "C" {

//...

void cudaD_add_mat(dim3 Gr, dim3 Bl, double alpha, const double *src,

double *dst, MatrixDim d, int src_stride, int A_trans);

void cudaF_add_mat(dim3 Gr, dim3 Bl, float alpha, const float *src, float *dst,

MatrixDim d, int src_stride, int A_trans);

//...

}

src/cudamatrix/cu-kernels.cu

void cudaD_add_mat(dim3 Gr, dim3 Bl, double alpha, const double* src,

double* dst, MatrixDim d, int src_stride, int A_trans) {

if (A_trans) {

_add_mat_trans<<<Gr,Bl>>>(alpha,src,dst,d,src_stride);

} else {

_add_mat<<<Gr,Bl>>>(alpha,src,dst,d,src_stride);

}

}

void cudaF_add_mat(dim3 Gr, dim3 Bl, float alpha, const float* src, float* dst,

MatrixDim d, int src_stride, int A_trans) {

if (A_trans) {

_add_mat_trans<<<Gr,Bl>>>(alpha,src,dst,d,src_stride);

} else {

_add_mat<<<Gr,Bl>>>(alpha,src,dst,d,src_stride);

}

}

template<typename Real>

__global__

static void _add_mat(Real alpha, const Real* src, Real* dst, MatrixDim d, int src_stride) {}

   

其中_add_mat是一个模板函数，cudaD_add_mat和cudaF_add_mat可视为_add_mat对不同浮点类型的模板实例化。

   

类似的，根据数据类型的不同，cudaD_add_mat和cudaF_add_mat被cudamatrix/cu-kernels.h的以下函数调用：

inline void cuda_add_mat(dim3 Gr, dim3 Bl, double alpha, const double *src,

double *dst, MatrixDim d, int src_stride,

int A_trans) {

cudaD_add_mat(Gr, Bl, alpha, src, dst, d, src_stride, A_trans);

}

inline void cuda_add_mat(dim3 Gr, dim3 Bl, float alpha, const float *src,

float *dst, MatrixDim d, int src_stride, int A_trans) {

cudaF_add_mat(Gr, Bl, alpha, src, dst, d, src_stride, A_trans);

}

   

而重载的cuda_add_mat函数被模板类的函数CuMatrixBase<Real>::AddMat调用：

src/cudamatrix/cu-matrix.cc

template<typename Real>

void CuMatrixBase<Real>::AddMat(Real alpha, const CuMatrixBase<Real>& A,MatrixTransposeType transA) {

...

cuda_add_mat(dimGrid, dimBlock, alpha, A.data_,

data_, Dim(), A.Stride(),

(transA == kTrans ? 1 : 0));

...

}

   

总结

这样，相当于经过了：

cu-matrix.cc -> cu-kernel.h -> cu-kernel.cu

模板->模板实例化->模板

的一个过程。

   

原因

CUDA中有根据函数名获取函数地址的API：

__global__ void exampleKernel(float** data) { ... }

//调用成功

cuModuleLoad(&cuModule, modulePath);

//调用失败

cuModuleGetFunction(&cuFunction, cuModule, "exampleKernel");

其中exampleKernel函数将被视为C++函数：

$ cat <<EOF > a.cc

void exampleKernel(float** data){}

EOF

$ gcc -c a.cc -o a.o

$ nm a.o

0000000000000000 T _Z13exampleKernelPPf

编译后将得到被"修饰"的函数名。

   

为了使得上述API可用，一般将CUDA函数声明为C函数：

extern "C"

__global__ void exampleKernel(float** data) { ... }

//调用成功

cuModuleLoad(&cuModule, modulePath);

//调用成功

cuModuleGetFunction(&cuFunction, cuModule, "exampleKernel");

   

$ cat <<EOF > a.cc

extern "C"

void exampleKernel(float** data){}

EOF

$ gcc -c a.cc -o a.o

$ nm a.o

0000000000000000 T exampleKernel