斯坦福“雷猫”：GPU 性能的秘密武器 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 的潜力。 雷猫的未来展望： 跨平台支持： 研究人员计划将雷猫扩展到 AMD 硬件平台，进一步扩大其应用范围。 更深度的硬件融合： 雷猫团队将继续研究如何更深入地利用 GPU 硬件特性，例如利用 H100 的分布式共享内存，以进一步提升性能。 AI 设计理念的革新： 研究人员认为，未来 AI 的设计理念应该更加贴近硬件，例如循环状态的大小应该与 SM 的容量相匹配，计算密度应该不低于硬件的需求。 雷猫的出现，为 AI 计算领域带来了新的希望。通过深入理解硬件特性，并设计与之匹配的 AI，我们可以进一步提升 GPU 的性能，释放 AI 的潜力，为 AI 技术的进步贡献力量。 参考文献： [1] https://hazyresearch.stanford.edu/blog/2024-05-12-tk 斯坦福大学的研究人员开发了一款名为“雷猫”的新工具，它被称为GPU性能的秘密武器。这个工具可以提高GPU的运行效率，比使用FlashAttention-2（FA2）还要快。通过使用雷猫，研究人员成功地将H100 GPU的性能提升了30%[1]。 以下是关于斯坦福“雷猫”工具的一些关键信息： 工具简介： 雷猫是一个嵌入式CUDA DSL工具，它可以简化AI内核的编写，并充分利用底层硬件能力[1]。 雷猫的主要抽象是寄存器和共享内存中的小型张量块（tile），与GPU中对小矩阵乘法的优化相匹配[1]。 工具优势： 使用雷猫实现的注意力机制内核，代码量少且能实现很高的硬件利用率，性能超过直接使用底层库（如Cutlass）[1]。 雷猫可以帮助开发者相对简单地编写代码，充分利用张量核心、异步数据传输和共享内存等硬件特性[1]。 H100 GPU的优化方法： H100 GPU是斯坦福研究人员探讨的一个例子，用于讨论GPU的优化方法[1]。 H100 GPU具有80GB的HBM3内存、50MB的L2缓存和132个流式多处理器（SM）等硬件细节[1]。 为了充分发挥H100的计算能力，关键是保持张量核心持续运算[1]。 保持张量核心持续运算的注意事项： WGMMA指令是保持矩阵乘法运行的关键，但使用起来较为麻烦[1]。 共享内存的速度可能不如预期快，需要额外注意[1]。 生成地址的成本较高，需要注意控制[1]。 保持高占用率对于提升性能是有益的，寄存器至关重要[1]。 斯坦福大学的研究人员通过这项研究和开发的“雷猫”工具，成功提高了GPU的运行效率，为AI算力资源紧张的时代提供了一种新的解决方案。这项研究对于进一步优化GPU性能和提高AI算力的利用率具有重要意义。 Learn more: 斯坦福让“GPU高速运转”的新工具火了，比FlashAttention2更快创事记新浪科技_新浪网 斯坦福让”GPU高速运转”的新工具火了，比FlashAttention2更快 | 量子位 优化架构，降低频率，骁龙8gen2高性能、低功耗的秘密武器之一|高通|gpu|处理器|cpu|骁龙+移动平台_网易订阅
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 的潜力。
雷猫的未来展望：
雷猫的出现，为 AI 计算领域带来了新的希望。通过深入理解硬件特性，并设计与之匹配的 AI，我们可以进一步提升 GPU 的性能，释放 AI 的潜力，为 AI 技术的进步贡献力量。
参考文献：
[1] https://hazyresearch.stanford.edu/blog/2024-05-12-tk
斯坦福大学的研究人员开发了一款名为“雷猫”的新工具，它被称为GPU性能的秘密武器。这个工具可以提高GPU的运行效率，比使用FlashAttention-2（FA2）还要快。通过使用雷猫，研究人员成功地将H100 GPU的性能提升了30%[1]。
以下是关于斯坦福“雷猫”工具的一些关键信息：
斯坦福大学的研究人员通过这项研究和开发的“雷猫”工具，成功提高了GPU的运行效率，为AI算力资源紧张的时代提供了一种新的解决方案。这项研究对于进一步优化GPU性能和提高AI算力的利用率具有重要意义。
Learn more: