基于 NVIDIA Ampere GPU 架構(gòu)的新型 NVIDIA A100 GPU 在加速計(jì)算方面實(shí)現(xiàn)了最大的一代飛躍。 A100 GPU 具有革命性的硬件功能，我們很高興宣布 CUDA 11 與 A100 結(jié)合使用。

CUDA 11 使您能夠利用新的硬件功能來加速 HPC 、基因組學(xué)、 5G 、渲染、深度學(xué)習(xí)、數(shù)據(jù)分析、數(shù)據(jù)科學(xué)、機(jī)器人技術(shù)和更多不同的工作負(fù)載。

CUDA 11 包含了從平臺(tái)系統(tǒng)軟件到您開始開發(fā) GPU 加速應(yīng)用程序所需的所有功能。本文概述了此版本中的主要軟件功能：

支持 NVIDIA Ampere GPU 架構(gòu)，包括新的 NVIDIA A100 GPU ，用于加速 AI 和 HPC 數(shù)據(jù)中心的擴(kuò)展和擴(kuò)展；具有 NVSwitch 結(jié)構(gòu)的多 GPU 系統(tǒng)，如 DGX A100 型 和 HGX A100 型 。

多實(shí)例 GPU （ MIG ）分區(qū)功能，這對(duì)云服務(wù)提供商（ csp ）提高 GPU 利用率特別有利。

新的第三代張量核心加速混合精度，不同數(shù)據(jù)類型的矩陣運(yùn)算，包括 TF32 和 Bfloat16 。

用于任務(wù)圖、異步數(shù)據(jù)移動(dòng)、細(xì)粒度同步和二級(jí)緩存駐留控制的編程和 API 。

CUDA 庫中用于線性代數(shù)、 fft 和矩陣乘法的性能優(yōu)化。

Nsight 產(chǎn)品系列的更新，用于跟蹤、分析和調(diào)試 CUDA 應(yīng)用程序。

全面支持所有主要的 CPU 體系結(jié)構(gòu)，跨 x86 趶 64 、 Arm64 服務(wù)器和電源體系結(jié)構(gòu)。

一篇文章不能公正地反映 CUDA 11 中提供的每一個(gè)特性。在這篇文章的最后，有一些鏈接到 GTC 數(shù)字會(huì)議，這些會(huì)議提供了對(duì)新的 CUDA 特性的深入探討。

CUDA 和 NVIDIA Ampere微體系結(jié)構(gòu) GPUs

NVIDIA Ampere GPU 微體系結(jié)構(gòu)采用 TSMC 7nm N7 制造工藝制造，包括更多的流式多處理器（ SMs ）、更大更快的內(nèi)存以及與第三代 NVLink 互連帶寬，以提供巨大的計(jì)算吞吐量。

A100 的 40 GB （ 5 站點(diǎn)）高速 HBM2 內(nèi)存的帶寬為 1 。 6 TB /秒，比 V100 快 1 。 7 倍以上。 A100 上 40 MB 的二級(jí)緩存幾乎比 Tesla V100 大 7 倍，提供 2 倍以上的二級(jí)緩存讀取帶寬。 CUDA 11 在 A100 上提供了新的專用二級(jí)緩存管理和駐留控制 API 。 A100 中的短消息包括一個(gè)更大更快的一級(jí)緩存和共享內(nèi)存單元（每平方米 192 千字節(jié)），提供 1.5 倍于沃爾塔 V100GPU 的總?cè)萘俊?/p>

A100 配備了專門的硬件單元，包括第三代張量核心、更多視頻解碼器（ NVDEC ）單元、 JPEG 解碼器和光流加速器。所有這些都被各種 CUDA 庫用來加速 HPC 和 AI 應(yīng)用程序。

接下來的幾節(jié)將討論 NVIDIA A100 中引入的主要?jiǎng)?chuàng)新，以及 CUDA 11 如何使您充分利用這些功能。 CUDA 11 為每個(gè)人提供了一些東西，無論你是管理集群的平臺(tái) DevOps 工程師，還是編寫 GPU 加速應(yīng)用程序的軟件開發(fā)人員。

多實(shí)例

