基于通用計算平臺SM4-CTR 算法并行實現(xiàn)與優(yōu)化*

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

北京電子科技學院， 北京 100070

1 概述

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

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

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

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

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

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

(2) 實驗平臺通用性不夠;

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

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

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

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

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

2 SM4 的工作模式

2.1 SM4 密碼算法

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

SM4 密碼算法的密鑰長度為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計算公式如下:

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

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

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

SM4 算法的解密過程與加密過程基本一致， 只是輪密鑰使用順序不同. 加密輪密鑰使用順序為rk0，rk1，···，rk31， 解密輪密鑰是加密輪密鑰的逆序， 即rk31，rk30，···，rk0.

2.2 ECB 模式

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

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

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

(1) 簡單， 計算速度快;

(2) 支持并行計算;

(3) 運算錯誤不會被傳送， 無需知道明文長度.

ECB 工作模式缺點是:

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

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

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

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

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

2.3 CTR 模式

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

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

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

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

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

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

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

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

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

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

3 基于CTR 的GPU 并行SM4 算法

3.1 并行化策略

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

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

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

3.2 并行化實現(xiàn)

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

3.2.1 主機端初始化階段

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

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

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

(3) 生成一個隨機的計數(shù)器值counter;

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

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

(6) 為明文p、密文c、輪密鑰rk、計數(shù)器值counter 申請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)); //計數(shù)器值其中， size 代表明文的字節(jié)數(shù);

(7) 向GPU 發(fā)送明文數(shù)據(jù)text、輪密鑰數(shù)據(jù)rk、計數(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、計數(shù)器值counter:SM4_encrypt_ctr?＜blocks，threads?＞(p，c，rk，n，counter).

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

這一階段通過線程劃分、啟動進行加解密計算.

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

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

(3) 啟動線程， 并將線程所需要進行計算的數(shù)據(jù)存于該線程獨占的寄存器中: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ù)計算完成后， 之前為密文c申請的GPU 全局內(nèi)存已經(jīng)存儲了計算完成后的密文.

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

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

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

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

4 實驗及其分析

4.1 CUDA 編程對性能影響

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

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

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

(3) 線程塊大小一個線程塊內(nèi)要有足夠多的線程， 因為一個線程塊會占用一個SM，如果塊內(nèi)線程過少會導致SM內(nèi)多個Core 閑置.

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

CUDA 最大的特點就是其高度的編程自由性， 開發(fā)人員可以根據(jù)自己的需求對程序進行修改， 但也因為其自由性導致在CUDA 代碼編寫的時候要考慮多方因素， 任何一個微小的差別都可能對程序的執(zhí)行速度造成巨大的影響.

4.2 實驗環(huán)境

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

表1 實驗環(huán)境Table 1 Experiment environment

4.3 CTR 模式性能分析實驗結(jié)果

實驗主要通過CTR 模式下串行加解密和并行加解密時間的比值來衡量GPU 并行加速的效果， 其公式如下:

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

4.3.1 不同明文大小

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

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

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

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

4.3.2 不同線程塊大小

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

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

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

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

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

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

4.4 CTR 模式與ECB 模式對比實驗結(jié)果

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

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

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

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

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

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

5.1 SM4 不同運行方式對比

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

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

通過上述對比我們可以發(fā)現(xiàn)以下幾點:

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

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

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

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

5.2 對比結(jié)果分析

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

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

6 結(jié)束語

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

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