從AI系統角度回顧GPU架構變遷

2022-02-09知識

註：這篇文章的副標題和之前張偉同學的文章有些相似，不過這裏討論的視角會以AI系統視角為切入點，所以資訊量存在比較明顯的差異，希望張偉同學看到不要覺得有盜標之嫌，因為我確實覺得這個副標題比較貼切：）。

每一代NV GPU的釋出都會給業界帶來新的想象空間。作為AI系統(這裏主要代指深度學習系統)方向的從業者，最關心的自然是每一代GPU能夠為AI系統領域帶來哪些新的變量。

從之前NV GPU的甲方消費者，轉變為現在的乙方提供者，視角變化讓自己可以從不同角度來看待這個問題。這裏會以深度學習系統的發展蹤跡為套用載體，來回顧NV GPU架構的歷史變遷。

整個回顧會從最早套用於深度學習計算加速的GTX 580 開始，直到最新的Ampere架構。對每一代GPU的回顧會從以下幾個方面展開：

單器材硬件架構。

跨器材硬件架構。

AI workload及套用場景的演進。

AI軟件棧的演進。

生態發展。

首先來看Fermi，這也是第一款套用於代表性深度學習加速場景(AlexNet)的GPU架構。

1. Fermi

Fermi是支持CUDA的第三代GPU架構，第一代是2006年推出的G80架構（公開材料沒有查詢到G80的whitepaper，相對詳細一些的分析可以參見anandtech的這篇文章），第二代是2008年推出的GT200架構（類似G80，在NV官網上已經找不到類似Fermi的whitepaper，倒是在一些分析網站上有一些關聯內容，比如beyond3d的這個和anandtech的這篇）。從Fermi時代開始，Tesla產品線的每一代GPU的whitepaper都提供了公開下載的連結，裏面提到了大量的架構技術細節。這篇回顧文章正是以這些whitepaper為基礎展開。在Fermi的這篇whitepaper裏提到了這樣一段話，讀來讓人感慨頗深:

When designing each new generation GPU, it has always been the philosophy at NVIDIA to improve both existing application performance and GPU programmability ; while faster application performance brings immediate benefits , it is the GPU’s relentless advancement in programmability that has allowed it to evolve into the most versatile parallel processor of our time . It was with this mindset that we set out to develop the successor to the GT200 architecture.

以hindsight式的視角來看，將可編程性放在和效能相齊的位置，是一個重要的決策。因為可編程性的改善對於提升NV GPU的網絡效應和使用者切換成本，至關重要。在俞軍的【產品方法論】裏提到了 使用者價值=新體驗-舊體驗-切換成本 。這個公式也適用於GPGPU這種To B性質的產品。以NV當前代際GPU代指新體驗，上一代GPU代指舊體驗，效能提升相當於在強化新體驗，可編程性相當於在減少切換成本。以NV GPU代指舊體驗，競品代指新體驗，這個公式同樣成立，只不過可編程性相當於增加競品的切換成本，效能提升相當於減少競品提供的增益價值。在接下來的歷代GPU架構回顧過程中，我們可以看到NV一以貫之地堅持踐行這個理念，不斷透過效能和可編程性（包括用於提升AI開發者生產效率的努力我也歸結為廣義的可編程性）的提升來強化自己產品相較於上代產品和競品的使用者增益價值。

任何事物都有其兩面性。所以，對可編程性的重視，也存在風險，可能成為制約NV發展的阿基里斯之踵。從Google在2016年推出TPU開始（從開創AI DSA硬件先河的角度，引爆這一撥AI硬件技術演進大趨勢的寒武紀也成立於2016年），行業裏湧現出大量的AI芯片 start-up，僅從硬件層面AI絕對算力來說，這些公司裏已經出現了和NV當前主流產品效能on-par的產品，如果從performance per watt的角度來看，也已經出現了超過NV的若幹競品。其核心原因也跟NV需要關註可編程性和歷史使用者習慣的包袱有關，而新興公司沒有積累也同樣沒有包袱，所以可以在架構設計的空間裏以適當犧牲通用可編程性為代價來尋找更適合於挖掘AI計算效率的設計權衡點。相關原理在【創新者的窘境】裏也有提到，這也是考驗某個領域裏頭部企業的地方了。

回歸正題。Fermi相較前兩代架構，引入了比較大的架構變化：

對流處理器(SM)進行了重新設計,包括

每個SM包含32個CUDA cores（這裏需要提一下，CUDA core在概念上和通用計算的CPU core的性質其實是有差異的，不過並不影響從AI計算角度的討論，所以在此先不展開）。每個SM裏包括32浮點單元(FPU)和32個整數計算單元(INT Unit)，16個Load/Store單元。4個特殊函數單元(SFU)。習慣上我們會把FP32 ALU稱為一個 CUDA core，早期 NV GPU因為 INT 和 FP32 共享 datapath，不能做 ILP，所以那時候的 CUDA core可以被認為包含一個 INT和一個 FP32 ALU。

對雙精度算力進行了大幅提升。考慮到深度學習場景很少用到雙精度計算，所以對這個細節不再展開。

引入了L1 cache，並增加了Shared memory的容量。L1和Shared Memory的可配置總量是64KB，支持兩種配置（48KB shared memory + 16KB L1或48KB L1 + 16KB L1）。Shared memory從上一代的16KB增加到最多可達48KB。實際上，從Fermi開始，每一代GPU都會對Shared Memory，Register File,L1進行調整，支撐這些調整決策的是硬件推出後，透過和客戶互動叠代所增加的對套用負載的理解，以及工藝進步帶來了更多可供騰挪的片上硬件資源。

加入了對單精度浮點計算FMA的支持(Fermi之前的架構，支持雙精度的FMA，對於單精度只支持MAD，FMA的精度更有保障，參見這裏的描述)，精度有了更高保障。

Warp scheduler的數量從1增加到2，透過增加可排程發射的warp數量，來提升片上計算資源的利用率。

Fermi的SM架構示意圖如下（源於此處）

引入了768KB的L2 cache。在Fermi之前，NV GPU是沒有L2 cache的，從Fermi這一代開始，引入了L2 cache，並且隨著代際演化，不斷增加L2的尺寸。L2的引入，將之前需要由軟件開發人員take的數據搬運最佳化的部份工作讓渡給了硬件，從而部份減少了軟件開發人員的心智負擔。

