1
| 本文作者： AI研習(xí)社 | 2017-05-17 11:22 |
編者按：5 月 11 日，在加州圣何塞舉辦的的 2017 年度 GPU 技術(shù)大會(huì)上，英偉達(dá)發(fā)布了 Tesla V100，號(hào)稱史上最強(qiáng)的 GPU 加速器。發(fā)布之后，英偉達(dá)第一時(shí)間在官方開發(fā)者博客放出一篇博文，詳細(xì)剖析了包括 Tesla V100，GV100 GPU，Tensor Core，以及 Volta 架構(gòu)等在內(nèi)的各項(xiàng)新特性/新產(chǎn)品的技術(shù)內(nèi)涵，雷鋒網(wǎng)編譯如下。
眾所周知，目前無論是語音識(shí)別，還是虛擬個(gè)人助理的訓(xùn)練；路線探測(cè)，還是自動(dòng)駕駛系統(tǒng)的研發(fā)，在這些人工智能領(lǐng)域，數(shù)據(jù)科學(xué)家們正在面對(duì)越來越復(fù)雜的 AI 挑戰(zhàn)。而為了更好地實(shí)現(xiàn)這些頗具未來感的強(qiáng)大功能，就必須在實(shí)踐中引入一些指數(shù)級(jí)的更加復(fù)雜的深度學(xué)習(xí)模型。
另一方面，HPC（高性能計(jì)算）在現(xiàn)代科學(xué)研究中一直起著至關(guān)重要的作用。無論是預(yù)測(cè)天氣，新藥物的研究，或是探索未來能源，科研人員每天都需要利用大型計(jì)算系統(tǒng)對(duì)現(xiàn)實(shí)世界做各種各樣的仿真和預(yù)測(cè)。而通過引入 AI 技術(shù)，HPC 就可以顯著提升科研人員進(jìn)行大數(shù)據(jù)分析的效率，并得到一些此前通過傳統(tǒng)的仿真和預(yù)測(cè)方法無法得到新結(jié)論。
為了進(jìn)一步推動(dòng) HPC 和 AI 領(lǐng)域的相關(guān)發(fā)展，英偉達(dá)近期發(fā)布了新一代 Tesla V100 GPU 加速器。它基于最新的 NVIDIA Volta GV100 GPU 平臺(tái)和各種突破性技術(shù)創(chuàng)新，可以為各種超級(jí)計(jì)算系統(tǒng)提供一個(gè)強(qiáng)大的運(yùn)算平臺(tái)，不論在以科學(xué)仿真為主要手段的計(jì)算科學(xué)領(lǐng)域，還是在以洞悉數(shù)據(jù)奧秘為目標(biāo)的數(shù)據(jù)科學(xué)領(lǐng)域，Tesla V100 都能為相關(guān)應(yīng)用提供強(qiáng)大的算力支持。
下面，我們會(huì)通過這篇博客對(duì) Tesla V100 的核心：Volta 架構(gòu)做一個(gè)深度剖析，同時(shí)幫助開發(fā)者了解它在實(shí)際開發(fā)中具體帶來了哪些優(yōu)勢(shì)。
NVIDIA Tesla V100 是目前世界上最高性能的并行處理器，專門用于處理需要強(qiáng)大計(jì)算能力支持的密集型 HPC、AI、和圖形處理任務(wù)。
Tesla V100 加速器的核心是 GV100 GPU 處理器?；谂_(tái)積電專門為 NVIDIA 設(shè)計(jì)的最新 12nm FFN 高精度制程封裝技術(shù)，GV100 在 815 平方毫米的芯片尺寸中，內(nèi)部集成了高達(dá) 211 億個(gè)晶體管結(jié)構(gòu)。相較于上一代產(chǎn)品，也就是 Pascal 系列 GPU，GV100 不但在計(jì)算性能上有了長(zhǎng)足的進(jìn)步，同時(shí)還增加了許多令人眼前一亮的新特性。包括進(jìn)一步精簡(jiǎn)的 GPU 編程和應(yīng)用部署流程，以及針對(duì) GPU 資源利用情況的深度優(yōu)化。其結(jié)果是，GV100 在提供強(qiáng)大計(jì)算性能的同時(shí)還非常省電，下圖顯示了 Tesla V100 加速器和上代產(chǎn)品 Tesla P100 加速器在 ResNet-50 模型訓(xùn)練和推理中的性能對(duì)比，可以看到最新的 V100 要遠(yuǎn)超上一代 P100。
Tesla V100 的關(guān)鍵特性總結(jié)如下：
● 針對(duì)深度學(xué)習(xí)優(yōu)化的流式多處理器（SM）架構(gòu)。作為 GPU 處理器的核心組件，在 Volta 架構(gòu)中 NVIDIA 重新設(shè)計(jì)了 SM，相比之前的 Pascal 架構(gòu)而言，這一代 SM 提高了約 50% 的能效，在同樣的功率范圍內(nèi)可以大幅提升 FP32（單精度浮點(diǎn)）和 FP64（雙精度浮點(diǎn)）的運(yùn)算性能。專為深度學(xué)習(xí)設(shè)計(jì)的全新 Tensor Core 在模型訓(xùn)練場(chǎng)景中，最高可以達(dá)到 12 倍速的 TFLOP（每秒萬億次浮點(diǎn)運(yùn)算）。另外，由于全新的 SM 架構(gòu)對(duì)整型和浮點(diǎn)型數(shù)據(jù)采取了相互獨(dú)立且并行的數(shù)據(jù)通路，因此在一般計(jì)算和尋址計(jì)算等混合場(chǎng)景下也能輸出不錯(cuò)的效率。Volta 架構(gòu)新的獨(dú)立線程調(diào)度功能還可以實(shí)現(xiàn)并行線程之間的細(xì)粒度同步和協(xié)作。最后，一個(gè)新組合的 L1 高速數(shù)據(jù)緩存和共享內(nèi)存子系統(tǒng)也顯著提高了性能，同時(shí)大大簡(jiǎn)化了開發(fā)者的編程步驟。
● 第二代 NVLink。第二代 NVIDIA NVLink 高速互連技術(shù)為多 GPU 和多 GPU/CPU系統(tǒng)配置提供了更高的帶寬，更多的連接和更強(qiáng)的可擴(kuò)展性。GV100 GPU 最多支持 6 個(gè) NVLink 鏈路，每個(gè) 25 GB/s，總共 300 GB/s。NVLink 還支持基于 IBM Power 9 CPU 服務(wù)器的 CPU 控制和高速緩存一致性功能。另外，新發(fā)布的 NVIDIA DGX-1V 超級(jí) AI 計(jì)算機(jī)也使用了 NVLink 技術(shù)為超快速的深度學(xué)習(xí)模型訓(xùn)練提供了更強(qiáng)的擴(kuò)展性。
● HBM2 內(nèi)存：更快，更高效。Volta 高度優(yōu)化的 16GB HBM2 內(nèi)存子系統(tǒng)可提供高達(dá) 900 GB/s 的峰值內(nèi)存帶寬。相比上一代 Pascal GP100，來自三星的新一代 HBM2 內(nèi)存與 Volta 的新一代內(nèi)存控制器相結(jié)合，帶寬提升 1.5 倍，并且在性能表現(xiàn)上也超過了 95% 的工作負(fù)載。
● Volta 多處理器服務(wù)（Multi-Process Service，MPS）。Volta MPS 是 Volta GV100 架構(gòu)的一項(xiàng)新特性，可以提供 CUDA MPS 服務(wù)器關(guān)鍵組件的硬件加速功能，從而在共享 GPU 的多計(jì)算任務(wù)場(chǎng)景中顯著提升計(jì)算性能、隔離性和服務(wù)質(zhì)量（QoS）。Volta MPS 還將 MPS 支持的客戶端最大數(shù)量從 Pascal 時(shí)代的 16 個(gè)增加到 48 個(gè)。
● 增強(qiáng)的統(tǒng)一內(nèi)存和地址轉(zhuǎn)換服務(wù)。Volta GV100 中的 GV100 統(tǒng)一內(nèi)存技術(shù)實(shí)現(xiàn)了一個(gè)新的訪問計(jì)數(shù)器，該計(jì)數(shù)器可以根據(jù)每個(gè)處理器的訪問頻率精確調(diào)整內(nèi)存頁的尋址，從而大大提升了處理器之間共享內(nèi)存的使用效率。另外，在 IBM Power 平臺(tái)上，新的地址轉(zhuǎn)換服務(wù)（Address Translation Services，ATS）還允許 GPU 直接訪問 CPU 的存儲(chǔ)頁表。
● Cooperative Groups（協(xié)作組）和新的 Cooperative Launch API（協(xié)作啟動(dòng) API）。Cooperative Groups 是在 CUDA 9 中引入的一種新的編程模型，用于組織通信線程組。Cooperative Groups 允許開發(fā)人員表達(dá)線程之間的溝通粒度，幫助他們更豐富、更有效地進(jìn)行并行分解（decompositions）。Kepler 系列以來，所有的 NVIDIA GPU 都支持基本 Cooperative Groups 特性。Pascal 和 Volta 系列還支持新的 Cooperative Launch API，通過該 API 可以實(shí)現(xiàn) CUDA 線程塊之間的同步。另外 Volta 還增加了對(duì)新的同步模式的支持。
● 最大性能和最高效率兩種模式。顧名思義，在最高性能模式下，Tesla V100 極速器將無限制地運(yùn)行，達(dá)到 300W 的 TDP（熱設(shè)計(jì)功率）級(jí)別，以滿足那些需要最快計(jì)算速度和最高數(shù)據(jù)吞吐量的應(yīng)用需求。而最高效率模式則允許數(shù)據(jù)中心管理員調(diào)整 Tesla V100 的功耗水平，以每瓦特最佳的能耗表現(xiàn)輸出算力。而且，Tesla V100 還支持在所有 GPU 中設(shè)置上限功率，在大大降低功耗的同時(shí)，最大限度地滿足機(jī)架的性能要求。
● 針對(duì) Volta 優(yōu)化的軟件。各種新版本的深度學(xué)習(xí)框架（包括 Caffe2，MXNet，CNTK，TensorFlow 等）都可以利用 Volta 大大縮短模型訓(xùn)練時(shí)間，同時(shí)提升多節(jié)點(diǎn)訓(xùn)練的性能。各種 Volta 優(yōu)化版本的 GPU 加速庫（包括 cuDNN，cuBLAS 和 TensorRT 等）也都可以在 Volta GV100 各項(xiàng)新特性的支持下，為深度學(xué)習(xí)和 HPC 應(yīng)用提供更好的性能支持。此外，NVIDIA CUDA Toolkit 9.0 版也加入了新的 API 和對(duì) Volta 新特性的支持，以幫助開發(fā)者更方便地針對(duì)這些新特性編程。
搭載 Volta GV100 GPU 的 NVIDIA Tesla V100 加速器是當(dāng)今世界上性能最強(qiáng)的并行計(jì)算處理器。其中，GV100 GPU 具有一系列的硬件創(chuàng)新，為深度學(xué)習(xí)算法和框架、HPC 系統(tǒng)和應(yīng)用程序，均提供了強(qiáng)大的算力支持。其中在 HPC 領(lǐng)域的性能表現(xiàn)如下圖所示，在各種 HPC 任務(wù)中，Tesla V100 平均比 Tesla P100 快 1.5 倍（基于 Tesla V100 原型卡）。
Tesla V100擁有業(yè)界領(lǐng)先的浮點(diǎn)和整型運(yùn)算性能，峰值運(yùn)算性能如下（基于 GPU Boost 時(shí)鐘頻率）：
● 雙精度浮點(diǎn)（FP64）運(yùn)算性能：7.5 TFLOP/s；
● 單精度（FP32）運(yùn)算性能：15 TFLOP/s;
● 混合精度矩陣乘法和累加：120 Tensor TFLOP/s。
和之前的 Pascal GP100 一樣，GV100 也由許多圖形處理集群（Graphics Processing Cluster，GPC）、紋理處理集群（Texture Processing Cluster，TPC）、流式多處理器（Streaming Multiprocessor，SM）以及內(nèi)存控制器組成。一個(gè)完整的 GV100 GPU 由 6 個(gè) GPC、84 個(gè) Volta SM、42 個(gè) TPC（每個(gè) TPC 包含 2 個(gè) SM）和 8 個(gè) 512 位的內(nèi)存控制器（共 4096 位）。其中，每個(gè) SM 有 64 個(gè) FP32 核、64 個(gè) INT32 核、32 個(gè) FP64 核與 8 個(gè)全新的 Tensor Core。同時(shí)，每個(gè) SM 也包含了 4 個(gè)紋理處理單元（texture units）。
更具體地說，一個(gè)完整版 Volta GV100 中總共包含了 5376 個(gè) FP32 核、5376 個(gè) INT32 核、2688 個(gè) FP64 核、672 個(gè) Tensor Core 以及 336 個(gè)紋理單元。每個(gè)內(nèi)存控制器都鏈接一個(gè) 768 KB 的 2 級(jí)緩存，每個(gè) HBM2 DRAM 堆棧都由一對(duì)內(nèi)存控制器控制。整體上，GV100 總共包含 6144KB 的二級(jí)緩存。下圖展示了帶有 84 個(gè) SM 單元的完整版 Volta GV100，需要注意的是，不同的產(chǎn)品可能具有不同的配置，比如Tesla V100 就只有 80 個(gè) SM。
下表展示了 Tesla V100 與過去五年歷代 Tesla 系列加速器的參數(shù)對(duì)比。
為了提供更高的性能，Volta SM 具有比舊版 SM 更低的指令和緩存延遲，并且針對(duì)深度學(xué)習(xí)應(yīng)用做了特殊優(yōu)化。其主要特性如下：
● 為深度學(xué)習(xí)矩陣計(jì)算建立的新型混合精度 FP16/FP32 Tensor Core；
● 為更高性能、更低延遲而強(qiáng)化的 L1 高速數(shù)據(jù)緩存；
● 為簡(jiǎn)化解碼和縮短指令延遲而改進(jìn)的指令集；
● 更高的時(shí)鐘頻率和能效。
下圖顯示了 Volta GV100 SM 單元的基本結(jié)構(gòu)。
全新的 Tensor Core 是 Volta GV100 架構(gòu)中最重要的一項(xiàng)新特性，在訓(xùn)練超大型神經(jīng)網(wǎng)絡(luò)模型時(shí)，它可以為系統(tǒng)提供強(qiáng)勁的運(yùn)算性能。Tesla V100 的 Tensor Core 可以為深度學(xué)習(xí)相關(guān)的模型訓(xùn)練和推斷應(yīng)用提供高達(dá) 120 TFLOPS 的浮點(diǎn)張量計(jì)算。具體來說，在深度學(xué)習(xí)的模型訓(xùn)練方面，相比于 P100 上的 FP32 操作，全新的 Tensor Core 可以在 Tesla V100 上實(shí)現(xiàn)最高 12 倍速的峰值 TFLOPS。而在深度學(xué)習(xí)的推斷方面，相比于 P100 上的 FP16 操作，則可以實(shí)現(xiàn)最高 6 倍速的峰值 TFLOPS。Tesla V100 GPU 一共包含 640 個(gè) Tensor Core，每個(gè)流式多處理器（SM）包含 8 個(gè)。
眾所周知，矩陣乘法運(yùn)算是神經(jīng)網(wǎng)絡(luò)訓(xùn)練的核心，在深度神經(jīng)網(wǎng)絡(luò)的每個(gè)連接層中，輸入矩陣都要乘以權(quán)重以獲得下一層的輸入。如下圖所示，相比于上一代 Pascal 架構(gòu)的 GP100，Tesla V100 中的 Tensor Core 把矩陣乘法運(yùn)算的性能提升了至少 9 倍。
如本節(jié)小標(biāo)題所述，Tensor Core 不僅是一個(gè)全新的高效指令集，還是一種數(shù)據(jù)運(yùn)算格式。
在剛發(fā)布的 Volta 架構(gòu)中，每個(gè) Tensor Core 都包含一個(gè) 4x4x4 的矩陣處理隊(duì)列，來完成神經(jīng)網(wǎng)絡(luò)結(jié)構(gòu)中最常見的 D=AxB+C 運(yùn)算。其中 A、B、C、D 是 4 個(gè) 4×4 的矩陣，因此被稱為 4x4x4。如下圖所示，輸入 A、B 是指 FP16 的矩陣，而矩陣 C 和 D 可以是 FP16，也可以是 FP32。
按照設(shè)計(jì)，Tensor Core 在每個(gè)時(shí)鐘頻率可以執(zhí)行高達(dá) 64 次 FMA 混合精度浮點(diǎn)操作，也就是兩個(gè) FP16 輸入的乘積，再加上一個(gè) FP32。而因?yàn)槊總€(gè) SM 單元都包含 8 個(gè) Tensor Core，因此總體上每個(gè)時(shí)鐘可以執(zhí)行 1024 次浮點(diǎn)運(yùn)算。這使得在 Volta 架構(gòu)中，每個(gè) SM 單元的深度學(xué)習(xí)應(yīng)用吞吐量相比標(biāo)準(zhǔn) FP32 操作的 Pascal GP100 大幅提升了 8 倍，與Pascal P100 GPU相比，Volta V100 GPU的吞吐量總共提高了 12 倍。下圖展示了一個(gè)標(biāo)準(zhǔn)的 Volta GV100 Tensor Core 流程。
在程序執(zhí)行期間，多個(gè) Tensor Cores 通過 warp 單元協(xié)同工作。warp 中的線程同時(shí)還提供了可以由 Tensor Cores 處理的更大的 16x16x16 矩陣運(yùn)算。CUDA 將這些操作作為 Warp-Level 級(jí)的矩陣運(yùn)算在 CUDA C++ API 中公開。通過 CUDA C++ 編程，開發(fā)者可以靈活運(yùn)用這些開放 API 實(shí)現(xiàn)基于 Tensor Cores 的乘法、加法和存儲(chǔ)等矩陣操作。
Volta SM 的 L1 高速數(shù)據(jù)緩存和共享內(nèi)存子系統(tǒng)相互結(jié)合，顯著提高了性能，同時(shí)也大大簡(jiǎn)化了開發(fā)者的編程步驟、以及達(dá)到或接近最優(yōu)系統(tǒng)性能的系統(tǒng)調(diào)試成本。
值得強(qiáng)調(diào)的是，Volta 架構(gòu)將數(shù)據(jù)高速緩存和共享內(nèi)存功能組合到單個(gè)內(nèi)存塊中的做法，在整體上為兩種類型的內(nèi)存訪問均提供了最佳的性能。組合后的內(nèi)存容量達(dá)到了 128 KB/SM，比老版的 GP100 高速緩存大 7 倍以上，并且所有這些都可以配置為不共享的獨(dú)享 cache 塊。另外，紋理處理單元也可以使用這些 cache。例如，如果共享內(nèi)存被設(shè)置為 64KB，則紋理和加載/存儲(chǔ)操作就可以使用 L1 中剩余的 64 KB 容量。
總體上，通過和共享內(nèi)存相互組合的獨(dú)創(chuàng)性方式，使得 Volta GV100 L1 高速緩存具有比過去 NVIDIA GPU 的 L1 高速緩存更低的延遲和更高的帶寬。一方面作為流數(shù)據(jù)的高吞吐量管道發(fā)揮作用，另一方面也可以為復(fù)用度很高的數(shù)據(jù)提供高帶寬和低延遲的精準(zhǔn)訪問。
下圖顯示了 Volta 和 Pascal 的 L1 緩存性能對(duì)比。
GV100 GPU 支持英偉達(dá)全新的 Compute Capability 7.0。下表顯示了 NVIDIA GPU 不同架構(gòu)之間的計(jì)算能力對(duì)比。
Volta 架構(gòu)相較之前的 NVIDIA GPU 顯著降低了編程難度，用戶可以更專注于將各種多樣的應(yīng)用產(chǎn)品化。Volta GV100 是第一個(gè)支持獨(dú)立線程調(diào)度的 GPU，也就是說，在程序中的不同線程可以更精細(xì)地同步和協(xié)作。Volta 的一個(gè)主要設(shè)計(jì)目標(biāo)就是降低程序在 GPU 上運(yùn)行所需的開發(fā)成本，以及線程之間靈活的共享機(jī)制，最終使得并行計(jì)算更為高效。
此前的單指令多線程模式（SIMT MODELS）
在 Pascal 和之前的 GPU 中，可以執(zhí)行由 32 個(gè)線程組成的 group，在 SIMT 術(shù)語里也被稱為 warps。在 Pascal 的 warp 里，這 32 個(gè)線程使用同一個(gè)程序計(jì)數(shù)器，然后由一個(gè)激活掩碼（active mask）標(biāo)明 warp 里的哪些線程是有效的。這意味著不同的執(zhí)行路徑里有些線程是“非激活態(tài)”的，下圖給出了一個(gè) warp 里不同分支的順序執(zhí)行過程。在程序中，原始的掩碼會(huì)先被保存起來，直到 warps 執(zhí)行結(jié)束，線程再度收斂，掩碼會(huì)被恢復(fù)，程序再接著執(zhí)行。
從本質(zhì)上來說，Pascal 的 SIMT 模式通過減少跟蹤線程狀態(tài)所需的資源和積極地恢復(fù)線程將并行效率最大化。這種對(duì)整個(gè) warps 進(jìn)行線程狀態(tài)跟蹤的模式，其實(shí)意味著當(dāng)程序出現(xiàn)并行分支時(shí)，warps 內(nèi)部實(shí)際上是順序執(zhí)行的，這里已經(jīng)喪失了并行的意義，直到并行分支的結(jié)束。也就是說，不同 warp 里的線程的確在并行執(zhí)行，但同一 warp 里的分支線程卻在未恢復(fù)之前順序執(zhí)行，它們之間無法交互信息和共享數(shù)據(jù)。
舉個(gè)例子來說，要求數(shù)據(jù)精準(zhǔn)共享的那些算法，在不同的線程訪問被鎖和互斥機(jī)制保護(hù)的數(shù)據(jù)塊時(shí)，因?yàn)椴淮_定遇到的線程是來自哪個(gè) warp，所以很容易導(dǎo)致死鎖。因此，在 Pascal 和之前的 GPU 里，開發(fā)者們不得不避免細(xì)粒度同步，或者使用那些不依賴鎖，或明確區(qū)分 warp 的算法。
Volta 架構(gòu)的單指令多線程模式
Volta 通過在所有線程間（不管是哪個(gè) warp 的）實(shí)施同等級(jí)別的并發(fā)性解決了這一問題，對(duì)每個(gè)線程，包括程序計(jì)數(shù)器和調(diào)用棧，Volta 都維護(hù)同一個(gè)執(zhí)行狀態(tài)，如下圖所示。
Volta 的獨(dú)立線程調(diào)配機(jī)制允許 GPU 將執(zhí)行權(quán)限讓步于任何一個(gè)線程，這樣做使線程的執(zhí)行效率更高，同時(shí)也讓線程間的數(shù)據(jù)共享更合理。為了最大化并行效率，Volta 有一個(gè)調(diào)度優(yōu)化器，可以決定如何對(duì)同一個(gè) warp 里的有效線程進(jìn)行分組，并一起送到 SIMT 單元。這不僅保持了在 NVIDIA 之前的 GPU 里較高的 SIMT 吞吐量，而且靈活性更高：現(xiàn)在，線程可以在 sub-warp 級(jí)別上分支和恢復(fù)，并且，Volta 仍將那些執(zhí)行相同代碼的線程分組在一起，讓他們并行運(yùn)行。
下圖展示了 Volta 多線程模式的一個(gè)樣例。這個(gè)程序里的 if/else 分支現(xiàn)在可以按照時(shí)序被間隔開來，如圖12所示?？梢钥吹?，執(zhí)行過程依然是 SIMT 的，在任意一個(gè)時(shí)鐘周期，和之前一樣，同一個(gè) warp 里的所有有效線程，CUDA 核執(zhí)行的是同樣的指令，這樣依然可以保持之前架構(gòu)中的執(zhí)行效率。重點(diǎn)是，Volta 的這種獨(dú)立調(diào)度能力，可以讓程序員有機(jī)會(huì)用更加自然的方式開發(fā)出復(fù)雜且精細(xì)的算法和數(shù)據(jù)結(jié)構(gòu)。雖然調(diào)度器支持線程執(zhí)行的獨(dú)立性，但它依然會(huì)優(yōu)化那些非同步的代碼段，在確保線程收斂的同時(shí)，最大限度地提升 SIMT 的高效性。
另外，上圖中還有一個(gè)有趣的現(xiàn)象：Z 在所有的線程中都不是同一時(shí)刻執(zhí)行的。這是因?yàn)?Z 可能會(huì)輸出其它分支進(jìn)程需要的數(shù)據(jù)，在這種情況下，強(qiáng)制進(jìn)行收斂并不安全。但在之前的架構(gòu)中，一般認(rèn)為 A，B，X，Y 并不包含同步性操作，因此調(diào)度器會(huì)認(rèn)定在 Z 上收斂是安全的。
在這種情況下，程序可以調(diào)用新的 CUDA 9 中的 warp 同步函數(shù) __syncwarp() 來強(qiáng)制進(jìn)行線程收斂，如下圖所示。這時(shí)分支線程可能并不會(huì)同步執(zhí)行 Z，但是通過調(diào)用 __syncwarp() 函數(shù)，同一個(gè) warp 里的這些線程的所有執(zhí)行路徑將會(huì)在執(zhí)行到 Z 語句之前完備。類似的，在執(zhí)行 Z 之前，如果調(diào)用一下 __syncwarp() 函數(shù)，則程序?qū)?huì)在執(zhí)行 Z 之前強(qiáng)制收斂。如果開發(fā)者能提前確保這種操作的安全性，無疑這會(huì)在一定程度上提升 SIMT 的執(zhí)行效率。
Starvation-Free 算法
Starvation-free 算法是獨(dú)立線程調(diào)度機(jī)制的一個(gè)重要模式，具體是指：在并發(fā)計(jì)算中，只要系統(tǒng)確保所有線程具有對(duì)競(jìng)爭(zhēng)性資源的恰當(dāng)訪問權(quán)，就可以保證其正確執(zhí)行。例如，如果嘗試獲取互斥鎖（mutex）的線程最終成功獲得了該鎖，就可以在 starvation-free 算法中使用互斥鎖（或普通鎖）。在不支持 starvation-free 算法的系統(tǒng)中，可能會(huì)出現(xiàn)一個(gè)或多個(gè)線程重復(fù)獲取和釋放互斥鎖的情況，這就有可能造成其他線程始終無法成功獲取互斥鎖的問題。
下面看一個(gè)關(guān)于 Volta 獨(dú)立線程調(diào)度的實(shí)例：在多線程應(yīng)用程序中將節(jié)點(diǎn)插入雙向鏈表。
__device__ void insert_after(Node *a, Node *b)
{
Node *c;
lock(a); lock(a->next);
c = a->next;
a->next = b;
b->prev = a;
b->next = c;
c->prev = b;
unlock(c); unlock(a);
}
在這個(gè)例子中，每個(gè)雙向鏈表的元素至少含有 3 個(gè)部分：一個(gè)后向指針，一個(gè)前向指針，以及一個(gè) lock（只有 owner 才有權(quán)限更新結(jié)點(diǎn)）。下圖展示了在 A 和 C 之間插入 B 結(jié)點(diǎn)的過程。
Volta 這種獨(dú)立線程調(diào)度機(jī)制可以確保即使線程 T0 目前鎖住了結(jié)點(diǎn) A，同一個(gè) warp 里的另一個(gè)線程 T1 依然可以成功地等到其解鎖，而不影響 T0 的執(zhí)行。不過，值得注意的一點(diǎn)是，因?yàn)橥粋€(gè) warp 下的有效線程是一起執(zhí)行的，所以等解鎖的線程可能會(huì)讓鎖住的線程性能降低。
同樣需要重視的是，如此例中這種針對(duì)每個(gè)結(jié)點(diǎn)上鎖的用法對(duì) GPU 的性能影響至關(guān)重要。傳統(tǒng)上，雙向鏈接表的創(chuàng)建可能會(huì)用粗粒度 lock（對(duì)應(yīng)前面提到的細(xì)粒度 lock），粗粒度 lock 會(huì)獨(dú)占整個(gè)結(jié)構(gòu)（全部上鎖），而不是對(duì)每一個(gè)結(jié)點(diǎn)分別予以保護(hù)。由于線程間對(duì) lock 的爭(zhēng)奪，因此這種方法可能會(huì)導(dǎo)致多線程代碼的性能下降（Volta 架構(gòu)最多允許高達(dá) 163,840 個(gè)并發(fā)線程）。這時(shí)可以嘗試在每個(gè)節(jié)點(diǎn)采用細(xì)粒度 lock 的辦法，這樣除了在某些特定節(jié)點(diǎn)的插入操作之外，大型列表中平均每個(gè)節(jié)點(diǎn)的 lock 競(jìng)爭(zhēng)效應(yīng)就會(huì)大大降低。
上述這種具備細(xì)粒度 lock 的雙向鏈接表只是個(gè)非常簡(jiǎn)單的例子，我們想通過這個(gè)例子傳達(dá)的信息是：通過獨(dú)立的線程調(diào)度機(jī)制，開發(fā)者們可以用最自然的方式在 NVIDIA GPU 上實(shí)現(xiàn)熟悉的算法和數(shù)據(jù)結(jié)構(gòu)。
NVIDIA Tesla V100 無疑是目前世界上最先進(jìn)的數(shù)據(jù)中心 GPU，專門用于處理需要強(qiáng)大計(jì)算能力支持的密集型 HPC、AI、和圖形處理任務(wù)。憑借最先進(jìn)的 NVIDIA Volta 架構(gòu)支持，Tesla V100 可以在單片 GPU 中提供 100 個(gè) CPU 的運(yùn)算性能，這使得數(shù)據(jù)科學(xué)家、研究人員和工程師們得以應(yīng)對(duì)曾經(jīng)被認(rèn)為是不可能的挑戰(zhàn)。
搭載 640 個(gè) Tensor cores，使得 Tesla V100 成為了目前世界上第一款突破 100 TFLOPS 算力大關(guān)的深度學(xué)習(xí) GPU 產(chǎn)品。再加上新一代 NVIDIA NVLink 技術(shù)高達(dá) 300 GB/s 的連接能力，現(xiàn)實(shí)場(chǎng)景中用戶完全可以將多個(gè) V100 GPU 組合起來搭建一個(gè)強(qiáng)大的深度學(xué)習(xí)運(yùn)算中心。這樣，曾經(jīng)需要數(shù)周時(shí)間的 AI 模型現(xiàn)在可以在幾天之內(nèi)訓(xùn)練完成。而隨著訓(xùn)練時(shí)間的大幅度縮短，未來所有的現(xiàn)實(shí)問題或許都將被 AI 解決。
雷鋒網(wǎng)(公眾號(hào)：雷鋒網(wǎng))相關(guān)閱讀：
2小時(shí)、5大AI 新品、英偉達(dá)股價(jià)暴漲17%，GTC大會(huì)上黃仁勛都講了些啥？（內(nèi)附PPT） | GTC 2017
GTC大會(huì)第二日亮點(diǎn)：NVIDIA將推出多用戶VR系統(tǒng)，計(jì)劃培養(yǎng)100000名開發(fā)人員 | GTC 2017
雷峰網(wǎng)版權(quán)文章，未經(jīng)授權(quán)禁止轉(zhuǎn)載。詳情見轉(zhuǎn)載須知。
