編者按：深度學(xué)習(xí)的飛速發(fā)展離不開硬件技術(shù)的突破，而 GPU 的崛起無(wú)疑是其中最大的推力之一。但你是否曾好奇過(guò)，為何一行簡(jiǎn)單的“.to('cuda')”代碼就能讓模型的訓(xùn)練速度突飛猛進(jìn)？本文正是為解答這個(gè)疑問(wèn)而作。

作者以獨(dú)特的視角，將復(fù)雜的 GPU 并行計(jì)算原理轉(zhuǎn)化為通俗易懂的概念。從 CPU 與 GPU 的設(shè)計(jì)哲學(xué)對(duì)比，到 CUDA 編程的核心要素，再到具體的代碼實(shí)現(xiàn)，文章循序漸進(jìn)地引領(lǐng)讀者把握 GPU 并行計(jì)算的精髓。特別是文中巧妙的比喻 —— 將 CPU 比作法拉利，GPU 比作公交車，這一比喻生動(dòng)形象地詮釋了兩種處理器的特性。

這篇文章不僅回答了"為什么"，更指明了"如何做"，在當(dāng)前人工智能技術(shù)飛速發(fā)展的背景下，理解底層技術(shù)原理的重要性不言而喻。這篇文章雖為入門級(jí)別的技術(shù)內(nèi)容介紹，但也提到了更高級(jí)的優(yōu)化技術(shù)和工具庫(kù)，指明了進(jìn)一步的學(xué)習(xí)方向，具有一定的學(xué)習(xí)和參考價(jià)值。

作者 | Lucas de Lima Nogueira

編譯 | 岳揚(yáng)

Image by the author with the assistance of AI (??https://copilot.microsoft.com/images/create??)

現(xiàn)如今，當(dāng)我們提及深度學(xué)習(xí)時(shí)，人們自然而然地會(huì)聯(lián)想到通過(guò) GPU 來(lái)增強(qiáng)其性能。

GPU（圖形處理器，Graphical Processing Units）起初是為了加速圖像（images）及 2D、3D 圖形（graphics）的渲染而生。但憑借其強(qiáng)大的并行運(yùn)算能力，GPU 的應(yīng)用范圍迅速拓展，已擴(kuò)展至深度學(xué)習(xí)（deep learning）等應(yīng)用領(lǐng)域。

GPU 在深度學(xué)習(xí)模型中的應(yīng)用始于 2000 年代中后期，2012 年 AlexNet 的橫空出世更是將這種趨勢(shì)推向高潮。 AlexNet，這款由 Alex Krizhevsky、Ilya Sutskever 和 Geoffrey Hinton 共同設(shè)計(jì)、研發(fā)的卷積神經(jīng)網(wǎng)絡(luò)，在 2012 年的 ImageNet Large Scale Visual Recognition Challenge (ILSVRC) 上一鳴驚人。這一勝利具有里程碑式的意義，它不僅證實(shí)了深度神經(jīng)網(wǎng)絡(luò)在圖像分類領(lǐng)域（image classification）的卓越性能，同時(shí)也彰顯了使用 GPU 訓(xùn)練大型模型的有效性。

在這一技術(shù)突破之后，GPU 在深度學(xué)習(xí)模型中的應(yīng)用愈發(fā)廣泛，PyTorch 和 TensorFlow 等框架應(yīng)運(yùn)而生。

如今，我們只需在 PyTorch 中輕敲 .to("cuda")，即可將數(shù)據(jù)傳遞給 GPU，從而加速模型的訓(xùn)練。但在實(shí)踐中，深度學(xué)習(xí)算法究竟是如何巧妙地利用 GPU 算力的呢？讓我們一探究竟吧！

深度學(xué)習(xí)的核心架構(gòu)，如神經(jīng)網(wǎng)絡(luò)、CNNs、RNNs 和 transformer，其本質(zhì)都圍繞著矩陣加法（matrix addition）、矩陣乘法（matrix multiplication）以及對(duì)矩陣應(yīng)用函數(shù)（applying a function a matrix）等基本數(shù)學(xué)操作展開。因此，優(yōu)化這些核心運(yùn)算，便是提升深度學(xué)習(xí)模型性能的關(guān)鍵所在。

那么，讓我們從最基礎(chǔ)的場(chǎng)景說(shuō)起。想象一下，你需要對(duì)兩個(gè)向量執(zhí)行相加操作 C = A + B。

可以用 C 語(yǔ)言簡(jiǎn)單實(shí)現(xiàn)這一功能：

不難發(fā)現(xiàn)，傳統(tǒng)上，計(jì)算機(jī)需逐一訪問(wèn)向量中的各個(gè)元素（elements），在每次迭代中按順序?qū)γ繉?duì)元素進(jìn)行加法運(yùn)算。但有一點(diǎn)需要注意，各對(duì)元素間的加法操作互不影響，即任意一對(duì)元素的加法不依賴于其它任何一對(duì)。那么，若我們能同時(shí)執(zhí)行這些數(shù)學(xué)運(yùn)算，實(shí)現(xiàn)所有元素對(duì)（pairs of elements）的并行相加，效果會(huì)如何呢？

