項目地址:github.com/hijkzzz/neu…html
做爲深度學習領域的初學者,不少人會好奇TensorFlow和PyTorch等深度學習框架的工做原理是什麼。毫無例外,這些系統都使用了CUDA進行並行計算加速,在此我用CUDA實現了一個簡單的CNN網絡,方便你們學習和理解,並在無扭曲的MNIST數據集上實現了99.23%的準確率。git
這裏推薦《CUDA編程極簡入門教程》github
推薦知乎大神分享的《矩陣求導術》算法
在實現神經網絡以前,咱們須要設計一個存儲類,用於保存GPU上的參數和數據。這裏稱爲Storage類,爲了方便實現,咱們直接使用CUDA提供的thrust::device_vector(相似於std::vector)管理顯存上的動態數組。並增長一個std::vector保存Storage的形狀,能夠理解爲TensorFlow中Tensor的形狀。數組
神經網絡的實現大量用到矩陣乘法,因此CUDA並行加速的一個關鍵在於實現高效的並行矩陣乘法。這裏我直接使用了《CUDA編程極簡入門教程》中Shared Memory加速的矩陣乘法。實際上還能夠繼續優化,使效率大大提高。bash
設X爲輸入數據矩陣,其中每一行爲一個樣本。W爲參數矩陣,b爲偏置向量,L爲樣本平均損失。* 表示矩陣乘法,而非逐元素相乘,^T表示轉置:網絡
全鏈接
前向傳播
Y = X * W
反向傳播
dL/dX = dL/dY * W^T
dL/dW = X^T * dL/dY
偏置
前向傳播
Y = X + b
反向傳播
dL/db = sum(dL/dY, 0) 逐樣本梯度求和
複製代碼
爲了方便用矩陣乘法實現卷積,我參考了Caffe的卷積原理,即im2col:函數
基本的思想是把卷積運算展成矩陣乘法,因此能夠用並行加速的矩陣乘法高效實現卷積。設F爲卷積核參數,且形狀爲:channel_out*channel_in*kernel_width*kernel_height
,X爲一個輸入樣本形狀爲channel_in*width*height
,b爲偏置向量。
卷積
前向傳播
col = im2col(im) 根據im2col展開輸入圖
Y = F * col
反向傳播
dL/dF = dL/dY * col^T
dL/d_col = F^T * dL/dY
dL/d_im = col2im(dL/d_col)
偏置
前向傳播
Y = X + b 逐通道相加
反向傳播
dL/db = sum(sum(X, 2), 1) 對整個通道進行規約
複製代碼
Maxpool的反向傳播須要記錄池化前元素的位置,而後把反向梯度直接傳回
激活函數的前向反向傳播都是同樣的
ReLU
前向傳播
Y = relu(X)
反向傳播
dL/dX = relu'(X) element_mul dL/dY 逐元素相乘 其中relu'(x) = 1 if x > 0 else 0
Sigmoid
前向傳播
Y = sigmoid(X)
反向傳播
dL/dX = sigmoid'(X) element_mul dL/dY 逐元素相乘 其中 sigmoid'(x) = sigmoid(x) * (1 - sigmoid(x))
複製代碼
在工程實現上:爲了防止Softmax的分母溢出,通常使用LogSoftmax代替。設定1_n爲全爲1的列向量
Logsoftmax
正向傳播
Y = log_softmax(X) = x - log(exp(X) * 1_n) * 1_n^T
由前言中矩陣求導的方法可得
反向傳播
dL/dX = dL/dY - (dL/dY * 1_n * exp(x)) / (exp(x) * 1_n)
複製代碼
NLLLoss是平均負的對數似然損失,爲了配合LogSoftmax使用而實現。設Y爲樣本標籤矩陣,每一行爲一個樣本。N爲樣本數量
前向傳播
L = mean(sum(-log_P element_mul Y, 1), 0)
反向傳播
用矩陣乘法,L可表示爲 L = 1_n^T * ((-log_P element_mul Y) * 1_k) / N
由矩陣求導術可得
dL/d(log_P) = -Y / N
NLLLoss+LogSoftmax爲咱們常見的Softmax損失
將dL/d(log_P)帶入LogSoftmax梯度中可得softmax損失的梯度: softmax(X) - Y
複製代碼
爲了實現單獨的優化器,咱們須要在反向傳播的時候把梯度保存下來,而後用RMSProp算法進行統一的滑動平均計算新梯度。同理能夠很方便的實現Adam等優化器。
src
cuda CUDA源碼
minist MNIST DEMO
test
cuda CUDA源碼單元測試
CMakeLists.txt CMake編譯腳本
複製代碼
因爲篇幅有限,因此這裏只能去看GitHub上的實際代碼。每一個層都封裝爲了一個類,而且可調用connect函數鏈接層與層。
能夠經過CUDA提供的Visual Profiler能夠很方便的看出程序的性能瓶頸。
在個人實驗中發現80%的執行時間都在等待顯卡I/O,因此經過Pinned Memory以及合併傳輸/內存分配等方式使運行效率提高了數十倍。 其次是矩陣乘法還有較大的優化空間,不過總的來講在GTX1070上數十秒即可以跑完MNIST的6W個樣本,基本實現了個人目標。
編程一天,調試兩天,Debug是開發的一個困難並且重要的環節,掌握適當的工具的方法將事半功倍。CUDA提供的Nsight、cuda-memcheck都是很好的工具。固然printf+註釋大法也是屢試不爽。
網絡結構
conv 1 32 5 relu
maxpool 2
conv 32 64 5 relu
maxpool 2
conv 64 128 3 relu
fc 4 * 128 128 relu
fc 128 10 relu
softmax
nllloss
調參
shuffle = true
batch_size = 128
learning_rate = 0.003
L2 = 0.0001
beta = 0.99
準確率
1 epoch 93%
10 epochs 99.12%
30 epochs 99.23%
10s / epoch(GTX1070)
複製代碼