top of page
検索

なぜAIは「掛け算」に命をかけるのか──GPU行列演算カーネルの深淵

シリーズ: 知新察来


◆今回のピックアップ記事：Aleksa Gordić, "Inside NVIDIA GPUs: Anatomy of high performance matmul kernels" (Aleksa Gordić's blog, 2025年9月29日）

  • 概要：NVIDIA GPU（特にHopper H100）のアーキテクチャを基礎から解説し、行列乗算カーネルの最適化手法を段階的に紹介。ナイーブな実装からwarp-tilingによる同期カーネル、さらにTMA（Tensor Memory Accelerator）やテンソルコアを活用した非同期カーネルまで、cuBLASを超える性能を達成するまでの道筋を詳述している。



ChatGPTやClaudeのようなAIが一つの質問に答えるたび、裏側では何十億もの「掛け算」と「足し算」が行われています。その計算を担うのがGPU。でも、同じGPUでも、プログラムの書き方一つで性能が10倍以上変わることがあります。たった1行のコードを入れ替えただけで、処理速度が13分の1に落ちることすらあるのです。


今回取り上げるのは、GPUエンジニアAleksa Gordić氏による技術解説記事。NVIDIA GPUの物理的な構造から、最新のHopper世代で使われる最先端の行列演算カーネルまでを丁寧に解きほぐした、まさに「GPU職人の仕事場」を覗くような内容です。


「行列演算」と聞くと、大学の線形代数を思い出して身構える方もいるかもしれません。でも、この記事の核心は意外とシンプル。つまり、いかに「データを効率よく運び」「計算ユニットを遊ばせない」か。その工夫の積み重ねが、AIの性能と電力消費を左右しているのです。


富良野とPhronaが、この複雑な世界を少しずつ紐解いていきます。




GPUって結局、何をしているの？


富良野： 今回の素材、なかなか濃いですね。NVIDIA GPUの内部構造から、行列乗算のカーネル最適化まで。正直、僕も全部は追いきれていないんですが。


Phrona： 私も途中から図を眺めながら「なるほど、なるほど」って唸ってました。でも、根っこの部分は意外とシンプルですよね。GPUがやっていることって、つまり二つだけ。


富良野： ああ、そうですね。記事の冒頭にもありました。データを運ぶことと、データを使って計算すること。この二つ。


Phrona： で、その二つがうまく噛み合わないと、どんなに高価なGPUでも宝の持ち腐れになる。


富良野： まさにそこが肝心で。記事の中で印象的だったのは、たった一行のコードを書き換えただけで、性能が13分の1に落ちた例。3171 GFLOP/sが243 GFLOP/sになったって話。


Phrona： 一行で？


富良野： 割り算と余りを計算する演算子を入れ替えただけ。見た目は些細な変更なんだけど、それでメモリへのアクセスパターンが変わって、データの「運び方」が台無しになった。


Phrona： それって、高速道路を使えば10分で着くところを、わざわざ細い路地を通って2時間かけちゃうようなもの？


富良野： うん、そのイメージに近い。GPUのメモリは階層構造になっていて、一番遠い「デバイスメモリ」──HBMって呼ばれる高帯域メモリなんですが──からデータを取ってくるのはすごく時間がかかる。だから、なるべく近くの「共有メモリ」や「レジスタ」にデータを置いておきたい。


Phrona： 近いところに倉庫を持っておいて、必要なものはそこから取り出す。


富良野： そう。そして、その「近い倉庫」は容量が小さい。だから、何をいつどこに置くかの戦略が、ものすごく重要になってくる。



メモリの「銀行」問題


Phrona： 記事の中で「バンクコンフリクト」って言葉が出てきましたよね。これ、最初ピンとこなかったんですけど。


富良野： 共有メモリの話ですね。共有メモリは32個の「バンク」に分かれていて、それぞれのバンクが独立してデータを出し入れできる。だから、32人が同時に別々のバンクにアクセスすれば、全員が一斉にデータを取れる。


Phrona： 銀行の窓口が32個あって、全部別々の窓口に並べば待ち時間ゼロ、みたいな？


富良野： そうそう。でも、もし32人が同じ窓口に殺到したら、一人ずつ順番に処理するしかない。これがバンクコンフリクト。最悪の場合、32倍の時間がかかる。


Phrona： 行列を扱うとき、列方向にアクセスすると、同じバンクに当たりやすいんですよね。


富良野： その通り。行列の要素が行優先で並んでいると、同じ列の要素は離れた場所に格納されている。で、その離れた場所がたまたま同じバンクに属していると、衝突が起きる。


Phrona： じゃあ、どうするんですか？


富良野： 一つの解決策が「スウィズリング」。データの配置をXOR演算で入れ替えて、列方向にアクセスしても異なるバンクに散らばるようにする。記事の図がすごく分かりやすかったんですが、スウィズリング後は、元の行列の「列」がメモリ上では「対角線」に並ぶ。だから一度に全部読める。


Phrona： 銀行の窓口番号をシャッフルしておくようなものですね。お客さんは自分の番号で窓口を探すけど、シャッフルされているから結果的にバラバラの窓口に分散する。


富良野： うまいたとえだ。



なぜ行列演算がAIの心臓なのか


Phrona： ところで、なんで行列の掛け算がこんなに重要なんですか？AIって、もっと複雑なことをしているイメージがあるんですけど。


富良野： 実は、Transformerと呼ばれる今のAIの中核アーキテクチャ、その計算のほとんどが行列演算なんです。具体的には、入力データに「重み」行列を掛ける操作。これが何層も重なっている。


Phrona： つまり、どれだけ速く正確に行列を掛け算できるかが、AIの性能を左右する。


富良野： その通り。しかも、行列演算は「並列化しやすい」という特性がある。出力行列の各要素は、他の要素と独立に計算できる。だからGPUみたいな並列プロセッサと相性がいい。


Phrona： でも、単純に並列化すればいいってわけじゃないんですよね。


富良野： そう。ここで「算術密度」という概念が出てくる。メモリから1バイト読み込むあたり、何回の計算ができるか。この比率が低いと、計算ユニットがデータ待ちで遊んでしまう。


Phrona： データを運ぶ速度がボトルネックになる。


富良野： 記事では「ルーフラインモデル」という図で説明されていました。縦軸が性能、横軸が算術密度。算術密度が低いうちは、どんなに頑張ってもメモリ帯域幅で性能が頭打ちになる。密度がある閾値を超えて初めて、計算能力がフルに活かせる。


Phrona： じゃあ、算術密度を上げる工夫が必要になる。


富良野： その一つが「各スレッドに複数の出力を計算させる」こと。そして、出力タイルをなるべく正方形に近づける。細長いタイルより正方形のほうが、同じデータを再利用できる回数が増えるんです。



タイリングという発想


Phrona： 「タイリング」って、床のタイルを敷き詰めるイメージですか？


富良野： まさにそう。大きな行列をそのまま扱うのではなく、小さな「タイル」に分割して処理する。各タイルを共有メモリに読み込んで、そこで計算を済ませてから結果を書き出す。


Phrona： 全体を一度に処理するのではなく、部分ごとに区切って。


富良野： で、ここで重要なのが、行列積の性質。ドット積は「部分和の合計」に分解できる。だから、タイル単位で計算した結果を足し合わせれば、最終的な答えになる。


Phrona： なるほど。タイルごとの計算は独立しているから、並列に進められる。


富良野： そして、記事で紹介されている「warp-tiling」という手法では、さらに細かく「ワープ」──32スレッドの塊──単位でタイルを割り当てる。各ワープが自分の担当領域を計算する。


Phrona： でも、これもまだ同期的な処理なんですよね。


富良野： そう。次のステップは、計算とデータ転送を「非同期」に行うこと。Hopper世代のGPUには、TMAという専用ハードウェアがあって、これがデータの移動を自動的にやってくれる。CPUで言うところのDMAに近い概念かな。



パイプラインで「待ち」をなくす


Phrona： 非同期って、具体的にはどういうことですか？


富良野： 従来の同期的なカーネルでは、データをメモリから読み込んで、読み込みが終わったら計算して、計算が終わったら書き出す、という順番。各ステップが終わるまで次に進めない。


Phrona： 料理で言えば、野菜を切り終わってから炒め始めて、炒め終わってから盛り付ける。


富良野： そう。でも非同期なら、野菜を切りながら前の野菜を炒め、炒めながら盛り付ける──みたいな流れ作業ができる。


Phrona： パイプライン処理。


富良野： 記事の図がすごく分かりやすくて。同期的な処理だと、TMA（データ転送）が動いているときテンソルコア（計算）は遊んでいる。逆にテンソルコアが動いているときTMAは遊んでいる。でもパイプライン化すれば、両方を同時に動かせる。


Phrona： 「プロデューサー」と「コンシューマー」という役割分担が出てきましたね。


富良野： ええ。128スレッドのワープグループを2つ用意して、一方がデータを供給し続け、もう一方がそれを消費して計算する。両者の間には「循環バッファ」があって、プロデューサーはバッファが空いたら書き込み、コンシューマーはバッファにデータがあれば読み取る。


Phrona： 工場の生産ラインみたい。


富良野： そのたとえ、的確ですね。そして、この協調を可能にしているのが「バリア」と呼ばれる同期プリミティブ。Hopper世代では、バリアが単にスレッドの到着を数えるだけでなく、「何バイトのデータが到着したか」も追跡できる。だから、TMAの転送完了を正確に検知できる。



テンソルコアの魔法

Phrona： テンソルコアって、普通の計算ユニットと何が違うんですか？


富良野： 通常のCUDAコアは、1回の操作で1組の浮動小数点演算──たとえば「a×b＋c」──を実行する。でもテンソルコアは、小さな行列の積を一発で計算できる。たとえば「64×16の行列」と「16×64の行列」を掛けて「64×64の結果」を得る、みたいな。


Phrona： 桁違いの効率。


富良野： しかも、4つのワープが協調して一つのテンソルコア命令を実行する。これが「wgmma」──ワープグループ行列積和命令。Hopper世代で導入された、非同期のテンソルコア命令です。


Phrona： 4つのワープで128スレッド。だからブロックサイズが128スレッドなんですね。


富良野： そう。で、記事の最適化では、このテンソルコア命令を連続して発行しつつ、TMAでのデータ転送を並行して走らせる。結果として、32 TFLOP/sだった性能が最終的に764 TFLOP/sまで上がった。約24倍。


Phrona： それでもまだ理論上の上限には届いていないんですよね。


富良野： ええ。記事によると、H100 SXM5のbf16での理論性能は約1979 TFLOP/s。実際には電力制限やその他の要因で、実効性能はもっと低くなる。でも、cuBLASという高度に最適化されたライブラリを超える性能──107%──を達成したのは、かなりの偉業。



スケジューリングの妙


Phrona： 最後のほうで出てきた「ヒルベルト曲線」、あれは何ですか？


富良野： どのSMがどの出力タイルを処理するか、その順番の話です。最もナイーブな方法は、出力行列を左上から右下へ順番に処理していく。でもこれだと、L2キャッシュの利用効率が悪い。


Phrona： なぜ？


富良野： たとえば、横一列を処理していくと、Bの行はキャッシュに残っていても、Aの列は次々と入れ替わってしまう。でも、もし処理順をブロック単位でまとめれば、隣接するタイルが同じAやBのチャンクを再利用できる。


Phrona： それがヒルベルト曲線の役割？


富良野： ヒルベルト曲線は「空間充填曲線」の一種で、2次元平面を1本の線で埋め尽くすんですが、その特徴は、曲線上で近い点は2次元空間でも近い、ということ。だから、この曲線の順番でタイルを処理すると、隣接タイルが続けて処理されやすくなり、キャッシュ効率が上がる。


Phrona： へえ、数学的な構造が実用的な最適化につながるんですね。


富良野： 面白いですよね。あと、「クラスター」という概念も重要で。Hopper世代では、複数のSMをまとめて一つのクラスターとして扱える。クラスター内のSMは、互いの共有メモリに直接アクセスできる。だから、大きな「スーパータイル」を協調して処理できる。



1%の積み重ね


Phrona： 記事の著者が言っていた「O(NR)」って表現、印象的でした。


富良野： NRは「原子炉」の略。大規模AIの訓練には何万台ものGPUが使われていて、1%の効率改善が原子炉何基分もの電力節約につながる、という話。


Phrona： アルゴリズム的な大発見はもう尽きていて、残っているのは細かい最適化の積み重ね。


富良野： でも、その積み重ねが馬鹿にならない。記事の最後の表を見ると、個々の最適化は数パーセントずつしか寄与しないものもある。バリアの呼び出しを減らして数TFLOP/s、レジスタの割り当てを調整して数TFLOP/s。


Phrona： でも、全部足し合わせると24倍になる。


富良野： そう。しかも、これらの最適化は相互に依存している。パイプライン化しなければ、非同期ストアの恩恵は得られない。テンソルコアを使わなければ、そもそも計算スループットが足りない。


Phrona： 一つ一つのピースが全体像を構成している。


富良野： 記事の結びの言葉が良かったですね。「コンピュータは理解できる」って。複雑に見えても、一つずつ紐解いていけば、必ずロジックがある。


Phrona： でも、その「紐解く」作業に何年もかかることもある。


富良野： それはまあ、そうなんですけどね。



誰がAIを支えるのか


Phrona： この記事を読んで、私が考えたのは、こういう最適化を誰がやるのか、という問題です。


富良野： どういうことですか？


Phrona： ここまで細かいハードウェアの知識が必要な最適化は、ごく一部の専門家にしかできない。AIの民主化と言われる一方で、その基盤を支えるインフラは、どんどん専門化・集中化していっている。


富良野： 確かに。cuBLASのような最適化されたライブラリがあるから、多くの研究者はそれを使うだけで済む。でも、その中身を理解している人は限られている。


Phrona： そして、その「中身」が電力消費や環境負荷に直結している。最適化が1%進むか進まないかで、データセンターの消費電力が変わる。


富良野： うーん、それは確かに考えさせられますね。技術の恩恵を受ける人と、その技術を支える人の非対称性。


Phrona： 記事の著者は、こういう知識をオープンに共有しようとしている。それ自体は素晴らしいことだと思います。でも、この記事を読んで「よし、自分もやってみよう」と思える人がどれだけいるか。


富良野： 少なくとも、「何が起きているか」を知ることには意味があると思う。全員がカーネルを書く必要はないけど、AIの裏側でこういう営みがあるということを知っておくことは、技術との付き合い方を考える上で大事かもしれない。



 

ポイント整理


  • GPUの二つの仕事

    • GPUは「データを運ぶ」ことと「データで計算する」ことの二つを行う。この二つがうまく協調しないと、性能が大幅に低下する。

  • メモリ階層の重要性

    • GPUのメモリはデバイスメモリ（HBM）→L2キャッシュ→共有メモリ（SMEM）→レジスタという階層構造を持つ。上位（遠い）ほど容量が大きいが遅く、下位（近い）ほど高速だが容量が小さい。頻繁にアクセスするデータは計算ユニットの近くに置く戦略が重要。

  • コアレス化とバンクコンフリクト

    • デバイスメモリへのアクセスは連続したアドレスを同時にアクセスする「コアレスアクセス」が効率的。共有メモリでは32個のバンクに分散されており、異なるスレッドが同じバンクの異なるアドレスにアクセスすると「バンクコンフリクト」が発生し、処理がシリアル化される。

  • スウィズリング

    • バンクコンフリクトを回避するため、データ配置をXOR演算で入れ替える手法。行列の列方向アクセスでも、異なるバンクに分散されるようになる。TMA（Tensor Memory Accelerator）が自動的にスウィズリングを適用する。

  • 算術密度とルーフラインモデル

    • 性能は「メモリ帯域幅に制約される領域」と「計算能力に制約される領域」に分かれる。1バイトあたりの計算回数（算術密度）が低いと、どんなに計算能力があってもメモリがボトルネックになる。

  • タイリング

    • 大きな行列を小さなタイルに分割し、タイル単位で共有メモリに読み込んで計算する手法。行列積は「部分積の和」に分解できるため、この分割が可能。正方形に近いタイルほどデータ再利用率が高くなる。

  • warp-tilingからテンソルコアへ

    • 従来のwarp-tiling手法は通常のCUDAコアを使った同期的な処理。Hopper世代ではテンソルコア（専用の行列演算ユニット）とTMA（非同期データ転送ユニット）を活用することで、10倍以上の性能向上が可能。

  • 非同期パイプライン処理

    • プロデューサー（データ転送担当）とコンシューマー（計算担当）のワープグループに分け、循環バッファを介して並行動作させる。これによりTMAとテンソルコアの両方を同時に活用できる。

  • ヒルベルト曲線スケジューリング

    • 出力タイルの処理順序を空間充填曲線（ヒルベルト曲線）に従って決めることで、L2キャッシュの再利用率を向上させる。

  • クラスター機能

    • Hopper世代では複数のSMをクラスターとしてグループ化し、クラスター内で共有メモリを相互アクセスできる（DSMEM）。これにより、より大きなタイルを協調処理できる。

  • 最適化の累積効果

    • 個々の最適化は数%ずつの改善でも、すべてを組み合わせることで32 TFLOP/sから764 TFLOP/sへと約24倍の性能向上を達成。大規模AI訓練では、1%の効率改善が膨大なエネルギー節約につながる。



キーワード解説


【matmul（行列乗算）】

二つの行列を掛け合わせる演算。Transformerなど現代のAIモデルの中核的な処理であり、計算量のほとんどを占める。


【SM（Streaming Multiprocessor）】

GPUの基本的な計算ユニット。H100には132個のSMが搭載されている。


【SMEM（Shared Memory）】

SM内の高速なプログラマ管理メモリ。L1キャッシュと物理ストレージを共有している。


【HBM（High Bandwidth Memory）】

GPUのデバイスメモリとして使用される高帯域積層メモリ。大容量だがアクセス遅延が大きい。


【warp（ワープ）】

32スレッドのグループ。GPUの命令実行の基本単位。


【warp-group（ワープグループ）】

4つのワープ（128スレッド）からなるグループ。Hopper世代のテンソルコア命令の実行単位。


【テンソルコア】

小さな行列積を一発で計算できる専用ユニット。通常のCUDAコアより桁違いに高いスループットを持つ。


【TMA（Tensor Memory Accelerator）】

Hopper世代で導入されたハードウェアユニット。グローバルメモリと共有メモリ間の非同期データ転送を担当し、スウィズリングも自動適用する。


【wgmma（Warp Group Matrix Multiply-Accumulate）】

Hopper世代のテンソルコア命令。4ワープが協調して非同期に行列積和演算を実行する。


【コアレスアクセス】

ワープ内のスレッドが連続したメモリアドレスにアクセスするパターン。メモリ効率が高い。


【バンクコンフリクト】

共有メモリの同じバンクに複数スレッドが異なるアドレスでアクセスした際に発生する競合。処理がシリアル化される。


【スウィズリング】

XOR演算によってデータ配置を入れ替え、バンクコンフリクトを回避する手法。


【算術密度（Arithmetic Intensity）】

メモリから読み込む1バイトあたりの演算回数。この値が低いとメモリ帯域幅がボトルネックになる。


【ルーフラインモデル】

算術密度と性能の関係を可視化するモデル。メモリ律速領域と計算律速領域を識別できる。


【パイプライン処理】

データ転送と計算を並行して行い、待ち時間を削減する手法。


【persistent kernel（永続カーネル）】

SM数と同じ数のブロックを起動し、各ブロックが内部ループで複数タイルを処理する手法。カーネル起動オーバーヘッドを削減できる。


【ヒルベルト曲線】

2次元平面を1本の線で埋め尽くす空間充填曲線。曲線上で近い点は平面上でも近いという性質があり、キャッシュ効率の良いスケジューリングに利用される。


【DSMEM（Distributed Shared Memory）】

クラスター内のSM間で共有メモリを相互アクセスできる機能。Hopper世代で導入。


【PTX】

NVIDIAの仮想ISA（命令セットアーキテクチャ）。GPUの世代間で互換性がある。


【SASS】

NVIDIAの実機ISA。GPUが実際に実行する命令。世代ごとに異なる。



本稿は近日中にnoteにも掲載予定です。
ご関心を持っていただけましたら、note上でご感想などお聞かせいただけると幸いです。
bottom of page