李曉東， 胡一鳴， 池亞平， 錢 榕， 張健毅

北京電子科技學(xué)院， 北京 100070

1 概述

隨著大數(shù)據(jù)、云計(jì)算、5G 等新技術(shù)的快速發(fā)展和廣泛應(yīng)用， 在高速網(wǎng)絡(luò)系統(tǒng)中如何保證數(shù)據(jù)傳輸?shù)陌踩煽扛咝С蔀槟壳把芯康臒狳c(diǎn)之一. 密碼算法一直是確保國(guó)家和經(jīng)濟(jì)生產(chǎn)活動(dòng)中數(shù)據(jù)傳輸安全的關(guān)鍵， 尤其是國(guó)產(chǎn)密碼算法的研究日益重要. SM4 算法[1]是目前國(guó)產(chǎn)密碼算法中廣泛應(yīng)用于金融、工業(yè)生產(chǎn)等領(lǐng)域的重要算法之一.

GPU 的概念是由NVIDIA 公司于1999 年提出的， 它是顯卡上的一塊芯片， 外接在PCI-E 接口， 主要用作輔助處理. GPU 自誕生起面對(duì)的就是類型高度統(tǒng)一的、相互無(wú)依賴的大規(guī)模數(shù)據(jù)和不需要被打斷的純凈的計(jì)算環(huán)境. 在2007 年CUDA 發(fā)布后， GPU 編程不需要再借用圖形API， 開發(fā)者只需要通過(guò)C語(yǔ)言就可以進(jìn)行編程， 從而大大降低了開發(fā)的門檻. 此后， GPU 不再只能處理圖像， 同時(shí)還可以通過(guò)調(diào)整線程塊大小、存儲(chǔ)空間使用、數(shù)據(jù)傳輸模式等手段對(duì)計(jì)算密集型和易于并行的程序進(jìn)行優(yōu)化， 使其運(yùn)行速度更快、效率更高[2].

分組密碼的工作模式可以對(duì)原有的分組密碼算法進(jìn)行加強(qiáng)， 使分組密碼能夠應(yīng)用于實(shí)際加密中， 且不同的工作模式其性能和安全性也完全不同， 合適的工作模式可以彌補(bǔ)分組密碼算法的缺陷. 1980 年，美國(guó)國(guó)家標(biāo)準(zhǔn)局(現(xiàn)在的NIST) 公布了DES 的4 種工作模式: ECB (電子密碼本， electronic code book)、CBC (密文分組鏈接， cipher block chaining)、CFB (密文反饋， cipher feedback)、OFB (輸出反饋， output feedback)， 之后在2001 年， 新增了由Diffie 和Hellman 在1979 年提出的CTR (計(jì)數(shù)器，counter) 模式[3]. 現(xiàn)有的關(guān)于分組密碼算法工作模式的研究便是在這五種基礎(chǔ)工作模式上進(jìn)行改進(jìn). 在CUDA 發(fā)布之初的2007 年， Manavski[4]就首次實(shí)現(xiàn)了利用CUDA 并行加速AES 密碼算法， 并比較了不同的GPU 內(nèi)存使用對(duì)AES 算法計(jì)算速度的影響， 為后續(xù)其他團(tuán)隊(duì)的研究提供了借鑒. 之后， Harrison等人[5]實(shí)現(xiàn)了基于CUDA 的AES 算法CTR 模式運(yùn)行; Nishikawa 等人[6]更是將ECB 模式下的AES密碼算法提升到了605.9 Gbps. Kim 等人[7]實(shí)現(xiàn)RSA 算法的GPU 加速且加速比可達(dá)10 以上. 費(fèi)雄偉等人[8]研究了明文大小、線程分塊、存儲(chǔ)選擇對(duì)CTR 模式的AES 算法并行運(yùn)算速度的影響， 所得到的綜合加速比高達(dá)31.59. 可以說(shuō)， 現(xiàn)在針對(duì)對(duì)稱密碼算法和非對(duì)稱密碼算法的CUDA 并行實(shí)現(xiàn)和加速已經(jīng)有了較多的研究.