MIG 功能可以將單個(gè) A100 GPU 物理劃分為多個(gè) GPUs 。它允許多個(gè)客戶機(jī)（如 vm 、容器或進(jìn)程）同時(shí)運(yùn)行，同時(shí)在這些程序之間提供錯(cuò)誤隔離和高級(jí)服務(wù)質(zhì)量（ QoS ）。

圖 1 A100 中的新 MIG 功能。

A100 是第一款 GPU 可以通過 NVLink 擴(kuò)展到完整的 GPU ，也可以通過降低每個(gè) GPU 實(shí)例的成本，使用 MIG 擴(kuò)展到許多用戶。 MIG 支持多個(gè)用例來提高 GPU 的利用率。這可能是 CSP 租用單獨(dú)的 GPU 實(shí)例，在 GPU 上運(yùn)行多個(gè)推理工作負(fù)載，托管多個(gè) Jupyter 筆記本會(huì)話進(jìn)行模型探索，或者在組織中的多個(gè)內(nèi)部用戶（單個(gè)租戶、多用戶）之間共享 GPU 的資源。

nvidia-smi 對(duì) CUDA 是透明的，現(xiàn)有的 CUDA 程序可以在 MIG 下運(yùn)行，以減少編程工作量。 CUDA 11 允許使用 NVIDIA 管理庫（ NVML ）或其命令行界面 NVIDIA （ nvidia-smi MIG 子命令）在 Linux 操作系統(tǒng)上配置和管理 MIG 實(shí)例。

使用啟用了 MIG 的 集裝箱工具包 和 A100 ，您還可以使用 Docker 運(yùn)行 GPU 容器（使用從 Docker 19 。 03 開始的 --gpus 選項(xiàng)）或使用 NVIDIA 設(shè)備插件 擴(kuò)展 Kubernetes 容器平臺(tái)。

下面的命令顯示了使用 nvidia-smi 進(jìn)行的 MIG 管理：

# List gpu instance profiles:
# nvidia-smi mig -i 0 -lgip
+-------------------------------------------------------------------------+
| GPU instance profiles: |
| GPU Name ID Instances Memory P2P SM DEC ENC |
| Free/Total GiB CE JPEG OFA |
|=========================================================================|
| 0 MIG 1g.5gb 19 0/7 4.95 No 14 0 0 |
| 1 0 0 |
+-------------------------------------------------------------------------+
| 0 MIG 2g.10gb 14 0/3 9.90 No 28 1 0 |
| 2 0 0 |
+-------------------------------------------------------------------------+
| 0 MIG 3g.20gb 9 0/2 19.81 No 42 2 0 |
| 3 0 0 |
+-------------------------------------------------------------------------+
| 0 MIG 4g.20gb 5 0/1 19.81 No 56 2 0 |
| 4 0 0 |
+-------------------------------------------------------------------------+
| 0 MIG 7g.40gb 0 0/1 39.61 No 98 5 0 |
| 7 1 1 |
+-------------------------------------------------------------------------+ 

系統(tǒng)軟件平臺(tái)支持

為企業(yè)級(jí)應(yīng)用程序 NVIDIA 引入了新的恢復(fù)功能，避免了 NVIDIA A100 對(duì)內(nèi)存使用的影響。先前架構(gòu)上不可糾正的 ECC 錯(cuò)誤將影響 GPU 上所有正在運(yùn)行的工作負(fù)載，需要重置 GPU 。

在 A100 上，影響僅限于遇到錯(cuò)誤并被終止的應(yīng)用程序，而其他正在運(yùn)行的 CUDA 工作負(fù)載不受影響。 GPU 不再需要重置來恢復(fù)。 NVIDIA 驅(qū)動(dòng)程序執(zhí)行動(dòng)態(tài)頁面黑名單以標(biāo)記頁面不可用，以便當(dāng)前和新的應(yīng)用程序不會(huì)訪問受影響的內(nèi)存區(qū)域。

當(dāng) GPU 被重置時(shí)，作為常規(guī) GPU / VM 服務(wù)窗口的一部分， A100 配備了一種稱為行重映射的新硬件機(jī)制，它用備用單元替換內(nèi)存中降級(jí)的單元，并避免在物理內(nèi)存地址空間中造成任何漏洞。

