摘  要： 基于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é)果對于提高神經(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卓越的性能對開發(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è)備，對所進(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]對這些應(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)元對應(yīng)10個(gè)數(shù)字 ，其中對應(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行得到線程塊號和線程號，第4行的templet數(shù)組是為了在輸入層中選擇25個(gè)值的模版。第5、6行的循環(huán)計(jì)算出激活值，第7行對該值進(jìn)行激活運(yùn)算，第8行將結(jié)果送到第二層的神經(jīng)元值數(shù)組Gn2。
    上述函數(shù)還可以修改的效率更高些，第4行的templet數(shù)組是默認(rèn)分配在全局內(nèi)存中的，而訪問全局內(nèi)存需要相對較長的時(shí)間，因此有必要消除這種訪存時(shí)間消耗。通過將第5、6行的for循環(huán)展開，templet[i]的值直接寫在代碼里，就既可以在程序中不要templet數(shù)組，又減少了循環(huán)判斷的時(shí)間代價(jià)。
    對函數(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ì)算量相對較小，體現(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í)間對比，這是經(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倍。

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

    (1)GPU算法前期處理時(shí)要把一些數(shù)據(jù)從CPU裝入GPU，還需要在GPU上給一些數(shù)據(jù)分配內(nèi)存空間，這些前期處理都要花費(fèi)時(shí)間，所以如果只做一次神經(jīng)網(wǎng)絡(luò)前向傳播運(yùn)算，GPU上的實(shí)現(xiàn)可能比CPU上的還要慢，但由于前期處理只做一次，所以多次神經(jīng)網(wǎng)絡(luò)前向傳播運(yùn)算能體現(xiàn)出GPU運(yùn)算的好處。這種情況可能出現(xiàn)在多個(gè)測試數(shù)據(jù)的驗(yàn)證或計(jì)算神經(jīng)網(wǎng)絡(luò)后向傳播時(shí)。
    (2)當(dāng)神經(jīng)網(wǎng)絡(luò)規(guī)模較小時(shí)，可能GPU方法會(huì)比CPU方法慢，比如本網(wǎng)絡(luò)的最后一層的規(guī)模。
    (3)有些應(yīng)用問題在GPU上計(jì)算時(shí)速度能提高上百倍[3]，而神經(jīng)網(wǎng)絡(luò)在GPU上運(yùn)算速度提高有限，這是因?yàn)樯窠?jīng)網(wǎng)絡(luò)的并行性主要體現(xiàn)在同一層，總體并行性不太高。要想充分發(fā)揮GPU的運(yùn)算能力，那些應(yīng)用問題應(yīng)該具有計(jì)算復(fù)雜度高、數(shù)據(jù)傳輸量少的特點(diǎn)，或者能修改原來的算法，使得計(jì)算-內(nèi)存訪問比提高。
參考文獻(xiàn)
[1] NVIDIA Corporation.NVIDIA CUDA programming Guide Version 2.2[EB/OL]，2009-04.http：//developer.nvidia.com/cuda.
[2] RYOO S，RODRIGUES C I，BAGHSORKHI S S，et al. Optimization principles and application performance evaluation of a multithreaded GPU using CUDA[J].In Proceedings  of ACM SIGPLAN Symposium on Principles and Practice of  Parallel Programming，2008：73-82.
[3] HWU W M，RODRIGUES C，RYOO S，et al.Compute  unified device architecture application suitability[J].Computing in Science and Engineering，2009，11(3)：16-26.
[4] LECUN Y，BOTTOU L，BENGIO Y，et al.Gradient-based learning applied to document recognition[J].Proceedings  of the IEEE，1998，86（11）：2278-2324.
[5] SIMARD P Y，STEINKRAUS D，PLATT J.Best practices  for convolutional neural networks applied to visual document analysis[C].International Conference on Document Analysis and Recognition(ICDAR)，IEEE Computer Society，Los Alamitos，2003：958-962.
[6] Mike O′Neill.Neural network for recognition of handwritten digits[EB/OL]，2006-12.http：//www.codeproject.com/KB/library/NeuralNetRecognition.aspx.