機(jī)器之心報(bào)道

機(jī)器之心編輯部

幾個(gè)小時(shí)前，NVIDIA CUDA Toolkit 13.1 正式發(fā)布，英偉達(dá)官方表示：「這是 20 年來(lái)最大的一次更新。」

這個(gè)自 2006 年 CUDA 平臺(tái)誕生以來(lái)規(guī)模最大、最全面的更新包括：

NVIDIA CUDA Tile 的發(fā)布，這是英偉達(dá)基于 tile 的編程模型，可用于抽象化專用硬件，包括張量核心。Runtime API exposure of green contexts（是指把所謂的 Green Context「指輕量級(jí)的、可并發(fā)調(diào)度的上下文或執(zhí)行環(huán)境」暴露給外部調(diào)用者使用。）NVIDIA cuBLAS 中的雙精度和單精度仿真。一本完全重寫的 CUDA 編程指南 ，專為 CUDA 新手和高級(jí)程序員設(shè)計(jì)。

下面我們就來(lái)具體看看。

CUDA Tile

CUDA Tile 是 NVIDIA CUDA Toolkit 13.1 最核心的更新。它是一種基于 tile 的編程模型，能夠以更高的層次編寫算法，并抽象化專用硬件（例如張量核心）的細(xì)節(jié)。

解讀 CUDA Tile 的核心概念

英偉達(dá)博客解釋說(shuō)：CUDA Tile 可讓開發(fā)者在高于 SIMT（單指令多線程）的層級(jí)編寫 GPU 核函數(shù)。

在目前的 SIMT 編程中，開發(fā)者通常通過劃分?jǐn)?shù)據(jù)并定義每個(gè)線程的執(zhí)行路徑來(lái)指定核函數(shù)。

而借助 CUDA Tile，開發(fā)者可以提升代碼的抽象層級(jí)，直接指定被稱為「Tile」的數(shù)據(jù)塊。只需指定要在這些 Tile 上執(zhí)行的數(shù)學(xué)運(yùn)算，編譯器和運(yùn)行時(shí)環(huán)境會(huì)自動(dòng)決定將工作負(fù)載分發(fā)到各個(gè)線程的最佳方式。

這種 Tile 模型屏蔽了調(diào)用 Tensor Core 等專用硬件的底層細(xì)節(jié)，并且 Tile 代碼將能夠兼容未來(lái)的 GPU 架構(gòu)。

CUDA 13.1 包含兩個(gè)用于 Tile 編程的組件：

CUDA Tile IR：一種用于 NVIDIA GPU 編程的全新虛擬指令集架構(gòu)（ISA）。cuTile Python：一種新的領(lǐng)域特定語(yǔ)言（DSL），用于在 Python 中編寫基于數(shù)組和 Tile 的核函數(shù)。

編譯的 Tile 路徑可以融入完整的軟件棧，與 SIMT 路徑對(duì)應(yīng)。

這是該軟件的首個(gè)版本，其包含以下注意事項(xiàng)：

CUDA Tile 僅支持 NVIDIA Blackwell（計(jì)算能力 10.x 和 12.x）系列產(chǎn)品。未來(lái)的 CUDA 版本將擴(kuò)展對(duì)更多架構(gòu)的支持。目前的開發(fā)重點(diǎn)聚焦于 AI 算法的 Tile 編程。英偉達(dá)表示在未來(lái)的 CUDA 版本中將持續(xù)增加更多特性、功能并提升性能。英偉達(dá)計(jì)劃在即將發(fā)布的 CUDA 版本中引入 C++ 實(shí)現(xiàn)。

為什么要為 GPU 引入 Tile 編程？

CUDA 向開發(fā)者提供了單指令多線程（SIMT）硬件和編程模型。這種模式要求（同時(shí)也允許）開發(fā)者以最大的靈活性和針對(duì)性，對(duì)代碼的執(zhí)行方式進(jìn)行細(xì)粒度控制。然而，編寫高性能代碼往往需要付出巨大的心力，尤其是在需要適配多種 GPU 架構(gòu)的情況下。

盡管已有許多庫(kù)（如 NVIDIA CUDA-X 和 NVIDIA CUTLASS）旨在幫助開發(fā)者挖掘性能，但CUDA Tile 引入了一種比 SIMT 層級(jí)更高的新型 GPU 編程方式。

