arXiv:2410.20399cs.LGcs.AI2024-10被引 45

用简单抽象框架实现高效AI核函数,性能超越现有方案。

ThunderKittens: Simple, Fast, and Adorable AI Kernels

  • 用三层次抽象简化GPU核函数编写:瓦片、线程块、网格层级分别优化计算
  • 在GEMM和注意力推理上达到CuBLAS与FlashAttention-3水平,部分任务快14倍
  • 适合追求高性能又不想手写复杂内核的开发者使用

将AI架构映射到GPU硬件面临严重瓶颈。尽管投入大量精力,手写自定义核函数仍难以达到理论性能,即使在成熟操作如线性注意力上也是如此。我们探索是否可用少量关键抽象显著简化过程。提出ThunderKittens(TK)框架,在保持易用性和可维护性的同时实现高性能核函数。其抽象覆盖GPU三层结构:(1)瓦片级提供16×16矩阵块与类似PyTorch的并行运算;(2)线程块级支持跨瓦片异步操作重叠;(3)网格级隐藏块启动/销毁与内存开销。实验表明,所生成核函数在多种AI操作中匹配或超越已有方案:在GEMM与注意力推理上等效于CuBLAS与FlashAttention-3,注意力反向传播快10–40%,状态空间模型快8倍,线性注意力快14倍。

原文摘要 · Abstract (English)

The challenge of mapping AI architectures to GPU hardware is creating a critical bottleneck in AI progress. Despite substantial efforts, hand-written custom kernels fail to meet their theoretical performance thresholds, even on well-established operations like linear attention. The diverse hardware capabilities of GPUs might suggest that we need a wide variety of techniques to achieve high performance. However, our work explores whether a small number of key abstractions can drastically simplify the process. We present ThunderKittens (TK), a framework for writing performant AI kernels while remaining easy to use and maintain. Our abstractions map to the three levels of the GPU hierarchy: (1) at the warp-level, we provide 16x16 matrix tiles as basic data structures and PyTorch-like parallel compute operations over tiles, (2) at the thread-block level, we provide a template for overlapping asynchronous operations across parallel warps, and (3) at the grid-level, we provide support to help hide the block launch and tear-down, and memory costs. We show the value of TK by providing kernels that match or outperform prior kernels for a range of AI operations. We match CuBLAS and FlashAttention-3 on GEMM and attention inference performance and outperform the strongest baselines by $10-40\%$ on attention backwards, $8\times$ on state space models, and $14\times$ on linear attention.

GPU优化AI核函数高性能计算

Thank you to arXiv for use of its open access interoperability. PaperDance 不是 arXiv 官方产品;中文卡片由大模型生成,请以原文为准。