CUTLASS

3 秒看懂

CUTLASS 是 NVIDIA 開源的 C++ 模板庫，提供在 NVIDIA GPU 上實現高效能矩陣乘法（GEMM）及相關線性代數運算的可組合建置模組。 它不是面向終端使用者的庫（如 cuBLAS），而是面向核心開發者——讓你像搭積木一樣，從暫存器→共享視訊記憶體→全域性視訊記憶體逐層組裝出 Tensor Core 加速的高效 GEMM 核心。可以說，CUTLASS 是 NVIDIA GPU 高效能運算核心的”零件工廠”。

3 分鐘產業解釋

為什麼 CUTLASS 重要？

在 AI 推論/訓練的算力棧中，矩陣乘法（GEMM / GEMV）佔據了絕大部分計算量。NVIDIA 自家的 cuBLAS 是閉源黑盒，無法定製。當產業界需要：

針對特定精度（FP8、INT4、FP4 等新格式）快速適配 GEMM 核心；
針對特定運算元融合（GEMM + bias + activation）最佳化；
針對非標準矩陣形狀（如 MoE 中稀疏專家的小 batch GEMM）定製排程；

就需要在 cuBLAS 之下的層級進行程式設計。CUTLASS 正是這個層級的標準工具。

產業位置

┌──────────────────────────────────────────┐
│  應用層: PyTorch / TensorFlow / JAX      │
├──────────────────────────────────────────┤
│  架構層: cuDNN / cuBLAS / TensorRT       │  ← NVIDIA 官方閉源庫
├──────────────────────────────────────────┤
│  核心庫層: CUTLASS / Triton / FlashAttention │  ← 開源,可定製
├──────────────────────────────────────────┤
│  程式設計層: CUDA C++ / PTX                  │
├──────────────────────────────────────────┤
│  硬體層: SM / Tensor Core / TMA          │
└──────────────────────────────────────────┘

關鍵洞見：FlashAttention、FlashDecoding、各種 FP8 GEMM kernel、乃至 NVIDIA TensorRT 內部的部分核心，其底層都大量複用了 CUTLASS 的建置模組。它實質上是 NVIDIA GPU 上高效能線性代數核心的”標準零件庫”。

15 分鐘專家深入

一、設計哲學：分層可組合（Hierarchical Composability）

CUTLASS 的核心設計理念是將 GPU GEMM 的計算對映到儲存層次結構的每一層，每一層提供模板化的建置塊：

層級	對應硬體	CUTLASS 建置塊（2.x 術語）	職責
執行緒級 (Thread)	暫存器檔案	`Mma::Operator`	一條指令完成的小矩陣乘（對映到 Tensor Core 指令）
Warp 級 (Warp)	暫存器 + warp 內同步	`Mma::Policy` 的 warp tile	組織多個執行緒級 operator，形成 warp 級 tile
CTA 級 (CTA / Threadblock)	共享視訊記憶體 (SMEM)	`Mma::Base` / `Epilogue`	從全域性視訊記憶體載入 tile 到 SMEM，執行 warp 級 MMA，寫回結果
Device 級 (Device)	全域性視訊記憶體	`Gemm::Device` / 隱式 grid 排程	將整個矩陣分 tile，排程多個 CTA 並行執行

每個層級都可以被替換——你可以換一種資料載入策略、換一種 epilogue 融合、換一種 tile 排程，而不需要重寫整個核心。這就是”可組合”的含義。

二、CUTLASS 2.x vs 3.x（CuTe）

維度	CUTLASS 2.x	CUTLASS 3.x（含 CuTe）
版面配置抽象	基於預定義的 Layout 模板（RowMajor / ColumnMajor 等）	CuTe 張量版面配置代數：用 layout algebra 組合描述任意資料排列
程式設計風格	較重的模板巢狀，“選型”式配置	更輕量的 composable 式程式設計，layout + algorithm 解耦
Hopper 特性支援	有限	完整支援：TMA、warpgroup MMA、cluster scheduling
程式碼量	相對成熟穩定	CuTe 元件引入後，核心程式碼量顯著縮減（社群反饋約縮減 30-50%）[社群經驗估算]
代表版本	CUTLASS 2.x（約 2020-2022）	CUTLASS 3.x（約 2023 起）

