FPGA如何能讓程序員們更好友好的開發(fā)

簡介

對于一個軟件開發(fā)人員，可能聽說過 FPGA，甚至在大學(xué)課程設(shè)計中，可能拿FPGA做過計算機(jī)體系架構(gòu)相關(guān)的驗證，但是對于它的第一印象可能覺得這是硬件工程師干的事兒。

目前，隨著人工智能的興起，GPU 借助深度學(xué)習(xí)，走上了歷史的舞臺，并且正如火如荼的跑者各種各樣的業(yè)務(wù)，從 training 到 inference 都有它的身影。FPGA 也借著這股浪潮，慢慢地走向數(shù)據(jù)中心，發(fā)揮著它的優(yōu)勢。所以接下來就講講 FPGA 如何能讓程序員們更好友好的開發(fā)，而不需要寫那些煩人的 RTL 代碼，不需要使用 VCS，Modelsim 這樣的仿真軟件，就能輕輕松松實現(xiàn) unit test。

實現(xiàn)這一編程思想的轉(zhuǎn)變，是因為 FPGA 借助 OpenCL 實現(xiàn)了編程，程序員只需要通過 C/C++ 添加適當(dāng)?shù)膒ragma就能實現(xiàn) FPGA 編程。為了讓您用 OpenCL 實現(xiàn)的 FPGA 應(yīng)用能夠有更高的性能，您需要熟悉如下介紹的硬件。另外，將會介紹編譯優(yōu)化選項，有助于將您的 OpenCL 應(yīng)用更好的實現(xiàn) RTL 的轉(zhuǎn)換和映射，并部署到 FPGA 上執(zhí)行。

FPGA 概覽

FPGA 是高規(guī)格的集成電路，可以實現(xiàn)通過不斷的配置和拼接，達(dá)到無限精度的函數(shù)功能，因為它不像 CPU 或者 GPU 那樣，基本數(shù)據(jù)類型的位寬都是固定的，相反 FPGA 能夠做的非常靈活。在使用 FPGA 的過程中，特別適合一些 low-level 的操作，比如像 bit masking、shifting、addition 這樣的操作都可以非常容易的實現(xiàn)。

為了達(dá)到并行化計算，F(xiàn)PGA 內(nèi)部包含了查找表（LUTs），寄存器（register），片上存儲（on-chip memory）以及算術(shù)運算硬核（比如數(shù)字信號處理器 (DSP) 塊）。這些 FPGA 內(nèi)部的模塊通過網(wǎng)絡(luò)連接在一起，通過編程的手段，可以對連接進(jìn)行配置，從而實現(xiàn)特定的邏輯功能。這種網(wǎng)絡(luò)連接可重配的特性為 FPGA 提供了高層次可編程的能力。（FPGA的可編程性就體現(xiàn)在改變各個模塊和邏輯資源之間的連接方式）

舉個例子，查找表（LUTs）體現(xiàn)的 FPGA 可編程能力，對于程序猿來說，可以等價理解為一個存儲器（RAM）。對于 3-bits 輸入的 LUT 可以等價理解為一個擁有 3位地址線并且 8 個 1-bit 存儲單元的存儲器（一個8長度的數(shù)組，數(shù)組內(nèi)每個元素是 1bit）。那么當(dāng)需要實現(xiàn) 3-bits 數(shù)字按位與操作的時候，8長度數(shù)組存的是 3-bits 輸入數(shù)字的按位與結(jié)果，一共是 8 種可能性。當(dāng)需要實現(xiàn) 3-bits 按位異或的時候，8長度數(shù)組存的是 3-bits 輸入數(shù)字的按位異或結(jié)果，一共也是 8 種可能性。這樣，在一個時鐘周期內(nèi)，3-bits 的按位運算就能夠獲取到，并且實現(xiàn)不同功能的按位運算，完全是可編程的（等價于修改 RAM 內(nèi)的數(shù)值）。

3-bits 輸入 LUT 實現(xiàn)按位與（bit-wise AND）示例：

注： 3-bits 輸入 LUT 查找表

我們看到的三輸入的按位與操作，如下所示，在 FPGA 內(nèi)部，可通過 LUT 實現(xiàn)。

如上展示了 3輸入，1輸出的 LUT 實現(xiàn)。當(dāng)將 LUT 并聯(lián)，串聯(lián)等方式結(jié)合起來后就可以實現(xiàn)更加復(fù)雜的邏輯運算了。

傳統(tǒng) FPGA 開發(fā)

傳統(tǒng) FPGA 與軟件開發(fā)對比

