基于CUDA的SM4加密算法高速實現(xiàn)

王德民 陳 達

（桂林電子科技大學信息與通信學院 廣西桂林 541004）

基于CUDA的SM4加密算法高速實現(xiàn)

王德民 陳 達

（桂林電子科技大學信息與通信學院 廣西桂林 541004）

傳統(tǒng)的SM4加密運算是在CPU上實現(xiàn)的，為了提高加密速度以處理大規(guī)模的加密運算，根據(jù)分組密碼SM4的結(jié)構(gòu)和特點，實現(xiàn)了一個利用GPU的通用計算能力，在統(tǒng)一計算設(shè)備架構(gòu)（CUDA）平臺上運行的SM4并行算法。通過兩個實驗平臺，對SM4并行算法（基于CUDA）和串行算法（基于CPU）的性能進行了對比和實驗驗證。結(jié)果表明，該并行SM4算法在平臺1上最高能達到40.6倍的加速比和85.4%的加速效率，在平臺2上最高能達到64.7倍的加速比和49.5%的加速效率。

SM4 CUDA 并行計算

1 引言

隨著無線局域網(wǎng)（Wireless Local Area Network, WLAN）的日益普及，其傳輸信息的安全性面臨嚴峻的挑戰(zhàn)。SM4算法[1]，是由國家商用密碼管理辦公室于2006年1月公布的用于無線局域網(wǎng)產(chǎn)品的分組對稱密碼算法，該算法在大量需要安全通信的應用中擔當加密、解密任務，具有安全、高效和簡明等特性。但是SM4運算復雜度極高，過程冗長，非常不適應于網(wǎng)絡(luò)傳輸和視頻加密等對速度要求較高的場合。業(yè)界許多基于FPGA設(shè)計的被提出[2,3]。但采用FPGA這類專用硬件設(shè)計的SM4算法實現(xiàn)不夠靈活，存在如升級維護難、開發(fā)周期長，成本高等缺點。因此，開發(fā)快速運行的并行SM4加解密算法，從而實現(xiàn)高速數(shù)據(jù)流的實時加解密功能，是亟待解決的重要問題。

基于以上背景，為了提升加解密性能，本文在分析SM4算法的基礎(chǔ)上，提出基于GPU并行計算的SM4加解密基本架構(gòu)。并在不同的實驗品平臺上對不同的明文數(shù)據(jù)分塊大小進行性能比較，得到不同分塊下的加速比與加速效率。

2 SM4算法描述

SM4 是一種Feistel 結(jié)構(gòu)的分組密碼算法，其分組長度和密鑰長度均為128bit。加解密算法與密鑰擴張算法都采用32輪非線性迭代結(jié)構(gòu)[4]。解密算法與加密算法的結(jié)構(gòu)相同，只是輪密鑰的使用順序相反，即解密算法使用的輪密鑰是加密算法使用的輪密鑰的逆序。

2.1 加解密算法描述

定義 為e比特的向量集，<<

現(xiàn)對SM4解密算法進行具體介紹。128bit的明文或密文輸入經(jīng)初始變換分成4個字節(jié)，與輪密鑰經(jīng)過輪函數(shù)的運算經(jīng)過32 輪的迭代完成加解密運算。，第二輪使用

SMS解密過程與加密變換結(jié)構(gòu)相同，不同的僅是輪密鑰的使用順序，第一輪使用，依此類推。

2.2 密鑰擴展算法

該算法中加密算法的輪密鑰由加密密鑰通過密鑰擴展算法生成。

3 CUDA模型及并行算法設(shè)計

3.1 CUDA模型

CUDA（Compute Unified Device Architecture）即統(tǒng)一計算架構(gòu)[6]，由Nvidia公司提出。CUDA通過調(diào)用一個核函數(shù)（kernel），來支持并行計算。核函數(shù)由CPU調(diào)用而由設(shè)備（GPU）執(zhí)行。CUDA使用塊（block）和線程（thread）的概念來表示算法的并行性[7]。核心的執(zhí)行在調(diào)用時通過kernal<<>>（parameters）修飾符提供塊數(shù)量和每塊的線程數(shù)，其中blocks表示塊大小，而threads表示每塊線程數(shù)，parameters表示函數(shù)的參數(shù)列表[8]。

3.2 并行算法設(shè)計

圖1 SM4并行算法整體框架

SM4算法的分組加密結(jié)構(gòu)適合進行并行處理，因為明文分組之間不存在數(shù)據(jù)依賴，加密時按明文分組進行并行加密，能充分利用 GPU 多核的處理能力。SM4并行算法實現(xiàn)如圖1所示。

