NVIDIA 宣布 CUDA 開發(fā)環(huán)境的最新版本 CUDA 11 . 5 。 CUDA 11 . 5 專注于增強您的 CUDA 應(yīng)用程序的編程模型和性能。 CUDA 繼續(xù)推動 GPU 加速的邊界，并為 HPC 、可視化、 AI 、 ML 和 DL 中的新應(yīng)用打下基礎(chǔ)，和數(shù)據(jù)科學(xué)。

CUDA 11 . 5 有幾個重要特性。這篇文章概述了關(guān)鍵功能：

• CUDA 編程模型增強
  • 掃描合作小組中的集體
  • 標準化整數(shù)格式
  • 塊壓縮格式
  • C ++中可配置的緩存提示
• MPS 增強功能（客戶端內(nèi)存限制）
• WSL 上的 CUDA 驅(qū)動程序更新
• CUDA Python GA
• NVIDIA 開普勒驅(qū)動器的棄用
• CUDA C ++（有關(guān)更多信息，請參閱使用 CUDA C ++編譯輔助工具減少應(yīng)用程序構(gòu)建時間）
• NSight 計算/系統(tǒng)工具

CUDA 11 . 5 附帶 R510 驅(qū)動程序，該驅(qū)動程序是一個長期支援科. CUDA 11 . 5 可供下載。

CUDA 編程模型增強

此版本引入了關(guān)鍵的增強功能，以提高 CUDA 圖形的可用性和性能，而無需對應(yīng)用程序進行任何修改或任何其他用戶干預(yù)。它還提高了多進程服務(wù)（ MPS ）的易用性。我們在 CUDA 編程指南中對異步編程模型進行了形式化。

掃描合作小組中的集體

與reductions和障礙，前綴和（也稱為scans）一起，它們是并行計算的基石。掃描操作采用二進制運算符（通常為加法），并在輸入數(shù)組上累積應(yīng)用該運算符進行迭代。掃描可以是inclusive，包括所有元素x[0]…x[n]，或在范圍{0，x[0]…x[n-1]}上迭代的exclusive。

例如，使用輸入數(shù)組[3 1 7 0 4 1 6 3]的+運算符進行獨占掃描將導(dǎo)致以下結(jié)果：

          [0 3 4 11 11 15 16 22]

CUDA 11 . 5 添加了一個新的頭，它在協(xié)作組名稱空間中定義了四個新函數(shù)來實現(xiàn)這些操作。

在所有情況下，返回類型必須與輸入值類型匹配。

規(guī)范化整數(shù)數(shù)據(jù)類型

標準化有符號和無符號 8 位和 16 位數(shù)據(jù)類型是 GPU 編程語言最廣泛支持的一些紋理格式。 CUDA 一段時間以來一直支持將這些格式用于紋理對象，但在 11 . 5 版本中，我們擴展了對這些數(shù)據(jù)類型的現(xiàn)有支持，使與其他外部 API 的互操作更加直觀。

我們在驅(qū)動程序和運行時 API 中引入了新的 CUDA 數(shù)組格式。驅(qū)動程序 API 公開了 12 種新的數(shù)組格式，如下所示：

CU_AD_FORMAT_UNORM_INT8X{1|2|4} 
CU_AD_FORMAT_UNORM_INT16X{1|2|4}
CU_AD_FORMAT_SNORM_INT8X{1|2|4}
CU_AD_FORMAT_SNORM_INT16X{1|2|4}

這些可用于創(chuàng)建 1 、 2 或 4 通道 CUDA 陣列。運行時 API 同樣公開了 12 種新的等效通道格式：

cudaChannelFormatKindUnsignedNormalized8X{1|2|4}
cudaChannelFormatKindSignedNormalized8X{1|2|4}
cudaChannelFormatKindUnsignedNormalized16X{1|2|4}
cudaChannelFormatKindSignedNormalized16X{1|2|4}

這些還可用于創(chuàng)建 1 、 2 或 4 通道、 8 位或 16 位通道寬度 CUDA 陣列。此外，您現(xiàn)在可以從外部 API （如 DirectX12 / 11 或 Vulkan ）導(dǎo)入匹配的格式化紋理，并將其映射為 CUDA 數(shù)組。使用資源視圖創(chuàng)建紋理對象時，格式 texel 大小必須與數(shù)組 texel 大小匹配。

對于紋理對象，可以創(chuàng)建和訪問它們，如以下代碼示例所示：

