随着人工智能和加速计算持续拓展技术边界，NVIDIA最新发布的CUDA版本带来了前所未有的性能与规模提升。
NVIDIA CUDA 13.5在多个重要新平台上实现了正式发布（GA）。NVIDIA Vera Rubin支持从早期开发者访问阶段升级为正式发布，已可用于企业生产环境。自CUDA 13.4起正式发布的Windows on Arm支持，如今进一步扩展至NVIDIA DGX Station for Windows。
CUDA 13.5还通过局部性域（locality domains）和NVIDIA多进程服务（MPS）引入了更多GPU任务调度的定制化选项。此外，CUDA 13.5在库、编译器、NVIDIA CUDA Tile以及开发者工具方面也带来了诸多创新。
本文将展示CUDA 13.5的主要亮点。
CUDA 13.5的增强功能
CUDA 13.5带来的一些值得关注的增强功能包括：
局部性域增强
CUDA 13.5为应用程序提供了选择权：可以在各局部性域之间使用统一的流式多处理器（SM）分区，也可以使用包含所有可用SM的分区。默认情况下，本地化的绿色上下文分区在每个域中使用相同的保证SM数量，额外的SM则不在这些分区之内。
新增的cudaDevSmResourceGroupLocalityDomainSupplementalSms标志与cudaDevSmResourceGroupLocalityDomainId结合使用，可将这些额外的SM纳入每个选定的域中。这使得应用程序在跨局部性域进行分区时能够使用设备的所有SM，尽管最终分区的大小可能会有所不同。
应用程序可以使用cudaDevSmResourceSplit() API中的发现模式来获取每个分区的SM数量并相应分配工作。这样一来，开发者既可以选择每个局部性域统一且可预测的SM规模，也可以选择让每个局部性域获得最大数量的SM，尽可能充分利用GPU硬件。更多信息请参阅CUDA运行时API文档。
NVIDIA MPS
MPS v3让开发者对工作负载如何共享GPU拥有更多控制权，涵盖从单个服务器的调度到客户端组的资源分配等多个方面。
加权时间片
加权时间片功能允许为承载需要更长运行时间工作负载的服务器分配更长的计算时间片。开发者可以设置基础时间片，按MPS服务器进行调整，并独立调节GPU在工作负载间切换的积极程度。
SM分区负载均衡
SM分区负载均衡功能让工作负载在资源可用时能够使用超出其基础分区的计算能力。通过--max-chunks启用借用功能，并使用--shareable=true允许某个分区在没有连接客户端时出借其SM资源。这些设置相互独立，因此一个分区可以同时具备借用和出借的能力。MPS会在重新分配资源之前先回收被借用的资源，而应用程序在此过程中始终能看到稳定的基础分区大小。
基于Cgroup的路由
基于Cgroup的路由功能将工作负载启动器的进程分组与MPS资源策略连接起来。开发者可将cgroup v2子树映射到某个MPS命名空间，通过守护进程主套接字连接的客户端会自动路由到该命名空间。
SM分区分配
按命名空间进行的SM分区分配功能让开发者只需为每个设备选择一次分区，该选择便会应用于路由到该命名空间的所有客户端，不再需要为这些客户端单独配置CUDA_MPS_SM_PARTITION。命名空间分配具有优先级，而未分配命名空间的设备则保留环境变量的默认设置。
基于轮询的内存统计回退机制
当cgroup阈值通知缺失或不可靠时，基于轮询的内存统计提供了一种回退方案。在进行驱逐决策时，MPS可以直接读取当前使用情况和限制，而不依赖缓存的通知状态。要强制启用轮询，可在内存统计开始前运行nvidia-cuda-mps-control memacct set force_polling=true命令。这会为每次驱逐决策增加一定的工作量，但消除了对通知传递的依赖。
CUDA图中的批量一维复制
CUDA 13.5为批量一维内存复制增加了CUDA图支持。应用程序现在可以在流捕获中包含现有的cudaMemcpyBatchAsync调用，而无需更改批量复制调用的语法，从而使多个传输能够作为可复用图工作流的一部分被捕获和重放。应用程序还可以显式添加批量复制节点，并在不改变批处理结构的情况下更新兼容的复制地址和大小。
NVIDIA Rubin上的Windows驱动模式
从NVIDIA Rubin架构开始，CUDA支持Windows显示驱动模型（WDDM）和微软计算驱动模型（MCDM），其中MCDM提供纯计算驱动模式。Tesla计算集群（TCC）模式在Rubin及后续GPU上不再受支持。假定使用TCC模式的应用程序和部署工具，应针对目标系统切换至受支持的WDDM或MCDM配置。
CUDA Tile更新
CUDA Tile中的多GPU通信
CUDA 13.5新增了在CUDA Tile内核内部直接进行多GPU通信的支持。单播指针允许内核访问通过NVLink连接的另一GPU上的内存，而多播指针和专用的多内存操作则支持在配备NVIDIA NVLink SHARP功能的NVIDIA NVSwitch架构下跨多GPU进行归约和存储操作。应用程序在启动前准备好共享内存映射，随后在内核中将通信与计算结合起来，从而支持诸如融合矩阵乘法和跨GPU归约等工作负载。
Tile与SIMT互操作性（公开预览）
CUDA 13.5允许CUDA Tile内核直接调用SIMT库的设备函数。此前，Tile代码与SIMT库代码只能跨内核边界组合，数据需要通过全局内存在二者之间传递。在此预览版本中，使用cuTile Python或CUDA Tile C++编写的内核可以在自身代码中间调用SIMT库的设备函数，直接处理已存在于寄存器或共享内存中的数据。这些库包括NVIDIA NVSHMEM、NVIDIA CCCL Coop、NVIDIA Warp和NVIDIA MathDx（涵盖NVIDIA cuFFTDx、NVIDIA cuBLASDx、NVIDIA cuSolverDx和NVIDIA nvCOMPDx），这些库即将各自启用此项互操作支持。
Tile库（.tilelib）为SIMT库函数打包了接口和实现信息。编译器会按照库所期望的方式在各线程间划分每个Tile，传递任何共享内存和同步需求，并生成单一内核。要在C++中启用此预览功能，需使用nvcc --tile-preview simt进行编译，并将tilelib传递给Tile编译器。在cuTile Python中，可使用cuda.tile_preview包提供的cuda.tile_preview.foreign_call。
在13.5版本中，Tile-SIMT互操作预览仅涵盖库函数。从Tile代码调用自定义的__device__函数的功能计划在未来版本中推出。更多信息请参阅相关文档。
NVIDIA CUDA Python
NVIDIA CUDA Python持续扩展对CUDA编程的Python化访问能力。CUDA Python核心库cuda.core 1.2引入了批量内存复制、程序化依赖启动（PDL）、更简单的运行时编译以及更灵活的CUDA图更新功能。
批量内存复制
新的cuda.core.utils.copy_batch()函数通过单次CUDA驱动调用提交多个缓冲区复制操作。CopyOptions提供了对源访问顺序、内存位置提示以及重叠行为的控制，使应用程序对数据移动拥有更多控制权。
以下示例使用现有的源和目标缓冲区提交两个独立的复制操作：
（代码示例）
Python中的PDL
LaunchConfig现在通过programmatic_stream_serialization选项支持PDL。在计算能力为9.0或更高的GPU上，相互协作的内核可以在同一CUDA流中重叠执行，使得依赖内核能够在前驱内核完成之前就开始独立的工作。这使Python应用程序能够访问内核间重叠执行能力，同时保留对启动配置的显式控制。详见LaunchConfig文档。
（代码示例）
更灵活的CUDA图更新
应用程序可以在未来实例化之前更新图定义节点，并通过graph[node]视图修改可执行图节点，而无需重新实例化整个图。支持的更新内容包括内核参数、内存复制、内存设置、子图、事件以及主机回调。这些功能简化了在应用程序参数变化时复用图的流程。
（代码示例）
CCCL
CUDA Toolkit 13.5带来了CCCL 3.5版本，新增了CUB的公开调优API、可复现的浮点扫描功能、新的批量Top-K选择操作，以及可直接从GPU内存读取问题规模的归约操作。这些新增功能为开发者提供了更多算法执行控制权，并简化了处理数据依赖型工作负载的流程。
通过公开调优API定制CUB性能
CCCL 3.5为CUB的设备级算法引入了公开调优策略。开发者可以通过执行环境中的cuda::execution::tune(...)来定制线程块大小、每线程处理项数、向量化方式以及算法选择等参数。
以下示例展示了策略选择器如何根据GPU计算能力选择调优参数。该选择器为设备级归约操作返回cub::ReducePolicy；扫描操作的选择器则会返回cub::ScanPolicy。未提供自定义策略的应用程序将继续使用CUB的默认调优方案。更多示例请参阅CUB调优文档。
（代码示例）
批量Top-K选择
Top-K选择是AI工作负载中的重要操作，包括文本生成过程中的Token选择以及稀疏注意力机制中相关Token的筛选。新的cub::DeviceBatchedTopK可在众多数据段中独立选取最小或最大的K个元素，支持键值对、可变数据段大小以及可按数据段变化的K值。NVIDIA在开发此算法时充分考虑了AI工作负载的需求。CUB的批量Top-K功能已集成到FlashInfer中——这是一个用于大语言模型推理的GPU内核库。
（代码示例）
可复现的浮点扫描
cub::DeviceScan新增了对浮点求和运行间可复现性的支持。现在开发者可以通过执行环境请求固定的归约顺序，在相同GPU、相同输入和相同构建配置下产生可重复的结果。
以下示例展示了如何为包含性求和请求可复现的结果。详见cub::DeviceScan文档。
（代码示例）
让数据依赖型归约操作保留在GPU上
CUB归约操作现在可以使用cuda::args::deferred来消费存储在设备可访问内存中的元素计数。前置内核或选择操作可以生成该计数，而归约操作能够按流顺序读取该计数，无需将其传输至主机。这同时也使得被捕获的归约操作可以在每次CUDA图重放时处理不同的元素计数，而无需更新或重新捕获该图。
例如，当某个选择操作将其输出计数写入d_count后，归约操作可直接使用该计数：
（代码示例）
详见cub::DeviceReduce文档。
全新开发者工具
NVIDIA Nsight Perf SDK for AI——公开预览
CUDA Toolkit 13.5推出了NVIDIA Nsight Perf SDK for AI的公开预览版，为长时间运行的训练和推理工作负载带来了常驻性能分析能力。该SDK能够保留原始运行过程中选定的GPU与CPU性能指标及执行上下文，帮助开发者排查那些难以复现的间歇性性能问题。开发者可在Perfetto中查看收集到的数据。详情请点击此处了解更多。
其他新增功能
NVIDIA Nsight Compute 2026.4在NVIDIA Nsight Compute中引入了专门的时间线页面，使开发者能够更方便地探究GPU工作负载随时间变化的性能表现。该页面将时间分辨的性能分析数据汇集到一个中心视图中，帮助开发者识别聚合指标可能掩盖的工作负载阶段、瞬时停顿以及硬件利用率变化。开发者可以在不离开该页面的情况下切换时间线并访问集成的函数统计分析功能，使相关性能洞察保持集中呈现。详情页面提供的直接链接简化了跳转至相应时间线的操作，而内置指引功能则在报告不包含时间线数据时帮助用户快速上手。
CUDA Toolkit 13.5在CUDA-GDB中新增了基于CUDA驱动日志事件停止执行的支持，帮助开发者在错误发生时及时进行排查。可配置的规则能够按线程、严重程度或消息模式过滤事件，使开发者能够专注于相关的诊断信息。当匹配的事件触发执行停止时，开发者可以检查主机端调用堆栈，追溯触发代码并确定根本原因。
此版本还支持在启用CUPTI跟踪的同时生成GPU核心转储。开发者可以在异常发生时捕获GPU状态，并随后使用受支持的调试器检查该转储，而无需在禁用跟踪的情况下重现该故障。
全新的NVIDIA Compute Sanitizer智能体技能可帮助AI编码智能体选择恰当的诊断方式，并解读CUDA内存错误、竞态条件错误以及同步错误。
NVIDIA Nsight Cloud现已支持通过Nsight Operator为NVIDIA Dynamo部署提供OpenTelemetry跟踪集成和性能分析配置，帮助开发者大规模排查AI工作负载性能问题。
数学库
CUDA 13.5中的核心CUDA数学库在Vera Rubin架构上实现了性能提升。cuBLAS将Ozaki-II方案扩展至FP64对称秩k更新操作。其他新功能和值得关注的性能提升包括：
NVIDIA cuBLAS
NVIDIA Vera Rubin上更快的GEMM运算：经过更新的启发式算法和内核调优，在Vera Rubin架构上开箱即用的NVFP4和MXFP8矩阵乘法性能得到提升，涵盖混合专家（MoE）大语言模型训练场景。训练和推理工作负载无需针对特定形状进行自动调优即可获得更高吞吐量。
FP64对称秩k更新的Ozaki-II方案：Ozaki-II方案现已在NVIDIA Ampere及更新架构的GPU上支持{D,Z}SYR{K,2K,KX}和ZHER{K,2K,KX}运算。因此，依赖cuSOLVER构建的Cholesky分解和最小二乘求解器相关应用程序，如今这些运算也能获得Tensor Core加速。
NVIDIA Blackwell上更快的原生FP64对称更新：在NVIDIA B200上，双精度和双精度复数的SYR2K、HER2K、SYRKX和HERKX运算速度提升最高达1.5倍。使用原生FP64对称更新的科学计算代码（例如特征值求解器和矩阵分解）无需修改代码即可获得加速。
针对非对齐形状的优化BF16内核：在NVIDIA Hopper及更新架构的GPU上，维度不是8的倍数的BF16矩阵乘法现在使用与FP16相同的优化内核。包括通过NVIDIA cuDNN调度的卷积运算在内，具有不规则层大小的模型在选择BF16而非FP16时不再需要承担性能损耗。
Blackwell GPU的Windows on Arm支持：cuBLAS新增了对Blackwell GPU（包括GB300工作站）在Windows on Arm平台上的功能支持。基于Arm架构的Windows系统开发者无需切换平台即可运行cuBLAS加速的应用程序。
NVIDIA cuSPARSE
更快的FP32和FP64 CSR SpMM（ALG3算法）：在数据中心GPU上实现1.65倍至2.27倍的性能提升，加速稀疏矩阵与密集矩阵的乘法运算。
全新SpSV ALG2算法：相较于ALG1，几何平均求解性能提升2.7倍，加速了需要反复求解稀疏三角系统的工作负载。
更快的SpSV分析阶段：分析阶段的几何平均速度提升2.79倍，缩短了稀疏三角求解的准备时间。
更低的SpSV内存占用：降低了稀疏三角求解的内存占用，为应用程序数据留出更多可用的GPU内存。
扩展的混合索引支持：在SpMV（CSR、SELL）、SpMM（CSR）、SpSV（SELL）以及稀疏-密集/密集-稀疏转换（CSR、CSC）中使用64位偏移量配合32位索引，从而支持更大规模的稀疏矩阵，同时保持比完全64位索引更低的索引内存占用。
混合精度CSR SpMV：在cusparseSpMV中结合FP16或BF16矩阵值与FP32输入/输出向量及FP32计算，在保留FP32累加精度的同时降低了矩阵存储需求。
全新实验性cusparseSpMVGeneric API：对对称矩阵执行SpMV运算时仅需存储其上三角或下三角部分，通过避免重复存储非对角元素来降低存储需求。
全新实验性cusparseSpMatConvert API：实现CSR与CSC格式之间的转换，使稀疏矩阵布局更易于适配不同的运算操作。
csrcolor实现更快的图着色：几何平均性能提升约3倍，减少了稀疏矩阵着色所耗费的时间。
nvJPEG
GPU混合解码吞吐量：根据图像大小以及是否使用标准霍夫曼表的不同，NVJPEG_BACKEND_GPU_HYBRID后端的解码吞吐量提升了10%至600%。
自动后端选择：使用nvjpegDecode API时，NVJPEG_BACKEND_GPU_HYBRID现已纳入自动后端选择范围（此前仅考虑NVJPEG_BACKEND_HYBRID和NVJPEG_BACKEND_HARDWARE）。
编译器更新
使用nvprune在打包阶段预编译设备代码
CUDA 13.5在nvprune中引入了--precompile选项，允许开发者在剪枝步骤中将中间设备代码（PTX、TileIR）编译为原生SASS/ELF格式。
这消除了那些向已知目标架构部署胖二进制文件的下游用户（例如云服务提供商、软件包分发商）所面临的运行时JIT编译开销。
开发者应选择与自身部署环境相匹配的目标架构，并在计划支持的各GPU架构上验证最终生成的产物。
扩展CompileIQ支持
CUDA 13.5将高级控制文件（ACF）支持扩展至PTX编译器和NVVM库。这使得基于这些库构建的编译工作流能够应用由CompileIQ生成的编译器控制指令。
更精细的TMA寻址控制与更快的基于周期的三角函数运算
CUDA 13.5扩展了PTX和设备内置函数的能力，帮助开发者优化数据移动和数学运算。
在内核内部配置TMA越界寻址
平铺计算经常会遇到张量边界情况，即某个数据块的部分区域超出了张量的维度范围。高效处理这些不规则数据块对于MoE、矩阵乘法和卷积等工作负载非常重要。
CUDA 13.5扩展了PTX的tensormap.replace指令，使内核能够直接更新张量映射的张量内存加速器（TMA）越界寻址模式。这使开发者在复用张量映射描述符时能够控制寻址行为，并可消除受影响工作负载中针对非法地址故障的性能代价较高的变通方案。
基于周期输入的快速近似单精度三角函数计算
CUDA 13.5新增了__sin2pif、__cos2pif和__sincos2pif设备内置函数，用于对以周期表示的输入进行快速近似的单精度三角函数计算。
这项功能对于可以使用快速近似三角函数结果的基于相位的计算非常有用。
相应的PTX中sin和cos指令的.2pi变体，为已经使用基于周期相位的工作负载避免了额外的角度转换运算。
使用cuobjdump更轻松地读取CUDA二进制符号
CUDA 13.5改进了cuobjdump中的ELF符号表输出，使开发者更容易理解CUDA二进制文件中包含的函数和变量。
此前，符号信息和CUDA特有的属性以原始十六进制字段形式呈现，开发者需要手动解码。更新后的输出以四列形式呈现描述性数值，遵循readelf -s用户熟悉的惯例。
这使得区分内核入口点与设备函数、检查与全局、常量、共享和托管变量相关的属性变得更加容易。当多个CUDA属性同时适用时，输出会将它们的名称一并列出。无法识别的属性值仍会以UNKNOWN(0xNN)的形式显示。
此项增强功能适用于现有的--dump-elf和--dump-elf-symbols选项。例如，开发者可以通过以下命令检查CUDA二进制文件：
cuobjdump --dump-elf-symbols application.cubin
无需任何额外选项即可获得这种描述性输出。现有信息——包括符号名称、数值、大小和段索引——仍然保持可用。
如果您使用脚本解析这些符号表转储数据，请更新脚本解析器以适配新增的列和命名属性值。
开始使用CUDA Toolkit 13.5
CUDA Toolkit 13.5为NVIDIA Vera Rubin GPU架构提供了正式发布支持，更新了GPU资源管理和通信能力，并在CUDA Python、CCCL、CUDA Tile、NVIDIA Nsight开发者工具以及核心数学库等方面带来了诸多增强功能。
下载CUDA Toolkit 13.5，并查阅CUDA Toolkit 13.5发行说明以获取完整的功能列表、支持平台及兼容性信息。
Jackson Marusarz、Mahender Hari、Mridula Prakash、Daniel Rodriguez和Sandeep Hiremath为本文的撰写做出了贡献。
