基于CUDA的SM4加密算法高速實現

王德民 陳 達

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

基于CUDA的SM4加密算法高速實現

王德民 陳 達

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

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

SM4 CUDA 并行計算

1 引言

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

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

2 SM4算法描述

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

2.1 加解密算法描述

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

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

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

2.2 密鑰擴展算法

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

3 CUDA模型及并行算法設計

3.1 CUDA模型

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

3.2 并行算法設計

圖1 SM4并行算法整體框架

SM4算法的分組加密結構適合進行并行處理，因為明文分組之間不存在數據依賴，加密時按明文分組進行并行加密，能充分利用 GPU 多核的處理能力。SM4并行算法實現如圖1所示。

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

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

4 實驗結果與分析

4.1 實驗環境

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

4.2 實驗結果

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

4.3 實驗分析與討論

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

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

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

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

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

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

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

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

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

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

5 總結

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

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

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

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

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

[4]李大為,趙旭鑫,武萌,SM S4密碼算法的高速流水線實現[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]蘇統華,李東等譯.CUDA并行程序設計GPU編程指南[M].北京,機械工業出版社,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-）男，海南臨高人，學士，研究方向集成電路。