帶有 CUDA 11 的 NVIDIA 驅(qū)動(dòng)程序現(xiàn)在報(bào)告了與帶內(nèi)（使用 NVML / NVIDIA -smi ）和帶外（使用系統(tǒng) BMC ）行重映射相關(guān)的各種度量。 A100 包括新的帶外功能，包括更多可用的 GPU 和 NVSwitch 遙測(cè)、控制和改進(jìn)的 GPU 和 BMC 之間的總線傳輸數(shù)據(jù)速率。

為了提高多 GPU 系統(tǒng)（如 DGX A100 和 HGX A100 ）的彈性和高可用性，系統(tǒng)軟件支持禁用發(fā)生故障的 GPU 或 NVSwitch 節(jié)點(diǎn)的能力，而不是像前幾代系統(tǒng)中那樣禁用整個(gè)基板。

CUDA 11 是第一個(gè)為 Arm 服務(wù)器添加生產(chǎn)支持的版本。通過將 Arm 的 節(jié)能 CPU 體系結(jié)構(gòu) 與 CUDA 相結(jié)合， Arm 生態(tài)系統(tǒng)將受益于 GPU ——一種針對(duì)各種用例的加速計(jì)算：從邊緣、云、游戲到超級(jí)計(jì)算機(jī)。 CUDA 11 支持 Marvell 的高性能 基于 ThunderX2 服務(wù)器，并與 Arm 和生態(tài)系統(tǒng)中的其他硬件和軟件合作伙伴密切合作，以快速支持 GPUs 。

第三代多精度張量核

NVIDIA A100 中每個(gè) SM 的四個(gè)大張量核（總共 432 個(gè)張量核）為所有數(shù)據(jù)類型提供更快的矩陣乘法累加（ MMA ）操作： Binary 、 INT4 、 INT8 、 FP16 、 Bfloat16 、 TF32 和 FP64 。

您可以通過不同的深度學(xué)習(xí)框架、 CUDA C ++模板抽象來訪問張量核， CUTLASS 、 CUDA 庫 cuBLAS 、 CuSOLVER 、 CursSuror 或 TensorRT 。

CUDA C ++利用張量矩陣（ WMMA ） API 使張量核可用。這種便攜式 API 抽象公開了專門的矩陣加載、矩陣乘法和累加運(yùn)算以及矩陣存儲(chǔ)操作，以有效地使用 CUDA C ++程序的張量核。 WMMA 的所有函數(shù)和數(shù)據(jù)類型都可以在 nvcuda：：wmma 命名空間中使用。您還可以使用 mma_sync PTX 指令直接訪問 A100 的張量核心（即具有計(jì)算能力 compute _及更高版本的設(shè)備）。

CUDA 11 增加了對(duì)新的輸入數(shù)據(jù)類型格式的支持： Bfloat16 、 TF32 和 FP64 。 Bfloat16 是另一種 FP16 格式，但精度降低，與 FP32 數(shù)值范圍匹配。它的使用會(huì)降低帶寬和存儲(chǔ)需求，從而提高吞吐量。 BFLAT16 作為一個(gè)新的 CUDA C ++ __nv_bfloat16 數(shù)據(jù)類型在 CUDA α BF16 。 H 中通過 WMMA 公開，并由各種 CUDA 數(shù)學(xué)庫支持。

TF32 是一種特殊的浮點(diǎn)格式，用于張量核心。 TF32 包括一個(gè) 8 位指數(shù)（與 FP32 相同）、 10 位尾數(shù)（精度與 FP16 相同）和一個(gè)符號(hào)位。這是默認(rèn)的數(shù)學(xué)模式，允許您在不改變模型的情況下，獲得超過 FP32 的 DL 訓(xùn)練加速。最后， A100 為 MMA 操作帶來了雙精度（ FP64 ）支持， WMMA 接口也支持這種支持。

圖 2 矩陣運(yùn)算支持的數(shù)據(jù)類型、配置和性能表。

編程 NVIDIA Ampere架構(gòu) GPUs

為了提高 GPU 的可編程性和利用 NVIDIA A100 GPU 的硬件計(jì)算能力， CUDA 11 包括內(nèi)存管理、任務(wù)圖加速、新指令和線程通信結(jié)構(gòu)的新 API 操作。下面我們來看看這些新操作，以及它們?nèi)绾问鼓軌蚶?A100 和 NVIDIA Ampere微體系結(jié)構(gòu)。