直接做法是借助 CPU 的多線程功能，并行執(zhí)行所有數(shù)學(xué)運(yùn)算。但在深度學(xué)習(xí)領(lǐng)域，我們需要處理的向量規(guī)模巨大，往往包含數(shù)百萬(wàn)個(gè)元素。通常情況下，普通 CPU 只能同時(shí)處理十幾條線程。此時(shí)，GPU 的優(yōu)勢(shì)便凸顯出來(lái)！目前的主流 GPU 能夠同時(shí)運(yùn)行數(shù)百萬(wàn)個(gè)線程，極大地提高了處理大規(guī)模向量中數(shù)學(xué)運(yùn)算的效率。

01 GPU vs. CPU comparison

雖然從單次運(yùn)算（single operation）的處理速度來(lái)看，CPU 或許略勝 GPU 一籌，但 GPU 的優(yōu)勢(shì)在于其卓越的并行處理能力。究其根源，這一情況源于兩者設(shè)計(jì)初衷的差異。CPU 的設(shè)計(jì)側(cè)重于高效執(zhí)行單一序列的操作（即線程（thread）），但一次僅能同時(shí)處理幾十個(gè)；相比之下，GPU 的設(shè)計(jì)目標(biāo)是實(shí)現(xiàn)數(shù)百萬(wàn)個(gè)線程的并行運(yùn)算，雖有所犧牲單個(gè)線程的運(yùn)算速度，卻在整體并行性能上實(shí)現(xiàn)了質(zhì)的飛躍。

打個(gè)比方，你可以將 CPU 視作一輛炫酷的法拉利（Ferrari）跑車，而 GPU 則如同一輛寬敞的公交車。倘若你的任務(wù)僅僅是運(yùn)送一位乘客，毫無(wú)疑問(wèn)，法拉利（CPU）是最佳選擇。然而，如若當(dāng)前的運(yùn)輸需求是運(yùn)送多位乘客，即使法拉利（CPU）單程速度占優(yōu)，公交車（GPU）卻能一次容納全部乘客，其集體運(yùn)輸效率遠(yuǎn)超法拉利多次單獨(dú)接送的效率。由此可見(jiàn)，CPU 更適于處理連續(xù)性的單一任務(wù)，而 GPU 則在并行處理大量任務(wù)時(shí)展現(xiàn)出色的效能。