在CUDA 實(shí)現(xiàn)SM4 密碼算法并行方面， 楊海云等人[9]基于CUDA 分別實(shí)現(xiàn)了AES 和SM4 算法并行; 王德民等人[10]從明文大小角度研究CUDA 并行對(duì)SM4 算法運(yùn)行速度的影響， 當(dāng)明文過(guò)小時(shí)，CPU 和GPU 之間的數(shù)據(jù)傳輸耗時(shí)較長(zhǎng)， 無(wú)法起到很好的加速作用， 當(dāng)明文較大時(shí)， 加速比顯著增加后會(huì)趨于平穩(wěn)， 不會(huì)無(wú)限制的擴(kuò)大; 張才賢[11]針對(duì)SM4 認(rèn)證加密方面設(shè)計(jì)出一套新的GCM 工作模式用于高速認(rèn)證中的加密模塊; Cheng 等人[12]基于GPU 設(shè)計(jì)了一種高性能并行對(duì)稱密碼服務(wù)器， 通過(guò)GPU內(nèi)存選擇和優(yōu)化CPU 和GPU 之間的數(shù)據(jù)傳輸， 數(shù)據(jù)加密速度可達(dá)到15.96 Gbps， 是現(xiàn)有對(duì)稱加密服務(wù)器速度的23 倍. 其他的研究工作包括SM4 的專用硬件實(shí)現(xiàn)[13]、多模式實(shí)現(xiàn)[14]以及針對(duì)SM4 實(shí)現(xiàn)的攻擊[15-17].

綜上所述， 目前針對(duì)SM4 算法的研究大多是采用傳統(tǒng)的ECB 模式在CPU 或GPU 上進(jìn)行實(shí)現(xiàn)和優(yōu)化. 而除去傳統(tǒng)的ECB 工作模式， 多數(shù)團(tuán)隊(duì)針對(duì)SM4 算法的研究集中于在特定使用場(chǎng)景設(shè)計(jì)新的工作模式或利用軟件實(shí)現(xiàn)算法的快速運(yùn)行. 由此可見， 已有的研究存在如下問(wèn)題:

(1) 采用傳統(tǒng)工作模式在安全性方面有所欠缺;

(2) 實(shí)驗(yàn)平臺(tái)通用性不夠;

(3) 算法優(yōu)化后針對(duì)運(yùn)行效果討論不足.

本文在全面分析GPU 下并行計(jì)算的特點(diǎn)的基礎(chǔ)上， 在通用計(jì)算機(jī)平臺(tái)上對(duì)國(guó)密分組密碼SM4 算法的CTR 模式在GPU 上并行實(shí)現(xiàn)策略進(jìn)行了深入討論， 本文具體貢獻(xiàn)有:

(1) 系統(tǒng)地分析核函數(shù)設(shè)計(jì)、明文大小、線程塊大小等對(duì)算法的影響. 在設(shè)計(jì)出SM4-CTR 算法的核函數(shù)后， 對(duì)線程網(wǎng)格、線程塊大小以及內(nèi)存類型選擇進(jìn)行優(yōu)化， 實(shí)現(xiàn)了并行加解密;

(2) 實(shí)現(xiàn)了SM4-CTR 在不同架構(gòu)GPU 下的并行化算法， 并對(duì)比了不同架構(gòu)GPU 的加速效果及CTR 相對(duì)于傳統(tǒng)工作模式的加速效果. 選擇不同架構(gòu)的新舊兩個(gè)GPU 實(shí)驗(yàn)環(huán)境， 分別研究?jī)?yōu)化后的SM4-CTR 運(yùn)行效果， 并與相同條件下利用CPU 串行運(yùn)行和基于傳統(tǒng)工作模式并行運(yùn)行的結(jié)果進(jìn)行對(duì)比. 實(shí)驗(yàn)發(fā)現(xiàn)， GPU 平臺(tái)面對(duì)較大明文時(shí)， 相比于計(jì)算能力有限的CPU 具有較大的性能優(yōu)勢(shì); CTR 模式在安全性能提高的同時(shí)， 其并行運(yùn)行的速率也優(yōu)于傳統(tǒng)工作模式;

(3) 系統(tǒng)性地對(duì)比了SM4-CTR 的CPU、GPU 以及快速軟件優(yōu)化的運(yùn)行效率， 為SM4 算法的應(yīng)用以及未來(lái)相關(guān)的研究奠定基礎(chǔ). 利用實(shí)驗(yàn)結(jié)果與之前團(tuán)隊(duì)的研究成果進(jìn)行對(duì)比， 結(jié)果顯示本文提出的方案在通用計(jì)算機(jī)平臺(tái)的運(yùn)行效果大幅優(yōu)于傳統(tǒng)工作模式下利用CPU、GPU 和快速軟件優(yōu)化的運(yùn)行效果.

2 SM4 的工作模式

2.1 SM4 密碼算法

SM4 是2006 年中國(guó)國(guó)家密碼管理局發(fā)布的用于商用領(lǐng)域的分組密碼算法， 在2016 年成為國(guó)家標(biāo)準(zhǔn).SM4 密碼算法為非平衡Feistel 結(jié)構(gòu)分組密碼算法， 分組長(zhǎng)度和密鑰長(zhǎng)度都為128 位， 加解密過(guò)程和密鑰擴(kuò)展都采用32 輪非線性迭代結(jié)構(gòu)， 加密算法和解密算法結(jié)構(gòu)相同， 只有輪密鑰的使用順序相反. SM4 密碼算法由密鑰擴(kuò)展算法和加解密算法組成， 加解密算法中最為重要的就是輪函數(shù)F.

