汽車長翅膀：GPU 是如何加速深度學習模型的訓練和推理過程的？ 原創 精華
編者按：深度學習的飛速發展離不開硬件技術的突破，而 GPU 的崛起無疑是其中最大的推力之一。但你是否曾好奇過，為何一行簡單的“.to('cuda')”代碼就能讓模型的訓練速度突飛猛進？本文正是為解答這個疑問而作。
作者以獨特的視角，將復雜的 GPU 并行計算原理轉化為通俗易懂的概念。從 CPU 與 GPU 的設計哲學對比，到 CUDA 編程的核心要素，再到具體的代碼實現，文章循序漸進地引領讀者把握 GPU 并行計算的精髓。特別是文中巧妙的比喻 —— 將 CPU 比作法拉利，GPU 比作公交車，這一比喻生動形象地詮釋了兩種處理器的特性。
這篇文章不僅回答了"為什么"，更指明了"如何做"，在當前人工智能技術飛速發展的背景下，理解底層技術原理的重要性不言而喻。這篇文章雖為入門級別的技術內容介紹，但也提到了更高級的優化技術和工具庫，指明了進一步的學習方向，具有一定的學習和參考價值。
作者 | Lucas de Lima Nogueira
編譯 | 岳揚
Image by the author with the assistance of AI (??https://copilot.microsoft.com/images/create??)
現如今，當我們提及深度學習時，人們自然而然地會聯想到通過 GPU 來增強其性能。
GPU（圖形處理器，Graphical Processing Units）起初是為了加速圖像（images）及 2D、3D 圖形（graphics）的渲染而生。但憑借其強大的并行運算能力，GPU 的應用范圍迅速拓展，已擴展至深度學習（deep learning）等應用領域。
GPU 在深度學習模型中的應用始于 2000 年代中后期，2012 年 AlexNet 的橫空出世更是將這種趨勢推向高潮。 AlexNet，這款由 Alex Krizhevsky、Ilya Sutskever 和 Geoffrey Hinton 共同設計、研發的卷積神經網絡，在 2012 年的 ImageNet Large Scale Visual Recognition Challenge (ILSVRC) 上一鳴驚人。這一勝利具有里程碑式的意義，它不僅證實了深度神經網絡在圖像分類領域（image classification）的卓越性能，同時也彰顯了使用 GPU 訓練大型模型的有效性。
在這一技術突破之后，GPU 在深度學習模型中的應用愈發廣泛，PyTorch 和 TensorFlow 等框架應運而生。
如今，我們只需在 PyTorch 中輕敲 .to("cuda")，即可將數據傳遞給 GPU，從而加速模型的訓練。但在實踐中，深度學習算法究竟是如何巧妙地利用 GPU 算力的呢？讓我們一探究竟吧！
深度學習的核心架構，如神經網絡、CNNs、RNNs 和 transformer，其本質都圍繞著矩陣加法（matrix addition）、矩陣乘法（matrix multiplication）以及對矩陣應用函數（applying a function a matrix）等基本數學操作展開。因此，優化這些核心運算，便是提升深度學習模型性能的關鍵所在。
那么，讓我們從最基礎的場景說起。想象一下，你需要對兩個向量執行相加操作 C = A + B。
可以用 C 語言簡單實現這一功能：
不難發現，傳統上，計算機需逐一訪問向量中的各個元素（elements），在每次迭代中按順序對每對元素進行加法運算。但有一點需要注意，各對元素間的加法操作互不影響，即任意一對元素的加法不依賴于其它任何一對。那么，若我們能同時執行這些數學運算，實現所有元素對（pairs of elements）的并行相加，效果會如何呢？
直接做法是借助 CPU 的多線程功能，并行執行所有數學運算。但在深度學習領域，我們需要處理的向量規模巨大，往往包含數百萬個元素。通常情況下，普通 CPU 只能同時處理十幾條線程。此時，GPU 的優勢便凸顯出來！目前的主流 GPU 能夠同時運行數百萬個線程，極大地提高了處理大規模向量中數學運算的效率。
01 GPU vs. CPU comparison
雖然從單次運算（single operation）的處理速度來看，CPU 或許略勝 GPU 一籌，但 GPU 的優勢在于其卓越的并行處理能力。究其根源，這一情況源于兩者設計初衷的差異。CPU 的設計側重于高效執行單一序列的操作（即線程（thread）），但一次僅能同時處理幾十個；相比之下，GPU 的設計目標是實現數百萬個線程的并行運算，雖有所犧牲單個線程的運算速度，卻在整體并行性能上實現了質的飛躍。
打個比方，你可以將 CPU 視作一輛炫酷的法拉利（Ferrari）跑車，而 GPU 則如同一輛寬敞的公交車。倘若你的任務僅僅是運送一位乘客，毫無疑問，法拉利（CPU）是最佳選擇。然而，如若當前的運輸需求是運送多位乘客，即使法拉利（CPU）單程速度占優，公交車（GPU）卻能一次容納全部乘客，其集體運輸效率遠超法拉利多次單獨接送的效率。由此可見，CPU 更適于處理連續性的單一任務，而 GPU 則在并行處理大量任務時展現出色的效能。
Image by the author with the assistance of AI (??https://copilot.microsoft.com/images/create??)
為了實現更出色的并行計算能力，GPU 在設計上傾向于將更多晶體管資源（transistors）投入到數據處理中，而非數據緩存（data caching）和流控機制（flow contro），這與 CPU 的設計思路大相徑庭。CPU 為了優化單一線程的執行效率和復雜指令集的處理，特意劃撥了大量的晶體管來加強這些方面的性能。
下圖生動地描繪了 CPU 與 GPU 在芯片資源分配上的顯著差異。
Image by the author with inspiration from CUDA C++ Programming Guide
(??https://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf??)
CPU 配備了高性能內核（powerful cores）與更為精妙的緩存內存架構（cache memory architecture）（消耗了相當多的晶體管資源），這種設計方案能夠極大地優化順序任務的執行速度。而圖形處理器（GPU）則著重于內核（cores）數量，以實現更高的并行處理能力。
現在已經介紹完這些基礎知識，那么在實際應用中，我們應如何有效利用并行計算的優勢呢？
02 Introduction to CUDA
當我們著手構建深度學習模型時，很可能會傾向于采用諸如 PyTorch 或 TensorFlow 這類廣受歡迎的 Python 開發庫。盡管如此，一個不爭的事實是，這些庫的核心代碼都是 C/C++ 代碼。另外，正如我們先前所提及的，利用 GPU 加快數據的處理速度往往是一種主流優化方案。此時，CUDA 的重要作用便凸顯出來！CUDA 是統一計算設備架構（Compute Unified Device Architecture）的縮寫，是英偉達（NVIDIA）為使 GPU 能夠在通用計算領域大放光彩而精心打造的平臺。與 DirectX 被游戲引擎用于圖形運算（graphical computation）不同，CUDA 使開發人員能夠將英偉達（NVIDIA）的 GPU 計算能力集成到通用軟件中，而不僅僅局限于圖形渲染。
為了實現這一目標，CUDA 推出了一款基于 C/C++ 的簡易接口（CUDA C/C++），幫助開發者調用 GPU 虛擬指令集（virtual intruction se）及執行特定操作（specific operations）（如在 CPU 與 GPU 間傳輸數據）。
在繼續深入技術細節之前，我們有必要澄清幾個 CUDA 編程的基礎概念和專業術語：
- host：特指 CPU 及其配套內存；
- device：對應 GPU 及其專屬內存；
- kernel：指代在設備（GPU）上運行的函數代碼；
因此，在一份使用 CUDA 撰寫的基本代碼（basic code）中，程序主體在 host (CPU) 上執行，隨后將數據傳遞給 device (GPU) ，并調用 kernels (functions) 在 device (GPU) 上并行運行。這些 kernels 由多條線程同時執行。運算完成后，結果再從 device (GPU) 回傳至 host (CPU) 。
話說回來，讓我們再次聚焦于兩組向量相加這個具體問題：
借助 CUDA C/C++，編程人員能夠創建一種被稱為 kernels 的 C/C++ 函數；一旦這些 kernels 被調用， N 個不同的 CUDA 線程會并行執行 N 次。
若想定義這類 kernel，可運用 ??__global__??? 關鍵字作為聲明限定符（declaration specifier），而若欲設定執行該 kernel 的具體 CUDA 線程數目，則需采用 ??<<<...>>>?? 來完成：
每個 CUDA 線程在執行 kernel 時，都會被賦予一個獨一無二的線程 ID，即 threadIdx，它可以通過 kernel 中的預設變量獲取。上述示例代碼將兩個長度（size）均為 N 的向量 A 和 B 相加，并將結果保存到向量 C 中。值得我們注意的是，相較于循環逐次處理成對加法的傳統串行方式，CUDA 的優勢在于其能夠并行利用 N 個線程，一次性完成全部加法運算。
不過，在運行上述這段代碼前，我們還需對其進行一次修改。切記，kernel 函數的運行環境是 device (GPU) ，這意味著所有相關數據均須駐留于 device 的內存之中。 要達到這一要求，可以借助 CUDA 提供的以下內置函數：
直接將變量 A、B 和 C 傳入 kernel 的做法并不適用于本情況，我們應當使用指針。在 CUDA 編程環境下，host 數組（比如示例中的 A、B 和 C）無法直接用于 kernel 啟動（<<<...>>>）。鑒于 CUDA kernels 的工作空間為 device 的內存（device memory），故需向 kernel 提供 device 指針（device pointers）（d_A、d_B 和 d_C），以確保其能在 device 的內存上運行。
除此之外，我們還需通過調用 cudaMalloc 函數在 device 上劃分內存空間，并運用 cudaMemcpy 實現 host 和 device 之間的數據傳輸。
至此，我們可在代碼中實現向量 A 和 B 的初始化，并在程序結尾處清理 CUDA 內存（cuda memory）。
另外，調用 kernel 后，務必插入 ??cudaDeviceSynchronize();?? 這一行代碼。該函數的作用在于協調 host 線程與 device 間的同步，確保 host 線程在繼續執行前，device 已完成所有先前提交的 CUDA 操作。
此外，CUDA 的錯誤檢測機制同樣不可或缺，這種檢測機制能協助我們及時發現并修正 GPU 上潛在的程序缺陷（bugs）。倘若忽略此環節，device 線程（CPU）將持續運行，而 CUDA 相關的故障排查則將變得異常棘手，很難識別與 CUDA 相關的錯誤。
下面是這兩種技術的具體實現方式：
要編譯和運行 CUDA 代碼，首先需要確保系統中已裝有 CUDA 工具包（CUDA toolkit）。緊接著，使用 nvcc —— NVIDIA CUDA 編譯器完成相關代碼編譯工作。
然而，當前的代碼尚存優化空間。在前述示例中，我們處理的向量規模僅為 N = 1000，這一數值偏小，難以充分展示 GPU 強大的并行處理能力。特別是在深度學習場景下，我們時常要應對含有數以百萬計參數的巨型向量。然而，倘若嘗試將 N 的數值設為 500000，并采用 <<<1, 500000>>> 的方式運行 kernel ，上述代碼便會拋出錯誤。因此，為了完善代碼，使之能順利執行此類大規模運算，我們亟需掌握 CUDA 編程中的核心理念 —— 線程層級結構（Thread hierarchy）。
03 Thread hierarchy（線程層級結構）
調用 kernel 函數時，采用的是 <<<number_of_blocks, threads_per_block>>> 這種格式（notation）。因此，在上述示例中，我們是以單個線程塊的形式，啟動了 N 個 CUDA 線程。然而，每個線程塊所能容納的線程數量都有限制，這是因為所有處于同一線程塊內的線程，都被要求共存于同一流式多處理器核心（streaming multiprocessor core），并共同使用該核心的內存資源。
欲查詢這一限制數量的具體數值，可通過以下代碼實現：
就作者當前使用的 GPU 而言，其單一線程塊最多能承載 1024 個線程。因此，為了有效處理示例中提及的巨型向量（massive vector），我們必須部署更多線程塊，以實現更大規模的線程并發執行。 同時，這些線程塊被精心布局成網格狀結構（grids），如下圖所展示：
??https://handwiki.org/wiki/index.php?curid=1157670?? (CC BY-SA 3.0)
現在，我們可以通過以下途徑獲取線程 ID：
于是，該代碼腳本更新為：
04 性能對比分析
下表展示了在處理不同大小向量的加法運算時，CPU 與 GPU 的計算性能對比情況。
Image by the author
顯而易見，GPU 的處理效能優勢，唯有在處理大規模向量時方能得以凸顯。此外，切勿忽視一件事，此處的時間對比僅僅考量了 kernel/function 的執行耗時，而未將 host 和 device 間數據傳輸所需的時間納入考慮范圍。盡管在大多數情況下，數據傳輸的時間開銷微不足道，但就我們目前僅執行簡易加法運算（simple addition operation）的情形而言，這部分時間消耗卻顯得相對可觀。因此，我們應當銘記，GPU 的計算性能，僅在面對那些既高度依賴計算能力又適合大規模并行處理的任務時，才能得以淋漓盡致地展現。
05 多維線程處理（Multidimensional threads）
現在，我們已經知道如何提升簡單數組操作（simple array operation）的性能了。然而，在處理深度學習模型時，必須要處理矩陣和張量運算（matrix and tensor operations）。在前文的示例中，我們僅使用了內含 N 個線程的一維線程塊（one-dimensional blocks）。然而，執行多維線程塊（multidimensional thread blocks）（最高支持三維）同樣也是完全可行的。因此，為了方便起見，當我們需要處理矩陣運算時，可運行一個由 N x M 個線程組成的線程塊。還可以通過 row = threadIdx.x 來確定矩陣的行索引，而 col = threadIdx.y 則可用來獲取列索引。此外，為了簡化操作，還可以使用 dim3 變量類型定義 number_of_blocks 和 threads_per_block。
下文的示例代碼展示了如何實現兩個矩陣的相加運算。
此外，我們還可以將此示例進一步拓展，實現對多個線程塊的處理：
此外，我們也可以用同樣的思路將這個示例擴展到三維運算（3-dimensional operations）操作的處理。
上文已經介紹了處理多維數據（multidimensional data）的方法，接下來，還有一個既重要又容易理解的概念值得我們學習：如何在 kernel 中調用 functions。 一般可以通過使用 ??__device__??? 聲明限定符（declaration specifier）來實現。這種限定符定義了可由 device （GPU）直接調用的函數（functions）。因此，這些函數僅能在 ??__global__??? 或其他 ??__device__?? 函數中被調用。下面這個示例展示了如何對一個向量進行 sigmoid 運算（這是深度學習模型中極其常見的一種運算方式）。
至此，我們已經掌握了 CUDA 編程的核心概念，現在可以著手構建 CUDA kernels 了。對于深度學習模型而言，其實質就是一系列涉及矩陣（matrix）與張量（tensor）的運算操作，包括但不限于求和（sum）、乘法（multiplication）、卷積（convolution）以及歸一化（normalization ）等。舉個例子，一個基礎的矩陣乘法算法，可以通過以下方式實現并行化：
我們可以注意到，在 GPU 版本的矩陣乘法算法中，循環次數明顯減少，從而顯著提升了運算處理速度。下面這張圖表直觀地展現了 N x N 矩陣乘法在 CPU 與 GPU 上的性能對比情況：
Image by the author
我們會發現，隨著矩陣大小（matrix size）的增大，GPU 在處理矩陣乘法運算時的性能提升幅度更大。
接下來，讓我們聚焦于一個基礎的神經網絡模型，其核心運算通常表現為 y = σ(Wx + b)，如下圖所示：
Image by the author
上述運算主要涉及矩陣乘法（matrix multiplication）、矩陣加法（matrix addition）以及對數組施加函數變換（applying a function to an array）。如若你已掌握這些并行化處理技術，意味著你現在完全具備了從零構建、并在 GPU 上構建神經網絡的能力！
06 Conclusion
本文我們探討了通過 GPU processing （譯者注：使用 GPU進行數據處理和計算。）提升深度學習模型效能的入門概念。不過，有一點還需要指出，本文所介紹的內容僅僅是皮毛，背后還隱藏著很多很多更深層次的東西。PyTorch 和 Tensorflow 等框架實現了諸多高級性能優化技術，涵蓋了 optimized memory access、batched operations 等復雜概念（其底層利用了基于 CUDA 的 cuBLAS 和 cuDNN 等庫）。 但愿這篇文章能夠讓各位讀者對使用 .to("cuda") 方法，在 GPU 上構建、運行深度學習模型時的底層原理，有個初步的了解。
Thanks so much for reading! ??
Lucas de Lima Nogueira
??https://www.linkedin.com/in/lucas-de-lima-nogueira/??
END
原文鏈接：
