基于CUDA的BP算法并行化与实例验证

孙香玉，冯百明，杨鹏斐

1.西北师范大学 计算机科学与工程学院，兰州 730070 2.中国科学院 计算技术研究所 计算机体系结构国家重点实验室，北京 100190

基于CUDA的BP算法并行化与实例验证

孙香玉，冯百明，杨鹏斐

1.西北师范大学 计算机科学与工程学院，兰州 730070 2.中国科学院 计算技术研究所 计算机体系结构国家重点实验室，北京 100190

1 引言

BP神经网络是人工神经网络的一种，它被广泛应用于各个领域，但是BP神经网络的训练过程需要进行大量费时的矩阵运算，所以经常并行化BP算法以加快计算。

传统的BP算法并行方法[1]需要使用大型并行机或网络并且并行程序的编写较为复杂，所以没有在普通人群中推广开来。文献[2-3]把神经网络和BP神经网络的训练过程转化为GPU纹理的渲染过程加快了神经网络训练过程，但是需要使用图形学API编程，要求编程人员对图形学硬件和编程接口有深入了解，开发难度大。文献[4-5]在CUDA下并行化BP算法并进行了语音识别，加速比是在CPU上的6倍。文献[6-7]在CUDA下并行化BP算法并分别用于文本检测和图像压缩，加速比分别是在CPU上的8倍和9倍。CUDA不需要借助图形学API，可直接用大多数人熟悉的C或C++语言编写在GPU上执行的程序，用起来较为简便，基于CUDA模型实现的BP神经网络并行算法已成功用于语音识别，文本检测，图像压缩，图像分割等方面。文献[8]用CUDA在GPU上加速了手写数字识别，手写数字的训练过程是在CPU上用神经网络训练实现的。针对当训练样本集过大时神经网络成批训练受限的问题提出了一种采用单样本训练的方式与CUDA模型结合实现BP神经网络算法并行化的方法，有效节省了GPU存储空间，并将该方法用于手写数字训练，验证了它的有效性。

2 CUDA架构

CUDA模型采用单指令流，多数据流（Single Instruction Multiple Data，SIMD）执行模式，将CPU作为主机（Host），GPU作为设备（Device），CPU与GPU协同工作，CPU、GPU各自拥有相互独立的存储地址空间。

运行在GPU上的CUDA并行计算函数称为内核函数（kernel），它不是一个完整的程序，而是整个CUDA程序中的一个可以并行执行的步骤。一个完整的CUDA程序是由一系列设备端kernel函数并行步骤和主机端的串行步骤共同组成的。这些步骤会按照程序中相应语句的顺序执行，满足顺序一致性。一个kernel函数中存在两个层次的并行[9]，即线程格（grid）中的线程块（block）间的并行和block中的线程（thread）间的并行。两层并行模型是CUDA最重要的创新之一。CUDA编程模型[10]如图1。

图1CUDA编程模型

3 串行BP算法并行化

3.1 BP算法

BP神经网络是一种按误差逆传播算法训练的多层前馈网络，它由一个输入层，一个或多个隐含层和一个输出层组成，含有一个隐含层的BP神经网络结构如图2所示。各层神经元仅与相邻层神经元之间相互全连接，同层内神经元间无连接，各层神经元间无反馈连接。每个输出单元取前一层所有单元输出的加权和作为输入，将一个非线性（激励）函数作用于加权输入得到输出，如此下去，直至得到输出层的输出。

图2 多层前馈（BP）神经网络

采用单样本训练方式的BP算法[11]描述如下：

（1）初始化网络的权重数组W和偏倚数组b。

（2）对于隐含层或输出层的每个单元j：

关于前一层i，计算单元j的输入：

//对于两层神经网络，前一层即为输入层计算单元j的输出：

（5）达到预先指定的训练次数或误差精度要求则输出权值和偏倚结果，否则继续（2）至（4）。

3.2 串行实现

根据3.1节对BP算法的描述得其串行伪码如下：

