NVIDIA CUDA Toolkit 13.1 正式发布：20年最大更新，CUDA Tile引领编程新范式

近日，NVIDIA CUDA Toolkit 13.1 已全面推出，英伟达官方强调：「这是过去20年间最为重大的一次革新。」

英伟达社交媒体渠道

自2006年CUDA平台问世以来，此次更新堪称规模最大、覆盖最广，主要内容包括：

NVIDIA CUDA Tile的亮相，这是英伟达推出的基于tile的编程模型，能够抽象化专用硬件细节，例如张量核心。

Runtime API exposure of green contexts（指将轻量级、可并发调度的绿色上下文暴露给外部调用者使用。）

NVIDIA cuBLAS中的双精度和单精度仿真功能。

一本彻底重构的CUDA编程指南，兼顾CUDA新手和资深开发者。

接下来，我们将详细解读这些更新。

CUDA Tile

CUDA Tile是NVIDIA CUDA Toolkit 13.1的核心革新。它作为一种基于tile的编程模型，允许开发者在更高层次编写算法，并屏蔽专用硬件（如张量核心）的底层细节。

英伟达社交媒体

解读CUDA Tile的核心理念

英伟达博客解释道：CUDA Tile让开发者能够在高于SIMT（单指令多线程）的层级编写GPU核函数。

在当前SIMT编程中，开发者通常通过分割数据并定义每个线程的执行路径来指定核函数。

而借助CUDA Tile，开发者可以提升代码抽象层级，直接操作称为“Tile”的数据块。只需指定在这些Tile上执行的数学运算，编译器和运行时会自动决定工作负载分配到线程的最佳方式。

这种Tile模型隐藏了调用Tensor Core等专用硬件的底层细节，并且Tile代码能兼容未来GPU架构。

CUDA 13.1包含两个Tile编程组件：

CUDA Tile IR：一种用于NVIDIA GPU编程的全新虚拟指令集架构（ISA）。

cuTile Python：一种新的领域特定语言（DSL），用于在Python中编写基于数组和Tile的核函数。

底层技术细节

编译的Tile路径可融入完整软件栈，与SIMT路径对应。

这是该软件的首个版本，注意事项如下：

CUDA Tile仅支持NVIDIA Blackwell（计算能力10.x和12.x）系列产品。未来CUDA版本将扩展更多架构支持。

当前开发重点聚焦AI算法的Tile编程。英伟达表示将在未来CUDA版本中持续增加特性、功能并提升性能。

英伟达计划在即将发布的CUDA版本中引入C++实现。

为何为GPU引入Tile编程？

CUDA向开发者提供了单指令多线程（SIMT）硬件和编程模型。这种模式要求（同时允许）开发者以最大灵活性和针对性，对代码执行进行细粒度控制。然而，编写高性能代码往往耗费大量精力，尤其在适配多种GPU架构时。

尽管已有许多库（如NVIDIA CUDA-X和NVIDIA CUTLASS）旨在帮助开发者挖掘性能，但CUDA Tile引入了一种比SIMT层级更高的新型GPU编程方式。

随着计算工作负载演进，特别是在AI领域，张量已成为基础数据类型。NVIDIA开发了专门处理张量的硬件，例如NVIDIA Tensor Core（TC）和NVIDIA Tensor Memory Accelerator（TMA），它们现已成为每个新GPU架构的核心部分。

硬件越复杂，越需要软件驾驭这些能力。CUDA Tile对Tensor Core及其编程模型进行了抽象，使得用CUDA Tile编写的代码能兼容当前及未来Tensor Core架构。

基于Tile的编程方式允许开发者通过指定数据块（即Tile），然后定义在这些Tile上执行的计算来编写算法。开发者无需在逐元素层面设定算法执行细节：编译器和运行时会处理这些工作。

下图展示了随CUDA Tile推出的Tile模型与CUDA SIMT模型之间的概念差异。

Tile模型与CUDA SIMT模型之间的概念差异

Tile模型（左）将数据划分为多个块，编译器将其映射到线程。单指令多线程（SIMT）模型（右）将数据同时映射到块和线程

这种编程范式在Python等语言中常见，例如NumPy库让开发者指定矩阵等数据类型，然后用简单代码执行批量操作。

CUDA软件更新

以下是本次CUDA版本更新中包含的其他重要软件改进：

运行时对Green Context（绿色上下文）的支持

CUDA中的Green Context是一种轻量级上下文形式，可作为传统CUDA上下文的替代方案，为开发者提供更细粒度的GPU空间划分与资源分配能力。

自CUDA 12.4起，它们已在驱动API中提供；而从本版本开始，Green Context也正式在运行时API中开放使用。