Image by the author with the assistance of AI (??https://copilot.microsoft.com/images/create??)

為了實(shí)現(xiàn)更出色的并行計(jì)算能力，GPU 在設(shè)計(jì)上傾向于將更多晶體管資源（transistors）投入到數(shù)據(jù)處理中，而非數(shù)據(jù)緩存（data caching）和流控機(jī)制（flow contro），這與 CPU 的設(shè)計(jì)思路大相徑庭。CPU 為了優(yōu)化單一線程的執(zhí)行效率和復(fù)雜指令集的處理，特意劃撥了大量的晶體管來(lái)加強(qiáng)這些方面的性能。

下圖生動(dòng)地描繪了 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 配備了高性能內(nèi)核（powerful cores）與更為精妙的緩存內(nèi)存架構(gòu)（cache memory architecture）（消耗了相當(dāng)多的晶體管資源），這種設(shè)計(jì)方案能夠極大地優(yōu)化順序任務(wù)的執(zhí)行速度。而圖形處理器（GPU）則著重于內(nèi)核（cores）數(shù)量，以實(shí)現(xiàn)更高的并行處理能力。

現(xiàn)在已經(jīng)介紹完這些基礎(chǔ)知識(shí)，那么在實(shí)際應(yīng)用中，我們應(yīng)如何有效利用并行計(jì)算的優(yōu)勢(shì)呢？

02 Introduction to CUDA

當(dāng)我們著手構(gòu)建深度學(xué)習(xí)模型時(shí)，很可能會(huì)傾向于采用諸如 PyTorch 或 TensorFlow 這類廣受歡迎的 Python 開發(fā)庫(kù)。盡管如此，一個(gè)不爭(zhēng)的事實(shí)是，這些庫(kù)的核心代碼都是 C/C++ 代碼。另外，正如我們先前所提及的，利用 GPU 加快數(shù)據(jù)的處理速度往往是一種主流優(yōu)化方案。此時(shí)，CUDA 的重要作用便凸顯出來(lái)！CUDA 是統(tǒng)一計(jì)算設(shè)備架構(gòu)（Compute Unified Device Architecture）的縮寫，是英偉達(dá)（NVIDIA）為使 GPU 能夠在通用計(jì)算領(lǐng)域大放光彩而精心打造的平臺(tái)。與 DirectX 被游戲引擎用于圖形運(yùn)算（graphical computation）不同，CUDA 使開發(fā)人員能夠?qū)⒂ミ_(dá)（NVIDIA）的 GPU 計(jì)算能力集成到通用軟件中，而不僅僅局限于圖形渲染。

為了實(shí)現(xiàn)這一目標(biāo)，CUDA 推出了一款基于 C/C++ 的簡(jiǎn)易接口（CUDA C/C++），幫助開發(fā)者調(diào)用 GPU 虛擬指令集（virtual intruction se）及執(zhí)行特定操作（specific operations）（如在 CPU 與 GPU 間傳輸數(shù)據(jù)）。

在繼續(xù)深入技術(shù)細(xì)節(jié)之前，我們有必要澄清幾個(gè) CUDA 編程的基礎(chǔ)概念和專業(yè)術(shù)語(yǔ)：

• host：特指 CPU 及其配套內(nèi)存；
• device：對(duì)應(yīng) GPU 及其專屬內(nèi)存；
• kernel：指代在設(shè)備（GPU）上運(yùn)行的函數(shù)代碼；

因此，在一份使用 CUDA 撰寫的基本代碼（basic code）中，程序主體在 host (CPU) 上執(zhí)行，隨后將數(shù)據(jù)傳遞給 device (GPU) ，并調(diào)用 kernels (functions) 在 device (GPU) 上并行運(yùn)行。這些 kernels 由多條線程同時(shí)執(zhí)行。運(yùn)算完成后，結(jié)果再?gòu)?device (GPU) 回傳至 host (CPU) 。

話說(shuō)回來(lái)，讓我們?cè)俅尉劢褂趦山M向量相加這個(gè)具體問(wèn)題：

借助 CUDA C/C++，編程人員能夠創(chuàng)建一種被稱為 kernels 的 C/C++ 函數(shù)；一旦這些 kernels 被調(diào)用， N 個(gè)不同的 CUDA 線程會(huì)并行執(zhí)行 N 次。

若想定義這類 kernel，可運(yùn)用 ??__global__??? 關(guān)鍵字作為聲明限定符（declaration specifier），而若欲設(shè)定執(zhí)行該 kernel 的具體 CUDA 線程數(shù)目，則需采用 ??<<<...>>>?? 來(lái)完成：