隨著計(jì)算工作負(fù)載的演進(jìn)，特別是在 AI 領(lǐng)域，張量已成為一種基礎(chǔ)數(shù)據(jù)類型。NVIDIA 開發(fā)了專門用于處理張量的硬件，例如 NVIDIA Tensor Core（TC）和 NVIDIA Tensor Memory Accelerator（TMA），它們現(xiàn)已成為每個(gè)新 GPU 架構(gòu)中不可或缺的組成部分。

硬件越復(fù)雜，就越需要軟件來(lái)幫助駕馭這些能力。CUDA Tile 對(duì) Tensor Core 及其編程模型進(jìn)行了抽象，使得使用 CUDA Tile 編寫的代碼能夠兼容當(dāng)前及未來(lái)的 Tensor Core 架構(gòu)。

基于 Tile 的編程方式允許開發(fā)者通過指定數(shù)據(jù)塊（即 Tile），然后定義在這些 Tile 上執(zhí)行的計(jì)算來(lái)編寫算法。開發(fā)者無(wú)需在逐元素的層面上設(shè)定算法的執(zhí)行細(xì)節(jié)：編譯器和運(yùn)行時(shí)將處理這些工作。

下圖展示了隨 CUDA Tile 推出的 Tile 模型與 CUDA SIMT 模型之間的概念差異。

Tile 模型（左）將數(shù)據(jù)劃分為多個(gè)塊，編譯器將其映射到線程。單指令多線程（SIMT）模型（右）將數(shù)據(jù)同時(shí)映射到塊和線程

這種編程范式在 Python 等語(yǔ)言中很常見，在這些語(yǔ)言中，像 NumPy 這樣的庫(kù)可以讓開發(fā)者指定矩陣等數(shù)據(jù)類型，然后用簡(jiǎn)單的代碼指定并執(zhí)行批量操作。

CUDA 軟件更新

以下是本次 CUDA 版本更新中包含的其他重要軟件改進(jìn)：

運(yùn)行時(shí)對(duì) Green Context（綠色上下文）的支持

CUDA 中的 Green Context 是一種輕量級(jí)的上下文形式，可作為傳統(tǒng) CUDA 上下文的替代方案，為開發(fā)者提供更細(xì)粒度的 GPU 空間劃分與資源分配能力。

自 CUDA 12.4 起，它們已在驅(qū)動(dòng) API 中提供；而從本版本開始，Green Context 也正式在運(yùn)行時(shí) API 中開放使用。

Green Context 使用戶能夠定義和管理 GPU 資源的獨(dú)立分區(qū)，主要是 Streaming Multiprocessors（SM）。你可以將特定數(shù)量的 SM 分配給某個(gè)特定的 Green Context ，然后在該 context 所擁有的資源范圍內(nèi)啟動(dòng) CUDA kernel 并管理只在此 context 內(nèi)運(yùn)行的 stream。

一個(gè)典型的應(yīng)用場(chǎng)景是：你的程序中有部分代碼對(duì)延遲極為敏感，并且需要優(yōu)先于其他所有 GPU 工作執(zhí)行。通過為這段代碼單獨(dú)創(chuàng)建一個(gè) Green Context 并分配 SM 資源，而將剩余的 SM 分配給另一個(gè) Green Context 處理其他任務(wù)，你就能確保始終有可用的 SM 供高優(yōu)先級(jí)計(jì)算使用。

CUDA 13.1 還引入了更加可定制的 split () API。開發(fā)者可以通過這一接口構(gòu)建此前需要多次 API 調(diào)用才能完成的 SM 分區(qū)，并且可以配置工作隊(duì)列，從而減少不同 Green Context 之間提交任務(wù)時(shí)產(chǎn)生的偽依賴（false dependencies）。

有關(guān)這些功能及 Green Context 的更多信息，請(qǐng)參見 CUDA Programming Guide。

CUDA 編程指南地址：https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/green-contexts.html

CUDA 多進(jìn)程服務(wù)（MPS）更新

CUDA 13.1 為多進(jìn)程服務(wù)帶來(lái)了多項(xiàng)新特性和功能。有關(guān)這些新功能的完整信息，請(qǐng)參閱 MPS 文檔。以下是部分亮點(diǎn)內(nèi)容：

內(nèi)存局部性優(yōu)化分區(qū)

內(nèi)存局部性優(yōu)化分區(qū)（Memory locality optimization partition，MLOPart）是 NVIDIA Blackwell 系列（計(jì)算能力 10.0 和 10.3，為架構(gòu)版本號(hào)）及更新 GPU 上提供的一項(xiàng)特性。

