Blog Tags GitHub
All tags

efficient-attention

98 posts tagged with "efficient-attention"

Optimal Software Pipelining and Warp Specialization for Tensor Core GPUs（Twill：面向 Tensor Core GPU 的最优软件流水线与 Warp 特化）

arXiv 2512.18134，来自 CudaDMA/Singe 同一血统的团队（Bauer/Aiken + Stanford/NVIDIA/Toronto）。首次把软件流水线（Software Pipelining, SWP）与 warp 特化（Warp Specialization, WS）表述为一个可交给现成约束求解器求解的联合优化问题。核心洞见：SWP 与 WS 不应分开处理——WS 恰恰是实现 Tensor Core GPU 上 SWP schedule 所必需的代码生成手段（解决多 warp 协作发射 TC、寄存器工作集、变长延迟操作、阻塞同步中断四大难题）。方法上先用模调度（modulo scheduling，ℤLP 最优求解）得到初始 initiation interval I 与 modulo schedule M，再将 M/I 展开为一段直线程序 Q，在其上叠加 SWP 约束（Uniqueness/Consistency/Completion/Dependence/Capacity）、内存感知约束（活跃变量 live 传播 + 内存容量/SSA）、warp 分配约束（Warp Uniqueness/Variable Latency/Register Limit/Cross-Warp Spills/Concurrency），交给 SMT 求解器（Yices2 QFLIA）联合求解得到 M*、I*、A*。关键实现技巧：成本归一化（cycle count 比值保持的 ℤLP，令问题可解）、变长流式操作（TMA load）零延迟建模 + 深度作为 autotuning 参数、不可满足时单调增大 I/L 重试。系统 Twill 从 Triton TTGIR 抽取依赖图，无启发式、易扩展新架构、对单层循环保证最优。评估在 H100(Hopper)/B200(Blackwell) 上对 Flash Attention 前向/后向传播自动重新发现 FA3（ping-pong scheduling）与 FA4（三组 warp 策略）专家手工方案，性能达 cuDNN/FA 手工实现的 1%~2% 以内。

gpu-kernel hpc efficient-attention warp-specialization code-generation

Veda: Scalable Video Diffusion via Distilled Sparse Attention（蒸馏式稀疏注意力）

面向大规模视频扩散 Transformer（DiT）的稀疏注意力加速框架。核心洞见：生成质量并非由稀疏率本身决定，而取决于稀疏掩码与全注意力 tile 级几何结构的对齐程度。Veda 把 tile 选择建模为对全注意力的显式重建问题：(1) 蒸馏式 tile 打分——用轻量估计器从全注意力 backbone 蒸馏 tile 级分数，配合 TripPool（Avg/Max/Min 三元统计）与逐头 MLP 投影降低估计误差，训练时对 backbone 特征做 stop-gradient 解耦掩码学习与特征学习；(2) 逐头 tiling 搜索（Head-aware Tiling）——为每层每头分配 (p_t,p_h,p_w) 分块配置以最小化稀疏输出与全注意力输出的 Frobenius 误差；(3) 硬件高效 tile-skipping kernel（ThunderKittens/Hopper TMA+Warp Specialization，达 FA3 80% MFU）。在 Waver-T2V-12B 720P 10 秒视频上取得 5.1× 端到端、10.5× 自注意力加速，注意力开销从 92% 降至 50%，且序列越长收益越大。

diffusion-model video-generation llm-inference efficient-attention sparse-attention knowledge-distillation

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×）。

gpu-kernel hpc efficient-attention warp-specialization compiler

CompactAttention: Accelerating Chunked Prefill with Block-Union KV Selection

面向 chunked prefill 的稀疏注意力机制，将 2D 块稀疏掩码视为「KV 选择信号」而非「稀疏内核执行计划」，通过 Q-block union + intra-group union 构造 GQA 感知的 per-group KV 块表，实现零拷贝分页注意力执行，在 LLaMA-3.1-8B 上 128K 上下文取得 2.72x 注意力加速且精度接近 dense

sparse-attention efficient-attention kv-cache long-context llm-inference inference-serving

DashAttention: Differentiable and Adaptive Sparse Hierarchical Attention

用可微、自适应稀疏的 α-entmax 变换替代分层稀疏注意力中的 top-k 路由，让第一阶段按 query 自适应选择可变数量的 KV 块并为第二阶段 softmax 提供 prior，使整个层级端到端可微且非弥散（non-dispersive）。75% 稀疏度下精度媲美 full attention，Pareto 前沿优于 NSA/InfLLMv2，推理较 FlashAttention-3 最高 3.36x 加速

sparse-attention efficient-attention long-context llm-training kv-cache