# CUTLASS ## 概要 CUTLASS(CUDA Templates for Linear Algebra Subroutines)は[[NVIDIA]]が提供するC++テンプレートベースのGEMM(一般行列積)・Tensor Core最適化ライブラリである。開発者はGEMM形状や精度型をテンプレートパラメータとして指定するだけで、共有メモリタイリング・非同期メモリコピー(`cp.async`/TMA)・ダブルバッファリング・ワープ特化・スレッドブロッククラスタ/ペアといった高度な最適化が自動的に適用される。cuBLASなど他の高性能ライブラリの内部実装にも使われており、PyTorchの`torch.compile`(TorchInductor)が生成する融合アテンションカーネルなどでも間接的に利用される。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 9 Increasing CUDA Kernel Efficiency and Arithmetic Intensity]] §Using CUTLASS for Optimal Arithmetic Intensity and Tensor Core Performance) ## 主な特徴 - **テンプレート1つでの最適化適用**: `cutlass::gemm::device::Gemm<...>`のようなテンプレート型を宣言し、行列形状・ストライドを渡すだけでTensor Coreカーネルが生成される。 - **手書きMMAカーネルに匹敵する性能**: 書籍(ch.9)の例示的な比較(表9-1)では、Tensor Core利用率98%で手書きMMAカーネルと同等、レジスタ使用量(約52→60)・共有メモリ使用量(2KB→4KB)がやや多い程度の差にとどまる。 - **TMEM/TMA対応**: BlackwellのTensor Memory(TMEM、256KB/SM)とTensor Memory Accelerator(TMA)を用いたタイルのステージング・アキュムレータ管理を透過的に行う。 - **精度対応**: FP4・FP8・FP16・TF32のオペランド型をテンプレートでサポートし、バイアス加算や活性化関数などの後処理も同一カーネルに融合できる。 ## 横断的知見 - **CUTLASSは「演算強度を上げる技術」の実装を自動化するレイヤーである**: [[Rooflineモデル]]・[[カーネルフュージョン]]で語られるタイリング・ダブルバッファリング・フュージョンといった技術は、本来は開発者が手作業で実装する必要があるが、CUTLASSはこれらをテンプレート宣言1つに圧縮する。書籍(ch.9)はCUTLASSを「標準的なGEMM/融合アテンションではまず選ぶべき手段」と位置づけ、独自データレイアウトや特殊な融合パターンが必要な場合のみ手書きMMAカーネルを検討すべきとする。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 9 Increasing CUDA Kernel Efficiency and Arithmetic Intensity]] §Using CUTLASS for Optimal Arithmetic Intensity and Tensor Core Performance) ## 未解決の問い - CUTLASSが選択するタイル形状(128×128、256×128等)の探索は、どこまで自動チューニング(auto-tuning)されており、どこまでNVIDIAによる手動のGPU世代別チューニングに依存しているか。 - Triton・CuTe-DSL([[カーネルフュージョン]]の横断的知見を参照)とCUTLASSは、どちらもTensor Core最適化を高レベルに抽象化するが、両者の使い分け(パフォーマンス上限・開発コスト・対応形状の広さ)はどう整理されるか。 - PyTorchのTorchInductorがCUTLASSカーネルをディスパッチする条件(形状・精度・演算パターン)は、書籍のどの章でより詳しく扱われるか(ch.13/14のPyTorchコンパイラ章を確認する必要がある)。 ## 関連 - 概念: [[Rooflineモデル]] / [[カーネルフュージョン]] / [[CUDA]] - エンティティ: [[NVIDIA]] - ソース: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 9 Increasing CUDA Kernel Efficiency and Arithmetic Intensity]] ## 出典 - Chris Fregly, *AI Systems Performance Engineering*, O'Reilly Media, 2025, Chapter 9, §Using CUTLASS for Optimal Arithmetic Intensity and Tensor Core Performance.