對于一個(gè)軟件開發(fā)人員，可能聽說過FPGA，甚至在大學(xué)課程設(shè)計(jì)中，可能拿FPGA做過計(jì)算機(jī)體系架構(gòu)相關(guān)的驗(yàn)證，但是對于它的第一印象可能覺得這是硬件工程師干的事兒。 ? 目前，隨著人工智能的興起，GPU 借助深度學(xué)習(xí)，走上了歷史的舞臺，并且正如火如荼的跑者各種各樣的業(yè)務(wù)，從 training 到 inference 都有它的身影。FPGA 也借著這股浪潮，慢慢地走向數(shù)據(jù)中心，發(fā)揮著它的優(yōu)勢。所以接下來就講講 FPGA 如何能讓程序員們更好友好的開發(fā)，而不需要寫那些煩人的 RTL 代碼，不需要使用 VCS，Modelsim 這樣的仿真軟件，就能輕輕松松實(shí)現(xiàn) unit test。 ? 實(shí)現(xiàn)這一編程思想的轉(zhuǎn)變，是因?yàn)?FPGA 借助 OpenCL 實(shí)現(xiàn)了編程，程序員只需要通過 C/C++ 添加適當(dāng)?shù)?pragma 就能實(shí)現(xiàn) FPGA 編程。為了讓您用 OpenCL 實(shí)現(xiàn)的 FPGA 應(yīng)用能夠有更高的性能，您需要熟悉如下介紹的硬件。另外，將會介紹編譯優(yōu)化選項(xiàng)，有助于將您的 OpenCL 應(yīng)用更好的實(shí)現(xiàn) RTL 的轉(zhuǎn)換和映射，并部署到 FPGA 上執(zhí)行。 ?

FPGA 概覽

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

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

對于 3-bits 輸入的 LUT 可以等價(jià)理解為一個(gè)擁有 3位地址線并且 8 個(gè) 1-bit 存儲單元的存儲器（一個(gè)8長度的數(shù)組，數(shù)組內(nèi)每個(gè)元素是 1bit）。那么當(dāng)需要實(shí)現(xiàn) 3-bits 數(shù)字按位與操作的時(shí)候，8長度數(shù)組存的是 3-bits 輸入數(shù)字的按位與結(jié)果，一共是 8 種可能性。當(dāng)需要實(shí)現(xiàn) 3-bits 按位異或的時(shí)候，8長度數(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)部，可通過 LUT 實(shí)現(xiàn)。 ?

? 如上展示了 3輸入，1輸出的 LUT 實(shí)現(xiàn)。當(dāng)將 LUT 并聯(lián)，串聯(lián)等方式結(jié)合起來后就可以實(shí)現(xiàn)更加復(fù)雜的邏輯運(yùn)算了。 ? 傳統(tǒng) FPGA 開發(fā) ? ▍傳統(tǒng) FPGA 與軟件開發(fā)對比 ? 對于傳統(tǒng)的 FPGA 開發(fā)與軟件開發(fā)，工具鏈可以通過下表簡單對比： ?

注：傳統(tǒng) FPGA 與軟件開發(fā)對比表 ? 重點(diǎn)介紹一下，編譯階段的 Synthesis (綜合)，這部分與軟件開發(fā)的編譯有較大的不同。一般的處理器 CPU、GPU等，都是已經(jīng)生產(chǎn)出來的 ASIC，有各自的指令集可以使用。但是對于 FPGA，一切都是空白，有的只是零部件，什么都沒有，但是可以自己創(chuàng)造任何結(jié)構(gòu)形式的電路，自由度非常的高。這種自由度是 FPGA 的優(yōu)勢，也是開發(fā)過程中的劣勢。 ? 寫到這里，讓我想起了最近 《神秘的程序員們》中的一個(gè)梗：

?

?

?

注：漫畫來源《神秘的程序員們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í)序電路，輸出邏輯。通過組合邏輯獲取下一個(gè)狀態(tài)是什么，時(shí)序邏輯用于存儲當(dāng)前狀態(tài)，輸出邏輯混合組合、時(shí)序電路，得到最終輸出結(jié)果。 ?

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

? 實(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

.|-- 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

__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)

// 聲明本地存儲，暫存數(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;

#?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

$?aocl?program??matrix_mult.aocx

$ ./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