簡(jiǎn)介

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

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

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

FPGA 概覽

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

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

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

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

注：3-bits 輸入 LUT 查找表

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

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

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

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

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

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

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

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

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

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

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

實(shí)現(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è)計(jì)，對(duì)于程序員簡(jiǎn)直就是噩夢(mèng)啊，夢(mèng)啊，啊～～～工具鏈完全不同，開發(fā)思路完全不同，還要分析時(shí)序，一個(gè) Clock 節(jié)拍不對(duì)，就要推翻重來(lái)，重新驗(yàn)證，一切都顯得太底層，不是很方便。那么，這些就交給專業(yè)的 FPGAer 吧，下面介紹的 OpenCL 開發(fā) FPGA，有點(diǎn)像 25 歲的 Linux 了。有了高層次的抽象。用起來(lái)自然也會(huì)更加方便。

基于 OpenCL 的 FPGA 開發(fā)

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

OpenCL 是個(gè) what？

注：引用自 wiki

Open Computing Language?(OpenCL) is a framework for writing programs that execute across?heterogeneousplatforms consisting of?central processing units?(CPUs),?graphics processing units?(GPUs),?digital signal processors(DSPs),?field-programmable gate arrays?(FPGAs) and other processors or?hardware accelerators. OpenCL specifies aprogramming language?(based on?C99) 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 using?task-based?and?data-based parallelism.

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

OpenCL 案例分析

這里采用 Altera 官網(wǎng)的矩陣乘法案例進(jìn)行分析?？梢酝ㄟ^(guò)如下鏈接下載案例：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ù)會(huì)通過(guò)編譯器生成 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 傳過(guò)來(lái)的數(shù)據(jù)，存放到全局內(nèi)存中，可以是 FPGA 片上存儲(chǔ)資源，DDR，QDR 等，這個(gè)視 FPGA 的 OpenCL BSP 驅(qū)動(dòng)，會(huì)有所區(qū)別。num_simd_work_items 用于指明 SIMD 的寬度。reqd_work_group_size 指明了工作組的大小。這些概念，可以參考 OpenCL 的使用手冊(cè)。

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

// 聲明本地存儲(chǔ)，暫存數(shù)組的某一個(gè) 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); ?// 計(jì)算部分，將計(jì)算單元并行展開，形成乘法加法樹
?#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

對(duì)其進(jìn)行仿真，不需要 programer 關(guān)心具體的時(shí)序是怎么走的，只需要驗(yà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

上述腳本中，通過(guò)?-march=emulator 設(shè)置創(chuàng)建一個(gè)可用于 CPU debug 的設(shè)備可執(zhí)行文件。-g 添加調(diào)試 flag?！猙oard 用于創(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 OpenCL
Platform: Altera SDK for OpenCL
Using 8 device(s)
?EmulatorDevice : Emulated Device
?...
?EmulatorDevice : Emulated Device
Using AOCX: matrix_mult.aocx
Generating input matrices
Launching for device 0 (global size: 512, 64)
...
Launching for device 7 (global size: 512, 64)

Time: 5596.620 ms
Kernel time (device 0): 5500.896 ms
...
Kernel time (device 7): 5137.931 ms

Throughput: 0.05 GFLOPS

Computing reference output
Verifying
Verification: PASS

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

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

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

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

這個(gè)過(guò)程比較慢，一般需要幾個(gè)小時(shí)到10幾個(gè)小時(shí)，視 FPGA 上資源大小而定。（目前這部分時(shí)間太長(zhǎng)暫時(shí)無(wú)法解決，因?yàn)檫@里的編譯，其實(shí)是在行程一個(gè)能夠正常工作的電路，軟件會(huì)進(jìn)行布局布線等工作）

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

燒寫的命令如下：

$ aocl program matrix_mult.aocx

這時(shí)候，大功告成，可以運(yùn)行 host 端程序了：

$ ./host -ah=512 -aw=512 -bw=512Matrix sizes:
?A: 512 x 512
?B: 512 x 512
?C: 512 x 512Initializing OpenCL
Platform: Altera SDK for OpenCL
Using 1 device(s)
? : Altera OpenCL QPI FPGA
Using AOCX: matrix_mult.aocx
Generating input matrices
Launching for device 0 (global size: 512, 512)

Time: 2.253 ms
Kernel time (device 0): 2.191 ms

Throughput: 119.13 GFLOPS

Computing reference output
Verifying
Verification: PASS

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

小結(jié)

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

　　審核編輯：湯梓紅?