哈希算法異構(gòu)可重構(gòu)高能效計(jì)算系統(tǒng)研究

鄭博文，聶 一，柴志雷

1.江南大學(xué)物聯(lián)網(wǎng)工程學(xué)院，江蘇無錫214122

2.江南大學(xué)人工智能與計(jì)算機(jī)學(xué)院，江蘇無錫214122

3.江蘇省模式識別與計(jì)算智能工程實(shí)驗(yàn)室，江蘇無錫214122

哈希算法一直在數(shù)字簽名和消息驗(yàn)證等領(lǐng)域具有重要的地位。哈希算法應(yīng)用場景日益多元化，這就要求哈希算法具有更高的計(jì)算性能和能效，純軟件的方式已難以滿足需求。越來越多的研究開始探索使用FPGA 以硬件方式突破計(jì)算性能瓶頸。通過FPGA 實(shí)現(xiàn)哈希算法，目前主要的優(yōu)化方法有預(yù)計(jì)算[1]、優(yōu)化關(guān)鍵路徑[2]和設(shè)計(jì)流水線結(jié)構(gòu)[3]。文獻(xiàn)[4] 對SHA-1 算法使用預(yù)計(jì)算和優(yōu)化關(guān)鍵路徑的優(yōu)化方法，設(shè)置四級流水線，提升了單位邏輯資源的利用效率。文獻(xiàn)[5-6] 通過對SHA-256 算法進(jìn)行預(yù)計(jì)算，縮短計(jì)算的關(guān)鍵路徑，并采用多周期劃分流水線以提升硬件工作頻率，加速SHA-256 算法計(jì)算。

在不同場景中使用多種哈希算法的組合也有明確的需求[7-8]。雖然對個(gè)別哈希算法進(jìn)行硬件優(yōu)化，可以滿足算法的性能需求，但是，有些應(yīng)用場景往往會涉及多個(gè)算法。因此，也有文獻(xiàn)嘗試設(shè)計(jì)多個(gè)哈希算法共用的硬件結(jié)構(gòu)。文獻(xiàn)[9] 設(shè)計(jì)了SHA-1/SHA256/SM3 算法的IP 復(fù)用電路，通過電路實(shí)現(xiàn)了算法中壓縮運(yùn)算的循環(huán)展開，并設(shè)計(jì)了流水線結(jié)構(gòu)。比單獨(dú)實(shí)現(xiàn)三種算法，節(jié)約了電路面積并提升了算法吞吐量。文獻(xiàn)[10] 在同一個(gè)硬件方案中實(shí)現(xiàn)了SHA3-224、SHA3-256、SHA3-384、SHA3-512 算法的優(yōu)化，在數(shù)據(jù)填充和壓縮計(jì)算中使用流水線設(shè)計(jì)，降低了關(guān)鍵路徑的延遲。雖然FPGA 硬件實(shí)現(xiàn)哈希算法和復(fù)用IP 的電路設(shè)計(jì)，提升算法性能的同時(shí)也具備了一定的通用性，但是已經(jīng)設(shè)計(jì)好的電路結(jié)構(gòu)無法擴(kuò)展，當(dāng)需要使用其他哈希算法時(shí)依舊需要重新設(shè)計(jì)。近年來，也有文獻(xiàn)使用OpenCL 以及HLS 等高層次綜合工具在FPGA 上實(shí)現(xiàn)了哈希算法，文獻(xiàn)[11-13] 分別實(shí)現(xiàn)了MD5、SHA-1 以及SHA-256 算法，但是效果比起其他文獻(xiàn)中Verilog 的設(shè)計(jì)還存在差距。

針對上述問題，本文提出了一種異構(gòu)且可重構(gòu)的哈希算法高能效計(jì)算系統(tǒng)，整體系統(tǒng)采用異構(gòu)方式，適合硬件加速的部分采用專用硬件的方式進(jìn)行加速；適合軟件計(jì)算的部分由CPU執(zhí)行。同時(shí)，硬件加速部分利用可重構(gòu)技術(shù)，針對不同算法通過重構(gòu)進(jìn)行有針對性的硬件重構(gòu)。從而獲得整體系統(tǒng)的最佳計(jì)算能效。通過在FPGA SoC 上實(shí)現(xiàn)該架構(gòu)，在可編程邏輯端實(shí)現(xiàn)哈希算法硬件加速，并利用OpenCL 實(shí)現(xiàn)異構(gòu)設(shè)備間的交互，控制哈希算法在FPGA 上的動態(tài)重構(gòu)。驗(yàn)證了上述方法的有效性。

