65.9K
CodeProject 正在变化。 阅读更多。
Home

GPU上的神经网络

billconan
starIconstarIconstarIconstarIcon
emptyStarIcon
starIcon

4.90/5 (69投票s)

2008年3月14日

CC (ASA 2.5)

9分钟阅读

viewsIcon

430414

downloadIcon

20125

本文介绍了使用CUDA实现神经网络。

引言

人工神经网络是一种信息处理方法，其灵感来源于生物神经系统（如大脑）处理信息的方式。它由大量高度互连的处理单元（神经元）组成，这些单元协同工作以解决特定问题。神经网络已被广泛用于“类比”信号分类，包括手写、语音和图像识别。神经网络还可以用于电子游戏，使游戏能够从玩家行为中自适应学习。这种技术已被用于赛车游戏，例如由计算机控制的对手车辆可以学习如何像人类玩家一样驾驶。

由于神经网络需要大量的向量和矩阵运算才能获得结果，因此非常适合在并行编程模型中实现并在图形处理单元（GPU）上运行。我们的目标是利用并释放GPU的强大功能，以提高神经网络解决手写识别问题的性能。

该项目最初是我们的图形架构课程项目。我们在GPU上运行了Mike O'Neill在其精彩文章“用于识别手写数字的神经网络”中所描述的同一个神经网络。

关于神经网络

神经网络由两种基本元素组成：神经元和连接。神经元通过连接相互连接形成网络。这是一个对人脑的简化理论模型。

神经网络通常有多个层；某个层的神经元以某种方式连接到下一层的神经元。它们之间的每个连接都被分配一个权重值。起初，输入数据被馈送到第一层的神经元，通过计算所有连接的第一层神经元的加权和，我们可以得到第二层神经元的神经元值，依此类推。最后，我们可以到达最后一层，即输出层。操作神经网络所涉及的所有计算都是一系列点积。

神经网络的秘密在于权重值。正确的权重值使其完美。然而，一开始我们不知道这些值。因此，我们需要用样本输入来训练我们的网络，并将结果与我们期望的答案进行比较。某些算法可以以错误作为输入并修改网络权重。如果足够耐心，神经网络可以被训练到达到高精度。

IllustrationNeuralNet.gif

我们实现的神经网络是一个称为卷积神经网络的5层网络。这种网络已被证明适合识别手写数字。有关更多理论细节，请查阅Mike的文章和他列出的参考文献。

我们的神经网络的前三层由多个特征图组成。其中每个特征图都是从前一层收缩而来的。我们的输入是一个29*29的数字图像。因此，第一层有29*29=841个神经元。第二层是一个具有6个特征图的卷积层。每个特征图是一个13*13的图像，是从第一层采样的。特征图中的每个像素/神经元是输入层的5*5卷积核。因此，这一层有13*13*6 = 1014个节点/神经元，（5*5+1)*6 = 156个权重，1014*(5*5+1) = 26364个连接到第一层。

第三层也是一个卷积层，但具有50个较小的特征图。每个特征图的大小为5*5，这些特征图中的每个像素都是前一层所有6个特征图的相应区域的5*5卷积核。因此，这一层有5*5*50 = 1250个神经元，(5*5+1)*6*50 = 7800个权重，以及1250*26 = 32500个连接。

第四层是一个具有100个神经元的全连接层。由于它是全连接的，这一层中的每个100个神经元都连接到前一层的所有1250个神经元。因此，它有100个神经元，100*(1250+1) = 125100个权重，以及100x1251 = 125100个连接。

第五层是最终的输出层。这一层也是一个具有10个单元的全连接层。这一层中的每个10个神经元都连接到前一层的所有100个神经元。第五层有10个神经元，10*(100+1) = 1010个权重，以及10x101 = 1010个连接。

正如你所见，尽管结构简单，这个神经网络是一个庞大的数据结构。

之前的GPU实现

Fast Neural Network Library (FANN) 在GPU上有一个非常简单的神经网络实现，使用GLSL。每个神经元被表示为纹理像素的单个颜色通道。这个网络非常具体；神经元的取值范围为0到1，精度只有8位。这个实现利用硬件加速的点积函数来计算神经元。神经元和权重都存储在纹理图中。

这个实现是直接且简单的，但也是有限的。首先，在我们的神经网络中，每个神经元需要32位浮点精度。由于我们的网络有五层，第一层损失的精度会累积并改变最终结果。而且，由于手写识别系统需要足够敏感以检测不同输入之间的细微差别，因此使用仅8位来表示神经元是不可接受的。其次，普通的神经网络将神经元值映射到0到1的范围。然而，在我们的程序中，专门为手写识别设计的神经网络有一个特殊的激活函数，将每个神经元的值映射到-1到1的范围。因此，如果神经元像FANN库那样由单个颜色值表示，我们的神经元将进一步损失精度。最后，FANN方法使用点积来计算神经元，这适用于全连接的神经网络。在我们的实现中，神经网络是部分连接的。在我们神经网络上执行的计算涉及大向量的点积。

我们的实现

由于上述GLSL的所有不便之处，我们最终选择了CUDA。神经网络之所以适合GPU，是因为神经网络的训练和执行是两个独立的过程。一旦正确训练，在使用神经网络时就不需要写访问。因此，没有需要解决的同步问题。此外，同一网络层上的神经元是完全隔离的，因此神经元值计算可以实现高度并行化。

