Tawa: Automatic Warp Specialization for Modern GPUs with Asynchronous References(Tawa:面向现代 GPU 的异步引用自动 Warp 特化编译器)
CGO 2026 论文,Cornell + NVIDIA 合作。提出 Tawa——首个面向现代 NVIDIA GPU(Hopper+)的**全自动 warp 特化编译流程**,从未经修改、无标注的 Triton 程序出发,自动分区为 producer/consumer warp group,自动管理软件流水线。核心创新是**异步引用(asynchronous reference, aref)**——一种形式化语义保证的 IR 抽象,将 warp 间通信建模为带硬件 mbarrier 的单槽信道(put/get/consumed),封装 TMA 描述符配置、mbarrier 初始化/相位管理等底层细节。编译器三步变换:① 基于语义标记的任务感知分区(iteration stmt → producer, tile stmt → consumer,依赖闭包 + 计算复制);② 基于 aref 的多槽环形缓冲自动构建跨 warp 数据流;③ 多级软件流水线(细粒度:CUDA 地址计算与 MMA 重叠;粗粒度:CUDA Core 变换与 TC 流水线装配线)。额外优化:协作计算 warp group(多 WG 共享 tile 寄存器预算)、持久化 kernel。在 H100 上:GEMM 达 cuBLAS 水平(FP16 +1.01×,FP8 +1.06×);Attention 达 FA3 手工优化内核的 96%,超 Triton 1.21×。ablation 显示 +Auto WS 提升 3.78×,最终 GEMM 达 718 TFLOPs/s(baseline 的 ~7×)。