1 哈希算法的結(jié)構(gòu)分析

1.1 哈希函數(shù)的一般結(jié)構(gòu)

哈希函數(shù)可以將無限定義域的輸入映射為有限定義域的哈希值，這種壓縮映射關(guān)系是通過壓縮函數(shù)實(shí)現(xiàn)的。Merkle-Damgard 結(jié)構(gòu)[14]是這種映射關(guān)系的典型代表。常用的哈希算法，如SHA-1、SHA-256 以及MD5 等算法均使用了此種結(jié)構(gòu)。Merkle-Damgard 結(jié)構(gòu)可以表示為

式中：H為哈希函數(shù)；M為哈希的輸入數(shù)據(jù)；CV 為鏈接變量；IV 為初始向量；f為壓縮函數(shù)。M可以被劃分為L個(gè)分組數(shù)據(jù)M0,M1,M2,···,ML-1，分別對應(yīng)L輪壓縮函數(shù)f的迭代計(jì)算。第i輪壓縮函數(shù)計(jì)算的輸入數(shù)據(jù)為分組數(shù)據(jù)Mi-1和上一次壓縮函數(shù)的計(jì)算結(jié)果CVi。哈希函數(shù)會設(shè)置初始向量IV，作為CV0用于首次壓縮函數(shù)的計(jì)算。當(dāng)壓縮函數(shù)計(jì)算完所有的分組數(shù)據(jù)，最后的計(jì)算結(jié)果CVL即為哈希值。哈希函數(shù)的結(jié)構(gòu)如圖1 所示。

圖1 哈希函數(shù)結(jié)構(gòu)Figure 1 Hash function structure

許多哈希算法都使用了Merkle-Damgard 結(jié)構(gòu)，如SHA-1、SHA-256 和MD5 等，表1 列舉了它們的一些結(jié)構(gòu)特點(diǎn)。這些哈希算法在結(jié)構(gòu)上具有一定的相似性，但是壓縮函數(shù)f中的邏輯運(yùn)算存在著很大的差別，比如SHA-256 算法的壓縮函數(shù)使用了64 輪相同的迭代計(jì)算，RIPEMD-160 算法的壓縮函數(shù)則包括了兩組各80 輪不完全相同的迭代計(jì)算。

表1 部分哈希算法對比Table 1 Comparison of some Hash algorithms

1.2 算法并行性分析

哈希算法的核心部分是壓縮函數(shù)f的計(jì)算，分析其并行性有助于針對性地進(jìn)行算法優(yōu)化。因篇幅有限，本文以較常用的SHA-256 算法[15]為例進(jìn)行分析。SHA-256 算法的壓縮函數(shù)f需要進(jìn)行64 輪相同的迭代計(jì)算，如圖2 所示。

圖2 SHA-256 算法的迭代計(jì)算Figure 2 SHA-256 algorithm single-step iterative calculation

壓縮函數(shù)f輸入CV 的長度為256 bits，可以拆分為8 個(gè)32 bits 的數(shù)據(jù)，分別對應(yīng)圖2中A-H。分組數(shù)據(jù)的長度為512 bits，被拆分為16 個(gè)消息字W0,W1,W2,···,W15，并且由式(2)～(6) 計(jì)算得到更多的消息字W16,W17,W18,···,W63，對應(yīng)迭代計(jì)算中的Wt。K0,K1,K2,···,K63為64 個(gè)常量關(guān)鍵字，對應(yīng)迭代計(jì)算中的Kt。迭代計(jì)算中的邏輯函數(shù)描述見式(7)～(10)。