內(nèi)存管理

最大化 GPU 內(nèi)核性能的優(yōu)化策略之一是最小化數(shù)據(jù)傳輸。如果內(nèi)存駐留在全局內(nèi)存中，將數(shù)據(jù)讀入二級(jí)緩存或共享內(nèi)存 MIG ht 的延遲需要幾百個(gè)處理器周期。

例如，在 GV100 上，共享內(nèi)存提供的帶寬比全局內(nèi)存快 17 倍，比 L2 快 3 倍。因此，一些具有生產(chǎn)者 – 消費(fèi)者范式的算法可以觀察到在內(nèi)核之間的 L2 中持久化數(shù)據(jù)的性能優(yōu)勢(shì)，從而獲得更高的帶寬和性能。

在 A100 上， CUDA 11 提供了 API 操作，以留出 40 MB 二級(jí)緩存的一部分來持久化對(duì)全局內(nèi)存的數(shù)據(jù)訪問。持久訪問優(yōu)先使用二級(jí)緩存的這一預(yù)留部分，而全局內(nèi)存的正?；蛄魇皆L問只能在持久訪問未使用的情況下使用這部分二級(jí)緩存。

L2 持久性可以設(shè)置為在 CUDA 流或 CUDA 圖內(nèi)核節(jié)點(diǎn)中使用。在留出二級(jí)緩存區(qū)域時(shí)，需要考慮一些問題。例如，在不同的流中并發(fā)執(zhí)行的多個(gè) CUDA 內(nèi)核在具有不同的訪問策略窗口時(shí)，共享二級(jí)預(yù)留緩存。下面的代碼示例顯示為持久性預(yù)留二級(jí)緩存比率。

cudaGetDeviceProperties( &prop, device_id);
// Set aside 50% of L2 cache for persisting accesses size_t size = min( int(prop.l2CacheSize * 0.50) , prop.persistingL2CacheMaxSize );
cudaDeviceSetLimit( cudaLimitPersistingL2CacheSize, size); // Stream level attributes data structure cudaStreamAttrValue attr ;? attr.accessPolicyWindow.base_ptr = /* beginning of range in global memory */ ;? attr.accessPolicyWindow.num_bytes = /* number of bytes in range */ ;? // hitRatio causes the hardware to select the memory window to designate as persistent in the area set-aside in L2 attr.accessPolicyWindow.hitRatio = /* Hint for cache hit ratio */ // Type of access property on cache hit attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;?
// Type of access property on cache miss
attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming; cudaStreamSetAttribute(stream,cudaStreamAttributeAccessPolicyWindow,&attr);

虛擬內(nèi)存管理 API 操作 已經(jīng)擴(kuò)展到支持對(duì)固定的 GPU 內(nèi)存的壓縮，以減少 L2 到 DRAM 的帶寬。這對(duì)于深度學(xué)習(xí)訓(xùn)練和推理用例非常重要。使用 cuMemCreate 創(chuàng)建可共享內(nèi)存句柄時(shí)，將向 API 操作提供分配提示。

算法（如 3D 模板或卷積）的有效實(shí)現(xiàn)涉及內(nèi)存復(fù)制和計(jì)算控制流模式，其中數(shù)據(jù)從全局內(nèi)存?zhèn)鬏數(shù)骄€程塊的共享內(nèi)存，然后進(jìn)行使用該共享內(nèi)存的計(jì)算。全局到共享內(nèi)存的復(fù)制擴(kuò)展為從全局內(nèi)存讀取到寄存器，然后寫入共享內(nèi)存。

CUDA 11 允許您利用一種新的異步復(fù)制（ async-copy ）范式。它本質(zhì)上是將數(shù)據(jù)從全局復(fù)制到共享內(nèi)存與計(jì)算重疊，并避免使用中間寄存器或一級(jí)緩存。異步復(fù)制的好處是：控制流不再遍歷內(nèi)存管道兩次，不使用中間寄存器可以減少寄存器壓力，增加內(nèi)核占用。在 A100 上，異步復(fù)制操作是硬件加速的。

