基于 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 遙測、控制和改進(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)勢，從而獲得更高的帶寬和性能。

在 A100 上， CUDA 11 提供了 API 操作，以留出 40 MB 二級(jí)緩存的一部分來持久化對(duì)全局內(nèi)存的數(shù)據(jù)訪問。持久訪問優(yōu)先使用二級(jí)緩存的這一預(yù)留部分，而全局內(nèi)存的正常或流式訪問只能在持久訪問未使用的情況下使用這部分二級(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ù)制的簡單示例。生成的代碼雖然性能更好，但可以通過對(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ù)木龋铀佻F(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)測的 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)存訪問和競爭條件提供運(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ì)算消毒劑示例。

//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ù)。

審核編輯：郭婷