壓縮函數(shù)f完成64 輪迭代計(jì)算后，計(jì)算得到的A-H加上最初輸入的CV，作為下一次壓縮迭代輸入的CV，直至所有計(jì)算完成后作為哈希值輸出。SHA-256 算法不僅每次壓縮函數(shù)f的計(jì)算存在數(shù)據(jù)依賴，壓縮函數(shù)f中的64 輪迭代計(jì)算也存在數(shù)據(jù)依賴，使得SHA-256 算法只能串行執(zhí)行，極大地限制了SHA-256 算法的性能。除了SHA-256 算法，其余常見的哈希算法也存在數(shù)據(jù)依賴，無法通過并行計(jì)算提升性能。Merkle-Damgard 結(jié)構(gòu)哈希算法的核心部分是壓縮函數(shù)，且不同算法的壓縮函數(shù)結(jié)構(gòu)相似。通過硬件加速壓縮函數(shù)的計(jì)算，設(shè)計(jì)一種通用的硬件優(yōu)化方案，有助于提升此類哈希算法的性能，在FPGA 上實(shí)現(xiàn)不同哈希算法的加速設(shè)計(jì)。

2 哈希算法加速器設(shè)計(jì)

2.1 算法硬件設(shè)計(jì)

Merkle-Damgard 結(jié)構(gòu)哈希算法中的數(shù)據(jù)依賴是硬件加速的瓶頸，在FPGA 中以流水線結(jié)構(gòu)實(shí)現(xiàn)哈希算法能夠提升數(shù)據(jù)依賴部分的計(jì)算性能，并且對數(shù)據(jù)帶寬的利用效率更高。哈希算法的壓縮函數(shù)由消息字?jǐn)U展和迭代計(jì)算兩個(gè)部分組成，兩者都存在著數(shù)據(jù)依賴關(guān)系，可以分別使用流水線結(jié)構(gòu)實(shí)現(xiàn)。在SHA-256 算法中，雖然消息字?jǐn)U展和迭代計(jì)算在各自的結(jié)構(gòu)中存在數(shù)據(jù)依賴，但是兩者之間的數(shù)據(jù)依賴僅存在于單步迭代計(jì)算中，第t輪迭代計(jì)算和消息字Wt之間存在數(shù)據(jù)依賴關(guān)系，與其他消息字則沒有。因此壓縮函數(shù)中的消息字?jǐn)U展和迭代計(jì)算可以通過并行的流水線結(jié)構(gòu)實(shí)現(xiàn)，圖3 展示了流水線結(jié)構(gòu)的串行執(zhí)行和并行執(zhí)行。

圖3 流水線執(zhí)行方式Figure 3 Pipeline execution mode

2.2 針對流水線結(jié)構(gòu)的優(yōu)化

流水線結(jié)構(gòu)對于算法的性能提升是以硬件資源的消耗為代價(jià)的，可以通過硬件資源的復(fù)用對流水線結(jié)構(gòu)進(jìn)行優(yōu)化。在SHA-256 算法中，壓縮函數(shù)共需要64 個(gè)32 bits 的消息字，占用了2 048 bits 的存儲空間，當(dāng)使用流水線結(jié)構(gòu)實(shí)現(xiàn)時(shí)，占用了大量的存儲資源。由式(6) 可知，構(gòu)造Wt時(shí)，依賴于Wt-2,Wt-7,Wt-15和Wt-16，并且最初的16 個(gè)消息字由分組數(shù)據(jù)拆分得到，因此通過16 個(gè)寄存器的復(fù)用實(shí)現(xiàn)消息字的存儲，根據(jù)迭代計(jì)算的輪數(shù)不斷更新，如圖4 所示。前15 輪迭代計(jì)算中不更新寄存器中的消息字，從第16 輪迭代計(jì)算開始，消息字?jǐn)U展和迭代計(jì)算同步執(zhí)行，新的消息字更新至之后不會被使用的消息字的寄存器中。

圖4 消息字?jǐn)U展的存儲優(yōu)化Figure 4 Memory optimization of extension

消息字?jǐn)U展和壓縮函數(shù)的流水線結(jié)構(gòu)，以并行執(zhí)行的方式共同組成了哈希算法的全流水線結(jié)構(gòu)。對于第i次壓縮函數(shù)計(jì)算，輸入數(shù)據(jù)包括鏈接變量CVi和分組數(shù)據(jù)Mi。壓縮函數(shù)中第j輪迭代計(jì)算的輸入數(shù)據(jù)是常量Kj，消息字Mj以及上一輪迭代計(jì)算的結(jié)果Aj-Hj，其中首輪迭代計(jì)算中的A0-H0就是鏈接變量CVi。當(dāng)?shù)?3 輪迭代計(jì)算完成后，計(jì)算結(jié)果A64-H64和鏈接變量CVi相加得到了鏈接變量CVi+1用于下一輪壓縮函數(shù)計(jì)算。當(dāng)所有分組數(shù)據(jù)都由壓縮函數(shù)計(jì)算完成后，最后的鏈接變量作為哈希值輸出。由式(2)～(10) 可知，消息字?jǐn)U展的邏輯運(yùn)算比起迭代計(jì)算更加簡單，計(jì)算的需要時(shí)鐘周期更少，保證了每輪迭代計(jì)算和消息字?jǐn)U展能夠在同一個(gè)時(shí)鐘周期啟動，實(shí)現(xiàn)了壓縮函數(shù)的全流水線計(jì)算，圖5 展示了SHA-256 算法的全流水線結(jié)構(gòu)。