SM4 密碼算法的密鑰長(zhǎng)度為128 位， 表示為: MK = (MK0，MK1，MK2，MK3)， 其中MKi(i=0，1，2，3) 為32 位. 輪密鑰表示為: rk0，rk1，···，rk31， 其中rki(i=0，1，···，31) 由128 位MK 經(jīng)32 輪迭代生成.

輪密鑰rki計(jì)算公式如下:

輪函數(shù)F是SM4 算法加解密過(guò)程的重要組成部分， 它主要由非線性變換τ和線性變換L組成， 先將128 位的明文x分成4 個(gè)32 位的Xi(i= 0，1，2，3)， 再經(jīng)32 輪迭代， 公式(其中i= 0，1，···，31) 如下:

經(jīng)過(guò)32 輪迭代后， 明文x新生成了32 個(gè)位Xi， 記為:Xi(i= 0，1，···，35)， 則密文y表示為:y=(Y0，Y1，Y2，Y3)=(X35，X34，X33，X32). 輪函數(shù)F和完整SM4 算法加密流程如圖1 所示.

圖1 輪函數(shù)F 計(jì)算流程和SM4 加密流程Figure 1 Round function F calculation process and SM4 encryption process

SM4 算法的解密過(guò)程與加密過(guò)程基本一致， 只是輪密鑰使用順序不同. 加密輪密鑰使用順序?yàn)閞k0，rk1，···，rk31， 解密輪密鑰是加密輪密鑰的逆序， 即rk31，rk30，···，rk0.

2.2 ECB 模式

ECB 模式是最簡(jiǎn)單直接的也是最傳統(tǒng)常用的工作模式， 直接將明文分成固定大小的塊， 每個(gè)塊獨(dú)立地進(jìn)行加密， 在加密時(shí)需要將明文填充至塊大小的整數(shù)倍. ECB 模式下加解密公式為:

其中，C代表密文，P代表明文，Ek代表加密，代表解密.

ECB 工作模式優(yōu)點(diǎn)有:

(1) 簡(jiǎn)單， 計(jì)算速度快;

(2) 支持并行計(jì)算;

(3) 運(yùn)算錯(cuò)誤不會(huì)被傳送， 無(wú)需知道明文長(zhǎng)度.

ECB 工作模式缺點(diǎn)是:

(1) 加密長(zhǎng)度必須是分組長(zhǎng)度的整數(shù)倍;

(2) 加密過(guò)程是高度確定的， 只要密鑰不發(fā)生改變， 同樣的明文加密后密文也相同， 易受到攻擊;

(3) 無(wú)法檢測(cè)密文分組的順序， 無(wú)法抵抗分組的重放、嵌入、刪除攻擊.

這些缺點(diǎn)導(dǎo)致ECB 模式非常不安全， 實(shí)用性不強(qiáng)， 現(xiàn)在ECB 模式僅用于一些較短明文的加解密.ECB 模式加解密流程如圖2 所示.

圖2 ECB 模式加解密Figure 2 ECB mode encryption and decryption

2.3 CTR 模式

CTR 模式， 是一種類似于流密碼的模式， 加密算法只是用來(lái)產(chǎn)生密鑰流與明文分組異或. 該模式加密的輸入是一個(gè)計(jì)數(shù)器值， 加密后再與明文數(shù)據(jù)進(jìn)行異或得到密文， 其中計(jì)數(shù)器的值應(yīng)該是互不相同的.

CTR 模式下的加解密公式如下:

其中， ctr 代表計(jì)數(shù)器值，C代表密文，P代表明文，Ek代表加密.

CTR 工作模式優(yōu)點(diǎn)有:

(1) CTR 模式的加密和解密使用了完全相同的結(jié)構(gòu)， 因此在程序?qū)崿F(xiàn)上比較容易;

(2) CTR 模式中能夠以任意順序處理分組， 這就意味著能夠?qū)崿F(xiàn)并行計(jì)算， 在執(zhí)行并行計(jì)算的系統(tǒng)中CTR 模式的速度是非常快的;

(3) CTR 模式中加解密相同， 安全性好.

CTR 工作模式缺點(diǎn)有: 需要加解密雙方要保持同步.

CTR 模式加解密流程如圖3 所示.

圖3 CTR 模式加解密Figure 3 CTR mode encryption and decryption

3 基于CTR 的GPU 并行SM4 算法

3.1 并行化策略

