NVIDIA/cutlass

⭐ 8,042 🍴 1,300 C++ (CUDA C++)📄 BSD-3-Clause（Python CuTe DSLはNVIDIA EULA） GitHubで見る →

主な特徴

**マルチアーキテクチャ対応**: Volta、Turing、Ampere、Ada、Hopper、Blackwellアーキテクチャをサポート
**混合精度演算**: FP64、FP32、TF32、FP16、BF16、FP8（e5m2、e4m3）、INT8、INT4、バイナリ型など多様なデータ型をサポート
**Tensor Core最適化**: 各世代のNVIDIA Tensor Coreに最適化された実装
**ヘッダーオンリーライブラリ**: テンプレートベースの設計で簡単に統合可能
**階層的抽象化**: スレッド、ワープ、スレッドブロック、デバイスレベルの演算を提供
**CuTe DSL**: Python APIによる高速なプロトタイピングとカーネル開発（v4.0以降）
**包括的な最適化**: タイリング、スウィズリング、パイプライニングなどの高度な最適化技術

リポジトリ解析: NVIDIA/cutlass

基本情報

リポジトリ名: NVIDIA/cutlass
主要言語: C++ (CUDA C++)
スター数: 8,042
フォーク数: 1,300+（推定）
最終更新: v4.1.0（2025年7月）
ライセンス: BSD-3-Clause（Python CuTe DSLはNVIDIA EULA）
トピックス: cuda, gpu, linear-algebra, gemm, tensor-cores, high-performance-computing, matrix-multiplication, cutlass, cute, convolution, nvidia-gpu

概要

一言で言うと

CUTLASSは、NVIDIA GPUで高性能な行列積（GEMM）と関連する線形代数演算を実装するためのC++テンプレートライブラリで、Tensor Coreを活用した最適化された演算を提供する。

詳細説明

CUTLASS（CUDA Templates for Linear Algebra Subroutines）は、2017年から開発されているNVIDIA公式のGPU向け線形代数ライブラリである。階層的な分解とデータ移動の戦略を組み込み、これらの「動く部品」を再利用可能でモジュラーなソフトウェアコンポーネントと抽象化に分解している。並列化階層の異なるレベルのプリミティブは、カスタムタイリングサイズ、データ型、その他のアルゴリズムポリシーによって特殊化およびチューニングできる。CUTLASS 4.0からはPython DSL（CuTe DSL）も追加され、C++の深い専門知識なしに高性能カーネルを記述できるようになった。

主な特徴

マルチアーキテクチャ対応: Volta、Turing、Ampere、Ada、Hopper、Blackwellアーキテクチャをサポート
混合精度演算: FP64、FP32、TF32、FP16、BF16、FP8（e5m2、e4m3）、INT8、INT4、バイナリ型など多様なデータ型をサポート
Tensor Core最適化: 各世代のNVIDIA Tensor Coreに最適化された実装
ヘッダーオンリーライブラリ: テンプレートベースの設計で簡単に統合可能
階層的抽象化: スレッド、ワープ、スレッドブロック、デバイスレベルの演算を提供
CuTe DSL: Python APIによる高速なプロトタイピングとカーネル開発（v4.0以降）
包括的な最適化: タイリング、スウィズリング、パイプライニングなどの高度な最適化技術

使用方法

インストール

前提条件

Python 3.10以上（CuTe DSL使用時）
C++17対応コンパイラ（GCC 7.5+、Clang 7.0+）
CUDA Toolkit 11.4以上（11.8以上推奨、最新機能には12.8以上が必要）
CMake 3.19以上
アーキテクチャ: Volta（compute capability 7.0）以上

インストール手順

# 方法1: ヘッダーオンリーの直接利用（ビルド不要）
# include/ディレクトリをインクルードパスに追加

# 方法2: ソースからビルド
git clone https://github.com/NVIDIA/cutlass.git
cd cutlass
export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc
mkdir build && cd build
cmake .. -DCUTLASS_NVCC_ARCHS=80  # Ampereアーキテクチャ向け
# または
cmake .. -DCUTLASS_NVCC_ARCHS="90a"  # Hopperアーキテクチャ向け
# または
cmake .. -DCUTLASS_NVCC_ARCHS="100a"  # Blackwellアーキテクチャ向け
make -j

