1、首先了解Intel FPGA SDK for OpenCL實現OpenCL的設計組件，包括:
kernels, global memory interconnect, local memory, loops 以及channels
(1) Kernels
Loops一般是Kernel優化的重點，尤其是nested loops。
OpenCL系統中每個kernel是通過一系列block表示的。Block主要由三部分：輸入或循環輸入節點，一組指令以及一個分支節點。若block中沒有分支則沒有輸入與分支節點，輸入或循環輸入節點根據分支在block中的起始位置決定變量的初始值。block的余下部分包含可停頓以及不可停頓指令以及clusters。優化好的模型應該包含最少數量的可停頓指令，例如I/O或內存訪問。
在block中不可停頓的指令被分組成多個clusters，以減小可停頓指令的握手開銷。cluster包含入口/出口節點，且cluster是不可停頓的。可以在cluster的出口節點找到出口FIFO信息，在這種情況下，分支節點會通知下一個要跳轉到的block。
只有當Kernel中不使用get_global_id()以及get_local_id()這種內置的work-item時，SDK才會將其編譯為single work-item kernel，否則將其編譯為NDRange kernel。
而offline compiler不能將NDRange中的loops進行pineline，但這些loops可以同時接受多個work-item。一個kernel可能包括多個循環，且每個都有嵌套循環。若將每個外部loop的nested loop的迭代總數制表，kernel的吞吐量一般會因此降低。而要有效率地執行NDRange Kernel，通常需要巨大數量的線程。
(2) Global memory interconnect
為讀寫訪問最大化內存帶寬的能力對高性能計算非常重要。OpenCL中存在用於讀寫全局存儲器的各種類型的模塊，稱為負載儲存單元(load-store units, LSUs)。
與GPU不同，FPGA可以構建最適合應用程序編譯內存訪問模式的任意自定義的LSU。選擇最理想的LSU類型會為應用程序顯著地提高設計性能。
(3) Local memory
Local Memory比較復雜，不同於GPU結構具有不同級別緩存的架構，FPGA在內部使用專用的內存塊實現local memory。
Local Memory有以下特性：
對local memory的每次讀寫訪問都要映射到一個端口；
可以將local memory中的內容划分為一個或多個存儲體bank，以使每個存儲體包含該local memory的數據子集。
一個存儲體由一個或多個副本組成。Bank中的副本與其他的副本包含相同的數據。創建副本是為了有效的支持local memory的多次訪問。每個副本都具有一個寫端口與讀端口，從而可以同時訪問。如果local memory是double dumped，那么每個副本有四個物理端口，可以最多有三個讀端口。
在kernel代碼中，用 local 類型來聲明local memory：
local int lmem[1024];
Intel SDK能夠自己設定local memory的width, deph, banks, replication, private copies的數量, interconnect等。離線編譯器可以分析訪問模式然后優化local memory最小化訪問競爭。
實現Kernel高效工作的關鍵是不停歇地訪問內存。離線編譯器始終嘗試為Kernel中所有的local memory找到不停歇訪問的配置方法，但由於kernel比較復雜，離線編譯器可能沒有足夠的信息推斷內存訪問是否有沖突，那么就需要local interconnect仲裁寄存器來仲裁內存訪問，會降低性能。
Local Memory Banks and Private Copis
默認local memory的存儲體只在最小維度上起作用，多個存儲體允許同時寫入。在下面的例子中，循環中每個local memory的訪問都有單獨的的地址。離線編譯器可以推斷出訪問模式，從而創建四個單獨的存儲體bank，四個獨立bank允許對4個lmem同時訪問，從而實現了無停頓的程序配置。此外，離線編譯器為lmem創建了兩個private copies，從而允許兩個work-groups同時pipline運行。
#define BANK_SIZE 4 __attribute__(reqd_work_group_size(8, 1, 1)) kernel void bank_arb_consecutive_multidim(global int* restrict in, global int* restrict out) { local int lmem[1024][BANK_SIZE]; int gi = get_global_id(0); int gs = get_global_size(0); int li = get_local_id(0); int ls = get_local_size(0); int res = in[gi]; #progma unroll for(int i = 0; i < BANK_SIZE; i++) { lmem[((li + i ) & 0x7f)][i] = res + i; res = res >> 1; } int rdata = 0; barrier(CLK_GLOBAL_MEM_FENCE); #progma unroll for(int i = 0; i < BANK_SIZE; i++) { rdata ^= lmem[(li + i) & 0x7f][i]; } out[gi] = rdata; return; }
如果私有副本顯著增加了設計area，考慮減少kernel中barrier的數量或增加max_work_group_size的值來減少SDK推斷出的private copies的數量。
可以使用 __attribute__((numbanks(N)) 指定bank的數量。
若不希望在最小維度上進行存儲，使用 bank_bits 指定存儲字節。而通過使用 bank_bits 能夠將memory中的數據分開存儲在多個bank中，同時可以指定使用哪些地址位選擇bank。下例中，使用第七位和第八位而不是最低的兩個維度來進行bank存儲。
#define BANK_SIZE 4 kernel void bank_arb_consecutive_multidim_origin (global int* restrict in, global int* restrict out) { local int a[BANK_SIZE][128] __attribute__((bank_bits(8,7),bankwidth(4))); int gi = get_global_id(0); int li = get_local_id(0); int res = in[gi]; #pragma unroll for (int i = 0; i < BANK_SIZE; i++) { a[i][((li+i) & 0x7f)] = res + i; res = res >> 1; } int rdata = 0; barrier(CLK_GLOBAL_MEM_FENCE); #pragma unroll for (int i = 0; i < BANK_SIZE; i++) { rdata ^= a[i][((li+i) & 0x7f)]; } out[gi] = rdata; return; }
Local Memory Replication
為了實現無停頓的程序配置，編譯器可能復制一個local memory系統來增加可讀端口的數量。對local memory的每個存儲操作都是在每個副本上同時執行的，梅格副本都具有相同相同的數據，可以獨立讀取。這樣增加了local memory系統可支持的同時讀取操作的數量。
Double Pumping
通常來說，每個local memory 副本具有兩個物理端口，而double pumping是的每個副本最多支持四個物理端口。
優點：可以增加可使用的物理端口；可能會通過減少復制操作從而減少RAM的使用。
缺點：與single pumping相比邏輯更復雜、延遲更高；有可能會降低kernel時鍾頻率
可以使用 __attribute __（（singlepump）） 和 __attribute __（（doublepump）） 來控制本地內存系統的pump配置。
以下代碼示例說明了具有三個讀取端口和三個寫入端口的本地內存的實現。 脫機編譯器啟用了double pumping功能，並復制了本地內存三次，以實現無停頓內存配置。
#define BANK_SIZE 4 kernel void bank_arb_consecutive_multidim_origin (global int* restrict in, global int* restrict out) { local int a[BANK_SIZE][128]; int gi = get_global_id(0); int li = get_local_id(0); int res = in[gi]; #pragma unroll 1 for (int i = 0; i < BANK_SIZE; i++) { a[i][li+i] = res + i; a[gi][li+i] = res + i; a[gi+i][li] = res + i; res = res >> 1; } int rdata = 0; barrier(CLK_GLOBAL_MEM_FENCE); #pragma unroll 1 for (int i = 0; i < BANK_SIZE; i++) { rdata ^= a[i][li+i]; rdata += a[gi+i][li+i]; rdata += a[gi][li]; } out[gi] = rdata; return; }
(4) Nested Loop
Intel® FPGA SDK for OpenCL的離線編譯器由於循環迭代無法infer pipeline的執行，會導致外循環的迭代相對於隨后的內循環可能是亂序的，因為對於不同的外循環迭代，內循環的迭代次數可能會有所不同。要解決無序的外循環迭代問題，需要設計上下限在外循環迭代之間不變的內循環。
通過使用 loop_coalesce 減少嵌套循環消耗的面積。當循環嵌套層數在三層以上時，要消耗更多的面積area，利用loop coalescing可以減小延遲，減小area。
(5) Loops in a Single Work-item Kernel
SDK離線編譯器通過pipeline循環來優化單work-item 的kernel的性能。
單個work-item的kernel中循環的數據路徑可以包含多個正在運行的迭代。 此行為與NDRange內核中的循環不同，NDRange內核的循環包含正在運行的多個工作項（而不是循環迭代）。 在最佳pipeline循環中，每個時鍾周期都會啟動一個新的循環迭代。 每個時鍾周期啟動一次循環迭代可最大化流水線效率並產生最佳性能。
一次循環迭代與下一次之間的時鍾周期數稱為循環的啟動間隔（initiation interval, II）。 最佳流水線循環的II值為1，因為每個時鍾周期都會啟動一個新的循環迭代。而SDK可能無法流水線化內核中的每個循環。 如果未對循環進行流水線處理，則直到上一次迭代執行完畢，循環迭代才能開始。 在這種情況下，一次在循環的數據路徑中只有一個循環迭代處於活動狀態。
在II啟動間隔與最大頻率之間trade-off權衡
有些情況下，離線編譯器會以減小最大頻率為代價使II啟動間隔達到1。
循環依賴會影響循環的II啟動間隔
一些情況下，loop是pipelined，但II啟動間隔不是1。通常是loop的數據依賴或內存依賴引起的。
數據依賴是指循環迭代使用依賴於先前迭代的變量的情況。內存依賴是指在循環迭代中的內存訪問無法完成，直到完成來自先前循環迭代的內存訪問為止的情況。
循環推測
循環推測是一種優化技術，它通過允許在確定循環是否已經退出之前啟動將來的迭代來啟用更有效的循環pipeline。
循環融合
循環融合是一種編譯器轉換，其中兩個相鄰循環在相同索引范圍內合並為一個循環。 此轉換通常用於減少循環開銷並提高運行時性能。
循環融合的條件：
能夠融合的循環必須相鄰；
循環必須具有相同的迭代計數；
每個循環必須具有一個入口和一個出口。 例如，不考慮包含break語句的循環進行融合。
循環必須沒有負距離依賴關系。
(6) Channels
Intel® FPGA SDK for OpenCL™的通道實現提供了數據從一個kernel傳遞到另一個kernel的靈活方法，以提高性能。在kernel的代碼中添加channel的關鍵字，例如：
channel long16 myCh __attribute__((depth(16)));
(7) Load-Store Units
SDK會生成許多不同類型的負載存儲單元（LSU）。 對於某些類型的LSU，編譯器可能會根據內存訪問模式和其他內存屬性來修改LSU行為和屬性。
LSU類型(Best Practice Guide-3.7.1)：
- Burst-Coalesced Load-Store Units
- Prefetching Load-Store Units
- Pipelined Load-Store Units
- Constant-Pipelined Load-Store Units
- Atomic-Pipelined Load-Store Units
LSU修飾符：
Cached
Write-Acknowledge(write-ack)
Nonaligned
Never-Stall
LSU的控制：
離線編譯器允許通過一組內置調用來控制為全局存儲器訪問而生成的LSU的類型，可以使用這些調用加載和存儲到全局存儲器。
Load Built-ins
Load Built-in Arguments
Store Built-ins
Store Built-in Arguments
Example:
kernel void oclTest(global int * restrict in, global int * restrict out) { int i = get_global_id(0); int a1 = __pipelined_load(in + 3*i+0); // Uses a pipelined LSU // Uses a burst-coalesced LSU with a cache of size 1024 bytes int a2 = __burst_coalesced_cached_load(&in[3*i+1], 1024); int a3 = __prefetching_load(&in[3*i+2]); // Uses a prefetching LSU __burst_coalesced_store(&out[3*i+0], a3); // Uses a burst-coalesced LSU }
Note:編譯器不允許選擇可能在請求上下文中導致功能上不正確的結果的LSU；Prefeching LSU在Intel®Stratix®10設備上不可用。
對於LSU使用時的選擇：
可以根據對加載/存儲站點的訪問模式的了解或根據硅面積要求來決定LSU的類型。 以下是LSU按其面積需求的升序排列：
Pipelined LSU (load/store):區域有效，但可能比其他LSU慢。 如果受區域限制或訪問方式不一定是連續的，則應使用此LSU。
Prefetching LSU (only for loads): 也是區域有效的，但對於完全連續的訪問模式來說是完美的。 將其用於非連續訪問模式會導致吞吐量下降，因此，僅當知道訪問的地址嚴格連續時，才使用它。
Burst-coalesced LSU (load/store): 需求面積很大，但可以非常有效地處理連續的訪問模式。 檢查訪問模式是否連續要付出代價。 如果可能，這種LSU動態地嘗試將多個內核請求組合為一個跨越多個內存字的大burst。
Burst-coalesced cached LSU (only for loads): 最消耗面積，因為它包含LSU本地的額外緩存。 如果打算多次讀取內存中的同一位置，尤其是在多個ND范圍線程中，則使用這種LSU可以提高吞吐量。
2、OpenCL Kernel Design Best Practices設計的最佳做法
使用Intel® FPGA SDK for OpenCL™ 離線編譯器，不需要調整kernel代碼便可以將其最佳的適應於固定的硬件設備，而是離線編譯器會根據kernel的要求自適應調整硬件的結構。
通常來說，應該先優化針對單個計算單元的kernel，之后累哦通過增加計算單元數量來拓展硬件以填充FPGA其余的部分，從而提升性能。Kernel的使用面積與硬件編譯所需要的時間有關，因此為了避免硬件編譯時間過長，首先要專注於優化kernel在單個計算單元上的性能。
要優化kernel的性能，主要包括數據處理以及內存訪問優化。
a. 通過SDK的channel 或pipe來傳輸數據。為了提高kernel之間的數據傳輸效率，在kernel程序中使用channel通道拓展。 如果想利用通道功能，又想使用其他SDK運行kernel，則使用OpenCL pipes。
b. 展開循環。
c. 優化浮點運算。對於浮點操作，可以手動引導SDK的離線編譯器進行優化，從而在硬件中創建更有效的pipeline結構並減少總體硬件使用率。
d. 分配對齊的內存。再分配與FPGA進行數據傳輸的主機端存儲器時，存儲器至少是64字節對齊的。
e. 使用或不用Padding來對齊結構。
f. 保持向量元素的相似結構。如果更新了向量的一個元素，那么更新這個向量的所有元素。
g. 避免指針混淆。盡量在指針參數中插入strict關鍵字。
h. 避免開銷大的函數/功能。有些函數在FPGA中實現開銷很大，可能會減低kernel的性能，或是需要大量硬件來實現。
i. 避免依賴於work-item id的后向分支。避免在kernel中包括任何與工作項ID相關的向后分支（即，循環中發生的分支），因為這會降低性能。