該功能允許用戶創(chuàng)建專門優(yōu)化內(nèi)存局部性的 CUDA 設(shè)備。MLOPart 設(shè)備基于同一塊物理 GPU 派生而來(lái)，但呈現(xiàn)為多個(gè)獨(dú)立設(shè)備，每個(gè)設(shè)備擁有更少的計(jì)算資源和更小的可用內(nèi)存。

在計(jì)算能力 10.0 和 10.3 的 GPU 上，每塊 GPU 都包含兩個(gè)分區(qū)。

當(dāng)在 GPU 上啟用 MLOPart 時(shí)，每個(gè)分區(qū)都會(huì)作為一個(gè)獨(dú)立的 CUDA 設(shè)備出現(xiàn)，并具有其對(duì)應(yīng)的計(jì)算與內(nèi)存資源。

目前，MLOPart 僅支持 NVIDIA B200 與 NVIDIA B300 系列產(chǎn)品。未來(lái)的 CUDA 發(fā)布版本將加入對(duì) NVIDIA GB200 與 NVIDIA GB300 系列的支持。

靜態(tài)流式多處理器（SM）分區(qū)

作為 MPS 中現(xiàn)有的動(dòng)態(tài)執(zhí)行資源供給（provisioning）的一種替代方案，靜態(tài)流式多處理器（SM）分區(qū)是針對(duì) NVIDIA Ampere 架構(gòu)（計(jì)算能力 8.0）及更新 GPU 的一項(xiàng)特性，它為 MPS 客戶端提供了一種創(chuàng)建獨(dú)占 SM 分區(qū)的方法。

該模式通過使用 -S 或 --static-partitioning 標(biāo)志啟動(dòng) MPS 控制守護(hù)進(jìn)程來(lái)啟用，其主要目的是提供確定性的資源分配，并改善 MPS 客戶端之間的隔離性。分區(qū)的基本單位是一個(gè)「Chunk」（塊），其大小根據(jù) GPU 架構(gòu)而異 —— 例如，在 Hopper（計(jì)算能力 9.0）及更新的獨(dú)立 GPU 上，一個(gè) Chunk 包含 8 個(gè) SM。

cuBLAS 中的雙精度和單精度模擬

雖然嚴(yán)格來(lái)說(shuō)這不屬于 CUDA 13.1 的更新，但 NVIDIA CUDA Toolkit 13.0 中的 cuBLAS 更新引入了新的 API 和實(shí)現(xiàn)，旨在提升雙精度（FP64）矩陣乘法（matmul）的性能。

這是通過在 NVIDIA GB200 NVL72 和 NVIDIA RTX PRO 6000 Blackwell Server Edition 等 GPU 架構(gòu)的 Tensor Core 上進(jìn)行浮點(diǎn)（FP）模擬來(lái)實(shí)現(xiàn)的。

開發(fā)者工具

開發(fā)者工具是 CUDA 平臺(tái)的重要組成部分。此次發(fā)布帶來(lái)了多項(xiàng)創(chuàng)新和功能增強(qiáng)，包括：

CUDA Tile 核函數(shù)性能分析工具

在摘要頁(yè)新增「Result Type」（結(jié)果類型）列，用于區(qū)分 Tile 核函數(shù)與 SIMT 核函數(shù)。詳情頁(yè)新增「Tile Statistics」（Tile 統(tǒng)計(jì)）部分，總結(jié) Tile 維度和重要管線（pipeline）的利用率。源碼頁(yè)支持將指標(biāo)映射到高層級(jí)的 cuTile 核函數(shù)源碼。

Nsight Compute 分析，重點(diǎn)展示了分析輸出中的 Tile Statistics 部分

此次發(fā)布的 Nsight Compute 還增加了對(duì)設(shè)備端啟動(dòng)的圖（device-launched graphs）中 CUDA 圖節(jié)點(diǎn)的分析支持，并改進(jìn)了源碼頁(yè)導(dǎo)航，為編譯器生成和用戶生成的標(biāo)簽提供了可點(diǎn)擊的鏈接。

編譯時(shí)修補(bǔ)

