費雄偉,李肯立,陽王東
(1.湖南城市學院信息科學與工程學院,湖南益陽413000;2.湖南大學信息科學與工程學院,長沙410008)
基于CUDA的AES并行算法優化
費雄偉1,2,李肯立2,陽王東1,2
(1.湖南城市學院信息科學與工程學院,湖南益陽413000;2.湖南大學信息科學與工程學院,長沙410008)
為提升高級加密標準(AES)的加密性能,利用顯卡的通用計算能力,在統一計算設備架構(CUDA)平臺上實現AES的128位、192位和256位3個版本的GPU并行算法,并提出優化的AES并行算法。在考慮塊內線程數量、共享存儲器容量和總塊數的基礎上,根據分塊最優值的經驗數據指導AES算法在GPU上的最優分塊。實驗結果表明,與未優化的AES并行算法相比,該算法的3個版本在Nvidia Geforce G210顯卡上的加密速度分別提高5.28%,14.55%和12.53%,而在 Nvidia Geforce GTX460顯卡上的加密速度分別提高 12.48%,15.40%和15.84%,且能更好地對SSL數據進行加密。
分塊;經驗數據;并行算法;優化;高級加密標準;統一計算設備架構
隨著電子商務和電子金融的用戶群體擴大和對安全的要求更高,加密及其加密性能成為亟待處理的問題。對加密算法,首先要考慮的是它的安全性能,但在實際應用時還需考慮它的性能和成本。高級加密標準(Advanced Encryption Standard,AES)算法在大量需要安全通信的應用中擔當加密、解密任務,很好地協調了安全、高效和簡明等特性。但隨著網絡用戶的迅速增加,對大規模加密任務應用服務器(如電子商務或電子金融)的加密性能提出了更高的要求。為了提升AES的加密性能,最近許多基于FPGA和基于GPU加速的AES算法被提出。基于FPGA的加速AES主要應用于無線和藍牙通信環境的加速加密,需要考慮能耗和大小等問題。比如文獻[1]在Xilinx Spartan-III和Xilinx Spartan-II這2類設備上對AES進行了速度和面積大小方面的對比研究;文獻[2]提出了適應無線局域網環境的AES高速FPGA實現;文獻[3]在FPGA上實現了AES-128、AES-192和AES-256,并進行優化,達到了節能的要求,適合于移動的設備;文獻[4]提出了高吞吐量的AES引擎,采用FPGA實現了面積優化36.2%的AES算法。但采用FPGA這類專用硬件設計的AES算法實現不夠靈活、難以升級維護,而且開發周期長、開發成本高,還存在擴展性差等缺點,并不適合于電子銀行或者電子金融這類應用。最近的顯卡設備(如NVIDIA Geforce G210)提供了通用計算能力,支持整數和異或運算,可實現密集的并行加密計算,可以滿足AES高強度的加解密需求,同時也提供了通用的軟件開發環境——CUDA,降低開發的難度?;贕PU的加速AES可彌補FPGA的不足,有低成本(不需額外增加設備)、易開發(提供CUDA平臺)等優良特性,特別適用于服務器執行高密度的加密任務。比如文獻[5]探索了利用GPU設備進行并行的數據加密;文獻[6]在AVR微控制器和NVIDIA顯卡等多個平臺上實現了AES,實驗顯示GPU并行的AES算法速度更快且代碼更短;文獻[7]在GPU上對小型內存文件和大型硬盤文件均進行了并行AES加速,得到了很好的加速比;文獻[8]使用CUDA、OpenMP、OpenCL對AES進行并行加速,結果顯示CUDA獲得的加速性能最好,其次是OpenCL,最差的是OpenMP;文獻[9]在配置了一個多核處理器和一個GPU設備的計算平臺上,采用擴展一次加密數據塊大小的方法對AES的CTR模式進行加速處理。
基于CUDA的并行AES算法的優化有如下進展:文獻[10]提出適當地在不同的GPU存儲器空間安排數據能帶來性能提升,在GeForce 9200MGS顯卡上能達到6.4 GB/s的吞吐量。文獻[11]采用字節分片(byte-sliced)的方法在ARM微處理器和Nvidia顯示卡對AES-128的加密及解密同時進行優化處理,并采用平臺相關的技術來提升AES-128的性能,在NVIDIA 8800 GTX顯卡上運行的優化AES-128能達到8.8 GB/S的吞吐量。文獻[12]將線程塊的明文和輪密鑰存儲在共享存儲區來優化AES,但沒有與未優化的AES進行對比。文獻[13]主要對4種不同的并行粒度(分別按明文1字節/線程、4字節/線程、8字節/線程和16字節/線程的粒度)進行了對比研究,得出在CUDA上16字節/線程的AES并行算法的性能最好,此外也從內存分配方面對AES并行算法進行了優化,提升了2%~20%的性能。
國內學者在AES算法優化方面的進展有:文獻[14]利用流水線、并行處理和可重構技術,提出了一種可重構體系結構,實現的AES吞吐率在110 MHz工作頻率下分別可達到1.4 Gb/s。文獻[15]采用指令集架構(ISA)擴展的方式對AES密碼算法進行優化,提出一個高效AES專用指令處理器(AES-ASIP)模型,最終實現于FPGA中。對比ARM處理器指令集架構,AES-ASIP以增加少許硬件資源為代價,提高了算法58.4%的執行效率,并節省了47.4%的指令代碼存儲空間。
基于以上背景,為了提升加密性能,本文在統一計算設備架構(Compute Unified Device Architecture, CUDA)平臺上實現AES 3個版本的GPU并行算法,并對不同的分塊大小進行性能比較,得到分塊最優值的經驗數據。在實際中能根據明文大小選擇最優的分塊來執行AES加密,分別在NVIDIA Geforce G210和Geforce GTX 460顯卡上運行,對比分塊優化的并行AES和未經過優化的并行AES的性能。
AES也叫Rijndael密碼,1998年由Daemen J和Rijmen V提出,2001被美國標準技術局(NIST)選為新一代對稱密碼標準。它是一個分組對稱密碼,以128位為一分組,加密、解密使用同一密鑰,密鑰可以使用128位、192位或256位密鑰。AES采用完全對稱的結構,主要有異或、查找表和移位等操作,能夠高效執行。AES安全性強,能對抗差分攻擊、線性攻擊和序列等已知明文攻擊。
AES通過對明文分組進行若干輪與密鑰的運算得到密文分組。每輪包括輪密鑰加、字節替換、行移位和列混淆4個階段。AES-128、AES-192和AES-256均可以128位為一分組,分別使用128位、192位或256位密鑰進行多輪迭代轉換得到128位的密文。輪數由密鑰位數決定:128位密鑰需10輪,192位密鑰需12輪,而256位密鑰需14輪。除了首輪有一次輪密鑰相加和最后一輪少一次列混淆外,其余每輪處理方式相同。每輪加密的對象為一個16 Byte(128位)的分組,可以看做1個4×4大小的二維數組(表),稱為狀態(state)。每輪的加密處理可分為4個階段:輪密鑰加(AddRoundKey,對輪密鑰和狀態進行異或運算,進行盲化),字節替換(SubBytes,使用S盒進行非線性替換),行移位(ShiftRows,對每行分別進行0, 1 Byte,2 Byte,3 Byte的循環移位,以在字內混淆),列混淆(MixColumns,線性變換每列以混淆數據)。整個AES-192的加密流程如圖1所示。
在圖1中,AES-192的192位(24 Byte)密鑰需擴展為13個16 Byte(共208 Byte)的密鑰分組。第1輪之前的輪密鑰加時使用的是密鑰的本身值w[0..3],表示密鑰的前4個字。接下來的12輪過程都一樣,包括字節替換、行移位、列混淆和輪密鑰加4個階段,使用的輪密鑰是由原密鑰擴展成的4字組,分別為w[4..7]~w[48..51]。特別注意第12輪少了一個列混淆的階段。密鑰擴展只需執行一次,故無需并行處理,設計為在CPU上串行執行1次。這樣每輪訪問密鑰都沒有數據依賴關系。置換表(S盒,Sbox)為16×16 Byte的二維表,值保持不變,每輪均按此表進行替換。行移位對第i行循環左移i個字節(0≤i<4)。列混淆采用的是有限域GF(28)上的乘法,數據針對的是狀態??梢娒荑€擴展后,每個分組之間的加密過程都不存在數據依賴,能夠以并行方式設計和實現AES-192。AES-128和AES-256的不同體現在:加密的輪數分別為10輪和14輪,相應的密鑰位數和密鑰擴展有所不同,密鑰位數分別為128位和256位,密鑰擴展后的長度分別為44 Byte和60 Byte。AES-128與AES-256的密鑰擴展方法相同,但與AES-256的密鑰擴展算法有細微差別。由于密鑰擴展串行執行一次,因此并不影響AES不同位數版本的并行算法的設計和實現。
3.1 統一計算設備架構模型
統一計算設備架構(CUDA)由NVIDIA公司為它的G80、G92和GT200等圖形處理器提出的通用計算平臺。CUDA的編程模型能集成主機和GPU代碼到同一個C/C++源文件中,降低了開發的難度。模型的結構如圖2所示。
CUDA通過調用一個函數,稱之為核函數(kernel)來支持并行計算。核函數由CPU調用而由設備(GPU)執行。CUDA使用塊(block)和線程(thread)的概念來表示算法的并行性。塊是一組并發的線程單位,能獨立執行。塊與塊之間的執行順序可以是串行的,也可以是并行的。塊內線程是并行執行的。同一塊線程需要同步時,調用__syncthreads來實現。同一塊線程通信時,可以使用共享存儲器(shared memory),共享存儲器只對一個線程塊內部的線程提供訪問功能。一塊與另一塊線程通信則必須使用全局存儲器(global memory)。

