作者：Lawliet
翻譯：仿佛若有光
前言：
幾個月前，我根據 Simoncelli 2016 年的論文編寫了自己的自動編碼器，用于研究目的。一開始，我想使用一些流行的深度學習架構（例如 Tensor Flow、Caffe2 或 MXNet）來做我的實驗。然而，在對所有這些架構進行了幾周的調查之後，我發現了一個非常令人頭疼的問題——可擴充性。我不是說這些架構設計得不好，而是不允許使用者開發第三方算子，就像寫一個插件一樣，你給我一個沒有任何參數的函數。那麼改變函數行為的唯一方法就是修改源代碼，由于文檔組織不善，這無疑是一個巨大的工程。（這似乎是開源軟體的通病。）是以，由于不常見的算子 GDN 并未包含在所有這些架構中，是以設計一個新架構似乎是唯一的解決方案。
點個關注，專注于計算機視覺的技術總結和分享
GDN
這個算子是這個理論中的核心非線性函數，表達式如下（公式不重要，如果你不喜歡這些該死的符号，你可以直接跳過這一節。）：
上标(k)和(k+1)表示層數，w和u是多通道圖像的輸入和輸出，下标i是通道數。β 和 γ 是我要訓練的參數。假設我們有 N 個通道，那麼 γ 是一個 N × N 矩陣，β 是一個 N × 1 向量。乍一看，這個功能與 cudnn 和所有深度學習架構都很好地支援的批量歸一化 (BN) 或局部響應歸一化 (LRN) 非常相似。但相信我，不要讓你的眼睛欺騙你。這是非常不同的。（注意大除法是元素除法。）
前向不會消耗太多計算能力，而後向會消耗我 GPU 的大部分能量。現在讓我們看看後面。我需要計算 3 個梯度，∇β、∇γ 和 ∇u。
我知道人們第一次看到這個的感覺，因為我第一次看到這個怪物時也想自殺。 但如果我能為所有這些狗屎畫一幅畫，你會感覺更舒服。
首先，我們可以很容易地注意到輸入可以看作是一個長度為 m x n 的向量。其次，(blabla...)^(-3/2) 出現在所有這些梯度中。這意味着我們可以隻計算該術語 1 次，并将它們緩存以備後用。我們稱其為“(blabla...)^(-1/2)”矩陣 D 。最後，δ 是傳播到前一層的誤差。
Fig 1. Computation of γ
經過一些簡化，它更清楚了，對吧？ 我知道仍然需要一些解釋。 對于等式的右側，每個矩形都是由我們上面提到的矩陣堆疊而成的向量。 D 是 GDN 公式中的分母項，還記得我們剛剛提到的“(blabla...)^(-1/2)”嗎？
與一些進階算法不同，這種計算對大多數人來說非常直覺，我們可以輕松編寫 CPU 程式來處理它。隻要稍微了解一下 CUDA，每個人都可以将他們的 CPU 代碼移植到 GPU。但是，如果您可以選擇不同的組織來啟動核心，則速度會有很大的不同。
1. 不僅僅是天真的算法。
我稱這種方法“不隻是天真”是因為這是我用過的第一種方法。即使使用小尺寸圖像作為輸入，它也幾乎耗盡了我所有的 GPU 記憶體，并實作了最慢的性能。沒有利用任何記憶體重用，我隻是垂直和水準複制所有這些小矩形以獲得更大的矩陣，如下圖所示，并啟動許多一維組織的核心。然後将它們相加。
Fig 2. Less than naive Algo.
該算法唯一的優點是不需要在每個CUDA線程中計算索引，因為線程id隻是唯一對應的記憶體索引。是以你需要做的就是一些乘法，然後使用 cublas 将每個小彩色矩形與 1 向量（一個充滿所有 1 的向量）的點積相加。但是正如你所看到的，矩形的大小并不像我這裡畫的那麼小，大小和圖像一樣。對于這張圖檔中的每個向量，大小将為 N x N x imageSize x batchSize。很明顯，我們浪費了 (N-1) x N x imageSize x batchSize x 4 個位元組，更不用說浪費在通路所有這些備援全局記憶體上的時間了。
2. 樸素算法。
對于第一種算法，我每次疊代隻能在我的網絡中訓練不到 4 張大小為 128 x 128 的圖像，時間幾乎為 2 秒。（我的 GPU 是 GTX 1080。）這個現實迫使我改進我的算法，否則，我必須等待近 2 個月才能得到我的結果。
因為我需要啟動的核心數量肯定比我GPU中的CUDA核心多很多，是以不管我用什麼方法，cuda驅動都會把這些任務序列化。然後我決定不複制所有這些記憶。相反，我将啟動 N x 一維組織的 N x imageSize 核心 N 次（N 是通道總數）。
Fig 3. Without memory replication
可以看出，改進是顯而易見的。因為，我們不再需要大量複制資料。 GPU 中的全局記憶體通路非常昂貴。記憶體通路模式也很簡單，因為當您獲得線程 id 時，隻需使用一個 mod 操作就可以獲得記憶體索引（記憶體索引 = 線程 id % imageSize）。但是，在這種方法中，由于核心仍然是一維組織的，并且我們使用for循環來啟動所有這些核心，那麼我們可能無法從GPU更智能的排程算法中受益，盡管我已經嘗到了血的滋味.現在，通過這個小小的改變，2 個月的訓練時間可以縮短到将近 2 周。
3. 更智能的組織算法。
到目前為止，我還沒有考慮過共享記憶體的威力，因為對我來說，通常設計一個好的核心模式是枯燥和頭痛的。顯然，一維核心模式是最容易編寫的代碼。然而，更好的性能值得更仔細的設計。令我驚訝的是，本節中的算法實作了第二個算法的 3 倍速度。
回到圖 1，可以看到前 3 個右側矩陣的第一行 δ0、w0 和 D0 是相同的。是以，我們可以在一個塊中計算一行 γ，對于每個塊我們可以啟動 imageSize 個線程，并且對于每個線程我們可以使用 for 循環計算所有通道。
Fig 5. Computation in one block
是以從圖 5 來看，将 δ0、w0 和 D0 放在共享記憶體中是非常直覺的，而對于線程 i，它從 0 到 N-1 讀取 N 個通道中的一個像素與 δ0、w0 和 D0 相乘 分享回憶。僞代碼如下：
blockId = blockIdx.x; threadId = threadIdx.x;shareDelta <- delta[blockId]; shareW <- W[blockId]; shareD <- D[blockId]; _synchronize();for(i = 0; i < N-1; i++) { result[threadIdx i*imgSize] = shareDelta[threadId] * shareW[threadId] * shareD[threadId] * W[threadId + i*imgSize]; } Algo 2 選擇行主計算而不是列主計算是因為在一個網格中計算一行，我們可以共享 3 個向量 δ0、w0 和 D0。但是如果我們像在 Algo 中那樣計算一列，我們隻能共享 1 個向量 w0。（再次參見圖 1。）。
在這段代碼片段中，沒有 if ... else ... 塊。這在并行計算中非常重要。因為所有線程都是并行運作的，理想的情況是所有這些線程同時完成它們的工作。但是如果有 if ... else ... 阻塞，分支會讓這些線程做不同的任務，以便它們在不同的時間完成。然後計算時間将由最慢的線程決定。
無索引計算也是一個優勢。通過設計一維模式，我們必須使用線程id來計算記憶體索引，但這裡不需要将blockId和threadId轉換為一維記憶體索引來通路資料。
最後，因為我的資料存儲在列major中，這意味着，像向量δ0一樣，這個向量中的所有元素都是連續存儲的。是以它受益于全局記憶體合并機制。全局記憶體也是cuda中的一個重要概念。
在硬體方面，16個cuda核心被組織在一個warp中。當其中一個線程通路資料時，例如上圖中的 a1，資料總線不僅會傳輸 a1，還會将 a1~a32 傳輸到緩存中，以加速其他 15 個核心的資料通路。是以，當我讀取全局資料以共享記憶體時，每 32 個位元組我隻讀取一次，所有其他位元組都從緩存中讀取，速度快了數百。多虧了時空局域性理論。
4. 多一點改進
今天突然發現其實我不需要共享記憶體，但是可以使用const記憶體。因為對于向量δ0、w0和D0，一個block中的每個線程隻需要通路一次。是以在for循環之前，我們實際上可以将元素緩存在const記憶體中。另一個糖是因為每個線程隻通路一個元素，不需要線程同步。
代碼如下：
blockId = blockIdx.x; threadId = threadIdx.x;const float constDelta = delta[blockId * imgSize + threadId]; const float constW = W[blockId * imgSize + threadId]; const float constD = D[blockId * imgSize + threadId];for(i = 0; i < N-1; i++) { result[threadIdx + i*imgSize] = constDelta * constW * constD * W[threadId + i*imgSize]; } 從上面的代碼可以看出，constDelta、constW、constD可以從本地記憶體中重複使用N次，本地記憶體總是存儲在本地寄存器中。是以，帶寬大于共享記憶體。
Reduce Operation
我講的所有算法都沒有完成，因為我從上述算法中得到的實際上都是原始γ，如下所示：
我需要在左側累積每個向量以獲得一個元素。第一個選擇是 cublas API，cublasSsbmv。此函數将進行矩陣向量乘法。是以我們可以把左邊的向量看成一個矩陣，将它與一個全1向量相乘，得到γ的一行梯度。并重複N次以獲得最終結果。但我注意到還有其他 API cublasSgemmBatched。此函數可以進行批量矩陣向量乘法。然後我做了一個實驗來測試哪個更快：
N 個矩陣向量乘法 VS 批處理矩陣向量乘法的 for 循環。
結果表明for循環要快得多。但是我不知道原因，也許是因為我這裡的 N 太小（N = 256）。
我不會展示如何計算 ∇β 和 ∇u，因為它們類似于 ∇γ。我知道必須有比我更進一步的優化或更好的設計。CUDA 優化對于不深入了解 GPU 組織的人來說通常是困難的。熟悉 CPU 的程式員總是受益于現代作業系統和強大的編譯器。然而，GPU 在編寫足夠的代碼方面與 CPU 有很大不同和複雜性，盡管它比以前使用圖形着色器進行計算要友善得多。生态環境的完善還需要幾年時間。
原文連結：
https://medium.com/@Lawliet0320/ramble-in-cuda-optimization-8fbbcf81e7c5
本文來源于公衆号 CV技術指南 的論文分享系列。
歡迎關注公衆号 CV技術指南 ，專注于計算機視覺的技術總結、最新技術跟蹤、經典論文解讀。
在公衆号中回複關鍵字 “技術總結” 可擷取以下文章的彙總pdf。
其它文章
計算機視覺中的自注意力
經典論文系列--膠囊網絡：新的深度學習網絡
綜述專欄 | 姿态估計綜述
漫談CUDA優化
為什麼GEMM是深度學習的核心
使用深度神經網絡為什麼8位足夠？
經典論文系列 | 目标檢測--CornerNet & 又名 anchor boxes的缺陷
如何看待人工智能的泡沫
使用Dice loss實作清晰的邊界檢測
PVT--無卷積密集預測的多功能backbone
CVPR2021 | 開放世界的目标檢測
Siamese network總結
視覺目标檢測和識别之過去，現在及可能
在做算法工程師的道路上，你掌握了什麼概念或技術使你感覺自我提升突飛猛進？
計算機視覺專業術語總結(一)建構計算機視覺的知識體系
欠拟合與過拟合技術總結
歸一化方法總結
論文創新的常見思路總結
CV方向的高效閱讀英文文獻方法總結
計算機視覺中的小樣本學習綜述
知識蒸餾的簡要概述
優化OpenCV視訊的讀取速度
NMS總結
損失函數技術總結
注意力機制技術總結
特征金字塔技術總結
池化技術總結
資料增強方法總結
CNN結構演變總結（一）經典模型
CNN結構演變總結（二）輕量化模型
CNN結構演變總結（三）設計原則
如何看待計算機視覺未來的走向
CNN可視化技術總結（一）特征圖可視化
CNN可視化技術總結（二）卷積核可視化
CNN可視化技術總結（三）類可視化
CNN可視化技術總結（四）可視化工具與項目