cudaArray_t array;
cudaChannelFormatDesc formatDesc = {8, 8, 0, 0, cudaChannelFormatKindUnsignedNormalized8X2};
 
cudaMallocArray(&array, formatDesc, width, height);
 
cudaTextureDesc texDesc = {0};
texDesc.addressMode[0] = texDesc.addressMode[1] = cudaAddressModeClamp ;
 
// (3) Create CUDA texture object
cudaResourceDesc resDesc = {0};
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array = array;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

// Read from texture object in a kernel as follows:
// float4 texel = tex2D(texture, x, y);
 
// (6) Release all resources
cudaDestroyTextureObject(texObj); 
cudaFreeArray(array);

同樣，對于曲面對象：

cudaArray_t array;
cudaChannelFormatDesc formatDesc = {16, 16, 16, 16, cudaChannelFormatKindSignedNormalized16X4};
 
cudaMallocArray(&array, formatDesc, width, height);
 
// (3) Create CUDA surface object
cudaResourceDesc resDesc = {0};
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array = array;
cudaCreateSurfaceObject(&surfObj, &resDesc);
 
// Read/Write to/from surface object in a kernel as follows
// Read:
// short4 texel = surf2DRead(surface, xInBytes, y);
// Unformatted stores:
// surf2DWrite(texel, surface, xInBytes, y);
// Formatted stores: (Formatted surface stores are currently not exposed in CUDA runtime device functions)
// sust.p.2d.v4.b32
 
// (6) Release all resources
cudaDestroySurfaceObject(surfObj); 
cudaFreeMipmappedArray(array);

對塊壓縮數(shù)據(jù)類型的支持

在所有圖形編程語言和框架中，用于減小紋理大小的最常見有損壓縮技術(shù)之一是使用塊壓縮（ BC ）紋理格式。使用這些格式可以顯著節(jié)省紋理的內(nèi)存占用。有幾種 BC 格式，每種格式都有其獨特的優(yōu)點和缺點，通常稱為 BCn 格式。

NVIDIA GPU 體系結(jié)構(gòu)本機支持 BCn 格式，并且通過紋理資源視圖在 CUDA 中的支持有限?，F(xiàn)在，我們在驅(qū)動程序和運行時 API 中引入新的 BC CUDA 數(shù)組格式。

這些格式可用于使用cuArray[3D]Create運行時 API 或cuArray[3D]Create驅(qū)動程序 API 創(chuàng)建 BCn 格式的cuMipmappedArrayCreate數(shù)組。類似地，可以使用 CUDA 運行時 API 或cuMipmappedArrayCreate驅(qū)動程序 API 創(chuàng)建 CUDA mipmapped 數(shù)組。使用這些格式創(chuàng)建 CUDA 陣列時，陣列范圍必須是壓縮塊大小的倍數(shù)（ 2D 為 4×4 ， 3D 為 4x4x1 ）。這些陣列還可用于創(chuàng)建紋理對象。

C ++中可配置的緩存提示

在在 CUDA 11 . 4 中發(fā)現(xiàn)新功能中，我們引入了一個 PTX ISA 擴展，為駐留在 GPU 上的數(shù)據(jù)向編譯器和運行時提供緩存提示。在 CUDA 11 . 5 中，我們將此能力擴展到帶有注釋指針的 C ++。它們充當普通指針，并應(yīng)用附加屬性。

注釋指針是使用中定義的函數(shù)創(chuàng)建的，緩存駐留提示定義為以下內(nèi)容之一：

cuda::access_property::normal		(evict_normal_demote)
cuda::access_property::streaming        (evict_first)
cuda::access_property::persisting       (evict_last)
cuda::access_property::shared           (shared memory)
cuda::access_property::global           (evict_normal)

例如，在內(nèi)核代碼中，聲明和使用帶注釋的指針可能類似于以下代碼：

static __device__
void my_kernel(int * in, int * out) { 
cuda::access_property ap(cuda::access_property::persisting{}); 

// Retrieve global id 
int i = blockIdx.x * blockDim.x + threadIdx.x; 

cuda::annotated_ptr in_ann{in, ap}; 
cuda::annotated_ptr out_ann{out, ap};
 
...
}

MPS 增強功能（客戶端內(nèi)存限制）

當 GPU 的計算能力超過任何單個應(yīng)用程序時，運行共享同一 GPU 硬件的多個應(yīng)用程序進程可能會很有吸引力。多進程服務(wù)（ MPS ）運行時體系結(jié)構(gòu)控制多個獨立進程同時使用單個 GPU 。

