基于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-）男，海南临高人，学士，研究方向集成电路。