輪密鑰只需要產(chǎn)生一次，可在整個加密過程中循環(huán)使用。首先，在 CPU執(zhí)行一次密鑰擴展，生成輪密鑰。然后，將輸入數(shù)據(jù)和擴展密鑰存儲到GPU全局存儲空間內(nèi)。在加密的初始階段，輸入數(shù)據(jù)將被分成每塊1024字節(jié)的數(shù)據(jù)分組，并行地完成加密和解密運算。在整個加密和解密過程中，每個CUDA線程塊負責計算一個輸出塊，而每個CUDA線程塊由256個GPU線程組成。加密完成后，輸出數(shù)據(jù)的結(jié)果將再次被寫到全局設(shè)備存儲器內(nèi)。隨后，CPU程序?qū)⑤敵鰯?shù)據(jù)結(jié)果從GPU 的全局設(shè)備存儲器內(nèi)取回，整個加密過程完成。解密過程類似。

為了達到最好的加速性能，采用每塊256個線程[8]。支持塊的大小(256)和網(wǎng)格的大小定義如下int BK=256;Dim3 dmblock(BK);Dim3 dmgrid(size/BK/16)。核函數(shù)的調(diào)用方式為 cryptKernel <<< dmgrid, dmblock >>>(d_sk, d_input, d_output);dmgrid表示網(wǎng)格的三維數(shù)據(jù)，dmblock表示塊的三維結(jié)構(gòu)，d_sk表示擴展密鑰，d_input表示明文，d_output表示加密后的密文。

4 實驗結(jié)果與分析

4.1 實驗環(huán)境

實驗設(shè)計在2個不同的實驗平臺上進行，實驗平臺1 和實驗平臺2 的配置有相同的主要環(huán)境參數(shù)。。實驗平臺配置參數(shù)如下：開發(fā)環(huán)境-Windows 7,Visual Studio 2013,CUDA Toolkit6.5；GPU- GeForce GT240M 1.21 GHz 1GB顯存，48個流處理器；CPU- Intel Core T6600 2.20GHz 2G內(nèi)存。

4.2 實驗結(jié)果

實驗的方法是在兩個平臺上利用單個CPU核編程實現(xiàn)的串行SM4加密算法與CUDA編程平臺下的并行SM4加密算法進行比較，測量的數(shù)據(jù)長度從16B到64MB以2的倍數(shù)遞增，測量的時間包括CPU串行加密時間，CUDA并行加密片上運行時間，CPU和GPU之間的數(shù)據(jù)傳輸時間以及并行加密的總時間。每種明文大小的算法均運行10次，取平均值以減少數(shù)據(jù)誤差，具體實驗數(shù)據(jù)如表1所列。

4.3 實驗分析與討論

表1給出了實驗平臺1的完整數(shù)據(jù)，給出了明文大小從64B到32MB的串行SM4和并行SM4的時間和性能比較情況。用 表示串行SM4算法在CPU上的加密時間。用 表示在GPU上執(zhí)行并行SM4算法的總時間，包括并行數(shù)據(jù)加密及傳輸明密文時間。

表1 實驗平臺1并行算法和串行算法性能比較表 （單位：ms）

加速比表示并行SM4相比串行SM4的性能提升，計算公式如式（9）所示：

隨著明文大小成倍增加，串行SM4的計算時間也成倍增加，而并行SM4的執(zhí)行總時間以緩慢的速度增長，如圖2所示。

實驗平臺1和平臺2下，明文大小在分別小于8K和512B，并行SM4算法執(zhí)行時間要比串行SM4執(zhí)行時間多，這是因為一方面本身單線程CPU的執(zhí)行效率比GPU高，另一方面從CPU到GPU傳輸數(shù)據(jù)也要一定的時間。當加密數(shù)據(jù)較少時（平臺1≤8K，平臺2≤512B），使用GPU進行并行加密算法相比于CPU串行加密的優(yōu)勢還沒有體現(xiàn)出來，加速比只有1左右；當加密數(shù)據(jù)的長度超過一定值時（平臺1>8K，平臺2>512B），使用GPU進行加密的速度優(yōu)勢就顯現(xiàn)出來了，無論從是SM4并行加密算法在GPU片上時間還是總時間上都與串行算法總時間逐漸拉開。

圖2 (a)串行SM4與并行SM4執(zhí)行時間比較；(b)并行SM4算法加速比

