cutlass
CUTLASS
彎刀 3.6.0 - 2024 年 10 月
CUTLASS 是 CUDA C++ 範本抽象的集合，用於在 CUDA 內的所有層級和規模上實現高效能矩陣-矩陣乘法 (GEMM) 和相關運算。它包含類似於用於實現 cuBLAS 和 cuDNN 的分層分解和資料移動策略。 CUTLASS 將這些「移動部件」分解為由 C++ 模板類別抽象化的可重複使用、模組化軟體元件。 概念並行化層次結構不同層級的原語可以透過自訂平鋪大小、資料類型和其他演算法策略進行專門化和調整。由此產生的靈活性簡化了它們在自訂核心和應用程式中作為構建塊的使用。
為了支援各種應用，CUTLASS 為混合精度計算提供了廣泛的支持，為半精度浮點 (FP16)、BFloat16 (BF16)、Tensor Float 32 (TF32)、單精度浮點(FP32)、透過張量核心指令進行FP32 模擬、雙精確度浮點(FP64) 類型、整數資料型別（4b 和8b）以及二進位資料型別(1b)。 CUTLASS 示範了針對由 NVIDIA Volta、Turing、Ampere 和 Hopper 架構實現的可編程、高吞吐量Tensor Core 的曲速同步矩陣乘法運算。
請參閱快速入門指南以快速開始。
請參閱功能列表，以了解執行模型層次結構的每個層級支援的操作列表。
CUTLASS 3.0 引入了一個新的核心庫 CuTe，用於描述和操作執行緒和資料的張量。 CuTe 是 C++ CUDA 範本抽象的集合，用於定義和操作執行緒和資料的分層多維佈局。 CuTe 提供了Layout和Tensor對象，它們緊湊地封裝了資料的類型、形狀、記憶體空間和佈局，同時為使用者執行複雜的索引。這讓程式設計師可以專注於演算法的邏輯描述，而 CuTe 則為他們進行機械記帳。借助這些工具，我們可以快速設計、實作和修改所有密集線性代數運算。
CuTe 的核心抽像是分層多維佈局，可以用資料數組組成來表示張量。佈局的表示足夠強大，足以表示我們實現高效密集線性代數所需的幾乎所有內容。佈局還可以透過功能組合來組合和操作，在此基礎上我們建立了大量常見操作，例如平鋪和分區。
CUTLASS 3.0 及更高版本在其模板的整個 GEMM 層次結構中採用 CuTe。這極大地簡化了設計並提高了程式碼的可組合性和可讀性。更多特定於 CuTe 的文檔可以在其專用文件目錄中找到。
除了 GEMM 之外，CUTLASS 透過隱式 GEMM 演算法實現高效能卷積。隱式 GEMM 是將卷積運算表述為 GEMM，從而利用 CUTLASS 的模組化 GEMM 管道。這使得 CUTLASS 能夠透過重複使用高度最佳化的 GEMM 元件來建立卷積。
CUTLASS 3.6.0 是 CUTLASS 的更新，新增了：
Hopper 結構化稀疏 GEMM。
FP16
FP8
INT8
TF32
對 CUTLASS 3.x 卷積kernel::ConvUniversal API 進行重構，使其與gemm::GemmUniversal保持一致。現在 3.x 卷積 API 不再被視為 beta API。
改進的混合輸入 GEMM 和INT4 x FP8僅縮放模式的查找表實作。
用於 Top-K 選擇的 EVT 節點以及使用這些節點的 softmax 和 GEMM 範例。
程式化依賴啟動 (PDL)，利用新的 Hopper 功能來加速兩個背對背核心及其對應的文件。
一個新的偵錯工具，synclog，用於將核心內的所有同步事件轉儲到檔案中。有關詳細信息，請參閱同步日誌文件。
分組 GEMM 的新的支援 TMA 的尾聲可帶來顯著的性能改進以及 EVT 支援。
支援 SIMT 的指標數組尾聲。
用於分組 GEMM 的新 Ping-Pong 核心調度和一些其他最佳化。
CUTLASS 分析器核心的新實例化策略以及 CUTLASS 分析器中實例化層級的改進文件。
對cutlass::bfloat16_t的比較和計算的新硬體支持
修復了 Windows 上 isnan 對half_t的使用。
最低要求：
建築：沃特
編譯器：必須至少支援 C++17
CUDA 工具包版本：11.4
從 CUTLASS 3.0 開始，CUTLASS 刪除了對以下內容的支援：
Maxwell 與 Pascal GPU 架構
烏班圖16.04
CUDA 10.2
C++ 語言版本低於 17。
有關版本和更新的詳細列表，請參閱變更日誌。
CUTLASS 原語非常有效率。 當用於建立設備範圍的 GEMM 核心時，它們在標量 GEMM 計算方面表現出與 cuBLAS 相當的峰值性能。上圖顯示了自 CUTLASS 3.1 以來 NVIDIA H100（NVIDIA Hopper 架構）上持續的 CUTLASS 效能改進。 CUTLASS 3.5.1 使用 CUDA 12.5u1 工具包編譯。 Tensor Core 運算是使用 CUDA 的 mma 和 wgmma 指令實現的。
當使用 CUTLASS 構建塊構建設備範圍的隱式 gemm（Fprop、Dgrad 和 Wgrad）核心時，在 NVIDIA A100 上運行 Resnet-50 層時，CUTLASS 性能也與 cuDNN 相當，如上圖所示。 Tensor Core 運算是使用 CUDA 的 mma 指令實現的。
CUTLASS 需要 C++17 主機編譯器，並且在使用CUDA 12.4 工具包建置時效能最佳。它也相容於 CUDA 11.4、CUDA 11.5、CUDA 11.6、CUDA 11.7、CUDA 11.8、CUDA 12.0、CUDA 12.1、CUDA 12.2.2、CUDA 12.3.1 和 CUDA 12.3.2。
我們測試了以下環境。
| 作業系統 | 編譯器 |
|---|---|
| 烏班圖18.04 | 海灣合作委員會7.5.0 |
| 烏班圖20.04 | 海灣合作委員會10.3.0 |
| 烏班圖22.04 | 海灣合作委員會 11.2.0 |
| 烏班圖22.04 | 鏗鏘10.0.0 |
| 烏班圖22.04 | 鏗鏘14.0.6 |
| 烏班圖22.04 | 鏗鏘17.0.6 |
| 視窗10.0 | Visual Studio 2019 v16.11.27 |
注意：GCC 8.5.0 具有有關折疊表達式和重載運算子的已知回歸。建議使用 GCC 7.5.0 或（首選）GCC >= 9。
CUTLASS 在以下 NVIDIA GPU 上成功運行，並預計在基於 Volta、Turing、Ampere、Ada 和 Hopper 架構的 NVIDIA GPU 上也能有效運作。
| 圖形處理器 | CUDA運算能力 | CUTLASS-3 所需的最低 CUDA 工具包 |
|---|---|---|
| NVIDIA V100 張量核心 GPU | 7.0 | 11.4 |
| NVIDIA泰坦V | 7.0 | 11.4 |
| NVIDIA GeForce RTX 2080 TI、2080、2070 | 7.5 | 11.4 |
| 英偉達T4 | 7.5 | 11.4 |
| NVIDIA A100 張量核心 GPU | 8.0 | 11.4 |
| 英偉達A10 | 8.6 | 11.4 |
| NVIDIA GeForce RTX 3090 | 8.6 | 11.4 |
| NVIDIA GeForce RTX 4090 | 8.9 | 11.8 |
| 英偉達L40 | 8.9 | 11.8 |
| NVIDIA H100 張量核心 GPU | 9.0 | 11.8 |
一般來說，為一種目標架構產生的 PTX 程式碼可以在未來的架構上運行（即，它是向前相容的）。 然而，CUDA 12.0引入了「架構加速功能」的概念，其PTX沒有前向相容性保證。有些 Hopper PTX 指令屬於此類架構加速功能，因此需要sm_90a目標架構（請注意附加的「a」）。有關此指令和其他架構加速指令的更多詳細信息，請參閱 CUDA 文件。
目標架構資訊透過 cmake 標誌CUTLASS_NVCC_ARCHS傳遞到 CUTLASS。為了最大化Hopper GH100的效能，使用者需要建構以90a為目標架構的CUTLASS。如果使用者意外地使用 SM90 目標（注意缺少「a」）以及 CUDA Toolkit 12 或 11.8 建構了使用 SM90a 功能（例如 Hopper Tensor Core 指令）的內核，則內核預計會因執行時間而失敗錯誤。
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
請參閱功能文檔，以了解有關哪些核心需要哪些目標架構的詳細資訊。
以下文件和隨附的 Doxygen 文件中描述了 CUTLASS。
快速入門指南 - 建置與執行 CUTLASS
功能 - 總結 CUTLASS 中可用的功能
CUDA 中的高效 GEMM - 描述如何在 CUDA 中高效實現 GEMM 內核
CUTLASS 3.x 設計 - 描述 CUTLASS 3.x 設計、其優點以及 CuTe 如何使我們能夠編寫更多可組合元件
GEMM API 3.x - 描述 CUTLASS 3.x GEMM 模型與 C++ 範本概念
GEMM API 2.x - 描述 CUTLASS 2.x GEMM 模型與 C++ 範本概念
隱式 GEMM 卷積 - 描述 CUTLASS 中的 2-D 和 3-D 卷積
程式碼組織 - 描述 CUTLASS 專案的組織和內容
術語 - 描述代碼中使用的術語
程式設計指南 - 撰寫高效能現代 CUDA C++ 的指南
基本型別 - 描述 CUTLASS 中用於表示數值量和陣列的基本 C++ 類
佈局 - 描述記憶體中矩陣和張量的佈局
Tile Iterators - 描述了在記憶體中迭代矩陣區塊的 C++ 概念
CUTLASS Profiler - 命令列驅動的分析應用程式
CUTLASS 實用程式 - 用於促進快速開發的附加模板
相關內核啟動 - 描述了 Hopper 中的一項新功能，該功能允許在同一流中重疊相關內核，以及如何在 CUTLASS 中使用它。
我們也在 2018 年 GPU 技術大會的演講中描述了高效能 GEMM 的結構。
CUTLASS：CUDA 內所有等級和規模的密集線性代數的軟體基元
開發 CUDA 核心將 Tensor Core 推向 NVIDIA A100 的絕對極限
在 CUTLASS 中使用張量核心加速卷積
透過增加 CUTLASS 中張量核心的利用率來加速後向資料梯度
CUTLASS：Python API、增強功能和 NVIDIA Hopper
CUTLASS 是一個僅包含頭檔的範本庫，不需要建置即可供其他專案使用。客戶端應用程式應將 CUTLASS 的include/目錄作為其包含路徑中的目標。
CUTLASS 單元測試、範例和實用程式可以使用 CMake 建置。快速入門指南中給出了 CMake 的最低版本。確保CUDACXX環境變數指向系統上安裝的 CUDA 工具包中的 NVCC。
$ 匯出 CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc在 CUTLASS 專案中建立一個建置目錄，然後執行 CMake。預設情況下，CUTLASS 將為 CUDA 架構版本 5.0、6.0、6.1、7.0、7.5、8.0、8.6、8.9 和 9.0 建置核心。為了減少編譯時間，您可以透過變更 CMake 設定設定CUTLASS_NVCC_ARCHS來指定建置 CUTLASS 的體系結構。
$ mkdir 構建 && cd 構建 $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # 針對 NVIDIA 的 Ampere 架構進行編譯
在build/目錄中，透過使用 make 建立目標test_unit來編譯並執行 CUTLASS 單元測試。
單元測試被組織為幾個二進位文件，鏡像 CUTLASS 的頂級命名空間，並且它們可以透過 make 的-j命令列參數並行執行。
$ 使 test_unit -j …… …… …… [----------] 全域測試環境拆解 [==========] 運行了來自 57 個測試案例的 946 個測試。 （總共 10812 毫秒） [通過] 946 項測試。
所有測試都應在支援的平台上通過，但測試的確切數量可能會隨著時間的推移而變化。
CUTLASS 與實用程式、工具、範例和單元測試一起被安排為僅包含頭檔的程式庫。 Doxygen 文件提供了 CUTLASS 專案中定義的文件、類別和範本概念的完整清單。
原始碼組織的詳細說明可以在 CUTLASS 文件中找到，但下面總結了幾個主要元件。
include/ # client applications should target this directory in their build's include paths cutlass/ # CUDA Templates for Linear Algebra Subroutines and Solvers - headers only arch/ # direct exposure of architecture features (including instruction-level GEMMs) conv/ # code specialized for convolution epilogue/ # code specialized for the epilogue of gemm/convolution gemm/ # code specialized for general matrix product computations layout/ # layout definitions for matrices, tensors, and other mathematical objects in memory platform/ # CUDA-capable Standard Library components reduction/ # bandwidth-limited reduction kernels that do not fit the "gemm" model thread/ # simt code that can be performed within a CUDA thread transform/ # code specialized for layout, type, and domain transformations * # core vocabulary types, containers, and basic numeric operations cute/ # CuTe Layout, layout algebra, MMA/Copy atoms, tiled MMA/Copy algorithm/ # Definitions of core operations such as copy, gemm, and operations on cute::tuples arch/ # Bare bones PTX wrapper structs for copy and math instructions atom/ # Meta-information either link to or built from arch/ operators mma_atom.hpp # cute::Mma_Atom and cute::TiledMma copy_atom.hpp # cute::Copy_Atom and cute::TiledCopy *sm*.hpp # Arch specific meta-information for copy and math operations * # Core library types such as Shape, Stride, Layout, Tensor, and associated operations
CUTLASS SDK範例應用CUTLASS模板來實現基本計算。
tools/ library/ # CUTLASS Instance Library - contains instantiations of all supported CUTLASS templates include/ cutlass/ library/ profiler/ # CUTLASS Profiler - command-line utility for executing operations in the # CUTLASS Library util/ # CUTLASS Utilities - contains numerous helper classes for include/ # manging tensors in device memory, reference cutlass/ # implementations for GEMM, random initialization util/ # of tensors, and I/O.
test/unit/目錄包含使用 Google Test 實現的單元測試，演示了核心 API 元件的基本用法和 CUTLASS GEMM 計算的完整測試。
快速入門指南中描述了建置和運行單元測試的說明。
tools/profiler/目錄包含用於啟動每個 GEMM 核心的命令列實用程式。它可以建構如下：
$ make cutlass_profiler -j16
預設情況下，只會為每種資料類型、數學指令和佈局實例化一個圖塊大小。若要實例化所有內容，請在從空build/目錄執行 CMake 時設定下列環境變數。請注意，這會導致數以萬計的核心和較長的建置時間。這也會導致二進位大小過大，並且在某些平台上連結器無法建立庫。因此，強烈建議僅產生內核的子集，如下面小節所示。
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=全部 …… $ make cutlass_profiler -j16
為了嚴格編譯一個內核或一小組內核，可以使用帶有通配符的逗號分隔的內核名稱列表來減少內核集。以下範例展示了為 NVIDIA Ampere 和 Turing 架構建立一個或一個核心子集：
若要編譯針對 NVIDIA Ampere 和 Turing 架構的具有 FP32 累積和 FP16 輸入的 Tensor Core GEMM 核心子集，請使用下列 cmake 命令列：
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 …… $ make cutlass_profiler -j16
用於分析 Tensor Core GEMM 核心子集的範例命令列如下：
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
……
===============================
問題編號：1
提供者： CUTLASS
操作類型：gemm
操作：cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
狀態：成功
驗證：開
處置：透過
參考設備：透過
cuBLAS：透過
參數： --gemm_kind=通用 --m=3456 --n=4096 --k=4096 --A=f16:列 --B=f16:行 --C=f32:列 --alpha=1
--beta=0 --split_k_slices=1 --batch_count=1 --op_class=tensorop --accum=f32 --cta_m=256 --cta_n=128
--cta_k=32 --stages=2 --warps_m=4 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=8 --min_cc=75
--max_cc=1024
位元組：118489088 位元組
失敗次數：115992428544次失敗
運行時間：1.55948 毫秒
記憶體：70.7616 GiB/秒
數學：74378.8 GFLOP/s
===============================
……要編譯一個針對 NVIDIA Ampere 和 Turing 架構的 SGEMM 內核，請使用下列 cmake 命令列：
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 …… $ make cutlass_profiler -j16
用於分析單一 SGEMM CUDA 核心的範例命令列如下：
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
===============================
問題編號：1
提供者： CUTLASS
操作類型：gemm
操作：cutlass_simt_sgemm_128x128_8x2_nn_align1
狀態：成功
驗證：開
處置：透過
cuBLAS：透過
參數： --m=3456 --n=4096 --k=4096 --A=f32:列 --B=f32:列 --C=f32:列 --alpha=1 --beta=0 -- split_k_slices=1
--batch_count=1 --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4
--warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024
位元組：180355072 位元組
失敗次數：115992428544次失敗
運行時間：6.73655 毫秒
記憶體：24.934 GiB/秒
數學：17218.4 GFLOP/s
===============================若要編譯針對 NVIDIA Ampere 和 Turing 架構的具有 FP32 累積和 FP16 輸入的前向傳播 (fprop) 的 Tensor 核心卷積核子集，請使用下列 cmake 命令列：
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 …… $ make cutlass_profiler -j16
用於分析 Tensor Core 卷積核子集的範例命令列如下：
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*fprop_optimized_f16 --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3
……
===============================
問題編號：1
提供者： CUTLASS
操作類型：conv2d
操作：cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
狀態：成功
驗證：開
處置：透過
參考設備：透過
參數： --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q =224 --pad_h=1 --pad_w=1
--stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --啟動=f16:nhwc --Filter=f16:nhwc --輸出=f32:nhwc
--conv_mode=交叉 --iterator_algorithm=優化 --alpha=1 --beta=0 --split_k_mode=序列 --split_k_slices=1
--eq_gemm_provider=無 --op_class=tensorop --accum=f32 --cta_m=128 --cta_n=128 --cta_k=32 --stages=5
--warps_m=2 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=16 --min_cc=80 --max_cc=1024
位元組：1130659840 位元組
失敗次數：118482796544 次失敗
運行時間：0.711496 毫秒
記憶體：1479.99 GiB/秒
數學：166526 GFLOP/s
===============================
……要編譯並執行一個針對 NVIDIA Ampere 和 Turing 架構的具有 F32 累積和 FP32 輸入的前向傳播 (fprop) 的 CUDA Core 卷積內核，請使用以下 cmake 命令列：
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc …… $ make cutlass_profiler -j16
用於分析一個 CUDA Core 卷積內核的範例命令列：
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=33
===============================
問題編號：1
提供者： CUTLASS
操作類型：conv2d
操作：cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
狀態：成功
驗證：開
處置：透過
參考設備：透過
參數： --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q =224 --pad_h=1 --pad_w=1
--stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --啟動=f32:nhwc --Filter=f32:nhwc --輸出=f32:nhwc
--conv_mode=交叉 --iterator_algorithm=優化 --alpha=1 --beta=0 --split_k_mode=序列 --split_k_slices=1
--eq_gemm_provider=無 --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4
--warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024
位元組：2055798784 位元組
失敗次數：118482796544 次失敗
運行時間：7.34266 毫秒
記憶體：260.752 GiB/秒
數學：16136.2 GFLOP/s
===============================請點擊以下鏈接，以獲取有關選擇性編譯 CUTLASS 核心的更多 CMake 範例：
GEMM CMake 範例
隱式 GEMM 卷積 CMake 範例
有關 CUTLASS Profiler 的更多詳細資訊請參閱此處。
CUTLASS 是 NVIDIA Corporation 根據 3 個「新」BSD 許可證作為開源軟體發布的。
CUTLASS 開發者和貢獻者的官方清單可在此處找到：貢獻者。
版權所有 (c) 2017 - 2024 NVIDIA 公司及附屬公司。版權所有。 SPDX 許可證識別碼：BSD-3 條款
Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: 1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. 2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. 3. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission. THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
