神經(jīng)網(wǎng)絡(luò)前向傳播在GPU上的實(shí)現(xiàn)

摘 要： 基于CUDA架構(gòu)在GPU上實(shí)現(xiàn)了神經(jīng)網(wǎng)絡(luò)前向傳播算法，該算法利用神經(jīng)網(wǎng)絡(luò)各層內(nèi)神經(jīng)元計(jì)算的并行性，每層使用一個(gè)Kernel函數(shù)來并行計(jì)算該層神經(jīng)元的值，每個(gè)Kernel函數(shù)都根據(jù)神經(jīng)網(wǎng)絡(luò)的特性和CUDA架構(gòu)的特點(diǎn)進(jìn)行優(yōu)化。實(shí)驗(yàn)表明，該算法比普通的CPU上的算法快了約7倍。研究結(jié)果對(duì)于提高神經(jīng)網(wǎng)絡(luò)的運(yùn)算速度以及CUDA的適用場合都有參考價(jià)值。
關(guān)鍵詞： 神經(jīng)網(wǎng)絡(luò)；CUDA；GPU

GPU（圖形處理器）是處理計(jì)算機(jī)圖形的專用設(shè)備。近十年來，由于高清晰度復(fù)雜圖形實(shí)時(shí)處理的需求，GPU發(fā)展成為高并行度、多線程、多核的處理器。目前主流GPU的運(yùn)算能力已超過主流通用CPU，從發(fā)展趨勢上看將來差距會(huì)越來越大。GPU卓越的性能對(duì)開發(fā)GPGPU（使用GPU進(jìn)行通用計(jì)算）非常具有吸引力。最初GPGPU需要將非圖形應(yīng)用映射為圖形應(yīng)用的方式，這個(gè)處理過程與圖形硬件緊密相關(guān)，程序?qū)崿F(xiàn)非常困難。近年來，GPU的主要供應(yīng)商N(yùn)VIDIA提出了新的GPGPU模型，稱為CUDA（統(tǒng)一計(jì)算設(shè)備架構(gòu)）[1]。
CUDA是一種在NVIDIA公司的GPU上進(jìn)行計(jì)算的新型的硬件和軟件架構(gòu)，可以將GPU視為一個(gè)并行數(shù)據(jù)計(jì)算的設(shè)備，對(duì)所進(jìn)行的計(jì)算給予分配和管理。在CUDA的架構(gòu)中，這些計(jì)算不再像過去的GPGPU架構(gòu)那樣必須將計(jì)算映射到圖形API中，開發(fā)者不需要去學(xué)習(xí)艱澀的顯示芯片的指令或是特殊的結(jié)構(gòu)，因此開發(fā)門檻大大降低。CUDA的GPU編程語言基于標(biāo)準(zhǔn)的C語言，開發(fā)CUDA的應(yīng)用程序較為方便。
在CUDA架構(gòu)下，一個(gè)程序分為兩個(gè)部分：host端和device端。host端是指在CPU上執(zhí)行的部分，而device端則是在GPU上執(zhí)行的部分。device端的程序又稱為Kernel。通常host端程序會(huì)將數(shù)據(jù)準(zhǔn)備好后，復(fù)制到顯卡的內(nèi)存中，再由GPU執(zhí)行device端程序，完成后再由host端程序?qū)⒔Y(jié)果從顯卡的內(nèi)存中取回。
在CUDA架構(gòu)下，GPU執(zhí)行時(shí)的最小單位是線程（thread）。幾個(gè)線程可以組成一個(gè)線程塊（block），每個(gè)線程都有自己的一份寄存器（register）和本地內(nèi)存（local memory）空間。同一個(gè)線程塊中的每個(gè)線程則可以共享一份共享內(nèi)存(shared memory)。此外，所有的線程都共享一份全局內(nèi)存(global memory)、常量內(nèi)存(constant memory)和紋理內(nèi)存（texture memory）。
自從CUDA發(fā)布以來,大量應(yīng)用問題使用CUDA技術(shù)取得了令人驚奇的效果，加速十幾倍甚至上百倍的實(shí)例很多。參考文獻(xiàn)[2]對(duì)這些應(yīng)用有一個(gè)綜合介紹。
最適合利用 CUDA 處理的問題，是可以大量并行化的問題，這樣能有效利用顯示芯片上的大量執(zhí)行單元。使用 CUDA 時(shí)，上千個(gè)線程同時(shí)執(zhí)行是很正常的。如果問題不能大量并行化，使用 CUDA 就不能達(dá)到最好的效率。參考文獻(xiàn)[3]通過許多成功應(yīng)用CUDA加速的實(shí)例探討了CUDA的高效實(shí)現(xiàn)策略及應(yīng)用范圍問題。
1 本文使用的神經(jīng)網(wǎng)絡(luò)簡介
本文實(shí)驗(yàn)選用的是一個(gè)較為復(fù)雜的用于手寫漢字識(shí)別的神經(jīng)網(wǎng)絡(luò)[4-6]。該網(wǎng)絡(luò)是一種五層卷積神經(jīng)網(wǎng)絡(luò)，結(jié)構(gòu)如圖1所示。選擇這樣一個(gè)神經(jīng)網(wǎng)絡(luò)基于兩點(diǎn)考慮：其一，過于簡單的神經(jīng)網(wǎng)絡(luò)結(jié)構(gòu)不易體現(xiàn)GPU計(jì)算的優(yōu)勢；其二，此網(wǎng)絡(luò)的結(jié)構(gòu)比較全面，在它上面實(shí)現(xiàn)的算法有一定的代表性，做一些修改就可以用在其他不同的神經(jīng)網(wǎng)絡(luò)結(jié)構(gòu)上。