CuTe (CUDA Templates) 是 CUTLASS 3.x 中引入的版面配置代數核心。它把 GPU 上的張量資料版面配置抽象為數學上的 layout function（offset → coordinate 對映），支援 layout 的 composition、complement、divide 等代數運算，使得開發者可以用極少量程式碼描述複雜的資料搬運與分塊邏輯。

三、關鍵核心流程（以 GEMM C = A × B 為例）

┌─────────────────────────────────────────────────────────┐
│                    Device-level Grid                      │
│  將 M×K (A) 和 K×N (B) 分成多個 tile                     │
│  每個 CTA 負責一個 M_tile × N_tile 的輸出 tile            │
│                                                          │
│  ┌─────────────────────────────────────────────────┐    │
│  │              CTA / Threadblock                    │    │
│  │                                                   │    │
│  │  1. Prologue: 從 Global → SMEM 載入 A_tile, B_tile│    │
│  │     (Ampere+: cp.async; Hopper+: TMA)            │    │
│  │                                                   │    │
│  │  2. Mainloop: 沿 K 維度迭代                       │    │
│  │     ┌──────────────────────────────────────┐     │    │
│  │     │  * Ampere及之前:                        │     │    │
│  │     │    SMEM → Registers (ldmatrix)         │     │    │
│  │     │    Tensor Core MMA (mma.sync)          │     │    │
│  │     │  * Hopper+:                            │     │    │
│  │     │    SMEM 直接由 wgmma 指令讀取          │     │    │
│  │     │    累加到暫存器中的 accumulator           │     │    │
│  │     │  預取下一輪 tile 到 SMEM (double buffer)│     │    │
│  │     └──────────────────────────────────────┘     │    │
│  │                                                   │    │
│  │  3. Epilogue: 暫存器 → Global 寫回                 │    │
│  │     可融合 bias、activation、type conversion 等   │    │
│  └─────────────────────────────────────────────────┘    │
└─────────────────────────────────────────────────────────┘

關鍵最佳化手段：

Double/Triple Buffering：SMEM 載入與計算重疊，隱藏全域性視訊記憶體延遲。
cp.async（Ampere+）：非同步複製，繞過暫存器直達 SMEM。
TMA（Hopper+）：硬體張量記憶體加速器，自動處理多維 tile 的地址計算與非同步搬運。
Warpgroup MMA（Hopper+）：4 個 warp 協作執行一條 wgmma 指令，輸出更大的 tile。
Cluster-level Scheduling（Hopper+）：CTA Cluster 內共享 SMEM，實現跨 CTA 資料複用。

四、支援的資料型別與 Tensor Core 對映

資料型別	首次支援的 GPU 架構	Tensor Core 指令族	備註
FP16 (FP16×FP16→FP32)	Volta (SM70)	mma.sync	最成熟路徑
BF16	Ampere (SM80)	mma.sync
TF32	Ampere (SM80)	mma.sync	FP32 輸入截斷為 19-bit
INT8	Turing (SM75)	mma.sync	推論量化常用
INT4	Turing (SM75)	mma.sync	CUTLASS 2.0 已支援
FP8 (E4M3 / E5M2)	Hopper (SM90)	wgmma	訓練+推論關鍵精度
FP4（實驗性）	Blackwell (SM100)	wgmma / 未充分公開	CUTLASS 3.x 逐步適配

以上”首次支援”指 CUTLASS 庫中的適配情況，與 NVIDIA GPU 硬體 Tensor Core 指令首次出現的代際一致，但具體版本號以 [NVIDIA CUTLASS GitHub Release Notes] 為準。

五、CUTLASS 3.x 在 Hopper（SM90）上的關鍵特性

