1. 十年坐标：从手写CUDA到块级DSL，算子开发方式的三次换挡
1.1 CUDA时代：kernel是一门手艺，性能靠"经验积累"慢慢喂
十年前你要让我写一个高性能GPU算子，流程基本是这样的：打开CUDA，先把grid、block、thread想清楚，再手动安排每个线程处理哪些数据，然后绞尽脑汁处理shared memory、做同步、避免bank conflict、尽量让访存合并。一个稍微复杂的softmax或者矩阵乘，从思路成型到性能达标，往往要折腾好几天。
那个年代，算子开发非常依赖个人经验。不是说CUDA写不出来，而是"写出来"和"写快"之间隔着一条巨大的鸿沟。我记得当年调一个elementwise的kernel，为了把global memory的读写换成vectorized load，把float4类型安排得明明白白，性能提升了将近一倍。这种优化手段，说穿了就是你是否熟悉硬件的行为细节。
如果把那个时期的算子开发方式做个总结，我会说：程序员被迫站在硬件的角度看问题。你要想的不是"这个计算是什么"，而是"这个计算在硬件上该怎么被切分、怎么被搬运、怎么被隐藏延迟"。GPU的能力很强，但驾驭它的门槛也很高，这也是为什么很长一段时间里，能写CUDA kernel的人在团队里都算稀缺资源。
1.2 DSL与编译器尝试：声明式描述一度很热闹，但落地远没想象中顺利
大概是五六年前，行业里冒出一波"用编译器替代手工优化"的尝试。Halide提出把计算描述和调度策略分开，TVM想把神经网络的算子统一到一个IR里做自动优化，XLA则在TensorFlow里直接干掉手动写kernel的诉求。这套思路的逻辑很诱人：你只要描述清楚"算什么"，剩下的"怎么算"交给编译器去搜索和生成。
我确实用过TVM写过几个算子，也参加过XLA相关的技术分享。坦白说，方向上是对的，但它有一个绕不开的问题：抽象层级选得很尴尬。TVM的调度语法虽然比CUDA友好，但你仍然要懂"绑定了线程、向量化、流水线"这些底层概念，而且要花不少时间去学它的调度原语。等于说，门槛降低了一点，但学习路径换了一套，没解决根本问题。
更麻烦的是，这种"完整图级别"的编译器一旦塌缩到某几个异构硬件，代码生成的效率很难追得上资深工程师手写CUDA。你调优半天，可能还是差那么几个百分点。于是业内慢慢形成一个共识：大规模、静态shape的模型可以靠编译器，但灵活的动态shape、特殊算子，还是得有人能手动写点东西。所以那几年的实际状态是——编译器很热闹，但真正全自动替代手工的场景少之又少。
1.3 Triton作为转折点：像写NumPy一样写算子，编译器接管硬件细节
Triton在这条时间线上的位置其实非常有意思。它既不是传统意义上"从头写CUDA"，也不是TVM那套"先描述再调度"的复杂语法。OpenAI在做它的时候，选了一个更聪明的中间粒度：块级编程。你写的还是普通的Python函数，只不过在里面用triton.language提供的一系列原语来操作"块"级别的数据。
我第一次打开Triton文档的时候印象很深：一个softmax kernel，代码量也就三四十行，读起来几乎和伪代码一样清晰。你不用去管线程id是多少，也不用自己算stride，只需要声明"我当前这个program要处理X个元素"，然后像NumPy一样做load、compute、store。剩下的所有东西，包括线程怎么组织、访存怎么合并、寄存器怎么分配，都是编译器帮你完成的。
更关键的是，Triton的抽象精度卡在了一个很舒服的位置。它放弃了"全图自动优化"这种宏大目标，专注把单个kernel做到足够好，然后和PyTorch生态打成一片。等到PyTorch 2.0把Triton作为torch.compile的默认后端之一，大量的研究者忽然发现，自己用Python写的高性能算子竟然真的能跑出接近手写CUDA的效果。从这一刻起，算子开发的方式才算真正换了一代。
2. Triton核心抽象：为什么"块"是人与硬件都能接受的编程粒度
2.1 人和硬件对算子的需求完全不同
要理解Triton的设计，得先搞清楚一个矛盾。从人的角度看，我希望写算子的时候只关注"数据长什么样、怎么变形、怎么规约"，最好连指针和线程都不出现；从硬件角度看，GPU就是一台极度依赖并行和局部性的机器，它需要知道：每一块数据放在哪、由多少线程一起算、中间变量是不是能放进shared memory。
这两个需求天然冲突。你给程序员太多自由，就容易回到CUDA的复杂度；你给编译器太多自由，又容易变成TVM那样需要另一套领域知识才能驾驭。Triton给出的答案是：别让程序员管线程，但让程序员管"块"。一个块包含多个元素，程序员站在块的视角写计算，编译器则在块内部完成最终到线程的映射。
举个生活化的例子：CUDA像是直接给你一张城市地图，让你自己规划从A到B怎么避开拥堵；Triton则像是告诉你"你去这一片区域，把那里的快递都处理了"，至于你开车走哪条路、中间在哪个路口停一下，由导航系统替你决定。对于大部分算子开发者来说，后者显然更省心。
2.2 一个简单的softmax kernel，看Triton如何表达块计算
我直接贴一个最经典的softmax实现，你看一眼就明白我在说什么：
import torch import triton import triton.language as tl @triton.jit def softmax_kernel( output_ptr, input_ptr, input_row_stride, output_row_stride, n_rows, n_cols, BLOCK_SIZE: tl.constexpr, ): row_start = tl.program_id(0) col_offsets = tl.arange(0, BLOCK_SIZE) input_ptrs = input_ptr + row_start * input_row_stride + col_offsets mask = col_offsets < n_cols row = tl.load(input_ptrs, mask=mask, other=-float('inf')) row_minus_max = row - tl.max(row, axis=0) numerator = tl.exp(row_minus_max) denominator = tl.sum(numerator, axis=0) softmax_output = numerator / denominator output_ptrs = output_ptr + row_start * output_row_stride + col_offsets tl.store(output_ptrs, softmax_output, mask=mask)这段代码里没有threadIdx.x，没有blockIdx.x，没有shared memory，也没有同步屏障。你只需要知道三件事：每个program处理一行；tl.arange生成列的偏移量；tl.max和tl.sum沿着指定的轴做规约。编译器会自动把这一行数据切到足够多的线程上，必要时把中间结果塞进shared memory做规约，然后再写回去。
我第一次跑这个kernel的时候最惊讶的是，性能居然完全能打。对比同样配置下手写CUDA，Triton版本通常只慢个百分之几，有些情况下甚至因为自动向量化做得好而反超。这就是块级抽象的价值：你保留了"这一个块里面的数据我要怎么变换"的控制权，却把"块内怎么组织执行"交给了编译器。
2.3 编译器在块之上做的事：访存合并、shared memory与调度
有人会问：编译器凭什么能做得比人更好？答案是它在一个非常狭窄但关键的范围内做deep optimization。Triton的编译器拿到块级操作后，会分析每一段访存的shape、步长、对齐方式，然后自动生成"合并访问"模式，尽量让相邻线程访问相邻地址，提高global memory带宽利用率。这一点和手写CUDA里最重要的优化思路完全对齐，只不过由编译器推理完成。
同理，当你做softmax这种需要跨元素规约的操作时，编译器会识别出"块内归约"模式，自动插入shared memory分配和同步逻辑。你在Python代码里一行同步都没写，但生成的PTX里该有的bar.sync一条不少。这种能力背后就是LLVM/MLIR那一整套现代编译器基础设施，它会在Triton IR层"自我发现"哪些块操作可以映射成硬件原语，哪些必须拆成更小的步骤。
这其实就是很多人提到的"算子自发现"：不是人去找算子的可优化机会，而是编译器在IR层面系统性地扫描访存模式、计算模式、复用距离，然后自动套用优化规则。Triton的聪明之处在于，它把搜索空间限制在块内，避免了TVM那样在全图上暴力搜索调度方案，所以编译速度快、生成代码质量也稳定。
3. 一个Triton算子从Python到GPU执行的全流程
3.1 jit装饰与AST捕获：Python函数如何变成编译器IR
很多人第一次看Triton，会以为@triton.jit装饰过的函数还是普通Python函数。实际上，当你用指定的指针参数去调用它时，Triton并不会真的逐行执行Python，而是捕获这个函数的AST，做类型推导，然后翻译到自己的中间表示Triton IR。
这个过程有点像"把Python语法糖拆掉，还原成一个数据流图"。循环、if、算术操作都会被转成静态的、带形状信息的中间表示。副作用就是Python里很多动态特性（比如任意对象方法、闭包捕获复杂对象）在kernel里是不能用的。能用的是tl.arange、tl.load、tl.store这些建立在形状推断基础上的核心原语，以及普通整数、浮点运算。
当初写的时候我还不习惯，老想在kernel里调用Python库函数，结果编译直接报错。后来才明白，Triton要求你以"块"为单位思考，你写的是一段会被编译器整体分析和变换的代码，而不是一行行解释执行的脚本。这个心智模型转过来之后，很多东西就顺了。
3.2 Triton IR到LLVM/PTX：自动向量化、寄存器分配与代码生成
AST捕获完，接下来的链路是：Triton IR -> MLIR/LLVM方言 -> LLVM IR -> PTX -> SASS（最终由NVIDIA驱动做JIT编译）。很多人把Triton简单理解成"Python到PTX的翻译器"，其实中间的优化多得很。
在Triton IR层，编译器会根据块的形状和访存对齐信息做向量化。如果你处理的是连续的float数据，它会优先生成128-bit的向量load；如果对齐条件不满足，则回退到更窄的访问。这一层的决策直接决定你最终能不能吃满显存带宽。接着是shared memory的liveness分析，编译器会安排中间缓冲区，减少global memory的反复读写。
到了LLVM阶段，常规的循环不变量外提、指令合并、寄存器分配全部生效。这也是为什么Triton生成的kernel性能往往不弱于手写CUDA的核心原因之一：它同时吃到了"高层块级语义带来的精准优化机会"和"底层LLVM十几年打磨的优化能力"。我做过一次实验，同一个GELU算子，Triton编译出来的PTX里面居然有自动生成的__nv_bfloat16类型转换指令，比我手写CUDA时更懂得利用硬件特性。
3.3 launch与硬件调度：从grid循环到SM上的执行器
代码生成完毕，剩下的就是launch了。Triton在Python端拿到你传入的grid大小，为每个program instance生成对应的起始地址和参数组合。GPU的调度器会把这些block按顺序分发到流多处理器上，每个SM内部再按照warp调度器把线程束放行。
这里有个容易忽略的点：当你的grid特别大时，一个program instance并不是独享一整块连续的执行时间。SM会在空闲时插入别的block，形成一种硬件级别的overlap。Triton的编译器其实无法直接控制这个层面，但它会尽量让每个block内部有足够的独立工作量，用更多并行度来掩盖访存延迟。这跟手写CUDA的道理一样，只不过你在Triton侧通过调整num_warps间接影响block的线程规模，而不是手写一切。
配合torch.cuda上的launch机制，整个过程在Python侧看起来就是一个普通的函数调用。你要是只关心功能，完全不用管中间发生了什么；但如果你后面要排查性能问题，就必须要能顺着"Python wrapper -> Triton kernel -> grid设置 -> 实际SM执行"这条链路去定位瓶颈。
3.4 验证与性能分析：别让"能跑"骗了你
跑通一个Triton算子不算难事，难的是确认它真的对。我见过很多新手在对比kernel输出和PyTorch参考实现时直接assert，结果遇到float精度问题，一比对就崩，然后怀疑Triton代码写错了。正确做法是先用torch.allclose设定合适的atol和rtol，特别是涉及exp、reduce这类容易累积误差的操作。
性能分析也不要只盯time.time()。Triton自带的triton.testing.do_bench比手动计时靠谱得多，它会做warmup，自动算平均和标准差。更进一步，你可以在Nsight Compute里看到Triton生成的kernel名字，然后逐个看访存吞吐、计算吞吐和warp occupancy。我第一次用Nsight看Triton的kernel时还特意找了一下名字，确认它确实是被正常编译和执行的，而不是走了什么magic路径。事实证明，它就是普通kernel，只是生成过程自动化了而已。
4. 大量算子叠加，真正挑战从来不是单个kernel
4.1 单算子优化和整图性能之间的鸿沟
很多做算子开发的人一开始都会陷入一个误区：把单个kernel优化到极致，模型整体性能就该很好了。但当你真正去分析一个PyTorch模型时，会发现里面可能有几百个算子。其中很多算子本身很简单，比如加一个bias、过一个激活、做一次layout转换，单独看每个也就几十微秒，但架不住数量多。
每个kernel的launch overhead、参数传递、两次kernel之间的全局内存写回，都会成为性能的黑洞。更麻烦的是，kernel和kernel之间的数据来回倒腾，往往会把本来能驻留在寄存器或shared memory里的数据逼到global memory里走一遭，白花带宽。这其实是"大量使用算子对硬件性能的挑战"最本质的部分：你优化的不是数学运算，而是数据搬运的路径。
我看过一个很典型的模型trace，里面光elementwise类型的算子就有三十多个。如果每个都单独跑，GPU的算力利用率可能连10%都不到。这就是为什么无论编译器怎么演进，算子融合都是绕不开的话题。不做融合，堆再多的优化技巧都是在给搬运工加工资。
4.2 访存瓶颈与算子融合：为什么matmul+prelu这类融合成为标配
当模型里出现大算子和小算子交替时，最划算的做法就是让小算子"贴"在大算子身上，让中间结果留在片上。一个非常典型的例子就是matmul + prelu：矩阵乘法的输出如果先写回global memory，再由下一个kernel读出来做prelu，等于多了一整轮DDR流量；如果直接在矩阵乘的epilogue阶段把prelu应用在寄存器里的结果上，访存开销几乎可以忽略。
这类融合算子之所以越来越流行，是因为它不仅出现在GPU上，在NPU上也一样重要。我查过一些昇腾上的算子优化案例，AscendC里做matmul+prelu融合的思路，和CUDA上做FMA融合本质上没有区别：把矩阵乘的累加结果留在AI Core内部，紧接着做激活，最后一次性写回。换的是硬件名字，不变的是"减少片外搬运"这条原则。
在Triton里做这种融合也非常自然。你可以在load完一个tile、做完tile级别的矩阵累加后，直接对寄存器中的结果做tl.where之类的激活操作，然后再store。编译器会帮你把中间状态锁定在片上，不会额外分配大块global memory。我写过一个带LeakyReLU的matmul，融合后整体算子比分开跑快了将近30%，而且代码只多了三行。
4.3 Triton实战调优：num_warps、num_stages和autotune的取舍
到了调优环节，Triton提供了几个核心旋钮：num_warps、num_stages、以及块大小。很多人一上来就照着默认值跑，性能不太行就放弃了，其实这几个参数对性能影响非常大。
num_warps控制一个program内部的线程束数量，直接影响并行度。块太小、元素太少时，开太多warp反而会因为同步开销增大而变慢；块很大、计算很重时，warp太少又不足以掩盖访存延迟。num_stages用于软件流水线：在load下一块数据的同时计算当前块，用shared memory做缓冲，隐藏global memory访问延迟。这个参数对访存密集算子尤其关键。
我通常的习惯是把关键算子交给triton.autotune，让它在少数几组配置里自动搜索。别搞几百组，冷启动会怀疑人生；4到8组配置足够了。我实测下来，num_warps=4配合num_stages=3对很多elementwise和softmax算子都是不错的起点，而matmul类更偏向更大的块和更深的流水线。
4.4 编译时间、缓存与动态shape：工程化中最容易翻车的地方
Triton虽然运行快，但编译一个kernel不是完全零成本的。每次你改变launch配置、块大小、数据类型、甚至某些tl.constexpr参数的值，编译器都可能要重新走一遍优化链路。在训练循环里频繁换shape，最常见的结果就是"卡顿一下然后速度正常"，那一卡就是编译的开销。
Triton本身有kernel缓存，默认会按源码、参数、后端等信息做哈希，落到磁盘上，所以同样配置重复执行不会重复编译。但动态shape会让缓存命中率直线下降。如果某个维度在训练中频繁变化，我建议干脆把它当成BLOCK_SIZE的constexpr处理时设置一个覆盖所有shape的公共块大小，避免为每一个shape都生成一份kernel；或者用padding把形状对齐到同一档位，收益远大于浪费的那点计算。
还有个小坑是环境里的TRITON_CACHE_DIR，默认目录空间如果被占满，编译可能会失败。离线部署场景我一般会显式指定一个可写的缓存路径，并在镜像里预编译一遍关键算子，把缓存一起打包进去，这样推理服务器起来后就不用现场编译了。
5. 从CUDA到NPU：Triton如何应对多硬件时代的算子开发
5.1 NPU带来的新变量：AI Core、带宽层级与调度方式
过去做高性能算子开发，大家默认目标就是NVIDIA GPU。但近两年NPU的声音越来越大，尤其是昇腾这类AI芯片，核心计算单元叫AI Core，很多设计逻辑和GPU的SM类似，但又不完全一样。你在CUDA上习惯的那些假设——比如统一的shared memory大小、固定的warp宽度、L2缓存的分配方式——在NPU上可能全部要重设。
NPU上算子开发的难点不完全在"算力不够"，而在搬运和调度的自由度不同。有些NPU对局部内存的容量和Bank组织有自己的讲究，对多核任务切分也有一套固定的同步原语。直接在AscendC里写算子，有点像当年手写CUDA：性能天花板高，但开发链路长、调试工具相对少。这时候Triton这类高级抽象就有机会了，它把"块级计算"变成一种硬件无关的描述，再由不同的后端映射到具体架构。
5.2 MLIR/LLVM作为中间桥梁：Triton IR映射到AscendC等后端
Triton能被多个硬件后端盯上，核心原因在于它的IR设计比较干净，没有过早绑定PTX。社区里其实已经有很多人在做Triton到非NVIDIA后端的适配，基本路径是：先复用Python前端和Triton IR，然后把块级操作lower到MLIR方言，再由各硬件厂商自己的编译器后端继续做调度和代码生成。
对应到昇腾生态，你会看到两条路线：一条是直接用AscendC手写融合算子，面向量产和极致性能；另一条是把Triton kernel翻译成类似的IR，再映射到AscendC的C代码结构上。对我来说，Triton这条路线更大的价值在于快速原型验证：我先用Triton把算法逻辑写清楚，跑通shape和数值，再决定要不要花力气用AscendC手工精调。
当然，能编译不等于高效。不同NPU的向量单元宽度、矩阵单元形态、内存层级差异很大，Triton的自动优化规则在GPU上学到的那套经验，直接搬过去不一定合适。所以多后端适配的现实状态是：前端通用，后端逐硬件精调。这也是为什么各家在做落地时，都不会只依赖Triton一条链，而是把Triton当成"易于生成、易于验证"的入口。
5.3 多后端适配的现状与坑：能编译不等于高效
真实踩过坑的人会告诉你，跑通NPU后端和跑出高性能之间隔着一个巨大的"调优深渊"。跨后端之后，原本在GPU上不需要关心的细节会突然冒出来：NPU的core之间通信方式可能不是全局同步，block的切分方式会影响片上数据复用，甚至不同shape的padding策略都会导致计算效率翻倍变化。
另外，Triton社区的PR和更新基本还是以NVIDIA backend为主，非NVIDIA后端的维护活跃度、bug修复速度都没法和CUDA路径比。如果你所在团队要长期做某款NPU的算子库，我建议不要把整个技术栈押在社区维护的第三方后端上，而是把它作为辅助验证工具，核心算子还是要有自己的性能回退方案。
从更长远的视角看，算子开发正在走向"一套块级语义，多个硬件后端"的格局。将来一个模型要在GPU、NPU、甚至CPU上跑，开发者需要的不是为每个硬件单独写一遍算子，而是把算子的逻辑描述好，然后让编译器针对目标硬件生成最合适的执行方案。Triton在这方面至少开了一个很实际的头。
6. 实验记录：搭建Triton环境、跑通第一个算子并避坑
6.1 安装与版本匹配：pip一条龙？没那么简单
很多人上来就pip install triton，装完跑demo，大概率会撞上版本和PyTorch对不上的问题。Triton和PyTorch的CUDA runtime走得比较近，版本错位时经常出现Illegal instruction或者奇怪的CUDA error。
我现在的做法是：优先使用PyTorch自带匹配的Triton，而不是自己额外装一份最新的。如果你用pip安装PyTorch官方wheel，它通常已经捆绑了一个兼容版本的Triton；需要单独升级Triton时，再去triton的release页面找对应的wheel包。源码编译也是一种方式，但耗时较长，还要装好LLVM依赖，除非你想改编译器的源码做二次开发，否则不建议一上来就跳进源码编译的坑。
装完可以跑一个最简单的加法kernel验证环境。重点看它能不能正常生成PTX，以及triton.testing.do_bench能不能正常计时。如果能跑通，说明CUDA环境、Python版本、编译链都OK了。
6.2 从vector add到softmax：上手只需要两个晚上
学习Triton最快路径，我建议直接从vector add开始，别看太多入门教程。第一晚：写一个vector_add_kernel，传入两个指针和大小，用tl.arange和mask处理边界，跑通和PyTorch结果一致。第二晚：改写成softmax，加上tl.max和tl.sum做行规约，顺便试一下num_warps=4和num_warps=8的性能差异。
这两个例子足够覆盖80%的常用语法：load/store、mask、arange、reduce、constexpr。之后再接触矩阵乘法的tl.dot、原子操作的tl.atomic_add、多级流水的num_stages，都会非常快。我甚至见过一个完全没碰过CUDA的研究者，用两周时间把一套Transformer里的自定义算子全部换成Triton实现，跑出来的性能和原来手写CUDA版本基本持平。这放在十年前是不可想象的。
6.3 我踩过的三个坑：版本兼容、边界shape、自动调优冷启动
坑一，版本兼容。有一回我把Triton升到大版本，结果旧代码里tl.load的eviction_policy参数行为变了，编译不报错，但性能掉了20%。排查了半天才定位到是这个语义变化。所以升级Triton之后，一定要对着release note核对一遍你用了哪些高级参数，别默认它们是稳定的。
坑二，边界shape。block size和实际shape不一致时，mask必须有。但mask太多会影响访存效率，因为编译器没办法确定哪些地址一定能合并访问。我的经验是尽量让shape对齐到block size的整数倍，实在不行再开mask。比如做softmax时，把一行元素padding到2的幂，经常比用精确mask跑得更快。
坑三，自动调优冷启动。autotune配置太多组，程序第一次启动会逐个编译候选kernel，几十个组合下来，可能等了好几分钟才正式开始跑。这在本地开发还能忍，在线上服务或大规模评测时会非常致命。我现在都严格控制候选集数量，并把缓存目录稳定下来；如果benchmark的场景很固定，还不如直接手动选一组最合适的配置，把编译时间压到一次。
最后再分享一点我的个人感受：做算子开发这十年，最明显的变化是"写算子的门槛在快速下降"，但"判断一个好算子需要什么"的能力反而更值钱了。Triton帮你省掉了手工管理线程的体力活，可要写出能在不同硬件上都稳定高效的算子，你还是得理解访存、融合、调度这些底层逻辑。工具在迭代，基本功不会过时。如果你正准备入坑算子开发，我的建议就是别怕底层的那些概念，先上手写几个Triton算子，再回过头去看CUDA，你会发现很多曾经需要死记硬背的优化细节，现在都变得顺理成章了。