# Python CuTe DSLのインストール
pip install cutlass-core  # またはuv add cutlass-core

基本的な使い方

Hello World相当の例

// examples/00_basic_gemm/basic_gemm.cuからの抽出
#include "cutlass/gemm/device/gemm.h"

// 単精度GEMMの定義
using ColumnMajor = cutlass::layout::ColumnMajor;
using CutlassGemm = cutlass::gemm::device::Gemm<
    float, ColumnMajor,  // A行列
    float, ColumnMajor,  // B行列
    float, ColumnMajor>; // C行列

// GEMM実行
CutlassGemm gemm_operator;
CutlassGemm::Arguments args(
    {M, N, K},                    // 問題サイズ
    {A, lda}, {B, ldb},          // 入力行列
    {C, ldc}, {C, ldc},          // 出力行列
    {alpha, beta}                // スカラー
);
cutlass::Status status = gemm_operator(args);

実践的な使用例

// Tensor Coreを使用した半精度GEMM（Ampereアーキテクチャ）
#include "cutlass/gemm/device/gemm_universal.h"

using ElementA = cutlass::half_t;
using ElementB = cutlass::half_t;
using ElementC = float;

using Gemm = cutlass::gemm::device::GemmUniversal<
    ElementA, cutlass::layout::RowMajor,
    ElementB, cutlass::layout::ColumnMajor,
    ElementC, cutlass::layout::RowMajor,
    float,  // accumulator type
    cutlass::arch::OpClassTensorOp,  // Tensor Core使用
    cutlass::arch::Sm80,             // Ampere
    cutlass::gemm::GemmShape<128, 256, 64>,   // CTA tile
    cutlass::gemm::GemmShape<64, 64, 64>,     // Warp tile  
    cutlass::gemm::GemmShape<16, 8, 16>       // MMA shape
>;

高度な使い方

# CuTe DSLを使用したGEMMの実装例
from cutlass import cutlass, cutlass_tensor
from cutlass import layouts, datatypes
import cute

# Hopperアーキテクチャ向けGEMMカーネル
@cutlass.kernel
def gemm_kernel(A, B, C, alpha, beta):
    # CuTeのレイアウトとテンソル操作を使用
    tA = cute.make_tensor(A, cute.make_layout(shape_A, stride_A))
    tB = cute.make_tensor(B, cute.make_layout(shape_B, stride_B))
    tC = cute.make_tensor(C, cute.make_layout(shape_C, stride_C))
    
    # Tensor Core MMA操作
    cute.gemm(tiled_mma, tCrA, tCrB, tCrC)
    
    # エピローグでアルファ/ベータ適用
    cute.axpby(alpha, tCrC, beta, tCrD)

ドキュメント・リソース

公式ドキュメント

README.md: プロジェクト概要、インストール方法、基本的な使い方
docs.nvidia.com/cutlass: 包括的な公式ドキュメントサイト
Doxygenドキュメント: APIリファレンス（nvidia.github.io/cutlass）
CHANGELOG.md: バージョン履歴と新機能

サンプル・デモ

examples/00_basic_gemm: 基本的なGEMMの実装例
examples/48_hopper_warp_specialized_gemm: Hopperアーキテクチャ向け最適化GEMM
examples/70_blackwell_gemm: Blackwellアーキテクチャ向けGEMM
examples/77_blackwell_fmha: Blackwell向けFused Multi-Head Attention
examples/python/CuTeDSL: Python DSLを使用したカーネル実装例

チュートリアル・ガイド

Quick Start Guide: CUTLASSの基本的なビルドと実行方法
CuTe DSL Quick Start: Python APIの入門ガイド
Efficient GEMM in CUDA: CUDAでの効率的なGEMM実装方法
CUTLASS 3.x Design: CuTeを使用した新しい設計アプローチ
GTCプレゼンテーション: GPU Technology Conferenceでの発表資料