CTR 模式與傳統(tǒng)SM4 算法采用的ECB 模式相似， 加解密過(guò)程都支持并行， 但CTR 模式中引入了計(jì)數(shù)器值， 而計(jì)數(shù)器值本身也會(huì)占用一定量的GPU 全局內(nèi)存空間， 所以在設(shè)計(jì)核函數(shù)時(shí)需要針對(duì)這一點(diǎn)進(jìn)行相關(guān)的改進(jìn). 并行方式如圖4 所示.

圖4 CTR 模式加解密并行Figure 4 CTR mode encryption and decryption parallel

其中，P是128 位明文塊，C是128 位密文塊， rk 是加密輪密鑰，E是加密過(guò)程，n是線程總數(shù). 加密過(guò)程和解密過(guò)程完全一樣， 只有輸入不同， 加密時(shí)輸入明文輸出密文， 解密時(shí)輸入密文輸出明文. 每個(gè)線程獨(dú)立完成自己的計(jì)算， 各線程之間互不干擾.

3.2 并行化實(shí)現(xiàn)

主機(jī)端負(fù)責(zé)CTR 模式下的SM4 算法的初始化和接收階段， 設(shè)備端負(fù)責(zé)加解密計(jì)算， 并行化實(shí)現(xiàn)主要包括主機(jī)端初始化階段、設(shè)備端并行計(jì)算階段、主機(jī)端結(jié)束階段這三個(gè)階段.

3.2.1 主機(jī)端初始化階段

這一階段的任務(wù)主要是將密碼算法所用到的密鑰、明文、密文等數(shù)據(jù)傳輸?shù)紾PU 內(nèi)存中.

(1) 利用fopen() 函數(shù)讀取明文數(shù)據(jù)text， 并生成密鑰key;

(2) 生成S 盒數(shù)據(jù)存于GPU 內(nèi)存中;

(3) 生成一個(gè)隨機(jī)的計(jì)數(shù)器值counter;

(4) 根據(jù)密鑰key 生成加密輪密鑰rk 和解密輪密drk:SM4Key(key， rk， 0);SM4Key(key， drk， 1);

(5) 設(shè)置運(yùn)行GPU， 單GPU 系統(tǒng)不需要這一步， 多GPU 系統(tǒng)默認(rèn)使用第0 號(hào)GPU (在之后的內(nèi)容中不再提及此步驟):cudaSetDevice(0);

(6) 為明文p、密文c、輪密鑰rk、計(jì)數(shù)器值counter 申請(qǐng)GPU 全局內(nèi)存空間:cudaMalloc((void**)&p，size*sizeof(muint8)); //明文cudaMalloc((void**)&c，size*sizeof(muint8)); //密文cudaMalloc((void**)&rk，32*sizeof(muint32)); //輪密鑰cudaMalloc((void**)&counter， 16*sizeof(muint8)); //計(jì)數(shù)器值其中， size 代表明文的字節(jié)數(shù);

(7) 向GPU 發(fā)送明文數(shù)據(jù)text、輪密鑰數(shù)據(jù)rk、計(jì)數(shù)器值counter:cudaMemcpy(p， text， size*sizeof(muint8)， cudaMemcpyHostToDevice); //明文cudaMemcpy(rk， key， 32*sizeof(muint32)， cudaMemcpyHostToDevice); //輪密鑰cudaMemcpy(counter， Counter， 16*sizeof(muint8)， cudaMemcpyHostToDevice); //輪密鑰

(8) 向GPU 發(fā)送明文數(shù)據(jù)text、輪密鑰數(shù)據(jù)rk、計(jì)數(shù)器值counter:SM4_encrypt_ctr?＜blocks，threads?＞(p，c，rk，n，counter).

3.2.2 設(shè)備端并行計(jì)算階段

這一階段通過(guò)線程劃分、啟動(dòng)進(jìn)行加解密計(jì)算.

(1) 確認(rèn)線程id: inti=(blockIdx.x*blockDim.x)+threadIdx.x;

(2) CTR 模式中算法輸入為計(jì)數(shù)器值， 每個(gè)計(jì)數(shù)器值都不相同:

(3) 啟動(dòng)線程， 并將線程所需要進(jìn)行計(jì)算的數(shù)據(jù)存于該線程獨(dú)占的寄存器中:SM4_ctr_encrypt_thread(ca，&(c[i*16])，rk，&(p[i*16]));

圖5 CTR 模式SM4 線程內(nèi)算法Figure 5 SM4 in-thread algorithm in CTR mode

(4) 核函數(shù)計(jì)算完成后， 之前為密文c申請(qǐng)的GPU 全局內(nèi)存已經(jīng)存儲(chǔ)了計(jì)算完成后的密文.

3.2.3 主機(jī)端結(jié)束階段

這一階段主要將數(shù)據(jù)傳輸回主機(jī)端， 并釋放內(nèi)存.

(1) 將GPU 全局內(nèi)存中的密文數(shù)據(jù)傳回CPU:cudaMemcpy(cipher，c，size*sizeof(muint8)，cudaMemcpyDeviceToHost);