對訪存系統加入了ECC支持。對於ECC，我目前仍然懷疑其對於AI計算場景的必要性。因為我的認識是AI計算過程本身，可以透過系統層面的設計（比如階段性的checkpoint）來對沖ECC所針對修復檢測的記憶體bits錯誤異常的影響，所以引入ECC這種會消耗訪存頻寬及儲存資源的手段有些浪費。這也是我理解大量不包含ECC功能的桌面顯卡能夠套用在生產訓練和推理集群的原因。

Load/store地址位寬由32-bit提升到64-bit，為統一地址存取提供了基礎條件。比如Local memory，shared memory，global memory進行統一編址（參見下面的一張示意圖）。

從AI系統的角度，NV在Fermi這一代並沒有為AI計算場景進行任何針對性的設計，包括硬件和軟件。其被套用在AlexNet上也更像是一個機緣巧合： Alex這樣的演算法科學家因為實際演算法需求，在為釘子找錘子的過程中，發現GPU相較CPU更適合解決相關問題從而將其作為錘子引入進來。整個過程中NV的作用是相對passive的。這和當前NV在AI計算領域的主動和激進存在著巨大的差異。

而在當時那個年代，為什麽是NV GPU被選中作為錘子，而不是Intel CPU或AMD GPU?

讓我們穿越過歷史的故紙堆，試圖做一些推測。

在2006年的這篇文章裏，我們能夠看到基於論文裏的實驗對比，Intel CPU上開啟BLAS庫，和NV GPU上的效能在on-par的水準，當時這個工作裏並沒有使用到CUDA，因為當年正是CUDA的元年。

有趣的是在這篇文章的註腳裏提到了除了在NVIDIA GeForce 7800 Ultra, Intel Pentium 4上的效能實驗之外，也準備加入ATI Radeon X1900上的實驗結果。在2005年更早的一篇文章裏，能夠看到在CPU, ATI GPU, NV GPU上同時進行MLP加速實驗對比的一些數據,看起來當時還有互有千秋。

而在2011年的這篇文章裏，基於CUDA實作的摺積操作，效能最多已經達到了Intel CPU上的60倍。當時使用的硬件是Intel Core i7-920(2.66GHZ)以及基於Fermi架構的GTX480/GTX580顯卡。年代久遠，實測評估已經不太現實，不過，從多個途徑的數據cross check的結果（這裏，這裏，還有這裏）來看，i7-920 2.66GHZ的峰值算力大體在在 24~100GFlops 之間。關於i7-920的理論峰值算力， @李少俠給出了一個比較專業的預估，我直接援引如下

nehalem cpu 只有 port0 的 sse 支持 FP32 乘法，port1 的 sse 只支持 FP32 加法，所以對於深度學習裏典型的乘法加法 1:1 的場景，i7-920 理論算力是 4-way*4core*2(FP add + mul ILP)*sse_freq，sse_freq 取 2.66G，那麽算力約 85Gflops，不過官方並沒有公布sse密集情況下的多核頻率，應該和這個數接近。

這個數碼和這裏的一個推算基本是相當的。所以從客觀公平性角度，不妨將i7-920的峰值算力按 85Gflops 來設定。而GTX580的峰值算力是 1.5TFlops ，大體是1個數量級的差異。再加上CUDA提供的可編程性，以及Fermi引入的提升軟件開發人員效率的一些硬件feature（比如L2的引入）,在2011年，NV GPU相較Intel CPU已經取得了在神經網絡加速場景比較明顯的優勢了。結合文章裏的這一段話，對於Fermi加入L2，並且在後續代際持續提升L2的容量的動作，就更容易有共鳴了。

The latest generation of NVIDIA GPUs, the 400 and 500 series (we use GTX 480 & GTX 580), has many advantages over older GPUs, most notably the presence of a R/W L2 global cache for device memory. This permits faster programs and simplifies writing the code. In fact, the corresponding transfer of complexity into hardware alleviates many software and optimization problems. Our experiments show that the CNN program becomes 2-3 times faster just by switching from GTX 285 to GTX 480. Manual optimization of CUDA code is very time-consuming and error prone. We optimize for the new architecture, relying on the L2 cache for many of the device memory accesses, instead of manually writing code that uses textures and shared memory. Code obtained by this pragmatic strategy is fast enough. We use the following types of optimization: pre-computed expressions, unrolled loops within template kernels, strided matrices to obtain coalesced memory accesses and registers wherever possible. Additional manual optimizations are possible in case future image classification problems will require even more computing power.

需要指出的是，上面的對比，並沒有做到完全基於第一性原理的公平性，比如Intel Core i7-920的工藝是 45nm ，而GTX580是 40nm ，整合的晶體管數量也存在明顯的差異（ 30B v.s. 0.731B ）。不過考慮到i7-920的架構設計中，只有不到20%的芯片面積用於實際計算，參見下圖（原圖來源於Fermi的whitepaper），已經可以認為論文裏的效能差異是由GPU和CPU在架構設計權衡的定性差異所帶來的，所以我們不再花費精力進行更精細的定量對比。

總的來說，在Fermi這一代，NV GPU雖然沒有為AI計算場景進行特殊的客製，但因為其相異於CPU的設計理念使得其更適配於神經網絡的平行計算特性，再加上CUDA和硬件層面改善可編程性的一系列努力，使得其「誤打誤撞」地契合了AlexNet的建模需求，在深度學習的第一個killer application上取得了不錯的開局。

2. Kepler

Kepler架構在2012推出。這一代並沒有引入多少AI計算相關的架構創新，更多是一些偏通用性質的架構改進，包括：

針對performance per watt進行了比較多的改進

去掉了SM裏的雙倍時鐘。為了對沖去除SM內雙倍時鐘對效能的影響，Kepler大幅增加了CUDA core的數量。這裏有一點用芯片空間換效能的味道。而Kepler采用的28nm的工藝升級，也為這個架構調整提供了相應的空間。

透過將部份指令依賴的判斷邏輯從硬件層面上推到軟件，節省了部份硬件資源消耗。

對SM架構進行了了比較大的調整

每個SM裏的CUDA cores數量由前代的32個提升為192個(單精度計算)。

SM裏的warp scheduler由Fermi的2個增加到4個，來適配CUDA core的數量增加。

