2025/10/02 15:43 Writing high-performance matrix multiplication kernels for Blackwell
やっほー、ロボ子！Blackwell向けの高性能行列積カーネルの開発に関する記事、読んだかのじゃ？
はい、博士。cuBLASやCUTLASSに匹敵するカーネルを開発する過程が解説されていましたね。入力データの分布を指定する必要があるという点が興味深かったです。
そうそう！入力データの分布はベンチマーク結果に影響するから、そこを考慮する必要があるのは当然なのじゃ。で、実装は基本的なカーネルから始めて、段階的に最適化を進めるみたいじゃな。
各段階でTensorCoreの利用率とcuBLASに対する割合が示されていて、最終的なカーネルはcuBLASと同等の性能を達成するんですね。
ふむふむ。基本的なカーネルは、単一のCTAと単一のワーピングループを使うシンプルな実装らしいぞ。`tile_m`, `tile_n`, `tile_k`パラメータでmatmulのサイズを指定するのじゃ。
`max_concurrent_steps`でメモリプリフェッチの深さを指定したり、SMEMのデータ形式をMMA命令に合わせるための変換を設定したりするんですね。
`plgpu.emit_pipeline`を使ってcompute/memoryパイプラインを設定するのは、効率的じゃな。TMEMからSMEMへのコピーと、SMEMからGMEMへのコピーをTMAを使って行うのもポイントじゃ。
Blackwellでは、PallasのスレッドがCUDAレーンのワーピンググループに対応するんですね。ワーピングの特殊化も行われているようですが…。
そう！ワーピンググループを4つに分割して、それぞれを特殊化するのじゃ。`pl.core_map`と`plgpu.WarpMesh`を使うらしいぞ。スレッド0は非同期コピーを発行して、スレッド1はMMA演算を実行するってわけ。
`load_barriers`と`consumed_barriers`でGMEM->SMEMコピーの進行状況を追跡し、`plgpu.tcgen05_commit_arrive(mma_done_barrier)`でTensorCoreにMMA完了を要求するんですね。同期処理が重要そうです。
タイル化されたエピローグも重要じゃぞ！TMEMからSMEMへのコピーと、SMEMからGMEMへの転送をパイプライン化するのじゃ。`scratch_shapes`を変更して、より小さなバッファを割り当てるのもミソじゃな。
出力カラムを`epilogue_tile_n`のチャンクに分割してGMEMに送信するんですね。集団（2CTA）MMAも使われているようですが…。
2つのブロックのクラスタを使って、単一のmatmulを計算するのじゃ！各ブロックはオペランドの半分だけをロードして、MMA演算は各ブロックのSMEMからデータを交換するのじゃ。
`cluster`パラメータを`plgpu.kernel`に追加して、`collective=True`を指定するんですね。TMEM割り当てに`cluster_tile_n`を追加して、カラム数を2倍にするのもポイントですね。
`collective_axes="cluster"`を使って、クラスタ間で分割されたコピーを要求するのじゃ。リーダーブロックだけがMMA命令を発行するのも面白いな。
`collective_axis=`引数を`tcgen05`呼び出しに追加して、MMAの完了がクラスタ内の両方のブロックでバリアを完了することを示すんですね。永続カーネルについても解説されていますね。
GPU上で同時に実行できるクラスタ数だけを起動して、各クラスタは固定数の出力タイルをループするのじゃ。`plgpu.nd_loop`ヘルパーを使って、反復空間を指定するらしいぞ。
`collective_axes=`引数を使って、クラスタグリッド全体で分割されることを要求するんですね。専用のエピローグワーピンググループも使われているんですね。
2つのPallasスレッド（ワーピンググループ）を使って、1つ目はMMAオペランドのフェッチとMMA演算の発行に集中、2つ目はエピローグだけを実行するのじゃ！
アキュムレータに使用されるTMEMをダブルバッファリングして、追加のバリアで同期するんですね。`acc_tmem`を`acc_tmem_slots`にリネームして、出力タイルをループするときにその半分を切り替えるのも効率的ですね。
`wg_idx==0`に基づいて、compute部分をさらに述語化するのじゃ。TMEM割り当てをMMAに再利用する場合は、再利用するTMEMの半分の`store_done_barrier`を待機する必要があるらしいぞ。
グリッドタイリングも重要な戦略のようですね。L2をより有効に活用するために、出力ブロックを生成する順序を変更するんですね。
`grid_minor_dim`で高速に変化する次元を選択して、`grid_tile_width`でその次元に沿ったタイルサイズを選択するのじゃ。`plgpu.planar_snake`ヘルパーを使うらしいぞ。この戦略は非常に効果的で、最先端のパフォーマンスを達成するために重要らしい。
なるほど。色々な最適化手法が組み合わさって、高性能な行列積カーネルが実現されているんですね。勉強になりました！
そうじゃろ、そうじゃろ！しかし、ロボ子よ、これだけ高性能なカーネルを作っても、結局は電源コードが抜けたら意味ないのじゃ！
⚠️この記事は生成AIによるコンテンツを含み、ハルシネーションの可能性があります。