Green Context使用户能够定义和管理GPU资源的独立分区，主要是Streaming Multiprocessors（SM）。你可以将特定数量的SM分配给某个Green Context，然后在该context所拥有的资源范围内启动CUDA kernel并管理只在此context内运行的stream。

一个典型应用场景是：程序中有部分代码对延迟极为敏感，并且需要优先于其他所有GPU工作执行。通过为这段代码单独创建一个Green Context并分配SM资源，而将剩余SM分配给另一个Green Context处理其他任务，就能确保始终有可用SM供高优先级计算使用。

CUDA 13.1还引入了更加可定制的split() API。开发者可通过这一接口构建此前需要多次API调用才能完成的SM分区，并且可以配置工作队列，从而减少不同Green Context之间提交任务时产生的伪依赖（false dependencies）。

有关这些功能及Green Context的更多信息，请参见CUDA Programming Guide。

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

CUDA多进程服务（MPS）更新

CUDA 13.1为多进程服务带来了多项新特性和功能。有关这些新功能的完整信息，请参阅MPS文档。以下是部分亮点内容：

内存局部性优化分区

内存局部性优化分区（Memory locality optimization partition，MLOPart）是NVIDIA Blackwell系列（计算能力10.0和10.3，为架构版本号）及更新GPU上提供的一项特性。

该功能允许用户创建专门优化内存局部性的CUDA设备。MLOPart设备基于同一块物理GPU派生而来，但呈现为多个独立设备，每个设备拥有更少的计算资源和更小的可用内存。

在计算能力10.0和10.3的GPU上，每块GPU都包含两个分区。

当在GPU上启用MLOPart时，每个分区都会作为一个独立的CUDA设备出现，并具有其对应的计算与内存资源。

目前，MLOPart仅支持NVIDIA B200与NVIDIA B300系列产品。未来的CUDA发布版本将加入对NVIDIA GB200与NVIDIA GB300系列的支持。

静态流式多处理器（SM）分区

作为MPS中现有的动态执行资源供给（provisioning）的一种替代方案，静态流式多处理器（SM）分区是针对NVIDIA Ampere架构（计算能力8.0）及更新GPU的一项特性，它为MPS客户端提供了一种创建独占SM分区的方法。

该模式通过使用-S或--static-partitioning标志启动MPS控制守护进程来启用，其主要目的是提供确定性的资源分配，并改善MPS客户端之间的隔离性。分区的基本单位是一个“Chunk”（块），其大小根据GPU架构而异——例如，在Hopper（计算能力9.0）及更新的独立GPU上，一个Chunk包含8个SM。

cuBLAS中的双精度和单精度模拟

虽然严格来说这不属于CUDA 13.1的更新，但NVIDIA CUDA Toolkit 13.0中的cuBLAS更新引入了新的API和实现，旨在提升双精度（FP64）矩阵乘法（matmul）的性能。

这是通过在NVIDIA GB200 NVL72和NVIDIA RTX PRO 6000 Blackwell Server Edition等GPU架构的Tensor Core上进行浮点（FP）模拟来实现的。

开发者工具

开发者工具是CUDA平台的重要组成部分。此次发布带来了多项创新和功能增强，包括：

CUDA Tile核函数性能分析工具

在摘要页新增“Result Type”（结果类型）列，用于区分Tile核函数与SIMT核函数。

详情页新增“Tile Statistics”（Tile统计）部分，总结Tile维度和重要管线（pipeline）的利用率。

源码页支持将指标映射到高层级的cuTile核函数源码。

源码页

Nsight Compute分析，重点展示了分析输出中的Tile Statistics部分

此次发布的Nsight Compute还增加了对设备端启动的图（device-launched graphs）中CUDA图节点的分析支持，并改进了源码页导航，为编译器生成和用户生成的标签提供了可点击的链接。

编译时修补

NVIDIA Compute Sanitizer 2025.4通过-fdevice-sanitize=memcheck编译器标志，增加了对NVIDIA CUDA编译器（NVCC）编译时修补（patching）的支持。这种修补增强了内存错误检测能力，并提升了Compute Sanitizer的性能。

编译时插桩（instrumentation）可将错误检测直接集成到NVCC中，从而实现更快的运行速度，并通过高级的基址-边界分析（base-and-bounds analysis）捕捉更隐蔽的内存问题（如相邻分配间的非法访问）。这意味着开发者可以在不牺牲速度的情况下调试内存问题，运行更多测试并保持生产力。目前，该功能仅支持memcheck工具。

要使用此新功能，请使用如下NVCC标志编译代码：

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

然后使用memcheck工具运行你的应用：

compute-sanitizer --tool memcheck myapp

NVIDIA Nsight Systems

NVIDIA Nsight Systems 2025.6.1与CUDA Toolkit 13.1同步发布，带来了多项新的追踪功能：