2.3 查表法降低流水線延遲

查找表（look-up-table,LUT）是FPGA 上的一種基礎(chǔ)結(jié)構(gòu)，本質(zhì)上是一個(gè)RAM。消息字?jǐn)U展和迭代計(jì)算中的數(shù)據(jù)只和迭代輪數(shù)相關(guān)，各輪迭代計(jì)算中需要訪存的寄存器相對固定，使用查找表記錄各輪迭代計(jì)算中需要訪存的寄存器地址和操作，可以避免生成負(fù)載的仲裁邏輯判斷寄存器訪存，節(jié)約硬件資源并降低訪存效率，提升了全流水線結(jié)構(gòu)的時(shí)鐘頻率。圖6 展示了通過查表法實(shí)現(xiàn)的數(shù)據(jù)訪存。

圖6 查表法Figure 6 Look-up table method

寄存器中存儲了迭代計(jì)算的結(jié)果A-H以及消息字W，隨著迭代計(jì)算的進(jìn)行不斷更新。第i輪迭代計(jì)算執(zhí)行時(shí)，查找表中已經(jīng)記錄的需要讀寫的寄存器地址和順序，消息字?jǐn)U展和迭代計(jì)算部分根據(jù)記錄讀取邏輯運(yùn)算中需要使用的數(shù)據(jù)并寫入計(jì)算結(jié)果。

3 基于OpenCL 實(shí)現(xiàn)異構(gòu)計(jì)算系統(tǒng)

3.1 針對OpenCL 的數(shù)據(jù)存儲優(yōu)化

OpenCL(open computer language) 是一種面向異構(gòu)系統(tǒng)的計(jì)算標(biāo)準(zhǔn)，使用OpenCL 的異構(gòu)計(jì)算系統(tǒng)包括一個(gè)主機(jī)端（Host），一個(gè)或多個(gè)設(shè)備端（Device），主機(jī)端和設(shè)備端通過共享內(nèi)存的機(jī)制實(shí)現(xiàn)數(shù)據(jù)交互。計(jì)算任務(wù)以內(nèi)核（Kernel）的形式運(yùn)行于設(shè)備端，由主機(jī)端負(fù)責(zé)調(diào)度和管理。OpenCL 為了提升開發(fā)效率，抽象了4 種內(nèi)存模型：全局內(nèi)存（global memory）、常量內(nèi)存（constant memory）、局部內(nèi)存（local memory）和私有內(nèi)存（private memory）。各個(gè)內(nèi)存模型具有不同的讀取速度和存儲空間[16]。表2 展示了FPGA 上不同內(nèi)存模型的性能參數(shù)，常量內(nèi)存中為只讀數(shù)據(jù)，存儲在全局內(nèi)存中，當(dāng)內(nèi)核執(zhí)行時(shí)加載到FPGA 的片上緩存中。

表2 不同內(nèi)存類型的參數(shù)Table 2 Parameters of different memory types

OpenCL 在FPGA 上運(yùn)行時(shí)，各個(gè)內(nèi)存模型以不同的硬件資源實(shí)現(xiàn)。全局內(nèi)存以DDR內(nèi)存實(shí)現(xiàn)，存儲空間最大，作為主機(jī)端和設(shè)備端的共享內(nèi)存。局部內(nèi)存由FPGA 的塊內(nèi)存實(shí)現(xiàn)，存儲空間小于全局內(nèi)存，但是訪存速度更快。私有內(nèi)存由FPGA 的片上寄存器實(shí)現(xiàn)，訪存速度最快，存儲空間最小。哈希算法處理的計(jì)算任務(wù)通常為密集計(jì)算，合理使用OpenCL 的內(nèi)存模型十分重要。