該神經(jīng)網(wǎng)絡(luò)第一層輸入層是手寫數(shù)字的灰度圖像，圖像大小為29×29像素，這樣輸入層有29×29=841個(gè)神經(jīng)元。第二層是卷積層，由6個(gè)特征圖組成，每個(gè)特征圖的尺寸是13×13個(gè)神經(jīng)元，每個(gè)神經(jīng)元的值是由輸入圖像的13×13個(gè)位置(輸入圖像水平和垂直方向的29個(gè)位置間隔選擇13個(gè))上作5×5的卷積得到。一個(gè)特征圖的卷積核相同，各特征圖的卷積核不同。這樣第二層有13×13×6=1014個(gè)神經(jīng)元，有(5×5+1)×6=156個(gè)權(quán)值(每個(gè)5×5多加一個(gè)偏向權(quán)值，后面各層中權(quán)值數(shù)量計(jì)算公式中＋１都是加一個(gè)偏向權(quán)值)。另外，因?yàn)? 014個(gè)神經(jīng)元每個(gè)都有26個(gè)與輸入層的連接，所以從輸入層到第二層共有1014×26=26364個(gè)連接。這里就能看出卷積神經(jīng)網(wǎng)絡(luò)共享權(quán)值的好處了：雖然有26 364個(gè)連接，但只需要156個(gè)權(quán)值來控制。第三層也是一個(gè)卷積層，它有50個(gè)特征圖構(gòu)成，每個(gè)特征圖是5×5大小，是從第二層的6個(gè)13×13的特征圖的相應(yīng)區(qū)域作5×5卷積得到。這樣第三層就有5×5×50=1250個(gè)神經(jīng)元，(5×5+1)×6×50=7800個(gè)權(quán)值，1250×26=32500個(gè)與第二層的連接。第四層是有100個(gè)神經(jīng)元的全連接層，因?yàn)槭侨B接，本層的每個(gè)神經(jīng)元都與第三層的全部1 250個(gè)神經(jīng)元連接，這樣所有的連接就有100×(1250+1)=125100個(gè)，有125 100個(gè)權(quán)值。第五層是輸出層，該層有10個(gè)全連接的單元，每個(gè)單元都與第四層的所有100個(gè)神經(jīng)元連接，這樣本層就有10×(100+1)=1010個(gè)連接，有1010個(gè)權(quán)值。這10個(gè)神經(jīng)元對(duì)應(yīng)10個(gè)數(shù)字 ，其中對(duì)應(yīng)于識(shí)別結(jié)果的一個(gè)神經(jīng)元的輸出值為+1，其他9個(gè)神經(jīng)元的輸出值為-1。
整個(gè)網(wǎng)絡(luò)總共有3 215個(gè)神經(jīng)元，134 066個(gè)權(quán)值，184 974個(gè)連接。輸入層神經(jīng)元數(shù)據(jù)和各層神經(jīng)網(wǎng)絡(luò)的權(quán)值已知。