對于傳統(tǒng)的 FPGA 開發(fā)與軟件開發(fā)，工具鏈可以通過下表簡單對比：

注：傳統(tǒng) FPGA 與軟件開發(fā)對比表

重點介紹一下，編譯階段的 Synthesis (綜合)，這部分與軟件開發(fā)的編譯有較大的不同。一般的處理器 CPU、GPU等，都是已經(jīng)生產(chǎn)出來的 ASIC，有各自的指令集可以使用。但是對于 FPGA，一切都是空白，有的只是零部件，什么都沒有，但是可以自己創(chuàng)造任何結(jié)構(gòu)形式的電路，自由度非常的高。這種自由度是 FPGA 的優(yōu)勢，也是開發(fā)過程中的劣勢。

寫到這里，讓我想起了最近 《神秘的程序員們》中的一個梗：

注：漫畫來源《神秘的程序員們56》by 西喬

傳統(tǒng)的FPGA開發(fā)就像10歲時候的 Linux，想吃一個蛋糕，需要自己從原材料開始加工。FPGA 正是這種狀態(tài)，想要實現(xiàn)一個算法，需要寫 RTL，需要設(shè)計狀態(tài)機(jī)，需要仿真正確性。

傳統(tǒng) FPGA 開發(fā)方式

復(fù)雜系統(tǒng)，需要使用有限狀態(tài)機(jī)（FSM），一般就需要設(shè)計下圖包含的三部分邏輯：組合電路，時序電路，輸出邏輯。通過組合邏輯獲取下一個狀態(tài)是什么，時序邏輯用于存儲當(dāng)前狀態(tài)，輸出邏輯混合組合、時序電路，得到最終輸出結(jié)果。

然后，針對具體算法，設(shè)計邏輯在狀態(tài)機(jī)中的流轉(zhuǎn)過程：

實現(xiàn)的 RTL 是這樣的：