技術的詳細

アーキテクチャ

全体構造

CUTLASSは階層的な抽象化を提供し、デバイスレベルのGEMMからスレッドレベルの操作までをカバーする。各レベルはテンプレートパラメータでカスタマイズ可能で、タイリングサイズ、データ型、レイアウト、アルゴリズムポリシーを選択できる。CuTeライブラリはコアの抽象化を提供し、テンソルとレイアウトの操作を簡素化する。

ディレクトリ構成

cutlass/
├── include/           # ヘッダーファイル（クライアントがインクルードする）
│   ├── cutlass/      # CUTLASSコアライブラリ
│   │   ├── arch/    # アーキテクチャ固有の機能（MMA命令等）
│   │   ├── conv/    # 畳み込み用特化コード
│   │   ├── epilogue/# エピローグ演算
│   │   ├── gemm/    # GEMM演算特化コード
│   │   └── layout/  # 行列レイアウト定義
│   └── cute/         # CuTeライブラリ
│       ├── algorithm/# テンソルアルゴリズム
│       ├── arch/    # アーキテクチャ固有のCopy/MMA atom
│       └── atom/    # 基本演算プリミティブ
├── examples/          # 使用例（80+のサンプル）
├── tools/             # ツール群
│   ├── library/      # CUTLASS Instance Library
│   └── profiler/     # パフォーマンスプロファイラ
├── test/              # ユニットテスト
└── python/CuTeDSL/    # Python DSL実装

主要コンポーネント

Device GEMM: デバイス全体のGEMMカーネル
- 場所: cutlass/gemm/device/
- 依存: Kernel, Threadblock, Warp, Thread各レベル
- インターフェース: Arguments構造体、operator()
Collective MMA: 集団MMA演算モジュール
- 場所: cutlass/gemm/collective/
- アーキテクチャ毎の実装（sm70〜sm120）
- Tensor Core命令の活用
CuTe Tensor/Layout: テンソルとレイアウトの抽象化
- 場所: cute/tensor.hpp、cute/layout.hpp
- 機能的なレイアウト合成、タイリング、パーティショニング

技術スタック

コア技術

言語: C++17（CUDA C++、テンプレートメタプログラミングを活用）、Python 3.10+（CuTe DSL）
フレームワーク: ヘッダーオンリーライブラリ、CUDAカーネルテンプレート
主要ライブラリ/機能:
- CUDA Toolkit (11.4+): GPUカーネル実行、Tensor Coreアクセス
- CuTe: テンソルとレイアウトの抽象化
- Google Test: ユニットテストフレームワーク

開発・運用ツール

ビルドツール: CMake 3.19+、NVCC（CUDAコンパイラ）
テスト: Google Testベースの包括的なユニットテスト
CI/CD: GitHub Actions（リント、型チェック、ユニットテスト）
デプロイ: ヘッダーオンリーのため、includeパスの設定のみで利用可能

設計パターン・手法

テンプレートメタプログラミング: コンパイル時に最適化されたカーネルを生成
階層的分解: Device > Kernel > Threadblock > Warp > Threadの階層構造
タイリング: 各階層での効率的なデータ分割とアクセス
パイプライン化: メモリレイテンシの隠蔽
ソフトウェアプリフェッチ: グローバルメモリから共有メモリへの非同期コピー

データフロー・処理フロー

グローバルメモリから共有メモリへのタイルロード: 非同期コピーを使用
共有メモリからレジスタへのフラグメントロード: 各スレッドがデータを取得
Tensor Core/SIMD命令による演算: MMA（Matrix Multiply-Accumulate）操作
アキュムレータの書き戻し: エピローグでのスケーリングと出力
結果のグローバルメモリへの書き込み: 最終結果の保存

API・インターフェース

公開API

Device GEMM API

目的: デバイス全体でのGEMM演算を実行
使用例:

