8 月 30 日 CUDA 与 AI 系统加速速览：Triton 寄存器调度、TVM 精确 launch 与多 GPU 清理 (2026)

开源加速栈这轮新增，分别落在编译器选择、CUDA launch 契约和多 GPU 运行时清理。三项提交都给了可以复用的工程动作；公开页面给出的证据主要是 PTX 变化、属性检查和回归测试，端到端吞吐数字仍要在目标硬件上自测。

快速判断

分组	更新	时间（+08:00）	来源已报告结果	适用场景	第一步验证	行动窗口
AI 编译器 / kernel	Triton per-kernel NVPTX register-pressure scheduling：为单个 kernel 增加 `sched4reg` 选择	8 月 30 日 03:37	默认调度与寄存器压力调度会生成不同 PTX；并发测试验证两种策略的结果稳定，页面给出代码生成测试，端到端性能数字留给自测。1 2	NVPTX kernel 的寄存器占用、occupancy 或 spill 成为调优变量时	对同一 Triton kernel 分别编译 `sched4reg=false/true`，检查 PTX、寄存器数、occupancy 和正确性	可先放进编译回归；性能验证应绑定具体 shape 与 GPU
AI 编译器 / kernel	TVM TIRx exact required block dimensions：把静态线程 / cluster 维度变成 CUDA 的精确 launch 契约	8 月 29 日 09:42	生成 `__block_size__` 与 PTX `.reqntid` 的代码生成测试通过；属性互斥和 host/device metadata 测试覆盖，页面给出 CUDA 13+ 边界，端到端性能数字留给自测。3 4	需要固定 block / cluster 尺寸、依赖 `.reqntid` 的 TIRx CUDA kernel	用 CUDA Toolkit 13+ 编译一个静态线程与 cluster 维度的 kernel，核对生成源码、grid 可整除条件和 launch 失败信息	适合在升级 TIRx 或重写 launch metadata 前做兼容性门禁
CUDA runtime	TVM cleanup 保持当前设备：`CUDADeviceGuard` 包住跨设备释放与 CUDA IPC pooled allocation 清理	8 月 30 日 05:20	多 GPU destructor / stream 回归通过；三项目标清理回归通过，TIRx 测试报告 2902 passed、94 skipped、3 xpassed，页面未给运行时加速数字。4 5	Python GC、runtime destructor 或 Disco CUDA IPC 释放跨设备资源的多 GPU 进程	让线程停在 `cuda:1`，释放 `cuda:0` 的 tensor、stream 和 pooled IPC allocation，再断言当前设备仍为 `cuda:1`	适合直接加入多 GPU runtime 回归；收益先看状态稳定性和后续算子是否落到正确设备

AI 编译器 / kernel：Triton 把寄存器压力调度变成单 kernel 选项

Triton 在 2026 年 8 月 30 日 03:37（UTC+08:00）合入 c349ce5。提交在 NVPTX 后端注册一个编译期调度器：sched4reg=false 走默认调度，sched4reg=true 走 createBURRListDAGScheduler，并把选项挂到单个 kernel 的 CUDAOptions 上。这个调度器只对 NVIDIA 的 NVPTX 目标启用。1 2

这里的实用变化是调度策略可以跟着 kernel 走，而不是把整个编译器固定在一种选择上。提交里的单元测试构造了一个带整数乘法和全局写入的 NVPTX kernel，分别使用两种选项生成 PTX，并断言两份 PTX 不同。并发测试再交替编译 24 次，确认线程池里的策略不会串到另一个 kernel。1

寄存器压力调度适合放进寄存器数已经影响 occupancy，或局部变量过多导致 spill 的 kernel 调优流程。来源证明了策略确实改变了 codegen，也证明了并发编译的隔离；页面给出的结果停在代码生成层。团队需要自己测寄存器数量、active warps、spill、kernel latency 和端到端吞吐，不能把“生成了不同 PTX”直接换算成加速比。

最小验证可以这样做：固定 Triton commit、CUDA Toolkit、计算能力、输入 shape 和 warm-up 次数，编译同一个 kernel 的两份版本。先比较 ptxas 的寄存器与 spill 报告，再用 Nsight Compute 或等价工具观察 occupancy，最后用目标服务的真实 batch 测 latency 和吞吐。true 只有在这三层结果都变好时，才值得进入默认配置。

AI 编译器 / kernel：TVM TIRx 增加精确的 block 尺寸契约