TMA（Tensor Memory Accelerator）：硬體單元，通過描述符（descriptor）一次配置即可非同步搬運任意多維 tile，取代 Ampere 時代軟體迴圈 cp.async。
wgmma（Warpgroup MMA）：一條指令讓 4 個 warp（128 執行緒）協作完成一個大 tile 矩陣乘，輸出 tile 尺寸可達 64×256（取決於資料型別）。
Distributed Shared Memory（DSMEM）：Cluster 內 CTA 可直接讀取彼此的 SMEM，無需經過全域性視訊記憶體，用於減少資料搬運。
非同步流水線（Async Pipeline）：硬體級 barrier 支援 TMA load → MMA compute → epilogue store 的多級流水線完全非同步執行。

技術原理（最深）

核心機制：Threadblock / CTA Tile 的計算-訪存對映

CUTLASS 核心設計的核心問題是：如何將一個 M×N×K 的 GEMM 對映到 GPU 的儲存層次，使得 Tensor Core 始終有資料可算。

Tile 分解示意

矩陣 A: M × K              矩陣 B: K × N
┌────┬────┬────┐           ┌────┬────┐
│ T00│ T01│ T02│           │ T00│ T01│
├────┼────┼────┤           ├────┼────┤
│ T10│ T11│ T12│           │ T10│ T11│
├────┼────┼────┤           ├────┼────┤
│ T20│ T21│ T22│           │ T20│ T21│
└────┴────┴────┘           └────┴────┘
  每個 tile: M_tile × K_tile   每個 tile: K_tile × N_tile

CTA (i,j) 負責輸出 tile C[i][j] = Σ_k A[i][k] × B[k][j]
        沿 K 維度累加, 每次載入一對 (A_tile, B_tile)

Warp 級 Tile 與執行緒級 Tile

在一個 CTA 內部，warp tile 和 thread tile 進一步細分：

CTA tile (例: 128×256×32, 以 FP16 為例)
  ├── Warp tile 0 (64×64×32)    [4 個 warp]
  ├── Warp tile 1
  ├── Warp tile 2
  └── Warp tile 3
        每個 warp tile 由多個執行緒級 operator 組成
        執行緒級 operator 對映到 Tensor Core mma.sync 指令
        (如 m16n8k16 for FP16 on Ampere)

關鍵引數選擇的影響

引數	增大效果	減小效果	約束
CTA Tile M/N	更高計算/訪存比（算術強度↑）	SMEM 佔用↓，可並行 CTA 數↑	受 SMEM 容量限制
CTA Tile K	每輪迭代資料複用↑	SMEM 佔用↑	受 SMEM 容量限制
Warp Tile	減少 warp 間同步	單 warp 暫存器壓力↑	受暫存器檔案大小限制
Stages（流水線深度）	更好隱藏記憶體延遲	SMEM 佔用↑	受 SMEM 容量限制

“Arithmetic Intensity”（算術強度） 是 GEMM 效能的核心衡量：

Arithmetic Intensity = 2×M×N×K / (資料搬運位元組數)
理想 GEMM: 接近 2×M×N×K / ((M×K + K×N + M×N) × bytes_per_element)
當 M,N 足夠大時趨近於計算/頻寬比

當算術強度 > GPU 的 FLOPS/Bandwidth 比時，GEMM 為 compute-bound；否則為 memory-bound。CUTLASS 的 tile 選擇與 double buffering 策略就是為了儘可能將核心保持在 compute-bound 區域。

Hopper wgmma 指令的微架構細節

wgmma.mma_async.sync.aligned.m64nNk{16,32}.f16.f16 ...
  - 執行 warp: 4 個 warp (warpgroup, 128 執行緒)
  - 運算元來源: SMEM（A 和 B 均可直接從 SMEM 讀取）
  - 輸出: 暫存器（distributed across warpgroup）
  - 非同步執行: 發射後可繼續執行後續指令，通過 barrier 同步
  - 關鍵優勢: 不再需要 ldmatrix 將 A/B 從 SMEM 裝入暫存器再做 MMA
              → 直接從 SMEM 讀運算元，減少暫存器壓力

這就是 CUTLASS 3.x Hopper 核心能實現更大 tile、更少暫存器溢位的硬體基礎。

技術演進史