圖2 CUDA模型
塊的大小由程序員指定,通過一個三維結構(dim3)表示,分成了x,y,z3個坐標。多個塊組成了一個更大的單位——網格(Grid)。網格的大小也由1個三維結構(dim3)表示,分成了x,y,z3個坐標(對NVIDIA Geforce G210z=1)。如圖2中,網格1 (Grid1)具有維度(x,y)=(3,2),第三維缺省為1,即z=1。網格1具有6個塊,分別編號為(0,0),(0, 1),…,(2,1)。這些塊都具有相同的大小,對塊(1, 1)而言,它的維度為(x,y)=(2,3),z缺省為1。所以具有 6線程,分別編號為(0,0),(0,1),…, (1,2)。
核函數通過kernel<<<blocks,threads>>>(parameters)修飾符的方式調用,即在普通C++函數調用格式中增加塊數量(blocks)和每塊的線程數(threads)這樣的參數。CUDA的執行模型如圖3所示。線程由線程處理器執行,但不是獨立執行的,而是以線程塊的形式在流多處理器上執行。如果多個線程塊運行時所需的共享存儲器和寄存器等流多處理器的資源能滿足,那么一個流多處理器可安排多個線程塊。流多處理器執行一個線程塊時,按照線程束為單位執行。線程束的大小為32線程,所以一個線程塊可以包含多個線程束。由線程束調度器選擇準備好的線程束執行。線程束按單指令多數據流(SIMD)的方式執行,一個線程束中的所有線程同時執行相同的指令。核函數以一個線程網格的形式啟動,對GPU設備而言一次只允許一個核函數執行。

