久久久久久精品免费无码,黄色视频在线观看www,任你躁任我躁

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

目前，隨著人工智能的興起，GPU 借助深度學(xué)習(xí)，走上了歷史的舞臺(tái)，并且正如火如荼的跑者各種各樣的業(yè)務(wù)，從 training 到 inference 都有它的身影。FPGA 也借著這股浪潮，慢慢地走向數(shù)據(jù)中心，發(fā)揮著它的優(yōu)勢(shì)。所以接下來(lái)就講講 FPGA 如何能讓程序員們更好友好的開(kāi)發(fā)，而不需要寫(xiě)那些煩人的 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ù)類(lèi)型的位寬都是固定的，相反 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 開(kāi)發(fā)

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

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

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

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

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

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

▍基于 OpenCL 的 FPGA 開(kāi)發(fā)

OpenCL 對(duì)于 FPGA 開(kāi)發(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)開(kā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ì)算單元并行展開(kāi)，形成乘法加法樹(shù)
 #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)行：

#Togeneratea.aocxfilefordebuggingthattargetsaspecificacceleratorboard$aoc-march=emulatordevice/matrix_mult.cl-obin/matrix_mult.aocx--fp-relaxed--fpc--no-interleavingdefault--board#GenerateHostexe.$make#Toruntheapplication$envCL_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

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

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

燒寫(xiě)的命令如下：

$aoclprogrammatrix_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é)

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