下面的代碼示例顯示了一個(gè)使用異步復(fù)制的簡(jiǎn)單示例。生成的代碼雖然性能更好，但可以通過對(duì)多批異步復(fù)制操作進(jìn)行管道化來進(jìn)一步優(yōu)化。這種額外的流水線可以消除代碼中的一個(gè)同步點(diǎn)。

異步復(fù)制是在 CUDA 11 中作為實(shí)驗(yàn)特性提供的，并且使用協(xié)作組集合公開。 CUDA C ++程序設(shè)計(jì)指南 包括在 A100 中使用多級(jí)流水線和硬件加速屏障操作的異步復(fù)制的更高級(jí)示例。

//Without async-copy using namespace nvcuda::experimental;
__shared__ extern int smem[];? // algorithm loop iteration
?while ( ... ) {? __syncthreads(); // load element into shared mem for ( i = ... ) {? // uses intermediate register // {int tmp=g[i]; smem[i]=tmp;} smem[i] = gldata[i]; ? }?

//With async-copy using namespace nvcuda::experimental;
__shared__ extern int smem[];? pipeline pipe; // algorithm loop iteration
?while ( ... ) {? __syncthreads(); // load element into shared mem for ( i = ... ) {? // initiate async memory copy memcpy_async(smem[i], gldata[i], pipe); ? }? // wait for async-copy to complete pipe.commit_and_wait(); __syncthreads();? /* compute on smem[] */? }?

任務(wù)圖加速

CUDA 10 中引入的 CUDA 圖 代表了使用 CUDA 提交工作的新模型。圖由一系列操作組成，如內(nèi)存拷貝和內(nèi)核啟動(dòng)，它們由依賴關(guān)系連接，并在執(zhí)行時(shí)單獨(dú)定義。

圖形支持定義一次重復(fù)運(yùn)行的執(zhí)行流。它們可以減少累積的啟動(dòng)開銷并提高應(yīng)用程序的整體性能。對(duì)于深度學(xué)習(xí)應(yīng)用程序來說尤其如此，這些應(yīng)用程序可能會(huì)啟動(dòng)多個(gè)內(nèi)核，任務(wù)大小和運(yùn)行時(shí)會(huì)減少，或者任務(wù)之間可能存在復(fù)雜的依賴關(guān)系。

從 A100 開始， GPU 提供任務(wù)圖硬件加速來預(yù)取網(wǎng)格啟動(dòng)描述符、指令和常量。與以前的 GPUs 如 V100 相比，在 A100 上使用 CUDA 圖改善了內(nèi)核啟動(dòng)延遲。

CUDA Graph API 操作現(xiàn)在有了一個(gè)輕量級(jí)的機(jī)制，可以支持對(duì)實(shí)例化圖的就地更新，而不需要重新構(gòu)建圖。在圖的重復(fù)實(shí)例化過程中，節(jié)點(diǎn)參數(shù)（如內(nèi)核參數(shù)）通常會(huì)在圖拓?fù)浔３植蛔兊那闆r下發(fā)生變化。 Graph API 操作提供了一種更新整個(gè)圖形的機(jī)制，您可以使用更新的節(jié)點(diǎn)參數(shù)提供拓?fù)湎嗤?cudaGraph_t 對(duì)象，或者顯式更新單個(gè)節(jié)點(diǎn)。

另外， CUDA 圖現(xiàn)在支持協(xié)同內(nèi)核啟動(dòng)（ cuLaunchCooperativeKernel ），包括與 CUDA 流進(jìn)行奇偶校驗(yàn)的流捕獲。

線集體

以下是 CUDA 11 添加到 合作小組 中的一些增強(qiáng)，這些增強(qiáng)在 CUDA 9 中引入。協(xié)作組是一種集體編程模式，旨在使您能夠顯式地表達(dá)線程可以在其中通信的粒度。這使得 CUDA 內(nèi)的新的協(xié)作并行模式得以實(shí)現(xiàn)。

在 CUDA 11 中，協(xié)作組集體公開了新的 A100 硬件特性，并添加了幾個(gè) API 增強(qiáng)。

A100 引入了一個(gè)新的 reduce 指令，它對(duì)每個(gè)線程提供的數(shù)據(jù)進(jìn)行操作。這是一個(gè)新的使用協(xié)作組的集合，它提供了一個(gè)可移植的抽象，也可以在舊的體系結(jié)構(gòu)上使用。 reduce 操作支持算術(shù)（例如 add ）和邏輯（例如和）操作。下面的代碼示例顯示 reduce 集合。

