# Tensor Memory Accelerator ## 概要 Tensor Memory Accelerator(TMA)は、NVIDIAのHopper/Blackwell世代GPUに搭載されたオンチップのDMA(direct memory access)エンジンであり、グローバルメモリと共有メモリ(SMEM)/テンソルメモリ(TMEM)の間で1D〜5Dの多次元バルクコピーを、SMの命令パイプラインをブロックせずに非同期実行する。CUDAの`cuda::memcpy_async()`と`cuda::pipeline` APIから利用でき、コアレッシング・ストライド転送・多次元転送を含むバルクコピーをハードウェアが自動処理する。ダブルバッファリング(ping-ponging)と組み合わせることで、SMの計算ユニットが前段のタイルを処理している間に次のタイルをTMAが先読みし、数百サイクルに及ぶDRAM往復レイテンシを計算の裏に隠蔽する。 ## 経歴・背景 - CUDA C++の`cuda::pipeline` APIと共同設計されたハードウェア機能で、APIが公開する抽象化(`producer_acquire`/`producer_commit`/`consumer_wait`/`consumer_release`)がTMAの転送能力に直接対応する。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 7 Profiling and Tuning GPU Memory Access Patterns]]) - Hopper/Blackwellでは`cp.async.bulk.*`系のPTX命令族としてコンパイルされ、アラインメントと転送方向(グローバル⇔共有メモリ)の条件を満たす場合に自動的に活用される。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 7 Profiling and Tuning GPU Memory Access Patterns]]) - 第10章(イントラカーネルパイプライニング・ウォープ特化)でスレッドブロッククラスタと組み合わせた発展的な利用が扱われる予定。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 7 Profiling and Tuning GPU Memory Access Patterns]]) ## 関連プロジェクト・製品 - [[NVIDIA]] — 開発元 - [[wiki/entities/AI Systems Performance Engineering|AI Systems Performance Engineering]] — 本エンティティの一次出典書籍 ## 関連 - [[@2025__OReilly__AI Systems Performance Engineering - Chapter 7 Profiling and Tuning GPU Memory Access Patterns]] - [[メモリコアレッシング]] ## 出典 - Chris Fregly, *AI Systems Performance Engineering*, O'Reilly Media, 2025, Chapter 7 (§Asynchronous Memory Prefetching and Tensor Memory Accelerator).