孫香玉，馮百明，楊鵬斐

1.西北師范大學(xué) 計算機科學(xué)與工程學(xué)院，蘭州 730070 2.中國科學(xué)院 計算技術(shù)研究所 計算機體系結(jié)構(gòu)國家重點實驗室，北京 100190

基于CUDA的BP算法并行化與實例驗證

孫香玉，馮百明，楊鵬斐

1.西北師范大學(xué) 計算機科學(xué)與工程學(xué)院，蘭州 730070 2.中國科學(xué)院 計算技術(shù)研究所 計算機體系結(jié)構(gòu)國家重點實驗室，北京 100190

1 引言

BP神經(jīng)網(wǎng)絡(luò)是人工神經(jīng)網(wǎng)絡(luò)的一種，它被廣泛應(yīng)用于各個領(lǐng)域，但是BP神經(jīng)網(wǎng)絡(luò)的訓(xùn)練過程需要進行大量費時的矩陣運算，所以經(jīng)常并行化BP算法以加快計算。

傳統(tǒng)的BP算法并行方法[1]需要使用大型并行機或網(wǎng)絡(luò)并且并行程序的編寫較為復(fù)雜，所以沒有在普通人群中推廣開來。文獻[2-3]把神經(jīng)網(wǎng)絡(luò)和BP神經(jīng)網(wǎng)絡(luò)的訓(xùn)練過程轉(zhuǎn)化為GPU紋理的渲染過程加快了神經(jīng)網(wǎng)絡(luò)訓(xùn)練過程，但是需要使用圖形學(xué)API編程，要求編程人員對圖形學(xué)硬件和編程接口有深入了解，開發(fā)難度大。文獻[4-5]在CUDA下并行化BP算法并進行了語音識別，加速比是在CPU上的6倍。文獻[6-7]在CUDA下并行化BP算法并分別用于文本檢測和圖像壓縮，加速比分別是在CPU上的8倍和9倍。CUDA不需要借助圖形學(xué)API，可直接用大多數(shù)人熟悉的C或C++語言編寫在GPU上執(zhí)行的程序，用起來較為簡便，基于CUDA模型實現(xiàn)的BP神經(jīng)網(wǎng)絡(luò)并行算法已成功用于語音識別，文本檢測，圖像壓縮，圖像分割等方面。文獻[8]用CUDA在GPU上加速了手寫數(shù)字識別，手寫數(shù)字的訓(xùn)練過程是在CPU上用神經(jīng)網(wǎng)絡(luò)訓(xùn)練實現(xiàn)的。針對當(dāng)訓(xùn)練樣本集過大時神經(jīng)網(wǎng)絡(luò)成批訓(xùn)練受限的問題提出了一種采用單樣本訓(xùn)練的方式與CUDA模型結(jié)合實現(xiàn)BP神經(jīng)網(wǎng)絡(luò)算法并行化的方法，有效節(jié)省了GPU存儲空間，并將該方法用于手寫數(shù)字訓(xùn)練，驗證了它的有效性。

2 CUDA架構(gòu)

CUDA模型采用單指令流，多數(shù)據(jù)流（Single Instruction Multiple Data，SIMD）執(zhí)行模式，將CPU作為主機（Host），GPU作為設(shè)備（Device），CPU與GPU協(xié)同工作，CPU、GPU各自擁有相互獨立的存儲地址空間。

運行在GPU上的CUDA并行計算函數(shù)稱為內(nèi)核函數(shù)（kernel），它不是一個完整的程序，而是整個CUDA程序中的一個可以并行執(zhí)行的步驟。一個完整的CUDA程序是由一系列設(shè)備端kernel函數(shù)并行步驟和主機端的串行步驟共同組成的。這些步驟會按照程序中相應(yīng)語句的順序執(zhí)行，滿足順序一致性。一個kernel函數(shù)中存在兩個層次的并行[9]，即線程格（grid）中的線程塊（block）間的并行和block中的線程（thread）間的并行。兩層并行模型是CUDA最重要的創(chuàng)新之一。CUDA編程模型[10]如圖1。

圖1CUDA編程模型

3 串行BP算法并行化

3.1 BP算法