哈希算法的輸入數(shù)據(jù)長度可變，由主機(jī)端程序按照一定規(guī)則填充后，劃分為固定長度的分組數(shù)據(jù)M存儲在全局內(nèi)存中，由內(nèi)核程序計(jì)算。當(dāng)FPGA 以固定數(shù)據(jù)位寬進(jìn)行計(jì)算時(shí)，可以減少處理數(shù)據(jù)位寬的資源開銷，提升電路的工作頻率。依據(jù)哈希算法中各個(gè)數(shù)據(jù)的存儲需求與訪存頻率，選擇不同的內(nèi)存模型實(shí)現(xiàn)。表3 展示了對SHA-256 算法的數(shù)據(jù)位寬調(diào)整和存儲優(yōu)化。

表3 SHA-256 算法存儲資源Table 3 SHA-256 algorithm storage resources

其中，M使主機(jī)端程序填充后的512 bits 分組數(shù)據(jù)，由于數(shù)據(jù)量大，所以在全局內(nèi)存中存儲。IV 是初始向量，K是關(guān)鍵字，兩者均為常量且會被頻繁讀取，使用常量內(nèi)存進(jìn)行實(shí)現(xiàn)，提升數(shù)據(jù)的訪問速度。W是消息字，CV 是鏈接變量，在算法計(jì)算過程中需求的存儲空間較大，使用局部內(nèi)存進(jìn)行實(shí)現(xiàn)可以獲得較低的訪存延遲。H是哈希值，和M一樣存儲在全局內(nèi)存中，便于在主機(jī)端和設(shè)備端之間進(jìn)行交互。

3.2 多線程管理

命令隊(duì)列對于OpenCL 的使用是不可或缺的。主機(jī)端程序通過命令隊(duì)列可以自由地操作各個(gè)OpenCL 對象，如內(nèi)存對象的讀寫、內(nèi)核對象的執(zhí)行等。當(dāng)計(jì)算系統(tǒng)中需要執(zhí)行多組計(jì)算任務(wù)，或者不同哈希算法的計(jì)算任務(wù)，大量的命令隊(duì)列就會使得系統(tǒng)的運(yùn)行十分混亂。因此主機(jī)端程通過多線程來管理命令隊(duì)列，以及通過命令隊(duì)列提交不同的哈希算法內(nèi)核到FPGA執(zhí)行，圖7 展示了多線程管理命令隊(duì)列的流程。

圖7 多線程管理命令隊(duì)列Figure 7 Multithreaded management command queue

主機(jī)端程序根據(jù)計(jì)算任務(wù)啟動線程，由線程創(chuàng)建一系列的命令隊(duì)列，負(fù)載OpenCL 執(zhí)行時(shí)的內(nèi)存操作以及內(nèi)核執(zhí)行。并且線程還會對計(jì)算需要的輸入數(shù)據(jù)進(jìn)行預(yù)處理，存儲于數(shù)據(jù)緩沖區(qū)中，等待內(nèi)核執(zhí)行時(shí)傳輸至FPGA 上的全局內(nèi)存。主機(jī)端程序在創(chuàng)建數(shù)據(jù)緩沖區(qū)時(shí)，會根據(jù)線程數(shù)量對緩沖區(qū)進(jìn)行內(nèi)存劃分，避免多線程同時(shí)訪存引起的沖突。內(nèi)核執(zhí)行時(shí)，需要提交命令隊(duì)列，當(dāng)上一個(gè)計(jì)算任務(wù)還未執(zhí)行完畢時(shí)，命令隊(duì)列會排隊(duì)等待直到隊(duì)列為空。各個(gè)線程間在命令隊(duì)列等待時(shí)實(shí)現(xiàn)同步。線程在對計(jì)算數(shù)據(jù)進(jìn)行預(yù)處理時(shí)，主要對數(shù)據(jù)按照算法規(guī)則進(jìn)行填充分組。內(nèi)核在FPGA 的全局內(nèi)存中獲取的數(shù)據(jù)可以直接被用于壓縮迭代計(jì)算，使得內(nèi)核的執(zhí)行效率更高，發(fā)揮了異構(gòu)計(jì)算的優(yōu)勢。

