Helion brings PyTorch-style kernel authoring to TPUs with 838 TFLOPs flash attention

Helion, PyTorch’s domain-specific language for writing performance-portable ML kernels, has partnered with Google to build a TPU backend that compiles Helion kernels to Pallas, Google’s low-level TPU DSL. On a flash attention workload, the Helion-generated kernel achieves 838 TFLOPs—roughly 79% of the theoretical peak of one TPU v7 tensor core. The compiler also autotunes over different code-generation strategies (e.g., outer-loop pipelining via pallas_call, inner-loop choices between emit_pipeline and full unrolling) to select the optimal pipelining scheme for each input shape, maximizing utilization of TPU’s explicit memory hierarchy (HBM→VMEM). Helion TPU targets three key use cases: performance-critical workloads that benefit from autotuning, teams new to TPU kernel development who want to avoid Pallas’s steep learning curve, and cross-hardware users who prefer maintaining a single kernel codebase across TPUs and GPUs.

This development lowers the barrier for deploying on Google’s TPU v7 (Ironwood), which offers comparable BF16 compute and HBM bandwidth to NVIDIA B200 at potentially lower TCO. By letting PyTorch developers write kernels using familiar constructs (e.g., hl.tile) and automatically generating pipelined Pallas code, Helion makes TPU acceleration accessible without deep hardware expertise. The autotuner further ensures that kernels remain efficient across varying problem sizes, addressing a common pain point in manual kernel optimization. As the ML industry shifts toward heterogeneous hardware, Helion’s TPU backend provides a practical path to portable, high-performance code.