在我们的代码中，第一层的权重存储为数组，这些输入被复制到设备。对于每个网络层，都有一个CUDA函数处理该层的神经元值计算，因为并行化只能在一个层内实现，并且层与层之间的连接是不同的。神经网络的连接在CUDA函数中通过下一层神经元计算的方程隐式定义。我们的代码中没有显式的连接数据结构。这是我们的代码与Mike的CPU版本之间的一个主要区别。

cuda.PNG

例如，第二层神经元的每个神经元值是25个第一层神经元和一个偏置的加权和。第二层神经元由6个特征图组成；每个特征图的大小为13*13。我们为每个特征图分配一个blockID，为特征图上的每个神经元分配一个threadID。每个特征图由一个块处理，图上的每个像素由一个线程处理。

这是计算第二层网络的CUDA函数

__global__ void executeFirstLayer
    (float *Layer1_Neurons_GPU,float *Layer1_Weights_GPU,float *Layer2_Neurons_GPU)
{
    int blockID=blockIdx.x;
    int pixelX=threadIdx.x;
    int pixelY=threadIdx.y;

    int kernelTemplate[25] = {
        0,  1,  2,  3,  4,
        29, 30, 31, 32, 33,
        58, 59, 60, 61, 62,
        87, 88, 89, 90, 91,
        116,117,118,119,120 };

    int weightBegin=blockID*26;
    int windowX=pixelX*2;
    int windowY=pixelY*2;

    float result=0;

    result+=Layer1_Weights_GPU[weightBegin];

    ++weightBegin;

    for(int i=0;i<25;++i)
    {
        result+=Layer1_Neurons_GPU
            [windowY*29+windowX+kernelTemplate[i]]*Layer1_Weights_GPU[weightBegin+i];
    }

    result=(1.7159*tanhf(0.66666667*result));

    Layer2_Neurons_GPU[13*13*blockID+pixelY*13+pixelX]=result;
}

所有其他层都以相同的方式计算；唯一的区别在于计算神经元的方程。

program.PNG

主程序首先将所有输入数据传输到GPU，然后按顺序调用每个CUDA函数，最后得到答案。

recod.jpg

用户界面是使用C#的一个独立程序。用户可以用鼠标在输入板上绘制一个数字，然后程序会生成一个29*29的图像并调用内核神经网络程序。内核将读取输入图像并将其馈送到我们的神经网络。结果也通过文件返回，然后由用户界面读回。

这是截图。绘制一个数字后，我们可以得到最后一个网络层的所有10个神经元值。最大神经元值的索引是最可能的数字。我们根据其可能性以不同深度的红色对候选者进行着色。

右侧，用户界面将打印出前三层的特征图。

请注意，Windows XP下的C#存在分辨率问题。我们在120dpi下测试了我们的程序。96dpi的分辨率设置可能会导致输入图像移位，从而严重影响精度。

我们的GPU实现中不包含训练部分。我们使用Mike的代码来训练所有权重，并将它们缓存到文件中。

结果

准确性

我们的神经网络可以达到95%的准确率。我们用来训练网络的数据库称为MNIST，包含来自不同人的60000个手写样本。LeCun博士报告说，这个网络可以在大约25次训练后收敛。这个数字也得到了我们的测试证实。在60000个输入中，我们只实现了大约1400个识别错误样本。

另外请注意，Mike的代码中存在一个bug。这是初始化第二层的更正后的代码。

for ( fm=0; fm<50; ++fm)
{
  for ( ii=0; ii<5; ++ii )
  {
    for ( jj=0; jj<5; ++jj )
    {
      // iNumWeight = fm * 26;  // 26 is the number of weights per feature map
      iNumWeight = fm * 156;  // 156 is the number of weights per feature map
      NNNeuron& n = *( pLayer->m_Neurons[ jj + ii*5 + fm*25 ] );

      n.AddConnection( ULONG_MAX, iNumWeight++ );  // bias weight

      for ( kk=0; kk<25; ++kk )
      {
        // note: max val of index == 1013, corresponding to 1014 neurons in prev layer
        n.AddConnection(       2*jj + 26*ii + kernelTemplate2[kk], iNumWeight++ );
        n.AddConnection( 169 + 2*jj + 26*ii + kernelTemplate2[kk], iNumWeight++ );
        n.AddConnection( 338 + 2*jj + 26*ii + kernelTemplate2[kk], iNumWeight++ );
        n.AddConnection( 507 + 2*jj + 26*ii + kernelTemplate2[kk], iNumWeight++ );
        n.AddConnection( 676 + 2*jj + 26*ii + kernelTemplate2[kk], iNumWeight++ );
        n.AddConnection( 845 + 2*jj + 26*ii + kernelTemplate2[kk], iNumWeight++ );
      }
    }
  }
}

有关此bug的详细信息，请参阅此处。

我们的GPU实现基于正确版本，但精度方面没有太大差异。

性能

performance.PNG

使用GPU计算神经网络的主要原因是实现鲁棒性。与CPU实现相比，结果是令人鼓舞的。如上表所示，比较了GPU版本、EmuRelease版本和CPU版本在处理单个输入样本时的执行时间。GPU版本比CPU版本快270倍，比EmuRelease版本快516.6倍。为了更精确，我们还考虑了GPU版本的IO时间消耗。正如我们所见，即使考虑了IO时间，我们的方法也快了10倍。在实际使用中，权重值只需要加载到设备一次。

历史

  • 2008年3月14日：首次发布
GPU上的神经网络 - CodeProject - 代码之家
© . All rights reserved.