3.3 FPGA 動態(tài)可重構(gòu)

通過FPGA 可重構(gòu)可以實(shí)現(xiàn)計(jì)算系統(tǒng)對不同哈希算法的加速，主機(jī)端程序可以靈活部署不同的哈希算法到FPGA。在多線程管理中，主機(jī)端程序根據(jù)使用的哈希算法創(chuàng)建線程，當(dāng)計(jì)算任務(wù)需要使用不同的算法時(shí)，會啟動新的線程并創(chuàng)建命令隊(duì)列。主機(jī)端程序通過OpenCL的API 函數(shù)clEnqueueNDRangeKernel 根據(jù)命令隊(duì)列的內(nèi)容提交內(nèi)核到FPGA 執(zhí)行，并且在這一過程中實(shí)現(xiàn)對FPGA 的可重構(gòu)。實(shí)現(xiàn)哈希算法的FPGA 鏡像文件存儲在主機(jī)端的DDR 內(nèi)存中，通過PCI Express 總線配置到FPGA 上[18]。

3.4 哈希算法高能效計(jì)算系統(tǒng)

計(jì)算系統(tǒng)由主機(jī)端和設(shè)備端兩個(gè)部分組成，選擇CPU 作為主機(jī)端，F(xiàn)PGA 作為設(shè)備端。主機(jī)端程序負(fù)載OpenCL 運(yùn)行環(huán)境的配置，并且通過多線程管理實(shí)現(xiàn)不同哈希算法在設(shè)備端的加速。FPGA 根據(jù)主機(jī)端程序的命令執(zhí)行配置的哈希算法硬件電路。CPU 和FPGA 之間通過PCI Express 總線進(jìn)行數(shù)據(jù)交互，圖8 展示了計(jì)算系統(tǒng)的總體架構(gòu)。計(jì)算系統(tǒng)在運(yùn)行時(shí)，由主機(jī)端程序?qū)penCL 平臺進(jìn)行初始化配置，包括查詢OpenCL 平臺和設(shè)備的信息。對需要執(zhí)行的計(jì)算任務(wù)通過多線程管理，預(yù)處理計(jì)算數(shù)據(jù)，等待FPGA 執(zhí)行對應(yīng)的哈希算法硬件電路。FPGA 執(zhí)行內(nèi)核時(shí)，將預(yù)處理后的計(jì)算數(shù)據(jù)由全局內(nèi)存搬運(yùn)至局部內(nèi)存以全流水線結(jié)構(gòu)執(zhí)行。執(zhí)行完畢后得到的哈希值會由局部內(nèi)存搬運(yùn)至全局內(nèi)存，通過PCI Express 總線傳輸回主機(jī)端的內(nèi)存中存儲，并在之后由外部程序使用。FPGA 中豐富的邏輯資源支持實(shí)例化多個(gè)全流水線結(jié)構(gòu)，通過并行計(jì)算提升性能。在內(nèi)核中實(shí)例化的硬件電路被OpenCL 稱為計(jì)算單元（compute unit,CU）。

圖8 計(jì)算系統(tǒng)整體架構(gòu)Figure 8 Computing system architecture

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

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

本文選取Intel Xeon E5-2650 V2 CPU 和Intel Stratix 10 GX2800 FPGA 搭建OpenCL異構(gòu)計(jì)算平臺，其中FPGA 的邏輯資源ALMs 為933120，寄存器資源為3732480。選取Intel Core I7-10700 CPU 作為CPU 計(jì)算平臺，NVIDIA GTX 1650 SUPER 作為GPU 計(jì)算平臺進(jìn)行對比驗(yàn)證。

4.2 測試結(jié)果及分析

本文為了驗(yàn)證計(jì)算系統(tǒng)對不同哈希算法的性能提升，在計(jì)算系統(tǒng)中實(shí)現(xiàn)了多個(gè)哈希算法的內(nèi)核設(shè)計(jì)，內(nèi)核中設(shè)置單個(gè)計(jì)算單元時(shí)算法的性能測試如表4 所示，計(jì)算性能的單位是百萬哈希/秒。

表4 不同哈希算法的性能Table 4 Performance of different Hash algorithms

