AI写算子时代：手写CUDA工程师的抽象阶梯与转型路径

1. 从刘胜的告别说起一个时代的分水岭刘胜是我认识快八年的一位老哥早年在某芯片大厂做高性能计算后来跳去一家自动驾驶公司带算子团队。上个月他发了一条朋友圈大意是干了十二年手写CUDA今天正式转岗去做AI编译器方向手写kernel这条路我认输了。底下评论区炸了一半人在感慨一半人在问是不是真的没得做了。这条朋友圈之所以值得拿出来聊不是因为某个人的职业选择而是它精准地戳中了一个正在发生的结构性变化计算机抽象阶梯正在往上爬一层而这一层恰好踩在了算子这个位置上。过去十几年算子开发是高性能计算领域最硬核、最吃经验、也最不容易被替代的工种之一。你得懂硬件架构、懂内存层级、懂指令流水、懂并行规约还得有那种看一眼profile就知道瓶颈在哪的直觉。一个成熟的CUDA工程师培养周期动辄三五年薪资也一直居高不下。但现在AI开始学会写算子了——不是写那种玩具级的demo而是能跑进生产环境、性能逼近甚至超过手写版本的算子。这件事的意义远比某个岗位要消失了要大得多。它意味着抽象阶梯的下一层正在被AI接管而人类工程师的位置正在被推向更高一层的抽象。我写这篇东西不是想贩卖焦虑也不是想鼓吹什么。我只是想把这几年在算子开发、AI辅助编程、编译器工具链这几个交叉地带踩过的坑、看过的变化、想明白的逻辑系统地梳理一遍。如果你正在做算子开发或者正在考虑要不要入这行或者你是个技术管理者在思考团队方向这篇内容应该能给你一些参考。先明确一下讨论范围。这里说的算子指的是深度学习框架和科学计算里那种可被调度、可被融合、可被编译的计算单元比如矩阵乘、卷积、归一化、激活函数、各种element-wise操作。不是数学意义上那种纯符号的算子也不是Halcon那种图像处理算子库里的算子——虽然它们共享算子这个名字但工程语境完全不同。热词里出现的sobel、canny、拉普拉斯这些属于传统图像处理算子和深度学习算子有交集但不等同后面会单独聊。2. 算子到底是什么从数学符号到GPU上的一行行指令2.1 算子的三层含义别搞混了算子这个词在不同语境下指的东西差别很大我见过太多新人一上来就懵。简单分三层数学层算子是一个从函数空间到函数空间的映射。拉普拉斯算子、梯度算子、傅里叶变换算子都是这个层面的东西。这一层是纯理论和工程实现无关。框架层在PyTorch、TensorFlow这类框架里算子是一个有名字、有输入输出签名、有梯度定义的计算节点。比如torch.nn.functional.relu、aten::addmm。这一层是逻辑概念描述要算什么。硬件层在GPU上算子最终要被翻译成一组线程如何组织、数据如何在寄存器/共享内存/全局内存之间流动、指令如何发射的具体方案。这一层才是真正决定性能的地方也是手写CUDA工程师的主战场。刘胜说的手写kernel指的是第三层。AI学会写算子目前主要冲击的也是第三层——但正在快速向上侵蚀第二层。2.2 一个算子从框架到GPU的完整旅程拿一个最简单的例子C A B bias也就是带偏置的矩阵乘。在PyTorch里你写一行torch.addmm(bias, A, B)背后发生了什么第一步框架的dispatcher根据输入张量的设备、数据类型、形状选择一个具体的kernel实现。这一步叫算子分发。第二步如果开了图模式比如torch.compile或者TensorRT这个算子和它前后的算子可能被融合成一个更大的kernel减少kernel launch开销和内存往返。第三步选定的kernel实现被编译成PTXParallel Thread ExecutionNVIDIA的中间表示再由驱动编译成SASS实际的GPU机器码。第四步GPU的线程调度器把成千上万个线程分配到SM流多处理器上每个线程按照kernel里写的逻辑从全局内存读数据、在寄存器里算、写回结果。这四步里第一步和第二步是框架和编译器的事第三步和第四步是kernel实现的事。手写CUDA工程师的价值集中在第三步和第四步——怎么写出一个让SM利用率高、内存带宽打满、指令流水不 stall 的kernel。2.3 为什么手写kernel这么难难在你同时要管三件事而且这三件事互相冲突。第一件是正确性。矩阵乘的边界处理、非对齐内存访问、bank conflict、warp divergence任何一个细节错了结果就是错的而且往往错得很隐蔽。第二件是性能。你要考虑tiling策略、shared memory的使用、寄存器压力、指令级并行、occupancy和ILP的权衡。这些参数之间是耦合的调一个动全身。第三件是可移植性。A100上跑得好的kernelH100上未必好CUDA上跑得好的迁移到其他加速器平台可能要重写。热词里出现的cuda迁移ascendc融合算子就是这个痛点的体现。一个成熟的CUDA工程师脑子里同时装着这三件事靠的是大量实战积累出来的直觉。这种直觉很难被文档化也很难被传授——这正是AI切入的缝隙。3. AI写算子的三条技术路线以及它们各自的天花板AI学会写算子不是一个单点技术而是三条路线在同时推进。搞清楚这三条路线的区别你才能判断自己所在的岗位到底受多大影响。3.1 路线一基于模板的搜索与调优这是最早成熟、也最务实的一条路。代表就是TVM的AutoTVM、Ansor以及各种auto-tuning框架。核心思路是人类专家写好一个参数化的kernel模板把tiling大小、线程块维度、向量化宽度、unroll因子这些做成可调参数然后让搜索算法遗传算法、模拟退火、贝叶斯优化去自动找最优参数组合。这条路线的本质是把人类的经验固化成搜索空间让机器在空间里找最优解。它不需要AI理解kernel在干什么只需要AI会搜索。优点是稳、可控、可解释。缺点是天花板明显——搜索空间是人定的人没想到的优化搜索算法也找不到。而且调优过程很慢一个复杂算子调几天是常事。我个人的判断这条路线的价值在于把重复劳动自动化它替代的是调参这个环节而不是设计kernel这个环节。所以它对高级工程师的威胁不大对初级调参工程师的威胁比较大。3.2 路线二基于大模型的代码生成这是这两年最热的方向。用GPT类模型或者专门训练的代码模型直接根据算子描述生成CUDA代码或者Triton代码。代表工作不少思路大同小异给模型看大量的CUDA代码、PTX代码、性能数据让它学会什么样的算子描述对应什么样的实现。这条路线的想象空间很大但目前的实际效果说实话在简单算子上已经可用在复杂算子上还差得远。我实测过几个主流方案生成element-wise算子比如各种激活函数、逐元素运算基本一次过性能也能到手写版本的七八成。但一旦涉及reduction、矩阵乘、卷积这种有复杂数据依赖和内存访问模式的算子生成的代码要么性能拉胯要么直接编译不过。原因也不难理解大模型擅长的是模式匹配和统计规律而高性能kernel的核心是对硬件微架构的精确控制。后者需要的是确定性的推理不是概率性的生成。3.3 路线三基于编译器的自动生成这条路线的代表是TVM的TensorIR、Halide、以及各种polyhedral编译技术。核心思路是把算子用高层声明式语言描述比如这是一个矩阵乘tile大小是32x32然后编译器自动完成调度、优化、代码生成。这条路线的理论最优雅因为它把写kernel变成了写调度。人类只需要描述要什么编译器负责怎么实现。但它的现实困境是调度空间太大编译器找不到最优解。而且很多硬件特性比如Tensor Core的特定指令、shared memory的bank结构很难用通用的调度原语精确表达。热词里出现的llvm算子自发现就属于这个方向的前沿探索——让编译器自动发现算子之间的等价变换和融合机会。这个方向很有前途但离完全替代手写还有相当距离。3.4 三条路线的对比路线核心机制优势天花板对工程师的冲击模板搜索参数化模板自动调优稳定可控受限于模板设计替代调参环节大模型生成代码模型直接生成泛化性好复杂算子性能差替代简单算子编写编译器生成声明式描述自动调度理论优雅调度空间爆炸替代调度设计看清楚这张表你就明白为什么刘胜会认输了——他做的那些重复性的、模式化的kernel编写工作正在被这三条路线从不同角度蚕食。但他真正值钱的那部分能力——对硬件行为的深刻理解、对性能瓶颈的直觉判断、对系统级权衡的把控——恰恰是这三条路线目前都还搞不定的。4. 抽象阶梯的下一层工程师的位置往哪挪4.1 抽象阶梯是什么为什么它一直在往上爬计算机科学的发展史本质上就是一部抽象阶梯不断往上爬的历史。最早的程序员用机器码编程后来有了汇编再后来有了C再后来有了Python和各种DSL。每一层抽象都让程序员离硬件更远但让开发效率更高。被抽象掉的那部分工作要么被编译器接管要么被框架接管要么被工具链接管。算子开发目前处在一个尴尬的中间位置它比框架层低比汇编层高它需要硬件知识但又不像驱动开发那样贴近硅片。这种位置最容易被抽象掉——因为它的很多工作是可以被形式化、被自动化的。AI写算子本质上是把从算子描述到硬件实现这一段抽象过程自动化了。这就像当年编译器把从C代码到汇编这一段自动化了一样。4.2 被替代的是什么没被替代的是什么我观察下来被替代的主要是这几类工作模式化的kernel编写。比如各种element-wise算子、简单的reduction、标准的卷积实现。这些算子的实现模式高度固定AI学一学就能写。参数调优。auto-tuning工具已经做得很好了人工调参的价值越来越低。代码迁移。从一个硬件平台迁移到另一个平台很多工作是机械的AI可以辅助完成。热词里的cuda迁移就是这个场景。没被替代的或者说短期内很难被替代的是这几类新硬件特性的首次适配。当一款新GPU发布它的新指令、新内存结构、新并行模型需要人去摸索、去试错、去总结最佳实践。AI没有先验知识只能等人喂数据。极端性能场景的手工优化。当你要把某个算子的性能压榨到硬件极限需要的是对微架构的精确理解和大量实验这不是AI擅长的。系统级的架构设计。算子怎么切分、怎么融合、怎么调度这些决策涉及整个计算图的全局优化需要的是系统思维不是局部代码生成。跨层的问题定位。当一个性能问题可能出在框架层、编译器层、kernel层、驱动层、甚至硬件层时能快速定位到根因的人价值极高。4.3 工程师该往哪挪我的判断是往抽象阶梯的上游挪往系统层挪往问题定义层挪。具体来说几个方向值得考虑算子编译器方向。不是写kernel而是写生成kernel的编译器。这个方向需要你既懂硬件又懂编译原理门槛高但护城河也深。性能工程方向。不是写单个算子而是做端到端的性能优化。你需要理解整个模型的计算图、内存访问模式、通信开销然后找到系统级的优化机会。AI辅助工具方向。不是被AI替代而是去造替代别人的AI工具。这需要你既懂算子开发又懂机器学习是个交叉地带。领域专用架构方向。针对特定领域比如推荐、科学计算、图计算设计专用的算子库和编译方案。领域知识是AI短期内学不会的。5. 实操用AI辅助算子开发我踩过的坑和总结的流程光聊趋势没意思说点能落地的。这部分我分享一套自己实际在用的AI辅助算子开发流程以及踩过的坑。5.1 环境准备别在环境上浪费时间热词里一堆cuda安装wsl2安装cudacuda多版本安装的问题我太理解了。环境配置是算子开发最烦人的环节之一尤其是多版本共存的时候。我的建议是用容器别在裸机上折腾。NVIDIA的官方CUDA镜像已经做得很好了nvidia/cuda:12.x-devel系列镜像开箱即用。如果你非要在WSL2里搞注意WSL2的CUDA支持需要Windows端的驱动足够新而且WSL2里的CUDA版本不能超过Windows驱动支持的上限。多版本共存的话用update-alternatives管理或者干脆每个项目一个容器。我见过太多人因为环境问题卡了几天这纯粹是浪费生命。提示如果你在编译CUDA代码时遇到in file included from之类的头文件报错九成是CUDA路径没配对检查CUDA_HOME和LD_LIBRARY_PATH。5.2 用AI生成算子代码的正确姿势直接让AI写一个矩阵乘kernel结果通常不能看。正确的姿势是分步骤、给约束、要解释。我的流程是这样的第一步让AI生成参考实现。不要一上来就要高性能版本先要一个正确但朴素的版本。提示词大概是用CUDA写一个矩阵乘kernel要求正确性优先不考虑性能优化每个线程计算一个输出元素。第二步让AI解释这个实现的问题。这一步很关键。让AI自己分析这个实现有哪些性能问题它会列出内存访问不合并、没有用shared memory、occupancy低等问题。这些分析不一定全对但能帮你快速建立优化清单。第三步逐步优化每步验证。让AI针对某个具体问题优化比如用shared memory做tilingtile大小32x32。每优化一步都要实际跑一遍验证正确性和性能。第四步让AI生成测试代码。这一步经常被忽略但极其重要。让AI生成和PyTorch对比的测试代码确保数值正确。这套流程下来一个中等复杂度的算子从零到可用大概能比纯手写快两三倍。但注意最后的性能调优还是得靠人AI生成的代码通常离硬件极限还有距离。5.3 踩过的坑坑一AI生成的代码有隐蔽的race condition。尤其是涉及shared memory和同步的时候AI经常漏掉__syncthreads()或者放错位置。这种bug在小规模测试时可能不出现一上大规模就炸。所以测试一定要覆盖各种形状包括边界情况。坑二AI对硬件特性的理解是过时的。比如你问它Tensor Core怎么用它可能给你一个Volta时代的答案而你的卡是Hopper。所以涉及新硬件特性时一定要自己查官方文档确认。坑三AI生成的性能数据不可信。它可能会说这个实现能达到X TFLOPS但那是它编的。性能必须自己实测。坑四过度依赖AI导致自己能力退化。这个最危险。如果你长期只做提示词工程不去深入理解硬件几年后你会发现自己除了调AI什么都不会。我的做法是AI生成的每一行代码我都要能解释为什么这么写。5.4 一个具体的例子ReLU算子的AI辅助开发拿最简单的ReLU举例说说完整流程。朴素实现是每个线程处理一个元素读、判断、写。AI一次就能写对。但性能上这个实现的问题是内存带宽利用率低因为每个线程只做一次读写没有向量化。让AI优化它会建议用float4向量化一次处理四个元素。这个优化是对的能显著提升带宽利用率。再让AI优化它会建议用grid-stride loop让线程数不依赖于数据规模。这个也对能提升occupancy。到这一步性能基本能到手写版本的九成以上。剩下的那一成可能需要考虑更激进的优化比如用__ldg走只读缓存、调整block大小等这些AI也能给建议但需要你判断是否适用。整个过程AI帮我省掉了写样板代码的时间但判断哪些优化值得做、哪些是过度优化还是得靠人。6. 那些被热词带偏的方向澄清几个常见误解热词列表里混了不少东西有些和算子开发相关有些是噪音。我挑几个容易误导人的澄清一下。6.1 算子不等于图像处理算子热词里的sobel、canny、拉普拉斯、halcon算子都是传统图像处理领域的概念。它们和深度学习算子有交集比如卷积但工程语境差别很大。图像处理算子通常是确定性的、固定参数的、单次执行的优化重点是内存访问和指令效率。深度学习算子则是数据驱动的、参数可变的、在计算图中被调度的优化重点还包括算子融合、图级优化、自动微分等。如果你是从图像处理转过来的要注意这个思维差异。6.2 算子开发不等于CUDA编程CUDA只是算子开发的一种实现手段。现在主流的还有Triton、OpenCL、SYCL以及各家加速器平台自己的编程模型热词里的ascendc就是其中之一。而且随着编译器技术的发展直接用CUDA写算子的场景正在减少更多是用高层DSL描述让编译器生成底层代码。所以如果你只会CUDA不会用Triton或者TVM职业选择会窄很多。6.3 AI写算子不等于AI取代算子工程师这个前面已经聊过了但值得再强调一次。AI替代的是算子开发中的重复性劳动不是算子工程师这个职业。就像编译器没有取代程序员只是让程序员的工作层次提高了。真正危险的是那些只会做重复性劳动、不愿意往上走的工程师。如果你一直在做把论文里的算子用CUDA实现一遍这种工作那确实要警惕。6.4 关于无禁词AI聊天这类热词热词列表里混进来一些和算子开发完全无关的东西比如ai无禁词聊天网页版无限制ai对话聊天之类的。这些是搜索噪音和本文主题无关直接忽略。做技术的人要有信息过滤能力别被这些带偏。7. 给不同阶段从业者的具体建议7.1 如果你刚入行别一上来就死磕手写CUDA。先把框架层和编译器层的东西搞明白PyTorch的计算图是怎么组织的、torch.compile做了什么、TVM的调度原语怎么用。这些是更高层的抽象也是未来更有价值的能力。同时一定要动手写。哪怕AI能生成你也要自己写一遍理解每一行在干什么。这是建立直觉的唯一途径。7.2 如果你有几年经验这是最需要警惕的阶段。你的经验值钱但你的经验里有一部分正在贬值。我的建议是主动做减法把那些重复性的、模式化的工作交给AI和工具把省下来的时间投入到系统级的能力建设上。具体来说去学编译器原理、去学系统性能分析、去学新硬件架构。这些是AI短期内替代不了的。7.3 如果你是技术管理者别急着裁算子团队。但你要重新定义这个团队的职责从写kernel转向设计算子方案把控性能建设工具链。同时要主动引入AI辅助工具提升团队效率。一个会用AI的算子工程师产出能顶过去两三个。这不是危言耸听是我实际观察到的。7.4 一个通用的能力模型不管你在哪个阶段我觉得未来算子方向的核心能力是这四块硬件理解知道数据在硬件上怎么流动知道瓶颈在哪。编译原理知道代码怎么被翻译和优化知道怎么用编译器工具。系统思维知道单个算子在整个系统中的位置知道怎么做全局优化。AI协作知道怎么用AI提升效率也知道AI的边界在哪。这四块里前两块是基础第三块是进阶第四块是必备。缺哪块补哪块。8. 我个人的一些判断和体会聊了这么多最后说点我自己的判断不一定对供参考。第一算子开发这个岗位不会消失但会分化。一部分人往上走做编译器、做系统优化、做架构设计一部分人往下走做硬件适配、做驱动、做极端性能优化中间那部分做重复性工作的会被压缩。第二AI写算子的进展会比很多人想象的快但比很多人想象的慢。快是因为简单算子已经基本被攻克了慢是因为复杂算子的性能优化涉及太多隐性知识这些知识很难被数据化。第三未来最值钱的算子工程师是那些能定义问题的人。AI能解决问题但定义什么问题值得解决、用什么标准衡量解决得好不好这些还是人的事。第四别把宝押在单一技术上。CUDA、Triton、TVM、各家加速器平台都要懂一点。技术栈越宽抗风险能力越强。刘胜的告别不是终点是一个信号。它告诉我们抽象阶梯又往上爬了一层而我们要做的是跟着往上走而不是站在原地感慨。我在实际带团队的过程中发现那些主动拥抱AI工具、同时坚持深入理解硬件的工程师反而是成长最快的。他们没有被AI替代而是被AI放大了。这大概就是抽象阶梯往上爬时个体最好的应对方式——不是对抗是借力。最后分享一个我一直在用的小习惯每学一个新算子我都会先让AI生成一版然后自己手写一版对比两者的差异。差异最大的地方往往就是最值得深入理解的地方。这个习惯帮我省了很多时间也帮我保持了对硬件的敏感度。