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 影響程式的效能甚巨

Lookup Table 在 NEON 中的處理

在 SIMD Programming 中由於希望能夠每個 lane 有一致的行為, 因此有一些事情是不容易達到的 而 Lookup Table (LUT) 即是其中之一 但若是特定條件之下, 還是有可能透過 NEON 加速 而這個 直接前提是 8bit LUT (當然...