從圖2，實驗平臺1的明文數(shù)據(jù)在8K到16M之間和實驗平臺2的明文數(shù)據(jù)在4K到256K之間，加速比迅速提升。實驗平臺1，當明文數(shù)據(jù)量繼續(xù)增大時，加速比雖然仍然提升但是提升變得緩慢，最終在40.6倍左右停頓了下來。而實驗平臺2，在明文數(shù)據(jù)在256K是，加速比達到極限值94倍左右。當明文數(shù)據(jù)量繼續(xù)增大時，加速比卻下降，且慢慢趨近于64倍左右。

通過觀察圖2(a)，由于CPU性能的提升，平臺2串行執(zhí)行時間比平臺1少，但降低的幅度并不明顯。同時，由于GPU性能的提升，平臺2的加速比相對于平臺1提升了將近50%。由此可看出，GPU性能提升帶給SM4的并行算法執(zhí)行效率要比CPU性能提升帶給SM4的串行執(zhí)行效率高。

不同 GPU的核心數(shù)目不同，帶來的加速性能的評價不能單純地從加速比判斷。為了消除不同的GPU給并行算法帶來的加速比差異性，引入加速效率的評價標準。加速效率 表示并行程序提升加速比的性能，由加速比和GPU的核心數(shù)C(GT240M為C=48)之比表示，計算公式如式（10）所示。

通過計算得實驗出平臺1的加速效率峰值為85.4%，實驗平臺2的加速效率峰值為49.5%。說明該并行算法充分挖掘了實驗平臺1的GPU設(shè)備多核心計算能力，實驗平臺2的GPU設(shè)備多核心計算能力沒有充分挖掘。

5 總結(jié)

為了解決高速數(shù)據(jù)流的實時加解密為題，采用了基于CUDA的并行SM4算法來提升SM4的加密性能。本文對傳統(tǒng)CPU的串行SM4加密算法和該并行SM4算法進行了實現(xiàn)和比較。實驗結(jié)果表明，該并行SM4算法在平臺1上最高能達到40.6倍的加速比和85.4%的加速效率，在平臺2上最高能達到64.7倍的加速比和49.5%的加速效率。很明顯，在支持新的統(tǒng)一架構(gòu)的GPU上實現(xiàn)并行SM4算法獲得了更好的性能。

[1]國家密碼管理局.國家密碼管理局公告第23號[EB/OL]. (2012-03-21).

http://www.oscca.gov.cn/News/201204/News_1227.htm.

[2]程海,丁群,杜輝,黃春光.基于FPGA實現(xiàn)的SMS4算法研究[J].儀器儀表學報, 2011, 32(12):2845-2850.

[3]馮春雨,胡波,劉會忠.基于FPGA的SMS4密碼算法的高速實現(xiàn)[J].河北省科學院學報,2010,27(6):8-11.

[4]李大為,趙旭鑫,武萌,SM S4密碼算法的高速流水線實現(xiàn)[J].電子器件, 2007, 30( 2): 590-592.

[5] Gao Xianwei,Lu Erhong,Xian Liqin,et al. FPGA Implementation of the SMS4 Block Cipher in the Chinese WAPI Standard[C]//Proc. of International Conference on Embedded Software and Systems Symposia. [S. l.]: IEEE Press,2008,104-106 [6]NVIDIA Corporation. CUDA Technology[OL]. http://www.nvidia.com/CUDA

[7]Manavski S A. CUDA compatible GPU as an efficient hardware accelerator for AES cryptography[C]∥2007 IEEE International Conference on Signal Processing and Communications ( ICSPC 2007).Dubai,United Arab Emirates,November 2007

[8]蘇統(tǒng)華,李東等譯.CUDA并行程序設(shè)計GPU編程指南[M].北京,機械工業(yè)出版社,2014

High Speed Implementation of SM4 Encryption Algorithm Based on CUDA

WANG De-min CHEN Da
(School of Information & Communication Engineering, Guilin University of Electronic Technology Guilin 541004 China)

Traditional SM4 encryption algorithm is implemented on the CPU. In order to improve the speed of encryption to deal with large-scale cryptographic operations, according to the structure and characteristics of the block cipher SM4, a SM4 parallel algorithm that using of GPU computing capabilities and running on the Compute Unified Device Architecture (CUDA ) platform is achieved. Two experiments platform, SM4 parallel algorithm (based on CUDA) and serial algorithm (based on CPU) performance were compared with experiment. The results show that the SM4 parallel algorithm on the platform 1 can achieve the highest speed of 40.6 times and 85.4% speedup acceleration efficiency, the algorithm on the platform 2 can achieve the highest speed of 64.7 and 49.5% speedup acceleration efficiency.

SM4 CUDA parallel computing

A

1673-1816(2017)01-0059-05

2016-04-12

王德民（1991-）男，海南臨高人，學士，研究方向集成電路。