費(fèi)雄偉,李肯立,陽(yáng)王東
(1.湖南城市學(xué)院信息科學(xué)與工程學(xué)院,湖南益陽(yáng)413000;2.湖南大學(xué)信息科學(xué)與工程學(xué)院,長(zhǎng)沙410008)
基于CUDA的AES并行算法優(yōu)化
費(fèi)雄偉1,2,李肯立2,陽(yáng)王東1,2
(1.湖南城市學(xué)院信息科學(xué)與工程學(xué)院,湖南益陽(yáng)413000;2.湖南大學(xué)信息科學(xué)與工程學(xué)院,長(zhǎng)沙410008)
為提升高級(jí)加密標(biāo)準(zhǔn)(AES)的加密性能,利用顯卡的通用計(jì)算能力,在統(tǒng)一計(jì)算設(shè)備架構(gòu)(CUDA)平臺(tái)上實(shí)現(xiàn)AES的128位、192位和256位3個(gè)版本的GPU并行算法,并提出優(yōu)化的AES并行算法。在考慮塊內(nèi)線程數(shù)量、共享存儲(chǔ)器容量和總塊數(shù)的基礎(chǔ)上,根據(jù)分塊最優(yōu)值的經(jīng)驗(yàn)數(shù)據(jù)指導(dǎo)AES算法在GPU上的最優(yōu)分塊。實(shí)驗(yàn)結(jié)果表明,與未優(yōu)化的AES并行算法相比,該算法的3個(gè)版本在Nvidia Geforce G210顯卡上的加密速度分別提高5.28%,14.55%和12.53%,而在 Nvidia Geforce GTX460顯卡上的加密速度分別提高 12.48%,15.40%和15.84%,且能更好地對(duì)SSL數(shù)據(jù)進(jìn)行加密。
分塊;經(jīng)驗(yàn)數(shù)據(jù);并行算法;優(yōu)化;高級(jí)加密標(biāo)準(zhǔn);統(tǒng)一計(jì)算設(shè)備架構(gòu)
隨著電子商務(wù)和電子金融的用戶群體擴(kuò)大和對(duì)安全的要求更高,加密及其加密性能成為亟待處理的問(wèn)題。對(duì)加密算法,首先要考慮的是它的安全性能,但在實(shí)際應(yīng)用時(shí)還需考慮它的性能和成本。高級(jí)加密標(biāo)準(zhǔn)(Advanced Encryption Standard,AES)算法在大量需要安全通信的應(yīng)用中擔(dān)當(dāng)加密、解密任務(wù),很好地協(xié)調(diào)了安全、高效和簡(jiǎn)明等特性。但隨著網(wǎng)絡(luò)用戶的迅速增加,對(duì)大規(guī)模加密任務(wù)應(yīng)用服務(wù)器(如電子商務(wù)或電子金融)的加密性能提出了更高的要求。為了提升AES的加密性能,最近許多基于FPGA和基于GPU加速的AES算法被提出?;贔PGA的加速AES主要應(yīng)用于無(wú)線和藍(lán)牙通信環(huán)境的加速加密,需要考慮能耗和大小等問(wèn)題。比如文獻(xiàn)[1]在Xilinx Spartan-III和Xilinx Spartan-II這2類設(shè)備上對(duì)AES進(jìn)行了速度和面積大小方面的對(duì)比研究;文獻(xiàn)[2]提出了適應(yīng)無(wú)線局域網(wǎng)環(huán)境的AES高速FPGA實(shí)現(xiàn);文獻(xiàn)[3]在FPGA上實(shí)現(xiàn)了AES-128、AES-192和AES-256,并進(jìn)行優(yōu)化,達(dá)到了節(jié)能的要求,適合于移動(dòng)的設(shè)備;文獻(xiàn)[4]提出了高吞吐量的AES引擎,采用FPGA實(shí)現(xiàn)了面積優(yōu)化36.2%的AES算法。但采用FPGA這類專用硬件設(shè)計(jì)的AES算法實(shí)現(xiàn)不夠靈活、難以升級(jí)維護(hù),而且開(kāi)發(fā)周期長(zhǎng)、開(kāi)發(fā)成本高,還存在擴(kuò)展性差等缺點(diǎn),并不適合于電子銀行或者電子金融這類應(yīng)用。最近的顯卡設(shè)備(如NVIDIA Geforce G210)提供了通用計(jì)算能力,支持整數(shù)和異或運(yùn)算,可實(shí)現(xiàn)密集的并行加密計(jì)算,可以滿足AES高強(qiáng)度的加解密需求,同時(shí)也提供了通用的軟件開(kāi)發(fā)環(huán)境——CUDA,降低開(kāi)發(fā)的難度?;贕PU的加速AES可彌補(bǔ)FPGA的不足,有低成本(不需額外增加設(shè)備)、易開(kāi)發(fā)(提供CUDA平臺(tái))等優(yōu)良特性,特別適用于服務(wù)器執(zhí)行高密度的加密任務(wù)。比如文獻(xiàn)[5]探索了利用GPU設(shè)備進(jìn)行并行的數(shù)據(jù)加密;文獻(xiàn)[6]在AVR微控制器和NVIDIA顯卡等多個(gè)平臺(tái)上實(shí)現(xiàn)了AES,實(shí)驗(yàn)顯示GPU并行的AES算法速度更快且代碼更短;文獻(xiàn)[7]在GPU上對(duì)小型內(nèi)存文件和大型硬盤(pán)文件均進(jìn)行了并行AES加速,得到了很好的加速比;文獻(xiàn)[8]使用CUDA、OpenMP、OpenCL對(duì)AES進(jìn)行并行加速,結(jié)果顯示CUDA獲得的加速性能最好,其次是OpenCL,最差的是OpenMP;文獻(xiàn)[9]在配置了一個(gè)多核處理器和一個(gè)GPU設(shè)備的計(jì)算平臺(tái)上,采用擴(kuò)展一次加密數(shù)據(jù)塊大小的方法對(duì)AES的CTR模式進(jìn)行加速處理。
基于CUDA的并行AES算法的優(yōu)化有如下進(jìn)展:文獻(xiàn)[10]提出適當(dāng)?shù)卦诓煌腉PU存儲(chǔ)器空間安排數(shù)據(jù)能帶來(lái)性能提升,在GeForce 9200MGS顯卡上能達(dá)到6.4 GB/s的吞吐量。文獻(xiàn)[11]采用字節(jié)分片(byte-sliced)的方法在ARM微處理器和Nvidia顯示卡對(duì)AES-128的加密及解密同時(shí)進(jìn)行優(yōu)化處理,并采用平臺(tái)相關(guān)的技術(shù)來(lái)提升AES-128的性能,在NVIDIA 8800 GTX顯卡上運(yùn)行的優(yōu)化AES-128能達(dá)到8.8 GB/S的吞吐量。文獻(xiàn)[12]將線程塊的明文和輪密鑰存儲(chǔ)在共享存儲(chǔ)區(qū)來(lái)優(yōu)化AES,但沒(méi)有與未優(yōu)化的AES進(jìn)行對(duì)比。文獻(xiàn)[13]主要對(duì)4種不同的并行粒度(分別按明文1字節(jié)/線程、4字節(jié)/線程、8字節(jié)/線程和16字節(jié)/線程的粒度)進(jìn)行了對(duì)比研究,得出在CUDA上16字節(jié)/線程的AES并行算法的性能最好,此外也從內(nèi)存分配方面對(duì)AES并行算法進(jìn)行了優(yōu)化,提升了2%~20%的性能。
國(guó)內(nèi)學(xué)者在AES算法優(yōu)化方面的進(jìn)展有:文獻(xiàn)[14]利用流水線、并行處理和可重構(gòu)技術(shù),提出了一種可重構(gòu)體系結(jié)構(gòu),實(shí)現(xiàn)的AES吞吐率在110 MHz工作頻率下分別可達(dá)到1.4 Gb/s。文獻(xiàn)[15]采用指令集架構(gòu)(ISA)擴(kuò)展的方式對(duì)AES密碼算法進(jìn)行優(yōu)化,提出一個(gè)高效AES專用指令處理器(AES-ASIP)模型,最終實(shí)現(xiàn)于FPGA中。對(duì)比ARM處理器指令集架構(gòu),AES-ASIP以增加少許硬件資源為代價(jià),提高了算法58.4%的執(zhí)行效率,并節(jié)省了47.4%的指令代碼存儲(chǔ)空間。
基于以上背景,為了提升加密性能,本文在統(tǒng)一計(jì)算設(shè)備架構(gòu)(Compute Unified Device Architecture, CUDA)平臺(tái)上實(shí)現(xiàn)AES 3個(gè)版本的GPU并行算法,并對(duì)不同的分塊大小進(jìn)行性能比較,得到分塊最優(yōu)值的經(jīng)驗(yàn)數(shù)據(jù)。在實(shí)際中能根據(jù)明文大小選擇最優(yōu)的分塊來(lái)執(zhí)行AES加密,分別在NVIDIA Geforce G210和Geforce GTX 460顯卡上運(yùn)行,對(duì)比分塊優(yōu)化的并行AES和未經(jīng)過(guò)優(yōu)化的并行AES的性能。
AES也叫Rijndael密碼,1998年由Daemen J和Rijmen V提出,2001被美國(guó)標(biāo)準(zhǔn)技術(shù)局(NIST)選為新一代對(duì)稱密碼標(biāo)準(zhǔn)。它是一個(gè)分組對(duì)稱密碼,以128位為一分組,加密、解密使用同一密鑰,密鑰可以使用128位、192位或256位密鑰。AES采用完全對(duì)稱的結(jié)構(gòu),主要有異或、查找表和移位等操作,能夠高效執(zhí)行。AES安全性強(qiáng),能對(duì)抗差分攻擊、線性攻擊和序列等已知明文攻擊。
AES通過(guò)對(duì)明文分組進(jìn)行若干輪與密鑰的運(yùn)算得到密文分組。每輪包括輪密鑰加、字節(jié)替換、行移位和列混淆4個(gè)階段。AES-128、AES-192和AES-256均可以128位為一分組,分別使用128位、192位或256位密鑰進(jìn)行多輪迭代轉(zhuǎn)換得到128位的密文。輪數(shù)由密鑰位數(shù)決定:128位密鑰需10輪,192位密鑰需12輪,而256位密鑰需14輪。除了首輪有一次輪密鑰相加和最后一輪少一次列混淆外,其余每輪處理方式相同。每輪加密的對(duì)象為一個(gè)16 Byte(128位)的分組,可以看做1個(gè)4×4大小的二維數(shù)組(表),稱為狀態(tài)(state)。每輪的加密處理可分為4個(gè)階段:輪密鑰加(AddRoundKey,對(duì)輪密鑰和狀態(tài)進(jìn)行異或運(yùn)算,進(jìn)行盲化),字節(jié)替換(SubBytes,使用S盒進(jìn)行非線性替換),行移位(ShiftRows,對(duì)每行分別進(jìn)行0, 1 Byte,2 Byte,3 Byte的循環(huán)移位,以在字內(nèi)混淆),列混淆(MixColumns,線性變換每列以混淆數(shù)據(jù))。整個(gè)AES-192的加密流程如圖1所示。
在圖1中,AES-192的192位(24 Byte)密鑰需擴(kuò)展為13個(gè)16 Byte(共208 Byte)的密鑰分組。第1輪之前的輪密鑰加時(shí)使用的是密鑰的本身值w[0..3],表示密鑰的前4個(gè)字。接下來(lái)的12輪過(guò)程都一樣,包括字節(jié)替換、行移位、列混淆和輪密鑰加4個(gè)階段,使用的輪密鑰是由原密鑰擴(kuò)展成的4字組,分別為w[4..7]~w[48..51]。特別注意第12輪少了一個(gè)列混淆的階段。密鑰擴(kuò)展只需執(zhí)行一次,故無(wú)需并行處理,設(shè)計(jì)為在CPU上串行執(zhí)行1次。這樣每輪訪問(wèn)密鑰都沒(méi)有數(shù)據(jù)依賴關(guān)系。置換表(S盒,Sbox)為16×16 Byte的二維表,值保持不變,每輪均按此表進(jìn)行替換。行移位對(duì)第i行循環(huán)左移i個(gè)字節(jié)(0≤i<4)。列混淆采用的是有限域GF(28)上的乘法,數(shù)據(jù)針對(duì)的是狀態(tài)??梢?jiàn)密鑰擴(kuò)展后,每個(gè)分組之間的加密過(guò)程都不存在數(shù)據(jù)依賴,能夠以并行方式設(shè)計(jì)和實(shí)現(xiàn)AES-192。AES-128和AES-256的不同體現(xiàn)在:加密的輪數(shù)分別為10輪和14輪,相應(yīng)的密鑰位數(shù)和密鑰擴(kuò)展有所不同,密鑰位數(shù)分別為128位和256位,密鑰擴(kuò)展后的長(zhǎng)度分別為44 Byte和60 Byte。AES-128與AES-256的密鑰擴(kuò)展方法相同,但與AES-256的密鑰擴(kuò)展算法有細(xì)微差別。由于密鑰擴(kuò)展串行執(zhí)行一次,因此并不影響AES不同位數(shù)版本的并行算法的設(shè)計(jì)和實(shí)現(xiàn)。
3.1 統(tǒng)一計(jì)算設(shè)備架構(gòu)模型
統(tǒng)一計(jì)算設(shè)備架構(gòu)(CUDA)由NVIDIA公司為它的G80、G92和GT200等圖形處理器提出的通用計(jì)算平臺(tái)。CUDA的編程模型能集成主機(jī)和GPU代碼到同一個(gè)C/C++源文件中,降低了開(kāi)發(fā)的難度。模型的結(jié)構(gòu)如圖2所示。
CUDA通過(guò)調(diào)用一個(gè)函數(shù),稱之為核函數(shù)(kernel)來(lái)支持并行計(jì)算。核函數(shù)由CPU調(diào)用而由設(shè)備(GPU)執(zhí)行。CUDA使用塊(block)和線程(thread)的概念來(lái)表示算法的并行性。塊是一組并發(fā)的線程單位,能獨(dú)立執(zhí)行。塊與塊之間的執(zhí)行順序可以是串行的,也可以是并行的。塊內(nèi)線程是并行執(zhí)行的。同一塊線程需要同步時(shí),調(diào)用__syncthreads來(lái)實(shí)現(xiàn)。同一塊線程通信時(shí),可以使用共享存儲(chǔ)器(shared memory),共享存儲(chǔ)器只對(duì)一個(gè)線程塊內(nèi)部的線程提供訪問(wèn)功能。一塊與另一塊線程通信則必須使用全局存儲(chǔ)器(global memory)。
圖2 CUDA模型
塊的大小由程序員指定,通過(guò)一個(gè)三維結(jié)構(gòu)(dim3)表示,分成了x,y,z3個(gè)坐標(biāo)。多個(gè)塊組成了一個(gè)更大的單位——網(wǎng)格(Grid)。網(wǎng)格的大小也由1個(gè)三維結(jié)構(gòu)(dim3)表示,分成了x,y,z3個(gè)坐標(biāo)(對(duì)NVIDIA Geforce G210z=1)。如圖2中,網(wǎng)格1 (Grid1)具有維度(x,y)=(3,2),第三維缺省為1,即z=1。網(wǎng)格1具有6個(gè)塊,分別編號(hào)為(0,0),(0, 1),…,(2,1)。這些塊都具有相同的大小,對(duì)塊(1, 1)而言,它的維度為(x,y)=(2,3),z缺省為1。所以具有 6線程,分別編號(hào)為(0,0),(0,1),…, (1,2)。
核函數(shù)通過(guò)kernel<<<blocks,threads>>>(parameters)修飾符的方式調(diào)用,即在普通C++函數(shù)調(diào)用格式中增加塊數(shù)量(blocks)和每塊的線程數(shù)(threads)這樣的參數(shù)。CUDA的執(zhí)行模型如圖3所示。線程由線程處理器執(zhí)行,但不是獨(dú)立執(zhí)行的,而是以線程塊的形式在流多處理器上執(zhí)行。如果多個(gè)線程塊運(yùn)行時(shí)所需的共享存儲(chǔ)器和寄存器等流多處理器的資源能滿足,那么一個(gè)流多處理器可安排多個(gè)線程塊。流多處理器執(zhí)行一個(gè)線程塊時(shí),按照線程束為單位執(zhí)行。線程束的大小為32線程,所以一個(gè)線程塊可以包含多個(gè)線程束。由線程束調(diào)度器選擇準(zhǔn)備好的線程束執(zhí)行。線程束按單指令多數(shù)據(jù)流(SIMD)的方式執(zhí)行,一個(gè)線程束中的所有線程同時(shí)執(zhí)行相同的指令。核函數(shù)以一個(gè)線程網(wǎng)格的形式啟動(dòng),對(duì)GPU設(shè)備而言一次只允許一個(gè)核函數(shù)執(zhí)行。
圖3 CUDA執(zhí)行模型
3.2 不同分塊的并行AES算法設(shè)計(jì)
在設(shè)計(jì)不同分塊的并行AES的核函數(shù)時(shí),采用了如下設(shè)計(jì)原則:
(1)每個(gè)塊中線程數(shù)量是32的整數(shù)倍;
(2)考慮共享存儲(chǔ)器容量的限制;
(3)考慮總塊數(shù)的限制。
原則(1)的理由是充分利用流多處理器的并行能力。因?yàn)榱鞫嗵幚砥饕跃€程束為單位執(zhí)行線塊。如果塊大小不是32的整數(shù)倍,則存在某些線程束少于32個(gè)線程(缺少的線程不執(zhí)行),這樣的線程束執(zhí)行時(shí)占用與一個(gè)完整的線程束同樣的計(jì)算時(shí)間,從而造成任務(wù)沒(méi)有滿載。
原則(2)與具體的算法執(zhí)行時(shí)所需的共享存儲(chǔ)器多少有關(guān)。對(duì)AES算法而言,共享存儲(chǔ)器容量能達(dá)到同塊所有線程執(zhí)行AES的需要。
原則(3)說(shuō)明CUDA程序運(yùn)行的總塊數(shù)受限于計(jì)算能力,總塊數(shù)必須小于或等于最大支持的塊數(shù)。
根據(jù)以上3條原則,AES算法的分塊方案設(shè)計(jì)為:分別對(duì)AES-128、AES-192、AES-256設(shè)計(jì)了分塊分別為64線程/塊、96線程/塊、128線程/塊、160線程/塊、192線程/塊、224線程/塊、256線程/塊、288線程/塊、320線程/塊、352線程/塊、384線程/塊等11種分塊方法。當(dāng)塊大于384線程/塊時(shí),由于Geforce G210的共享存儲(chǔ)器受到64 KB容量的限制,無(wú)法滿足執(zhí)行AES算法的需要。當(dāng)明文大于16 MB時(shí),若采用32線程/塊的分塊方法,會(huì)超過(guò)Geforce G210的總塊數(shù)的限制,因此也不能采用32線程/塊的分塊方法。其他的11種分塊方法包括了所有大于32且小于等于384的所有32倍的分塊大小。
3.3 基于分塊優(yōu)化的并行AES算法
由于AES算法本身的特點(diǎn)、明文塊的大小、線程束的調(diào)度、顯卡寄存器等硬件特點(diǎn)決定了不同的分塊大小對(duì)并行AES算法的性能有著顯著的影響,因此基于3.2節(jié)提出的分塊方法,對(duì)不同的明文和顯卡,按分塊方法執(zhí)行,對(duì)執(zhí)行結(jié)果進(jìn)行分析,得到特定的算法對(duì)特定的明文大小、最優(yōu)的分塊大小的經(jīng)驗(yàn)數(shù)據(jù)。以此經(jīng)驗(yàn)數(shù)據(jù)在實(shí)際處理時(shí),自適應(yīng)地根據(jù)明文大小和加密強(qiáng)度的要求,選擇合適的分塊來(lái)提升并行AES加密算法的性能。
基于分塊優(yōu)化并行AES算法,根據(jù)顯卡環(huán)境首先需要執(zhí)行一次以獲取該環(huán)境的優(yōu)化分塊信息。然后下次在該環(huán)境執(zhí)行并行AES加密處理時(shí),根據(jù)加密服務(wù)要求(明文大小和加密強(qiáng)度),依據(jù)經(jīng)驗(yàn)數(shù)據(jù)選擇優(yōu)化的分塊進(jìn)行加密,從而得到加密速度的提升。算法的具體情況見(jiàn)算法1。
算法1 基于分塊優(yōu)化的并行AES算法
功能 按照分塊對(duì)并行AES-128、AES-192和AES-256進(jìn)行優(yōu)化
輸入 明文plainText,密鑰key,AES算法版本version,分塊經(jīng)驗(yàn)數(shù)據(jù)experience,執(zhí)行次數(shù)times
輸出 密文cipherText
4.1 實(shí)驗(yàn)環(huán)境
實(shí)驗(yàn)設(shè)計(jì)在2個(gè)不同的實(shí)驗(yàn)平臺(tái)上進(jìn)行,實(shí)驗(yàn)平臺(tái)1和實(shí)驗(yàn)平臺(tái)2的主要環(huán)境參數(shù)如表1和表2所示。
表1 實(shí)驗(yàn)平臺(tái)1的主要環(huán)境參數(shù)
表2 實(shí)驗(yàn)平臺(tái)2的主要環(huán)境參數(shù)
4.2 實(shí)驗(yàn)結(jié)果
實(shí)驗(yàn)分別在實(shí)驗(yàn)平臺(tái)1和實(shí)驗(yàn)平臺(tái)2上進(jìn)行。每個(gè)平臺(tái)的實(shí)驗(yàn)分3組進(jìn)行,每組針對(duì)不同版本(AES-128,AES-192和AES-256)的AES算法。每組執(zhí)行11種不同分塊的AES并行程序各1次,分塊大小由3.2節(jié)的分塊方案確定。每種分塊方案的AES算法執(zhí)行時(shí)都對(duì)大小從64 Byte~16 MB的明文并行加密,統(tǒng)計(jì)每次實(shí)際運(yùn)行時(shí)間。平臺(tái)1的3組實(shí)驗(yàn)結(jié)果分別如圖4、圖5和圖6所示。由于明文從64 Byte~16 MB變化,跨度比較大,每組實(shí)驗(yàn)的結(jié)果分成64 Byte~16 KB明文和32 KB~16 MB這2個(gè)圖表分別展示。圖4(a)是不同分塊的AES-128在明文從64 Byte~16 KB區(qū)間的運(yùn)行時(shí)間對(duì)比,圖4(b)則是不同分塊的AES-128在明文從32 KB~16 MB區(qū)間的運(yùn)行時(shí)間對(duì)比。橫坐標(biāo)軸按64T/B,96T/B,128T/B,…, 384T/B從左到右依次排列,其中,64T/B表示的是分塊大小為64線程/塊,其余類似,。圖5和圖6的含義同圖4,對(duì)比了AES-192和AES-256在不同分塊情況下的運(yùn)行時(shí)間。平臺(tái)2的3組實(shí)驗(yàn)結(jié)果分別如圖7~圖9所示。這3個(gè)圖對(duì)比了平臺(tái)2上AES-192、AES-128和AES-256在不同分塊情況下的運(yùn)行時(shí)間。
圖4 平臺(tái)1上并行AES-128不同分塊的運(yùn)行時(shí)間
圖5 平臺(tái)1上并行AES-192不同分塊的運(yùn)行時(shí)間
圖6 平臺(tái)1上并行AES-256不同分塊的運(yùn)行時(shí)間
圖7 平臺(tái)2上并行AES-128不同分塊的運(yùn)行時(shí)間
圖8 平臺(tái)2上并行AES-192不同分塊的運(yùn)行時(shí)間
圖9 平臺(tái)2上并行AES-256不同分塊的運(yùn)行時(shí)間
4.3 分析和討論
考慮沒(méi)有優(yōu)化的AES算法,可能選擇64線程/塊~384線程/塊的任意一種分塊方法。因此,對(duì)沒(méi)有優(yōu)化的AES算法采用的是對(duì)這11種分塊方案的平均來(lái)度量。經(jīng)過(guò)數(shù)據(jù)分析,優(yōu)化的AES算法帶來(lái)的性能情況如表3所示。
表3 優(yōu)化的AES并行算法與未優(yōu)化的AES并行算法性能比較
在平臺(tái)1上,優(yōu)化的AES算法的加密速度提升比例按照AES-128、AES-256和AES-192增加,并且AES-192和AES-256的提升比例非常接近。這是由于AES-192和 AES-256比 AES-128算法更復(fù)雜。優(yōu)化的AES-256或AES-192獲得的加密速度的峰值提升比例也很接近,均比優(yōu)化的AES-128多提升了30%以上。
在平臺(tái)2上,優(yōu)化的AES算法的加密速度提升比例按照AES-128、AES-192和AES-256增加,并且AES-192和AES-256的提升比例非常接近。這是由于AES-192和 AES-256比 AES-128算法更復(fù)雜。優(yōu)化的AES-256或AES-192獲得的加密速度的峰值提升比例也很接近,比優(yōu)化的AES-128提升了2%左右。
從表3中可以看出,采用優(yōu)化的AES算法在不同的平臺(tái)上均有加密平均速度提升和更高比例的加密峰值速度提升。由于平臺(tái)2采用的顯卡為Nvidia Geforce GTX460,相比平臺(tái)1采用的顯卡 Nvidia Geforce G210性能有很大提升,因此平臺(tái)2的并行AES加密速度平均和峰值提升比例均優(yōu)于平臺(tái)1的并行AES算法。
為了分析不同分塊在不同算法和明文大小下的性能提升情況,給出了優(yōu)化的AES算法提升比例,如圖10所示。其中,AES-128 pf1表示在平臺(tái)1上優(yōu)化的AES-128算法相比未優(yōu)化的AES-128算法的性能提升比例,其余類似;而AES-128 pf2表示在平臺(tái)2上優(yōu)化的AES-128算法相比未優(yōu)化的AES-128算法的性能提高比例,其余類似。優(yōu)化的AES算法性能提升較大的區(qū)域主要有 2塊,其中,1塊是64 Byte~256 Byte的明文大小區(qū)間,另 1塊是16 KB~256 KB之間。16 KB~256 KB明文大小的加速性能更是優(yōu)于64 Byte~256 Byte這塊的加密性能,且具有很好的現(xiàn)實(shí)意義。因?yàn)橐淮蔚湫偷闹缓琀TML內(nèi)容的SSL會(huì)話的傳輸數(shù)據(jù)量約為35 KB,而既有文本內(nèi)容又有圖像內(nèi)容的SSL會(huì)話的數(shù)據(jù)量約為150 KB[16]。這說(shuō)明優(yōu)化的AES算法使用當(dāng)前普通的或者高端的顯卡對(duì)SSL加密提供了較好的性能提升。
圖10 AES優(yōu)化算法性能提升比例
為了應(yīng)對(duì)服務(wù)器上加密瓶頸問(wèn)題,采用了基于CUDA的并行AES算法來(lái)提升AES的加密性能。由于AES算法特征(AES版本和需要的計(jì)算資源)、顯卡的硬件特征(寄存器數(shù)量,分塊總數(shù)和線程束的調(diào)度等)和任務(wù)情況(明文大小)的原因,造成了不同的分塊大小對(duì)并行AES性能有著較大的影響。因此,本文依據(jù)分塊大小的經(jīng)驗(yàn)數(shù)據(jù),設(shè)計(jì)和提出了基于分塊優(yōu)化的 AES并行算法,包括 AES-128、AES-192和 AES-256 3個(gè)版本。通過(guò)在平臺(tái) 1 (Nvidia Geforce G210)和平臺(tái) 2(Nvidia Geforce GTX460)2個(gè)平臺(tái)的實(shí)驗(yàn)運(yùn)行,結(jié)果顯示優(yōu)化的AES算法帶來(lái)了很大的性能提升。在平臺(tái)1上,優(yōu)化的 AES-128、AES-192和 AES-256分別達(dá)到了5.28%、14.55%和12.53%的加密速度平均提升比例和10.02%、42.49%和44.22%的加密速度峰值提升比例。在平臺(tái)2上,優(yōu)化的AES-128、AES-192和AES-256分別達(dá)到了12.48%、15.40%和15.84%的加密速度平均提升比例和 45.49%、47.87%和47.15%的加密速度峰值提升比例。優(yōu)化的AES算法性能提升較大的明文區(qū)域主要有2塊,其中一塊是64 Byte~256 Byte的明文大小區(qū)間,另外一塊在16 KB~256 KB之間,且16 KB~256 KB明文大小的加速性能優(yōu)于64 Byte~256 Byte這塊的加密性能,非常適合于對(duì)SSL加密的性能提升。
[1] Good T,Benaissa M.AES on FPGA from the Fastest to the Smallest[C]//Proc.of CHES'05.Berlin,Germany: Springer,2005:427-440.
[2] Sivakumar C,Velmurugan A.HighSpeed VLSI Design CCMP AES Cipher for WLAN(IEEE 802.11 i)[C]// Proc.of International Conference on Signal Processing, Communications and Networking.[S.l.]:IEEE Press, 2007:398-403.
[3] Thulasimani L,Madheswaran M.A Single Chip Design and Implementation of AES-128/192/256 Encryption Alogorithms[J].International Journal of Engineering Science and Technology,2010,2(5):1052-1059.
[4] Wang Yi,Ha Yajun.FPGA-based 40.9-Gbits/s Masked AES with Area Optimization for Storage Area Network [J].IEEE Transactions on Circuits and Systems,2013, 60(1):36-40.
[5] Cook D,KeromytisA D.Cryptographics:Exploiting Graphics Cards for Security[M].[S.l.]:Springer,2006.
[6] Bos J W,Osvik D A,Stefan D.Fast Implementations of AES on Various Platforms[C]//Proc.of IACR'09. Taormina,Italy:[s.n.],2009:501-515.
[7] Tomoiagaa R D,Stratulat M. Accelerating Solution Proposal of AES Using a Graphic Processor[J].Advances in Electrical and Computer Engineering,2011,11(4): 99-104.
[8] Duta C L,Michiu G,StoicaS,etal.Accelerating Encryption Alogorithms Using Parallelism[C]//Proc.of IEEE International Conference on Control Systems and Computer Science.[S.l.]:IEEE Press,2013:549-554.
[9] Nhat-Phuong T,Myungho L E E,Sugwon H,et al.High Throughput Parallelization of AES-CTR Alogorithm[J]. IEICE Transactions on Information and Systems,2013, 96(8):1685-1695.
[10] Mei Chonglei,Jiang Hai,Jenness J.CUDA-based AES Parallelization with Fine-tuned GPU Memory Utilization [C]//Proc.of IPDPSW'10.[S.l.]:IEEE Press,2010:1-7.
[11] Osvik D A,Bos J W,Stefan D,et al.FastSoftware AES Encryption[C]//Proc.of the 17th International Workshop on Fast Software Encryption.Berlin,Germany:Springer, 2010:75-93.
[12] Le Deguang,Chang Jinyi,Gou Xingdou,et al.Parallel AES Alogorithm for Fast Data Encryption on GPU [C]//Proc.of the 2nd International Conference on Computer Engineering and Technology.[S.l.]:IEEE Press,2010:61-66.
[13] Iwai K,Nishikaw N,Kurokawa T.Acceleration of AES encryption on CUDA GPU[J].International Journal of Networking and Computing,2012,2(1):131-145.
[14] 高娜娜,李占才,王 沁.一種可重構(gòu)體系結(jié)構(gòu)用于高速實(shí)現(xiàn)DES,3DES和AES[J].電子學(xué)報(bào),2006,34 (8):1386-1390.
[15] 夏 輝,賈智平,張 峰,等.AES專用指令處理器的研究與實(shí)現(xiàn)[J].計(jì)算機(jī)研究與發(fā)展,2011,48(8):1554-1562.
[16] Levering R,Cutler M.ThePortrait of a Common HTML Web Page[C]//Proc.of ACM Symposium on Document Engineering.New York,USA:ACM Press,2006:198-204.
編輯 任吉慧
Optimization of AES Parallel Alogorithm Based on CUDA
FEI Xiong-wei1,2,LI Ken-li2,YANG Wang-dong1,2
(1.Department of Information Science and Engineering,Hunan City University,Yiyang 413000,China;
2.College of Computer Science and Electronic Engineering,Hunan University,Changsha 410008,China)
In order to enhance the efficiency of Advanced Encryption Standard(AES)and make use of general computing ability of Graphics Processing Unit(GPU),all the three versions of GPU parallel AES,namely 128 bit version,192 bit version and 256 bit version,are implemented on Compute Unified Device Architecture(CUDA).Then,it proposes optimization alogorithms of parallel AES with 3 versions.These alogorithms first consider threads amount in a block,shared memory size and total blocks,then use the experience data of optimal value of block size to guide AES alogorithm's optimal block on GPU.Experimental results show that compared with unoptimized parral AES,these alogorithms can obtain encryption mean speedup by 5.28%,14.55%and 12.53%respectively on Nvidia Geforce G210 graphics card,while by 12.48%,15.40% and 15.84% on Nvidia Geforce GTX460 graphics card.In addition,these alogorithms are better at improving encrypting of Secure Socket Layer(SSL).
block;experiential data;parallel alogorithm;optimization;Advanced Encryption Standard(AES);Compute Unified Device Architecture(CUDA)
1000-3428(2014)09-0006-07
A
TP301.6
10.3969/j.issn.1000-3428.2014.09.002
國(guó)家自然科學(xué)基金資助重點(diǎn)項(xiàng)目(61133005);國(guó)家自然科學(xué)基金資助項(xiàng)目(90715029,61070057,60603053)。
費(fèi)雄偉(1980-),男,講師、碩士,主研方向:網(wǎng)絡(luò)與信息安全,并行計(jì)算;李肯立,教授、博士生導(dǎo)師;陽(yáng)王東,教授、博士研究生。
2014-01-16
2014-02-21E-mail:25616235@qq.com