以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
模板->模板实例化->模板
的一个过程。
? ?
原因
__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 |
?
Kaldi使用cudaD_*和cudaF_*对kernel进行封装的意义
原文：https://www.cnblogs.com/JarvanWang/p/11754624.html