// Simple Reduction Sum
#include  ... const int threadId = cta.thread_rank(); int val = A[threadId]; // reduce across tiled partition reduceArr[threadId] = cg::reduce(tile, val, cg::plus()); // synchronize partition cg::sync(cta); // accumulate sum using a leader and return sum

協(xié)作組提供了集體操作（ labeled_partition ），將父組劃分為一維子組，在這些子組中線程被合并在一起。這對(duì)于試圖通過條件語句的基本塊跟蹤活動(dòng)線程的控制流特別有用。

例如，可以使用 labeled_partition 在一個(gè) warp 級(jí)別組（不限制為 2 的冪）中形成多個(gè)分區(qū)，并在原子加法操作中使用。 labeled_partition API 操作計(jì)算條件標(biāo)簽，并將具有相同標(biāo)簽值的線程分配到同一組中。

下面的代碼示例顯示自定義線程分區(qū)：

// Get current active threads (that is, coalesced_threads())
cg::coalesced_group active = cg::coalesced_threads(); // Match threads with the same label using match_any() int bucket = active.match_any(value); cg::coalesced_group subgroup = cg::labeled_partition(active, bucket); // Choose a leader for each partition (for example, thread_rank = 0)
// if (subgroup.thread_rank() == 0) { threadId = atomicAdd(&addr[bucket], subgroup.size()); } // Now use shfl to transfer the result back to all threads in partition
return (subgroup.shfl(threadId, 0));

CUDA C ++語言及其編譯器改進(jìn)

CUDA 11 也是第一個(gè)正式將 CUB 作為 CUDA 工具包一部分的版本。 CUB 現(xiàn)在是支持的 CUDA C ++核心庫之一。

CUDA 11 的 nvcc 公司 的一個(gè)主要特性是支持鏈路時(shí)間優(yōu)化（ LTO ），以提高單獨(dú)編譯的性能。 LTO 使用 --dlink-time-opt 或 -dlto 選項(xiàng)，在編譯期間存儲(chǔ)中間代碼，然后在鏈接時(shí)執(zhí)行更高級(jí)別的優(yōu)化，例如跨文件內(nèi)聯(lián)代碼。

nvcc 公司 在 CUDA 11 中添加了對(duì) ISO C ++ 17 的支持，并支持 PGI 、 GCC 、 CLAN 、 ARM 和微軟 VisualStudio 上的新主機(jī)編譯器。如果您想嘗試尚未支持的主機(jī)編譯器， CUDA 在編譯構(gòu)建工作流期間支持一個(gè)新的 nvcc 公司 標(biāo)志。 --allow-unsupported-compiler 添加了其他新功能，包括：

改進(jìn)的 lambda 支持

依賴文件生成增強(qiáng)功能（ -MD ， -MMD 選項(xiàng)）

傳遞選項(xiàng)到主機(jī)編譯器

圖 4 CUDA 11 中的平臺(tái)支撐。

CUDA 庫

CUDA 11 中的庫通過使用線性代數(shù)、信號(hào)處理、基本數(shù)學(xué)運(yùn)算和圖像處理中常見的 api 中的最新和最好的 A100 硬件特性，繼續(xù)推動(dòng)性能和開發(fā)人員生產(chǎn)力的邊界。

在整個(gè)線性代數(shù)庫中，您將看到 A100 上提供的全系列精度的張量核心加速度，包括 FP16 、 Bfloat16 、 TF32 和 FP64 。這包括 cuBLAS 中的 BLAS3 運(yùn)算， cuSOLVER 中的因子分解和稠密線性解算器，以及 cuTENSOR 中的張量收縮。

除了提高精度范圍外，對(duì)矩陣維數(shù)和張量核心加速度對(duì)齊的限制也被取消。為了達(dá)到適當(dāng)?shù)木?，加速現(xiàn)在是自動(dòng)的，不需要用戶選擇加入。 cuBLAS 的啟發(fā)式算法在 A100 上使用 MIG 在 GPU 實(shí)例上運(yùn)行時(shí)自動(dòng)適應(yīng)資源。

