對(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,一切都是空白,有的只是零部件,什么都沒(méi)有,但是可以自己創(chuàng)造任何結(jié)構(gòu)形式的電路,自由度非常的高。這種自由度是 FPGA 的優(yōu)勢(shì),也是開發(fā)過(guò)程中的劣勢(shì)。
寫到這里,讓我想起了最近 《神秘的程序員們》中的一個(gè)梗:
注:漫畫來(lái)源《神秘的程序員們 56》by 西喬
傳統(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è)用于異構(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 《your-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。—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 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 《your-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 《your-board》 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)
《your-board》 : 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)步。
責(zé)任編輯:gt
評(píng)論
查看更多