每個(gè) CUDA 線程在執(zhí)行 kernel 時(shí)，都會(huì)被賦予一個(gè)獨(dú)一無(wú)二的線程 ID，即 threadIdx，它可以通過(guò) kernel 中的預(yù)設(shè)變量獲取。上述示例代碼將兩個(gè)長(zhǎng)度（size）均為 N 的向量 A 和 B 相加，并將結(jié)果保存到向量 C 中。值得我們注意的是，相較于循環(huán)逐次處理成對(duì)加法的傳統(tǒng)串行方式，CUDA 的優(yōu)勢(shì)在于其能夠并行利用 N 個(gè)線程，一次性完成全部加法運(yùn)算。

不過(guò)，在運(yùn)行上述這段代碼前，我們還需對(duì)其進(jìn)行一次修改。切記，kernel 函數(shù)的運(yùn)行環(huán)境是 device (GPU) ，這意味著所有相關(guān)數(shù)據(jù)均須駐留于 device 的內(nèi)存之中。 要達(dá)到這一要求，可以借助 CUDA 提供的以下內(nèi)置函數(shù)：

直接將變量 A、B 和 C 傳入 kernel 的做法并不適用于本情況，我們應(yīng)當(dāng)使用指針。在 CUDA 編程環(huán)境下，host 數(shù)組（比如示例中的 A、B 和 C）無(wú)法直接用于 kernel 啟動(dòng)（<<<...>>>）。鑒于 CUDA kernels 的工作空間為 device 的內(nèi)存（device memory），故需向 kernel 提供 device 指針（device pointers）（d_A、d_B 和 d_C），以確保其能在 device 的內(nèi)存上運(yùn)行。

除此之外，我們還需通過(guò)調(diào)用 cudaMalloc 函數(shù)在 device 上劃分內(nèi)存空間，并運(yùn)用 cudaMemcpy 實(shí)現(xiàn) host 和 device 之間的數(shù)據(jù)傳輸。

至此，我們可在代碼中實(shí)現(xiàn)向量 A 和 B 的初始化，并在程序結(jié)尾處清理 CUDA 內(nèi)存（cuda memory）。

另外，調(diào)用 kernel 后，務(wù)必插入 ??cudaDeviceSynchronize();?? 這一行代碼。該函數(shù)的作用在于協(xié)調(diào) host 線程與 device 間的同步，確保 host 線程在繼續(xù)執(zhí)行前，device 已完成所有先前提交的 CUDA 操作。

此外，CUDA 的錯(cuò)誤檢測(cè)機(jī)制同樣不可或缺，這種檢測(cè)機(jī)制能協(xié)助我們及時(shí)發(fā)現(xiàn)并修正 GPU 上潛在的程序缺陷（bugs）。倘若忽略此環(huán)節(jié)，device 線程（CPU）將持續(xù)運(yùn)行，而 CUDA 相關(guān)的故障排查則將變得異常棘手，很難識(shí)別與 CUDA 相關(guān)的錯(cuò)誤。

下面是這兩種技術(shù)的具體實(shí)現(xiàn)方式：

要編譯和運(yùn)行 CUDA 代碼，首先需要確保系統(tǒng)中已裝有 CUDA 工具包（CUDA toolkit）。緊接著，使用 nvcc —— NVIDIA CUDA 編譯器完成相關(guān)代碼編譯工作。

然而，當(dāng)前的代碼尚存優(yōu)化空間。在前述示例中，我們處理的向量規(guī)模僅為 N = 1000，這一數(shù)值偏小，難以充分展示 GPU 強(qiáng)大的并行處理能力。特別是在深度學(xué)習(xí)場(chǎng)景下，我們時(shí)常要應(yīng)對(duì)含有數(shù)以百萬(wàn)計(jì)參數(shù)的巨型向量。然而，倘若嘗試將 N 的數(shù)值設(shè)為 500000，并采用 <<<1, 500000>>> 的方式運(yùn)行 kernel ，上述代碼便會(huì)拋出錯(cuò)誤。因此，為了完善代碼，使之能順利執(zhí)行此類大規(guī)模運(yùn)算，我們亟需掌握 CUDA 編程中的核心理念 —— 線程層級(jí)結(jié)構(gòu)（Thread hierarchy）。