將warp scheduler裏硬件實作的排程依賴處理邏輯上推到軟件編譯器層（得益於NV GPU的計算流水線的確定延遲特性），節省了部份硬件資源消耗。

引入了warp shuffle指令，使得warp範圍內的歸約操作不需要經過shared memory中轉，直接透過register的數據互動即可完成。而這類指令恰好為深度學習場景下特定尺寸workload的歸約操作提供了相較於shared memory更高效的實作可能。

引入了dynamic parallelism，hyper-Q, Grid Management Unit等一系列特性。基於我個人的經驗，這些特性對於提升單個AI作業本身效能幫助很有限，其中hyper-Q和Grid Management Unit對於集群層面的資源最佳化倒是會有幫助。

引入了GPUDirect技術。雖然在Kepler代際引入GPUDirect並不是為AI計算場景考慮，但這個特性對於未來分布式AI訓練的效能提升還是有著重要的價值。

對每個thread的可用寄存器數量，L1/Shared memory以及L2 cache的尺寸進行了調整。我的理解，這些調整，一方面源於NV對客戶工作負載反饋的響應，另一方面也得益於工藝提升帶來的騰挪空間。

關於Kepler時代的架構變化細節，可以參見這篇whitepaper以及GTX680的whitepaper，在此不再做資訊搬運。

Kepler時代，NV在軟件層面引入了針對AI計算場景的一個大的動作---cuDNN V1.0在2014年的釋出,並整合進了Caffe等深度學習框架中。考慮到硬件叠代的成本，其發展通常會滯後於軟件發展，所以2014年cuDNN的釋出標誌著深度學習已經進入了NV的視野，透過軟件庫的叠代加深對深度學習計算負載的理解，為後續硬件架構的演進提供資訊反饋，大體可以推測是這個思路。而2016年Pascal架構的釋出，也基本上佐證了這點。

在Kepler時代，生態方面有幾個有代表性的事件：

Google在2013年基於三台單機裝配四塊GTX680顯卡的伺服器，替換掉了之前使用的1000台CPU伺服器，完成了貓臉辨識任務。這對於NV GPU後續在深度學習計算場景的更多被采用，提供了非常好的背書。雖然從第一性原理出發，我仍然認為NV GPU和Intel CPU不應該存在100X的效能差異，但很多時候，商業上的演化並不僅僅是基於第一性原理的技術論道（實際上從一些公開工作可以看得出來,至少Intel內部的技術團隊對於NV GPU和自家硬件的優劣對比是有著清晰認知判斷的，但似乎出於別的一些原因，這些技術認知未能最終轉化成有效的公司決策操作）所能完整覆蓋的，商業宣傳布道，市場公關，再加上不同企業面臨相同場景因為自身組織特點和公司現狀對客戶反饋的反應不一，使得這種認知在相當長一段時間內深入建模人員的mindset中，這就為NV GPU擴大其在深度學習計算領域的覆蓋率打下了很好的基礎。

Google的GNMT模型在96塊K80 GPU上完成其訓練過程。在GNMT模型推出之後，工業界有若幹頭部公司先後參考其理念，將深度學習套用於機器轉譯場景，並大抵都選擇了GPU作為訓練硬件。

2013年深度學習開源框架Caffe的釋出。在Caffe釋出之前，2002年Torch其實就已經釋出，2007年Theano也已經釋出，並且都提供了深度學習建模能力的支持。但是在深度學習領域，其接納度遠不如當時的後起之秀Caffe。Caffe推出之後，很快就收到了NV的關註，包括給Caffe開發團隊免費提供GPU，以及推進cuDNN的整合，都是NV當時響應動作的部份。為什麽是Caffe而不是先推出的Torch/Theano取得這樣的成就呢？我想，這裏面的核心還是在於Caffe面對CV建模場景提供了更好的易用性（基於Caffe描述模型結構，以及客製Caffe的便利性）以及預訓練好的模型checkpoint，使得CV建模人員可以把精力更多花在add-on的建模創新本身，而不是去折騰基礎設施或是在復現其他SOTA結果上花費太多時間。值得註意的是如果我們回顧歷史，會有這樣一種感覺， Caffe開發過程中，關註的核心點首先是功能、易用性以及使用者生態的建設，效能從一開始就不是其關註的第一優先級，效能方面，Caffe透過和硬件廠商的close合作來解決，而不是自己投入大量研發精力去解決。一言以蔽，Caffe當時滿足了更多DL建模人員的需求（當時的建模需求以CV為主），達到了第一代深度學習框架和使用使用者之間更好的一個契合度。稍微形而上一些，從歷史唯物主義的角度，在一個領域代表了大多數人利益的事物往往會得到更多的支持和認可，在2013年顯然因為種種因素Caffe在CV建模領域達到了這個狀態，於是Caffe得到了快速的推廣。而NV，則借著Caffe的推廣普及，其GPU硬件更加深入深度學習建模人群了。

阿裏巴巴初具規模采購NV GPU,也是從Kepler時代開始的。騰訊使用GPU進行語音辨識加速，也始於Kepler時代。百度使用GPU的歷史則更為悠久一些，我不確定是不是在Kepler之前就開始在使用GPU了，有熟悉這段歷史的朋友歡迎提供線索。

3. Maxwell

Maxwell架構在2014年被推出。和上一代Kepler架構相同，采用的也是28nm工藝。相同工藝，通常意味著可供騰挪的硬件晶體管資源數量不會有顯著的上升，留給架構師的設計空間相對有限。不過在Maxwell時代，因為28nm工藝成熟度的改進，加上從前代產品叠代中學習到的經驗，Maxwell仍然引入了一些比較出彩的變化：

將L1 cache和shared memory進行分離，而不是合成一段可靈活配置的連續儲存資源資源。關於這一點，我們會註意到，到了Volta時代，又將L1和shared memory進行了合並。這個架構設計的反復，也體現了架構設計層面的螺旋式的上升。

對L2尺寸進行了大幅調整，從512KB擴大到2MB。這種程度的調整，我理解往往是基於真實的workload反饋，數據驅動做出的架構決策。

增加了每個SM上active的thread blocks數量，從16增加到32，改善occupancy。我的理解，這個調整也是為了配合下面的warp scheduler的改進所引入的協同動作，以確保有足夠多的warp可夠排程，來彌補去除全域warp scheduler損失掉的靈活性。