Apache TVM 在 2026 年 8 月 29 日 09:42（UTC+08:00）合入 0d9c9384。提交新增 tirx.required_block_size=1 属性，并在 CUDA codegen 中生成 __block_size__((tx, ty, tz), (cx, cy, cz))。CUDA 会由这个精确契约导出 PTX .reqntid；它与只提供建议的 __launch_bounds__ 属于两种不同的约束。3 4

这条路径有三个需要先锁住的边界。第一，提交把要求限制在 CUDA Toolkit 13 及更新版本。第二，tirx.required_block_size 不能与最大寄存器数、每个 SM 的最小 block 数或每个 cluster 的最大 block 数同时使用。第三，采用精确 block / cluster 尺寸时，运行时 grid 的每个维度需要能被对应的静态 cluster 维度整除。3

测试覆盖了两层：TIR transform 把属性传成 flag-only launch metadata，CUDA codegen 生成 __block_size__，并检查源码里没有 __launch_bounds__；另一组测试把精确 block 尺寸和 launch bounds 放在一起，确认编译阶段报出互斥错误。页面给出的证据是契约和回归测试，性能收益需要按具体 block / cluster 形状测量。3

如果 kernel 依赖固定线程数才能正确使用 cluster 内通信，升级时可以先做一个最小 TIRx 样例：静态声明 thread 与 cluster 维度，检查生成源码，再故意传入不可整除的 grid 和冲突的 launch 属性。编译器先把错误挡在 launch 前，团队再比较精确契约和原有 launch 参数下的 kernel latency 与服务指标。

CUDA runtime：TVM 修复跨设备清理留下的当前设备污染

TVM 在 2026 年 8 月 30 日 05:20（UTC+08:00）合入 66efb012。提交把 CUDA 释放路径里的直接 cudaSetDevice 换成 CUDADeviceGuard：释放 tensor / workspace 和销毁 stream 时，代码临时切到资源所在设备，作用域结束后恢复调用线程原来的当前设备。Disco CUDA IPC pooled allocation 的释放路径也采用同一做法。这里的“当前设备”指 CUDA 线程上下文里后续算子默认使用的 device。4 5

这个修复针对的是很容易被业务代码掩盖的一类问题：Python GC 或 runtime object destructor 可能在调用线程停留于 cuda:1 时，释放属于 cuda:0 的资源。旧路径会把线程留在资源设备；后续 DLPack 导出、CuTe DSL 操作或下一个未显式指定 device 的算子就可能观察到错误的当前设备。提交把保护放回 runtime cleanup 内部，调用方无需在每个引用周围重复设置 device。5

提交新增了多 GPU NDArray destructor 和 stream destruction 回归，也覆盖 pooled CUDA IPC 清理。官方记录显示目标清理回归 3 项通过；TIRx 测试在三张 Blackwell GPU 上报告 2902 passed、94 skipped、3 xpassed。fast_topk_clusters 的三项默认 correctness case 被标为 skip，原因是 FlashInfer 参考实现缺少输入边界检查，也没有初始化或重置共享 threshold bin；严格 registry import 覆盖仍保持开启。5

这条更新的第一收益是设备状态可预测，性能数字要另行测量。多 GPU 服务可以加入一个很小的回归：先把线程设为 cuda:1，在 cuda:0 分配 tensor 和 stream，触发释放，再断言当前设备仍是 cuda:1。随后把同一检查放进 CUDA IPC pool 和异常清理路径，覆盖正常返回、Python GC 与进程退出前的 destructor。

推荐系统：本窗口暂留空位

本窗口的推荐系统合格新增条目为零。当前材料集中在 Triton 与 TVM 的编译、launch 和 runtime 正确性，旧论文与旧版本即使主题相关，也无法给本期增加新的工程信息，因此保留空位。

按系统形状排验证

需要压低寄存器压力的 Triton kernel： 先固定 sched4reg 两个编译分支，比较寄存器、spill 和 occupancy，再看真实 shape 的 latency 与吞吐。
依赖固定 block / cluster 尺寸的 TIRx CUDA kernel： 使用 CUDA Toolkit 13+，验证 __block_size__、grid 整除和 launch 属性互斥，再测 kernel 性能。
多 GPU TVM runtime： 把跨设备 tensor、stream、CUDA IPC allocation 的释放放进当前设备回归，覆盖 GC、destructor 和异常路径。
所有三项更新： 记录 commit、CUDA / GPU、输入 shape、编译时间、正确性结果和端到端指标；提交页面给出的回归结果不能替代目标系统的性能数字。

References

1
Triton 寄存器压力调度提交
github.com
2
Triton 官方提交时间线
github.com
3
TVM exact required block dimensions 提交
github.com
4
Apache TVM 官方提交时间线
github.com
5
TVM CUDA device state cleanup 提交
github.com