去年，NVIDIA (NV) 發(fā)布了Amper新架構(gòu)的GPU，NVIDIA GPU架構(gòu)的發(fā)展類似Intel的CPU，針對不同場景和技術(shù)革新，經(jīng)歷了不同架構(gòu)的演進(jìn)。
內(nèi)容轉(zhuǎn)自：智能計算芯世界
《NVIDIA A100 Tensor Core GPU技術(shù)白皮書》《NVIDIA Kepler GK110-GK210架構(gòu)白皮書》《NVIDIA Kepler GK110-GK210架構(gòu)白皮書》《NVIDIA Kepler GK110架構(gòu)白皮書》《NVIDIA Tesla P100技術(shù)白皮書》《NVIDIA Tesla V100 GPU架構(gòu)白皮書》《英偉達(dá)Turing GPU 架構(gòu)白皮書》- Kepler架構(gòu)里，F(xiàn)P64單元和FP32單元的比例是1:3或者1:24；K80。
- Maxwell架構(gòu)里，這個比例下降到了只有1:32；型號M10/M40。
- Pascal架構(gòu)里，這個比例又提高到了1:2(P100)但低端型號里仍然保持為1:32，型號Tesla P40、GTX 1080TI/Titan XP、Quadro GP100/P6000/P5000
- Votal架構(gòu)里，F(xiàn)P64單元和FP32單元的比例是1:2；型號有Tesla V100、GeForceTiTan V、Quadro GV100專業(yè)卡。
- Turing架構(gòu)里，一個SM中擁有64個半精度，64個單精度，8個Tensor core，1個RT core。
- Ampere架構(gòu)的設(shè)計突破，在8代GPU架構(gòu)中提供了該公司迄今為止最大的性能飛躍，統(tǒng)一了AI培訓(xùn)和推理，并將性能提高了20倍。A100是通用的工作負(fù)載加速器，還用于數(shù)據(jù)分析，科學(xué)計算和云圖形。
從Kelper、Maxwell、Pascal、Volta，GPU架構(gòu)的更新體現(xiàn)在SM、TPC的增加，最終體現(xiàn)在GPU浮點計算能力的提升。NVIDIA GPU架構(gòu)相關(guān)文章：
相比Pascal架構(gòu)，Volta GPU架構(gòu)的顯著特征是它的Tensor Core，新的Tensor Core是專門為深度學(xué)習(xí)設(shè)計的，有助于提高訓(xùn)練神經(jīng)網(wǎng)絡(luò)所需的性能。Tensor Core推動卷積和矩陣運算。集成了Tensor Core也是NVIDIA新一代GPU架構(gòu)Turing(圖靈)的一個重要特性。關(guān)于GPU架構(gòu)，NV 很有意思的是會用一些歷史上杰出的科學(xué)家的名字來命名自己的硬件架構(gòu)。總體上，NV GPU 用到的 SIMT 基本編程模型都是一致的，每一代相對前代基本都會在 SM 數(shù)量、SM 內(nèi)部各個處理單元的流水線結(jié)構(gòu)等等方面有一些升級和改動。這篇暫時不涉及到渲染管線相關(guān)的部分，其他諸如多少 nm 工藝、內(nèi)存頻率提升等等也都先略過，只關(guān)注計算相關(guān)的硬件架構(gòu)演進(jìn)。關(guān)于初代 GPU 的架構(gòu)，Tesla可以找到的資料不太多，基本上都是從 Fermi 開始的。1、Fermi
Compute Capability：2.0，2.1
Fermi架構(gòu)SM- 2 個 Warp Scheduler/Dispatch Unit
- 32 個 CUDA Core（分在兩條 lane 上，每條分別是 16 個）
- 每個 CUDA Core 里面是 1 個單精浮點單元（FPU）和 1 個整數(shù)單元（ALU），可以直接做 FMA 的乘累加
我的理解是做一個雙精 FMA 需要用到兩個 CUDA Core？所以是 32 / 2 = 16
2、Kepler
Compute Capability：3.0,，3.2，3.5， 3.7；這一代 SM 整體結(jié)構(gòu)上跟之前是一致的，只不過升級完了以后又往里面塞進(jìn)去了更多的運算單元，其他部分也沒有做太大的改動。- 4 個 Warp Scheduler，8 個 Dispatch Unit
- CUDA Core 增加到 192 個（4 * 3 * 16，每條 lane 上還是 16 個）
- 單獨分出來 64 個（4 * 16，每條 lane 上 16 個）雙精運算單元。
- SFU 和 LD/ST Unit 分別也都增加到 32 個
Kepler 是附近幾代在硬件上直接有雙精運算單元的架構(gòu)，不用通過單精單元去做雙精運算了，所以對比前后幾代的雙精浮點的性能話會發(fā)現(xiàn) Kepler 要高出一截。3、Maxwell
Compute Capability：5.0， 5.2， 5.3可能是覺得 Kepler 往一個 SM 里面塞了太多東西，其實最終效率也并沒有那么高，這一代的 SM 開始做減法了，每個 SM（SMM）中包含：- 4 個 Warp Scheduler，8 個 Dispatch Unit
- 32 個 SFU 和 LD/ST Unit（4 * 8）
Kepler 里面 192 這個數(shù)字也被詬病了（不是 2 的倍數(shù)）。這些硬件單元的流水線分布也不再是像 Kepler 那樣大鍋燉了，而是有點像是把 4 個差不多像是 Fermi 的 SM 拼在一起組成一個 SM：
- 1 個 Warp Scheduler 和 2 個 Dispatch Unit
圖上沒有看到之前 lane 的標(biāo)記，不過我猜應(yīng)該也還是 4 條，兩條 CUDA Core 的 lane，1條 SFU，1條 LD/ST Unit。應(yīng)該是工藝和頻率的提升，Maxwell 每個 CUDA Core 的性能相比 Kepler 提升了 1.4 倍，每瓦性能提升了 2 倍。對 CUDA Core 的詳細(xì)結(jié)構(gòu)沒有再介紹，姑且認(rèn)為從 Fermi 開始一直到以后 CUDA Core 內(nèi)部的結(jié)構(gòu)都沒有什么改變。另外一點是，前面說到的雙精單元在這一代上也移除了。也許是覺得認(rèn)為只有少數(shù) HPC 科學(xué)計算才用的上的雙精單元在這代上不太有必要吧。
4、Pascal
Compute Capability：6.0, 6.1, 6.2；這一代可以說是有了質(zhì)的飛躍，還是先從 SM 開始：
可以看到一個 SM 內(nèi)的部分作了進(jìn)一步的精簡，整體思路是 SM 內(nèi)部包含的東西越來越少，但是總體的片上 SM 數(shù)量每一代都在不斷增加，每個 SM 中包含：- 2 個 Warp Scheduler，4 個 Dispatch Unit
- 16 個 SFU 和 LD/ST Unit（2 * 8）
一個 SM 里面包含的 Process Block 數(shù)量減少到了 2 個，每個 Process Block 內(nèi)部的結(jié)構(gòu)倒是 Maxwell 差不多：- 1 個 Warp Scheduler 和 2 個 Dispatch Unit
單個 Process Block 的流水線增加到 6 條 lane 了？
- 面向 Deep Learning 做了一些專門的定制（CuDNN 等等）
- 除了 PCIE 以外，P100 還有 NVLink 版，單機(jī)卡間通信帶寬逆天了，多機(jī)之間也能通過 Infiniband 進(jìn)一步擴(kuò)展 NVLink（GPUDirect）
然后 NV 現(xiàn)在已經(jīng)把 Infiniband 行業(yè)的龍頭 Mellanox 給收購了…… 說不定那時候就已經(jīng)有這個想法了呢
- P100 上把 GDDR5 換成了 HBM2，Global Memory 的帶寬漲了一個數(shù)量級
- 16nm FinFET 工藝，性能提升一大截，功耗還能控制住不怎么增加
- Unified Memory，支持把 GPU 的顯存和 CPU 的內(nèi)存統(tǒng)一到一個相同的地址空間，驅(qū)動層自己會做好 DtoH 和 HtoD 的內(nèi)存拷貝，編程模型上更加友好了
CUDA Core 在這一代也終于有了升級，現(xiàn)在硬件上直接支持 FP16 的半精計算了，半精性能是單精的 2 倍，猜測應(yīng)該是一個單精單元用來算兩個半精的計算。5、Volta
Compute Capability：7.0， 7.2；又一個針對深度學(xué)習(xí)的質(zhì)變 Feature，Tensor Core！看到 SM 的時候我們會發(fā)現(xiàn)這一代除了多出了一個額外的 Tensor Core 的單元以外，怎么 SM 的體積看起來好像又加回去了，每個 SM 中包含：- 4 個 Warp Scheduler，4 個 Dispatch Unit（發(fā)現(xiàn)不需要配 2 個 Dispatch 給每個 Scheduler 了？白皮書里面倒是沒有對這個的解釋）
- 4 個 SFU（發(fā)現(xiàn)對特殊計算的需求減少了？）
事實上相比 Pascal 而言，單個 SM 中的單精運算單元數(shù)量是一致的，相當(dāng)于把 Pascal 中的每個 Process Block 進(jìn)一步地又拆成了 2 個，每個 Process Block 中包含：- 1 個 Warp Scheduler，1 個 Dispatch Unit
這里把原本的 CUDA Core 給拆開了，F(xiàn)P32 和 INT32 的兩組運算單元現(xiàn)在是獨立出現(xiàn)在流水線 lane 里面了，這一設(shè)計的好處是在前幾代架構(gòu)中 CUDA Core 同時只能處理一種類型的運算，而現(xiàn)在每個 cycle 都可以同時有 FP32 和 INT32 的指令在一起跑了。Pascal 中需要 6 個 cycles 來做一組 FMA，現(xiàn)在在 Volta 中只需要 4 個 cycles。另外每個 Warp Scheduler 還有了自己的 L0 指令 cache。這一代還改進(jìn)了一下MPS，現(xiàn)在從硬件上直接支持對資源的隔離，方便多任務(wù)共享 GPU。Tensor Core：最重大的改動不用說也知道是 Tensor Core 了。
Tensor Core 的思路從系統(tǒng)設(shè)計上還是相當(dāng)直接的，目前深度學(xué)習(xí)的 workload 中最主要的計算量都在矩陣的乘加上，因此為了專門去高效地支持這些 workload，就增加一些專用于矩陣運算的專用部件進(jìn)去。這個也是常見的 AI ASIC（比如 Google 的 TPU、其他廠商的各種 xPU 等等）通常采用的思路，只不過 ASIC 可以從一開始就是針對特定的 workload 去的，因此設(shè)計上可以更直接更激進(jìn)一些，直接上大量的 MMU（Matrix Multiply Unit），然后采用例如脈沖陣列這種設(shè)計去最大化它的 throughput。而 NV 的 GPU 畢竟還要用作其他一些通用的運算，所以只能往原本的 SM 流水線里面插進(jìn)去一些額外的專用部件 lane 了。開個腦洞，要是哪一天發(fā)現(xiàn)除了 FMA 以外還有其他另外一種形式的運算有大量的需求，未來的 GPU 設(shè)計里面說不定也會出現(xiàn)其他 Core。好在 FMA 除了深度學(xué)習(xí)以外在 HPC 的 workload 里面也是挺常見的，這個設(shè)計以后還是比較有用的。Tensor Core 4x4 Matrix Multiply and AccumulateMixed Precision Multiply and Accumulate in Tensor CoreTensor Core 這個部件直接從 SM 的寄存器里面取兩個 FP16 的矩陣作為輸入，進(jìn)行全精度的矩陣乘之后得到的結(jié)果可以是 FP16 或者 FP32 的，然后累加到 FP16/FP32 的 accumulator 里面去。數(shù)據(jù)類型選擇 FP16 作為輸入然后輸出 FP32 猜測可能是為了保證結(jié)果不溢出，然后在加速部件設(shè)計等等方面做了一些 trade off。
所以 FP16 in -> FP16 out 和 FP16 in -> FP32 out 哪一個性能更好呢…
我沒有測過，但是猜測可能默認(rèn)結(jié)果是 FP32 out 更快？反而是輸出 FP16 需要從 FP32 再轉(zhuǎn)一次？
接下來道理我們都懂了，那 Tensor Core 要怎么用呢？這個部件的編程模型在一開始接觸的時候可能會有一些坑。我們知道常規(guī)的 CUDA 代碼需要制定 grid 的結(jié)構(gòu)、block 的結(jié)構(gòu)，然后其實我們寫的 kernel 代碼都是針對每一個單獨的 thread 的，可以認(rèn)為是 thread level 的編程。對一個子矩陣的 FMA 運算存在比較多的數(shù)據(jù)重用機(jī)會，這時候如果只是一個 thread 算一個矩陣塊的 FMA 就比較浪費了，因此 Tensor Core 的設(shè)計是用一整個 warp 去共同完成一個 FMA 運算，一個 warp 中的 32 個 thread 可以復(fù)用寄存器里面的數(shù)據(jù)。CUDA 對 Tensor Core 的指南里面把這個叫做 “WMMA warp-wide macro-instructions”。所以 Tensor Core 的編程模型直接就是針對一整個 warp 寫的。事實上，Tensor Core 的代碼寫起來還是有相當(dāng)多的限制的，CUDA 給 Tensor Core 提供了C的API：PTX 的指令應(yīng)該更多一些，不過我沒有詳細(xì)看過。首先用來做乘加的矩陣都需要放在這個叫 wmma::fragment 的變量里面，這個本質(zhì)上就是定義了一個要放在 SM 寄存器上的存儲空間，但是需要提供詳細(xì)的 FMA 參數(shù)：- 第一個參數(shù) Use 是這個 fragment 在 FMA 運算里面的角色，可選項有：matrix_a、matrix_b和 accumulator，含義就是字面意思，也沒什么需要再解釋的了。
- m，n，k，T 是這一個 warp 里面要處理的的 FMA 子矩陣的形狀以及數(shù)據(jù)類型，不同的 Capability 能夠支持的組合還不太一樣，比如最基礎(chǔ)的就是 a、b 都是 __half，accumulator 是 float，然后 m、n、k 都是 16。
m、n、k 的組合不是任意的，能支持的種類跟 Capability 直接相關(guān)，比如 V100 和后來出的 T4 能夠支持的就不一樣，具體可以在 Programming Guide 里面查。 - 最后這個 Layout 可選項有兩個 row_major 和 col_major，代表這個 fragment 在內(nèi)存里面實際存儲的行列主序情況。
load_matrix_sync 和 store_matrix_sync 分別是把數(shù)據(jù)寫到 fragment 空間里面和從這里面取出來寫到別的地方去。fill_fragment 對 fragment 初始化。mma_sync 就是對整個 warp 調(diào)用 Tensor Core 去跑完這一個 FMA 運算了。常規(guī)的寫法也是先把矩陣 A、B 都 load 到 shared_memory 上，然后再從 shared_memory 里面取對應(yīng) FMA 塊大小的數(shù)據(jù)到 fragment 里面，mma_sync 跑完，最后從 fragment 里面把結(jié)果寫到外面去。這里的注意點是上面這些代碼（包括 fragment 定義以及下面幾個函數(shù)的調(diào)用）都是針對 warp 的，即我們在寫代碼的一開始就需要考慮到每個 block 里面的 thread 結(jié)構(gòu)，保證一個 warp 的 32 個 thread 執(zhí)行的代碼是完全相同的。相應(yīng)地，對矩陣的分塊也是需要在寫代碼的時候就考慮清楚，我們要保證每個 warp 處理的 a、b 矩陣的大小剛好是這個地方設(shè)定好的 m、n、k。看起來確實相當(dāng)麻煩，不過想想可能好像也還好，本來如果要寫出性能很好的 CUDA 代碼來，每個 warp 要算多少東西也是需要精細(xì)考慮清楚的。
Volta 這一代對 SIMT 的編程模型也做了改變。在之前的 SIMT 流水線中，如果一個 warp 的指令里面出現(xiàn)了分支，這些分支塊是不能被同時執(zhí)行的。所以一直以來寫 CUDA 代碼都會要有一個原則是不要在一個 warp 里面出現(xiàn)不同的分支，要不需要花費兩倍的時間去處理。這一代開始把 PC 和調(diào)用棧做成了每個線程獨立的：Volta Warp with Per-Thread Program Counter and Call Stack現(xiàn)在呢，每個分支里面的指令可以在更細(xì)粒度的層面上進(jìn)行混合調(diào)度了，也可以手動插入一些在 warp 層面同步的指令進(jìn)去：Programs use Explicit Synchronization to Reconverge Threads in a Warp白皮書后面給了一個可以從這個改動上得到收益的 Starvation-Free Algorithms 的示例，修改帶鎖的雙向鏈表的時候，不同 thread 可能會被 block 在鎖上，以前的架構(gòu)應(yīng)該基本上不太可能能處理得了這種 case，新架構(gòu)就保證了即使有些 thread 還在等待鎖，另外的 thread 也有可能先拉出來跑。可能也是因為這樣所以 1 個 Dispatch Unit 配 1 個 Warp Scheduler 了？因為線程指令的實現(xiàn)事實上更加復(fù)雜了。
所以其實最后還是同時只能執(zhí)行一個分支里面的一部分，這個 upgrade 我暫時還沒有想到具體的應(yīng)用場景會有多常出現(xiàn)（上面這個帶鎖雙向鏈表我覺得寫在 CUDA 里面就很不常見啊…），以及會具體有多少性能收益，說不定還是原本的那種簡單的設(shè)計更直接更高效一些呢。（期待一下未來的硬件里面會不會把這個恢復(fù)回去……）
以前 CUDA 編程原則里面不要寫分支的那條在新架構(gòu)下我覺得還是適用的，不寫分支就不會有這么多額外的麻煩要考慮了。
另外有一個 Cooperative Group 的新設(shè)計倒是看起來感覺更有用一些。原本的 __syncthreads( ) 是針對一個 block 里面的所有 thread 做同步的，現(xiàn)在可以對不同 block 的不同 thread 單獨定義同步組了，CUDA launch 的時候會把同一個組的一起 launch 上去，同步可以在一個更加細(xì)粒度的層面上完成。6、Turing
Compute Capability：7.5；TU102 GPU包含6個圖像處理集群（GPC）、36個紋理處理集群（TPC）和72個流式多元處理器（SM）。
Turing TU102 GPU內(nèi)部構(gòu)造每個GPC均包含一個專用的光柵化引擎和6個TPC，且每個TPC均包含兩個SM。每個SM包含：- 4個紋理單元以及96KB的L1或共享內(nèi)存
且我們可根據(jù)計算或圖形工作負(fù)載將這些內(nèi)存設(shè)置為不同容量。每個SM中的全新RT核心處理引擎負(fù)責(zé)執(zhí)行光線追蹤加速。Turing架構(gòu)采用全新SM設(shè)計，每個TPC均包含兩個SM，每個SM共有64個FP32核心和64個INT32核心。Turing SM支持并行執(zhí)行FP32與INT32運算，每個Turing SM還擁有8個混合精度Turing Tensor核心和1個RT核心。
7、Ampere
Compute Capability：8.0；NVIDIA A100在AI訓(xùn)練（半/單精度操作，F(xiàn)P16/32）和推理（8位整數(shù)操作，INT8）方面，GPU比Volta GPU強(qiáng)大20倍。在高性能計算（雙精度運算，F(xiàn)P64）方面，NVIDIA表示GPU的速度將提高2.5倍。
GA100 (SM) 內(nèi)部架構(gòu)
GA100 Full GPU with 128 SMs (A100 Tensor Core GPU has 108 SMs)A100采用是GA100 GPU的縮減版GA100架構(gòu)和規(guī)格，GA100滿配版和A100配置如下。NVIDIA Ampere GA100 GPU架構(gòu)滿配如下：- 8 TPCs/GPC, 2 SMs/TPC, 16 SMs/GPC, 128 SMs per full GPU
- 64 FP32 CUDA Cores/SM, 8192 FP32 CUDA Cores per full GPU
- 4第三代Tensor Cores/SM, 512第三代Tensor Cores per full GPU
- 6 HBM2 stacks, 12 512bit 內(nèi)存控制器
NVIDIA Ampere GA100 GPU架構(gòu)配置如下:- 2 SMs/TPC, up to 16 SMs/GPC, 108 SMs
- 64 FP32 CUDA Cores/SM, 6912 FP32 CUDA Cores
- 4第三代Tensor Cores/SM, 432第三代Tensor Cores
- 5 HBM2 stacks，10 512bit 內(nèi)存控制器
Ampere GA100是迄今為止設(shè)計的最大的7nm GPU。GPU完全針對HPC市場而設(shè)計，具有科學(xué)研究，人工智能，深度神經(jīng)網(wǎng)絡(luò)和AI推理等應(yīng)用程序。NVIDIA A100 是一項技術(shù)設(shè)計突破，在五項關(guān)鍵技術(shù)領(lǐng)域得到創(chuàng)新和突破：- NVIDIA Ampere架構(gòu) — A100的核心是NVIDIA Ampere GPU架構(gòu)，其中包含超過540億個晶體管，使其成為世界上最大的7納米處理器。
- 基于TF32的第三代張量核(Tensor Core)： Tensor核心的應(yīng)用使得GPU更加靈活，更快，更易于使用。TF32包括針對AI的擴(kuò)展，無需進(jìn)行任何代碼更改即可使FP32精度的AI性能提高20倍。此外， TensorCore 現(xiàn)在支持FP64，相比上一代，HPC應(yīng)用程序可提供多達(dá)2.5倍的計算量。
- 多實例（Multi-Instance）GPU — MIG是一項新技術(shù)功能，可將單個A100GPU劃分為多達(dá)七個獨立的GPU，因此它可以為不同大小的作業(yè)提供不同程度的計算，從而提供最佳利用率。
- 第三代NVIDIA NVLink —使GPU之間的高速連接速度加倍，可在服務(wù)器中提供有效的性能擴(kuò)展。
- 結(jié)構(gòu)稀疏性—這項新的效率技術(shù)利用了AI數(shù)學(xué)固有的稀疏特性來使性能提高一倍。
NVIDIA A100基于7nm Ampere GA100 GPU，具有6912 CUDA內(nèi)核和432 Tensor Core，540億個晶體管數(shù)，108個流式多處理器。采用第三代NVLINK，GPU和服務(wù)器雙向帶寬為4.8 TB/s，GPU間的互連速度為600 GB/s。另外，Tesla A100在5120條內(nèi)存總線上的HBM2內(nèi)存可達(dá)40GB。
NVIDIA GPU架構(gòu)白皮書系列，下載鏈接：NVIDIA GPU架構(gòu)白皮書
《NVIDIA A100 Tensor Core GPU技術(shù)白皮書》
《NVIDIA Kepler GK110-GK210架構(gòu)白皮書》
《NVIDIA Kepler GK110-GK210架構(gòu)白皮書》
《NVIDIA Kepler GK110架構(gòu)白皮書》
《NVIDIA Tesla P100技術(shù)白皮書》
《NVIDIA Tesla V100 GPU架構(gòu)白皮書》
《英偉達(dá)Turing GPU 架構(gòu)白皮書》
參考鏈接：https://jcf94.com/2020/05/24/2020-05-24-nvidia-arch/https://wccftech.com/nvidia-ampere-ga100-gpu-7nm-architecture-specifications-performance-deep-dive/
轉(zhuǎn)載申明：轉(zhuǎn)載本號文章請注明作者和來源，本號發(fā)布文章若存在版權(quán)等問題，請留言聯(lián)系處理，謝謝。
推薦閱讀
更多架構(gòu)相關(guān)技術(shù)知識總結(jié)請參考“架構(gòu)師全店鋪技術(shù)資料打包”相關(guān)電子書(35本技術(shù)資料打包匯總詳情可通過“閱讀原文”獲取)。
全店內(nèi)容持續(xù)更新，現(xiàn)下單“全店鋪技術(shù)資料打包(全)”，后續(xù)可享全店內(nèi)容更新“免費”贈閱，價格僅收188元(原總價290元)。
溫馨提示：
掃描二維碼關(guān)注公眾號，點擊閱讀原文鏈接獲取“架構(gòu)師技術(shù)全店資料打包匯總(全)”電子書資料詳情。