NVIDIA Compute Sanitizer 2025.4 通過 -fdevice-sanitize=memcheck 編譯器標(biāo)志，增加了對(duì) NVIDIA CUDA 編譯器（NVCC）編譯時(shí)修補(bǔ)（patching）的支持。這種修補(bǔ)增強(qiáng)了內(nèi)存錯(cuò)誤檢測(cè)能力，并提升了 Compute Sanitizer 的性能。

編譯時(shí)插樁（instrumentation）可將錯(cuò)誤檢測(cè)直接集成到 NVCC 中，從而實(shí)現(xiàn)更快的運(yùn)行速度，并通過高級(jí)的基址 - 邊界分析（base-and-bounds analysis）捕捉更隱蔽的內(nèi)存問題（如相鄰分配間的非法訪問）。這意味著開發(fā)者可以在不犧牲速度的情況下調(diào)試內(nèi)存問題，運(yùn)行更多測(cè)試并保持生產(chǎn)力。目前，該功能僅支持 memcheck 工具。

要使用此新功能，請(qǐng)使用如下 NVCC 標(biāo)志編譯代碼：

nvcc -fdevice-sanitize=memcheck -o myapp myapp.cu

然后使用 memcheck 工具運(yùn)行你的應(yīng)用：

compute-sanitizer --tool memcheck myapp

NVIDIA Nsight Systems

NVIDIA Nsight Systems 2025.6.1 與 CUDA Toolkit 13.1 同步發(fā)布，帶來(lái)了多項(xiàng)新的追蹤功能：

系統(tǒng)級(jí) CUDA 追蹤：--cuda-trace-scope 可開啟跨進(jìn)程樹或整個(gè)系統(tǒng)的追蹤。CUDA 主機(jī)函數(shù)追蹤：增加了對(duì) CUDA Graph 主機(jī)函數(shù)節(jié)點(diǎn)和 cudaLaunchHostFunc () 的追蹤支持，這些函數(shù)在主機(jī)上執(zhí)行并會(huì)阻塞流（stream）。CUDA 硬件追蹤：在支持的情況下，基于硬件的追蹤現(xiàn)在成為默認(rèn)模式；使用 --trace=cuda-sw 可恢復(fù)為軟件模式。Green Context 時(shí)間軸行現(xiàn)在會(huì)在工具提示中顯示 SM 分配情況，幫助用戶理解 GPU 資源利用率。

數(shù)學(xué)庫(kù)

核心 CUDA 工具包數(shù)學(xué)庫(kù)的新功能包括：

NVIDIA cuBLAS：一項(xiàng)全新的實(shí)驗(yàn)性 API，支持 Blackwell GPU 的分組 GEMM 功能，并兼容 FP8 和 BF16/FP16 數(shù)據(jù)類型。針對(duì)上述數(shù)據(jù)類型，支持 CUDA 圖的分組 GEMM 提供了一種無(wú)需主機(jī)同步的實(shí)現(xiàn)方式，其設(shè)備端形狀可實(shí)現(xiàn)最高 4 倍的加速，優(yōu)于 MoE 用例中的多流 GEMM 實(shí)現(xiàn)。NVIDIA cuSPARSE：一種新的稀疏矩陣向量乘法 (SpMVOp) API，與 CsrMV API 相比性能有所提升。該 API 支持 CSR 格式、32 位索引、雙精度以及用戶自定義的后綴。NVIDIA cuFFT：一套名為 cuFFT 設(shè)備 API 的全新 API，提供主機(jī)函數(shù)，用于在 C++ 頭文件中查詢或生成設(shè)備功能代碼和數(shù)據(jù)庫(kù)元數(shù)據(jù)。該 API 專為 cuFFTDx 庫(kù)設(shè)計(jì)，可通過查詢 cuFFT 來(lái)生成 cuFFTDx 代碼塊，這些代碼塊可以與 cuFFTDx 應(yīng)用程序鏈接，從而提升性能。

針對(duì)新的 Blackwell 架構(gòu)，現(xiàn)已推出性能更新。用戶可選擇關(guān)鍵 API 進(jìn)行更新，并查看性能更新詳情。

cuBLAS Blackwell 性能

CUDA Toolkit 12.9 在 NVIDIA Blackwell 平臺(tái)上引入了塊縮放的 FP4 和 FP8 矩陣乘法。CUDA 13.1 增加了對(duì)這些數(shù)據(jù)類型和 BF16 的性能支持。圖 2 顯示了在 NVIDIA Blackwell 和 Hopper 平臺(tái)上的加速比。

cuSOLVER Blackwell 性能