对N个样本依次执行第6～15行的训练，直至达到预定训练次数Pre_times。IN，HN，ON分别为输入层，隐含层，输出层神经元数。

6～7行表示从N个样本的输入层输入数组Study_DataI[N*IN]中得到当前训练样本m的输入数组I[IN]，8～9行表示从N个样本的输出层已知目标值Study_DataT[N*ON]中得到当前训练样本m的已知目标值数组T[ON]。

根据3.1节中的公式（1）（2），通过Compute_H_out（）可得隐含层输出H_out[HN]。同计算H_out的过程类似，通过Compute_O_out（）得到输出层输出O_out[ON]。

根据公式（3）和Compute_O_err（）可得输出层误差O_err[ON]。根据公式（4）和Compute_H_err（）可得隐含层误差H_err[HN]。

根据公式（5）（6）和第14行操作可得更新后的权重O_H_W和偏倚O_bias[ON]。同理可得更新后的权重H_I_W和偏倚H_bias[HN]。

采用串行方法训练BP神经网络，当训练样本数达到几万时，为了取得较好的训练结果用于识别需要耗费至少几个小时，所以考虑在GPU上利用CUDA模型加速训练过程。

3.3 并行实现

BP算法在CUDA模型下的并行伪码如下。

在GPU上训练开始前，通过cudaMemcpy（）把CPU端的输入、已知目标输出、所有权重和偏倚复制到GPU端，训练过程完全在GPU端并行实现，训练结束后再将训练结果（权重和偏倚）复制回CPU端，训练过程中不存在CPU端与GPU端的数据传输，减少了通信时间开销。

并行伪码中的9，10，11，12，13分别对串行伪码中的10，11～12，13，14，15并行化。下面详细介绍并行化实现过程。

BP并行伪码中的9计算隐含层输出。由于隐含层各神经元的输出只与输入层所有神经元的输出相关，与隐含层其他神经元的输出不相关，所以可并行计算它们。〈〈〈>>>中HN表示核函数共启动HN个块，threadsPerBlock= min{IN，blockDim.x}为每个块内的线程数，blockDim.x为块内最大线程数，取IN和blockDim.x二者中较小值作为threadsPerBlock，第三个参数IN表示shared memory中外部数组d_I1的大小，为方便说明计算过程，假设blockDim.x= 512，IN=2*blockDim.x，并行化的具体实现如下。

//启动HN个块，每个块内threadsPerBlock个线程并行计算点积和，将每个线程上计算的部分和存到它所在块的shared memory中的数组cache中

//每个块内进行并行规约（reduction）加法运算得到该块内所有线程点积运算的总和，存到cache[0]中

//HN个块并行计算隐含层输入

//HN个块并行计算隐含层输出

在上述伪码6～11的规约运算中用到了以下优化方法，一是GPU整数处理功能较弱，用位运算i>>=1代替i/=2，节省计算时间；二是相邻thread操作数组内的相邻元素，避免了bank conflict，使访问shared memory的速度与register的相同，都为global memory的100多倍。

BP并行伪码中的第10行除了计算输出层的输出d_O_out（与上面计算d_H_out的过程类似，不再详述）外，还将ON个块并行计算出了输出层的误差d_O_out（参照公式（3））。

BP并行伪码中的第11行，核函数启动HN个块，每个块内ON个线程，并行计算隐含层的误差，详细实现如下。

BP并行伪码中的第12行，核函数启动ON个块，每个块中HN个线程并行计算输出层权重d_O_H_W和偏倚d_O_bias的更新，具体实现如下（并行伪码中的第13行更新隐含层权重和偏倚的过程与之类似，不再说明）。

4 实验及分析

用第3章中提出的BP算法并行化方法进行脱机手写数字训练，用实例验证其有效性。

4.1 手写数字图像集来源

