本文へ移動
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で処理すべきなのか、また、もしTMAを使うならクラスタマルチキャストを併用すべきか否かという問題である。本研究では、H100を対象としたワークロードに基づく特性評価を通じて、このギャップに取り組む。我々はTMAの利用状況を体系的に整理し、多

AI 自動翻訳

要旨:NVIDIA Hopperで導入されたテンサー・メモリ・アクセラレーター(TMA)は、実際のプロダクション向け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)のhttp URLディスパッチについて、ページサイズ、データ型、モデル形状、実行フェーズの各要因にわたって検討を行う。得られた結果は、HopperにおいてTMAが有効となる場面や、マルチキャストが有利または不利となる条件を具体的かつ予測可能な形で示すものであり、さらにCUTLASS GEMMのB200に対する解析を通じて、Blackwell(SM100)サーバーへと拡張することで、定性的なディスパッチルールは世代間で継承される一方で、絶対的な性能の交差点はテンソルコアや帯域幅のスケーリングに伴い変化することも示している。

投稿履歴

論文へのアクセス:

PDFを表示

HTML(実験的)

TeXソース

追加機能

音声要約

現在の閲覧コンテキスト:

参考文献・引用

Google Scholar

Semantic Scholar

ブックマーク

機関

arXivLabsは、共同作業者が当社のウェブサイト上で新しいarXiv機能を直接開発・共有できるフレームワークです。

arXivLabsを利用する個人および組織は、オープン性、コミュニティ、卓越性、そしてユーザーデータのプライバシーという当社の価値観を積極的に受け入れ、尊重しています。arXivはこれらの価値観に固くコミットしており、それらを遵守するパートナーのみと協力しています。

arXivのコミュニティに価値をもたらすプロジェクトのアイデアをお持ちですか?arXivLabsについて詳しくご覧ください。

情報源

arXiv cs.DC
← すべてのニュースへ