時間	版本/事件	關鍵里程碑
2017-2018	CUTLASS 1.0	首次開源，Volta Tensor Core 支援（FP16），基礎 GEMM 模板
2019	CUTLASS 2.0	重大重構，支援 Turing INT8/INT4，引入 epilogue 融合
2020-2021	CUTLASS 2.x 逐步成熟	陸續加入 Ampere BF16/TF32/INT8 支援，cp.async 支援；社群廣泛採用，FlashAttention v1 首版使用 CUTLASS 建置塊
2023	CUTLASS 3.0 + CuTe	全新版面配置代數，Hopper SM90 完整支援（TMA、wgmma、cluster）
2023-2024	CUTLASS 3.x 持續迭代	FP8 GEMM 模板成熟，Blackwell SM100 初步適配，Grouped GEMM 最佳化
2024-2025	社群生態爆發	FlashAttention-3、各種 FP8 推論引擎、MoE 核心均基於 CUTLASS 3.x

具體版本號時間線以 [NVIDIA/cutlass GitHub releases] 為準，上表為近似標註。

技術路線對比（量化表）

維度	CUTLASS	cuBLAS（閉源）	Triton（OpenAI）	手寫 CUDA/PTX
程式設計抽象級別	中（C++ 模板）	低（API 呼叫）	中高（DSL）	最低
可定製性	⭐⭐⭐⭐⭐	⭐（黑盒）	⭐⭐⭐⭐	⭐⭐⭐⭐⭐
峰值效能	接近 cuBLAS（部分場景持平或超越）	⭐⭐⭐⭐⭐ 基準線	通常 80-95% cuBLAS [社群基準估算]	取決於開發者水平
開發效率	中等（模板超程式設計學習曲線陡峭）	高（一行 API）	高（Python 級別）	最低
新硬體適配速度	⭐⭐⭐⭐	⭐⭐⭐⭐⭐	⭐⭐⭐（依賴後端）	⭐⭐⭐⭐⭐
Tensor Core 利用	完整控制	完整但不透明	通過後端對映，控制力有限	完整控制
生態規模	大（GitHub 數千 star）	最大（NVIDIA 官方）	大且快速增長	分散
典型使用者	FlashAttention、TensorRT 核心開發者	一般使用者	研究者、快速原型	極少數專家
開源	✅ BSD 3-Clause	❌	✅ MIT	N/A

關鍵判斷：

極致效能 + 可定製 → CUTLASS
快速原型 + Python 生態 → Triton
不想折騰 → cuBLAS/cuDNN
CUTLASS 與 Triton 並非完全競爭：Triton 後端可生成 CUTLASS 風格的核心，兩者生態互補。

上下游

上游（CUTLASS 依賴什麼）

硬體:
  NVIDIA GPU (Volta/Turing/Ampere/Hopper/Blackwell)
    └── Tensor Core 指令集 (mma.sync / wgmma / ...)
    └── SMEM / 暫存器檔案 / TMA 硬體單元

軟體:
  CUDA Toolkit (nvcc / nvrtc / PTX)
  C++ 模板超程式設計 (C++17 標準)
  cmake 建置系統

下游（誰在用 CUTLASS）

直接消費者:
  ├── FlashAttention 系列 (Tri Dao)     ← 使用 CUTLASS 建置塊實現 fused attention
  ├── NVIDIA TensorRT (部分核心)        ← [NVIDIA 未充分公開具體比例]
  ├── xFormers (Meta)
  ├── vLLM / SGLang 等推論引擎          ← 通過 FlashAttention 間接依賴
  ├── 各種 FP8/INT4 量化核心            ← 社群開源專案廣泛使用
  └── 各大 AI 晶片公司對標參考           ← 軟體棧 benchmark baseline

間接消費者:
  PyTorch / JAX 使用者 → 通過上述架構間接使用 CUTLASS 核心

關鍵指標

