kernel函数

转自：https://m.2cto.com/kf/201701/583626.html

核函数

核函数的特点

cuda代码文件的后缀为”.cu”，由单独的编译器进行编译 核函数是cu文件中的一部分代码，是运行在显存中的程序代码，是实现并行计算的载体 核函数一般放在cu文件中的前面，函数定义之前需要添加__global__关键字，函数体中包含CUDA_KERNEL_LOOP循环体 CUDA_KERNEL_LOOP循环体有两个参数，第一个是迭代器，第二个是总迭代数 CUDA_KERNEL_LOOP循环体中的代码是并行执行的，是互不关联的可独立执行的程序

示例代码

template 
__global__ void kernel_statistic(const int num, const Dtype* bottom_data, Dtype* temp, 
const int label_num, const int nsim, Dtype* counter) {
CUDA_KERNEL_LOOP(index, num) {
Dtype count_iter(0.0);
for (int j = index + 1; j < num; ++j) {
Dtype result_dot(0.0);
for (int k = 0; k < label_num; ++k) {
result_dot += bottom_data[index * label_num + k] * bottom_data[j * label_num + k];
}
temp[index * num * 2 + j * 2] = result_dot;
if (result_dot >= Dtype(1.0))
count_iter++;
}
counter[index] = count_iter;
}
}

注意事项

核函数中不能出现__host__类型的函数，例如caffe中定义的caffe_gpu开头的函数、C++ 标准库中的函数 核函数中的数学计算由CUDA Math API完成 核函数一般不需要返回值 核函数的参数是所有CUDA_KERNEL_LOOP循环体公用的，对数据的修改应该是互不干扰的，示例代码中counter数组存储了各循环体代码计数的结果，传出后再进行累加运算得到总的统计结果。

传送门

CUDA Math API

Forward_gpu和Backward_gpu

注意事项

这两个函数需要在层的hpp文件中声明 cu文件编译生成后，cpp文件中的Forward_cpu函数和Backward_cpu函数将不会被调用 初始化层时，cpp文件中的LayerSetUp函数和Reshape函数也会被执行 对数组求和，可以用caffe_gpu_asum函数 数据在GPU和CPU之间的拷贝速度特别慢，在cu文件中慎用cpu_data函数和mutabel_cpu_data函数 GPU擅长处理大规模矩阵运算，核函数应简单简洁