BP神經(jīng)網(wǎng)絡(luò)是一種按誤差逆?zhèn)鞑ニ惴ㄓ?xùn)練的多層前饋網(wǎng)絡(luò)，它由一個輸入層，一個或多個隱含層和一個輸出層組成，含有一個隱含層的BP神經(jīng)網(wǎng)絡(luò)結(jié)構(gòu)如圖2所示。各層神經(jīng)元僅與相鄰層神經(jīng)元之間相互全連接，同層內(nèi)神經(jīng)元間無連接，各層神經(jīng)元間無反饋連接。每個輸出單元取前一層所有單元輸出的加權(quán)和作為輸入，將一個非線性（激勵）函數(shù)作用于加權(quán)輸入得到輸出，如此下去，直至得到輸出層的輸出。

圖2 多層前饋（BP）神經(jīng)網(wǎng)絡(luò)

采用單樣本訓(xùn)練方式的BP算法[11]描述如下：

（1）初始化網(wǎng)絡(luò)的權(quán)重數(shù)組W和偏倚數(shù)組b。

（2）對于隱含層或輸出層的每個單元j：

關(guān)于前一層i，計算單元j的輸入：

//對于兩層神經(jīng)網(wǎng)絡(luò)，前一層即為輸入層計算單元j的輸出：

（5）達到預(yù)先指定的訓(xùn)練次數(shù)或誤差精度要求則輸出權(quán)值和偏倚結(jié)果，否則繼續(xù)（2）至（4）。

3.2 串行實現(xiàn)

根據(jù)3.1節(jié)對BP算法的描述得其串行偽碼如下：

對N個樣本依次執(zhí)行第6～15行的訓(xùn)練，直至達到預(yù)定訓(xùn)練次數(shù)Pre_times。IN，HN，ON分別為輸入層，隱含層，輸出層神經(jīng)元數(shù)。

6～7行表示從N個樣本的輸入層輸入數(shù)組Study_DataI[N*IN]中得到當(dāng)前訓(xùn)練樣本m的輸入數(shù)組I[IN]，8～9行表示從N個樣本的輸出層已知目標(biāo)值Study_DataT[N*ON]中得到當(dāng)前訓(xùn)練樣本m的已知目標(biāo)值數(shù)組T[ON]。

根據(jù)3.1節(jié)中的公式（1）（2），通過Compute_H_out（）可得隱含層輸出H_out[HN]。同計算H_out的過程類似，通過Compute_O_out（）得到輸出層輸出O_out[ON]。

根據(jù)公式（3）和Compute_O_err（）可得輸出層誤差O_err[ON]。根據(jù)公式（4）和Compute_H_err（）可得隱含層誤差H_err[HN]。

根據(jù)公式（5）（6）和第14行操作可得更新后的權(quán)重O_H_W和偏倚O_bias[ON]。同理可得更新后的權(quán)重H_I_W和偏倚H_bias[HN]。

采用串行方法訓(xùn)練BP神經(jīng)網(wǎng)絡(luò)，當(dāng)訓(xùn)練樣本數(shù)達到幾萬時，為了取得較好的訓(xùn)練結(jié)果用于識別需要耗費至少幾個小時，所以考慮在GPU上利用CUDA模型加速訓(xùn)練過程。

3.3 并行實現(xiàn)

BP算法在CUDA模型下的并行偽碼如下。

在GPU上訓(xùn)練開始前，通過cudaMemcpy（）把CPU端的輸入、已知目標(biāo)輸出、所有權(quán)重和偏倚復(fù)制到GPU端，訓(xùn)練過程完全在GPU端并行實現(xiàn)，訓(xùn)練結(jié)束后再將訓(xùn)練結(jié)果（權(quán)重和偏倚）復(fù)制回CPU端，訓(xùn)練過程中不存在CPU端與GPU端的數(shù)據(jù)傳輸，減少了通信時間開銷。

并行偽碼中的9，10，11，12，13分別對串行偽碼中的10，11～12，13，14，15并行化。下面詳細介紹并行化實現(xiàn)過程。

BP并行偽碼中的9計算隱含層輸出。由于隱含層各神經(jīng)元的輸出只與輸入層所有神經(jīng)元的輸出相關(guān)，與隱含層其他神經(jīng)元的輸出不相關(guān)，所以可并行計算它們。〈〈〈>>>中HN表示核函數(shù)共啟動HN個塊，threadsPerBlock= min{IN，blockDim.x}為每個塊內(nèi)的線程數(shù)，blockDim.x為塊內(nèi)最大線程數(shù)，取IN和blockDim.x二者中較小值作為threadsPerBlock，第三個參數(shù)IN表示shared memory中外部數(shù)組d_I1的大小，為方便說明計算過程，假設(shè)blockDim.x= 512，IN=2*blockDim.x，并行化的具體實現(xiàn)如下。