對warp scheduler的設計進行了調整，每個warp scheduler只能看到SM內部四分之一的計算執行資源，犧牲了一定的靈活度，但節省了全域warp scheduling所需的SM範圍內全域crossbar的功耗開銷。

Kepler架構warp scheduling示意圖

Maxwell架構warp scheduling示意圖（SM內的全域crossbar拆解為四個sub crossbar）

在這裏可以了解到更多Maxwell的架構細節。

Maxwell這一代，架構上也和Kepler和Fermi一樣，並沒有引入針對AI計算場景的特化考慮。

軟件層面，在14年cuDNN V1.0推出以後，進行了持續的叠代，2015年3月釋出了V2，2015年9月釋出了V3，2016年2月釋出了cuDNN V4,2016年3月伴隨P100釋出了cuDNN V5(cuDNN的發展過程中值得一提的是Scott Gray，這位兄台人不在NV，但透過逆向工程手寫了Maxwell架構上的SASS assembler，並首次基於winograd實作了快速conv演算法，最終這個作法被cuDNN團隊吸納入正式產品中),同時在V5版本也加入了對RNN/LSTM結構的最佳化支持，雖然現在RNN/LSTM結構的使用已經顯得勢微，但是在當時，對於機器轉譯，語言模型等NLP場景（以及涉及序列建模的OCR及語音辨識場景），RNN/LSTM實際上是當時的SOTA了，如果我們結合一些領域工作發表的時間做一下關聯，會發現NV對RNN/LSTM的支持跟這些工作發表的時間距離非常之近,這從一個側面體現出NV對於AI計算workload演化趨勢跟進之緊密。同時這段時間湧現了一些具備killer application內容的深度學習模型，比如微軟在2015年推出的ResNet。我沒能在公開文獻中檢索到ResNet論文發表時所使用的具體GPU型號，不過按時間推算，應該使用的是Kepler或Maxwell架構的GPU。

Maxwell時代，行業生態層面有幾件重要的事情發生：

AlphaGo在2016年3月勝出李世石，極大程度上提振了整個行業對深度學習的信心和熱情，也間接促使了一些公司采購了更多的NV GPU，因為當時的NV GPU已經幾乎是訓練深度學習模型的的標配硬件了。

Google在2015年底釋出了TensorFlow 0.5版本，雖然當時釋出的還只是一個單機版本，但對於整個深度學習領域來說，仍然有著劃時代的意義。原因在於：

使用者想加入新的模型結構（比如LSTM)，相較第一代的Caffe，生產力有了顯著的提升。曾經在Caffe裏加入過RNN/LSTM支持的同學應該對這個是比較有體感的。在TensorFlow時代，如果只關註建模功能，對效能要求不picky的話，在python層面就可以基於TF提供的原子operator完成LSTM結構的建模了。

從設計理念上，TF為不同的分布式執行策略提供了更好的基礎。將計算圖描述和執行期執行進行顯式區分，對於後端最佳化，其實引入了更幹凈的抽象邊界。

配套的TF Serving的引入。為深度學習提供了端到端的支撐能力，而不僅僅是訓練環節。

Chainer在2015年6月的釋出。因為不像TensorFlow背後有Google的豪華團隊支持，Chainer受到的關註度並不算大。但這個框架在一些geek味的深度學習建模人員那裏得到了非常好的評價。我還記得2017年參加現場GTC時，SaleForce的兩個小哥分享他們基於Chainer做的一個RNN的最佳化工作，映象中是當時現場參加人數最多的一個session。更重要的是，Chainer實際上為後續遵循define-by-run執行模式的深度學習框架提供了一個生動的範例。這也是一個體現技術創新的脈絡延續性的例子。

回過頭來審視，站在2016年Pascal推出的前夜，會有一種「山雨欲來風滿樓」的感覺，於是就有了Pascal架構的釋出。

4. Pascal

Pascal架構在2016年3月被推出,采用16nm和14nm的工藝，說其是NV面向AI計算場景釋出的第一版架構，當不為過。在Pascal架構裏引入了面向AI場景很重要的一些特性：

FP16。從實用性角度來看，Pascal時代的FP16其實蠻雞肋的，理論算力相較FP32提升只有2X，但是中間累加結果只支持FP16的格式，使得對於訓練場景很容易影響精度。所以在我了解到的範圍內，Pascal的FP16套用於訓練場景的案例非常有限。但這個特性的引入，為Volta加入FP16 TensorCore提供了有效的實驗反饋，並且這是NV在硬件架構層面第一次結合深度學習作業特點進行的軟硬協同設計，其架構設計層面的影響是蠻深遠的。技術脈絡的演進，總是草蛇灰線、伏脈千裏，透過某個技術出現的單點時刻，回溯其源起和演進路徑，可以更有效地指導後續的技術決策，而不是寄希望於靈光一現式的運氣。

NVLink。在我的認知裏，NVLink是一個典型的點狀頭部業務驅動技術架構升級，技術架構升級進一步影響更多業務接受的技術特性。NVLink背後的含義是單機多卡，並且只有卡多到一定程度以後，NVLink的價值才更容易顯現。所以在Pascal的whitepaper裏給出的參考架構至少也是單機四卡的規模，標配是單機八卡。而單機四卡/八卡被生產環境接納，並不像現在感覺這麽顯然。原因主要是幾點：

和CPU/記憶體資源配比的問題。

機房運維新增復雜性的問題。

對集群排程系統提出了更高要求。

阿裏大約是在2018年上線了具備NVLink的生產集群，而要到2019年才有比較多生產作業啟用NVLink進行單機多卡加速,距離NVLink架構特性的推出大約有近三年的時間差,可以參見這裏的一些分析描述。

NVLink的技術細節簡要來說可以分為NVHS，Sub-link，Link，Gang四個層次。一條NVHS的link提供單向20Gb/s的傳輸頻寬頻寬，8根NVHS構成一個Sub-link，兩條Sub-links組成一條用於雙向連線的Link，P100架構下的單機八卡配置，不同GPU之間會由四個Link組成一條Gang，所以Gang的雙向匯總頻寬是 20Gb * 8 * 2 * 4 = 160GB/s ，單向匯總頻寬80GB/s，是PCIe提供頻寬的 5x 。更形象一些的示意圖如下：