module fsm_using_single_always (clock

, // clockreset

, // Active high, syn resetreq_0

, // Request 0req_1

, // Request 1gnt_0

, // Grant 0gnt_1 );//=============Input Ports=============================input

clock,reset,req_0,req_1; //=============Output Ports===========================output gnt_0,gnt_1;//=============Input ports Data Type===================wire

clock,reset,req_0,req_1;//=============Output Ports Data Type==================reg gnt_0,gnt_1;//=============Internal Constants======================parameter SIZE = 3

;parameter IDLE = 3'b001,GNT0 = 3'b010,GNT1 = 3'b100 ;//=============Internal Variables======================reg [SIZE-1:0]

state

;// Seq part of the FSMreg [SIZE-1:0]

next_state ;// combo part of FSM//==========Code startes Here==========================always @ (posedge clock)begin : FSMif (reset == 1'b1) begin

state <= #1 IDLE; ?gnt_0 <= 0; ?gnt_1 <= 0;end else case(state) ? IDLE : if (req_0 == 1'b1) begin ? ? ?

state <= #1 GNT0; ? ? ? ?

gnt_0 <= 1; ? ? ? ? ? ?

end else if (req_1 == 1'b1) begin

gnt_1 <= 1; ? ?

state <= #1 GNT1; ? ? ? ? ?

end else begin

state <= #1 IDLE; ? ? ?

end GNT0 : if (req_0 == 1'b1) begin

state <= #1 GNT0; ? ? ? ? ?

end else begin

gnt_0 <= 0; ? ? ? ? ?

state <= #1 IDLE; ? ? ?

end GNT1 : if (req_1 == 1'b1) begin

state <= #1 GNT1; ? ? ? ?

end else begin

gnt_1 <= 0; ? ? ? ? ? ?

state <= #1 IDLE; ? ? ?

end default : state <= #1 IDLE;endcaseendendmodule // End of Module arbiter

傳統(tǒng)的 RTL 設(shè)計，對于程序員簡直就是噩夢啊，夢啊，啊～～～工具鏈完全不同，開發(fā)思路完全不同，還要分析時序，一個 Clock 節(jié)拍不對，就要推翻重來，重新驗證，一切都顯得太底層，不是很方便。那么，這些就交給專業(yè)的 FPGAer 吧，下面介紹的 OpenCL 開發(fā) FPGA，有點像 25 歲的 Linux 了。有了高層次的抽象。用起來自然也會更加方便。

基于 OpenCL 的 FPGA 開發(fā)

OpenCL 對于 FPGA 開發(fā)，注入了新鮮的血液，一種面向異構(gòu)系統(tǒng)的編程語言，將 FPGA 最為異構(gòu)實現(xiàn)的一種可選設(shè)備。由 CPU Host 端控制整個程序的執(zhí)行流程，F(xiàn)PGA Device 端則作為異構(gòu)加速的一種方式。異構(gòu)架構(gòu)，有助于解放 CPU，將 CPU 不擅長的處理方式，下發(fā)到 Device 端處理。目前典型的異構(gòu) Device 有：GPU、Intel Phi、FPGA。

OpenCL 是個 what？

注：引用自 wiki

Open Computing Language(OpenCL) is a framework for writing programs that execute acrossheterogeneousplatforms consisting ofcentral processing units(CPUs),graphics processing units(GPUs),digital signal processors(DSPs),field-programmable gate arrays(FPGAs) and other processors orhardware accelerators. OpenCL specifies aprogramming language(based onC99) for programming these devices and application programming interfaces (APIs) to control the platform and execute programs on the compute devices. OpenCL provides a standard interface for parallel computing usingtask-basedanddata-based parallelism.

大意是說：OpenCL 是一個用于異構(gòu)平臺編程的框架，主要的異構(gòu)設(shè)備有 CPU、GPU、DSP、FPGA以及一些其它的硬件加速器。OpenCL 基于 C99 來開發(fā)設(shè)備端代碼，并且提供了相應(yīng)的 API 可以調(diào)用。OpenCL 提供了標(biāo)準(zhǔn)的并行計算的接口，以支持任務(wù)并行和數(shù)據(jù)并行的計算方式。

OpenCL 案例分析

這里采用 Altera 官網(wǎng)的矩陣乘法案例進(jìn)行分析?？梢酝ㄟ^如下鏈接下載案例：Altera OpenCL Matrix Multiplication

代碼結(jié)構(gòu)如下：

.|-- common| |-- inc| | `-- AOCLUtils| | |-- aocl_utils.h| | |-- opencl.h| | |-- options.h| | `-- scoped_ptrs.h| |-- readme.css| `-- src| `-- AOCLUtils| |-- opencl.cpp| `-- options.cpp`-- matrix_mult |-- Makefile |-- README.html |-- device | `-- matrix_mult.cl `-- host |-- inc | `-- matrixMult.h `-- src `-- main.cpp

其中，和 FPGA 相關(guān)的代碼是matrix_mult.cl，該部分代碼描述了 kernel 函數(shù)，這部分函數(shù)會通過編譯器生成 RTL 代碼，然后 map 到 FPGA 電路中。

kernel 函數(shù)的定義如下：

__kernel__attribute((reqd_work_group_size(BLOCK_SIZE,BLOCK_SIZE,1)))__attribute((num_simd_work_items(SIMD_WORK_ITEMS)))void matrixMult( __global float *restrict C, __global float *A, __global float *B, int A_width, int B_width)

模式比較固定，需要注意的是__global指明從 CPU 傳過來的數(shù)據(jù)，存放到全局內(nèi)存中，可以是 FPGA 片上存儲資源，DDR，QDR 等，這個視 FPGA 的 OpenCL BSP 驅(qū)動，會有所區(qū)別。num_simd_work_items用于指明 SIMD 的寬度。reqd_work_group_size指明了工作組的大小。這些概念，可以參考 OpenCL 的使用手冊。

函數(shù)實現(xiàn)如下：

// 聲明本地存儲，暫存數(shù)組的某一個 BLOCK__local float A_local[BLOCK_SIZE][BLOCK_SIZE];__local float B_local[BLOCK_SIZE][BLOCK_SIZE];// Block indexint block_x = get_group_id(0);int block_y = get_group_id(1);// Local ID index (offset within a block)int local_x = get_local_id(0);int local_y = get_local_id(1);// Compute loop boundsint a_start = A_width * BLOCK_SIZE * block_y;int a_end = a_start + A_width - 1;int b_start = BLOCK_SIZE * block_x;float running_sum = 0.0f;for (int a = a_start, b = b_start; a <= a_end; a += BLOCK_SIZE, b += (BLOCK_SIZE * B_width)){ ?// 從 global memory 讀取相應(yīng) BLOCK 數(shù)據(jù)到 local memory ?A_local[local_y][local_x] = A[a + A_width * local_y + local_x]; ?B_local[local_x][local_y] = B[b + B_width * local_y + local_x]; ?// Wait for the entire block to be loaded. ?barrier(CLK_LOCAL_MEM_FENCE); ?// 計算部分，將計算單元并行展開，形成乘法加法樹 ?#pragma unroll ?for (int k = 0; k < BLOCK_SIZE; ++k) ?{ ? ?running_sum += A_local[local_y][k] * B_local[local_x][k]; ?} ?// Wait for the block to be fully consumed before loading the next block. ?barrier(CLK_LOCAL_MEM_FENCE);}// Store result in matrix CC[get_global_id(1) * get_global_size(0) + get_global_id(0)] = running_sum;

采用 CPU 模擬仿真 FPGA

對其進(jìn)行仿真，不需要 programer 關(guān)心具體的時序是怎么走的，只需要驗證邏輯功能就可以，Altera OpenCL SDK 提供了 CPU 仿真 Device 設(shè)備的功能，采用如下方式進(jìn)行：

# To generate a .aocx file for debugging that targets a specific accelerator board$ aoc -march=emulator device/matrix_mult.cl -o bin/matrix_mult.aocx --fp-relaxed --fpc --no-interleaving default --board # Generate Host exe.$ make# To run the application$ env CL_CONTEXT_EMULATOR_DEVICE_ALTERA=8 ./bin/host -ah=512 -aw=512 -bw=512

上述腳本中，通過-march=emulator設(shè)置創(chuàng)建一個可用于 CPU debug 的設(shè)備可執(zhí)行文件。-g 添加調(diào)試 flag。—board用于創(chuàng)建適配該設(shè)備的 debugging 文件。CL_CONTEXT_EMULATOR_DEVICE_ALTERA為用于 CPU 仿真的設(shè)備數(shù)量。

當(dāng)執(zhí)行上述腳本后，輸出如下：

$ env CL_CONTEXT_EMULATOR_DEVICE_ALTERA=8 ./bin/host -ah=512 -aw=512 -bw=512Matrix sizes: A: 512 x 512 B: 512 x 512 C: 512 x 512Initializing OpenCLPlatform: Altera SDK for OpenCLUsing 8 device(s) EmulatorDevice : Emulated Device ... EmulatorDevice : Emulated DeviceUsing AOCX: matrix_mult.aocxGenerating input matricesLaunching for device 0 (global size: 512, 64)...Launching for device 7 (global size: 512, 64)Time: 5596.620 msKernel time (device 0): 5500.896 ms...Kernel time (device 7): 5137.931 msThroughput: 0.05 GFLOPSComputing reference outputVerifyingVerification: PASS

通過仿真時候設(shè)置 Device ＝ 8，模擬 8 個設(shè)備運行 (512, 512) * (512, 512) 規(guī)模的矩陣，最終驗證正確。接下來就可以將其真正編譯到 FPGA 設(shè)備上后運行。

FPGA 設(shè)備上運行矩陣乘

這個時候，真正要將代碼下載到 FPGA 上執(zhí)行了，這時候，只需要做一件事，那就是用 OpenCL SDK 提供的編譯器，將*.cl代碼適配到 FPGA 上，執(zhí)行編譯命令如下：

$ aoc device/matrix_mult.cl -o bin/matrix_mult.aocx --fp-relaxed --fpc --no-interleaving default --board

這個過程比較慢，一般需要幾個小時到10幾個小時，視 FPGA 上資源大小而定。（目前這部分時間太長暫時無法解決，因為這里的編譯，其實是在行程一個能夠正常工作的電路，軟件會進(jìn)行布局布線等工作）

等待編譯完成后，將生成的matrix_mult.aocx文件燒寫到 FPGA 上就 ok 啦。

燒寫的命令如下：

$ aocl program matrix_mult.aocx

這時候，大功告成，可以運行 host 端程序了：

$ ./host -ah=512 -aw=512 -bw=512Matrix sizes: A: 512 x 512 B: 512 x 512 C: 512 x 512Initializing OpenCLPlatform: Altera SDK for OpenCLUsing 1 device(s) : Altera OpenCL QPI FPGAUsing AOCX: matrix_mult.aocxGenerating input matricesLaunching for device 0 (global size: 512, 512)Time: 2.253 msKernel time (device 0): 2.191 msThroughput: 119.13 GFLOPSComputing reference outputVerifyingVerification: PASS

可以看到，矩陣乘法能夠在 FPGA 上正常運行，吞吐大概在 119GFlops 左右。

小結(jié)

從上述的開發(fā)流程，OpenCL 大大的解放了 FPGAer 的開發(fā)周期，并且對于軟件開發(fā)者，也比較容易上手。這是他的優(yōu)勢，但是目前開發(fā)過程中，還是存在一些問題，如：編譯器優(yōu)化不足，相比 RTL 寫的性能存在差距；編譯到 Device 端時間太長。不過這些隨著行業(yè)的發(fā)展，一定會慢慢的進(jìn)步。

另外，對 FPGA 感興趣，或者有用 FPGA 做方案的同學(xué)，歡迎一起探討。