GPU架構變遷之AI系統視角:從費米到安培

語言: CN / TW / HK

 

撰文 | 楊軍

每一代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的文章: https://www.anandtech.com/show/2116/5 ),第二代是2008年推出的GT200架構(類似G80,在NV官網上已經找不到類似Fermi的whitepaper,倒是在一些分析網站上有一些關聯內容,比如beyond3d的文章 https://www.beyond3d.com/content/reviews/51/8 和anandtech的文章 https://www.anandtech.com/show/2549 )。

從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.(來源: https://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf

以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的精度更有保障,可參見這裡的描述: https://en.wikipedia.org/wiki/Multiply%E2%80%93accumulate_operation ),精度有了更高保障。

     

  • warp scheduler的數量從1增加到2,通過增加可排程發射的warp數量,來提升片上計算資源的利用率。Fermi的SM架構示意圖如下(https://www.nvidia.com/content/PDF/fermi_white_papers/P.Glaskowsky_NVIDIA%27s_Fermi-The_First_Complete_GPU_Architecture.pdf

 

 

 

  • 引入了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年的這篇文章( https://hal.inria.fr/inria-00112631/document )裡,我們能夠看到基於論文裡的實驗對比,Intel CPU上開啟BLAS庫,和NV GPU上的效能在on-par的水準,當時這個工作裡並沒有使用到CUDA,因為當年正是CUDA的元年。

有趣的是在,這篇文章的腳註裡提到了除了在NVIDIA GeForce 7800 Ultra( https://www.techpowerup.com/gpu-specs/geforce-7800-gtx.c127 ),Intel Pentium 4( https://en.wikipedia.org/wiki/Pentium_4 )上的效能實驗之外,也準備加入ATI Radeon X1900( https://www.techpowerup.com/gpu-specs/radeon-x1900-xt.c454 )上的實驗結果。在2005年更早的一篇文章( https://ieeexplore.ieee.org/stamp/stamp.jsp?tp=&arnumber=1575717 )裡,能夠看到在CPU, ATI GPU, NV GPU上同時進行MLP加速實驗對比的一些資料,看起來當時還互有千秋。

而在2011年的這篇文章( https://people.idsia.ch/~juergen/ijcai2011.pdf )裡,基於CUDA實現的卷積操作,效能最多已經達到了Intel CPU上的60倍。當時使用的硬體是Intel Core i7-920(2.66GHZ)以及基於Fermi架構的GTX480/GTX580顯示卡。年代久遠,實測評估已經不太現實,不過,從多個途徑的資料cross check的結果( https://gadgetversus.com/processor/intel-core-i7-920-specs/, https://fiehnlab.ucdavis.edu/staff/kind/collector/benchmark/core-i7-vs-opteron, 還有https://www.overclock.net/threads/intel-and-amd-gflops-data-thread-looking-for-sandybridge-too.869018/ )來看,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密集情況下的多核頻率,應該和這個數接近。

這個數字和( https://progforperf.github.io/nehalem.pdf )的一個推算基本是相當的。所以從客觀公平性角度,不妨將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%的芯片面積用於實際計算,(參見下圖:https://www.nvidia.com/content/PDF/fermi_white_papers/P.Glaskowsky_NVIDIA%27s_Fermi-The_First_Complete_GPU_Architecture.pdf),已經可以認為,論文裡的效能差異是由GPU和CPU在架構設計權衡的定性差異所帶來的,所以我們不再花費精力進行更精細的定量對比。

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

2

Kepler

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

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

 

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

     

  • 通過將部分指令依賴的判斷邏輯從硬體層面上推到軟體,節省了部分硬體資源消耗。

 

2. 對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更高效的實現可能。

 

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

 

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

 

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

 

關於Kepler時代的架構變化細節,可以參見這篇whitepaper( https://www.nvidia.com/content/dam/en-zz/Solutions/Data-Center/tesla-product-literature/NVIDIA-Kepler-GK110-GK210-Architecture-Whitepaper.pdf )以及GTX680的whitepaper( https://www.nvidia.com/content/PDF/product-specifications/GeForce_GTX_680_Whitepaper_FINAL.pdf ),在此不再做資訊搬運。

 

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(https://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/achievedoccupancy.htm)。我的理解是,這個調整也是為了配合下面的warp scheduler的改進所引入的協同動作,以確保有足夠多的warp可夠排程,來彌補去除全域性warp scheduler損失掉的靈活性。

     

  • 對warp scheduler的設計進行了調整(https://www.anandtech.com/show/7764/the-nvidia-geforce-gtx-750-ti-and-gtx-750-review-maxwell/3),每個warp scheduler只能看到SM內部四分之一的計算執行資源,犧牲了一定的靈活度,但節省了全域性warp scheduling所需的SM範圍內全域性crossbar的功耗開銷。

     

  • Kepler架構warp scheduling示意圖。

 

  • Maxwell架構warp scheduling示意圖(SM內的全域性crossbar拆解為四個sub crossbar)

 

在這裡( https://www.microway.com/download/whitepaper/NVIDIA_Maxwell_GM204_Architecture_Whitepaper.pdf )可以瞭解到更多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( https://github.com/NervanaSystems/maxas ),並首次基於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時代,行業生態層面有幾件重要的事情發生:

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

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

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

     

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

     

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

3. 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場景很重要的一些特性:

1. FP16。從實用性角度來看,Pascal時代的FP16其實蠻雞肋的,理論算力相較FP32提升只有2X,但是中間累加結果只支援FP16的格式,使得對於訓練場景很容易影響精度。所以在我瞭解到的範圍內,Pascal的FP16應用於訓練場景的案例非常有限。

但這個特性的引入,為Volta加入FP16 TensorCore提供了有效的實驗反饋,並且這是NV在硬體架構層面第一次結合深度學習作業特點進行的軟硬協同設計,其架構設計層面的影響是蠻深遠的。 技術脈絡的演進,總是草蛇灰線、伏脈千里,透過某個技術出現的單點時刻,回溯其源起和演進路徑,可以更有效地指導後續的技術決策,而不是寄希望於靈光一現式的運氣。

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

  • 和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架構的一個突破性特性是首次將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的表現( https://developer.nvidia.com/blog/beyond-gpu-memory-limits-unified-memory-pascal/ ),當時的結論非常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在起步的時候,就考慮按一個開源專案的方式來運作,不確定會不會帶來更快的迭代速度(如果迭代速度是一項重要的商業度量)?這也是自己有時候會思考的問題之一——如果有一款全新硬體,需要為其規劃AI軟體棧的技術路線的話,NV的哪些做法是應該借鑑的,哪些是應該改良的?

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

  • 2016年5月,MxNet 0.7版本的釋出。其實很長一段時間,MxNet在業內都有著不錯的認可度。包括我知道國內的一些AI公司在早期都是更傾向於基於MxNet進行擴充套件定製,而不是TensorFlow。以及NV在新款GPU上進行深度學習模型效能優化時,往往會先在MxNet上開展,比如MLPerf。

     

    但一個發人深思的事實是,時至今日,基於MxNet發表的論文已經廖廖(參考Paperswithcode的一個趨勢統計)。相近的專案啟動時間,為什麼最後出現這樣的結果?我想,至少有一個資訊是我們可以解讀出來的,對於深度學習框架這個場景來說,效能可能並不是決定勝負的因素,因為至少在NV GPU上,MxNet的綜合性能是最好的,但最後勝出的不是MxNet。PyTorch核心團隊的一篇分享內容(https://soumith.ch/posts/2021/02/growing-opensource/),也許能部分回答這個問題。

     

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

 

5

Volta

在距離Pascal架構推出僅過去一年之後,2017年5月的GTC keynote( https://www.anandtech.com/show/11360/the-nvidia-gpu-tech-conference-2017-keynote-live-blog ),NV宣佈了下一代Volta架構的釋出。

考慮到Pascal代際引入了較大的架構升級,間隔這麼短又釋出了下一代的Volta架構,這並不是一個常規行為,實際上是源於Google TPU當時給NV帶來的壓力。

在Volta之前,面向AI計算場景,NV GPU相較TPU其實是存在技術上的代際差異的,這就嚴重威脅到了NV在AI計算領域的地位。如果技術上不能及時拉平代際差異,僅靠CUDA生態建立的使用者切換成本來進行對抗,很可能會出現《創新者的窘境》裡的狀況——被顛覆性的技術拉下馬來。這個技術代際差異在Volta架構通過引入第一代Tensor Core在訓練場景進行了拉平,隨後Turing架構的第二代Tensor Core在推理場景上進行了拉平,直到Ampere時代,NV才算再次鞏固了自己在AI計算領域的龍頭地位。

時隔四年,回顧這一段行業歷史,還是感覺精彩之至。比如,為什麼選擇了先支援FP16,而不是Google提及的BF16?為什麼考慮選擇了4x4尺寸的Tensor Core,而不是更大的尺寸?如何figure out出來和FP16配套的Loss scaling的訓練策略?Tensor Core的程式設計API如何對外暴露來儘量避免和現有CUDA體系形成撕裂?怎樣讓NV之外的開發者也具備在Tensor Core上開發程式的能力?這裡涉及到了大量技術、非技術因素的綜合權衡。

一些結論決策,當時看是合理的,但現在來看,已經被推翻或迭代更新了。比如從最早支援FP16,到支援BF16,及至引入TF32。從4x4尺寸,到更大的8x8。CUDA層面暴露的API粒度也在發生演進變化。如此等等。這裡重要的不是結論,而是探討這些結論產生的過程。因為這些過程,是可以遷移到下一撥workload,下一撥架構創新機會上,具體的結論則很可能未必。

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

  • Tensor Core的引入。已經有很多資料對Tensor Core進行過介紹,比如這篇介紹自動混合精度落地實踐的文章(https://zhuanlan.zhihu.com/p/56114254),對Tensor Core的基本概念進行了介紹,Andrew Kerr的這份slides(https://developer.download.nvidia.cn/video/gputechconf/gtc/2020/presentations/s21745-developing-cuda-kernels-to-push-tensor-cores-to-the-absolute-limit-on-nvidia-a100.pdf)對使用Tensor Core開發高效計算kernel進行了比較詳細的探討。這裡我想盡量減少一些不必要的重複。我想從兩個角度來展開對Tensor Core的討論,一個是硬體層面提供的基礎支援,另一個是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 Core數量差距如此之大,再加上Tensor Core對資料存取的需求和FP32 CUDA core存在明顯的差異,如何將Tensor Core計算能力對CUDA軟體開發人員暴露成為一個蠻考究的問題。

從Volta架構開始,PTX指令中引入了Warp level的矩陣計算及存取指令,並不斷擴充套件其靈活性,從只能使用wmma.load/wmma.store/wmma.mma按ISA約束的資料組織方式來進行存取,到使用mma指令由軟體開發人員根據需要進行顯式的資料組織排布(靈活性提升的同時也意味著程式設計複雜性的增加),以及為了簡化NV之外CUDA開發人員使用Tensor Core的負擔所推出的CUTLASS軟體庫。Tensor Core程式設計複雜性的根源在於:

  • 從原先相對符合樸素直覺的SIMT過渡到了Warp-level的程式設計正規化,同一個warp內的32個thread需要按照Tensor Core的數量分成8個分組(grouping),每個group內的四個thread需要協同完成shared memory到暫存器的資料載入,group與group之間也需要考慮協同。不同處理任務之間的依賴關係,增加了思考的複雜度度。

     

  • 考慮到效能,shared memory裡的資料組織方式需要注意規避bank conflict,這就引入了另一個反樸素直覺的維度。

     

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

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

  • 只支援FP16。雖然NV的技術團隊設計了精巧的loss scaling策略(https://arxiv.org/abs/1710.03740),並通過自動混合精度外掛化的手段完成了和主流深度學習框架的整合,但是對模型訓練精度仍然存在一定影響,導致其在實際業務中被啟用的比例並不算高。

     

  • 可程式設計性相較SIMT時代的CUDA差了很多。CUDA層面暴露的WMMA API的粒度很粗,靈活性較差。即便使用PTX層面的WMMA指令,也只能提供相似靈活度。使用HMMA SASS指令雖然足夠靈活,但又會影響到程式的可移植性,並且過於hacky。不過從PTX 6.4版本開始(對應於CUDA 10.1),在PTX層面對外暴露了MMA指令(https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-for-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之外的開發者開發Tensor Core程式提供了一個比較好的參考基礎。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裡的不少優化結果並不能直接遷移到生產環境裡(參考這裡的一些討論: https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf ),但其對於行業技術進步的促進作用還是毋庸置疑的。

在V100時代,隨著Tensor Core對Conv/GEMM這類計算密集運算元帶來顯著加速,訪存密集運算元以及kernel launch開銷對端到端效能的影響變得也越來越大,也是從V100時代開始,自動運算元融合技術開始受到更多關注。

在這方面,Google XLA應該是最早進行相關探索的團隊。Google在為TPU設計XLA編譯器的過程中,發現裡面的一些技術對於GPU和CPU也同樣能夠帶來收益,於是將XLA也使能到了GPU上(XLA CPU的投入一直乏善可陳,在此不提)。時至今日,無論是NV GPU還是新興的AI晶片公司,運算元融合技術已經是必備的了。

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( https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf )。

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

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

另一個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的問題。

8

小結

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

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

作者 注:這篇文章的副標題和之前 張偉同學的《GPU架構演進十年:從費米到安培》 有些相似,不過這裡討論的視角以AI系統視角為切入點,所以資訊量存在比較明顯的差異,希望張偉同學看到不要覺得有盜標之嫌,因為我確實覺得這個副標題比較貼切:

作者簡介:楊軍,目前在NVIDIA Compute Arch團隊從事AI軟硬協同設計工作。六分系統,四分演算法背景,致力於通過將系統與演算法相結合構建技術軍火庫,來推動產品業務的進步。

(本文經授權後釋出,原文:https://zhuanlan.zhihu.com/p/463629676)

其他人都在看

歡迎下載體驗OneFlow新一代開源深度學習框架:https://github.com/Oneflow-Inc/oneflow/

 


本文分享自微信公眾號 - OneFlow(OneFlowTechnology)。
如有侵權,請聯絡 [email protected] 刪除。
本文參與“OSC源創計劃”,歡迎正在閱讀的你也加入,一起分享。