本神经网络为3层结构,一个输入层(N0N个神经元)+一个隐层(N1N个神经元)+一个输出层(N2N个神经元)
cuda加速为神经网络的3大主要计算部分 神经元输入计算,神经元错误计算,神经权重调整
#include<iostream>
#include<cmath>
#include<stdio.h>
#include<stdlib.h>
using namespace std;
#define N0N 3
#define N1N 2
#define N2N 1
#define HANDLE_ERROR(x) {if(x!=0) printf("error is caught on gpu: %d\n",x);}
//训练用的参数
struct neuroNode
{
float I;
float O;
float Err;
float bias;
};
float pic[N0N]={0};
neuroNode node1[N1N];
neuroNode node2[N2N];
float w01[N0N][N1N]={0};
float w12[N1N][N2N]={0};
float l=0.9;
//cuda第一加速部分
__global__ void cuda_fun_mul(float *a,float *b,float *c)
{
int tid=threadIdx.x;
c[tid]=a[tid]*b[tid];
}
//cuda第二加速部分
__global__ void cuda_fun_mul2(float *a,float *b,float *cc,int *NxN,float *l)
{
int bid=blockIdx.x;
int tid=threadIdx.x;
cc[bid*(*NxN)+tid]=(*l)*a[bid]*b[tid];
}
//保存网络
void saveNet()
{
int i,j,k;
freopen("net.txt","w",stdout);
//保存权重
printf("biasnode1\n");
for(i=0;i<N1N;i++)
printf("%.2f ",node1[i].bias);
printf("\n");
printf("biasnode2\n");
for(i=0;i<N2N;i++)
printf("%.2f ",node2[i].bias);
printf("\n");
//保存偏量
printf("w01[N0N][N1N]\n");
for(i=0;i<N0N;i++)
{
for(j=0;j<N1N;j++)
{
printf("%.2f ",w01[i][j]);
}
printf("\n");
}
printf("w12[N1N][N2N]\n");
for(i=0;i<N1N;i++)
{
for(j=0;j<N2N;j++)
{
printf("%.2f ",w12[i][j]);
}
printf("\n");
}
fclose(stdout);
}
int main()
{
int i,j,k;
pic[0]=1;
pic[1]=0;
pic[2]=1;
w01[0][0]=0.2;
w01[0][1]=-0.3;
w01[1][0]=0.4;
w01[1][1]=0.1;
w01[2][0]=-0.5;
w01[2][1]=0.2;
w12[0][0]=-0.3;
w12[1][0]=-0.2;
node1[0].bias=-0.4;
node1[1].bias=0.2;
node2[0].bias=0.1;
int Nmax=N0N;
if(Nmax<N1N)
Nmax=N1N;
if(Nmax<N2N)
Nmax=N2N;
float a[Nmax];
float b[Nmax];
float c[Nmax];
float cc[Nmax*Nmax];
int NxN;
float *dev_a;
float *dev_b;
float *dev_c;
float *dev_cc;
HANDLE_ERROR(cudaMalloc((void**)&dev_a,sizeof(float)*Nmax));
HANDLE_ERROR(cudaMalloc((void**)&dev_b,sizeof(float)*Nmax));
HANDLE_ERROR(cudaMalloc((void**)&dev_c,sizeof(float)*Nmax));
HANDLE_ERROR(cudaMalloc((void**)&dev_cc,sizeof(float)*Nmax*Nmax));
float *dev_l;
int *dev_NxN;
HANDLE_ERROR(cudaMalloc((void**)&dev_l,sizeof(float)));
HANDLE_ERROR(cudaMalloc((void**)&dev_NxN,sizeof(int)));
HANDLE_ERROR(cudaMemcpy(dev_l,&l,sizeof(float),cudaMemcpyHostToDevice));
//训练神经网络
//正向过程
//1st层神经网络训练
for(j=0;j<N1N;j++)//1st层神经网络训练,every neuro unit
{
node1[j].I=node1[j].bias;
for(i=0;i<N0N;i++)//every colum of a pic,1层第j神经元 与 0层第i输入元的关系
{
a[i]=w01[i][j];
b[i]=pic[i];
}
HANDLE_ERROR(cudaMemcpy(dev_a,a,sizeof(float)*N0N,cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_b,b,sizeof(float)*N0N,cudaMemcpyHostToDevice));
cuda_fun_mul<<<1,N0N>>>(dev_a,dev_b,dev_c);
HANDLE_ERROR(cudaMemcpy(c,dev_c,sizeof(float)*N0N,cudaMemcpyDeviceToHost));
for(i=0;i<N0N;i++)//every colum of a pic,1层第j神经元 与 0层第i输入元的关系
{
node1[j].I+=c[i];//w01[j][i] means 0层第j输入元的关系 与 1层第i神经元
}
node1[j].O=1.0/(1.0+exp(-node1[j].I));
}
//1->2层训练
for(j=0;j<N2N;j++)//2nd层神经网络训练,every neuro unit
{
node2[j].I=node2[j].bias;
for(i=0;i<N1N;i++)//1层第i神经元与 2层第j输入元的关系
{
a[i]=w12[i][j];//w12[i][j]表示1st层第i输入单元到2nd层第j神经单元的权重
b[i]=node1[i].O;
}
HANDLE_ERROR(cudaMemcpy(dev_a,a,sizeof(float)*N1N,cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_b,b,sizeof(float)*N1N,cudaMemcpyHostToDevice));
cuda_fun_mul<<<1,N1N>>>(dev_a,dev_b,dev_c);
HANDLE_ERROR(cudaMemcpy(c,dev_c,sizeof(float)*N1N,cudaMemcpyDeviceToHost));
for(i=0;i<N1N;i++)
{
node2[j].I+=c[i];
}
node2[j].O=1.0/(1.0+exp(-node2[j].I));
}
//反向过程:
//cul err
int targetNum=0;
for(i=0;i<N2N;i++)//2nd层神经元错误计算
{
if(i==targetNum)
node2[i].Err=node2[i].O*(1-node2[i].O)*(1-node2[i].O);
else
node2[i].Err=node2[i].O*(1-node2[i].O)*(0-node2[i].O);
node2[i].bias+=l*node2[i].Err;
}
for(i=0;i<N1N;i++)//1st层神经元错误计算,对于1st层第i个神经元作计算
{
node1[i].Err=0;
for(j=0;j<N2N;j++)
{
a[j]=w12[i][j];
b[j]=node2[j].Err;
}
HANDLE_ERROR(cudaMemcpy(dev_a,a,sizeof(float)*N2N,cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_b,b,sizeof(float)*N2N,cudaMemcpyHostToDevice));
cuda_fun_mul<<<1,N2N>>>(dev_a,dev_b,dev_c);
HANDLE_ERROR(cudaMemcpy(c,dev_c,sizeof(float)*N2N,cudaMemcpyDeviceToHost));
float temp=node1[i].O*(1-node1[i].O);
for(j=0;j<N2N;j++)
{
node1[i].Err+=temp*c[j];
}
node1[i].bias+=l*node1[i].Err;
}
//change weight
//w12
for(i=0;i<N1N;i++)
a[i]=node1[i].O;
for(j=0;j<N2N;j++)
b[j]=node2[j].Err;
NxN=N2N;
HANDLE_ERROR(cudaMemcpy(dev_a,a,sizeof(float)*N1N,cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_b,b,sizeof(float)*N2N,cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_NxN,&NxN,sizeof(int),cudaMemcpyHostToDevice));
cuda_fun_mul2<<<N1N,N2N>>>(dev_a,dev_b,dev_cc,dev_NxN,dev_l);
HANDLE_ERROR(cudaMemcpy(cc,dev_cc,sizeof(float)*N2N*N1N,cudaMemcpyDeviceToHost));
for(i=0;i<N1N;i++)
for(j=0;j<N2N;j++)
w12[i][j]+=*(cc+i*N2N+j);
//w01
for(i=0;i<N0N;i++)
a[i]=pic[i];
for(j=0;j<N1N;j++)
b[j]=node1[j].Err;
NxN=N1N;
HANDLE_ERROR(cudaMemcpy(dev_a,a,sizeof(float)*N0N,cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_b,b,sizeof(float)*N1N,cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_NxN,&NxN,sizeof(int),cudaMemcpyHostToDevice));
cuda_fun_mul2<<<N0N,N1N>>>(dev_a,dev_b,dev_cc,dev_NxN,dev_l);
HANDLE_ERROR(cudaMemcpy(cc,dev_cc,sizeof(float)*N0N*N1N,cudaMemcpyDeviceToHost));
for(i=0;i<N0N;i++)
for(j=0;j<N1N;j++)
w01[i][j]+=*(cc+i*N1N+j);
saveNet();
}