指標	說明	參考量級
GEMM 效率	實際 TFLOPS / 硬體峰值 TFLOPS	高手最佳化後可達 80-95%+ [取決於矩陣形狀和精度]
SMEM 佔用	單 CTA 使用的共享視訊記憶體量	典型 48KB-228KB（取決於 tile size 和 stages）
暫存器壓力	每執行緒暫存器數	受限於 255 reg/thread（預設），可通過 maxrregcount 調節
Occupancy	SM 上活躍 warp / 最大 warp 比	CUTLASS 可精細調控，典型 50-100%
支援的矩陣形狀	M, N, K 範圍	從極小（如 MoE 的 M=1 token batch）到極大（數十萬維度）均有最佳化路徑
模板編譯時間	CUTLASS 核心的編譯耗時	較長（分鐘級），這是 C++ 模板超程式設計的固有代價
編譯產物大小	單核心 PTX/SASS 大小	取決於 tile 配置，通常數十 KB

供需與市場資料

CUTLASS 本身的”市場”

CUTLASS 是開源免費工具庫，沒有直接的商業營收。但其產業影響力體現在：

NVIDIA 生態護城河：CUTLASS 深度繫結 NVIDIA GPU 指令集（Tensor Core ISA），競爭對手（AMD ROCm / Intel oneAPI）需要類似工具但缺乏同等成熟度的開源替代品。AMD 的 composable_kernel (CK) 是類比專案，但社群採用度和生態成熟度與 CUTLASS 有差距 [定性判斷]。
間接影響的市場體量：全球 AI 推論/訓練市場中，絕大多數工作負載執行在 NVIDIA GPU 上。CUTLASS 作為底層核心基礎設施，其最佳化直接影響每 FLOP 的實際吞吐。保守估算，基於 NVIDIA GPU 的 AI 計算市場年規模在 [數百億美元級別，具體參考 NVIDIA Data Center 財報]，CUTLASS 的效能最佳化間接影響該市場中 GEMM 密集型工作負載的效率。

人才供需

CUTLASS 開發者極度稀缺：需要同時精通 CUDA 程式設計、GPU 微架構、C++ 模板超程式設計、線性代數演算法。全球能獨立編寫 CUTLASS 級別核心的工程師估計在 [千人量級，定性估算]。
薪資溢價：具備 CUTLASS 經驗的工程師在 AI 基礎設施領域的薪資通常顯著高於一般 CUDA 開發者 [行業經驗估算]。

代表公司與資本對映

維度	代表	說明
核心開發方	NVIDIA (NVDA)	CUTLASS 由 NVIDIA 硬體團隊維護，是其 CUDA 生態的關鍵組成部分
重度使用者 — 大廠	Meta (META)、Google (GOOGL)、Microsoft (MSFT)	內部推論最佳化團隊廣泛使用 CUTLASS 定製核心
重度使用者 — 推論創業	xAI、Together AI、Anyscale 等	使用/定製 CUTLASS 核心最佳化推論吞吐
開源 AI 推論架構	vLLM、SGLang、TensorRT-LLM	間接依賴 CUTLASS（通過 FlashAttention 等）
競爭對標	AMD (AMD) composable_kernel	AMD 的類比方案，生態成熟度較低
競爭對標	OpenAI Triton	更高層抽象的 DSL 路線，與 CUTLASS 互補

投資含義：

CUTLASS 強化了 NVIDIA CUDA 生態的開發者鎖定效應——當整個高效能核心生態都基於 CUTLASS 建置，切換到其他硬體平台的成本極高。
這種”軟生態壁壘”與 CUDA 的硬體指令集優勢疊加，構成 NVIDIA 在 AI 算力市場最深的護城河之一。

投資邏輯

核心論點

CUTLASS 是 NVIDIA AI 生態軟壁壘的重要組成部分：不是直接盈利產品，但它使得 NVIDIA GPU 上的高效能運算形成了”標準零件”生態，極大地增加了替代成本。
新精度格式（FP8/FP4）的快速適配能力是競爭武器：每當 NVIDIA 推出支援新資料型別的硬體（如 Hopper FP8、Blackwell FP4），CUTLASS 通常能快速提供高質量的 GEMM 核心模板，使得整個軟體棧能迅速跟進，而競爭對手需要更長時間適配。
MoE（Mixture of Experts）等新架構對 GEMM 核心的新需求：MoE 推論涉及大量小 batch GEMM + routing 邏輯，Grouped GEMM 等 CUTLASS 特性正好滿足此類需求。隨著 MoE 成為大型模型主流架構，CUTLASS 的 Grouped GEMM 能力愈發關鍵。