圖3 CUDA執行模型
3.2 不同分塊的并行AES算法設計
在設計不同分塊的并行AES的核函數時,采用了如下設計原則:
(1)每個塊中線程數量是32的整數倍;
(2)考慮共享存儲器容量的限制;
(3)考慮總塊數的限制。
原則(1)的理由是充分利用流多處理器的并行能力。因為流多處理器以線程束為單位執行線塊。如果塊大小不是32的整數倍,則存在某些線程束少于32個線程(缺少的線程不執行),這樣的線程束執行時占用與一個完整的線程束同樣的計算時間,從而造成任務沒有滿載。
原則(2)與具體的算法執行時所需的共享存儲器多少有關。對AES算法而言,共享存儲器容量能達到同塊所有線程執行AES的需要。
原則(3)說明CUDA程序運行的總塊數受限于計算能力,總塊數必須小于或等于最大支持的塊數。
根據以上3條原則,AES算法的分塊方案設計為:分別對AES-128、AES-192、AES-256設計了分塊分別為64線程/塊、96線程/塊、128線程/塊、160線程/塊、192線程/塊、224線程/塊、256線程/塊、288線程/塊、320線程/塊、352線程/塊、384線程/塊等11種分塊方法。當塊大于384線程/塊時,由于Geforce G210的共享存儲器受到64 KB容量的限制,無法滿足執行AES算法的需要。當明文大于16 MB時,若采用32線程/塊的分塊方法,會超過Geforce G210的總塊數的限制,因此也不能采用32線程/塊的分塊方法。其他的11種分塊方法包括了所有大于32且小于等于384的所有32倍的分塊大小。
3.3 基于分塊優化的并行AES算法
由于AES算法本身的特點、明文塊的大小、線程束的調度、顯卡寄存器等硬件特點決定了不同的分塊大小對并行AES算法的性能有著顯著的影響,因此基于3.2節提出的分塊方法,對不同的明文和顯卡,按分塊方法執行,對執行結果進行分析,得到特定的算法對特定的明文大小、最優的分塊大小的經驗數據。以此經驗數據在實際處理時,自適應地根據明文大小和加密強度的要求,選擇合適的分塊來提升并行AES加密算法的性能。
基于分塊優化并行AES算法,根據顯卡環境首先需要執行一次以獲取該環境的優化分塊信息。然后下次在該環境執行并行AES加密處理時,根據加密服務要求(明文大小和加密強度),依據經驗數據選擇優化的分塊進行加密,從而得到加密速度的提升。算法的具體情況見算法1。
算法1 基于分塊優化的并行AES算法
功能 按照分塊對并行AES-128、AES-192和AES-256進行優化
輸入 明文plainText,密鑰key,AES算法版本version,分塊經驗數據experience,執行次數times
輸出 密文cipherText