(2) 釋放GPU 上明文、密文、輪密鑰數(shù)據(jù):

4 實(shí)驗(yàn)及其分析

4.1 CUDA 編程對(duì)性能影響

GPU 上實(shí)現(xiàn)密碼算法的并行計(jì)算， 需要考量的方面非常多， 這些方面的合理與否， 直接影響著密碼算法的執(zhí)行速度. 它包括:

(1) 線程束分配線程束作為CUDA 最小的執(zhí)行單位， 要確保束內(nèi)有32 個(gè)線程， 線程塊內(nèi)有32 的整數(shù)倍個(gè)線程，線程束在進(jìn)行數(shù)據(jù)傳輸時(shí)也要考慮到內(nèi)存地址的影響， 避免出現(xiàn)全局內(nèi)存非合并的情況發(fā)生.

(2) 內(nèi)存選擇和使用不同內(nèi)存的大小、延遲、使用方法各不相同， 在使用的時(shí)候應(yīng)該考慮到各內(nèi)存的特點(diǎn). 比如對(duì)于S 盒這類小體量的數(shù)據(jù)存放在常量?jī)?nèi)存或共享內(nèi)存就比存放在全局內(nèi)存可得到更優(yōu)的密碼算法運(yùn)算速度. 同時(shí)， 在使用共享內(nèi)存的時(shí)候， 要盡量避免bank 沖突(bank-conflict)， 這會(huì)讓原本并行的對(duì)共享內(nèi)存的訪存操作變成串行從而極大地降低程序效率.

(3) 線程塊大小一個(gè)線程塊內(nèi)要有足夠多的線程， 因?yàn)橐粋€(gè)線程塊會(huì)占用一個(gè)SM，如果塊內(nèi)線程過(guò)少會(huì)導(dǎo)致SM內(nèi)多個(gè)Core 閑置.

(4) 線程數(shù)目線程數(shù)目的大小直接決定并行規(guī)模， 是最主要影響算法并行計(jì)算速度的因素， 應(yīng)盡可能的增加線程數(shù)量.

CUDA 最大的特點(diǎn)就是其高度的編程自由性， 開發(fā)人員可以根據(jù)自己的需求對(duì)程序進(jìn)行修改， 但也因?yàn)槠渥杂尚詫?dǎo)致在CUDA 代碼編寫的時(shí)候要考慮多方因素， 任何一個(gè)微小的差別都可能對(duì)程序的執(zhí)行速度造成巨大的影響.

4.2 實(shí)驗(yàn)環(huán)境

為了研究SM4 算法CTR 模式下GPU 并行的效果， 首先要讀取出實(shí)驗(yàn)計(jì)算機(jī)中的相關(guān)參數(shù)， 在這里可以通過(guò)編寫CheckGPU 程序， 確定自己電腦GPU 各項(xiàng)數(shù)據(jù)和實(shí)驗(yàn)環(huán)境， 其中要用到cudaGetDevice-Count、cudaGetDeviceProperties 等函數(shù). 為了全面分析本文設(shè)計(jì)實(shí)現(xiàn)的SM4-CTR 算法的并行方案的可行性和高效性， 本文搭建了兩個(gè)不同測(cè)試實(shí)驗(yàn)環(huán)境(如表1 所示) ， 并在這兩個(gè)實(shí)驗(yàn)環(huán)境中進(jìn)行了多角度實(shí)驗(yàn)數(shù)據(jù)的采集和分析. 總的來(lái)看， 雖然兩個(gè)實(shí)驗(yàn)環(huán)境的計(jì)算性能有很大差異， 不同平臺(tái)所得到的加速比也不同， 但是加速比的變化趨勢(shì)是一樣的. 其中實(shí)驗(yàn)環(huán)境1 為目前新的顯卡架構(gòu)， 屬于通用計(jì)算平臺(tái).

表1 實(shí)驗(yàn)環(huán)境Table 1 Experiment environment

4.3 CTR 模式性能分析實(shí)驗(yàn)結(jié)果

實(shí)驗(yàn)主要通過(guò)CTR 模式下串行加解密和并行加解密時(shí)間的比值來(lái)衡量GPU 并行加速的效果， 其公式如下:

其中，TC代表CPU 串行執(zhí)行加解密所需要的時(shí)間，TG代表GPU 并行執(zhí)行加解密所需要的時(shí)間，Tt代表在GPU 上分配內(nèi)存、釋放內(nèi)存以及CPU 與GPU 之間的數(shù)據(jù)傳輸所需要的時(shí)間，Tg代表GPU 并行進(jìn)行計(jì)算所需要的時(shí)間，Tt+Tg表示GPU 并行執(zhí)行加解密所需要的時(shí)間TG.