圖 6 用 cuBLAS 對(duì) A100 進(jìn)行混合精度矩陣乘法。

彎刀 、 CUDA C ++模板用于高性能 GEMM 的抽象，支持 A100 提供的各種精度模式。有了 CUDA 11 ， Cutless 現(xiàn)在達(dá)到了 95% 以上的性能與 cuBLAS 相當(dāng)。這允許您編寫您自己的自定義 CUDA 內(nèi)核來編程 NVIDIA GPUs 中的張量核心。

cuFFT 利用了 A100 中更大的共享內(nèi)存大小，從而在較大的批處理大小下為單精度 fft 帶來更好的性能。最后，在多個(gè) GPU A100 系統(tǒng)上，與 V100 相比， cuFFT 可擴(kuò)展并提供 2 倍于 GPU 的性能。

nvJPEG 是一個(gè)用于 JPEG 解碼的 GPU 加速庫。它與數(shù)據(jù)擴(kuò)充和圖像加載庫 NVIDIA DALI 一起，可以加速圖像分類模型，特別是計(jì)算機(jī)視覺的深度學(xué)習(xí)訓(xùn)練。圖書館加速了深度學(xué)習(xí)工作流的圖像解碼和數(shù)據(jù)擴(kuò)充階段。

A100 包括一個(gè) 5 核硬件 JPEG 解碼引擎， nvJPEG 利用硬件后端對(duì) JPEG 圖像進(jìn)行批處理。專用硬件塊的 JPEG 加速緩解了 CPU 上的瓶頸，并允許更好的 GPU 利用率。

選擇硬件解碼器是由給定圖像的 nvjpegDecode 自動(dòng)完成的，或者通過使用 nvjpegCreateEx init 函數(shù)顯式選擇硬件后端來完成的。 nvJPEG 提供基線 JPEG 解碼的加速，以及各種顏色轉(zhuǎn)換格式，例如 yuv420 、 422 和 444 。

圖 8 顯示，與僅使用 CPU 處理相比，這將使圖像解碼速度提高 18 倍。如果您使用 DALI ，您可以直接受益于這種硬件加速，因?yàn)?nvJPEG 是抽象的。

圖 9 nvJPEG 加速比 CPU 。

（第 128 批使用 Intel Platinum 8168 @ 2GHz 3 。 7GHz Turbo HT ；使用 TurboJPEG ）

在 CUDA 數(shù)學(xué)庫中，有許多特性比一篇文章所能涵蓋的還要多。

開發(fā)人員工具

CUDA 11 繼續(xù)為現(xiàn)有的開發(fā)人員工具組合添加豐富的功能，其中包括熟悉的 Visual Studio 插件和 NVIDIA Nsight 集成 for Visual Studio ，以及 Eclipse 和 Nsight Eclipse 插件版。它還包括獨(dú)立的工具，如用于內(nèi)核評(píng)測(cè)的 Nsight Compute 和用于系統(tǒng)范圍性能分析的 Nsight 系統(tǒng)。 Nsight Compute 和 Nsight 系統(tǒng)現(xiàn)在在 CUDA 支持的三種 CPU 架構(gòu)上都得到支持： x86 、 POWER 和 Arm64 。

Nsight Compute for CUDA 11 的一個(gè)關(guān)鍵特性是能夠生成應(yīng)用程序的屋頂線模型。 Roofline 模型是一種直觀直觀的方法，可以讓您通過將浮點(diǎn)性能、算術(shù)強(qiáng)度和內(nèi)存帶寬組合到二維圖中來了解內(nèi)核特性。

通過查看 Roofline 模型，可以快速確定內(nèi)核是受計(jì)算限制還是內(nèi)存受限。您還可以了解進(jìn)一步優(yōu)化的潛在方向，例如，靠近屋頂線的內(nèi)核可以優(yōu)化利用計(jì)算資源。

圖 11 Nsight 計(jì)算中的屋頂線模型。

CUDA 11 包括 Compute Sanitizer ，這是下一代的功能正確性檢查工具，它為內(nèi)存訪問和競(jìng)爭(zhēng)條件提供運(yùn)行時(shí)檢查。 Compute Sanitizer 旨在替代 cuda-memcheck 工具。

