2017年12月31日 星期日

Divergence 與 Convolution/Filtering 的近似在 NEON 的處理

在 2017 年的最後一天想用技術分享做一個結束

Divergence


會特別想提 NEON divergence handling 的原因是許多的 NEON 教學並沒有特別著墨這部份, 儘管並不困難, 但是在眾多 instruction 找出適當的指令也不是容易的事.

對於多數 SIMD instruction 而言 divergence (if-else, switch-cases)都是不容易處理的部份
在許多 modern SIMD ISA 的設計中都採用了 predication 的作法: 可以透過 per-lane flag 數值來個別控制每個 lane 是否執行該指令, 輸出結果.
然而 NEON 並沒有 predication 的設計, 因此對於 if-else 的作法採用的是對於結果做 selection 的方式, 而其中扮演關鍵角色的是 VBSL 這個指令
考量下列範例

unsigned char pix;
...

if(pix >= 192){
     pix += 10;
}else{
     pix +=5;
}
改以 NEON 實作則如下
uint8x16_t vpix;
...
//selection mask
uint8x16_t vsel = vcgeq_u8(vpix, vdupq_n_u8(192));
//for >= 192
vpix1 = vaddq_u8(vpix, vdupq_n_u8(10));
//for < 192
vpix2 = vaddq_u8(vpix, vdupq_n_u8(5));
//get correspond
16-bit Multiply-Accumulation ing result from each lane
vpix = vbslq_u8(vsel, vpix1, vpix2);

事實上 vbsl 的用途不僅如此因為他是 bit-selection, 可以處理 bit-wise operation

16-bit Multiply-Accumulation 

 

這篇第二個要分享的是 fixed point 的技巧, 適合 convolution 與 image filtering 的數值近似.
在 image filtering 與 NN 的 convolution/FC 中的計算, 有許多數值介於 -1.0 ~ 1.0 的浮點數與 整數相乘的處理, 一方面 floating point 本身為 32b 運算(需要 ARMv8.2 才有 NEON FP16 的支援), 再者 float 與 int16/32 的轉換也需要消化額外的指令, 另一方面轉為 int16/32 的 fixed point 處理需要使用更多位元數的 int32/64.

對於許多這類應用有許多乘加的運算, 16bit 的數值相乘需要 32bit 來存放, 然而對於輸出通常也都會到 16bit (也就是最後結果需要 right shift 16 bit), 考量這類應用可以有些微的誤差, 可以考慮一個特別的指令 - vqdmulh or vqrdmulh

vqdmulh 的輸出結果是兩個 16b 數值相乘的 "兩倍" 作 right shift 16 bit
vqrdmulh 相較 vqdmulh 會多做一個 rounding 的動作

對於浮點的處理可以先轉為 16b 的數值 (類似 1/65536 for unsigned or 1/32768 for signed), 而相乘 x2 的結果主要是做 0.5 的近似, 可以取得更好的累加近似結果(最後結果需要 right shift 1bit).  對於一些能事先轉為 16b 資料固定的 pattern (像是 filtering 與 NN 的 weight, 或是除法), 這兩個指令相當實用.
out = 0.14 * pix0 + 0.7 * pix1 + 0.16 * pix2;
轉為 NEON 可以計算近似為:
// 0.14 * pix0
vout = vqdmulh_s16(vpix0, vdupq_n_s16(4588));
// 0.7 * pix1

vout = vaddq_s16(vout, vqdmulh_u16(vpix1, vdupq_n_u16(22934)));
// 0.16 * pix2
vout = vaddq_s16(vout, vqdmulh_u16(vpix2, vdupq_n_u16(5243)));

2017年11月3日 星期五

Lookup Table 在 NEON 中的處理

在 SIMD Programming 中由於希望能夠每個 lane 有一致的行為,
因此有一些事情是不容易達到的
而 Lookup Table (LUT) 即是其中之一

但若是特定條件之下, 還是有可能透過 NEON 加速
而這個直接前提是 8bit LUT (當然藉此方法可以做 8bit 轉 16bit LUT)
在 ARMv7 中 NEON 就提供兩個指令 - VTBL 與 VTBX 作為有限度的 LUT功能使用
兩者提供了同時能夠處理 8筆 8bit 資料的 LUT 動作 (64bit NEON register)所使用的 table 最多為 4個 64bit NEON 暫存器
這也就是說 index 範圍為 < 8x4 = 32
而 VTBL 與 VTBX 兩者的差異是:
VTBL 對於 index 超出範圍的 lane 輸出值為 0
VTBX 對於 index 超出範圍的 lane 並不更動

值得慶幸的是在 ARMv8 中將這兩者寬度加倍
也就是一次可以查找 16筆 8bit 資料的 LUT 動作 ( 128bit NEON register)
table 也變為 4個 128bit NEON register
因此 index 合法範圍為 < 16x4 =64
以下就只針對重點部分演示:

