PyTorch 携手 Google 发布 Helion TPU 后端:硬件异构核编译新范式
更新背景
随着 Google TPU v7 (Ironwood) 的推出,其性能已媲美 NVIDIA B200,成为大规模训练和推理的重要选择。然而,TPU 的编程模型与 GPU 存在显著差异:TPU 采用顺序执行架构,依赖显式的内存层次结构(HBM 与片上 VMEM),而 GPU 则依赖大规模并行 SIMT 和隐式缓存。
传统上,编写高性能 TPU 内核需要深厚的 Pallas 底层 DSL 知识,学习曲线陡峭且代码复杂。为了降低这一门槛,PyTorch 推出了 Helion——一种用于编写性能可移植机器学习内核的高级领域特定语言 (DSL)。此次更新标志着 Helion 正式集成 TPU 后端,实现了从 PyTorch 风格代码到优化 TPU 代码的无缝编译。
核心突破与功能特性
1. PyTorch 风格的 TPU 内核开发
Helion 允许开发者使用直观的 PyTorch 语法定义内核,无需深入理解 Pallas 的细节。例如,一个简单的加法操作在 Helion 中仅需几行代码:
@helion.kernel
def add(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
out = torch.empty_like(x)
for tile in hl.tile(out.size()):
out[tile] = x[tile] + y[tile]
return out
编译器会自动将其转换为包含宿主端启动器(Launcher)和设备端函数的高效 Pallas 代码,自动处理数据在 HBM 与 VMEM 之间的传输与计算重叠。
2. 智能自动调优 (Autotuning)
Helion 内置强大的自动调优器,能够针对不同的输入形状探索多种代码生成策略。它自动选择最佳的流水线缓冲区大小 (_BLOCK_SIZE_0) 和调度方案,以最大化 TPU 的向量内存 (VMEM) 利用率,确保内存传输与浮点计算完美重叠。
3. 卓越的硬件性能
在 Flash Attention 工作负载上,Helion 生成的内核在 TPU v7 上实现了 838 TFLOPS 的峰值性能,达到单个张量核心 (Tensor Core) 最大浮点利用率 (MFU) 的 79%。这一性能水平证明了 Helion 在硬件异构环境下的编译效率。
实际应用价值
- 降低开发门槛:非 Pallas 专家也能快速上手 TPU 内核开发,加速 TPU 生态的普及。
- 提升跨硬件兼容性:开发者可以维护一套统一的 Helion 内核代码,同时适配 TPU 和 GPU,简化工程维护。
- 最大化硬件利用率:通过自动流水线优化,确保在 TPU v7 等新一代加速器上获得接近理论极限的性能。
"Helion 正在演变为编写 TPU 内核的诱人选项,它 bridged the gap between high-level portability and low-level hardware optimization."
关键亮点 (Key Highlights):
- Helion TPU Backend: 首个将 PyTorch 风格 DSL 编译为 Pallas 代码的官方后端,支持 TPU v7。
- 智能 Autotuning: 自动探索代码生成策略与流水线参数,针对特定输入形状优化性能。
- 高性能验证: 在 Flash Attention 任务中达到 838 TFLOPS,MFU 利用率高达 79%。
- 简化开发流程: 消除对 Pallas 底层知识的依赖,让开发者专注于算法逻辑。
关键指标 (Metrics):
版本: Helion TPU v1.0性能: 838 TFLOPS (Flash Attention 工作负载)MFU: ~79% (单张量核心)目标硬件: TPU v7 (Ironwood)