CUDA 13.1 繼續(xù)優(yōu)化用于特征分解的批處理 SYEVD 與 GEEV API，并帶來(lái)了顯著的性能增強(qiáng)。

其中，批處理 SYEV（cusolverDnXsyevBatched） 是 cuSOLVER 中 SYEV 例程的統(tǒng)一批處理版本，用于計(jì)算對(duì)稱／Hermitian 矩陣的特征值與特征向量，非常適合對(duì)大量小矩陣進(jìn)行并行求解的場(chǎng)景。

圖 3 展示了在批大小為 5,000（矩陣行數(shù) 24–256）的測(cè)試結(jié)果。與 NVIDIA L40S 相比，NVIDIA Blackwell RTX Pro 6000 Server Edition 實(shí)現(xiàn)了約 2 倍的加速，這與預(yù)期的內(nèi)存帶寬提升相吻合。

對(duì)于復(fù)數(shù)單精度和實(shí)數(shù)單精度兩類矩陣，當(dāng)行數(shù)N = 5時(shí)，加速比約為1.5×，并隨著行數(shù)增大逐漸提升，在N = 250 時(shí)達(dá)到 2.0×。

圖 4 顯示了 cusolverDnXgeev (GEEV) 的性能加速比，該函數(shù)用于計(jì)算一般（非對(duì)稱）稠密矩陣的特征值和特征向量。GEEV 是一種混合 CPU/GPU 算法。單個(gè) CPU 線程負(fù)責(zé)在 QR 算法中執(zhí)行高效的早期降階處理，而 GPU 則處理其余部分。圖中顯示了矩陣大小從 1,024 到 32,768 的相對(duì)性能加速比。

當(dāng)矩陣行數(shù)n = 5000時(shí)，加速比約為1.0，并隨著矩陣規(guī)模增大逐漸提升，在n = 30000 時(shí)達(dá)到約 1.7。

NVIDIA CUDA 核心計(jì)算庫(kù)

NVIDIA CUDA Core 計(jì)算庫(kù) (CCCL) 為 CUB 帶來(lái)了多項(xiàng)創(chuàng)新和增強(qiáng)功能。

確定性浮點(diǎn)運(yùn)算簡(jiǎn)化

由于浮點(diǎn)加法不具備結(jié)合律，cub::DeviceReduce 歷史上只能保證在同一 GPU 上每次運(yùn)行得到位上完全相同的結(jié)果。這被實(shí)現(xiàn)為一個(gè)兩遍算法。

作為 CUDA 13.1 的一部分， NVIDIA CCCL 3.1 提供了兩個(gè)額外的浮點(diǎn)確定性選項(xiàng)，您可以根據(jù)這些選項(xiàng)在確定性和性能之間進(jìn)行權(quán)衡。

不保證：使用原子操作進(jìn)行單次歸約。這不能保證提供位上完全相同的結(jié)果。GPU 間：基于 Kate Clark 在 NVIDIA GTC 2024 大會(huì)上演講中可復(fù)現(xiàn)的降維結(jié)果。結(jié)果始終逐位相同。

可以通過標(biāo)志位設(shè)置確定性選項(xiàng)，如下面的代碼所示。

更便捷的單相 CUB API

幾乎所有 CUB 算法都需要臨時(shí)存儲(chǔ)空間作為中間暫存空間。過去，用戶必須通過兩階段調(diào)用模式來(lái)查詢和分配必要的臨時(shí)存儲(chǔ)空間，如果兩次調(diào)用之間傳遞的參數(shù)不一致，這種模式既繁瑣又容易出錯(cuò)。

CCCL 3.1 為一些接受內(nèi)存資源的 CUB 算法添加了新的重載，從而用戶可以跳過臨時(shí)存儲(chǔ)查詢 / 分配 / 釋放模式。

CUDA Tile 資源鏈接：https://developer.nvidia.com/cuda/tileCUDA Toolkit 13.1 下載地址：https://developer.nvidia.com/cuda-downloads

https://developer.nvidia.com/blog/focus-on-your-algorithm-nvidia-cuda-tile-handles-the-hardware

https://developer.nvidia.com/blog/nvidia-cuda-13-1-powers-next-gen-gpu-programming-with-nvidia-cuda-tile-and-performance-gains

https://x.com/NVIDIAAIDev/status/1996976702732620271

https://developer.nvidia.com/blog/simplify-gpu-programming-with-nvidia-cuda-tile-in-python