使用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相關的向后分支(即,循環中發生的分支),因為這會降低性能。
1、通過SDK的channel 或pipe來傳輸數據
為了提高kernel之間的數據傳輸效率,在kernel程序中使用channel通道拓展。 如果想利用通道功能,又想使用其他SDK運行kernel,則使用OpenCL pipes。
有時,FPGA到global memory全局存儲器帶寬會限制內核之間的數據傳輸效率。 理論上FPGA到global memory全局存儲器的最大帶寬根據目標定制平台和板上可用的全局存儲器bank的數量而變化。 要確定主板的理論最大帶寬,要參考主板的文檔。
實際上,kernel無法實現最大可用全局內存帶寬的100%利用率。 利用率級別取決於算法的訪問模式。
如果全局內存帶寬是我們使用OpenCL內核的性能限制,首先嘗試將算法分解為多個較小的kernel。 其次,通過在內核之間實現SDK的channel或OpenCL的pipe進行數據傳輸來消除一些全局內存訪問。
(1) Channel與pipe的特性
a. Default Behavior
Channel默認行為是阻塞的,而pipe的默認行為是非阻塞的(nonblocking)。
b. 多個OpenCL內核的並發執行
可以同時執行多個OpenCL內核。 要啟用並發執行,要修改主機代碼以實例化多個命令隊列。 每個同時執行的kernel內核都與一個單獨的命令隊列關聯。
pipe的特別注意事項:Intel SDK 中OpenCL的pipe是允許在其他的OpenCL SDK上兼容的,但不能最大化kernel內核吞吐量。OpenCL 2.0中要求在進行pipe讀取前先進性pipe寫入,以免kernel在空pipe中讀取數據,因此kernel無法同時運行。由於Intel SDK支持並發執行,可以修改主機應用程序以及kernel程序來實現並發執行,從而提高吞吐量。但不能將kernel移植到其他SDK上。
要啟用並發執行包含pipe的內核,需要將內核代碼中的depth屬性替換為blocking屬性(即__attribute __((blocking)))。 blocking屬性在read_pipe和write_pipe函數調用時引入blocking行為。 調用點將阻止內核執行,直到管道的另一端准備好為止。
如果同時將blocking屬性和depth屬性添加到內核,則當管道為空時,read_pipe僅調用一個塊,而當管道為滿時,write_pipe僅調用一個塊。 blocking行為會導致內核之間的隱式同步,從而使得內核之間互鎖。
c. 隱式內核Kernel同步
通過blocking channel調用以及blocking pipe調用來隱式同步kernel。

channel int c0; __kernel void producer (__global int * in_buf) { for (int i = 0; i < 10; i++) { write_channel_intel (c0, in_buf[i]); } } __kernel void consumer (__global int * ret_buf) { for (int i = 0; i < 10; i++) { ret_buf[i] = read_channel_intel(c0); } }