4.1 實驗環境
實驗設計在2個不同的實驗平臺上進行,實驗平臺1和實驗平臺2的主要環境參數如表1和表2所示。

表1 實驗平臺1的主要環境參數

表2 實驗平臺2的主要環境參數
4.2 實驗結果
實驗分別在實驗平臺1和實驗平臺2上進行。每個平臺的實驗分3組進行,每組針對不同版本(AES-128,AES-192和AES-256)的AES算法。每組執行11種不同分塊的AES并行程序各1次,分塊大小由3.2節的分塊方案確定。每種分塊方案的AES算法執行時都對大小從64 Byte~16 MB的明文并行加密,統計每次實際運行時間。平臺1的3組實驗結果分別如圖4、圖5和圖6所示。由于明文從64 Byte~16 MB變化,跨度比較大,每組實驗的結果分成64 Byte~16 KB明文和32 KB~16 MB這2個圖表分別展示。圖4(a)是不同分塊的AES-128在明文從64 Byte~16 KB區間的運行時間對比,圖4(b)則是不同分塊的AES-128在明文從32 KB~16 MB區間的運行時間對比。橫坐標軸按64T/B,96T/B,128T/B,…, 384T/B從左到右依次排列,其中,64T/B表示的是分塊大小為64線程/塊,其余類似,。圖5和圖6的含義同圖4,對比了AES-192和AES-256在不同分塊情況下的運行時間。平臺2的3組實驗結果分別如圖7~圖9所示。這3個圖對比了平臺2上AES-192、AES-128和AES-256在不同分塊情況下的運行時間。

圖4 平臺1上并行AES-128不同分塊的運行時間

圖5 平臺1上并行AES-192不同分塊的運行時間

圖6 平臺1上并行AES-256不同分塊的運行時間

圖7 平臺2上并行AES-128不同分塊的運行時間

圖8 平臺2上并行AES-192不同分塊的運行時間

圖9 平臺2上并行AES-256不同分塊的運行時間
4.3 分析和討論
考慮沒有優化的AES算法,可能選擇64線程/塊~384線程/塊的任意一種分塊方法。因此,對沒有優化的AES算法采用的是對這11種分塊方案的平均來度量。經過數據分析,優化的AES算法帶來的性能情況如表3所示。