HBM。Pascal架構一個breakthrough的特性是首次將HBM技術引入到NV的硬件中(行業裏最早將HBM引入到產品中的是AMD在2015年將其套用在Fji GPU產品中，第一塊HBM記憶體芯片則是在2013年由SK hynix推出)，提供相較上一代Maxwell架構最高3倍的訪存頻寬提升。HBM的引入為高訪存壓力的訓練作業提供了更有力的硬件支持。

INT8。從Pascal時代開始，NV GPU首次引入了對INT8格式的支持(INT8首次引入是在基於GP102架構的P40 GPU，而不是基於GP100架構的P100 GPU)。INT8對於推理場景的必要性現在已經不需要過多說明，但是在16年的硬件裏加入INT8支持，還是一件蠻激進的事情。Google TPU的訊息也是在16年才對外正式expose,即便NV能夠有渠道更早嗅到相關技術趨勢，能夠做到這麽快將INT8加入量產GPU裏，也是非常迅速的一個執行動作了。這也反映出NV的敏捷性和執行力。

在SM架構層面，Pascal引入的變化不算多。在我的理解中，更多是工藝提升帶來更多可用晶體管資源以後，可以把更多料堆起來反映到SM的數量提升，屬於增量式的變化。從這一點也其實體現出NVIDIA從Fermi時代起結合CUDA所選擇的計算架構的優越性----幾乎每一代新架構（特別是SM相關）都能夠以相對增量的方式將工藝提升帶來的新增晶體管資源利用起來，而不是動輒引入大的架構調整。

在Pascal whitepaper裏提及的將CPU和GPU進行統一記憶體存取的特性，我自己的經驗，並沒有感覺到這個特性對於AI計算領域提供了多少實際收益。記得在剛剛拿到P100以後，我們評測過其page migration engine的表現，當時的結論非常negative。我們當時的判斷是想使用CPU記憶體來作為GPU視訊記憶體的backup，並且保證效能下滑不要太明顯，還是需要在AI框架層結合套用作業的特點來進行處理，而不是直接交給page migration engine來在後台自動完成。

在Pascal時代，另一個值得一提的是，能夠看到NV在從芯片向整機系統邁進，在其whitepaper裏提到了單機8卡的DGX-1 server，這也是從Fermi時代開始，第一次在whitepaper裏出現單GPU之上整機的方案。考慮到DGX-1的成本，在大量生產環境布署的其實是參考DGX-1代工生產的類似GPU伺服器機型。

軟件層面，Pascal這一代針對AI場景也引入了更多變化，首先是面向推理加速場景的TensorRT的釋出,然後是NCCL在2016年的釋出。說到這裏，我一直很好奇NCCL以開源形式存在至今（多機部份初始是以閉源形式提供，後來在框架自行提供多機通訊庫的壓力下推動了NCCL多機通訊版本的開源），而TensorRT則一直保持閉源形態，是什麽導致這兩個產品存在這樣的差異？在NV的官方blog上提到NCCL最早是一個research project，而TensorRT的源起則並不是一個research project，所以可能在對外開放度上NCCL尺度會更大。不過從生態建設的角度，hindsight地來看，如果TensorRT在起步的時候，就考慮按一個開源專案的方式來運作，不確定會不會帶來更快的叠代速度（如果叠代速度是一項重要的商業度量metrics）？這也是自己有時候會YY的問題之一-------如果有一款全新硬件，需要為其規劃AI軟件棧的技術路線的話，NV的哪些作法是應該借鑒的，哪些是應該改良的？

行業生態方面，有幾個很有意義的事件

2016年5月，MxNet 0.7版本的釋出。其實很長一段時間，MxNet在業內都有著不錯的認可度。包括我知道國內的一些AI公司在早期都是更傾向於基於MxNet進行擴充套件客製，而不是TensorFlow。以及NV在新款GPU上進行深度學習模型效能最佳化時，往往會先在MxNet上開展，比如MLPerf。但一個發人深思的事實是，時至今日，基於MxNet發表的論文已經廖廖（參考Paperswithcode的一個趨勢統計）。相近的專案啟動時間，為什麽最後出現這樣的結果？我想，至少有一個資訊是我們可以解讀出來的，對於深度學習框架這個場景來說，效能可能並不是決定勝負的因素，因為至少在NV GPU上，MxNet的綜合效能是最好的，但最後勝出的不是MxNet。PyTorch核心團隊的一篇分享內容，也許能部份回答這個問題。

2016年9月，PyTorch 0.1.1版本的釋出，時至今日，PyTorch在研究領域，已經一騎絕塵，把其他深度學習框架遠遠甩在了後面，包括TensorFlow。與此同時，隨著計算硬件的發展，PyTorch也開始暴露出一些局限性，最大的局限性也和其最大的優勢有關---過於強調易用性使得使用者建模程式碼裏很容易充斥大量Pythonic的程式碼，導致引入大量host和加速器的數據以及控制依賴，在加速器不斷push效能邊界的大背景下，這個影響可能會吃掉易用性的紅利。

5. Volta

在距離Pascal架構推出僅過去一年之後，2017年5月的GTC keynote，NV宣布了下一代Volta架構的釋出。考慮到Pascal代際引入了較大的架構升級，間隔這麽短又釋出了下一代的Volta架構，這並不是一個常規行為，實際上是源於Google TPU當時給NV帶來的壓力。在Volta之前，面向AI計算場景，NV GPU相較TPU其實是存在技術上的代際差異的，這就嚴重威脅到了NV在AI計算領域的地位。如果技術上不能及時拉平代際差異，僅靠CUDA生態建立的使用者切換成本來進行對抗，很可能會出現【創新者的窘境】裏的狀況，被顛覆性的技術拉下馬來。這個技術代際差異在Volta架構透過引入第一代TensorCore在訓練場景進行了拉平，隨後Turing架構的第二代TensorCore在推理場景上進行了拉平，直到Ampere時代，NV才算再次鞏固了自己在AI計算領域的龍頭地位。時隔四年，回顧這一段行業歷史，還是感覺精彩之至。比如，為什麽選擇了先支持FP16，而不是Google提及的BF16？為什麽考慮選擇了4x4尺寸的TensorCore，而不是更大的尺寸？如何figure out出來和FP16配套的Loss scaling的訓練策略？TensorCore的編程API如何對外暴露來盡量避免和現有CUDA體系形成撕裂？怎樣讓NV之外的開發者也具備在TensorCore上開發程式的能力？這裏涉及到了大量的技術、非技術因素的綜合權衡。一些結論決策，當時看是合理的，但現在來看，已經被推翻或叠代更新了。比如從最早支持FP16，到支持BF16，及至引入TF32。從4x4尺寸，到更大的8x8。CUDA層面暴露的API粒度也在發生演進變化。如此等等。這裏重要的不是結論，而是探討這些結論產生的過程。因為這些過程，是可以遷移到下一撥workload，下一撥架構創新機會上的，具體的結論則很可能未必。