using Gemm = cutlass::gemm::device::Gemm<...>;
Gemm::Arguments args({M, N, K}, {A, lda}, {B, ldb}, {C, ldc}, {D, ldd}, {alpha, beta});
Gemm gemm_op;
cutlass::Status status = gemm_op(args);

CuTe DSL API

@cutlass.kernel
def my_kernel(A, B, C):
    # CuTeのTensorとLayout操作
    tA = cute.make_tensor(A, layout_A)
    tB = cute.make_tensor(B, layout_B)
    
    # Collective MMA操作
    cute.gemm(mma, tA, tB, tC)

設定・カスタマイズ

CMake設定

# ターゲットアーキテクチャの指定
CUTLASS_NVCC_ARCHS="80;90a;100a"  # Ampere, Hopper, Blackwell

# カーネルの選択的ビルド
CUTLASS_LIBRARY_KERNELS="cutlass_tensorop_*gemm_f16_*"

# テストレベル
CUTLASS_TEST_LEVEL=0  # 0: Sanity, 1: Release, 2: Exhaustive

拡張・カスタマイズ

カスタムティルドコピーの実装（Copy Atom）
カスタムMMAアトムの定義
新しいエピローグフュージョンの追加
カスタムデータ型のサポート

パフォーマンス・スケーラビリティ

パフォーマンス特性

ベンチマーク結果:
- H100での理論ピーク性能の90-98%を達成（FP16/FP8 GEMM）
- Blackwell SM100で理論ピークの95%以上を達成
最適化手法:
- Tensor Core命令の活用（mma, wgmma, umma）
- 非同期コピーとパイプライン化
- ソフトウェアプリフェッチ
- 動的スケジューリング（Stream-K、Persistent Kernel）

スケーラビリティ

大規模行列に最適化（タイルサイズの調整が重要）
マルチGPUによる分散実行サポート
Grouped GEMMでのバッチ処理
ブロックスケールデータ型でのメモリ効率向上

制限事項

技術的な制限:
- 小さい行列ではオーバーヘッドが大きい
- 特定のデータ型はStructured Output対応のLLMが必要
- Windowsプラットフォームでのビルド問題（v3.x以降）
運用上の制限:
- アーキテクチャ毎のコンパイルが必要
- バイナリサイズが大きくなる可能性（全カーネルビルド時）

評価・所感

技術的評価

強み

最高レベルの性能: NVIDIA GPUの理論性能に近い実装を提供
包括的なアーキテクチャサポート: Voltaから最新のBlackwellまで対応
柔軟なカスタマイズ性: テンプレートパラメータで細かく調整可能
活発な開発: NVIDIA公式サポートで継続的な更新
Python DSL: プロトタイピングと学習の敷居を下げる

改善の余地

学習曲線の高さ: C++テンプレートメタプログラミングの習得が必要
ドキュメントの充実: APIリファレンスはあるが、詳細な設計指針が不足
ビルド時間: 全カーネルのビルドには非常に長い時間がかかる
プラットフォーム制限: Windowsサポートの問題

向いている用途

深層学習フレームワークの実装: PyTorch、TensorFlowのバックエンド
高性能コンピューティング: 科学計算、シミュレーション
AI推論エンジン: モデルサービングの最適化
カスタムCUDAカーネル開発: 特殊な線形代数演算の実装
教育・研究: GPUプログラミングの学習材料

向いていない用途

小規模行列演算: オーバーヘッドが大きい
CPUベースのアプリケーション: NVIDIA GPU専用
シンプルな行列演算のみ: cuBLASの直接利用で十分
メモリ制約が厳しい環境: 大きなバイナリサイズ

総評

CUTLASSは、NVIDIA GPUで最高性能の線形代数演算を実現するための最も強力なツールの一つである。特にTensor Coreを活用した深層学習用ワークロードでは、他の選択肢を圧倒する性能を発揮する。C++テンプレートの複雑さはあるものの、v4.0で導入されたPython DSLはより多くの開発者に門戸を開く可能性を秘めている。NVIDIAの公式サポートと継続的な開発により、今後もGPUコンピューティングの中核的な存在であり続けると考えられる。