4.3.1 不同明文大小

在實(shí)驗(yàn)環(huán)境1 中， SM4 算法的CTR 模式下同時(shí)利用CPU 串行和GPU 并行加解密一次不同大小的明文， 其中明文大小選取從16 KB 到256 MB， 每個(gè)線程塊中都是1024 個(gè)線程， 根據(jù)最后的結(jié)果計(jì)算加速比， 結(jié)果如圖6 所示.

圖6 不同明文大小的并行加速比Figure 6 Parallel speedup of different plaintext sizes

從實(shí)驗(yàn)結(jié)果中可以看出來(lái)， 當(dāng)明文數(shù)據(jù)小于64 KB 左右的時(shí)候， 利用GPU 對(duì)SM4 算法進(jìn)行并行執(zhí)行， 算法運(yùn)行速度還不如串行執(zhí)行， 但在之后加速比逐漸提高， 并且隨著明文的加大加速比也在提高， 最終加速比大致能達(dá)到40 倍. 同時(shí)在明文數(shù)據(jù)等于16 KB、32 KB、64 KB 的時(shí)候， GPU 并行計(jì)算時(shí)間TG幾乎相同， 這是因?yàn)榫€程塊的數(shù)目分別為1 個(gè)、2 個(gè)、4 個(gè)， 而實(shí)驗(yàn)環(huán)境中GPU 的SM (流式多處理器)有24 個(gè)， 也就是說(shuō)最大允許24 個(gè)線程塊同時(shí)參與計(jì)算， 當(dāng)明文數(shù)據(jù)較小時(shí)， 有部分SM 處于閑置狀態(tài)，這時(shí)提升線程塊的數(shù)目并不會(huì)拖慢GPU 計(jì)算速度.

可知雖然并行執(zhí)行SM4 算法可以提升算法執(zhí)行速度， 但不能無(wú)限制地增大， 一方面由于GPU 雖然可以實(shí)現(xiàn)上千個(gè)線程同時(shí)執(zhí)行但Core 數(shù)量也是有限的， 當(dāng)線程過(guò)多的時(shí)候， 多余的線程只能等待其他線程執(zhí)行完畢， 另一方面GPU 上SM4 算法的計(jì)算速度已經(jīng)快過(guò)數(shù)據(jù)傳輸速度， 算法并行性不再是限制算法執(zhí)行速度的瓶頸， 反而是總線傳輸數(shù)據(jù)的速度限制了算法的執(zhí)行速度.

4.3.2 不同線程塊大小

在實(shí)驗(yàn)環(huán)境2 中， 選取8 MB 大小的明文， 令SM4 算法的CTR 模式進(jìn)行并行加解密計(jì)算， 期間僅改變線程塊大小， 同時(shí)為了防止單次加解密的誤差， 這里進(jìn)行100 次計(jì)算， 實(shí)驗(yàn)結(jié)果如表2 所示.

表2 不同線程塊大小CTR 模式并行加解密速度Table 2 Parallel encryption and decryption speed in CTR mode with different thread block sizes

根據(jù)實(shí)驗(yàn)結(jié)果， CTR 模式當(dāng)線程塊大小較小的時(shí)候， 即使線程數(shù)目沒(méi)有改變， 也會(huì)影響算法的計(jì)算速度， 這是由于一個(gè)線程塊在進(jìn)行計(jì)算時(shí)就會(huì)占用一個(gè)SM， 無(wú)論線程塊的大小都肯定占用一個(gè)， 所以當(dāng)線程塊過(guò)小的時(shí)候SM 中會(huì)有大量的Core 不參與計(jì)算從而浪費(fèi)了GPU 資源; 當(dāng)線程塊的大小在128～768 之間時(shí)， SM4 算法的計(jì)算速度大體相同， 大約在256 時(shí)達(dá)到最高的運(yùn)算速度; 當(dāng)線程塊接近1024 時(shí)， 計(jì)算速度又有所減慢， 這是由于雖然一個(gè)SM 可以最大承載1024 線程同時(shí)參與運(yùn)算， 但在數(shù)據(jù)調(diào)用的時(shí)候還是會(huì)出現(xiàn)一定的沖突， 比如調(diào)用密鑰的時(shí)候， 導(dǎo)致計(jì)算速度有少許的減慢.

基于上述結(jié)論， 當(dāng)線程塊的大小在1～128 之間時(shí)， 隨著線程塊的增大， SM4 算法的計(jì)算速度應(yīng)該是不斷增加的; 當(dāng)線程塊大小在128～768 之間時(shí)， SM4 算法的計(jì)算速度應(yīng)該不會(huì)有太大的改變. 但表3 并不符合這一規(guī)定.

表3 線程塊劃分對(duì)CTR 模式加解密速度的影響Table 3 Influence of thread block division on speed of CTR mode encryption and decryption

