这一周主要用CUDA实现了BP前馈神经网络,但是一路也遇到了很多问题。
1. 批梯度下降时修改权值与偏置时候没有将累积的误差项/偏置项除以总样本数,导致每次修改值远远大于真实值,程序最后全1或全0. 我最后用matlab运行时候,结合李春光老师的神经计算课件找到了这个bug
2.CUDA运行多块多线程并行核函数时候,没有考虑到有密切前后运算逻辑关系又同时索引不严格相同的两个矩阵运算不能放在同一个核函数执行,因为每个线程执行的速度不一样。你可能在用到err值得时候,那个线程还没有把err算完,那么bias的更新就会出错。
//计算最后一层的err和bias
//O,NodeNum,SampleNum分别是最后一层的输出、节点数,样本数
__global__ void ErroBiasLastlayer(float* err, float*bias, float*Targets, float* O, const int NodeNum)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
for(int i = idx; i<NodeNum*SAMPLE_NUM; i=i+BLOCK_NUM*THREAD_NUM)
{
err[i] = (Targets[i]-O[i])*O[i]*(1-O[i]);
}
__syncthreads();
if(idx<NodeNum)
{
for(int i = 0; i<SAMPLE_NUM; ++i)
{
bias[idx] += ALFA*err[idx+i*NodeNum]/SAMPLE_NUM;
}
}
__syncthreads();
}
//计算内层的Err和Bias:
//err是这一层到下一层weight转置和下一层的err的矩阵积,是一个[NodeNum_this]*[SampleNum]的矩阵,更新这个矩阵
//O、NodeNum是这一层的输出和节点数
//err,bias是这一层需要计算的误差项、偏置项
__global__ void ErroBias(float* err, float*bias, float* O, const int NodeNum)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
for(int i = idx; i<NodeNum*SAMPLE_NUM; i=i+BLOCK_NUM*THREAD_NUM)
{
err[i] = err[i]*O[i]*(1-O[i]);
}
__syncthreads();
if(idx<NodeNum)
{
for(int i = 0; i<SAMPLE_NUM; ++i)
{
bias[idx] += ALFA*err[idx+i*NodeNum]/SAMPLE_NUM;
}
}
__syncthreads();
}
当线程idx计算 bias[idx]的时候,err的[idx+i*NodeNum]可能还没有更新完成就被使用了。解决办法就是把err和bias写成两个核函数。
3.当把数值gpuMse[BLOCK_NUM]传回CPU进行累加的时候,数据有时会传错。至今我都没有想明白这是为什么。
原文:http://www.cnblogs.com/huangshan/p/3945996.html