風險

Triton 的崛起：如果 Triton 在更多場景下能自動達到 CUTLASS 水平的效能，開發者可能轉向更易用的 DSL，削弱 CUTLASS 的直接影響力（但此時 Triton 的後端可能仍然產出類似 CUTLASS 風格的核心）。
AMD composable_kernel 的追趕：如果 AMD 硬體在價效比上具備競爭力且 CK 生態成熟，部分使用者可能遷移。

常見誤讀糾偏

誤讀 1：「CUTLASS 就是開源版 cuBLAS」

糾偏：兩者層級不同。

cuBLAS 是面向使用者的應用級庫——你呼叫 cublasGemmEx()，得到結果，不關心內部實現。
CUTLASS 是面向核心開發者的零件庫——你用它組裝出 GEMM 核心，可以精確控制 tile 大小、資料搬運策略、epilogue 融合邏輯等。
cuBLAS 的部分核心實現可能確實使用了 CUTLASS 元件 [NVIDIA 未充分公開具體比例]，但兩者不是等價關係。把 CUTLASS 比作”樂高積木”，cuBLAS 比作”用樂高搭好的成品車”更準確。

誤讀 2：「CUTLASS 效能不如 cuBLAS」

糾偏：在充分調優的場景下，CUTLASS 核心可以達到甚至超越 cuBLAS 的效能——因為它給了開發者更多的最佳化自由度（如自定義 epilogue 融合可以省去額外 kernel launch）。但”開箱即用”效能取決於開發者水平，cuBLAS 對常見形狀的自動調優（auto-tuning）更成熟。結論：CUTLASS 的理論效能上限 ≥ cuBLAS，但實際效能下限取決於使用者。

誤讀 3：「CUTLASS 只支援 GEMM」

糾偏：CUTLASS 同樣支援 卷積（Conv2d/Conv3d）、Grouped GEMM、Ranked K Reduction 等操作。但 GEMM 確實是最核心和最成熟的部分，社群使用中 GEMM 相關核心佔絕大多數。

誤讀 4：「CUTLASS 和 Triton 是競爭關係，二選一」

糾偏：兩者互補大於競爭。

Triton 擅長快速原型和Python 級別的自動核心生成；
CUTLASS 擅長極致效能調優和硬體特性深度利用；
實際專案中，常見策略是用 Triton 快速實現，對效能關鍵路徑的核心用 CUTLASS 重寫。
Triton 的後端在某些場景下也可能生成類似 CUTLASS 模式的 PTX 指令序列。

學習路徑

入門（1-2 周）

閱讀 NVIDIA CUTLASS GitHub Wiki 的 [Quick Start] 和 [Key Abstractions] 文件
跑通 CUTLASS 示例目錄中的 00_basic_gemm，理解 tile 配置
對比同一 GEMM 用 cuBLAS vs CUTLASS 的程式碼量和可調引數差異

進階（1-3 月）

學習 CUTLASS 2.x 的 Gemm::Device → Gemm::Kernel 層級，理解 threadblock tile 選擇
閱讀 CUTLASS Profiler 文件，理解不同 tile 配置的效能差異
研究 FlashAttention 原始碼中如何使用 CUTLASS 建置塊

專家（3-6 月+）

學習 CuTe 版面配置代數（Layout Algebra），理解 Layout::compose、Layout::complement 等操作
研究 CUTLASS 3.x Hopper 核心（SM90_64x128x32_F16... 系列），理解 TMA + wgmma 流水線
嘗試編寫自定義 Epilogue（如 GEMM + LayerNorm 融合）
對比 CUTLASS 與 AMD composable_kernel 的設計理念差異

一句話總結

CUTLASS 是 NVIDIA GPU 上高效能線性代數核心的”標準零件庫”——分層可組合、硬體深度繫結、開源可定製，構成 NVIDIA AI 算力生態的重要軟體護城河。