下面的代碼示例顯示了計(jì)算消毒劑檢查內(nèi)存訪問的示例。

//Out-of-bounds Array Access __global__ void oobAccess(int* in, int* out)
{ int bid = blockIdx.x; int tid = threadIdx.x; if (bid == 4) { out[tid] = in[dMem[tid]]; }
} int main()
{ ... // Array of 8 elements, where element 4 causes the OOB std::array hMem = {0, 1, 2, 10, 4, 5, 6, 7}; cudaMemcpy(d_mem, hMem.data(), size, cudaMemcpyHostToDevice); oobAccess<<<10, Size>>>(d_in, d_out); cudaDeviceSynchronize(); ... $ /usr/local/cuda-11.0/Sanitizer/compute-sanitizer --destroy-on-device-error kernel --show-backtrace no basic
========= COMPUTE-SANITIZER
Device: Tesla T4
========= Invalid __global__ read of size 4 bytes
========= at 0x480 in /tmp/CUDA11.0/ComputeSanitizer/Tests/Memcheck/basic/basic.cu:40:oobAccess(int*,int*)
========= by thread (3,0,0) in block (4,0,0)
========= Address 0x7f551f200028 is out of bounds

下面的代碼示例顯示了用于競(jìng)爭(zhēng)條件檢查的計(jì)算消毒劑示例。

//Contrived Race Condition Example __global__ void Basic()
{ __shared__ volatile int i; i = threadIdx.x;
} int main()
{ Basic<<<1,2>>>(); cudaDeviceSynchronize(); ... $ /usr/local/cuda-11.0/Sanitizer/compute-sanitizer --destroy-on-device-error kernel --show-backtrace no --tool racecheck --racecheck-report hazard raceBasic
========= COMPUTE-SANITIZER
========= ERROR: Potential WAW hazard detected at __shared__ 0x0 in block (0,0,0) :
========= Write Thread (0,0,0) at 0x100 in /tmp/CUDA11.0/ComputeSanitizer/Tests/Racecheck/raceBasic/raceBasic.cu:11:Basic(void)
========= Write Thread (1,0,0) at 0x100 in /tmp/CUDA11.0/ComputeSanitizer/Tests/Racecheck/raceBasic/raceBasic.cu:11:Basic(void)
========= Current Value : 0, Incoming Value : 1
=========
========= RACECHECK SUMMARY: 1 hazard displayed (1 error, 0 warnings)

最后，盡管 CUDA 11 不再支持在 macOS 上運(yùn)行應(yīng)用程序，但我們正在為 macOS 主機(jī)上的用戶提供開發(fā)工具：

使用 cuda-gdb 進(jìn)行遠(yuǎn)程目標(biāo)調(diào)試

NVIDIA Visual Profiler

Nsight Eclipse 插件

用于遠(yuǎn)程分析或跟蹤的 Nsight 工具系列

摘要

CUDA 11 為在 NVIDIA A100 上構(gòu)建應(yīng)用程序提供了基礎(chǔ)開發(fā)環(huán)境，用于構(gòu)建 NVIDIA Ampere GPU 體系結(jié)構(gòu)和強(qiáng)大的服務(wù)器平臺(tái)，用于 AI 、數(shù)據(jù)分析和 HPC 工作負(fù)載，這些平臺(tái)均用于內(nèi)部（ DGX A100 型 ）和云（ HGX A100 型 ）部署。

圖 12 得到 CUDA 11 的不同方法。

CUDA 11 現(xiàn)在可用。和往常一樣，您可以通過多種方式獲得 CUDA 11 ： 下載 本地安裝程序包，使用包管理器進(jìn)行安裝，或者從各種注冊(cè)中心獲取容器。對(duì)于企業(yè)部署， CUDA 11 還包括 RHEL8 的 驅(qū)動(dòng)程序打包改進(jìn) ，使用模塊化流來提高穩(wěn)定性并減少安裝時(shí)間。

關(guān)于作者

Pramod Ramarao 是 NVIDIA 加速計(jì)算的產(chǎn)品經(jīng)理。他領(lǐng)導(dǎo) CUDA 平臺(tái)和數(shù)據(jù)中心軟件的產(chǎn)品管理，包括容器技術(shù)。

審核編輯：郭婷