//啟動HN個塊，每個塊內(nèi)threadsPerBlock個線程并行計算點積和，將每個線程上計算的部分和存到它所在塊的shared memory中的數(shù)組cache中

//每個塊內(nèi)進行并行規(guī)約（reduction）加法運算得到該塊內(nèi)所有線程點積運算的總和，存到cache[0]中

//HN個塊并行計算隱含層輸入

//HN個塊并行計算隱含層輸出

在上述偽碼6～11的規(guī)約運算中用到了以下優(yōu)化方法，一是GPU整數(shù)處理功能較弱，用位運算i>>=1代替i/=2，節(jié)省計算時間；二是相鄰thread操作數(shù)組內(nèi)的相鄰元素，避免了bank conflict，使訪問shared memory的速度與register的相同，都為global memory的100多倍。

BP并行偽碼中的第10行除了計算輸出層的輸出d_O_out（與上面計算d_H_out的過程類似，不再詳述）外，還將ON個塊并行計算出了輸出層的誤差d_O_out（參照公式（3））。

BP并行偽碼中的第11行，核函數(shù)啟動HN個塊，每個塊內(nèi)ON個線程，并行計算隱含層的誤差，詳細實現(xiàn)如下。

BP并行偽碼中的第12行，核函數(shù)啟動ON個塊，每個塊中HN個線程并行計算輸出層權(quán)重d_O_H_W和偏倚d_O_bias的更新，具體實現(xiàn)如下（并行偽碼中的第13行更新隱含層權(quán)重和偏倚的過程與之類似，不再說明）。

4 實驗及分析

用第3章中提出的BP算法并行化方法進行脫機手寫數(shù)字訓(xùn)練，用實例驗證其有效性。

4.1 手寫數(shù)字圖像集來源

實驗用到的手寫樣本圖像集來自于MNIST（Modified National Institute of Standards and Technology）手寫數(shù)字數(shù)據(jù)庫[12]，所要訓(xùn)練和識別的圖片全部為.bmp格式的黑白圖片，像素為28×28，為了使其更適合在CUDA模型上并行計算，讀取文件時即根據(jù)位圖信息把每幅圖片的數(shù)字表示成32×32的點陣數(shù)列，原圖像中的黑白像素點在點陣數(shù)列中分別用0.0，1.0表示，即進行二值化處理。

4.2 BP網(wǎng)絡(luò)結(jié)構(gòu)

前面提到已經(jīng)把.bmp的圖像轉(zhuǎn)化為32×32的點陣數(shù)列，所以BP網(wǎng)絡(luò)的輸入層有IN=32×32=1 024個神經(jīng)元，并且每個輸入層單元的輸入為0.0或1.0。

采用3層BP網(wǎng)絡(luò)，根據(jù)Kolmogorov定理[13]和本實驗所用的GPU中流多處理器（Stream Multiprocessor，SM）的數(shù)量為4個，一個SM中最大活動block數(shù)為8這兩點，取隱含層神經(jīng)元數(shù)目HN=32。輸出有0～9共10個選擇，所以輸出層神經(jīng)元數(shù)ON定為10。當(dāng)已知目標(biāo)值T為0時，與之對應(yīng)的10個輸出層單元的已知值定為1.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，以此類推，T為9時，與之對應(yīng)的10個輸出層單元的已知值定為0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，0.0，1.0。

4.3 實驗環(huán)境和結(jié)果

實驗所用的硬件環(huán)境為Intel Xeon E5430四核CPU和NVIDIA Quadro FX 1700 GPU。執(zhí)行CUDA程序需要一個編譯CPU端代碼的編譯器和一個編譯GPU端代碼的編譯器，本實驗中分別用Visual Studio 2008和CUDA Toolkit 4.0。操作系統(tǒng)為64位Windows7操作系統(tǒng)。

在初始化權(quán)重、偏倚相同，學(xué)習(xí)率相同的情況下，BP神經(jīng)網(wǎng)絡(luò)分別在CPU和GPU上訓(xùn)練，當(dāng)取幾組相同的訓(xùn)練樣本且預(yù)先設(shè)定二者的訓(xùn)練次數(shù)使二者達到基本相同的訓(xùn)練誤差時，所用時間對比如圖3所示?？煽闯?，隨著訓(xùn)練樣本數(shù)的增多，與在CPU上訓(xùn)練一樣，在GPU上的訓(xùn)練時間也在增加，但是訓(xùn)練時間比CPU少很多，并且相比在CPU上訓(xùn)練加速比在不斷提高，訓(xùn)練樣本數(shù)為10 000時加速比為6.12，訓(xùn)練樣本數(shù)為60 000時就提高到了8.17，并行方法起到了顯著的加速作用。