前文提到， CUDA 中同一個(gè)線程塊的線程會(huì)自動(dòng)以32 為單位組成線程束， 這是CUDA 最小的調(diào)度單位， 也就是說(shuō)同一個(gè)束的線程要同時(shí)參與計(jì)算. 但是CUDA 的高自由性又允許程序員將線程塊的大小規(guī)定成1～1024 之間的任意整數(shù)， 出現(xiàn)線程塊的大小不為32 的整數(shù)倍的情況， 這就會(huì)造成最后一個(gè)線程束不滿32 個(gè)線程， 但它還是會(huì)占用一個(gè)完整線程束的運(yùn)算單元， 造成GPU 資源浪費(fèi). 這就導(dǎo)致線程塊在64 和512 時(shí)SM4 算法計(jì)算速度明顯要大于70 和520， 這一規(guī)則也適用于其他線程塊不為32 的整數(shù)的情況， 所以在規(guī)劃線程塊大小的時(shí)候應(yīng)盡量選擇32 的整數(shù)倍.

4.4 CTR 模式與ECB 模式對(duì)比實(shí)驗(yàn)結(jié)果

本次實(shí)驗(yàn)是在實(shí)驗(yàn)環(huán)境2 中， 通過(guò)公式計(jì)算出ECB 和CTR 工作模式下的SM4 算法加解密加速比，對(duì)比兩種模式之間的差別.

在實(shí)驗(yàn)環(huán)境2 中， SM4 算法的CTR 模式和ECB 模式下同時(shí)利用CPU 串行和GPU 并行加解密一次不同大小的明文， 其中明文大小選取從16 KB 到512 MB， 每個(gè)線程塊中都是1024 個(gè)線程， 根據(jù)最后的結(jié)果計(jì)算加速比， 結(jié)果如圖7 所示.

從圖7 中可以看出， 當(dāng)明文數(shù)據(jù)小于64 KB 左右的時(shí)候， 利用GPU 對(duì)SM4 算法進(jìn)行并行執(zhí)行， 算法運(yùn)行速度還不如串行執(zhí)行， 但在之后加速比逐漸提高， 在2 MB 左右的時(shí)候提升速度最快， 但明文數(shù)據(jù)再增大的時(shí)候， 加速比提升速度逐漸減慢， 最終ECB 模式SM4 算法停在18 左右， CTR 模式SM4 密碼算法停在21 左右.

圖7 不同明文ECB、CTR 模式加速比Figure 7 Acceleration ratios of different plaintext ECB and CTR modes

由此可以看出， 前文關(guān)于CTR 模式在并行執(zhí)行運(yùn)行的情況下速度會(huì)更快的推斷是正確的， 在明文越來(lái)越大的情況下CTR 模式加解密速度也越來(lái)越快， 同時(shí)實(shí)驗(yàn)2 的環(huán)境下GPU 架構(gòu)并不是最新的， 因此在目前現(xiàn)實(shí)環(huán)境中新的架構(gòu)下CTR 模式并行運(yùn)行的加解密速度會(huì)比ECB 模式快更多. 而基于前文分析，CTR 模式的安全性天然優(yōu)于ECB 模式， 因此CTR 模式在實(shí)際生活中應(yīng)用范圍會(huì)更廣， 尤其是有海量數(shù)據(jù)需要處理的情況下.

5 對(duì)其他相關(guān)工作對(duì)比分析

5.1 SM4 不同運(yùn)行方式對(duì)比

根據(jù)上一章在實(shí)驗(yàn)環(huán)境1 下進(jìn)行的實(shí)驗(yàn)結(jié)果計(jì)算后發(fā)現(xiàn)， 在通用計(jì)算平臺(tái)經(jīng)過(guò)上文方式實(shí)現(xiàn)和優(yōu)化后， CTR 模式下SM4 并行加密或解密速率能達(dá)到: 14 194.4 Mbps. 根據(jù)這個(gè)結(jié)果， 與之前開展過(guò)的相關(guān)實(shí)驗(yàn)進(jìn)行對(duì)比， 如表4 所示.

表4 對(duì)比結(jié)果Table 4 Comparing results

通過(guò)上述對(duì)比我們可以發(fā)現(xiàn)以下幾點(diǎn):

(1) 文獻(xiàn)[2] 中， 李秀瀅等人采用的是一種在舊的架構(gòu)環(huán)境下通過(guò)傳統(tǒng)模式將SM4 算法在GPU 上實(shí)現(xiàn)并行計(jì)算且優(yōu)化不明顯， 因此加解密速率大幅低于本文優(yōu)化后CTR 模式下在最新GPU 架構(gòu)的通用計(jì)算平臺(tái)并行運(yùn)行的結(jié)果;

