王德民 陳 達(dá)

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

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

王德民 陳 達(dá)

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

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

SM4 CUDA 并行計(jì)算

1 引言

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

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

2 SM4算法描述

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

2.1 加解密算法描述

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

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

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

2.2 密鑰擴(kuò)展算法

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

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

3.1 CUDA模型

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

3.2 并行算法設(shè)計(jì)

圖1 SM4并行算法整體框架

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

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

為了達(dá)到最好的加速性能，采用每塊256個(gè)線程[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表示擴(kuò)展密鑰，d_input表示明文，d_output表示加密后的密文。

4 實(shí)驗(yàn)結(jié)果與分析

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ù)。。實(shí)驗(yàn)平臺(tái)配置參數(shù)如下：開發(fā)環(huán)境-Windows 7,Visual Studio 2013,CUDA Toolkit6.5；GPU- GeForce GT240M 1.21 GHz 1GB顯存，48個(gè)流處理器；CPU- Intel Core T6600 2.20GHz 2G內(nèi)存。

4.2 實(shí)驗(yàn)結(jié)果

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

4.3 實(shí)驗(yàn)分析與討論

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

表1 實(shí)驗(yàn)平臺(tái)1并行算法和串行算法性能比較表 （單位：ms）

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

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

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

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

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

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

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

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

5 總結(jié)

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

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

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

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

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

[4]李大為,趙旭鑫,武萌,SM S4密碼算法的高速流水線實(shí)現(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è)計(jì)GPU編程指南[M].北京,機(jī)械工業(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-）男，海南臨高人，學(xué)士，研究方向集成電路。