GPU上神經(jīng)網(wǎng)絡(luò)前向傳播算法基本過程是逐層計(jì)算各層的所有神經(jīng)元的值。
輸入層神經(jīng)元值已知，其余每層有一個(gè)Kernel函數(shù)來計(jì)算該層的所有神經(jīng)元的值，上述的神經(jīng)網(wǎng)絡(luò)需要4個(gè)Kernel函數(shù)。并行計(jì)算只能體現(xiàn)在一層中，不同層之間沒有并行性。
首先將輸入層的神經(jīng)元值和每層的權(quán)值保存在5個(gè)數(shù)組中，并從host內(nèi)存?zhèn)鬟f到device內(nèi)存。由于每層的權(quán)值是不變的，所以可以將這些權(quán)值傳遞到device的常量內(nèi)存中，由于常量內(nèi)存有cache，這比放到全局內(nèi)存的存取速度要快很多。在device中為第二到第五層的神經(jīng)元值分配內(nèi)存空間，第一個(gè)Kernel函數(shù)根據(jù)輸入層的神經(jīng)元值和權(quán)值計(jì)算第二層神經(jīng)元值，第二個(gè)Kernel函數(shù)根據(jù)第二層的神經(jīng)元值和權(quán)值計(jì)算第三層神經(jīng)元值，如此往下，第四個(gè)Kernel函數(shù)計(jì)算出第五層即輸出層的值，然后將該值從device內(nèi)存?zhèn)鬟f到host內(nèi)存。神經(jīng)網(wǎng)絡(luò)的連接體現(xiàn)在每個(gè)Kernel函數(shù)處理計(jì)算過程里。
計(jì)算第二層神經(jīng)元值的基本Kernel函數(shù)的偽代碼如下：
1：bid=blockIdx.x;
2：tx=threadIdx.x; ty=threadIdx.y;
3：wt=ty*2*29+tx*2; result=0;
4：templet[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};
5：for(i=0;i<25;++i)
6： result+=Gn1[wt+templet[i]]*Gw1[bid*26+i+1];
7：result=(1.7159*tanhf(0.66666667*result));
8：Gn2[13*13*bid+ty*13+tx]=result;
該函數(shù)的輸入為第一層的神經(jīng)元值數(shù)組Gn1和權(quán)值數(shù)組Gw1，輸出為第二層的神經(jīng)元值數(shù)組Gn2。
因?yàn)榈诙佑?個(gè)特征圖構(gòu)成，每個(gè)特征圖有13×13個(gè)神經(jīng)元，所以這個(gè)函數(shù)在host端調(diào)用時(shí)，線程塊參數(shù)設(shè)置為6，線程參數(shù)設(shè)置為13×13，這樣一個(gè)線程塊處理一個(gè)特征圖，每個(gè)神經(jīng)元值由一個(gè)線程處理，由第一層中取25個(gè)神經(jīng)元乘以權(quán)值再加上一個(gè)偏向權(quán)值的累加得到。
這個(gè)函數(shù)的第1、2行得到線程塊號(hào)和線程號(hào)，第4行的templet數(shù)組是為了在輸入層中選擇25個(gè)值的模版。第5、6行的循環(huán)計(jì)算出激活值，第7行對(duì)該值進(jìn)行激活運(yùn)算，第8行將結(jié)果送到第二層的神經(jīng)元值數(shù)組Gn2。
上述函數(shù)還可以修改的效率更高些，第4行的templet數(shù)組是默認(rèn)分配在全局內(nèi)存中的，而訪問全局內(nèi)存需要相對(duì)較長的時(shí)間，因此有必要消除這種訪存時(shí)間消耗。通過將第5、6行的for循環(huán)展開，templet[i]的值直接寫在代碼里，就既可以在程序中不要templet數(shù)組，又減少了循環(huán)判斷的時(shí)間代價(jià)。
對(duì)函數(shù)的4～6行進(jìn)行修改，其他行不變。改變了的部分代碼如下：
4：__shared__ float s1[29*29];
5：//將Gn1數(shù)組元素送到s1數(shù)組；
6：result=s1[wt]*Gw1[bid*26+1]+
s1[wt+1]*Gw1[bid*26+2]+
s1[wt+2]*Gw1[bid*26+3]+…
//其他省略，一共有25個(gè)乘累加
第4行定義了一個(gè)在共享內(nèi)存中的數(shù)組s1。將Gn1數(shù)組中的數(shù)據(jù)先傳到s1中，第6行就不再需要引用Gn1數(shù)組，而是使用數(shù)組s1。由于存取共享內(nèi)存的速度比存取全局內(nèi)存快得多，這種改進(jìn)進(jìn)一步提高了速度。
計(jì)算第三層神經(jīng)元值的Kernel函數(shù)，基本結(jié)構(gòu)與上述的函數(shù)類似。也可以使用類似處理方式進(jìn)一步提高速度。該函數(shù)在host端調(diào)用時(shí)，線程塊數(shù)參數(shù)設(shè)置為50，線程參數(shù)設(shè)置為5×5。一個(gè)線程塊處理一個(gè)特征圖，每個(gè)神經(jīng)元值由一個(gè)線程處理。
第三層到第四層是全連接層，計(jì)算第四層的Kernel函數(shù)可以有兩種實(shí)現(xiàn)方案。方案1：由于第四層有100個(gè)神經(jīng)元，可以設(shè)計(jì)成有100個(gè)線程塊，每個(gè)線程塊只有一個(gè)線程，一個(gè)線程計(jì)算一個(gè)神經(jīng)元的值。這樣一個(gè)線程計(jì)算量還是很大的，大約需要作1 250個(gè)乘法和加法。另外，每個(gè)線程塊只有一個(gè)線程，無法充分利用CUDA的Warps切換能力[1]。這種方案沒有完全發(fā)揮GPU的能力，因而效率不太高。方案2：還是有100個(gè)線程塊，每個(gè)線程塊有250個(gè)線程，250個(gè)線程計(jì)算一個(gè)神經(jīng)元的值，每個(gè)線程作5次乘法和加法，最后由第一個(gè)線程將這些和累加，得到最終的值。這種方案效率較高。
計(jì)算第五層的神經(jīng)元值需要約1 010次浮點(diǎn)乘法和加法，實(shí)驗(yàn)證明該層由于計(jì)算量相對(duì)較小，體現(xiàn)不出在Kernel端計(jì)算的優(yōu)勢，計(jì)算時(shí)間甚至比在host端更長。
3 實(shí)驗(yàn)結(jié)果
本試驗(yàn)使用機(jī)器的CPU是Intel Core 2 Duo E7400，時(shí)鐘頻率為2.8 GHz，內(nèi)存為1 GB。GPU是NVIDIA GeForce GTX9800+，該GPU有128個(gè)頻率為1.836 GHz的流處理器，分為16個(gè)SM，512 MB顯存。
GPU上算法實(shí)現(xiàn)采用上文描述的最優(yōu)化的方案，即每層的權(quán)值放入device端的常量內(nèi)存，函數(shù)中的循環(huán)展開，利用共享內(nèi)存，計(jì)算第四層神經(jīng)元值時(shí)采用方案2。
試驗(yàn)結(jié)果如表1所示，分別給出了使用CPU和GPU計(jì)算本文神經(jīng)網(wǎng)絡(luò)前向傳播時(shí)各層所用時(shí)間對(duì)比，這是經(jīng)過三次試驗(yàn)求得的平均值。GPU實(shí)現(xiàn)采用最優(yōu)化的方案，CPU上的實(shí)現(xiàn)方法就是通常的神經(jīng)網(wǎng)絡(luò)前向傳播算法，這里不再贅述。CPU上計(jì)算整個(gè)神經(jīng)網(wǎng)絡(luò)前向傳播需要1.851 ms，GPU計(jì)算需要0.272 ms，比CPU上計(jì)算快了約7倍。

本文對(duì)神經(jīng)網(wǎng)絡(luò)的前向傳播在GPU上計(jì)算進(jìn)行了研究，得到了比較滿意的結(jié)果。使用本方法需要注意以下幾個(gè)方面：