圖3 兩種方法訓(xùn)練時間比較圖

用兩種方法得到的權(quán)重和偏倚數(shù)組分別對10 000幅測試集圖片進行識別，結(jié)果如表1所示。

表1 兩種方法訓(xùn)練后對測試集圖片的正確識別率（%）

從表1中可以看出，分別采用幾組不同的訓(xùn)練樣本進行訓(xùn)練，用GPU上訓(xùn)練結(jié)果得到的識別率比用CPU上訓(xùn)練結(jié)果得到的高出0.05%～0.22%，識別結(jié)果驗證了GPU上訓(xùn)練結(jié)果的正確性。

實驗結(jié)果表明提出的BP算法并行化方法是高效可行的。

5 總結(jié)

本文利用CUDA模型實現(xiàn)了BP算法并行化，并用于手寫數(shù)字訓(xùn)練，加速比為6.12～8.17，并用該訓(xùn)練結(jié)果對手寫數(shù)字測試集圖片進行識別進一步驗證了提出的并行方法的正確性。但仍存在兩點不足，一是提出的并行方法中用到的規(guī)約加法運算要求塊內(nèi)線程數(shù)是2的整數(shù)次冪，當(dāng)具體實驗數(shù)據(jù)不滿足這個要求時要對每個塊內(nèi)線程數(shù)和數(shù)據(jù)根據(jù)具體情況進行擴充（可以把擴充的數(shù)據(jù)都賦初值0），有時候這種擴充過大時就進行了大量的無用運算，這一點還需改進。二是本實驗用到的GPU流處理器（Stream Processor，SP）數(shù)量僅有32個，而目前一個GPU的最多的SP已達到幾百個，因此本實驗中CUDA模型的并行計算能力受到了一些限制。同時，CUDA模型對并行度不高的程序運行效率較低，所以選擇合適的程序并行化才能更好地發(fā)揮CUDA的性能優(yōu)勢。

[1]馮百明，洪遠麟，康繼昌.MIMD系統(tǒng)上成批訓(xùn)練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加速的神經(jīng)網(wǎng)絡(luò)BP算法[J].計算機應(yīng)用研究，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]韓曉雪.基于人工神經(jīng)網(wǎng)絡(luò)和GPU加速的手寫數(shù)字識別并行算法[D].遼寧大連：大連理工大學(xué)，2009.

[9]張舒，褚艷利，趙開勇，等.GPU高性能計算之CUDA[M].北京：中國水利水電出版社，2009.

[10]厲旭杰.GPU加速的圖像匹配技術(shù)[J].計算機工程與應(yīng)用，2012，48（2）：173-176.

[11]Han Jiawei，Kamber M.數(shù)據(jù)挖掘概念與技術(shù)[M].2版.范明，孟小峰，譯.北京：機械工業(yè)出版社，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].北京：電子工業(yè)出版社，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是應(yīng)用較廣的GPU通用計算模型，BP算法是目前應(yīng)用最廣泛的神經(jīng)網(wǎng)絡(luò)模型之一。提出了用CUDA模型并行化BP算法的方法。用該方法訓(xùn)練BP神經(jīng)網(wǎng)絡(luò)，訓(xùn)練開始前將數(shù)據(jù)傳到GPU，訓(xùn)練開始后計算隱含層和輸出層的輸入輸出和誤差，更新權(quán)重和偏倚的過程都在GPU上實現(xiàn)。將該方法用于手寫數(shù)字圖片訓(xùn)練實驗，與在四核CPU上的訓(xùn)練相比，加速比為6.12～8.17。分別用在CPU和GPU上訓(xùn)練得到的結(jié)果識別相同的測試集圖片，GPU上的訓(xùn)練結(jié)果對圖片的識別率比CPU上的高0.05%～0.22%。

向后傳播算法；并行化；計算統(tǒng)一設(shè)備架構(gòu)；手寫數(shù)字訓(xùn)練

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.

計算機體系結(jié)構(gòu)國家重點實驗室開放課題資助（No.CARCH201105）。

孫香玉（1989—），女，碩士研究生，主研方向：GPU高性能計算；馮百明（1966—），男，教授，碩士生導(dǎo)師，主研方向：分布式與并行計算、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