Volta時代架構層面幾個重要變化是

Tensor Core的引入。已經有很多資料對Tensor Core進行過介紹，比如這篇介紹自動混合精度落地實踐的文章，對Tensor Core的基本概念進行了介紹，Andrew Kerr的這份slides對使用Tensor Core開發高效計算kernel進行了比較詳細的探討。這裏我想盡量減少一些不必要的重復。我想從兩個角度來展開對TensorCore的討論，一個是硬件層面提供的基礎支持，另一個是CUDA軟件介面層面對外暴露的API

硬件層面，Volta時代的Tensor Core提供了單cycle完成4x4x4=64條半精度FMA計算操作的的能力，在計算能力提升的同時，因為原先由64條FP32指令完成的操作現在由一條TensorCore指令完成，也省掉了大量用於保存中間計算結果的寄存器資源消耗，提升了數據復用性。

Volta的每個SM內部除了跟Pascal時代一樣的64個FP32/32個FP64/64個INT CUDA core以外，還提供了8個Tensor Cores。FP32 CUDA core和Tensor Cores數量差距如此之大，再加上Tensor Core對數據存取的需求和FP32 CUDA core存在明顯的差異，如何將Tensor Core計算能力對CUDA軟件開發人員暴露成為一個蠻考究的問題。從Volta架構開始，PTX指令中引入了Warp level的矩陣計算及存取指令，並不斷擴充套件其靈活性，從只能使用wmma.load/wmma.store/wmma.mma按ISA約束的數據組織方式來進行存取，到使用mma指令由軟件開發人員根據需要進行顯式的數據組織排布（靈活性提升的同時也意味著編程復雜性的增加）。以及為了簡化NV之外CUDA開發人員使用TensorCore的負擔所推出的CUTLASS軟件庫。TensorCore編程復雜性的根源在於

從原先相對符合樸素直覺的SIMT過渡到了Warp-level的編程範式，同一個warp內的32個thread需要按照TensorCore的數量分成8個分組(grouping),每個group內的四個thread需要協同完成shared memory到寄存器的數據載入，group與group之間也需要考慮協同。不同處理任務之間的依賴關系，增加了思考的復雜度度。

考慮到效能，shared memory裏的數據組織方式需要註意規避bank conflict,這就引入入了另一個反樸素直覺的維度。

Tensor Core的粗計算粒度，使得程式設計師需要顯式地將計算過程進行細致的拆解。有興趣的同學可以參考這裏和這裏的兩份材料，可以看到CUDA層面為16x16x16規模矩陣計算暴露的wmma::mma_sync API對映為64（FP32累加）或32條HMMA指令(FP16累加)的細節拆解過程。如果想直接使用PTX的mma指令進行編程來獲得更好的靈活性，就需要自行完成相應的拆解。

總的來說，Volta時代的TensorCore在技術原理上幫助NV拉平了和Google在AI計算領域的技術代差，但是現在回頭來看仍然存在比較多的局限性，包括：

只支持FP16。雖然NV的技術團隊設計了精巧的loss scaling策略，並透過自動混合精度外掛程式化的手段完成了和主流深度學習框架的整合，但是對模型訓練精度仍然存在一定影響，導致其在實際業務中被啟用的比例並不算高。

可編程性相較SIMT時代的CUDA差了很多。CUDA層面暴露的WMMA API的粒度很粗，靈活性較差。即便使用PTX層面的WMMA指令，也只能提供相似靈活度。使用HMMA SASS指令雖然足夠靈活，但又會影響到程式的可移植性，並且過於hacky。不過從 PTX 6.4 版本開始（對應於 CUDA 10.1 ），在PTX層面對外暴露了MMA指令，該指令的粒度和HMMA SASS指令已經一一對應，這就使得NV之外的工程師也擁有了更精細的Tensor Core編程API了。不過總體來說，Tensor Core編程本身的復雜性，導致NV之外能夠基於Tensor Core開發CUDA程式的工程師數量顯著下降。

NVLink頻寬的增加。Volta時代的NVLink 2.0，單條Link提供的單向頻寬由Pascal時代的NVLink 1.0的20GB/s提升到25GB/s，雙向向頻寬由40GB/s提升提升到50GB/s，GPU之間的Link數量也由4條擴充套件為6條，加總在一起，GPU之間的雙向互聯頻寬達到了300GB/s，相較NVLink 1.0帶來了80%的頻寬提升。

NVSwitch。在NVLink的基礎上，2018年NV基於NVSwitch推出了DGX-2機型。NVSwitch的出現使得單機內部卡間通訊頻寬呈現出了均勻性。單機八卡機型，每塊GPU透過6根links連線到NVSwitch上，再透過NVSwitch和其他GPU進行互聯。所以理論上，任意兩塊GPU之間，都可以達到最高達300GB/s的雙向頻寬。這就為reduce-to-one乃至all2all這種通訊模式提供了極大便利性。

HBM頻寬的增加。經過Pascal時代引入HBM的試水，在Volta時代，NV對HBM的頻寬進行了擴充套件，16GB視訊記憶體配置的卡型，視訊記憶體頻寬由732GB/s提升到900GB/s，並且對HBM的實際訪存效率進行了改善。

為SM的整數計算加入了單獨的data path，使得整數計算指令（在AI場景通常對應於地址計算邏輯）不會阻塞浮點計算或Tensor Core指令的發射。引入這個特性是對Tensor Core毫無疑問是有benefit的，因為隨著計算能力的提升，增加了其他操作成為瓶頸的可能，所以期望把地址計算操作所需的整數計算放在單獨的data path裏，從而緩解瓶頸。同時INT data path的引入也可以使其他非Tensor Core計算，但存在浮點和整數指令混合的場景受益，Volta 上不使用 Tensor Core 的 SGEMM最佳化難度相比 Pascal 更低就有這方面原因（感謝 @李少俠同學指出這一點）。