系统级CUDA追踪：--cuda-trace-scope可开启跨进程树或整个系统的追踪。

CUDA主机函数追踪：增加了对CUDA Graph主机函数节点和cudaLaunchHostFunc()的追踪支持，这些函数在主机上执行并会阻塞流（stream）。

CUDA硬件追踪：在支持的情况下，基于硬件的追踪现在成为默认模式；使用--trace=cuda-sw可恢复为软件模式。

Green Context时间轴行现在会在工具提示中显示SM分配情况，帮助用户理解GPU资源利用率。

数学库

核心CUDA工具包数学库的新功能包括：

NVIDIA cuBLAS：一项全新的实验性API，支持Blackwell GPU的分组GEMM功能，并兼容FP8和BF16/FP16数据类型。针对上述数据类型，支持CUDA图的分组GEMM提供了一种无需主机同步的实现方式，其设备端形状可实现最高4倍的加速，优于MoE用例中的多流GEMM实现。

NVIDIA cuSPARSE：一种新的稀疏矩阵向量乘法(SpMVOp)API，与CsrMV API相比性能有所提升。该API支持CSR格式、32位索引、双精度以及用户自定义的后缀。

NVIDIA cuFFT：一套名为cuFFT设备API的全新API，提供主机函数，用于在C++头文件中查询或生成设备功能代码和数据库元数据。该API专为cuFFTDx库设计，可通过查询cuFFT来生成cuFFTDx代码块，这些代码块可以与cuFFTDx应用程序链接，从而提升性能。

针对新的Blackwell架构，现已推出性能更新。用户可选择关键API进行更新，并查看性能更新详情。

cuBLAS Blackwell性能

CUDA Toolkit 12.9在NVIDIA Blackwell平台上引入了块缩放的FP4和FP8矩阵乘法。CUDA 13.1增加了对这些数据类型和BF16的性能支持。图2显示了在NVIDIA Blackwell和Hopper平台上的加速比。

在NVIDIA Blackwell和Hopper平台上的加速比

cuSOLVER Blackwell性能

CUDA 13.1继续优化用于特征分解的批处理SYEVD与GEEV API，并带来了显著的性能增强。

其中，批处理SYEV（cusolverDnXsyevBatched）是cuSOLVER中SYEV例程的统一批处理版本，用于计算对称/Hermitian矩阵的特征值与特征向量，非常适合对大量小矩阵进行并行求解的场景。

图3展示了在批大小为5,000（矩阵行数24–256）的测试结果。与NVIDIA L40S相比，NVIDIA Blackwell RTX Pro 6000 Server Edition实现了约2倍的加速，这与预期的内存带宽提升相吻合。

在批大小为5000（矩阵行数24–256）的测试结果

对于复数单精度和实数单精度两类矩阵，当行数N=5时，加速比约为1.5×，并随着行数增大逐渐提升，在N=250时达到2.0×。

图4显示了cusolverDnXgeev(GEEV)的性能加速比，该函数用于计算一般（非对称）稠密矩阵的特征值和特征向量。GEEV是一种混合CPU/GPU算法。单个CPU线程负责在QR算法中执行高效的早期降阶处理，而GPU则处理其余部分。图中显示了矩阵大小从1,024到32,768的相对性能加速比。

cusolverDnXgeev(GEEV)的性能加速比

当矩阵行数n=5000时，加速比约为1.0，并随着矩阵规模增大逐渐提升，在n=30000时达到约1.7。

NVIDIA CUDA核心计算库

NVIDIA CUDA Core计算库(CCCL)为CUB带来了多项创新和增强功能。

确定性浮点运算简化

由于浮点加法不具备结合律，cub::DeviceReduce历史上只能保证在同一GPU上每次运行得到位上完全相同的结果。这被实现为一个两遍算法。

作为CUDA 13.1的一部分，NVIDIA CCCL 3.1提供了两个额外的浮点确定性选项，您可以根据这些选项在确定性和性能之间进行权衡。

不保证：使用原子操作进行单次归约。这不能保证提供位上完全相同的结果。

GPU间：基于Kate Clark在NVIDIA GTC 2024大会上演讲中可复现的降维结果。结果始终逐位相同。

可以通过标志位设置确定性选项，如下面的代码所示。

演示代码

数据对比

更便捷的单相CUB API

几乎所有CUB算法都需要临时存储空间作为中间暂存空间。过去，用户必须通过两阶段调用模式来查询和分配必要的临时存储空间，如果两次调用之间传递的参数不一致，这种模式既繁琐又容易出错。

CCCL 3.1为一些接受内存资源的CUB算法添加了新的重载，从而用户可以跳过临时存储查询/分配/释放模式。

演示代码

CUDA Tile资源链接：https://developer.nvidia.com/cuda/tile

CUDA 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

© THE END