表3 優化的AES并行算法與未優化的AES并行算法性能比較
在平臺1上,優化的AES算法的加密速度提升比例按照AES-128、AES-256和AES-192增加,并且AES-192和AES-256的提升比例非常接近。這是由于AES-192和 AES-256比 AES-128算法更復雜。優化的AES-256或AES-192獲得的加密速度的峰值提升比例也很接近,均比優化的AES-128多提升了30%以上。
在平臺2上,優化的AES算法的加密速度提升比例按照AES-128、AES-192和AES-256增加,并且AES-192和AES-256的提升比例非常接近。這是由于AES-192和 AES-256比 AES-128算法更復雜。優化的AES-256或AES-192獲得的加密速度的峰值提升比例也很接近,比優化的AES-128提升了2%左右。
從表3中可以看出,采用優化的AES算法在不同的平臺上均有加密平均速度提升和更高比例的加密峰值速度提升。由于平臺2采用的顯卡為Nvidia Geforce GTX460,相比平臺1采用的顯卡 Nvidia Geforce G210性能有很大提升,因此平臺2的并行AES加密速度平均和峰值提升比例均優于平臺1的并行AES算法。
為了分析不同分塊在不同算法和明文大小下的性能提升情況,給出了優化的AES算法提升比例,如圖10所示。其中,AES-128 pf1表示在平臺1上優化的AES-128算法相比未優化的AES-128算法的性能提升比例,其余類似;而AES-128 pf2表示在平臺2上優化的AES-128算法相比未優化的AES-128算法的性能提高比例,其余類似。優化的AES算法性能提升較大的區域主要有 2塊,其中,1塊是64 Byte~256 Byte的明文大小區間,另 1塊是16 KB~256 KB之間。16 KB~256 KB明文大小的加速性能更是優于64 Byte~256 Byte這塊的加密性能,且具有很好的現實意義。因為一次典型的只含HTML內容的SSL會話的傳輸數據量約為35 KB,而既有文本內容又有圖像內容的SSL會話的數據量約為150 KB[16]。這說明優化的AES算法使用當前普通的或者高端的顯卡對SSL加密提供了較好的性能提升。

圖10 AES優化算法性能提升比例
為了應對服務器上加密瓶頸問題,采用了基于CUDA的并行AES算法來提升AES的加密性能。由于AES算法特征(AES版本和需要的計算資源)、顯卡的硬件特征(寄存器數量,分塊總數和線程束的調度等)和任務情況(明文大小)的原因,造成了不同的分塊大小對并行AES性能有著較大的影響。因此,本文依據分塊大小的經驗數據,設計和提出了基于分塊優化的 AES并行算法,包括 AES-128、AES-192和 AES-256 3個版本。通過在平臺 1 (Nvidia Geforce G210)和平臺 2(Nvidia Geforce GTX460)2個平臺的實驗運行,結果顯示優化的AES算法帶來了很大的性能提升。在平臺1上,優化的 AES-128、AES-192和 AES-256分別達到了5.28%、14.55%和12.53%的加密速度平均提升比例和10.02%、42.49%和44.22%的加密速度峰值提升比例。在平臺2上,優化的AES-128、AES-192和AES-256分別達到了12.48%、15.40%和15.84%的加密速度平均提升比例和 45.49%、47.87%和47.15%的加密速度峰值提升比例。優化的AES算法性能提升較大的明文區域主要有2塊,其中一塊是64 Byte~256 Byte的明文大小區間,另外一塊在16 KB~256 KB之間,且16 KB~256 KB明文大小的加速性能優于64 Byte~256 Byte這塊的加密性能,非常適合于對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] 高娜娜,李占才,王 沁.一種可重構體系結構用于高速實現DES,3DES和AES[J].電子學報,2006,34 (8):1386-1390.
[15] 夏 輝,賈智平,張 峰,等.AES專用指令處理器的研究與實現[J].計算機研究與發展,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
國家自然科學基金資助重點項目(61133005);國家自然科學基金資助項目(90715029,61070057,60603053)。
費雄偉(1980-),男,講師、碩士,主研方向:網絡與信息安全,并行計算;李肯立,教授、博士生導師;陽王東,教授、博士研究生。
2014-01-16
2014-02-21E-mail:25616235@qq.com