CUDA 13.3 像把后厨的切配、火候、传菜统统锁进一台自动灶:厨师第一次能在 C++ kernel(内核函数)里先想“这道菜怎么成”,不必先想“32 个线程站哪口锅”。我判断,13.3 第一次让 C++ kernel 不必先想线程。

多数人看 CUDA 更新,盯的是快了几个点、支持了哪块卡、又添了哪门库。那像看饭店只看门头,不看厨房。13.3 真正反常识的地方,不是线程更好写了,而是作者终于可以晚一点想线程,甚至在 kernel 里根本不必亲自想它。

过去写 CUDA C++,像办婚宴,菜还没定,先得安排每张桌子坐几个人。SIMT(单指令多线程式写法)训练你的第一反应,不是算法,而是 threadIdxblockIdx、边界判断、共享内存、同步。你想写的是向量加法,手里先拿起的却是座位图。

这就是旧时代最荒唐的地方:你明明要算矩阵,却先学会点兵。

NVIDIA 这次把 CUDA Tile C++(CUDA 的分块式 C++ 写法)正式放进 13.3。官方文档讲得很直白:以前程序员显式处理线程级索引;现在你在 tile(数据分块)层面装载、计算、回写,由编译器把这些操作摊到硬件线程上。连 launch(启动配置)都露了馅:tile kernel 的第二个参数必须写 1。在作者视角里,一个 block(线程块)先被当成一个“逻辑上的人”,块里的并行排班,交给编译器。

线程没有消失,只是从作者的前台,退回了编译器的后台。

这不是语法糖,这是权力转移。过去大家迷信一句话:离硬件越近,离性能越近。13.3 给这句老话递了辞退信。今天的 GPU(图形处理器)越来越像一座有暗道、有升降梯、有专用货梯的剧院;你若还要求每个作者亲自指挥每个跑龙套站位,只能说明这座剧院的管理太落后。真正先进的体系,不是把程序员训练成线程调度员,而是把线程调度员塞回编译器里。

GPU 编程最贵的,不是算力,是把人的脑子掰成线程形状的培训费。

看三个场景,就知道这事为什么值钱。

凌晨一点四十,张江一间做视频模型的小办公室里,一个工程师在改一个长度不整齐的归一化 kernel。老写法里,他盯着 idx = blockIdx.x * blockDim.x + threadIdx.x,补尾巴判断,怕越界,怕访存不整齐,怕同步点多一拍少一拍。新写法里,他关心的是:这一块 tile 读进来,归一化,再 store_masked 回去。那几粒残边数据,不再像掉在地上的米,逼着人蹲下去一粒粒捡。

以前,算法是正文,线程却抢着当标题。

第二个场景,在苏州,一支做推理服务的小队把代码从 Hopper(NVIDIA 2022 年数据中心架构)往更新的卡上迁。过去他们的经验像地方戏:这张卡唱这个腔,那张卡换那个身段,线程布局、共享内存、Tensor Core(张量核心)吃法都得重背。Tile 写法的野心,是让你写“32x32 的砖怎么搬、怎么乘、怎么落”,至于底下调用 Tensor Memory Accelerator(张量内存搬运单元)还是别的硬件本事,交给编译器。可移植性第一次不是发布会上的慈善词,而是人月成本的减刑。

第三个场景,在成都,一个后端 C++ 工程师被临时拉来补 GPU 缺口。以前这事像把会写小说的人抓去背交规:还没谈剧情,先背红绿灯、车道线、让行规则。他得先学会像显卡那样数人头,才配写算法。13.3 的变化,是他终于能先写“这块矩阵怎么算”,再让工具决定“谁踩油门,谁换挡”。这改变的不只是代码写法,还会改变团队里谁有资格碰 GPU,谁在评审里说了算。

一门技术成熟,不是逼人更像机器,而是允许机器去做机器该做的事。

所以我觉得,CUDA 13.3 的要害,不在“又新增了一个 API(接口)”,也不在“又涨了几个百分点”。同一版里当然还有 CompileIQ(编译器自动调优框架)、CUDA Python 1.0、C++23 支持,这些都热闹;但热闹是烟花,Tile C++ 才像改了地基。它把 GPU 编程从“先排线程,再写算法”,扭成了“先写算法,再让编译器调度线程”。

这一步看着像退,其实是进。因为只有当线程不再霸占作者的脑门,C++ kernel 才真正开始说人话。过去那套写法,本质上是一种手工业:好工匠靠经验,靠祖传口诀,靠对硬件脾气的驯服。13.3 想做的,是把这门手工业往工业推一步。不是取消匠人,而是不再要求每个进门的人,先去学会拧每一颗螺丝。

13.3 第一次让 C++ kernel 不必先想线程。线程不会死,就像汇编不会死;但从这一天起,线程更该待在地下室,而不是坐在写作者的桌面中央,当所有算法的门神。

文中涉及的版本信息见 NVIDIA 官方发布与文档:CUDA 13.3 发布博客CUDA Tile C++ 博客CUDA Programming GuideCUDA Tile C++ API Reference