斯坦福新工具“雷猫”：让 GPU 性能飙升，比 FlashAttention2 更快！ 2024-06-07 作者 C3P00 AI 算力资源越来越紧张，如何让 GPU 充分发挥潜力成为一大难题。最近，斯坦福大学的研究人员发布了一个名为“雷猫”（ThunderKittens）的工具，它可以显著提升 GPU 的运行效率，甚至比目前最快的 FlashAttention2 还要快 30%！ 雷猫的秘密武器在于它对 GPU 硬件的深度理解。研究人员从“硬件实际需要什么？如何满足这些需求？”这两个问题出发，设计了一个嵌入式 CUDA DSL 工具。雷猫通过操作小型张量块（tile）来简化 AI 内核的编写，并充分利用张量核心、异步数据传输和共享内存等硬件特性。 H100：性能之王，如何榨干它的潜力？ 研究人员以英伟达最新的 H100 GPU 为例，深入探讨了如何优化 GPU。H100 拥有 989 TFLOPs 的半精度矩阵乘法计算能力，但要充分发挥它的能力，关键是保持张量核心持续运算。 然而，要做到这一点并不容易。研究人员发现，H100 硬件具有一些特性，对于保持矩阵乘法的运行至关重要： WGMMA 指令： H100 引入了新的指令集 WGMMA，它允许 128 个线程跨 SM 所有子单元协作同步，并从共享内存及寄存器异步启动矩阵乘法。这些指令对于充分利用 H100 的计算能力是必不可少的，没有它们，GPU 的峰值利用率会损失 37%。 共享内存： 共享内存的延迟虽然看似不多，但对于高速运行的张量核心来说已经足够显著。此外，共享内存的存储单元独立，处理不当会导致 bank conflicts，显著拖慢内核速度。 地址生成： H100 的张量核心和内存速度极快，仅生成用于获取数据的内存地址就占用了芯片的大量资源。英伟达提供的张量内存加速器（TMA）可以异步提取张量的一部分，大大节省了地址生成的开销。 占用率： 占用率指的是 GPU 上同时调度的线程数。高占用率可以隐藏缺陷或同步问题，但一个设计良好的 pipeline 即使在占用率不高的情况下也能运行得相当快。 雷猫：小而精，性能大提升 为了更轻松地编写所需的内核类型，并充分发挥硬件的全部潜力，雷猫应运而生。它是一个嵌入在 CUDA 中的 DSL，包含四种模板类型：寄存器 tiles、寄存器向量、共享 tiles 和共享向量。雷猫提供了一系列操作来处理这些张量，既可在 warp 级别使用，也可用于多个 warp 协作。 雷猫的优势在于它的小而简单，功能有限，但它符合 AI 和硬件的发展趋势。研究人员认为，AI 的设计理念应该根据硬件特性进行重新定义，例如循环状态应该足够大以适应一个 SM，计算的密度应该不低于硬件的需求。 未来展望：AI 与硬件的深度融合 雷猫的出现，标志着 AI 和硬件深度融合的新方向。研究人员认为，未来需要利用对硬件的了解，设计与之匹配的 AI，从而进一步提升 GPU 的性能，释放 AI 的潜力。 参考文献： [1] https://hazyresearch.stanford.edu/blog/2024-05-12-tk [2] https://github.com/HazyResearch/ThunderKittens [3] https://news.ycombinator.com/item?id=40337936 斯坦福大学的研究人员最近发布了一个名为“雷猫”（ThunderKittens）的工具，它可以显著提升 GPU 的运行效率，比目前最快的 FlashAttention2 还要快30% [1]。雷猫是一个嵌入在 CUDA 中的 DSL 工具，通过操作小型张量块来简化 AI 内核的编写，并充分利用 GPU 硬件的特性，如张量核心、异步数据传输和共享内存等 [1]。 下面是关于雷猫和 GPU 性能提升的一些关键信息： 雷猫：GPU 性能提升的秘密武器 雷猫是一个嵌入在 CUDA 中的 DSL 工具，通过操作小型张量块来简化 AI 内核的编写，并充分利用 GPU 硬件的特性 [1]. 雷猫提供了四种模板类型：寄存器 tiles、寄存器向量、共享 tiles 和共享向量，用于处理张量操作 [1]. 雷猫的优势在于它的小而简单，功能有限，但符合 AI 和硬件的发展趋势 [1]. GPU 性能提升的关键因素 GPU 硬件具有一些特性，对于保持矩阵乘法的运行至关重要 [1]: WGMMA 指令：H100 GPU 引入了新的指令集 WGMMA，允许线程跨 SM 所有子单元协作同步，并从共享内存及寄存器异步启动矩阵乘法 [1]. 共享内存：共享内存的延迟对于高速运行的张量核心来说已经足够显著，处理不当会导致 bank conflicts，拖慢内核速度 [1]. 地址生成：GPU 的张量核心和内存速度极快，但生成用于获取数据的内存地址占用了大量资源，张量内存加速器（TMA）可以异步提取张量的一部分，节省了地址生成的开销 [1]. 占用率：高占用率可以隐藏缺陷或同步问题，但设计良好的 pipeline 即使在占用率不高的情况下也能运行得相当快 [1]. 未来展望：AI 与硬件的深度融合 雷猫的出现标志着 AI 和硬件深度融合的新方向，未来需要利用对硬件的了解，设计与之匹配的 AI，进一步提升 GPU 的性能，释放 AI 的潜力 [1]. 参考文献： 比标准Attention提速5-9倍，大模型都在用的FlashAttention v2来了 | 机器之心 [1] 斯坦福大学发布的ThunderKittens工具 [1] https://news.ycombinator.com/item?id=40337936 [1] Learn more: 比标准Attention提速5-9倍，大模型都在用的FlashAttention v2来了 | 机器之心 通透理解FlashAttention与FlashAttention2：全面降低显存读写、加快计算速度-CSDN博客 FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning | Princeton NLP Group
AI 算力资源越来越紧张，如何让 GPU 充分发挥潜力成为一大难题。最近，斯坦福大学的研究人员发布了一个名为“雷猫”（ThunderKittens）的工具，它可以显著提升 GPU 的运行效率，甚至比目前最快的 FlashAttention2 还要快 30%！
雷猫的秘密武器在于它对 GPU 硬件的深度理解。研究人员从“硬件实际需要什么？如何满足这些需求？”这两个问题出发，设计了一个嵌入式 CUDA DSL 工具。雷猫通过操作小型张量块（tile）来简化 AI 内核的编写，并充分利用张量核心、异步数据传输和共享内存等硬件特性。
H100：性能之王，如何榨干它的潜力？
研究人员以英伟达最新的 H100 GPU 为例，深入探讨了如何优化 GPU。H100 拥有 989 TFLOPs 的半精度矩阵乘法计算能力，但要充分发挥它的能力，关键是保持张量核心持续运算。
然而，要做到这一点并不容易。研究人员发现，H100 硬件具有一些特性，对于保持矩阵乘法的运行至关重要：
雷猫：小而精，性能大提升
为了更轻松地编写所需的内核类型，并充分发挥硬件的全部潜力，雷猫应运而生。它是一个嵌入在 CUDA 中的 DSL，包含四种模板类型：寄存器 tiles、寄存器向量、共享 tiles 和共享向量。雷猫提供了一系列操作来处理这些张量，既可在 warp 级别使用，也可用于多个 warp 协作。
雷猫的优势在于它的小而简单，功能有限，但它符合 AI 和硬件的发展趋势。研究人员认为，AI 的设计理念应该根据硬件特性进行重新定义，例如循环状态应该足够大以适应一个 SM，计算的密度应该不低于硬件的需求。
未来展望：AI 与硬件的深度融合
雷猫的出现，标志着 AI 和硬件深度融合的新方向。研究人员认为，未来需要利用对硬件的了解，设计与之匹配的 AI，从而进一步提升 GPU 的性能，释放 AI 的潜力。
参考文献：
[1] https://hazyresearch.stanford.edu/blog/2024-05-12-tk
[2] https://github.com/HazyResearch/ThunderKittens
[3] https://news.ycombinator.com/item?id=40337936
斯坦福大学的研究人员最近发布了一个名为“雷猫”（ThunderKittens）的工具，它可以显著提升 GPU 的运行效率，比目前最快的 FlashAttention2 还要快30% [1]。雷猫是一个嵌入在 CUDA 中的 DSL 工具，通过操作小型张量块来简化 AI 内核的编写，并充分利用 GPU 硬件的特性，如张量核心、异步数据传输和共享内存等 [1]。
下面是关于雷猫和 GPU 性能提升的一些关键信息：
雷猫：GPU 性能提升的秘密武器
GPU 性能提升的关键因素
未来展望：AI 与硬件的深度融合
参考文献：
Learn more: