用微信扫码二维码
分享至好友和朋友圈
AI算力资源越发紧张的当下，斯坦福新研究将GPU运行效率再提升一波——
内核只有100行代码，让H100比使用FlashAttention-2，性能还要提升30% 。
怎么做到的？
研究人员从 “硬件实际需要什么？如何满足这些需求？” 这两个问题出发，设计了 一个嵌入式CUDA DSL工具，名为 ThunderKittens （暂且译为雷猫）。
雷猫可简化AI内核的编写，同时充分利用底层硬件能力。
具体来说，雷猫的主要抽象是寄存器和共享内存中的 小型张量块 （tile），和目前GPU中对小矩阵乘法的优化相匹配。
通过操作这些tile，开发者可相对简单地编写代码，充分利用张量核心、异步数据传输和共享内存等硬件特性。
使用雷猫实现的注意力机制内核，代码量少且能实现很高的硬件利用率，性能超过直接使用底层库（如Cutlass）。
详细讨论过程以及雷猫是怎么设计出的，研究人员以“GPUs Go Brrr”为题，发在了斯坦福Hazy Research的Blog网站上。
网友们对此讨论也十分热烈。
有网友表示读这篇Blog时，让他想起了初次了解超标量CPU架构时的惊讶感受：
还有网友表示：
H100里有什么？
斯坦福研究人员以H100为例，探讨了优化GPU的方法。
首先，回顾一下H100的硬件细节，这对于接下来的讨论非常重要。
一个H100 SXM GPU包含：
（1）80GB的HBM3内存，带宽为3TB/s（实际带宽略低）。
（2）50MB的L2缓存，带宽为12TB/s，在GPU上分为两个25MB的部分，通过交叉开关连接（这个交叉开关表现不佳）。
（3）132个流式多处理器（SM），每个包含：
除了这些，一个GPU还包括内存控制器、指令缓存……但对于这项研究而言不重要。
重要的是，所有的计算都发生在流式多处理器中， 大部分计算是在寄存器中 。
H100 GPU拥有989 TFLOPs的半精度矩阵乘法计算能力，以及约60 TFLOPs的“其他”计算能力。因此，每个周期内张量核心被使用时，至少能达到94%的硬件利用率。而张量核心不被使用时，硬件的利用率不会超过6%。
换句话说：
H100的利用率=张量核心活跃周期的百分比+/- 6%。
所以要充分发挥H100的能力， 关键是保持张量核心持续运算 。
榨干H100，要注意什么？
然鹅，要保持张量核心持续运行并不容易。
研究人员发现GPU硬件具有一些特性，对于保持矩阵乘法的运行非常重要：
这些特性在非H100 GPU上也有所适用，在H100上更加典型，就拿RTX 4090来说，相比H100处理起来简单得多。
所以接下来还是以H100为例，展开探讨这几点特性。
WGMMA指令
H100引入了一套新的指令集，名为“warp group matrix multiply accumulate”（在PTX中为wgmma.mma_async，在SASS中为HGMMA/IGMMA/QGMMA/BGMMA）。
要理解这些指令的特点，需回顾以往张量核心的使用方式。
早期GPU中的张量核心指令如wmma.mma.sync和mma.sync，要求SM一个子单元内的32个线程的一个warp同步传输数据块至张量核心并等待结果。
wgmma.mma_async指令则不同。它允许128个连续线程跨SM所有子单元协作同步，并从共享内存及寄存器（可选）异步启动矩阵乘法。这使得这些warp在等待矩阵乘法结果时可以处理其他任务。
研究人员通过微观基准测试，发现这些指令是充分发挥H100计算能力所必需的。没有这些指令，GPU的峰值利用率大约只有63%。
他们推测，这是由于张量核心需要从本地资源维持一个深度硬件pipeline。
然而，这些指令的内存布局极其复杂。未重排的共享内存布局合并性差，需要额外的L2带宽。重排的内存布局记录不准确，研究人员花费了大量时间才弄明白。
最终发现，这些布局只适用于特定矩阵形状，并与wgmma.mma_async指令的其他部分不兼容，例如硬件仅在未重排的布局下转置子矩阵。
此外，未重排的wgmma布局内存合并性差且有bank conflicts。尽管TMA和L2缓存在如flash attention这类内核上能较好地掩盖这些问题，但要充分利用硬件，必须精心控制内存请求的合并和避免bank conflicts。
尽管有这些问题，但这些指令对于充分利用H100是必不可少的。没有它们，GPU的潜在性能就损失了37%。
共享内存
共享内存的单次访问延迟约为30个周期（这也与研究人员观察的相符），这看似不多，但在这段时间内，SM的张量核心几乎能完成两次完整的32x32方阵乘法。
以前的研究，如Flash Attention，研究人员更多关注的是HBM-SRAM的瓶颈。但随着HBM速度的提升和张量核心的快速发展，即使是共享内存的相对较小延迟也变得尤为关键。
由于共享内存被分为32个独立的存储单元，处理不当可能会引发bank conflicts，即同一个内存bank同时被多个请求访问，这种情况会导致请求被序列化。研究人员实验后认为，这会显著拖慢内核速度，且wgmma与mma指令需要的寄存器布局容易受到bank conflicts的影响。
解决方法是通过各种“重排”模式调整共享内存的配置，避免bank conflicts，但细节要处理得当。
此外研究人员发现，尽可能避免在寄存器和共享内存之间的移动数据非常重要。可能的话，可使用内置硬件（如wgmma和TMA指令）进行异步数据传输。实在没法子了，再使用warp进行同步数据传输。
地址生成
H100还有一个有趣的特性，其张量核心和内存都足够快，以至于仅生成用于获取数据的内存地址就占用了芯片的大量资源，特别是加入复杂的交错或重排模式时，这种情况更为明显。
研究人员表示，英伟达提供了张量内存加速器（TMA），似乎就是已经意识到了这个问题。
TMA允许用户在全局和共享内存中指定多维张量布局，命令其异步提取张量的一部分，并在完成后触发一个屏障。这大大节省了地址生成的开销，并简化了pipelines的构建。
研究人员认为，TMA对于充分发挥H100的潜力至关重要，可能比wgmma.mma_async更为关键。
它不仅节省了寄存器资源和指令派发，还提供了如异步在全局内存上执行归约等实用功能——这在处理复杂的反向内核时尤其有用。
虽然TMA的重排模式解读有一定难度，需要进行一些逆向工程，但研究人员表示，相比之下，他们在这上面遇到的问题要少得多。
占用率
占用率指的是在GPU的相同执行硬件上同时调度的线程数。每个周期，SM的某一子单元的warp scheduler会尝试向准备就绪的warp线程发出指令。
研究人员认为，英伟达采用这种模型可以更容易地保持硬件的满负荷运行。例如，当一个线程warp等待执行矩阵乘法时，另一个可以被指派执行使用快速指数运算的指令。
在某些方面，H100对占用率的依赖程度低于前几代硬件。
它的异步特性使得即使单一指令流也能使多个硬件部分同时持续运行，包括读取内存、执行矩阵乘法、进行共享内存的归约，同时还能在寄存器上进行计算。
但高占用率容易隐藏缺陷或同步问题，一个设计良好的pipeline即使在占用率不高的情况下也能运行得相当快。
据研究人员观察，英伟达在设计GPU时确实考虑到了占用率。且由于存在足够多的同步操作和足够多的错误可能性，根据他们的经验，提高占用率通常能显著增加硬件的实际利用率。
此外，相比H100，A100和RTX 4090更依赖同步指令调度，占用率更重要。
用雷猫优化GPU
鉴于以上情况，如何才能更轻松地编写所需的内核类型，同时充分发挥硬件的全部潜力？
雷猫（ThunderKittens）登场了。
这是一个嵌入在CUDA中的DSL，本是斯坦福研究人员设计出来给自己内部使用的，后来发现还真挺好使。
Ps：起这么个名，一是他们觉得小猫很可爱，二来他们觉得大伙儿在代码中输入kittens::会很有趣。
具体来说，雷猫包含四种模板类型：
tiles通过高度、宽度和布局进行参数化；寄存器向量通过长度和布局进行参数化；而共享向量仅通过长度进行参数化，通常不会遇到bank conflicts问题。
此外，研究人员提供了一系列操作来处理这些张量，既可在warp级别使用，也可用于多个warp协作，包含初始化器，如将共享向量清零；一元操作，如exp；二元操作，如mul；行/列操作，例如行求和。
雷猫作为一个嵌入到CUDA中的库，其提供的抽象层在遇到不支持的功能时能够很好地处理。如果雷猫缺少某些功能，可以直接扩展它来实现你想要的效果。
以Tri的flash attention算法为例，在实际应用中，即使是使用英伟达的Cutlass库，实现起来也是相当复杂。
在RTX 4090上使用雷猫编写的简单flash attention内核。总共约60行CUDA代码，硬件利用率达到了75%。代码复杂性主要在于算法本身，而非交织模式或寄存器布局。
那么，它的表现如何？
这个内核只有100行代码，实际上它在H100上的性能比FlashAttention-2高出约30%。雷猫负责包装布局和指令，提供了一个可以在GPU上使用的迷你pytorch环境。
△FA2（通过Pytorch实现）与TK在H100 SXM上的多种配置比较
此外，研究人员还发布了基于线性注意力和其他新架构的内核。其中基于线性注意力的内核的运行速度可达215 TFLOPs，如果考虑到算法中固有的重计算，速度可超过300 TFLOPs。
尽管线性注意力在理论上效率更高，但此前在实际硬件上表现并不佳。因此，研究人员认为这可能促进一系列高吞吐量应用的发展。
small tile符合AI和硬件发展趋势
最后，雷猫研究团队总结了开发雷猫的一些思考。在他们看来，雷猫之所以有效，是因为它的目标并不是试图做所有事：
CUDA的确比雷猫表达能力更广，雷猫小而简单，功能有限。但雷猫的small tiles抽象设计符合AI和硬件的发展趋势。
虽然雷猫不支持小于16的维度，但研究人员认为这并不重要，因为硬件也不倾向于支持过小的维度。
从理论出发，研究人员认为需要进行一种框架转变。
“寄存器当然不应该像旧CPU那样32位字。CUDA使用的1024位宽向量寄存器确实是朝着正确方向迈出的一步。但对我们来说，寄存器是16x16的数据tile。我们认为AI需要这样的设计，毕竟，它仍然只是矩阵乘法、归约和重塑。我们认为硬件也需要这样的设计，小型矩阵乘法迫切需要超出系统级MMA的硬件支持。”
研究人员认为，应该根据硬件特性来重新定义AI的设计理念。例如，循环状态应该有多大？应该足够大以适应一个SM。计算的密度应该有多高？不应低于硬件的需求。
参考链接：[1]https://hazyresearch.stanford.edu/blog/2024-05-12-tk[2]https://github.com/HazyResearch/ThunderKittens[3]https://news.ycombinator.com/item?id=40337936
Notice: The content above (including the pictures and videos if any) is uploaded and posted by a user of NetEase Hao, which is a social media platform and only provides information storage services.
网易跟贴
注册
手机发跟贴
登录并发贴
网友评论仅供其表达个人看法，并不表明网易立场。
目前没有跟贴，欢迎你发表观点
李沐老师回归B站！带着大模型创业成果填坑来了
返回网易首页 下载网易新闻客户端学习机器学习有哪些好工具推荐
我是主推Python系的机器学习工具链的。 主要的理由是：1.现在的研究热点大都用Python实现。 可以编译到C(通过Cython)，所以可以很好的保护源码。 相对于Java的混淆字节码来说，编译到C之后的Python源码几乎不可能被分析。 是免费的，虽然用起来比matlab要麻烦。 但是有开源社区的支持，包括最近公布的numbapro可以把Python JIT到CUDA/GPU平台运行。 极大的简化了性能优化的工作。 在Linux集群上可以跑，但是matlab貌似是不行的。 作者：知乎用户来源：知乎
图形处理器的与DSP区别
GPU在几个主要方面有别于DSP(Digital Signal Processing，简称DSP，数字信号处理)架构。 其所有计算均使用浮点算法，而且此刻还没有位或整数运算指令。 此外，由于GPU专为图像处理设计，因此存储系统实际上是一个二维的分段存储空间，包括一个区段号（从中读取图像）和二维地址（图像中的X、Y坐标）。 此外，没有任何间接写指令。 输出写地址由光栅处理器确定，而且不能由程序改变。 这对于自然分布在存储器之中的算法而言是极大的挑战。 最后一点，不同碎片的处理过程间不允许通信。 实际上，碎片处理器是一个SIMD数据并行执行单元，在所有碎片中独立执行代码。 尽管有上述约束，但是GPU还是可以有效地执行多种运算，从线性代数和信号处理到数值仿真。 虽然概念简单，但新用户在使用GPU计算时还是会感到迷惑，因为GPU需要专有的图形知识。 这种情况下，一些软件工具可以提供帮助。 两种高级描影语言CG和HLSL能够让用户编写类似C的代码，随后编译成碎片程序汇编语言。 Brook是专为GPU计算设计，且不需要图形知识的高级语言。 因此对第一次使用GPU进行开发的工作人员而言，它可以算是一个很好的起点。 Brook是C语言的延伸，整合了可以直接映射到GPU的简单数据并行编程构造。 经 GPU存储和操作的数据被形象地比喻成“流”（stream），类似于标准C中的数组。 核心（Kernel）是在流上操作的函数。 在一系列输入流上调用一个核心函数意味着在流元素上实施了隐含的循环，即对每一个流元素调用核心体。 Brook还提供了约简机制，例如对一个流中所有的元素进行和、最大值或乘积计算。 Brook还完全隐藏了图形API的所有细节，并把GPU中类似二维存储器系统这样许多用户不熟悉的部分进行了虚拟化处理。 用Brook编写的应用程序包括线性代数子程序、快速傅立叶转换、光线追踪和图像处理。 利用ATI的X800XT和Nvidia的GeForce 6800 Ultra型GPU，在相同高速缓存、SSE汇编优化Pentium 4执行条件下，许多此类应用的速度提升高达7倍之多。 对GPU计算感兴趣的用户努力将算法映射到图形基本元素。 类似Brook这样的高级编程语言的问世使编程新手也能够很容易就掌握GPU的性能优势。 访问GPU计算功能的便利性也使得GPU的演变将继续下去，不仅仅作为绘制引擎，而是会成为个人电脑的主要计算引擎。
ATW技术是什么 异步时间扭曲技术是什么
ATW是一种生成中间帧的技术，当游戏不能保持足够帧率的时候，能产生中间帧，从而有效减少游戏画面的抖动，ATW技术让虚拟现实设备保持较低帧率运行看到了希望。 异步时间扭曲(Asynchronous Timewarp简称ATW)是一种生成中间帧的技术，当游戏不能保持足够帧率的时候，ATW能产生中间帧，从而有效减少游戏画面的抖动。 实现ATW是有挑战性的，主要有两个原因：1： 它需要GPU硬件支持合理的抢占粒度。 2： 它要求操作系统和驱动程序支持使GPU抢占。 让我们从抢占粒度开始，在90赫兹，帧之间的间隔大约是11ms(1/90)，这意味着为了使ATW有机生成一帧，它必须能够抢占渲染线程并且运行 时间少于11ms，然而11ms实际上不够好，如果ATW在一帧时间区间内任意随机点开始运行，那么起潜伏期(执行和帧扫描之间的时间)也将随机， 我们 需要确保我们不跳跃任何游戏渲染的帧。 我们真的期望ATW运行一直非常的短，短到在视频卡产生新的一帧之前结束，刚好有足够的时间来完成中间帧的生成，缺少自定义的同步ATW中断例程，我们可以获得高优先级抢占粒度和调度， 在最长2ms或更少的时间内。 原来，对现在的图形卡和驱动实现来说，2ms抢占是一个艰巨的任务，虽然许多GPU支持有限的形式的抢占，但执行存在显著差异。 1： 一些显卡实现厂商和驱动程序允许抢占任一批处理或回执调用粒度，虽然有帮助，但不是十分完美(举一个极端的例子，一个复杂的并包含很多绘制指令着色器可以很容易在10ms完成)。 2： 其他显卡实现厂商和驱动程序允许抢占计算着色器， 但需要特定扩展来支持。 如果抢占操作不是很快，则ATW将无法抢在画面同步之前生成中间帧。 这样，最后一帧将会再显示，将导致抖动，这意味着一个正确的实现应该能够抢占和恢复任意渲染操作，和管线状态。 理论上讲，甚至三角抢占(triangle-granularity) 不够好，因为我们不知道一个复杂着色器执行将花多长时间。 我们正与GPU制造商来实现更好的抢占，但是在这之前确实要因为这个问题花费一定时间。 另外一方面是操作系统对抢占的支持，在Windows8之前，Windiows显示驱动模型(WDDM)支持使用“批处理队列”粒度的有限抢占，对于内奸的图形驱动程序，很不幸，图形驱动程序趋向于大批量渲染效率， 导致支持ATW太粗糙。 对于Windows8，改善了WDDM1.2支持更细的抢占粒度，然而，这些抢占模式不被图形驱动程序普遍支持，渲染管线将在 Windows 10 或 DirectX12中得到显著提升。 这为开发人员提供了较低级别的渲染控制，这是一个好消息， 但直到Windows10变 为主流之前，我们还是没有标准的方式来支持渲染抢占， 造成的结果是，ATW需要特定显卡驱动的扩展。 ATW是有用的，但不是万能的。 一旦我们普遍实现了GPU渲染管线管理和任务抢占， ATW可能成为另一种工具来帮助开发人员提高性能和减少虚拟现实的抖动， 然而，由于我们这里 列出的挑战的问题，ATW不是万能的， VR的应用本身最好是维持较高的帧率，以提供最好的渲染质量。 最坏的情况，ATW生成的中间帧也可以导致用户有 不舒服的感受，换句话说，ATW无法根本解决这种不舒服。 根据生成中间帧的复杂性来说， ATW很显然表明， 甚至是位置时间扭曲， ATW不会成为一个完美的通用的解决方案，这意味着只有方向ATW和位 置ATW还算是可以的， 填充帧时偶尔会有跳跃。 为了产生一个舒适，令人信服的虚拟现实，开发人员仍然需要保持帧率在90赫兹。 试图支持传统显示器和VR双模式将会面临很大性能困难，这种巨大的性能要求是对引擎的伸缩性的考验，对于开发人员遇到的这种情况， ATW可能看起来很有吸引力， 如果达到90赫兹的频率，将使VR具有很好的舒适性，这是VR存在的真正魅力。