03 Thread hierarchy（線程層級(jí)結(jié)構(gòu)）

調(diào)用 kernel 函數(shù)時(shí)，采用的是 <<<number_of_blocks, threads_per_block>>> 這種格式（notation）。因此，在上述示例中，我們是以單個(gè)線程塊的形式，啟動(dòng)了 N 個(gè) CUDA 線程。然而，每個(gè)線程塊所能容納的線程數(shù)量都有限制，這是因?yàn)樗刑幱谕痪€程塊內(nèi)的線程，都被要求共存于同一流式多處理器核心（streaming multiprocessor core），并共同使用該核心的內(nèi)存資源。

欲查詢這一限制數(shù)量的具體數(shù)值，可通過(guò)以下代碼實(shí)現(xiàn)：

就作者當(dāng)前使用的 GPU 而言，其單一線程塊最多能承載 1024 個(gè)線程。因此，為了有效處理示例中提及的巨型向量（massive vector），我們必須部署更多線程塊，以實(shí)現(xiàn)更大規(guī)模的線程并發(fā)執(zhí)行。 同時(shí)，這些線程塊被精心布局成網(wǎng)格狀結(jié)構(gòu)（grids），如下圖所展示：

??https://handwiki.org/wiki/index.php?curid=1157670?? (CC BY-SA 3.0)

現(xiàn)在，我們可以通過(guò)以下途徑獲取線程 ID：

于是，該代碼腳本更新為：

04 性能對(duì)比分析

下表展示了在處理不同大小向量的加法運(yùn)算時(shí)，CPU 與 GPU 的計(jì)算性能對(duì)比情況。

Image by the author

顯而易見(jiàn)，GPU 的處理效能優(yōu)勢(shì)，唯有在處理大規(guī)模向量時(shí)方能得以凸顯。此外，切勿忽視一件事，此處的時(shí)間對(duì)比僅僅考量了 kernel/function 的執(zhí)行耗時(shí)，而未將 host 和 device 間數(shù)據(jù)傳輸所需的時(shí)間納入考慮范圍。盡管在大多數(shù)情況下，數(shù)據(jù)傳輸?shù)臅r(shí)間開銷微不足道，但就我們目前僅執(zhí)行簡(jiǎn)易加法運(yùn)算（simple addition operation）的情形而言，這部分時(shí)間消耗卻顯得相對(duì)可觀。因此，我們應(yīng)當(dāng)銘記，GPU 的計(jì)算性能，僅在面對(duì)那些既高度依賴計(jì)算能力又適合大規(guī)模并行處理的任務(wù)時(shí)，才能得以淋漓盡致地展現(xiàn)。

05 多維線程處理（Multidimensional threads）

現(xiàn)在，我們已經(jīng)知道如何提升簡(jiǎn)單數(shù)組操作（simple array operation）的性能了。然而，在處理深度學(xué)習(xí)模型時(shí)，必須要處理矩陣和張量運(yùn)算（matrix and tensor operations）。在前文的示例中，我們僅使用了內(nèi)含 N 個(gè)線程的一維線程塊（one-dimensional blocks）。然而，執(zhí)行多維線程塊（multidimensional thread blocks）（最高支持三維）同樣也是完全可行的。因此，為了方便起見(jiàn)，當(dāng)我們需要處理矩陣運(yùn)算時(shí)，可運(yùn)行一個(gè)由 N x M 個(gè)線程組成的線程塊。還可以通過(guò) row = threadIdx.x 來(lái)確定矩陣的行索引，而 col = threadIdx.y 則可用來(lái)獲取列索引。此外，為了簡(jiǎn)化操作，還可以使用 dim3 變量類型定義 number_of_blocks 和 threads_per_block。

下文的示例代碼展示了如何實(shí)現(xiàn)兩個(gè)矩陣的相加運(yùn)算。

此外，我們還可以將此示例進(jìn)一步拓展，實(shí)現(xiàn)對(duì)多個(gè)線程塊的處理：