实验用到的手写样本图像集来自于MNIST（Modified National Institute of Standards and Technology）手写数字数据库[12]，所要训练和识别的图片全部为.bmp格式的黑白图片，像素为28×28，为了使其更适合在CUDA模型上并行计算，读取文件时即根据位图信息把每幅图片的数字表示成32×32的点阵数列，原图像中的黑白像素点在点阵数列中分别用0.0，1.0表示，即进行二值化处理。

4.2 BP网络结构

前面提到已经把.bmp的图像转化为32×32的点阵数列，所以BP网络的输入层有IN=32×32=1 024个神经元，并且每个输入层单元的输入为0.0或1.0。

采用3层BP网络，根据Kolmogorov定理[13]和本实验所用的GPU中流多处理器（Stream Multiprocessor，SM）的数量为4个，一个SM中最大活动block数为8这两点，取隐含层神经元数目HN=32。输出有0～9共10个选择，所以输出层神经元数ON定为10。当已知目标值T为0时，与之对应的10个输出层单元的已知值定为1.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，以此类推，T为9时，与之对应的10个输出层单元的已知值定为0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，1.0。

4.3 实验环境和结果

实验所用的硬件环境为Intel Xeon E5430四核CPU和NVIDIA Quadro FX 1700 GPU。执行CUDA程序需要一个编译CPU端代码的编译器和一个编译GPU端代码的编译器，本实验中分别用Visual Studio 2008和CUDA Toolkit 4.0。操作系统为64位Windows7操作系统。

在初始化权重、偏倚相同，学习率相同的情况下，BP神经网络分别在CPU和GPU上训练，当取几组相同的训练样本且预先设定二者的训练次数使二者达到基本相同的训练误差时，所用时间对比如图3所示。可看出，随着训练样本数的增多，与在CPU上训练一样，在GPU上的训练时间也在增加，但是训练时间比CPU少很多，并且相比在CPU上训练加速比在不断提高，训练样本数为10 000时加速比为6.12，训练样本数为60 000时就提高到了8.17，并行方法起到了显著的加速作用。

图3 两种方法训练时间比较图

用两种方法得到的权重和偏倚数组分别对10 000幅测试集图片进行识别，结果如表1所示。

表1 两种方法训练后对测试集图片的正确识别率（%）

从表1中可以看出，分别采用几组不同的训练样本进行训练，用GPU上训练结果得到的识别率比用CPU上训练结果得到的高出0.05%～0.22%，识别结果验证了GPU上训练结果的正确性。

实验结果表明提出的BP算法并行化方法是高效可行的。

5 总结

本文利用CUDA模型实现了BP算法并行化，并用于手写数字训练，加速比为6.12～8.17，并用该训练结果对手写数字测试集图片进行识别进一步验证了提出的并行方法的正确性。但仍存在两点不足，一是提出的并行方法中用到的规约加法运算要求块内线程数是2的整数次幂，当具体实验数据不满足这个要求时要对每个块内线程数和数据根据具体情况进行扩充（可以把扩充的数据都赋初值0），有时候这种扩充过大时就进行了大量的无用运算，这一点还需改进。二是本实验用到的GPU流处理器（Stream Processor，SP）数量仅有32个，而目前一个GPU的最多的SP已达到几百个，因此本实验中CUDA模型的并行计算能力受到了一些限制。同时，CUDA模型对并行度不高的程序运行效率较低，所以选择合适的程序并行化才能更好地发挥CUDA的性能优势。

[1]冯百明，洪远麟，康继昌.MIMD系统上成批训练BP算法的并行划分[J].模式识别与人工智能，1998，11（1）：107-111.

[2]Su K，Jung K.GPU implementation of neural networks[J]. Elsevier，2004，37（6）：1311-1314.

[3]田绪江，江敏杰.GPU加速的神经网络BP算法[J].计算机应用研究，2009，26（5）：1679-1681.

[4]Scanzio S，Cumani S，Gemello R，et al.Parallel implementation of artificial neural network training[C]//IEEE International Conference on Acoustics Speech and Signal Processing，Dallas，TX，2010：4902-4905.