但是，當多個獨立進程共享 GPU 時，設(shè)置總體內(nèi)存分配限制通常很有用，以避免任何單個進程占用過多的可用 GPU 內(nèi)存。

在 CUDA 11 . 5 中，我們引入了一組新的控制機制，使您能夠限制 MPS 客戶端進程的固定內(nèi)存分配。您可以通過默認的全局限制層次結(jié)構(gòu)控制內(nèi)存分配。

默認全局限制

可以使用設(shè)備的set_default_device_pinned_mem_limit控制命令顯式啟用默認全局內(nèi)存限制。設(shè)置此命令將在將來生成的所有 MPS 服務(wù)器的所有 MPS 客戶端上強制執(zhí)行設(shè)備固定內(nèi)存限制。

$nvidia-cuda-mps-control set_default_device_pinned_mem_limit 0 2G

每服務(wù)器限制：對于內(nèi)存資源限制的細粒度控制，可以使用set_device_pinned_mem_limit控制命令在特定 MPS 服務(wù)器上選擇性地設(shè)置限制。設(shè)置此命令將在特定 MPS 服務(wù)器的所有 MPS 客戶端上強制執(zhí)行設(shè)備固定內(nèi)存限制。

$nvidia-cuda-mps-control set_device_pinned_mem_limit  1 1G

每個客戶限額：前面兩種控制機制為特定 MPS 服務(wù)器的所有 MPS 客戶端設(shè)置了一個總體限制。希望更好地控制資源限制的用戶；也就是說，在每個 MPS 客戶端的基礎(chǔ)上，可以通過為每個客戶端進程分別設(shè)置CUDA_MPS_PINNED_DEVICE_MEM_LIMIT環(huán)境變量來實現(xiàn)。

此環(huán)境變量與 CUDA \ u VISIBLE \ u 設(shè)備具有相同的語義。值字符串可以包含逗號分隔的設(shè)備序號和設(shè)備 UUID ，每個設(shè)備的內(nèi)存限制由等號（=）分隔。

$export CUDA_MPS_PINNED_DEVICE_MEM_LIMIT="0=1G,1=2G,GPU-7ce23cd8-5c91-34a1-9e1b-28bd1419ce90=1024M"

WSL 驅(qū)動程序上的 CUDA

用于英特爾 x86 體系結(jié)構(gòu)的 NVIDIA Windows GPU 驅(qū)動程序?qū)⒅С?WSL2 ，并可在 Windows 11 的 Windows Insider Preview （ WIP ）程序之外訪問。

CUDA Python

Python CUDA 為現(xiàn)有工具包和庫的驅(qū)動程序和運行時 API 提供Cython綁定和 Python 包裝，以簡化基于 GPU 的加速處理。 Python 是科學(xué)、工程、數(shù)據(jù)分析和深度學(xué)習應(yīng)用程序中最流行的編程語言之一。 CUDA Python 的目標是將 Python 生態(tài)系統(tǒng)統(tǒng)一為一組接口，提供 Python 對 CUDA 主機 API 的全面覆蓋和訪問。

庫開發(fā)人員可以直接從 Python 使用 CUDA Python 與 CUDA 的低級接口。我們很高興地宣布，從 11 . 5 版開始， CUDA Python 已普遍提供，可以使用 PIP 或 Conda 安裝。 CUDA 支持的所有平臺都支持該庫。

其他增強功能

除了前面討論的關(guān)鍵功能外， CUDA 11 . 5 中還有一些增強功能，有助于改進基線功能和可用性。

多線程提交吞吐量

在 11 . 5 中，我們減少了 CPU 線程之間 CUDA API 的序列化開銷。默認情況下，將啟用這些更改。但是，為了幫助對潛在更改可能導(dǎo)致的問題進行分類，我們提供了一個環(huán)境變量CUDA_REDUCE_API_SERIALIZATION，以關(guān)閉這些更改。這是前面討論的基礎(chǔ)更改之一，有助于 CUDA 圖的性能改進。

NVIDIA 開普勒驅(qū)動器的棄用

NVIDIA 開普勒微體系結(jié)構(gòu)于 2012 年首次引入，并已逐步淘汰。對于所有基于開普勒 NVIDIA 的 SKU ，我們已經(jīng)從 R495 版本開始取消了對驅(qū)動程序的支持。但是， CUDA 工具包開發(fā)工具和對 NVIDIA 開普勒 SKU 的支持將在未來的 CUDA 11 . x 版本中繼續(xù)。