由測試結(jié)果可知，在算法硬件實(shí)現(xiàn)時(shí)，計(jì)算邏輯越復(fù)雜占用的邏輯資源就越多。雖然5 種哈希算法的時(shí)鐘頻率都在370 MHz 左右，理論上以全流水結(jié)構(gòu)計(jì)算時(shí)每個(gè)時(shí)鐘周期可以完成一次哈希計(jì)算，但是算法性能和理論值存在差距。測試用時(shí)使用OpenCL 的API 函數(shù)在主機(jī)端程序獲取，包括了局部內(nèi)存與全局內(nèi)存、FPGA 和CPU 之間的數(shù)據(jù)通信時(shí)間，可能導(dǎo)致了計(jì)算性能的差距。5 種哈希算法的哈希值長度不同，其中MD5 算法的哈希值長度最短為128 bits，計(jì)算性能最高，SHA-512 算法的哈希值長度最長為512 bits 并且數(shù)據(jù)吞吐量達(dá)到了帶寬限制，所以實(shí)際計(jì)算性能與理論值差距較大，之后會在4.3 節(jié)中進(jìn)行驗(yàn)證。

SHA-256 算法是應(yīng)用最廣泛的哈希算法之一，為了進(jìn)一步驗(yàn)證本文設(shè)計(jì)哈希算法加速模塊的性能，也與相關(guān)文獻(xiàn)中的SHA-256 算法性能進(jìn)行對比。文獻(xiàn)中通常使用吞吐量表示算法性能，本文通過式(11) 將SHA-256 算法性能轉(zhuǎn)換為吞吐量進(jìn)行對比。

式中：block_size 為消息分組長度，SHA-256 算法的block_size 為512 bits；num 為哈希算法的運(yùn)行次數(shù)；time 為Kernel 執(zhí)行時(shí)間，實(shí)驗(yàn)結(jié)果如表5 所示。

表5 不同文獻(xiàn)SHA-256 對比Table 5 SHA-256 comparison of different documents

由SHA-256 算法的性能對比可知，本文的流水線設(shè)計(jì)可以使用較高的時(shí)鐘頻率，并且對算法進(jìn)行優(yōu)化，結(jié)合數(shù)據(jù)位寬調(diào)整、數(shù)據(jù)存儲優(yōu)化等方法降低每級流水線的延遲，對SHA-256算法的吞吐量提升明顯。上述文獻(xiàn)中的設(shè)計(jì)使用Verilog 進(jìn)行開發(fā)，耗時(shí)較長，而本文基于OpenCL 實(shí)現(xiàn)的哈希算法加速模塊設(shè)計(jì)，可以快速應(yīng)用于其他Merkle-Damgard 結(jié)構(gòu)哈希算法，取得不錯(cuò)的性能提升，極大地降低了哈希算法硬件設(shè)計(jì)的時(shí)間成本。不但保證了不同哈希算法計(jì)算性能，還可以通過異構(gòu)計(jì)算的方式實(shí)現(xiàn)FPGA 的動態(tài)重構(gòu)，根據(jù)不同的計(jì)算任務(wù)靈活組合哈希算法，為解決多哈希算法組合計(jì)算的需求提供了一種思路。

4.3 計(jì)算系統(tǒng)性能表現(xiàn)

加解密是哈希算法的主要應(yīng)用領(lǐng)域。哈希值具有不可逆性，只能通過不斷嘗試密碼組合，生成哈希值后與待破解哈希值對比，因此對哈希算法的計(jì)算性能要求很高。為了測試計(jì)算系統(tǒng)在應(yīng)用場景中的表現(xiàn)，本文將哈希計(jì)算系統(tǒng)應(yīng)用于哈希值的破解，在FPGA 上部署多個(gè)計(jì)算單元CU，測試暴力破解下各哈希算法的性能表現(xiàn)。Hashcat 自稱是目前世界上最快的密碼恢復(fù)程序，支持多種哈希算法。為了對比驗(yàn)證本文計(jì)算系統(tǒng)的性能，選擇CPU 和GPU 分別運(yùn)行Hashcat 進(jìn)行對照，實(shí)驗(yàn)結(jié)果如圖9 所示。

圖9 不同計(jì)算平臺加解密對比Figure 9 Comparison of encryption and decryption of different computing platforms

