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 correspond16-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)));

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 的效能

"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 或類似的實作)

"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 內部實作方式的說明

"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 中使用所需了解的基本型別, 介面, 執行方式以及範例

SIMD 加速的實例 - NEON acceleration of tiling in VC4 of RPi3

今日增加了 Matrix Multiplication 對於 GCC 的支援, 與相關效能數字後
(基本上 GCC 的 Vector 優化較 Clang 差)
一則關於NEON的新聞引起了我的注意
主要是 Broadcom 的開發者在OSS Driver 中使用 NEON 加速
其部分 Blog 原文引述如下:
"My first hope was that I could load a full cacheline with NEON's VLD.  VLD1 only loads up to 1 "quadword" (16 bytes) at a time, so that doesn't seem like much help.  VLD4 can load 64 bytes like we want, but it also turns AOS data into SOA in the process, and there's no corresponding "SOA-back-to-AOS store 8 or 16 bytes at a time" like we need to do to get things back into the CPU's strided representation.  I tried VLD4+VST4 into a temporary, then doing my old untiling path on the cached temporary, but that still left me a few percent slower on loads than not doing any of this work at all.

Finally, I hit on using the VLDM instruction.  It seems to be intended for stack loads/stores, but we can also use it to get 64 bytes of data in from memory untouched into NEON registers, and then I can use 4 (32bpp) or 8 (8 or 16bpp) VST1s to store it to the CPU side.  With this, we get a 208.256% +/- 7.07029% (n=10) improvement to GetTexImage performance at 1024x1024.  Doing the same NEON code for stores gave a 41.2371% +/- 3.52799% (n=10) improvement, probably mostly due to not calling into memcpy and having it go through its size/alignment-based memcpy path choosing process."

由於 NEON 中的 Vector load 指令 VLD1 "只能"一次讀入16bytes, 因此原本計劃採用了 VLD4 期望來一次大量的讀入資料, 但 VLD4 這個指令會做類似 AOS(Array of Structure) 轉 SOA (Structure of Array) 的動作(其實是 de-interleave), 而且並沒有一次 SOA轉回 AOS 的指令, 該開發者嘗試過 VLD4+VST4 然後再利用原有非 tiling 的作法, 但是還是比慢上了許多. (其實這樣使用 VLD4 的作法, 基本上有 register spilling 的疑慮, 再者 VLD4 基本上還是透過 L/S unit, 基本上由於 ARM 的 NEON unit 的實作, 基本上並不會比 VLD1 好)

最後他採用了 VLDM 指令(這個指令可以參考這篇), 基本上這指令帶來的好處是不用會 touch 到 NEON registers, 並且能夠處理的長度也較長, 最終該開發者在讀取的 GetTexImage 得到了 208.256% 的效能改善,  並且在寫入獲得了 41.2371% 的改善. 開發者認為這或許歸功於避免直接呼叫 memcpy, 並採取 size/alignment-based 的對應流程.