論文閱讀筆記：Espresso

摘要

這篇文章的主要貢獻是開源了一個基於C與CUDA的self-contained的神經網路量化的庫，利用這個庫能夠構建出在前向時實際能夠達到加速的程序。其中主要利用的技術是bit-packing、register-blocking以及基本的xnor,bit-count和bit-shift。在我實際測試中，能夠達到X10以上的加速比。這篇文章主要按照一下的順序介紹:

• 基礎知識介紹
• 論文要點回顧
• 實驗結果

由於原文提供的代碼CmakeList.txt文件有些問題，且部分測試代碼並未公布。目前我還在整理代碼，後續我會將測試代碼即更新過的CmakeList.txt文件分享到我的項目文件夾

基礎知識介紹

Cmake語法

• add_subdirectory():指定在指定的subdirectory中再繼續找CmakeList.txt
• add_library():指定編譯目標是編譯出一個庫
• find_package(CUDA):找某個包，例如CUDA

GEMM運算的定義

對於矩陣運算:

，其中alpha，beta是實數，A，B，C是矩陣

可以利用SGEMM( TRANSA, TRANSB, M, N, K, alpha, A, LDA, B, LDB, beta, C, LDC)完成相應的運算。這樣其實可以通用地定義了矩陣的點積和加法等多種運算。推薦擴展閱讀這篇文章函數（其實並不是函數，而是一個規範）的參數解釋如下:

• transa、transb:分別表示A，B是否需要轉置
• M:A和矩陣C的行數
• N:B和C的列數
• K:A和B的列數
• alpha,beta:實數係數
• A,B,C:矩陣A,B，一般用 T **A表示，T屬於{uint8,uint64,float}
• LDA,LDB,LDC:遞增步長，及第一維坐標增加1，相應的地址增量是多少，用於區別行列主序

CUDA中提供的高效率的函數/操作

• bit-count:
  __device__ ? int __popcll ( unsigned long long int x )

  Count the number of bits that are set to 1 in a 64 bit integer.

• 位運算
• >>:算術右移
• >>>:邏輯右移
• <<,<<<:也類似的是算術、邏輯左移
• ^:xnor即同或運算

論文要點回顧

實現量化的流程

• bit-packing:將64個參數量化（即按符號轉化為0/1）之後放入一個64位的整型，整型中的每一位表示一個參數。
• 基於bit-packing向量點積
• 可以證明兩個個量化後向量的點積運算可以等價與異或運算與bitcount運算的組合。
• 但需要滿足向量的長度是64的倍數
• 注意:向量運算完的結果還是利用浮點型存儲
• 量化GEMM:兩個矩陣相乘可以看成對應的行向量與列向量相乘，再求和。即可以轉化為之前定義的量化的向量積運算。

• 由此可以按照一般的方式定義dense、conv、pool等層級

Hybrid DNNs:

• 論文分別提供了CPU、GPU、 GPUOPT的代碼。
• CPU上的全精度的kernel
• GPU全精度的kernel
• GPU上的BDNN的kernel

Ｍemory layout:

• 在內存中張量的存儲形式，類似矩陣中對行主序和列主序的約定
• 在L維（channel）上進行bit-packing，在L==1時，在N維上
• 這樣的設計在lifting、UNROLL的時候會比較快

input-binarizartion:

• 即對輸入在bit-plane上展開，針對每個bit做運算，最終再合併
• 以[0,255]的顏色為例，像素值若是127
• 則變成0111,1111
• 相乘的結果再根據位置做weight
• 實際上沒有實現(參見src/kernel/pinputl.cu)

bit-packing帶來的問題:

• 若不滿足64位填不滿的填-1可能造成很嚴重的誤差比如
• 因此，嚴格要求進行矩陣運算的張量A(D,M,N,L)中M*N*L是64的倍數。由此對在此之上定義的層級、網路都有限制。但對一般的網路而言影響不大，因為一般都是64的倍數。

實驗結果

GEMM速度

• configuration :1060,8192*8192相乘,Batchsize=1,Average on 10
• 不包含設置時間，啟動時間
• 0.8us pgemm 得到float數,9us得到量化數
• 16.6us sgemm得到float數
• 注意，當重複測試10次時速度較快，隨重複測試次數增加速度減慢

在minist上inference的性能

• configuration: 1060,cuda 8.0、mnist、batch_size=512、average on 10000 iters、mlp（無卷積）
• 不包含啟動時間、不包含數據讀取時間
• pytorch上的全精度相同結構的對照:173.22 us/input
• espresso:13 us/input

正確性驗證

• 全精度全連接層（輸入1*2*3*1,權重2*6,輸出2*1）

輸入:35.000000 23.000000 34.000000 -43.000000 -45.000000 62.000000 | 113.000000 124.000000 121.000000 -7.000000 -21.000000 -46.000000 |權重:-108.000000 -109.000000 105.000000 | 98.000000 -83.000000 -47.000000 |輸出:-6110.000000 -9796.000000 |正確

• 量化的全連接層（對輸入的形狀是有要求的！）

經測試（在test.cu中）正確

推薦閱讀：