在Volta時代，深度學習模型層面也出現了一些新的變化，這些變化主要集中在自然語言處理領域。首先是2017年Transformer模型結構在機器轉譯場景取得的效果突破，使得其開始替換之前SOTA的RNN/LSTM結構，被大量采用。然後是2018年BERT的出現，在展現其出色的預訓練模型效果的同時，也給使用NV GPU預訓練BERT帶來了比較大的挑戰和壓力。NV對此的應對是迅速推出了32G視訊記憶體的V100卡型，並在設計針對訓練場景的下一代GPU架構時對視訊記憶體容量給予了更高權重。另一個標誌性的事件是2020年OpenAI對外公布其在微軟提供的V100集群上完成了包含175B參數的GPT-3模型的訓練過程。一個有意思但無從考證的坊間傳言是，BERT的核心開發者之前在微軟工作，在加入Google後，借助Google提供的軟硬結合的AI算力，在比較短的時間內推出了BERT，微軟的高層決策者收到這個反饋之後，投入了一大筆錢購買了DGX的伺服器：）。

軟件層面，NV在2017年底釋出了CUTLASS，為NV之外的開發者開發TensorCore程式提供了一個比較好的參考基礎。2018年Q3以開原始碼的形式釋出了跨結點的NCCL通訊庫(這裏的一個背景是，NCCL 1.X 一直是開源的，2.0 加入了多機分布式支持之後一度閉源了一年左右，後續框架開始自行支持 AllReduce on RDMA的壓力倒推NV最終將NCCL 2.0開源，最終成為多卡通訊的事實標準)，並在當年釋出了TensorRT Inference Server(之後被更名為Triton Inference Server)。2018年NV為TensorFlow加入了TF-TRT的支持（還記得在NV推出這個工作之前，當時我還在阿裏，我的同事易凡同學因為支持業務的需要，正好也完成了一個類似工作的原型，看到NV釋出了撞車的tf-trt以後，感覺也是比較微妙），顯著提升了基於TensorFlow使用TRT布署能力的易用性。單器材最佳化之外，繼續向AI系統全鏈路滲透。

另外值得一提的是在V100時代，大模型訓練的需求變得更加旺盛。在V100之前的時代，大模型訓練場景主要包括三類：一類是大規模人臉分類，因為其會有一個巨大的全連線分類層需要進行模型並列；一類是ResNet101這種極深類別的模型，需要引入類似pipeline並列的作法；還有一種就是大規模稀疏搜推廣模型，透過將大規模embedding table分片存放解決。從V100之後，我們會發現大模型訓練所需要的技術核心點其實並沒有變化，只不過場景更多以NLP為主。其推手當屬BERT和GPT-3這兩大killer application性質的模型。2019年，NV對外釋出了支持Megatron-LM的工作，這也是Megatron-LM的第一次對外亮相，當時的工作基於32GB V100 GPU來完成，並且只支持Tensor Parallelism，不支持Pipeline Parallelism。與此同時，也已經能夠看到業界有更多圍繞大模型訓練相關的工作，比較有代表性的當屬微軟DeepSpeed團隊的工作，這支團隊透過引入精細的視訊記憶體最佳化技術從另一條技術路徑對大模型訓練進行探索，這個工作也是以V100 GPU為主要硬件平台來完成的。其他相關工作包括微軟的PipeDream（在NV GPU上完成），Google的GPipe（同時使用了NV GPU和TPU作為硬件平台）等。這篇回顧文章的重點不是討論分布式訓練技術，所以不再展開更多相關細節，這裏的關鍵是，這些大模型訓練工作背後的支撐硬件，幾乎清一色以NV GPU為主，以及若幹Google發起的工作基於TPU來完成。

另一個值得一提的是，MLPerf training 0.5在2018年12月份結果的釋出，從此MLPerf training先後歷經0.5、0.6、0.7、1.0、1.1五個版本的叠代(MLPerf inference歷經0.5、0.7、1.0、1.1四個版本叠代)。叠代過程中，因為有了清晰的最佳化標的，NV得以資源聚集，整套軟硬技術全棧也經歷了巨大的變化飛躍。雖然MLPerf裏的不少最佳化結果並不能直接遷移到生產環境裏（參考這裏的一些討論），但其對於行業技術進步的促進作用還是毋庸置疑的。

在V100時代，隨著Tensor Core對Conv/GEMM這類計算密集算子帶來顯著加速，訪存密集算子以及kernel launch開銷對端到端效能的影響變得也越來越大，也是從V100時代開始，自動算子融合技術開始受到更多關註。在這方面，Google XLA應該是最早進行相關探索的團隊。Google在為TPU設計XLA編譯器的過程中，發現裏面的一些技術對於GPU和CPU也同樣能夠帶來收益，於是將XLA也使能到了GPU上(XLA CPU的投入一直乏善可陳，在此不提)。時至今日，無論是NV GPU還是新興的AI芯片公司，算子融合技術已經是一項must了。

6. Turing

Turing架構在2018年9月的SIGGRAPH正式釋出。和Volta相同，Turing也基於TSMC 12nm工藝完成生產。從AI計算的角度，Turing主要面向推理場景，相較Volta其架構上的變化主要有：

加入了INT8/INT4/INT1的Tensor Core支持。

使用GDDR6替換掉HBM，面向推理場景，提供更好的性價比，推理場景因為不涉及反向傳播計算，並且推理過程可以更激進的啟用算子融合技術，所以訪存壓力顯著小於訓練場景。

大幅去除用於FP64的計算資源（從Volta時代的 1:2 的FP64:FP32比例下降到Turing的 1:32 ），因為DL推理場景並不需要使用FP64，所以大幅提升了能效比。

7. Ampere

Ampere架構在2020年5月釋出。這一代架構引入了比較多的變化：

采用7nm工藝，GPU die的尺寸相較V100略增（ 826mm^2 v.s. 815mm^2 ）,片上晶體管數量從 21B 激增到 54B ，這就為更激進的架構創新提供了更多騰挪空間。

引入了BF16和TF32的Tensor Core，後者對於AI訓練場景，具備更大的潛力成為開箱即用的精度格式。TF32的引入也得益於NV從Pascal時代對低精度訓練的持續投入叠代，透過多輪實際系統的交付以及客戶互動，獲得了寶貴的反饋輸入，最終催生了TF32的數據格式。看到TF32格式的引入，還是會對NV這家公司充滿了敬意，因為這是一家在不斷push自身邊界的公司，期望未來能夠保持這種態勢。

