All tags

code-generation

3 posts tagged with "code-generation"

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% 以内。

Singe: Leveraging Warp Specialization for High Performance on GPUs(用 Warp 特化实现 GPU 高性能的 DSL 编译器)

PPoPP'14 论文,来自 CudaDMA 同一作者团队(Bauer/Treichler/Aiken, Stanford)。提出 Singe——一个面向燃烧化学(combustion chemistry)的 DSL 编译器,用 warp 特化(warp specialization)作为替代数据并行的编程模型,把一个 CTA 内的 warp 划分为不同子计算,通过硬件生产者-消费者命名屏障(named barriers, PTX bar.arrive/bar.sync)细粒度同步。核心解决燃烧 kernel 的三大难题:超大工作集(每点数百个双精度变量)、不规则计算(多阶段、部分可并行)、不规则访存。Singe 用 warp 特化把大工作集拆分到不同 warp 的片上寄存器/shared memory,避免寄存器溢出。三类 kernel 案例:Viscosity(按物种求和分块,Store 模式)、Diffusion(对 dij 矩阵按列分块,Mixed 模式)、Chemistry(反应速率存寄存器经 shared 交换,Buffer 模式 + QSSA DAG 分块 overlap)。编译器四阶段:分块→映射(FLOPS/寄存器/局部性三指标)→命名屏障放置调度(无死锁定理,16 个命名屏障=SSA 寄存器分配)→代码生成。代码生成三关键技术避免指令 cache 抖动:代码 overlay(同时遍历 AST 森林 + 位掩码/间接分支)、常量去重(Fermi shared 广播 / Kepler shuffle)、warp 索引(不规则访存)。在 Fermi C2070 与 Kepler K20c 上对 DME/Heptane 机制,相比已高度优化的数据并行 baseline 最高加速 3.75×。