__kernel void producer (__global int * in_buf, write_only pipe int __attribute__ ((blocking)) c0) { for (int i = 0; i < 10; i++) { write_pipe (c0, &in_buf[i]); } } __kernel void consumer (__global int * ret_buf, read_only pipe int __attribute__ ((blocking)) c0) { for (int i = 0; i < 10; i++) { int x; read_pipe (c0, &x); ret_buf[i] = x; } }
可以同步內核,以便在每次循環迭代期間,producer kernel寫入數據,consumer kernel讀取數據。如果producer中的write_channel_intel 或 write_pipe並沒有寫入數據,consumer將阻塞並在read_channel_intel或read_pipe調用處等待直到producer發送有效數據為止,反之亦然。d. 跨調用的數據持久性
在調用write_channel_intel將數據寫入channel或調用write_pipe將數據寫入pipe之后,數據在work-groups和NDRange調用之間保持不變。work-item寫入channel或pipe的數據將保留在該channel或pipe中,直到從另一個work-item讀取它為止。此外,channel或pipe中的數據順序等效於對該channel或pipe的寫操作順序,且該順序與執行寫操作的work-item無關。
例如,如果多個work-item嘗試同時訪問一個channel或pipe,則只有一個工作項可以訪問它。調用write_channel_intel或調用write_pipe分別將被稱為DATAX的特定的work-item數據寫入channel或pipe。同樣,訪問channel或pipe的第一個work-item將從中讀取DATAX。讀寫操作的這種順序順序使通道和管道成為內核之間共享數據的有效方法。
e. 強制的work-item順序
SDK強制work-item的順序與channel或pipe的讀寫操作保持一致性。
(2) Channel或pipe的執行順序
kernel程序中的每個channel或pipe的調用都會轉換為FPGA pipeline中執行的指令。如果通過pipeline執行有效的work-item,則會調用channel或pipe。但即使channel或pipe調用之間沒有控制或數據依賴性,它們的執行也可能無法在kernel的pipeline中實現最優的並行性。
(3) 優化channel或pipe的緩沖區推斷
除了手動添加緩沖的channel或pipe以外,Intel SDK的離線編譯器盡可能通過調整緩沖區大小來提高kernel的吞吐量。在編譯期間,離線編譯器計算交互的channel或pipe之間的調度不匹配,這些不匹配可能會導致讀寫操作之間的不平衡。離線編譯器自動執行緩沖區推斷優化來糾正不平衡性。
(4) Best Practices for Channel&Pipe
使用單線程kernel而不是多線程kernel。
考慮怎樣用前饋數據來表示設計模型,例如,back-to-back循環或離散處理步驟。確定是否應該將設計拆分為通過channel連接的多個kernel。
只有當kernel上的同一點使用整個數據時,才在channel上將數據聚合。
嘗試使每個kernel保持合理的channel數。
如果使用等待數據的循環結構,不要使用非阻塞(non-blocking)的channel或pipe。non-blocking的channel要比blocking的channel消耗更多的資源。
2、展開循環
使用 #pragma unroll 展開循環,注意盡量避免嵌套循環。
可以控制離線編譯器將kernel裝換位硬件資源的方式。如果kernel中包含循環迭代,可以通過展開循環來提高性能。循環的展開減少了離線編譯器執行的迭代次數,但代價是硬件資源消耗的增加。
如果有比較充分的硬件資源,直接在主循環中添加#progma unroll來展開循環。循環的展開會顯著地改變離線編譯器創建的計算單元的結構。
__kernel void example ( __global const int * restrict x,
__global int * restrict sum ) { int accum = 0; #pragma unroll for (size_t i = 0; i < 4; i++) { accum += x[i + get_global_id(0) * 4]; } sum[get_global_id(0)] = accum; }
#pragma unroll指令使離線編譯器完全展開循環的四個迭代。 為了完成展開,離線編譯器通過將加法運算數量增加三倍並加載四倍的數據來擴展pipeline。 移除循環后,計算單元將采用前饋結構。 結果,計算單元可以在初始加載操作和加法完成之后的每個時鍾周期存儲和的計算結果。 離線編譯器通過合並四個加載操作來進一步優化此kernel,以便計算單元可以加載所有必需的輸入數據從而在一個加載操作中計算出結果。
注意:
不要使用嵌套循環,盡可能地添加#progma unroll指令,實現大的單循環或展開內部循環。離線編譯器不一定能輕易地展開循環,如果循環嵌套層數較多較復雜,會導致編譯時間很長。
展開循環並合並全局內存中的加載操作,可使內核的硬件實現在每個時鍾周期執行更多操作。 通常,用於提高OpenCL的kernel性能的方法應達到以下結果:
增加並行操作數
增加實現的內存帶寬
增加內核可以在硬件中執行的每個時鍾周期的操作數
而在以下情況下,離線編譯器可能無法完全展開循環:
想要完全展開具有非常大數量迭代的存在數據依賴性的循環,kernel的硬件實現可能不適用於FPGA。
完全展開循環邊界不是常數的循環。
循環由負責的控制flow組成,例如包含復雜數組索引或退出條件(在編譯時未知)的循環。
要在這些情況下啟用循環展開,指定#pragma unroll <N>指令,其中<N>是展開因子。 展開因子限制離線編譯器展開的迭代次數。 例如,要阻止kernel中的循環展開,將指令#pragma unroll 1添加到該循環中。
3、優化浮點操作
對於浮點操作,可以手動指導離線編譯器進行優化,以在硬件中創建更有效的pipeline結構並減少總體硬件的使用率。這些優化可能會導致浮點操作的結果產生細小差別。
Tree Balancing
離線編譯器不會自動將程序優化成樹平衡結構運算,因為這會導致結果與真實結果有些許差別。如果不在意浮點結果的細微差別,並希望離線編譯器使用平衡樹結構優化浮點運算,則需要在aoc命令中包括 -fp-relaxed 選項:
aoc -fp-relaxed <kernel_filename>.cl
Rounding Operations
浮點運算的平衡樹結構包含多個舍入運算,這些舍入操作在某些應用程序中可能需要大量的硬件資源,離線編譯器不會自動減少舍入運算的數量。如果可以接受浮點運算結果的細小差異,可以使用aoc命令中的 -fpc 減少實現浮點運算所需的硬件數量:
aoc -fpc <kernel_filename>.cl
-fpc 選項指示離線編譯器執行以下任務:
盡可能的刪除浮點舍入運算及轉換。如果可能,離線編譯器僅在浮點運算樹的末尾舍入一次浮點運算。
攜帶其他尾數位以保持精度。離線編譯器在浮點計算過程中攜帶其他精度位,並在浮點運算樹的末尾刪除這些精度位。
這種類型的優化會導致硬件執行融合的浮點運算,這也是許多新硬件處理系統的功能。融合多個浮點運算可以最大程度地減少舍入步驟,從而獲得更准確的結果。
(1)浮點vs定點表示
FPGA包含大量用於實現浮點運算的邏輯,但只要有可能,就可以使用數據的定點表示來增加可用的硬件資源量。實現定點運算所需的硬件資源通常要小於等效地浮點運算。所以與等效的浮點運算相比,可以在FPGA中容納更多的定點運算。
OpenCL的標准不支持定點表示,必須使用整數數據類型實現。硬件開發者通常使用定點數據以節省硬件,並且僅保留執行計算所需的數據分辨率。必須使用8\16\32\64位標量數據類型,因為OpenCL標准支持這些數據分辨率。然而,可以在源代碼中包含適量的屏蔽操作,以便硬件編譯工具可以執行優化來節省硬件資源。
例如,算法使用17位數據的定點表示,那必須使用32位數據類型存儲該值。如果使用Intel SDK的離線比那一起將兩個17位的定點值加在一起,離線編譯器必須創建額外的硬件來處理多余的高15位。為了避免使用額外的硬件,可以使用靜態位掩碼來指示硬件編譯工具在硬件編譯期間忽略不必要的位。
4、分配對齊內存
再分配用於與FPGA進行數據傳輸的主機端存儲器時,該存儲器必須至少與64字節對齊。對齊主機端存儲器可以實現直接進出FPG的直接存儲器訪問(DMA)傳輸,並提高緩沖區傳輸效率。
注意:根據主機端內存的使用方式,Intel建議分配更嚴格的對齊方式。例如,如果使用 CL_MEM_USE_HOST_PTR 標志將分配的內存用於創建緩沖區,那么該內存也應該正確對齊以用於訪問內核中的緩沖區的數據類型。要設定對齊的內存分配,需要將以下代碼添加到HOST主機程序中:
對於Windows:
#define AOCL_ALIGNMENT 64 #include<malloc.h> void *ptr = _aligned_malloc(size, AOCL_ALIGNMENT);
對於Linux:
#define AOCL_ALIGNMENT 64 #include<stdlib.h> void *ptr = NULL; posix_memalign(&ptr, AOCL_ALIGNMENT, size);
如果要釋放內存,則使用 free(ptr) 。
5、對齊帶有或不帶有填充的結構
確保數據的正確對齊。要確保數據結構的四字節對齊,小於四字節的結構對其會導致硬件變大與變慢。
例如:
typedef struct{ char r, g, b, alpha; } __attribute__((aligned(4))) Pixel;
代碼將Pixel結構強制進行4字節對齊。
為了防止離線編譯器自動為struct加入padding填充,可以使用packed屬性,例如:
struct __attribute__((packed)) Mystruct { char a; int b; int c; }
這樣Mystruct結構大小是16位字節而不是12字節,要比沒有packed屬性的12字節更高效。
而如果同時使用aligned 與 packed, 如下例所示。
struct __attribute__((packed)) __attribute__((aligned(16))) Mystruct { char a; int b; int c; }
Mystruct結構大小是9位字節。而由於aligned(16),所以struct是以16字節的方式對齊存儲的,而且沒有padding,所以kernel數據訪問是高效的。
6、保持向量元素的相似結構
如果更新了向量類型的一個元素,那么更新這個向量的所有元素。
__kernel void update (__global const float4 * restrict in, __global const float4 * restrict out) { size_t gid = get_global_id(0); out[gid].x = process(in[gid].x); out[gid].y = process(in[gid].y); out[gid].z = process(in[gid].z); out[gid].w = 0; //Update w even if that variable is not required. }
7、避免指針混淆
盡量在指針參數中加入restrict。
8、避免代價高昂的函數/功能
在FPGA中代價高昂的函數包括:
整數除法以及取模(取余數)運算;
除了加、乘、絕對值與比較之外,絕大多數的浮點運算;
Atomic Function原子函數。
相反,代價較低的功能包括:
二進制的邏輯運算,AND/NAND/OR/XOR/NOR/XNOR等;
具有常數參數的邏輯運算;
常數位的位移;
整數乘或除以一個2的冪次常數。
如果一個代價高昂的函數為work-group中的每個work-item都創建了一個新數據,那么付出高昂代價是值得的,可以通過一次代價高昂的計算將計算結果傳遞到多個work-item中。
9、避免依賴於work-item id的后向分支
避免在kernel中包含任何與work-item id相關的后向分支(即循環中發生的分支),否則會降低性能。