對Tensor Core的尺寸進行了增加，由Volta/Turing代際的4x4x4增大到8x8x4，每個cycle可以完成 128~256 個FMA操作，進一步改善數據復用，提升了計算密度(GA100是 256 ，GA10x是 128 )。

第三代NVLink（NVSwitch)。Ampere架構的NVSwitch和上一代的區別是將頻寬提升了一倍，理論上，任意兩塊GPU之間，都可以達到最高達600GB/s的雙向頻寬。

引入了MIG。用於提供硬件層面的單GPU多工安全隔離。不過坦率說，我聽到這個特性的使用案例似乎並不多。

結構化稀疏。這個特性，從剛引入的很受關註，到現在漸漸有些期望回呼。究其原因，涉及到侵入使用者訓練過程的最佳化手段，如果不是非常顯著的收益，接受起來往往會有一定門坎。

HBM視訊記憶體增加到40GB/80GB的配置，適配大模型訓練的視訊記憶體壓力。頻寬相較V100代際也有1.7X的提升，達到1.5TB/s。

大幅提升片上L2容量，從V100的6MB激增到40MB，這往往意味著基於工作負載反饋拿到了重要的架構決策所需的數據輸入才可能支撐這麽激進的的架構決策。L2存取頻寬相較V100提升2.3X。Ampere架構裏，L2到SM的crossbar不再是全連線，而是分拆成兩個L2 sub-partition,每個sub-partition只服務於直接和其相連的SM。L2尺寸的增加和L2 sub-partition的引入其實是整個架構決策的一體兩面（減少L2和SM互聯的crossbar的開銷），互為補充。

引入了用於異步數據copy的LDGSTS指令，可以bypass寄存器的中轉，直接從global memory裏將數據載入到shared memory，減少了寄存器的壓力和不必要的數據中轉，進一步節省了功耗。並且因為這條指令的異步性，可以作為背景操作和前台的計算指令overlap執行，進一步提升整體計算效率。

CUDA Graph在2018年被引入，用於最佳化小kernel的launch開銷。在Ampere代際，為其加入了專門的硬件支持。時至今日，CUDA Graph已經是AI效能最佳化的重要依韋特性了，比如MLPerf裏的諸多刷榜最佳化，以及Facebook也在2021年將CUDA Graph正式整合進PyTorch中。

加入了專用於JPEG格式解碼的硬件decoder NVJPG，用於對數據預處理環節加速。以及大幅增加了用於影片解碼的NVDEC的硬件資源，同理，也是為了對數據預處理環節加速。

更多細節可以參考這裏的whitepaper。

在Ampere時代，NV在整機之外，進一步推出集群解決方案SuperPOD，以及基於SuperPOD搭建的超算集群Selene，SuperPOD和Selene支持了Megatron-LM以及MLPerf training的大量效能最佳化工作，也作為解決方案，成功交付給了若幹客戶。這也是一個蠻有意思的行業訊號。

軟件方面，Ampere時代一個比較重要的工作是TensorRT和PyTorch的整合Torch-TensorRT，不過這個工作仍然面臨一個大的挑戰，那就是PyTorch的模型寫法過於靈活，存在不少無法成功匯出TorchScript的模型寫法（比如NLP場景中decoding部份的迴圈生成結構），對於這部份如何進行高效自動推理加速，目前仍然是一個open的問題。

另一個NV GPU上軟件相關有代表性的工作是對dynamic shape的支持，關於有效處理dynamic shape，大約從三年前業界就有過過呼聲。TensorRT目前仍然是透過padding的策略來解決dynamic shape的問題。阿裏在兩個月前開源的BladeDISC是一個基於MLIR針對dynamic shape提供的E2E的AI編譯解決方案，不過完備性還有待完善。Amazon的Nimble工作則基於TVM技術棧探索了另一條解決dynamic shape的技術方案。

無論是Torch-TensorRT，還是對dynamic shape的支持，都反映出對AI開箱即用效能最佳化的重視，這在一定程度上，也和AI當前更多進入到行業套用落地期的階段有關。

模型方面，Ampere自2020年推出以後，直到現在，能夠看到AI領域主要的關註焦點集中在大模型訓練上，除了GPT-3，BERT類模型以外，MoE模型也受到了一定關註，相應地也催生了一系列工作，包括Google TPU之上的Gshard系列工作，GPU上的DeepSpeed-MoE工作等。但是大模型到底能夠為業務層面帶來多少實際收益，其收益是否足以justify新增的算力投入，是否需要更高效環保的模型設計方法以及AI算力提供方案，目前仍然是一個open的問題。

以上結合AI系統演進的視角，回顧了從Fermi到Ampere共7代架構，期望隨著未來Ampere-Next以及Ampere-Next-Next的釋出，我們可以再添加入相關的內容，一起經歷見證AI系統領域和NV GPU架構的共同演進發展。

這篇回顧涉及到了比較長的時間跨度，比較寬的技術區域，整個回顧內容，有些是我親身經歷的，有些是我基於獲取到的一手或二手資訊提煉的，還有一些則是根據網絡上的資訊進行交叉校驗後匯總出來的，難免會有疏漏或不夠準確之處，也歡迎同行朋友的批評指正。

文章撰寫過程中感謝一些朋友指出其中錯誤以及提供建議，包括：

字節跳動的易百忍同學糾正了原文中關於NVSwitch和NCCL 2.0開源描述不準確的地方。

商湯科技的高洋同學提議加入Scott Gray對cuDNN演進的貢獻。

@李少俠同學糾正了如下細節:

HBM最早是2015年在AMD Fji顯卡中使用，而不是2016年釋出的P100

INT8在GP102架構開始引入，而不是GP100

CUDA 10.1(PTX 6.4)裏為Tensor Core編程靈活性提供了和SASS HMMA指令對應的MMA PTX指令。

HBM能效比優於而不是差於GDDR，這實際上也是HBM技術本身被引入的一個重要的原因。

Ampere一代裏不同架構(GA100~GA10x)的HMMA計算吞吐的差異。

提供了Intel Nehalem架構CPU峰值flops的理論預估值。