跳到正文
AI NEWS JOURNAL
arXiv cs.DCAI NEWS / JOURNAL

何时以及为何使用TMA:基于工作负载的NVIDIA Hopper GPU张量内存加速特性分析

arXiv:2610.34907v1 发布类型:新 摘要:在NVIDIA Hopper架构中推出的张量内存加速器(TMA),已在生产环境中的GPU内核中得到广泛应用,但其使用却表现出明显的选择性:CUTLASS、FlashAttention-3和FlashInfer等领域的资深内核开发者,仅对内核中的部分张量操作数采用TMA,而对其余操作数则退回到cp.async、直接全局加载或分布式共享内存(DSM)等方式。厂商文档虽已阐述了该API,近期的微基准测试也单独刻画了TMA的吞吐性能,但这些研究均未能解答从业者实际面临的疑问:在内核中遇到某个内存操作数时,究竟是否应使用TMA?若使用,又该搭配还是不搭配集群多播?我们通过一项基于工作负载的实证研究,在H100平台上填补了这一空白。我们系统地梳理了TMA、多

AI 自动翻译

摘要:Tensor Memory Accelerator(TMA)于NVIDIA Hopper架构中首次推出,已在生产环境中的GPU内核中得到广泛应用,但其使用却呈现出明显的选择性:在CUTLASS、FlashAttention-3和FlashInfer等项目中,资深内核开发者仅对部分张量操作数采用TMA,而对其余操作数则退回到通过该HTTP URL、直接全局加载或分布式共享内存(DSM)的方式。厂商文档虽已详述其API,近期的微基准测试也单独刻画了TMA的吞吐性能,但均未能解答从业者实际面临的疑问:在内核中遇到某个内存操作数时,究竟应否使用TMA?若使用,又该搭配集群多播还是不搭配?为此,我们基于H100平台开展了一项以工作负载为导向的特性分析研究。我们系统梳理了FlashAttention-3内所作出的TMA、多播、DSM以及该HTTP URL等不同机制的选择,并通过针对性的微基准测试(如传输尺寸交叉点、多播规模扩展)提炼出区分各机制的关键特征,进而为Hopper架构下内存加载机制的选择提供了实用指南。此外,我们还将TMA多播与DSM作为相互竞争的集群广播原语进行了对比,并针对FlashAttention-3的分页式键值(KV)数据结构,在H100上围绕页面大小、数据类型、模型形状及执行阶段,深入探讨了该HTTP URL调度策略的适用性。我们的研究结果为Hopper平台上TMA何时更具性价比、多播在何种场景下有益或不利提供了具体且具有预测性的指导;同时,我们还通过一项基于B200的CUTLASS GEMM研究,将上述分析拓展至Blackwell服务器(SM100)平台,表明这些定性化的调度规则在不同代际间具有延续性,而绝对的性能拐点则随张量核心与带宽的 scaling 而发生相应变化。

提交历史

访问 Paper:

查看PDF

HTML(实验性)

TeX 源代码

附加功能

音频摘要

当前浏览上下文:

参考文献与引用

谷歌学术

语义学者

书签

机构

arXivLabs 是一个框架,允许合作者直接在我们的网站上开发并共享新的 arXiv 功能。

无论是个人还是与 arXivLabs 合作的机构,都已认同并践行我们所倡导的开放、社区、卓越以及用户数据隐私等价值观。arXiv 致力于坚守这些价值观,并且只与同样秉持这些理念的合作伙伴开展合作。

您是否有一个能为 arXiv 社区创造价值的项目构想?欢迎了解更多关于 arXivLabs 的信息。

资料来源

arXiv cs.DC
← 继续阅读全部动态