Warp Specialization（线程束专业化）是一种 GPU 内核编程策略：在同一个 kernel（核函数）里，让不同的 warp 固定扮演不同角色，有的专门搬数据，有的专门计算，有的专门做通信，用流水线把访存延迟藏起来。
GPU 的调度单位是 warp（线程束），通常包含 32 个线程。传统 SIMT（Single Instruction, Multiple Threads）模型里，所有 warp 执行同一段代码，只是处理不同数据。问题在于，全局内存访问很慢；如果计算 warp 亲自去取数据，取数时计算单元就空转。Warp Specialization 把“取”和“算”拆开：生产者（producer）warp 用异步拷贝或 TMA（Tensor Memory Accelerator）把数据从全局内存搬到共享内存（shared memory），消费者（consumer）warp 从共享内存取数计算，再用 mbarrier 等同步机制通知“数据已就绪”。当消费者算第 n 块时，生产者已经在预取第 n+1 块，延迟被重叠。
打个厨房比方：传统做法是每个厨师自己跑仓库拿菜、切菜、炒菜、传菜，跑仓库时灶台空着。Warp Specialization 后，备菜员专门从仓库拉货到操作台，厨师只炒，传菜员只送。仓库很远，但厨师不用等仓库，因为下一道菜的料已经提前放在台面。
与相邻概念的区别：
| 概念 | 谁在分叉 | 效果 |
|---|---|---|
| Warp divergence（线程束发散） | 同一个 warp 内的线程 | 因 if/else 走不同分支，通常串行化，是性能损失 |
| Warp Specialization | 不同 warp 之间 | 每个 warp 内部仍统一，但角色不同，是刻意设计的性能优化 |
| CUDA stream | 主机提交多个 kernel | 跨 kernel 并发，不改变单个 kernel 内部 |
| 流水线并行 | 设备或层之间 | 宏观吞吐，warp specialization 是单 GPU 内核内的微观流水线 |
对 AI 从业者来说，写高性能 kernel 时要考虑角色划分、缓冲区双缓冲、异步屏障和资源占用。CUTLASS、Triton 等库或编译器会封装部分模式，但理解后更容易调 tile、stage、warp group。对普通职场人来说，这是“专门的人做专门的事，并用流水线提前一步”的工程直觉：别让最贵的人干等搬运。需要注意的是，它会增加同步复杂度和共享内存压力，不是所有 kernel 都值得。具体 API 和硬件支持以官方页面为准。