以下的 code 即為常見的 LUT 方式
#define TEST_W 4096
#define TEST_H 3072
unsigned char *table = (unsigned char*)malloc(256); 
unsigned char *buf_0 = (unsigned char*)malloc(TEST_W*TEST_H); 
unsigned char *buf_1 = (unsigned char*)malloc(TEST_W*TEST_H);  
// buffer filled with random numbers
...
for(int i = 0; i < TEST_W*TEST_H; i++){
    buf_0[i] = table[buf_0[i]];

以下為 ARMv8 的實作方式

uint8x16x4_t vlut_tbl[4]; 
//table initialization
for(int i = 0; i < 4; i++){
    vlut_tbl[i].val[0] = vld1q_u8(table+i*64);
    vlut_tbl[i].val[1] = vld1q_u8(table+i*64+16);
    vlut_tbl[i].val[2] = vld1q_u8(table+i*64+32);
    vlut_tbl[i].val[3] = vld1q_u8(table+i*64+48);

//SIMD LUT
for(int i = 0; i < TEST_W*TEST_H; i+=16){ 
    uint8x16_t vlut_idx = vld1q_u8(buf_1 + i);
    uint8x16_t vlut_val = vqtbl4q_u8(vlut_tbl[0], vlut_idx); 
    vlut_val = vqtbx4q_u8(vlut_val, vlut_tbl[1], vsubq_u8(vlut_idx, vdupq_n_u8(64)));
    vlut_val = vqtbx4q_u8(vlut_val, vlut_tbl[2], vsubq_u8(vlut_idx, vdupq_n_u8(128)));
    vlut_val = vqtbx4q_u8(vlut_val, vlut_tbl[3], vsubq_u8(vlut_idx, vdupq_n_u8(192)));
    vst1q_u8(buf_1 + i, vlut_val);
}

透過手邊的 Xperia M4 (Snapdragon 615, Cortex-A53 octa cores)平台測試
C-LUT : 58271 us, NEON-LUT : 21702 us

得到了 2.7X 的效能


2017年11月2日 星期四

Android NN API 與 OpenVX

乍看之下會覺得很奇怪
一個是 Neutral Network API
另一個則是 CV 的 API
兩者如何能夠相提並論地比較?
但是這兩者其實有著很高的相似度

讓我們先從 Android Neural Network API 的流程圖看起
可以看到 Android NN API 有幾個元件組成
1. ANeuralNetworkModels
2. Operations
3. ANeuralNetworksMemory

而 Android NN API 分為3個步驟
1. Network Create
2. Network Compilation
3. Network Execution

接著來看 OpenVX, 下圖來自於 Khronos 於 2016 的官方 Tutorial T3
可以看到 OpenVX 有幾個元件組成
1. Graph
2. Node
3. Image

而 OpenVX 分為3個步驟
1. Graph Create
2. Graph Verify
3. Graph Process

這裡可以再對照 Android NN API 提供的 Graph 示意圖
是不是很類似的概念呢?

兩者流程與介面的對應上可以說是如出一轍, 而在 OpenVX Neural Network Extension 中, 做了一件事就是新增了 Tensor 資料封裝型別與定義了 Neural Network 中用到的 operations, 並且沿用既有的 Graph 模式

2017年10月21日 星期六

AlphaGo Zero 雜感




本文是 Facebook 上本人的貼文
==========
AlphaGo Zero 自學已成的強度帶給你多少的感想?
現今棋王柯潔 感嘆:"一个纯净、纯粹自我学习的alphago是最强的……对于alphago的自我进步来讲……人类太多余了。"
有人半開玩笑的說天網將至, 也有人說這是對人類知識的當頭棒喝
身為工程師的角色, 亦獻醜地統整一下最近對 AlphaGo Zero 的看法:
在 DeepMind 一路自以往的 AI 建立 AlphaGo 的研究, 而從歷史性的 AlphaGo 與李世乭對弈與勝出, 到 Master 勝過當今棋王柯潔, 最後經由 AlphaGo/Master 推進到即便棋王亦望之興嘆的 AlphaGo Zero, 這幾年來一步步的演進, 實則展示了建立任務型 NN 的開發該有的演進過程, 以知識的角度來看, 與其說人類多餘, 或是對人類知識與經驗的打臉或是暮鼓晨鐘, 這更是一個可以套用有階段方法性的推進過程
1. 既有資料的收集, 既有知識的分析, 許多機械學習科班出身的人都知道, prior 並不是個好方法, 但對於較為龐大的問題 prior 是個短時間快速有效的切入點
談人們下圍棋所圍繞的定石, 子效, 跳, 飛, 星, 目, 邊, 場, 局, 角, 形 等等判斷與對應的工具, 其實都是為了應付棋局變化繁複而嘗試以簡馭繁的工具, 真的細較會發現這些工具與規則在棋局中間衝突與矛盾亦多, 長遠的利益與損失難以計較, 人的思考能力是只能提綱挈領而無法盡數.
2. 建立正確的"網路模型", "評估優劣方式"與"可自我改善的機制"
有了資料並不是就解決所有問題了
AlphaGo 的論文早已公開, 照理說之後各家推出的 Go Machine 都能有 AlphaGo 強度, 應當殺得職業棋士大喊不如歸去, 但是這一年看來很明顯並非如此. 在這領域打滾的研究者都清楚, 有了資料後該如何建立有效的網路模型也是相當的難題, 除了學習速率外還得兼顧有效性, 此外還必須在未知中找到可能的改進方式. 台上的一步 30 秒棋以及與高段棋手的勝局, 這階段上代表的是對於圍棋問題在本質上更正確的認識(數學, 分析與工程上, 與職業棋手不同面向)
3. 透過 2. 建立起純資料訓練的 NN
在大規模且複雜的問題上, 沒有 1, 2. 的努力無法走到這步, , 透過 1. 2. 所建立起對於圍棋本身的了解與分析的方法, 回過頭去作做法上的優化, 這些包含的是 1. 的 prior data 亦或是 2. 的妥協設計與初期有效但不理解的架構, 甚至是根本架構與有效性的提升, 在 AlphaGo/Master 已是領先世界的情況下, DeepMind 依然有著知識與眼界上的高度, 再優化既有的 AlphaGo/Master 來達到更效的且令人為之一亮的 AlphaGo Zero.
這其實更是人類知識上的累積, 對問題理解層次的演進, 理論模型細膩度的提升

2017年9月15日 星期五

整合 MT-32 摹擬音源的 dosbox-patched

dosbox 對於許多老玩家並不陌生
提供了簡單的方式讓使用者能執行 DOS, Win 3.1/9x 的程式

相信多數用途是用來懷舊喜愛的遊戲
而音樂在遊戲中扮演相當重要的角色
在沒有強力計算與空間存放來支援 MP3/AAC 等 HQ 音樂
除了 sample rate 不高的 PCM 播放外
另外而靠的就是 MIDI (General MIDI 或是著名的音源器)
能透過少少的資料得到很棒的音樂輸出
(小時候不能理解同個遊戲在店家聽到的音樂跟我在家玩時怎麼差這麼多)

dosbox 本身並不俱備 MIDI 裝置的摹擬 (qemu, vbox 亦是)
加上現有的音效裝置多半移除了 MIDI 的硬體支援
(回到 20 年前, 內建高容量 WaveTable 的 SB AWE 可是一大訴求)
所以 dosbox 0.74 在 Linux 上想要聽美妙的 MIDI 音樂是很困難的
( Windows 上雖然輸出不佳, 但是 XP 以後至少內建了軟體 MIDI 音源)
通常 Linux 上的方案是設定 FluidSynth, Timidity 或是 Munt
千辛萬苦設定完之後, 會發現音樂播放中的切換可能造成沒有結尾的尾音
結果只能望遊戲而興嘆 (eg: 波斯王子1/2, A列車3/4, 三國志, Raptor, SimCity 2000 ...)

最近發現非官方版的 dosbox-patched ubuntu ppa
除了 code base 使用 SVN 上功能增加許多的新版外 (0.74 已經發佈多年)
最令人雀躍不已的就是直接整合了 MT32 的摹擬
(需要 Roland MT-32 ROM file, 不難找到)
維護者也很有心提供了相當多 ubuntu 版本的套件
著實方便了眾多老遊戲玩家

2017年9月4日 星期一

Kirin 970 NPU - 看圖說故事

Image Sensors World 看到了華為旗艦手機晶片 Kirin 970 中 NPU 的投影片
其中有趣的是第五張圖


"25X performance, 50X efficiency"
相較於 CPU 有著 25倍性能 與 50倍的效率
由於 Performance 有相當多比較的面向, 因此不清楚是用絕對的指標(total GFLOPs, or per core) 或是相對指標 (OPs/cycle/core), 甚至是實際應用的時間來做比對(eg. FPS)
但是若是 Efficiency 談的比較可能的是 Performance Per Watt (另外還有 area efficiency, 甚至一起綜合看)
Efficiency 比較上為了有較好的數字, 比較可能的是與大核等級的 CPU 相比, 因此這裡以 Kirin 970 中的 Cortex-A73 為對比
以 NEON 的 GFLOPS 來看2.4Ghz Cortex-A73 單核的 GFLOPs 為:
2.4Ghz * 4 lane * 2 mad * 2 unit = 38.4 GFLOPS 
(寫到這就知道 Performance 不是以絕對 GFLOPS 計算)
而若 TSMC 10nm 製程能將 2.4Ghz Cortex-A7x 最大功耗(Peak Power)壓低至 1 Watt/Core(這是個人經驗粗略推估數字), 那麼 50X 效率可以推得:
38.4 x 50 = 1920 GFLOPs/Watt =  1.92 TFlops/Watt
(寫到這覺得蠻剛好的)
又可以從投影片上得知華為對外宣稱其 Kirin 970 內的 NPU 有著高達 1.92 TFLOPS/s (FP16) computation throughput




那麼可以由上面的比較, 對於 NPU 的最大功耗能推估為 1 Watt
以此看推估觀察運作的功耗圖介於 0.3~0.7 Watt 看來也是合理的(畢竟應該有著 Thermal/Power 管理面向的 DVFS)


2017年5月12日 星期五

Android 軟體架構轉變的進行式

今日在 Google 的 Android Developer Blog 上貼出了篇名為的 "Here comes Treble: A modular base for Android" 貼文, 這是一件對於 Android 生態系統的大事, 也是 Google 方面嘗試來解決 Android 的碎片化與安全性上的問題.

這件事必須先從 Android 裝置的軟體是經過什麼樣的流程抵達每個使用者手上, 在這篇文章中即是一開始的五步驟圖:
  1. Android 團隊向世界發佈最新版本的原始程式碼
  2. 晶片製造公司為了讓他們各自的晶片在 Android 裝置上俱備更多優勢, 針對他們特定的硬體, 修改了來自 1. 的原始碼與增添驅動程式.
  3. 這些晶片製造商接著將這些修改提供給他們的客戶 - 也就是那些設計與生產 Android 裝置的品牌商. 這時裝置生產商再次針對他們的產品對 1. 的軟體做了修改.
  4. 裝置生產商與電信網路商測試與驗證新的版本.
  5. 裝置生產商與電信網路商提供新的版本給他們的使用者

 在以往 Android 裝置間是透過相容性文件 Compatibility Definition Document (CDD) 所規範的介面以及現在已經涵蓋上百萬對應測項的相容性測試套件 Compatibility Test Suite (CTS) 來避免軟體相容性上的問題. 然而面對來自上述 2. ~ 5. 項目的修改與限制, 造成了 Android 軟體層上碎片化不容易解決, 1. ~ 3. 表示裝置生產商會 follow 的不是 1. 的官方實作而是 2. 的 chip vendor 的 Android Framework 的修改, 此外 4. 與 5. 反映了 Android 裝置的安全性修正完全取決於裝置生產商自己的更新意願(特別是裝置生產商不願意透入成本在維護已發售許久的裝置).

為此 Google 在 Android O 開始要著手解決這樣的問題, 提出了 Project Treble, 其目的在於強化對於 Android 系統軟體這塊的控制與掌握, 這必須談到在以往即已經存在, 但是界線並沒有明確規範的 Vendor Implementation (廠商實作層), 現在 Google 將明確的對 Android Framework 與 Vendor Implementation 切出這一條線:
因此先前所採用的 CDD + CTS 方式, 將轉為 Vendor Interface + Vendor Test Suite(VTS), 這件事最大的意義在於 Chip Vendor(晶片商) (或是很有可能以後包含 Device Vendor(裝置生產商) )將不再能夠直接修改上述五步驟中 1. 裡面的 Android Framework 的原始碼, 而優化與性能提升的方式也被限縮在圖中灰黑色(這顏色是不是特別選過阿...) 的 Vendor Implementation 這一層, 介面也是所定義規範出的 Vendor Interface.
藉由明確規範 Vendor Interface 的好處如上圖, 也就是 Android OS framework 的更新將獨立於 Vendor Implementation, 而在該篇文章是這麼說的 "device makers can choose to deliver a new Android release to consumers by just updating the Android OS framework without any additional work required from the silicon manufacturers", 翻成中文就是裝置生產商能夠無需藉助於晶片商額外的處理下, 就能夠自行更新 Android OS framework 的部分(這或多或少都是裝置商對於更新推托的理由與藉口). 是的, 目前 Project Treble 並沒有強勢到一次將 Android Framework 客製的權力全部收回, 這或多或少是出於市場考量, 但是至少讓 Google 頭疼的源頭 - Chip Vendor 的客製已然受限, 當然這樣作法背後還是在於 Google 對於 Chip Vendor 的客製只能道德勸說並無其他的約束能力, 而 Chip Vendor 還是會嘗試對客戶提供他們自己客製的 Android Framework, 一方面這樣介面的建立會讓裝置系統商意識到不正確與不當的客製修改, 另一方面來說對於裝置生產商來說來自 Google 相關的軟體生態系(Google Play)的授權是相當重要的, 因此 Google 還是有一定其他要求與控制能力. 一旦架構建立了 Treble Architecture, Google 即更能順利的要求裝置廠商去提供裝置的軟體更新(而以往的理由會凸顯背後的問題, 讓 Google 更能第一時間掌握有問題的平台).

最後 Google 還在文末說道目前 Pixel 上運作的 Developer Preview of O 已經採用了這樣的架構方式.

2017年5月11日 星期四

有趣的 cache line 效應

在忙著處理一些個人事務的時候, Jim Huang 請我幫忙回覆學弟妹碰到的 一些問題
主要是 Jim Huang 修改了之前 Matrix Transpose 的例子作為範例給學弟妹們練功使用
而學弟妹碰到了一個問題, 就是對於 Matrix Transpose 的寬做調整時
執行時間與長寬度的調整的分佈圖如下:
如上圖所示, 隨著對寬作 4096 + (i*64)的調整時經分佈呈現鋸齒狀的變化!
依照文中敘述的實驗平台為俱備 6MB cache 的 Core-i5 6300HQ
以 Intel Skylake 架構而言, 為 64 bytes cache line, 8-way associative cache
基本上俱備 6MB/64/8 = 12288 sets (entries)
以這個 int matrix 而言, 原始的寬度為 4096x4=16384 bytes
16384/64 = 256
因此每增加一行記憶體基本上會跳躍 256 cache line index

接著再觀察 naive transpose 的實作:
https://github.com/yangyang95/prefetcher/blob/master/impl/naive_transpose.c
若 buffer 對應的 cache line index offset 為 X

首先考慮最簡單的 i == 0 的情況
對於第 n 行的 cache line index 計算為: (X + 256*n) % 12288
1. 12288 / 256 = 48 (自 cache line 0 ~ 12288, 每次跳 256 可以消耗 48 行)
2. 12288 % 256 = 0 (每次的 cache line index shift)
所以填完 output 第一行來說, 這也代表著當資料(二維座標以(x,y)表示)自 (0, 0), (0, 1) .... (0, 4095) 要讀 (1, 0) 時
相關的 cache line 早已被反覆地填入了 4096/48 = 85.33次 > 8
當 CPU 需要讀入 (1, 0) 時, 因為讀入 (0, 0) 所讀入的 cache line 有相當高的機會早已被 replace

接著當每次寬度增加 64 時(也就是 256 byte, 4 cache lines)
這裡我們先考慮 i == 1 的情況
對於第 n 行的 cache line index 計算為: (X + 260*n) % 12288
1. 12288 / 260 = 47.26 ...
2. 12288 % 260 = 68
260 與 68 兩者最小公倍數為 4420
4420/68 = 65  (表示反覆繞 cache line index 0 ~ 12288 到第 65 次才開始 cache line index 有所重疊)
而 4160/47.25 約為 88.02, 88.02 / 65 = 1.35 < 8
這表示在讀取 (1, 0) 時, 有很大的機會 (0, 0) 所填入的 cache line 還在

再考慮 i == 8 的情況
對於第 n 行的 cache line index 計算為: (X + 288*n) % 12288
1. 12288 / 288 = 42.6666....
2. 12288 % 288 = 192
而 288 與 192 兩者最小公倍數為 576
576/192 = 3 (表示反覆繞填 cache line index 0 ~ 12288 到第 3 次 cache line index 就有所重疊)
而 4608/42.666 約為 108, 108 / 3 = 36 > 8
這表示在讀取 (1, 0) 時, (0, 0) 所填入的 cache line 有相當高的機會早已被 replace

由類似的計算可以得知, 造成這樣的時間分佈的原因, 而 SSE 版本的成因也是基於相同道理
而這樣的效應更顯示程式的記憶體使用方式, data layout 與 cache 影響程式的效能甚巨

2017年5月5日 星期五

Blocks 初探與 multithreading 應用

近日由於個人一項工作的緣故, 為了能在短時間內能夠加速複雜程式的運作
因此現學現賣地採用了 OpenMP 做快速的優化實作, 最後得到 3倍左右的加速

儘管 OpenMP 表現不俗, 然而在 Android 上的實作必須仰賴已經被 deprecated 的 GCC
讓現今預設使用 clang 的環境必須特別撰寫 makefile
而儘管亦能夠使用先前撰文介紹的 Grand Central Dispatch 作為方案
然而一方面 libdispatch 的 Android port 已經過舊
而官方版必須限定 Android API-level 21 外 (另外現在編譯也有問題)
另一個問題是 GCD 的實作規模不可以說小

由於 Blocks 的使用其簡潔與動態的特性對於個人而言是充滿魅力的
基於如此的動機便開始構思結合 thread pool 與 Blocks 的方式
以此來簡化 multi-threading 程式的撰寫, 並且得到能夠有效修改的實作
為此目的必須先去了解 Blocks 是如何運作的
儘管 clang 提供了 "Language Specification for Blocks" 的頁面
然而個人認為 Apple 的 Blocks Programming - Introduction 撰寫的比較淺顯易懂

基本上 Blocks 提供了將 C/C++ 的大括號內的 code 區塊轉化作為類似下列型別作為"變數"的能力
void (*func)(void);
而 Blocks 中最有趣以及最為實用的功能是對於使用 global/local variable 上數值的"擷取"
以 local variable 為例, 一般的 serial code 毫無疑問會在 stack memory 中
若使用了 Blocks 並於 Blocks 中使用了該 Block 區間外的 global/local variable
這時的流程會產生了類似 process fork 的分歧, 理解上應為未明確寫出的 call by value
Blocks 中對於外部的 global/local variable 本身的修改是不具有寫回的效果
除此之外型別為 Blocks 的變數在使用上的概念其實與一般變數無異
程式撰寫的過程中同樣地必須考量與處理 variable lifetime 的問題
這時就必須藉由使用 Block_copy/Block_release 來手動複製與釋放 Blocks 所含的內容
對於程式中 local/global variable 處理概念上的不同是 OpenMP 與 Blocks 最大差異
而兩者所使用的方式, 在應用上來說真的是各有優缺

透過閱讀上述的參考資料建立概念後
接著就動手來做類似於 GCD 的 Blocks dispatching 的功能

首先是建立等同於 GCD 中使用來作為 task dispatch 的 dispatch_block_t 的型別
接著就是增加 thread pool 的 dispatch function
基本上是將原本的介面的參數自 function pointer 與型別為 (void *) 的 argument
改為直接使用 Block 型別
可不用再撰寫型別為 void* func(void*) 的 pthread glue code 的好處不用再多說了

在建立概念之後, 動手實作上就簡單多了
為了快速而採用了現成的 thread pool 實作 - C-Thread-Pool
而 Blocks 的操作是基於 BlocksRuntime (提供了 Block_copy / Block_release)
而初步的成果我暫且名為 gunshot
請 git clone 後記得 git submodule init/update

修改的 example.c 中, 嘗試比較填入一個 buffer 數值
    for(int pidx = 0; pidx < TEST_DEPTH; pidx++){
        int *plane = buf0 + pidx*TEST_W*TEST_H;
        for(int yidx = 0; yidx < TEST_H; yidx++){
            for(int xidx = 0; xidx < TEST_W; xidx++){
                plane[yidx*TEST_W + xidx] = pidx*4096 + (yidx + xidx);
            }
        }
    }
而 thread pool + Blocks 的版本可以寫為
    for(int pidx = 0; pidx < TEST_DEPTH; pidx++){
        int *plane = buf1 + pidx*TEST_W*TEST_H;
        thpool_add_block(thpool, ^{
            for(int yidx = 0; yidx < TEST_H; yidx++){
                for(int xidx = 0; xidx < TEST_W; xidx++){
                    plane[yidx*TEST_W + xidx] = pidx*4096 + (yidx + xidx);
                }
            }
        });
    }

在個人使用的 Quad-Core A8-5545M 平台得到了以下結果
single thread - 34832 us
4 threads - 13129 us

得到了 2.65 倍的加速

2017年4月19日 星期三

浮點數的美麗與哀愁

這幾年個人在影像處理程式優化的領域打滾, 如果問到感到棘手的工作, floating point 的處理應該可以排上很前面的名次

在許多演算來說由於同時對於 precision 與 dynamic range 的需求, 因此在計算過程中對於浮點數的使用是非常常見的 (若要避免使用會有很高的專業與困難度), floating 主要優點在於可以表示極大與極小值, 相較整數能大幅避免 overflow 與 underflow, 缺點是有效位數的減少, 而且現今多數的計算單元都俱備 floating 的支援, 已經讓一些人疏於了使用浮點的問題, (包含486與之前的時代 FPU是高檔貨, ARM 也自 ARMv7 才列標配)
然而若橫跨了 PC 與 手機, CPU 與 GPU, CPU 與 DSP, 甚至於三者 ( PC, 手機, GPU), floating point 就變成非常難以考量與處理的負擔, 而為了區分是程式錯誤或是誤差就必須耗費相當的心力

為了簡化問題, 因此文中談到 floating point 若無指名, 一律是指 32bit single precision, 但相同的問題 64bit double precision 中一樣存在

IEEE 754

首先必須要談的是問題的核心 - IEEE 754
對於計算機而言, 浮點數是以上圖的格式存放
fraction 一共 23 bits 存放一個介於 1~2 之間的數目, 這 b22 ~ b0 存放的是二進位小數以下的部分, 也就是說 fraction 所表示的數值為:
1 + b22*(2^-1) + b21*(2^-2) + ... + b0*(2^-23)
而 exponent 代表著指數, 一樣以 2 的只數次方表示, 具有 8bit, 因此值域為 0~255, 但是預設會減去 127, 所以即為 -127 ~ 128, 因此對於 exponent 本身所表示的數值為:
2^(exponent - 127)
而 sign 就不用多談了, 這是用以表示正負

然而 IEEE 754 定義的不僅僅只是 format 而已, 還有著 rounding mode, required operations 以及 exception handling, 符合了 IEEE 754 的規範下, 才可能有相同的輸出結果 (這當然只是一個最低門檻)

Format 本身的問題

扣除浮點數因格式問題不可能表示全部的數目外
格式本身最大的問題是因為 dynamic range 的移動, 像是 (A+B)+C A+(B+C) 單以代數考量這無疑是相等的, 但是若以 floating 格式去思考, 你就會意會到輸出結果很有可能會不同, 而原始演算實作所採用的累加或相乘的順序, 必須在優化實作上努力維持才能產生一致的輸出結果, 這樣的問題對於程式優化影響最大的部分是平行計算, 無論 TLP or ILP, 因為平行度優化考量, 都會有分割與不同面向個別累計的需求, 如此勢必都會產生一定的誤差

CPU 間的問題

或許有些人會認為只使用 CPU 那麼就不會碰到浮點數問題了, 這樣說只算對了一半, 而且你還是必須要只使用一種平台以及指令集, 對於多數演算法設計工程師而言, 他們很慣於使用 PC 平台, 甚至會使用 MATLAB, x86 PC 上的程式預設會使用 x87 浮點數協同單元 指令集, 而這是許多問題的開始 - x87 內部使用 80bit 浮點數表示
通常 x86 CPU 是在 x87 FPU 的內部以 80b 計算得到結果後, 再 truncate 為 IEEE-754 的 float(32b) or double(64b), 這就表示這與 IEEE 754 FPU 的輸出結果會有微小的差異, 目前常見的手機的 ARM 架構, 即為 IEEE 754 compatible FPU, 所以光是 ARM 與 x86 PC 相同的程式碼其輸出結果基本上就會有所不同, 而對於 ARM 與 x86 的部分, 就必須仰賴以 IEEE 754 設計的 SSE2 指令集, 若是 gcc 與 clang 很早就可透過 -mfpmath=sse2 的編譯參數來達到, 但是 Visual Studio 必須是 2013 版後才有正確的 code generation 實作, (也就是說 Windows 的使用者要安裝 VS 2012以後的版本 才有可能透過 /arch:sse2 有一致的輸出)

對於 ARM 與 x86 平台的一致性方案反而又揭露了另一個層面問題:
就算使用了單一CPU架構, 在 ISA 指令集間的支援還是會有所不一致!
類似於 x86 平台上 x87 指令與 SSE2 指令有著不同輸出, 同樣地 ARM VFP 指令(IEEE 754 相容)與 NEON 指令(非完全 IEEE 754 相容) 也可能會有輸出結果不同, 而這樣的問題還會再帶到 libmath 的實作方式, 讓要處理一致性的問題再度的變得更嚴重

GPU

GPU 本身有著龐大的浮點數計算能力, 但是通常為了能達到更高的吞吐量以及加速上的考量, 在計算結果與 IEEE-754 可能存在差異, 不同代的 GPU 或是不同架構都有可能有所不同. 像是 CUDA 是在 compute compatibilty v2.0 之後才完備了 IEEE 754 的支援, 除此之外許多硬體加速的數學函數的輸出上也不保證與 CPU 一致, 這點 Nvidia 在 2011 GTC 中給的 Floating Point and IEEE-754 Compliance for Nvidia GPUs 簡報中有很詳盡的說明. 對於其他GPU 以及各CPU/GPU 平台上的 OpenCL 中的 built-in functions 的實作與支援也有著相同的道理.

DSP

對於 DSP 而言這樣的痛苦並不在於 IEEE 754 本身, 而是多數的多媒體面向的 DSP 為了考量計算能力與面積, 結果多半是直接不俱備 floating 能力的, 像是 Qualcomm S82x 中的 Hexagon 680 HVX 就不俱備 floating 運算的 SIMD 指令, 而通常的處理作法是採用 fixed-point (quantization) 的浮點模擬, 然而若採用靜態位數的方式容易失真, 而動態的方式有著實作上的複雜度以及多餘計算的負擔. 而數學函數上的實作若難以避免則通常必須透過相當紆迴的方式.

Lookup-Table or Frame-based Parameters

對於跨裝置的正確性驗證, 由單一裝置輸出的 Lookup Table(單一的 math function 像是 sin, cos, log, exp 等等) 或是一整張預先透過單一裝置計算的 Frame-based Parameters(複雜的並結合多個 math function 的運算) , 是常用來確認誤差單純是由 floating 計算造成的技巧. 以此來確保實作上的流程與邏輯無誤.

延伸閱讀: The pitfalls of verifying floating-point computations

2017年4月2日 星期日

"ARM Compute Library for computer vision and machine learning" III - 總結

在實作上核心的實作是在各功能的 Kernel 類別實作中
因此若想瞭解可以去研讀各個繼承 IKernel 的 CL/NEON 實作
由系列文 II 多少可以了解 ARM Compute Library 是如何的工具
在官方的介紹也說明了這是 - a collection of low-level software functions
這樣的好處是設計簡單且易於使用, 若所需要功能不複雜, 其所提供的工具也相當堪用
但是若一個目的是需要使用多個 Kernel 的串連來達成
如此的應用就需要更進階的方式來作優化
以 OpenVX 來說即為其 Graph pipeline
基本上需要透過一個更為高階的抽象層
為問題帶入各個 stage 的分割與相關排程的分析
對於 ARM Compute Library 而言每個 function 需要 I/O image buffer
能接近這樣的方式在於兩個 stage 間以 Window + Thread 的 Tiling 方式
如此也僅是有限度地利用 data locality 的特性增加 cache 的有效性
況且對於 Load/Store 指令可是一個都沒能因此節省
(這需要能作 operation fusion 的 compiler)
這即是 ARM Compute Library 在效能與進階功能上的局限


然而若需要進一步解決上述的局限
會需要能針對 temporal/spatial scheduling 作 dynamic code generation 的 compiler以及 runtime
實作複雜度亦會大幅增加 (即 OpenVX/Halide 或類似的實作)

2017年3月28日 星期二

"ARM Compute Library for computer vision and machine learning" II - Framework 篇

ARM Compute Library 的使用上的 class 主要有二類分別是針對 data 以及 task/workload
data 類別為: Image/Tensor, TensorInfo
task/workload 類別為: Kernel, Window 與 Function
下列的內容主要為 ARM Compute Library: Documentation 所描述

並且搭配 source code 的內容做特定用途的說明
(取代文件內 MyKernel, MyFunction 的方式)

由於 ARM Computer Library 是做 Computer Vision 與 Machine Learning 應用的
因此主要處理的資料型別為 Image 及 Tensor
在 Compute Library 中基本上只是名稱不同而已

Image, Tensor, TensorInfo

在 NEON 下直接使用 Image
Image     src, dst;
而在 CL 下則使用 CLImage
CLImage   src, tmp, dst;
Image 宣告後並沒有實際的 buffer 空間, 必須進一步最配置的動作,配置的方法有二, 兩者都需要傳遞 TensorInfo 資訊, TensorInfo 基本上為 Image/Tensor 各維度的大小以及資料格式
配置的第一種方式為直接透過 Allocator 的 init() 方式
src.allocator()->init(TensorInfo(640, 480, Format::U8));
而第二種方式為先呼叫 configure() 在呼叫 Allocator 的 allocate()
TensorInfo dst_tensor_info(src.info()->dimension(0) / scale_factor, src.info()->dimension(1) / scale_factor, Format::U8);
dst.allocator()->init(dst_tensor_info);
dst.allocator()->allocate();

Kernel, Window & Function

空間配置好之後, 就必須透過各種各樣的 Kernel 來套用對應的功能來操作 Image/Tensor
使用上的核心 class 為 Kernel, 各個 Kernel 實作了 IKernel 相關的介面

使用的第一個步驟為宣告想使用的 kernel object, 假設我們想作 image scale
//Create a kernel object:
NEScaleKernel scale_kernel;
在使用之前必須對 kernel 作 input, output 使用參數以及 padding mode 作設定
// Initialize the kernel with the input/output and options you want to use:
Tensor offsets;
const TensorShape shape(dst.info()->dimension(0), dst.info()->dimension(1));
TensorInfo tensor_info_offsets(shape, Format::S32);
tensor_info_offsets.auto_padding();
offsets.allocator()->init(tensor_info_offsets);
scale_kernel.configure( &src, nullptr, nullptr, &offset, &dst, InterpolationPolicy::NEAREST_NEIGHBOR, BorderMode::UNDEFINED);

offsets.allocator()->allocate();
// compute offset 
...
這裡使用了 NEAREST Filter, 並且使用 UNDEFINED padding 方式 (對於邊界有缺所需資料的點不做處理)

最後就是呼叫使用該功能,即是呼叫 IKernel 的 run() 介面
// Retrieve the execution window of the kernel:
const Window& max_window = scale_kernel.window();
// Run the whole kernel in the current thread:
scale_kernel.run( max_window ); // Run the kernel on the full window
這些即為 Compute Library 基本的使用方法.
對於 CL Kernel 則有稍微不同的 flow (需要特別傳入以操作 cl::CommandQueue)

或許會覺得上面例子的 max_window 很多餘, 但它是有進階應用的
Window 用途在於對 Kernel 指定要套用執行的範圍描述
在官方說明文件是以 Multi-Threading 的方式來說明 Window 的用途
const Window &max_window = scale_kernel->window();
const int num_iterations = max_window.num_iterations(split_dimension);
int num_threads    = std::min(num_iterations, _num_threads);
for(int t = 0; t < num_threads; ++t){
    Window win = max_window.split_window(split_dimension, t, num_threads);
    win.set_thread_id(t);
    win.set_num_threads(num_threads);
    if(t != num_threads - 1){
        _threads[t].start(kernel, win);    }else{
        scale_kernel->run(win);    }
}
當下列所有的條件都符合後, Window 可以用來分割 workload 為多個子 Window
  • max[n].start() <= sub[n].start() < max[n].end()
  • sub[n].start() < sub[n].end() <= max[n].end()
  • max[n].step() == sub[n].step()
  • (sub[n].start() - max[n].start()) % max[n].step() == 0
  • (sub[n].end() - sub[n].start()) % max[n].step() == 0

至於 Function 的使用則是為了簡化繁雜的 Kernel, Window 的使用流程, Function 實作內部會自行配置所需的暫存  buffer, 甚至能透過上述的方式自行做 Multi-Threading
// Create and initialize a Scale function object:
NEScale scale;scale.configure(&src, &dst, InterpolationPolicy::NEAREST_NEIGHBOR, BorderMode::UNDEFINED);
// Run the scale operation:
scale.run();
若使用的為以 CL 實作的 Kernel 最後還有個確保執行完成的額外同步動作
// Make sure all the OpenCL jobs are done executing:
CLScheduler::get().sync();

如此一來 Function 比起直接使用 Kernel 簡化不少

在下一篇會進入 NEON 與 CL 內部實作方式的說明

2017年3月26日 星期日

簡報 - Video Compression Standards - History & Introduction

這份簡報是一年半前為了數位電視課程所準備, 發現 blog 沒有紀錄所以發文分享與 Link 一下


"ARM Compute Library for computer vision and machine learning" I - Overview 篇

日前 ARM 官方透過 github 並且以 MIT License 方式再次正式地釋出了 Compute Library 的原始碼(先前提供了 internal evaluation only 的 binary, 詳請請回顧當時的 release note), 這是個 low-level implementation, 而且是 pure function 的形式, Compute Library 提供了涵蓋下列的功能函式:
  • 基本的運算, 數學與布林運算函式 (Basic arithmetic, mathematical, and binary operator functions)
  • 色彩操作, 包含轉換, 頻道擷取與其他 (Color manipulation (conversion, channel extraction, and more))
  • 捲積濾波器 (Convolution filters (Sobel, Gaussian, and more))
  • Canny Edge, Harris corners, optical flow, and more
  • Pyramids (such as Laplacians)
  • HOG (Histogram of Oriented Gradients)
  • SVM (Support Vector Machines)
  • 半/全精準 通用矩陣乘法 (H/SGEMM (Half and Single precision General Matrix Multiply))
  • 捲積類神經網路建構功能區塊 (Convolutional Neural Networks building blocks (Activation, Convolution, Fully connected, Locally connected, Normalization, Pooling, Soft-max))
對於 Compute Library 來說, 它屬於個人介紹過的(請參考SIMD Introduction 簡報) SIMD Programming Model 中的 SIMD Optimized Library, 概念與與提供的效能上, 可以參考 ARM 官方釋出的介紹文,  本系列文會專注於 Compute Library 內部架構與更細節的如何使用, 所提供的能力以及, 以及探討使用這樣的 Library 依然存在什麼樣的限制.

首先 ARM Compute Library 其 github 位置為 https://github.com/ARM-software/ComputeLibrary , 而相關的原文文件在 source 與網站上各有一份

以目錄結構來說下列為主要較重要的目錄:
  • arm_compute/ - 放置所有 Compute Libraray 的 Headers
    • core/ - Core library 是由底層演算法的實作所組成
      • 基本共用資料型別 (Types, Window, Coordinates, Iterator, 等等)
      • 基本通用介面 (ITensor, IImage, 等等)
      • 物件 metadata 型別 (ImageInfo, TensorInfo, MultiImageInfo)
      • backend 目錄
    • runtime/ - Runtime library 是用來快速 prototyping 用途非常基本的 Core Library 的 wrapper (由於 CL/NEON 的 Programming Model, 這裡提供對應不同的 execution interface)
      • 基本通用物件介面的實作(Array, Image, Tensor, etc.)
    • 以上兩者, 內各自有 CL/CPP/NEON backedn 目錄, 提供對應 backend 定義的 kernel headers
  • documentation/ - Doxygen 所產生的文件
  • examples/ - 內有提供的 4 個範例程式
  • include/ - 基本上只放置 OpenCL 1.2 的 Headers
    • CL/
  • src/
    • core/ - 於 arm_compute/core/ 中定義的型別/介面的實作
    • runtime/  - 於 arm_compute/runtime/ 中定義的型別/介面的實作
    • 以上兩者, 內各自有 CL/CPP/NEON 目錄, 即為該 backend 實作相關原始碼

而值得一提的是在 Compute Library 中提供的功能中, 這些 Kernel 演算所使用的定義規範為 OpenVX 1.1 所制定的

以下為目前提供的 Kernel 列表, 若了解 image processing, DNN 該 function 名稱應解釋了其功用, 即不在此冗文解釋: (注明 NEON-only 表示目前尚未有 CL 實作)
AbsoluteDifferenceKernel
AccumulateKernel
ActivationLayerKernel
ArithmeticAdditionKernel
ArithmeticSubtractionKernel
BitwiseAndKernel
BitwiseNotKernel
BitwiseOrKernel
BitwiseXorKernel
Box3x3Kernel
CannyEdgeKernel
ChannelCombineKernel
ChannelExtractKernel
Col2ImKernel
ColorConvertKernel
ConvolutionKernel
ConvolutionLayerWeightsReshapeKernel
CumulativeDistributionKernel (NEON-only)
DepthConvertKernel
DerivativeKernel
DilateKernel
ErodeKernel
FastCornersKernel
FillArrayKernel (NEON-only)
FillBorderKernel
FillInnerBorderKernel (NEON-only)
Gaussian3x3Kernel
Gaussian5x5Kernel
GaussianPyramidKernel
GEMMInterleave4x4Kernel
GEMMLowpMatrixMultiplyKernel
GEMMMatrixAccumulateBiasesKernel
GEMMMatrixAdditionKernel
GEMMMatrixMultiplyKernel
GEMMTranspose1xWKernel
HarrisCornersKernel
HistogramKernel
HOGDescriptorKernel (NEON-only)
HOGDetectorKernel (NEON-only)
HOGNonMaximaSuppressionKernel (NEON-only)
Im2ColKernel
IntegralImageKernel
LKTrackerKernel
MagnitudePhaseKernel
MeanStdDevKernel
Median3x3Kernel
MinMaxLocationKernel
NonLinearFilterKernel
NonMaximaSuppression3x3Kernel
NormalizationLayerKernel
PixelWiseMultiplicationKernel
PoolingLayerKernel
RemapKernel
ScaleKernel
Scharr3x3Kernel
Sobel3x3Kernel
Sobel5x5Kernel
Sobel7x7Kernel
SoftmaxLayerKernel
TableLookupKernel
ThresholdKernel
TransposeKernel
WarpKernel (CL 細分為 WarpAffine, WarpPerspective 兩種)

下一篇將會介紹 Compute Library 中使用所需了解的基本型別, 介面, 執行方式以及範例

2017年3月19日 星期日

OpenCL Programming Tips for Qualcomm Adreno GPU 導讀

目前手機多半都有內建 OpenCL Runtime
(除了 Google 無謂地堅持 RenderScript, Do Evil 地阻礙標準的採用)
對於 OpenCL 有所了解的人, 多半清楚 OpenCL 是 function portability, 而無法做到 performance portability, 這其中的緣由主要還是在於各家的 GPU Architecture 的差異, 因此多半各家 GPU vendor 都會提供自家 OpenCL Optimization Guide, 以利開發者對自家平台優化應用性能

目前手機可能內建的 GPU 主要為三
  • ARM Mali
  • Imagination PowerVR
  • Qualcomm Adreno
由於個人 OpenCL 經驗多半與前兩者有關,  由於未曾接觸, 因此對於 Adreno 部分有很大的興趣, 因而選擇研讀與撰文. 這篇主要的內容主要來自下列公開的 Adreno OpenCL Programming Guide 文件. 有機會也會撰文介紹 ARM Mali 與 Imagination PowerVR 的 Programming Guide

Adreno OpenCL Programming Tips

Memory

其第一個章節即是 "Memory", 對於計算架構來說 Memory 幾乎是效能上的關鍵, 對於 GPU 亦不例外, 而在這份文件中 Memory 章節佔了一半的份量, 可見其重要性, 對於 Adreno GPU 而言, 其考量點有五:
  • Vectorization and coalescing 
對於 memory access 而言, Coalescing 是常用的方式, 儘管一些 compiler 會透過 auto-vectorization 嘗試去優化 workgroup 內更有效率的存取, 但是 programmer 自行作 vectorization 並控制 access pattern 對於後續的 finetune 是重要的. 一旦 vectorization 後, OpenCL 提供兩個介面做 vector loading, 一者為 vector-based pointer arithmetic, 另一為 vloadn 的方式, 這裡 Adreno 建議以 vloadn 並且不建議 n > 4. (這必定也只是一個 common rule, 程式的流程與記憶體的行為也會有差異)
  • Image vs. buffer memory objects
 OpenCL 的記憶體使用分為兩種 abstract object, 一者為 Buffer 另一為 Image Object, 對於 Image 所提供的優點如下:
主要是相較於 Buffer 多了 L1 cache (這是 Texture Unit 的特性所增加), 另外就是硬體加速的線性內插的計算, 以及最後是能夠透過硬體自動處理邊界的問題(以 Buffer 而言需要增加 GPU 所不擅長的 if/else 的 code)
  • Global memory (GM) vs. local memory (LM)
在 Global Memory(GM) 與 Local Memory(LM) 間應注意的事情有三, barrier 的使用, 再者為搬移資料需要考量的 cost, 最後是將資料存放在 LM 的條件
對於 Kernel 撰寫實作熟悉者, 應該對於 barrier 的使用會有相當的 overhead 不陌生, 但是對於一些有 data dependency stages 的實作這又是必要的, 減少 barrier 的使用幾乎是所有 GPU 平台一致要納入考量的點.
儘管 LM 有著較低的 latency 但對於 Adreno GPU 來說 GM 到 LM 的途徑中需要使用 GPU 內部的暫存器, 隨著需要搬移的資料數目耗費的暫存器可能會引起 register spilling 問題, 另外 Adreno GPU 對於 GM 有 L2 cache, 因此並不是直接將資料放置於 LM 就能獲得效益.
對於需要放置到 LM 的條件, 這裡建議 LM 存放介於兩個 stage 中間的資料, 或是會被使用至少三次的 input data.
  • Private memory
對於多數的 GPU 架構 private memory 對應到的地方是暫存器, 對於 Adreno 亦不例外, 由於 general register 數目有限, 若不良的 coding style 會造成 register spilling, 而一旦這樣的情況發生, Adreno 會嘗試自 LM 調度, 若 LM 不足則最後會調度到 GM, 這樣的流程若頻繁發生, 可以預期的是效能會大幅降低. 因此對於想要宣告為 array 的變數, 建議直接使用 LM. 其實除了這樣之外,  可以做 in-function 的 multi-stages 方式, 積極地將 kernel 透過多個 stage subroutine 來實作, 最後手段是透過精準的 variable life-scope 控制(也就是加入 {} 大括號, 主動提供 compiler 資訊) 來減少 register 的使用.
  • Constant memory
 Adreno 提供了相對特別的 constant memory (經驗上來說常看到架構會使用 LM 作為存放 constant 的地方), 然而大小為 3KB, 超過的部分系統會放在 GM 中, 對於透過 Kernel Argument 傳入的 constant buffer 需要透過屬性的設置來預期會被放置於 constant memory.(詳細請參考 1. 文件)

Zero memory copy

如同其他 OpenCL Runtime, 若需要同時 CPU/GPU 能夠存取, Buffer/Image 的配置要透過 CL_MEM_ALLOC_HOST_PTR 這個 flag 來取得能夠讓 CPU/GPU 無需 data copy 的空間, 在透過Map/Unmap 的方式來使用. 然而對於除了 CPU, GPU 外的硬體需要使用, 需適當地選擇 cl_ion_qcom_host_ptrcl_qcom_android_native_buffer_host_ptr 這兩個 flag 來使用

Work group size and shape

儘管多數 OpenCL 教科書建議在 Kernel 執行時輸入將 local work size 參數傳入 NULL, 讓 Runtime 自動配置最適當的 WorkGroup size, 但是 Adreno 還是提醒這樣的作法其實並不總是最佳的(事實上所有的平台都不是), programmer 應該嘗試尋找適當的 WorkGroup size.

Data type and bit width

資料型別的使用上 Adreno 上建議使用較短的資料型別, 除了能夠減少資料存放的大小與減少頻寬, 像是對於 half (16位元浮點數) 與 float 而言, half 還提供了兩倍的計算能力.

Math functions

這裡 Adreno 上應避免除法與餘數的計算, 此外若 mul24/mad24 (24位元乘法/乘加) 足夠精準度的話, 應該儘量使用, 這主因是一個32位元的乘法計算在 Adreno 內部是透過3個指令來組成的.



後續會再探討

2017年2月11日 星期六

OpenCL Programming for Intel FPGA 初探 - III

這篇主題進入到了 OpeCL Programming for FPGA 的優化
內容主要圍繞在官方文件 - Intel FPGA SDK for OpenCL Best Practices Guide

首先必須要知道在 FPGA 上程式撰寫的快慢與計算的多寡並無直接關係
其重點在於計算間是否存在 dependency
儘管計算煩雜, 但是多數能夠平行運作, 在 FPGA 上就能更快的計算
對於 OpenCL for FPGA 的思考在於 pipeline stage 的長度與使用的 resource
下圖為官方所使用的一個簡單的範例:
儘管例子中使用了 5個運算, 但是實際的運算方式會是如上圖
而再透過 pipeline 的處理即可達到 II =1 (initiation interval) 的情況:

如上圖所示, 這麼一來即如同於 5 instruction per cycle, 能夠這樣做的原因在於對於一般的處理器而言, 僅有著有限的 register, 因此 processor 必須透過 DRAM access 來達到處理較多的資料, 而甚至因為變數數目的關係, register spilling 效應會讓 processor 為了暫存資訊更頻繁的讀寫 DRAM, 因而造成程式效能的惡化. 但是 FPGA 可藉由 Flip-Flops (FF) 來在 pipeline 的 stages 中存放這些資訊.

但是由於 FPGA 這樣的特性, 在 Kernel 之間傳遞資料也需要更好更有效率的方式, 因此 Pipe/Channel 的使用是很重要的, 碰到演算上因記憶體頻寬限制所造成的效能問題, 基本上的處理原則即是將演算再切分為多個 kernel, 並且將 memory access 部分改以如下圖的方式使用 pipe/channel, 藉此來處理.


在 Kernel 中的 loop 處理上, 必須透過優化讓 loop 的 initiation interval 達到 1, 下圖為官方文件所提供的 optimization 策略的決策圖:

至於而何為 initiation interval, 為何它這麼的重要? 下圖凸顯了若 II 不為 1 時的差異, 主要在於 loop 若非 pipelined, 效率上的 throughput 會呈倍數差異.

上述的資訊,在 Intel FPGA SDK 環境中有提供 Loop Analysis Report, 能夠以此觀察 loop 優化的情況.
從上圖得知從 Loop Analysis Report 中可以得知 unrolled 程度, 是否 pipelined, II 以及 Bottleneck 的點, 對於 Loop 撰寫的優化有很大的幫助.

除此之外, SDK 中還提供了 Area Report, 以下為一個例子
在 Area Report 上面可以觀察到下列幾點的欄位, 也就是在系列文 II 中為何要了解 FPGA Architecture.
  • ALUTs : 即為 LE 中的 用以實作 4-input 任意函數的 LUT
  • FFs : 這部分是 Flip-Flop, 也屬於 LE 中, 通常消耗在 pipeline stage 中的 register
  • RAMs : 這是 memory load/store 的行為
  • DSPs : 這是使用到的 DSP block 數量
對於在 FPGA 平台上 OpenCL Kernel 上的撰寫除了 pipe/channel 與 loop 外的其他考量即是圍繞在上述幾個點, 因此官方以 Matrix Multiplication 為例子, 提供了優化的過程的 code 與對應使用的方式, 下圖為 Matrix Multiplication 優化中的上述幾個點的變化過程:

這裡並不做細部的探討, 有興趣者請閱讀官方文件章節 1.4.2 ~ 1.4.5 , 但是對於各版本中使用的優化技巧, 簡述如下:
  • v1 : 使用 local memory 方式降低 memory access
  • v2 : 以 loop counter + 條件計算 row 值, 減少使用 % 運算. 這運算同除法會消耗大量 ALUTs/FFs
  • v3 : 對於 offset 的計算改以在 loop 中累加的方式, 取代以乘法計算, 如此可降低 DSPs 的使用量
  • v4 : 使用 per-colume 的累加方式, 最後在計算, 來降低因 data dependency 造成的 II 為 8 的 loop (增加了一個 loop 與多使用了變數, 因此 ALUTs/FFs 增加了, 但是效能變好)

對於 local memory 的使用也並不是直接使用了就沒問題, 對於 local memory 除了需要指定大小外, 還有 bank 的相關參數來避免 access stall 的發生, 若沒有足夠的資訊的話 compiler 會產生 local memory LSU 的邏輯線路來控制與管理 local memory access, 然而 LSU 的使用可能導致效率的降低.
此外 local memory 的大小使用由於 FPGA 的 memory block 的因素, 使用的大小會 rounding 為相關大小的倍數, 因此在使用的精算中必須考量到這點.

在官方文件中後續 1.6 ~ 1.9 又詳述了針對下列幾個面向的 Optimization Strategy:
  1. Single-WorkItem: 這階段為 Optimization 的根本,  因此文件除了列出的三點也提供了下圖詳盡的 workflow:
    • 依照 analysis report 來改進
    • 移除 loop 間的因 memory access 造成的 dependency
    • 良好的設計練習 
  2. NDRange
    • 限定最大的 WorkGroup 大小, 或是要求特定的 WorkGroup 大小
    • Kernel Vectorization
    • 產生多個 Compute Unit
    • Kernel Vectorization + 多個 Compute Unit
    • 依照 resource 的使用面向來優化
  3. 改進 memory access 效率
    • 優化 Global Memory Access
    • 使用 constant, local, private memory 來做計算
    • 將變數存入 private memory array
  4. 優化 FPGA Area 的使用
    • compilation 的參數
    • memory access 的考量
    • 運算上的考量
    • 資料型別上的考量
除了上述依照 compilation report 的 optimization, 在 runtime 上可能還需要使用 profiling tool 觀察實際運作的情況, 官方文件 1.5 的部分即是介紹 profiling 相關的功能如何使用, 以及提供了哪些數據.

2017年2月10日 星期五

OpenCL Programming for Intel FPGA 初探 - II

在進入針對 OpenCL for FPGA 的優化過程之前
人們必須先了解 FPGA 是如何的裝置, 與 CPU 與 GPU 在本質上的不同之處

Architecture Overview

若以低階的 Altera Cyclone V 來看的, 我們可以從其 Handbook Volume 1 看到其架構組成
由上圖可以看出 FPGA 的主體主要由三種 block 組成
  • CoreLogic Fabric and MLABs : 主要為邏輯與簡易計算的 blocks, MLAB (memory logic array block) 基本上是通用的 SRAM array, 亦可以被設定作為 RAM/ROM 的用途
  • M10K internal Memory Blocks : dual port 的記憶體 blocks, 可以存放 processor code, 用以實作查表設計, 與實作較大的記憶體應用. (像是 OpenCL 中的 local_memory)
  • Variable-Precision DSP Blocks : 可變精準度的 DSP blocks, 透過使用數量不等的區塊可以串聯出整數/浮點數精準度不同的乘法與運算單元.

與 CPU 不同的是, FPGA 的性能主要取決於上述三種 resource 的數量, 另外應用上使用的 resource 若無法壓低, 可能無法塞入 FPGA 上執行, 這是要注意的事情, 所以應用的選擇上必須注意規格, 下圖即為 Cyclone V SE SoC 內 FPGA resource 的表格:
從上途中可以觀察出一些數字是彼此相關的, 像是 18x18 multiplier 的數目與 DSP blocks 是直接成倍數關係的, M10K memory 大小 與 M10 memory blocks 數目也是如此, 這兩個部分分別對應了三種區塊的 M10K memory 與 DSP blocks, 然而第一項的 CoreLogic Fabric 呢?

Logic Elements

這部分需要從 LEs 這個數字解釋起, LE為 Logic Elements 的縮寫, 從 Altera 的解釋, LE 為其 邏輯部分的最小單位, 提供了下列功能:
  •  4 組輸入的 lookup-table(LUT) 功能, 可以實作任何 4 變數的函數
  • 可程式的 registers
  • carry chain connection
  • register chain connection
  • 多種功能的 interconnection
  • register packing, feedback
而上述的部分必須又必須從 Adaptive Logic Module(ALM) 來解釋, 下圖為 ALM 的 block diagram:

對於 ALM 官方文件是這麼描述的 "The ALM uses an 8-input fracturable look-up table (LUT) with four dedicated registers to help improve timing closure in register-rich designs and achieve even higher design packing capability than previous generations.", 而 LE 官方文件的說明為 4 組輸入, 由此可得知 ALM 為 > 2 LE 的組合, 這推斷也與官方提供的表格數字間符合.

DSP blocks

在程式中所使用到的乘法與浮點運算必須使用到 DSP blocks, 因此對於 FPGA 而言是相當重要的運算資源, DSP 提供了18b 與高精準的兩種運算模式:
此外亦單一 DSP block 提供了單精準浮點數運算以及串連4個 DSP block 形成倍精準浮點數運算

了解這些 FPGA 內部資源與用途與後續的 optimization 有很大的關聯

2017年2月5日 星期日

OpenCL Programming for Intel FPGA 初探 - I

前一陣子許多新聞都提到使用 Intel FPGA 搭配 OpenCL 來加速應用的開發
像是中興電訊(ZTE)  透過 Intel FPGA 來開發人臉識別的加速
其實 Altera 相當早就開始將 OpenCL 導入到 FPGA Programming 這塊
在 2011 年 Altera 的簡報是這樣類比的 - VHDL/Verilog 很類似"組合語言"程式設計
而 OpenCL 讓軟體開發者能夠利用硬體加速

由於這些日子來的一些例子也令我有不同的思考
開始讓個人非常想了解這樣的 Programming Model 的不同, 以及其特性與優缺點為何
因此開始投入 Intel FPGA OpenCL Programming 的探索

現今相關的軟體 Intel 都有提供在網路上, 硬體上取決於規格價格落差也不小
台灣的友晶科技有販售相關的平台, 可以自行了解與選購
但目前先單以 Intel 提供的官方文件來討論 OpenCL Programming for Intel FPGA

Programming Overview

Intel FPGA SDK for OpenCL - Getting Starting Guide: 其中對於了解階段最重要的地方在於
它圖示了整個流程的 overview
可以看出 .cl 的 kernel 透過 SDK 分兩階段編譯為不同的 .aocx, 一者為 emulation 所使用, 另一為實質在板子上的運作時所使用, 而這文件其他篇幅著墨在環境的建立上. 一方面也可以了解, OpenCL for Intel SDK 應不俱備 Online Compiler 的方式, 可能皆以 Offline Compilation 方式開發.

Intel SDK for OpenCL - Programming Guide:

Schematic Diagram

這篇文件在流程上延伸了上述的圖, 提供了 host, kernel 與 custom part 的流程
在 Getting Start 的那張圖示, 其實在解說中間 "Kernel Code Path" 在不同開發時期的產生方式, 而這張圖是講述完整的開發流程中, Host/Kernel/Custom 相關的各部分是屬於從何而來. 這張圖也揭示了, 開發過程僅有 Offline Compiler.

文件中提供了以 PC + FPGA 卡的角度來看待這個流程:
這張凸顯了, 能夠搭配多張 FPGA 卡來負責不同的部分, 透過載入不同的 FPGA binary 就可以將功能的不同部分在不同的卡上運作.

Channel and Pipe

為了 mapping FPGA 的運作方式到 OpenCL, 提供了特別的方式: Channel 與 Pipe.
事實上兩者性質很類似, 但是 Channel 為 1.x 時針對 FPGA 部分, Altera 所提出的 extension, 而 Pipe 是 OpenCL 2.0 中標準有提供的 inter-kernel 間資料串連的方式.
而對於 FPGA 中, 使用 Channel/Pipe 要思考的運作方式如下:
也就是 Kernel 與 Kernel 間以 FIFO 方式來串連與傳遞資料, Host 並不介入資料的傳遞, Kernel 也不透過對於 RAM 的讀寫來傳遞.

Manager-Producer-Consumer Working Model

而對於 Kernel 間 Buffer 的傳遞, SDK 文件中建議以 Manager, Producer, Consumer 這3個角色的 Kernel 來實作.
而透過圖中描述的四步驟周期, 來彼此協調控制流程, 並作一些 resource 的 ping-pong 使用
  1. manager kernel 送出 token set 給 producer kernel 來告知哪些記憶體區塊是給 producer 使用
  2. 在 manager 配置好記憶體區塊後, producer 對該記憶體區塊的 ping-pong buffer 寫入資料
  3. 在 producer 完成寫入動作後, 它送出同步 token 給 consumer kernel, 來告知記憶體區塊有著資料待處理. consumer kernel 然後就自該 ping-pong buffer 區塊讀取資料
    • 必須注意的地方: 當 consumer 在執行讀取動作時, 由於 FPGA 同步運行著 producer, consumer 與 manager kernel 的緣故, producer 是能夠寫入資料到其他尚未使用的記憶體位置,
  4. 在 consumer 完成讀取動作, 它釋放了該記憶體區塊, 然後送 token 給 manager kernel. 而 manager kernel 就回收該區塊以提供給 producer 使用.

Memory Partition of Global Memory - Burst-Interleaved vs Separate

Global Memory 不管對於 CPU/GPU 而言都是相當遙遠的, 必須考量頻寬與 latency, 因此對於操作上有著不同的特性, CPU 有著 cache,  GPU 仰賴著 local memory 與 many-threads 方式. 在 FPGA 上, 為了的達到較高的頻寬使用, 預設使用 Burst-Interleaved 方式, 若要自行分割也提供選項關閉, 並且在 Host Coding 中自行指定使用的 Bank.

OpenCL Library

若有常會使用的功能, 或是以 OpenCL 的方式實作較無效率的功能, Intel FPGA SDK 也提供以 OpenCL 或是 HDL 的方式實作再以 OpenCL Library 的方式導入. 而 SDK 內已經內部預先做好的 library, 開發者也能自己實作屬於自己的 library.


使用的方法如上流程圖所顯示, 實作的 OpenCL Kernel 與 OpenCL Library 透過 SDK Offline Compiler 最後產生 .aocx 的 FPGA 執行檔案.

Parallel Execution Model of AOCL Pipeline Stages

而功能的實作上, 要考量的與 CPU/GPU 上的 OpenCL 的考量不同, 對於 CPU/GPU 的實作上, 通常考量的是 SIMD 與指令使用上的考量, 能夠 SIMD 化, Exec Unit 的數目與種類, 或是否具有專屬加速指令. 但而對於 FPGA 考量的是流程上的平行度以及最後完成的長度, 再加上所使用的 resource 與 path 長短(這部分會留待第2篇), 而這些其最主要的考量為實作後整個 function 的 latency. 如下圖的 Kernel, 其從頭到尾需要 4 cycles.

2017年1月30日 星期一

從 "Three tips for maximising your SoC performance" 看軟體優化

先前 ARM 官方提供了 Free Webinar "Three tips to Maximize your SoC performance"
目前簡報的錄影與投影片都已經上線, 填入一些個人資訊即可閱覽

從報告的內容可以看出這個簡報是為了其 "System Guidance" 服務提高能見度
但是從設計原則中也有許多軟體人員可以作為借鏡並思考的地方
而以簡報中對於 SoC 設計上三個對於效能的考量為:
  • 最短化自 CPU 到記憶體路徑 - 每個 clocks cycle 都會歸屬到 CPU latency
  • 最大化系統頻寬 -  確認系統記憶體已經對於效能優化
  • 依流量種類管理流量 - 在系統符合即時性的情況下, 儘可能提供 CPU 優先權

而從軟體優化的角度來看, 這三點反映了對於效能上, 實作上的考量為:
  • 評估各種類型記憶體的 memory latency: 了解 memory latency 才能評估效能, 並且透過理解 cache 管理的機制來避免不當的行為所造成的過長存取時間. 此外理解 CPU 的硬體設計是透過何種方式來克服這問題, 以此來分析軟體上對於硬體差異在效能落差的影響.
  • 量測平台系統的頻寬: 透過 profiling 與 behavior analysis 能了解何時程式達到了 bandwidth bound, 才能採取對應的優化方式
  • 了解應用上可能同時競爭頻寬的硬體: 應用上所仰賴的 GPU 與 硬體加速器都會消耗頻寬
對於 latency 問題主要採取的策略有上述三項, 這三者的意義分別是紙上架構推估, 以軟體模擬
推測, 實際硬體上的量測, 三個不同階段

上圖為以 SGM-773 (System Guidance for Mobile) 的推估 memory latency 的方式, 若僅考量單次的存取, 請注意 85.1ns 這個推估時間約為 209 CPU cycles (也可以看出 CPU 內 cache 為 15 CPU cycles, 而出了 CPU 到了 CCI 則會超過 50 CPU cycles ). 這樣的計算通常能夠以此推敲 best case 與 worst case (考量各階段的 queue depth)的狀況, 也能得知架構上的物理限制, 與合理的數值範圍.

此外 LMBench 也是不錯的工具套件, 以 latency 來說是其中 lat_mem_rd 這個工具,下圖為使用 SGM-773 平台透過 LMBench 去量測後的結果

有興趣者可以先於自己的 ubuntu 平台上安裝 lmbench 套件, 接著執行下列指令:
/usr/lib/lmbench/bin/x86_64-linux-gnu/lat_mem_rd 16 128
其中 16 代表為 16MB, 也就是 bench 的大小上限為 16MB
而 128 所代表的是 stride 大小, 執行後可以看到類似下列的 log:
"stride=128
0.00049 2.361
0.00098 2.361
0.00195 2.362
0.00293 2.361
0.00391 2.361
0.00586 2.361
0.00781 2.361
0.01172 2.361
0.01562 2.362
0.02344 11.838
0.03125 11.819
0.04688 11.827
0.06250 11.816
0.09375 11.818
0.12500 11.825
0.18750 11.822
0.25000 11.817
0.37500 12.190
0.50000 12.182
0.75000 12.188
1.00000 12.188
1.50000 12.249
2.00000 19.838
3.00000 37.352
4.00000 38.298
6.00000 39.394
8.00000 39.920
12.00000 40.455
16.00000 40.674
這裡每行的兩個數字, 前者為該次測試的大小, 後者為測出的 latency 以 ns 計
上面測試的為個人使用 A8-5545M 平台所測出的 (L1: 16KB, L2: 2MB)
可以觀察到對應兩個 size 的 latency 轉折


事實上, LMBench 也提供了 bandwidth 的估測工具 - bw_mem
簡易執行指令如下:
/usr/lib/lmbench/bin/x86_64-linux-gnu/bw_mem  16K rdwr
也就是測試以 16MB 為大小的 Read/Write 頻寬, 有興趣亦可以調整前面的 16K 的數字
輸出有兩個數字:
0.016000 9801.85
前者為所要求測試的大小 (16KB 為 0.016M), 後者為 MB/s (也就是測出為 9891.85MB/s)
了解系統的能力, 對於針對平台的軟體優化有相當的幫助 (像是 tiling 與批次演算作法, 其分割大小能夠以 L2 大小做考量)

2017年1月25日 星期三

解析 Qualcomm Hexagon 680 架構 III - Thread Model and Execution Model

Thread Model

最後讓我們講述 Hexagon 680 的 Thread Model
在 Qualcomm 開發 DSP 的過程中一直都有 Hardware Threads 的支援
因此在談 Hexagon 680 的 Thread Model 前, 讓我們先來看看 V5x 的 HW Threads

上圖為 Qualcomm 在 2013 年 Hotchips 會議上投影片的 Page 11
從這圖可以看到 Hexagon 680 的 V6 基本上與 V5 架構差異不大
也是 4-slot VLIW, 兩個是 Data Unit, 兩個是 eXecution Unit
同個 DSP 核心具有3個 Thread Context, 在硬體設計上已考慮簡化多工上的需求

在 2015 年度 Hotchips 中 Qualcomm 簡報的 Page 8:
由此可見, V6 Scalar DSP 以 2Ghz 頻率運作
俱備 4 HW Threads 支援, 每個 thread 以 500Mhz 運作
而 HVX coprocessor 以 1Ghz 頻率運作, 具有兩個 context, 一樣每個 thread 以 500 Mhz 運作
因此同一時間可以有兩個 Thread 控制著 HVX, 另外兩個處理 Scalar 工作
但光以上看投影片, 不知道是否有人懷疑過, 或許 V6 DSP 為 quad core, 而 HVX 為 dual core?

這裡我們回到 HVX 官方文件當中:
從這張圖可以看出, V6 DSP 具有的是 1C4T, 而 HVX 為 1C2T
但是, 官方在投影片與到此中所說明的, 並非 Hexagon 680 所有的模式
如果還有印象, 在系列文II當中留了一些不確定:
1. HVX 對於 L2 的總寬度為 512b x2, 而非 1024b x 1
2. 對於 VPF/VRF 的描述, 我加了 "基本上" 這個有所保留的詞

再談這之前, 先來介紹 HVX 的 Vector Modes:

也就是 HVX 中支援了兩種不同長度的 Vector
  • 64B 模式, Vector 寬度為 512 bits
  • 128B 模式, Vector 寬度為 1024 BITS
而這樣的模式並不是像 NEON 中分為長短 vector 並且可以混用, 而是如上圖, 直接將所有的單元(VX, VPF, VRF), 以 512b 寬度的方式等分為二

所以除了 1C2T(128B) 的方式外, 事實上 HVX 還存在另一種 HW thread 的型態 - 1C4T (64B)
在這個模式下, 4 個 V6 HW Threads 各自都可以控制一個 64B Vector Mode 的 VX Thread Context. 所以 Hexagon 680 提供多元且彈性的運作模式. 而這也說明了 系列文II 當中沒交代清楚的部分.

Execution Mode


Qualcomm 曾經官方撰文介紹過 Hexagon DSP 的軟體開發
在 Qualcomm 裝置上, 在 DSP 上執行工作需要幾個步驟:
  1. 下載與安裝 Qualcomm Hexagon SDK
  2. 將你於 DSP 上實作的功能 API 以其 IDL (Interface Description Language) 方式描述.
  3. SDK 會自動產生 Header files 與 stub (下圖中應用處理器端, 也就是 ARM 處理器端提供給 client 的呼叫介面, client 通常是其他的 library 或是 executable)及 skel (DSP 端去呼叫你實作功能的進入點) 等相關函式庫.
  4. 在 DSP 端實作你的 API, 並與第3點中自動產生的 skel 函式庫 link 成為一個 shared object 並放置於該平台裝置中.
  5. 將你的 native APP 與 link 到第3點自動產生的 stub 來呼叫你的 API.
  6. (選擇性) 建立你的 API 專屬的 Java bindings 好讓 APK 能夠使用.
而這些背後所仰賴的機制的核心即是上圖所描述的 FastRPC

一旦 Native APP 呼叫所需的 function, 便會執行自動產生的 Stub lib, 它會 透過 ARM 與 DSP 間橋接的 Driver-Framework, 載入先前所產生的 DSP 端的 shared object, 而 DSP 端上的 OS 會動態地載入該 object, 並且以 Skel 介面處理來自 ARM 處理器端的請求.

Chisel 學習筆記 - Scala 與 Chisel 基礎語法

標題為筆記, 但這篇比較屬於心得 延續 上一篇 的環境建立, 這次計劃藉由 Jserv 最新的 課程安排 來學習 Chisel, 當然個人目標是能夠按照 Jserv 的課程規劃在 期限之內 完成 Lab 3, 由於個人並非 digital designer (現在這年紀也算老貓學...