本文實(shí)現(xiàn)了5 種哈希算法的加速模塊設(shè)計(jì)，計(jì)算性能明顯高于CPU(I7-10700)、GPU(GTX 1650 SUPER)，但是FPGA 的功耗明顯低于GPU。性能提升最明顯的SHA-256 算法，F(xiàn)PGA的計(jì)算性能是GPU 的2 倍，能效比是5.6 倍。在加解密時(shí)，將待破解哈希值傳輸?shù)紽PGA，在FPGA 上設(shè)計(jì)對比驗(yàn)證模塊，降低了數(shù)據(jù)在全局內(nèi)存?zhèn)鬏數(shù)难舆t，使各個(gè)計(jì)算單元的計(jì)算性能和工作頻率持平，證明了哈希算法能夠以全流水線結(jié)構(gòu)執(zhí)行，每個(gè)時(shí)鐘周期可以完成一次哈希計(jì)算。表6 中展示了不同哈希算法內(nèi)核能夠設(shè)置的計(jì)算單元數(shù)量，以及每個(gè)計(jì)算單元的性能。

表6 加解密各計(jì)算單元性能Table 6 Computing performance of each computing unit for encryption and decryption

4.4 計(jì)算系統(tǒng)的可重構(gòu)機(jī)制

為了進(jìn)一步驗(yàn)證計(jì)算系統(tǒng)的可重構(gòu)機(jī)制的性能，本文設(shè)計(jì)了包含不同哈希算法的計(jì)算任務(wù)，分別在本文設(shè)計(jì)的計(jì)算系統(tǒng)中和CPU 計(jì)算平臺上進(jìn)行測試并對比結(jié)果，CPU 使用的測試軟件使用開源軟件庫OpenSSL 中開源的哈希算法代碼實(shí)現(xiàn)，保證測試結(jié)果的可靠性。選擇本文實(shí)現(xiàn)的5 種哈希算法進(jìn)行測試，測試內(nèi)容為連續(xù)執(zhí)行這五種算法各10 億次，測試結(jié)果如圖10 所示，其中FPGA 的測試結(jié)果包括了可重構(gòu)的時(shí)間。

圖10 不同哈希算法的計(jì)算時(shí)間Figure 10 Calculation time of different hash algorithms

本文在使用計(jì)算系統(tǒng)進(jìn)行性能測試時(shí)，對FPGA 硬件重構(gòu)的時(shí)間也進(jìn)行了測試。5 種哈希算法重構(gòu)的時(shí)間基本保持在180～200 ms 之間。從測試結(jié)果可以發(fā)現(xiàn)，對于5 種哈希算法，CPU 計(jì)算的時(shí)間都明顯大于FPGA。雖然FPGA 在重構(gòu)時(shí)需要花費(fèi)額外的時(shí)間，但是比起對于算法性能的提升，這明顯是值得的。

5 結(jié)語

本文針對哈希算法在實(shí)際應(yīng)用中對高速計(jì)算和多算法組合的需求，提出了一種異構(gòu)且可重構(gòu)的計(jì)算系統(tǒng)，由軟件和硬件協(xié)同處理，發(fā)揮異構(gòu)計(jì)算優(yōu)勢。硬件加速設(shè)計(jì)使用可重構(gòu)技術(shù)，提升系統(tǒng)整體計(jì)算能效。在FPGA 上實(shí)現(xiàn)計(jì)算系統(tǒng)，并應(yīng)用于加解密，實(shí)驗(yàn)結(jié)果表明對MD5、SHA-1、SHA-256、SHA-512、RIPEMD-160 等算法性能提升明顯，相較于CPU和GPU 計(jì)算平臺，性能分別提升了8.3～18.7 倍和1～2 倍，能效分別提升了15.5～34.7 倍和2.9～5.7 倍。雖然本文僅實(shí)現(xiàn)了Merkle-Damgard 結(jié)構(gòu)的5 種哈希算法，但是對于其他同結(jié)構(gòu)的哈希算法也可以推測具有性能提升。本文設(shè)計(jì)的計(jì)算系統(tǒng)，由于使用異構(gòu)計(jì)算實(shí)現(xiàn)，可以便捷地應(yīng)用于不同地場景中。除了本文已經(jīng)實(shí)現(xiàn)的加解密計(jì)算，哈希算法在數(shù)字貨幣、區(qū)塊鏈[25]以及數(shù)據(jù)驗(yàn)證等領(lǐng)域也有很廣泛的應(yīng)用空間，具有較高的實(shí)用價(jià)值。