(2) 文獻(xiàn)[18] 中， 張笑從等人是基于CPU 上的指令集同時(shí)利用比特切片技術(shù)對(duì)SM4 算法進(jìn)行優(yōu)化，其速率顯著低于本文利用GPU 模型運(yùn)行的結(jié)果;

(3) 文獻(xiàn)[19] 中， 郎歡等人利用SIMD 技術(shù)實(shí)現(xiàn)SM4 算法的快速軟件實(shí)現(xiàn)， 雖然其硬件性能和架構(gòu)優(yōu)于本文的實(shí)驗(yàn)平臺(tái)， 但速率依然顯著低于本文.

為使對(duì)比更加精確， 選取8 MB 大小的明文在文獻(xiàn)[2] 相同的實(shí)驗(yàn)環(huán)境進(jìn)行加解密實(shí)驗(yàn)， 得出基于本文提出的CTR 模式優(yōu)化后的并行加解密速率約為3904 Mbps， 依舊快于文獻(xiàn)[2] 中速率.

5.2 對(duì)比結(jié)果分析

經(jīng)過(guò)上文對(duì)比可以發(fā)現(xiàn)， 基于本文提出的優(yōu)化后的SM4-CTR 模式在基于CUDA 模型的最新通用計(jì)算平臺(tái)下加解密速率顯著快于傳統(tǒng)工作模式下利用CPU 優(yōu)化、GPU 優(yōu)化以及利用軟件優(yōu)化的SM4 算法加解密速率. 同時(shí)在舊平臺(tái)， 本文的優(yōu)化方案依舊快于傳統(tǒng)SM4 工作模式的加解密速率.

綜上所述， 本文提出的優(yōu)化后的基于CUDA 模型下SM4-CTR 模式在通用計(jì)算平臺(tái)中可以取得出色的加解密運(yùn)行速率.

6 結(jié)束語(yǔ)

本文針對(duì)分組密碼算法SM4 基于傳統(tǒng)工作模式運(yùn)行安全性低、利用CPU 串行運(yùn)行效率低、算法優(yōu)化多依靠專門硬件平臺(tái)通用性低的問(wèn)題， 提出基于通用計(jì)算機(jī)平臺(tái)的SM4-CTR 算法在CUDA 編程模型下利用GPU 并行運(yùn)行的方式提高算法的運(yùn)行安全性、運(yùn)算效率和算法應(yīng)用的范圍. 在研究CUDA 編程對(duì)加解密算法性能的影響的基礎(chǔ)上， 通過(guò)本文的優(yōu)化在最新的通用計(jì)算機(jī)GPU 架構(gòu)環(huán)境下， 可將SM4 算法的運(yùn)行速度提高到接近40 倍. 同時(shí)研究發(fā)現(xiàn)， 不同大小的明文數(shù)據(jù)塊和線程塊大小的劃分對(duì)加解密影響較大， 低于64 KB 的小明文利用GPU 并行無(wú)法起到提速的效果， 而線程塊劃分應(yīng)當(dāng)為32 的整倍數(shù)， 同時(shí)還需要根據(jù)明文的特點(diǎn)選擇合適的劃分， 否則會(huì)因?yàn)橛芯€程塊閑置而導(dǎo)致加解密算法并行運(yùn)行的效果不理想. 通過(guò)多項(xiàng)對(duì)比實(shí)驗(yàn)， 發(fā)現(xiàn)在本文給出的實(shí)驗(yàn)環(huán)境中CTR 模式在高于256 KB 明文的時(shí)候并行運(yùn)行的加解密的效果優(yōu)于傳統(tǒng)的ECB 模式; 對(duì)比之前研究團(tuán)隊(duì)在傳統(tǒng)工作模式下利用CPU 優(yōu)化、利用GPU優(yōu)化和利用軟件快速實(shí)現(xiàn)優(yōu)化的運(yùn)行效果， 本文運(yùn)算效率均大幅領(lǐng)先.

綜上所述， 本文提出的優(yōu)化后的CTR 模式下SM4 分組算法在通用計(jì)算平臺(tái)下經(jīng)過(guò)GPU 加速后并行運(yùn)行加解密效率有顯著提升， 在安全性、運(yùn)算速率提高的同時(shí)適用平臺(tái)也更加廣泛. 因此， 本文方案在實(shí)際生活中針對(duì)大數(shù)據(jù)和個(gè)人數(shù)據(jù)的安全保護(hù)中必將發(fā)揮巨大的作用. 未來(lái)， 我們將密切關(guān)注GPU 架構(gòu)和分組密碼算法的工作模式的發(fā)展， 本文提出的算法優(yōu)化方案也會(huì)根據(jù)實(shí)際情況不斷優(yōu)化迭代以適應(yīng)現(xiàn)實(shí)的需求.