此外，我們也可以用同樣的思路將這個(gè)示例擴(kuò)展到三維運(yùn)算（3-dimensional operations）操作的處理。

上文已經(jīng)介紹了處理多維數(shù)據(jù)（multidimensional data）的方法，接下來(lái)，還有一個(gè)既重要又容易理解的概念值得我們學(xué)習(xí)：如何在 kernel 中調(diào)用 functions。 一般可以通過(guò)使用 ??__device__??? 聲明限定符（declaration specifier）來(lái)實(shí)現(xiàn)。這種限定符定義了可由 device （GPU）直接調(diào)用的函數(shù)（functions）。因此，這些函數(shù)僅能在 ??__global__??? 或其他 ??__device__?? 函數(shù)中被調(diào)用。下面這個(gè)示例展示了如何對(duì)一個(gè)向量進(jìn)行 sigmoid 運(yùn)算（這是深度學(xué)習(xí)模型中極其常見(jiàn)的一種運(yùn)算方式）。

至此，我們已經(jīng)掌握了 CUDA 編程的核心概念，現(xiàn)在可以著手構(gòu)建 CUDA kernels 了。對(duì)于深度學(xué)習(xí)模型而言，其實(shí)質(zhì)就是一系列涉及矩陣（matrix）與張量（tensor）的運(yùn)算操作，包括但不限于求和（sum）、乘法（multiplication）、卷積（convolution）以及歸一化（normalization ）等。舉個(gè)例子，一個(gè)基礎(chǔ)的矩陣乘法算法，可以通過(guò)以下方式實(shí)現(xiàn)并行化：

我們可以注意到，在 GPU 版本的矩陣乘法算法中，循環(huán)次數(shù)明顯減少，從而顯著提升了運(yùn)算處理速度。下面這張圖表直觀地展現(xiàn)了 N x N 矩陣乘法在 CPU 與 GPU 上的性能對(duì)比情況：

Image by the author

我們會(huì)發(fā)現(xiàn)，隨著矩陣大?。╩atrix size）的增大，GPU 在處理矩陣乘法運(yùn)算時(shí)的性能提升幅度更大。

接下來(lái)，讓我們聚焦于一個(gè)基礎(chǔ)的神經(jīng)網(wǎng)絡(luò)模型，其核心運(yùn)算通常表現(xiàn)為 y = σ(Wx + b)，如下圖所示：

Image by the author

上述運(yùn)算主要涉及矩陣乘法（matrix multiplication）、矩陣加法（matrix addition）以及對(duì)數(shù)組施加函數(shù)變換（applying a function to an array）。如若你已掌握這些并行化處理技術(shù)，意味著你現(xiàn)在完全具備了從零構(gòu)建、并在 GPU 上構(gòu)建神經(jīng)網(wǎng)絡(luò)的能力！

06 Conclusion

本文我們探討了通過(guò) GPU processing （譯者注：使用 GPU進(jìn)行數(shù)據(jù)處理和計(jì)算。）提升深度學(xué)習(xí)模型效能的入門概念。不過(guò)，有一點(diǎn)還需要指出，本文所介紹的內(nèi)容僅僅是皮毛，背后還隱藏著很多很多更深層次的東西。PyTorch 和 Tensorflow 等框架實(shí)現(xiàn)了諸多高級(jí)性能優(yōu)化技術(shù)，涵蓋了 optimized memory access、batched operations 等復(fù)雜概念（其底層利用了基于 CUDA 的 cuBLAS 和 cuDNN 等庫(kù)）。 但愿這篇文章能夠讓各位讀者對(duì)使用 .to("cuda") 方法，在 GPU 上構(gòu)建、運(yùn)行深度學(xué)習(xí)模型時(shí)的底層原理，有個(gè)初步的了解。

Thanks so much for reading! ??

Lucas de Lima Nogueira

??https://www.linkedin.com/in/lucas-de-lima-nogueira/??

END

原文鏈接：

??https://towardsdatascience.com/why-deep-learning-models-run-faster-on-gpus-a-brief-introduction-to-cuda-programming-035272906d66??