C ++語言對 CUDA 的支持

作為本版本的一部分，有一些 CUDA 11 . 5 支持的關(guān)鍵 C ++語言增強。

• CUDA C ++編譯器支持 NVRTC 和 PTX 中的并發(fā)編譯，以提高編譯時間。編譯器現(xiàn)在還可以檢測并消除未使用的 CUDA 內(nèi)核，通過更好的代碼優(yōu)化減少編譯時間、二進制大小和總體性能。
• 對 128 位整數(shù)值的有限支持作為用戶反饋的預(yù)覽功能發(fā)布，同時引入了 NVRTC 的靜態(tài)庫版本，并擴展了主機編譯器支持以包括 Clang 12 . 0 。
• CUDA C++編譯器具有我們在[VZX57 ]帖子中深入研究的特性。
• NVIDIA C ++標準庫（ LIbCu +++） 1 . 5 . 0 被 CUDA 11 . 4 發(fā)布。
• 推力 1 . 12 . 0 具有新的推力：：通用_矢量 API ，使您能夠使用推力的 CUDA 統(tǒng)一內(nèi)存。

NSight 開發(fā)人員工具

NVIDIA NSight 開發(fā)工具： NSight System 2021 . 4 、 NSight Compute 2021 . 3 和 NSight Graphics 2021 . 4 . 2 現(xiàn)在提供了新版本，用于通過分析和調(diào)試 CUDA 代碼提高性能。

新發(fā)布的?NSight 系統(tǒng) 2021 . 4改進了對 Windows 、 Direct3D12 和 Vulkan 支持的評測。此版本添加了一些功能，以幫助更好地理解操作系統(tǒng)中斷時的進程執(zhí)行，并添加了數(shù)據(jù)捕獲以識別數(shù)據(jù)包隊列瓶頸。特色包括：

• Windows ISR 和 DPC 跟蹤
• GPU 基于硬件的調(diào)度跟蹤
• Windows Direct3D12
• Vulkan 與 WDDM 事件的相關(guān)性
• NVTX 事件分類支持
• 多種系統(tǒng)環(huán)境的多報告加載

NSight Compute 2021 . 3 添加了一些功能，以幫助用戶了解其 CUDA 內(nèi)核的性能。新的占用率計算器活動對 CUDA 內(nèi)核的資源利用率進行建模，以便您可以交互式地調(diào)整模型參數(shù)，以查看它們?nèi)绾斡绊懻加寐省Ｎ蓓斦劬€圖現(xiàn)在支持分層脫機，它表示除設(shè)備內(nèi)存之外的內(nèi)存層次中的其他級別。您可以查看內(nèi)核是否存在與緩存訪問請求相關(guān)的瓶頸。

還有其他改進，包括更多可配置的基線比較、從 CLI 查看源代碼級信息以及其他 SSH 功能。

圖 1 。占用率計算器

最新的 NSight 圖形 2021 . 4 . 2 現(xiàn)在包括對 Windows 11 的支持。這意味著您現(xiàn)在可以使用downloadDirect3D 和 Vulkan 的 NVIDIA 圖形調(diào)試器和探查器，在最前沿的 Windows 版本上創(chuàng)建驚人的 3D 圖形。

關(guān)于作者

Rob Armstrong 是 CUDA 工具包的主要技術(shù)產(chǎn)品經(jīng)理。 20 多年來，他一直專注于使用異構(gòu)硬件平臺加速軟件，并對計算機體系結(jié)構(gòu)和硬件/軟件交互特別感興趣。

Arthy Sundaram 是 CUDA 平臺的技術(shù)產(chǎn)品經(jīng)理。她擁有哥倫比亞大學(xué)計算機科學(xué)碩士學(xué)位。她感興趣的領(lǐng)域是操作系統(tǒng)、編譯器和計算機體系結(jié)構(gòu)。

Fred Oh 是 CUDA 、 CUDA on WSL 和 CUDA Python 的高級產(chǎn)品營銷經(jīng)理。弗雷德?lián)碛屑又荽髮W(xué)戴維斯分校計算機科學(xué)和數(shù)學(xué)學(xué)士學(xué)位。他的職業(yè)生涯開始于一名 UNIX 軟件工程師，負責將內(nèi)核服務(wù)和設(shè)備驅(qū)動程序移植到 x86 體系結(jié)構(gòu)。他喜歡《星球大戰(zhàn)》、《星際迷航》和 NBA 勇士隊。

審核編輯：郭婷