基于CUDA的BP算法并行化與實(shí)例驗(yàn)證

孫香玉，馮百明，楊鵬斐

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

基于CUDA的BP算法并行化與實(shí)例驗(yàn)證

孫香玉，馮百明，楊鵬斐

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

1 引言

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

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

2 CUDA架構(gòu)

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

運(yùn)行在GPU上的CUDA并行計(jì)算函數(shù)稱為內(nèi)核函數(shù)（kernel），它不是一個(gè)完整的程序，而是整個(gè)CUDA程序中的一個(gè)可以并行執(zhí)行的步驟。一個(gè)完整的CUDA程序是由一系列設(shè)備端kernel函數(shù)并行步驟和主機(jī)端的串行步驟共同組成的。這些步驟會(huì)按照程序中相應(yīng)語(yǔ)句的順序執(zhí)行，滿足順序一致性。一個(gè)kernel函數(shù)中存在兩個(gè)層次的并行[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ò)，它由一個(gè)輸入層，一個(gè)或多個(gè)隱含層和一個(gè)輸出層組成，含有一個(gè)隱含層的BP神經(jīng)網(wǎng)絡(luò)結(jié)構(gòu)如圖2所示。各層神經(jīng)元僅與相鄰層神經(jīng)元之間相互全連接，同層內(nèi)神經(jīng)元間無連接，各層神經(jīng)元間無反饋連接。每個(gè)輸出單元取前一層所有單元輸出的加權(quán)和作為輸入，將一個(gè)非線性（激勵(lì)）函數(shù)作用于加權(quán)輸入得到輸出，如此下去，直至得到輸出層的輸出。

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

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

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

（2）對(duì)于隱含層或輸出層的每個(gè)單元j：

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

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

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

3.2 串行實(shí)現(xiàn)

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

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

6～7行表示從N個(gè)樣本的輸入層輸入數(shù)組Study_DataI[N*IN]中得到當(dāng)前訓(xùn)練樣本m的輸入數(shù)組I[IN]，8～9行表示從N個(gè)樣本的輸出層已知目標(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]。同計(jì)算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ù)達(dá)到幾萬時(shí)，為了取得較好的訓(xùn)練結(jié)果用于識(shí)別需要耗費(fèi)至少幾個(gè)小時(shí)，所以考慮在GPU上利用CUDA模型加速訓(xùn)練過程。

3.3 并行實(shí)現(xiàn)

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

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

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

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

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

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

//HN個(gè)塊并行計(jì)算隱含層輸入

//HN個(gè)塊并行計(jì)算隱含層輸出

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

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

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

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

4 實(shí)驗(yàn)及分析

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

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

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

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

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

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

4.3 實(shí)驗(yàn)環(huán)境和結(jié)果

實(shí)驗(yàn)所用的硬件環(huán)境為Intel Xeon E5430四核CPU和NVIDIA Quadro FX 1700 GPU。執(zhí)行CUDA程序需要一個(gè)編譯CPU端代碼的編譯器和一個(gè)編譯GPU端代碼的編譯器，本實(shí)驗(yàn)中分別用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ù)使二者達(dá)到基本相同的訓(xùn)練誤差時(shí)，所用時(shí)間對(duì)比如圖3所示。可看出，隨著訓(xùn)練樣本數(shù)的增多，與在CPU上訓(xùn)練一樣，在GPU上的訓(xùn)練時(shí)間也在增加，但是訓(xùn)練時(shí)間比CPU少很多，并且相比在CPU上訓(xùn)練加速比在不斷提高，訓(xùn)練樣本數(shù)為10 000時(shí)加速比為6.12，訓(xùn)練樣本數(shù)為60 000時(shí)就提高到了8.17，并行方法起到了顯著的加速作用。

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

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

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

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

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

5 總結(jié)

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

[1]馮百明，洪遠(yuǎn)麟，康繼昌.MIMD系統(tǒng)上成批訓(xùn)練BP算法的并行劃分[J].模式識(shí)別與人工智能，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].計(jì)算機(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ù)字識(shí)別并行算法[D].遼寧大連：大連理工大學(xué)，2009.

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

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

[11]Han Jiawei，Kamber M.數(shù)據(jù)挖掘概念與技術(shù)[M].2版.范明，孟小峰，譯.北京：機(jī)械工業(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通用計(jì)算模型，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)練開始后計(jì)算隱含層和輸出層的輸入輸出和誤差，更新權(quán)重和偏倚的過程都在GPU上實(shí)現(xiàn)。將該方法用于手寫數(shù)字圖片訓(xùn)練實(shí)驗(yàn)，與在四核CPU上的訓(xùn)練相比，加速比為6.12～8.17。分別用在CPU和GPU上訓(xùn)練得到的結(jié)果識(shí)別相同的測(cè)試集圖片，GPU上的訓(xùn)練結(jié)果對(duì)圖片的識(shí)別率比CPU上的高0.05%～0.22%。

向后傳播算法；并行化；計(jì)算統(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.

計(jì)算機(jī)體系結(jié)構(gòu)國(guó)家重點(diǎn)實(shí)驗(yàn)室開放課題資助（No.CARCH201105）。

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