[5]Scanzio S，Cumani S，Gemello R，et al.Parallel implementation of artificial neural network training for speech recognition[J].Elsevier，2010，3（11）：1302-1309.

[6]Honghoon J，Anjin P，Keechul J.Neural network implementation using CUDA and OpenMP[C]//Digital Image Computing：Techniques and Application，Canberra，ACT，2008：155-161.

[7]Lin Jinxian，Lin Jianghong.Accelerating BP neural networkbased image compression by CPU and GPU cooperation[C]// IEEE International Conference on Multimedia Technology，2010：1-4.

[8]韩晓雪.基于人工神经网络和GPU加速的手写数字识别并行算法[D].辽宁大连：大连理工大学，2009.

[9]张舒，褚艳利，赵开勇，等.GPU高性能计算之CUDA[M].北京：中国水利水电出版社，2009.

[10]厉旭杰.GPU加速的图像匹配技术[J].计算机工程与应用，2012，48（2）：173-176.

[11]Han Jiawei，Kamber M.数据挖掘概念与技术[M].2版.范明，孟小峰，译.北京：机械工业出版社，2011：212-219.

[12]Lecun Y，Cortes C.The MNIST database of handwritten digits[EB/OL].[2012-05-10].http：//yann.lecun.com/exdb/mnist.

[13]Pandya A S，Macy R B.Pattern recognition with neural networks in C++[M].北京：电子工业出版社，1999.

SUN Xiangyu,FENG Baiming,YANG Pengfei

1.College of Computer Science and Engineering,Northwest Normal University,Lanzhou 730070,China 2.State Key Laboratory of Computer Architecture,Institute of Computing Technology,Chinese Academy of Sciences,Beijing 100190,China

CUDA is a generally used GPGPU（General Purpose Computing on GPU）model.BP algorithm is one of the most widely used neural network model at present.A method of parallelizing BP algorithm using CUDA is proposed in this paper. When this method are used to train BP neural network,data are transferred to GPU before training.Process of computing inputs, outputs,errors of hidden layer and output layer and updating weights,biases are realized on GPU.Training handwritten digital images with this method has speed-up ratio between 6.12 and 8.17 compared to training on four cores CPU.When this two results are respectively used to recognize the same test set,the recognition rate based on training result on GPU increases 0.05%～0.22%compared to that of CPU.

Back-Propagation（BP）algorithm;parallelization;Compute United Device Architecture（CUDA）;handwritten digits training

CUDA是应用较广的GPU通用计算模型，BP算法是目前应用最广泛的神经网络模型之一。提出了用CUDA模型并行化BP算法的方法。用该方法训练BP神经网络，训练开始前将数据传到GPU，训练开始后计算隐含层和输出层的输入输出和误差，更新权重和偏倚的过程都在GPU上实现。将该方法用于手写数字图片训练实验，与在四核CPU上的训练相比，加速比为6.12～8.17。分别用在CPU和GPU上训练得到的结果识别相同的测试集图片，GPU上的训练结果对图片的识别率比CPU上的高0.05%～0.22%。

向后传播算法；并行化；计算统一设备架构；手写数字训练

A

TP316.4

10.3778/j.issn.1002-8331.1210-0116

SUN Xiangyu,FENG Baiming,YANG Pengfei.Parallelization of BP algorithm and example verification based on CUDA. Computer Engineering and Applications,2013,49（23）：31-34.

计算机体系结构国家重点实验室开放课题资助（No.CARCH201105）。

孙香玉（1989—），女，硕士研究生，主研方向：GPU高性能计算；冯百明（1966—），男，教授，硕士生导师，主研方向：分布式与并行计算、GPU高性能计算；杨鹏斐（1985—），男，硕士研究生，主研方向：分布式与并行计算。E-mail：s_xiangyu@sina.cn

2012-10-12

2012-11-26

1002-8331（2013）23-0031-04

CNKI出版日期：2013-07-09 http://www.cnki.net/kcms/detail/11.2127.TP.20130709.1015.001.html