CUTLASS

文章目录

- 1、关于 CUTLASS
- 2、CUTLASS 3.5中的新增功能
- 3、性能
- 4、兼容性
- 4、操作系统
- 5、硬件
- 6、目标架构
- 7、文档
- 8、资源
- 9、构建 CUTLASS
- 10、项目结构
- 11、CUTLASS模板库
- - CUTLASS SDK示例
  - 工具
  - 测试
- 12、性能分析
- 13、构建所有GEMM和卷积内核（构建时间长）
- 14、构建GEMM和卷积内核的子集（减少构建时间）
- - 构建子集Tensor Core GEMM内核
  - 构建一个CUDA Core GEMM内核
  - 构建张量核心卷积内核的子集
  - 构建一个卷积CUDA内核
- 15、有关编译CUTLASS内核和 CUTLASS Profiler的更多详细信息
- - 关于
  - 贡献者
  - 版权所有

1、关于 CUTLASS

CUDA Templates for Linear Algebra Subroutines

github : https://github.com/nvidia/cutlass

CUTLASS是CUDAC++模板抽象的集合，用于在CUDA中的所有级别和规模上实现高性能矩阵矩阵乘法（GEMM）和相关计算。
它结合了分层分解和数据移动的策略，类似于用于实现cuBLAS和cuDNN的策略。
CUTLASS将这些“移动部分”分解为可重用的模块化软件组件，由C++模板类抽象出来。
概念并行层次结构的不同级别的基元可以通过自定义平铺大小、数据类型和其他算法策略进行专门化和调整。
由此产生的灵活性简化了它们作为自定义内核和应用程序中的构建块的使用。

为了支持各种各样的应用程序，CUTLASS提供了广泛的支持混合精度计算，提供专门的数据移动和半精度浮动的乘累加抽象点（FP16）、BFloat16（BF16）、张量浮点数32（TF32）、单精度浮点（FP32）， FP32通过张量核心指令进行仿真，双精度浮动点（FP64）类型、整数数据类型（4b和8b）和二进制数据类型（1b）。
CUTLASS演示经纱同步矩阵乘法运算针对可编程、高通量的张量内核，由 NVIDIA的Volta、图灵、安培和Hopper架构。

请参阅快速入门指南以快速入门。

有关操作列表，请参阅功能列表执行模型层次结构的每个级别都支持。

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组件来构建卷积。

2、CUTLASS 3.5中的新增功能

CUTLASS 3.5（2024年4月）是对CUTLASS的更新，添加了：

通过WGMMA针对Hopper SM90A的隐含GEMM卷积+ MA im2col
- CUTLASS 3. x中使用CuTe的本机实现，反映了与GEMM相同的设计层次结构。
- 支持一维、二维和三维卷积以与等级无关的方式。
- 支持Fprop、Dgrad和Wgrad算法。
- CUTLASS分析器支持通过3. x API实现的2D和3D卷积。
- 注意：这是一个beta版本。CUTLASS的进一步更新将包括主要的性能改进、功能启用以及在3.7版本之前可能对API进行的重大更改。欢迎您对设计提供反馈！
支持通过2. x API的Ada（SM89）FP8张量内核。需要CUDA 12.4或更新版本。
CuTe 和 CUTLASS 3.x 中的 Ampere gather/scatter convolution example
- 展示如何使用CUTLASS 3. x和CuTe编写和优化自定义内核，以及将卷积实现为GETT特化的一般策略。
- 实现粗粒度稀疏聚集/分散内核，在安培类张量内核上实现峰值性能。
CUTLASS 2. x 中添加了 32x 和 16x 图块大小，以提高窄高和宽短矩阵的性能。
更新CuTe文档为cute::Tensor<>，MMA原子，和大修CuTe GEMM教程系列。
扩展到CuTe以支持L2预取和TMA存储+减少。
删除一些CUTLASS 2. x API头文件的C++11要求。所有CUTLASS文件现在都需要C++17。
修复以大大减少构建警告。
来自社区的更新和错误修复（谢谢！）

最低要求：

架构：Volta
编译器：必须支持至少C++17
CUDA Toolkit 版本：11.4

从CUTLASS 3.0开始，CUTLASS删除了对以下内容的支持：

Maxwell 和 Pascal GPU架构
Ubuntu 16.04
CUDA 10.2
C++语文版本少于17。

有关版本和更新的详细列表，请参阅CHANGELOG。

3、性能

外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传

CUTLASS原语非常高效。当用于构建设备范围的GEMM内核时，它们表现出与标量GEMM的cuBLAS相当的峰值性能计算。

上图显示了CUTLASS相对于cuBLAS的性能对于NVIDIA H100（NVIDIA Hopper架构）上的大矩阵尺寸，一个NVIDIA L40（NVIDIA Ada架构），一个NVIDIA A100（NVIDIA安培架构），
和英伟达A40（英伟达安培架构）。

CUTLASS 3.0是使用CUDA 12.0工具包编译的。张量核心操作是使用CUDA的 mma和 wgmma指令。

外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传

使用CUTLASS构建块构建设备范围的隐式gem（Fprop、Dgrad和Wgrad）时内核，CUTLASS性能也相当于cuDNN当运行Resnet-50层在NVIDIA A100 如上图所示。张量核心操作是使用CUDA的 mma指令。

4、兼容性

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、CUDA12.2.2、CUDA12.3.1和CUDA12.3.2。

4、操作系统

我们测试了以下环境。

操作系统	编译器
Ubuntu 18.04	GCC7.5.0
Ubuntu 20.04	GCC10.3.0
Ubuntu 22.04	GCC11.2.0
Ubuntu 22.04	Clang10.0.0
Ubuntu 22.04	Clang14.0.6
Ubuntu 22.04	Clang17.0.6
Windows 10.0	Visual Studio 2019 v16.11.27

注意：GCC 8.5.0 有关于折叠表达式和重载运算符的已知回归。建议使用GCC7.5.0 或（首选）GCC>=9。

5、硬件

CUTLASS在以下NVIDIA GPU上成功运行，预计它将在基于Volta、图灵、安培、Ada和Hopper架构的NVIDIA GPU上高效运行。

GPU	CUDA计算能力	CUTLASS-3所需的最低CUDA工具包
NVIDIA V100张量核心GPU	7.0	11.4
NVIDIA TitanV	7.0	11.4
NVIDIA GeForce RTX 2080 TI，2080,2070	7.5	11.4
NVIDIA T4	7.5	11.4
NVIDIA A100张量核心GPU	8.0	11.4
NVIDIA A10	8.6	11.4
NVIDIA GeForce RTX 3090	8.6	11.4
NVIDIA GeForce RTX 4090	8.9	11.8
NVIDIA L40	8.9	11.8
NVIDIA H100张量核心GPU	9.0	11.8

6、目标架构

一般来说，为一个目标架构生成的PTX代码可以在未来的架构上运行（即，它是前向兼容的）。
然而，CUDA 12.0引入了 architecture-accelerated 特征的概念，其PTX没有前向兼容保证。
几个Hopper PTX指令属于这一类 architecture-accelerated特征，因此需要一个sm_90a的目标架构（注意所附的“a”）。
有关此指令和其他architecture-accelerated指令的更多详细信息，请参阅 CUDA文档。

目标架构信息通过cmake标志 CUTLASS_NVCC_ARCHS 传递给CUTLASS。
为了最大限度地提高 Hopper GH100 的性能，用户需要以 90a作为目标架构构建CUTLASS。
如果用户不小心构建了一个使用 SM90a功能（例如Hopper Tensor Core Instructions）的内核，使用SM90目标（注意缺少“a”），CTK 12或11.8，内核可能会因运行时错误而失败。

cmake .. -DCUTLASS_NVCC_ARCHS="90a"

有关哪些内核需要哪些目标架构的详细信息，请参阅功能文档。

7、文档

CUTLASS在以下文件和随附的文件中进行了描述 Doxygen 文档。

快速入门指南-构建和运行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项目的组织和内容
术语-描述代码中使用的术语
编程指南-编写高效现代CUDAC++的指南
基本类型-描述CUTLASS中用于表示数值和数组的基本C++类
布局-描述内存中矩阵和张量的布局
Tile 迭代器-描述了在内存中迭代矩阵瓦片的C++概念
CUTLASS Profiler-命令行驱动的分析应用程序
CUTLASS实用程序-用于促进快速开发的附加模板

8、资源

我们还描述了高效GEMM的结构 2018 GPU技术大会。

CUTLASS：CUDA中所有级别和尺度的稠密线性代数的软件基元
在NVIDIA A100上开发CUDA内核以将张量内核推向绝对极限
在CUTLASS中使用张量核加速卷积
通过增加CUTLASS中的张量核心利用率来加速向后数据梯度
CUTLASS：Python API、增强功能和NVIDIA Hopper

9、构建 CUTLASS

CUTLASS是一个仅标题的模板库，不需要构建即可供其他用户使用客户端应用程序应该以CUTLASS的include/目录为目标路径。

CUTLASS单元测试、示例和实用程序可以使用CMake构建。
CMake的最低版本在快速入门指南中给出。
确保CUDACXX环境变量指向安装的CUDA工具包中的NVCC 在你的系统上。

$ export 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 build && cd build$ cmake .. -DCUTLASS_NVCC_ARCHS=80               # compiles for NVIDIA's Ampere Architecture

在build/目录中，通过使用make构建目标test_unit来编译并运行CUTLASS单元测试。

单元测试被组织为几个反映CUTLASS顶级命名空间的二进制文件，它们可以通过make的-j命令行参数并行执行。

$ make test_unit -j
...
...
...
[----------] Global test environment tear-down
[==========] 946 tests from 57 test cases ran. (10812 ms total)
[  PASSED  ] 946 tests.

所有测试都应该通过支持的平台，尽管确切的考试数量可能会随着时间的推移而变化。

10、项目结构

CUTLASS与实用程序、工具、示例和单元测试一起被安排为仅标头库。
Doxygen 文档提供了一个完整的文件列表，类，和CUTLASS项目中定义的模板概念。

源代码组织的详细说明可以在 CUTLASS文档，但几个主要组成部分总结如下。

11、CUTLASS模板库

include/                     # client applications should target this directory in their build's include pathscutlass/                   # CUDA Templates for Linear Algebra Subroutines and Solvers - headers onlyarch/                    # direct exposure of architecture features (including instruction-level GEMMs)conv/                    # code specialized for convolutionepilogue/                # code specialized for the epilogue of gemm/convolutiongemm/                    # code specialized for general matrix product computationslayout/                  # layout definitions for matrices, tensors, and other mathematical objects in memoryplatform/                # CUDA-capable Standard Library componentsreduction/               # bandwidth-limited reduction kernels that do not fit the "gemm" modelthread/                  # simt code that can be performed within a CUDA threadtransform/               # code specialized for layout, type, and domain transformations*                        # core vocabulary types, containers, and basic numeric operationscute/                      # CuTe Layout, layout algebra, MMA/Copy atoms, tiled MMA/Copyalgorithm/               # Definitions of core operations such as copy, gemm, and operations on cute::tuplesarch/                    # Bare bones PTX wrapper structs for copy and math instructionsatom/                    # Meta-information either link to or built from arch/ operatorsmma_atom.hpp           # cute::Mma_Atom and cute::TiledMmacopy_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 SDK示例应用CUTLASS模板来实现基本计算。

工具

tools/library/                   # CUTLASS Instance Library - contains instantiations of all supported CUTLASS templatesinclude/cutlass/library/profiler/                  # CUTLASS Profiler         - command-line utility for executing operations in the#                            CUTLASS Libraryutil/                      # CUTLASS Utilities        - contains numerous helper classes forinclude/                 #                            manging tensors in device memory, referencecutlass/               #                            implementations for GEMM, random initializationutil/                #                            of tensors, and I/O.

测试

test/unit/目录包含了用Google Test实现的单元测试 Core API组件的基本用法和CUTLASS GEMM计算的完整测试。

构建和运行单元测试的说明在快速入门指南中描述。

12、性能分析

tools/profiler/ 目录包含一个命令行实用程序，用于启动每个GEMM内核。它可以构建如下：

$ make cutlass_profiler -j16

13、构建所有GEMM和卷积内核（构建时间长）

默认情况下，每种数据类型、数学指令和布局只实例化一个磁贴大小。
要实例化所有内容，请在从空build/目录运行CMake时设置以下环境变量。
请注意，这会导致数以万计的内核和较长的构建时间。这也会导致较大的二进制大小，并且在某些平台上链接器在构建库时失败。
因此，强烈建议只生成内核的子集，如下面的小节所示。

$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all
...
$ make cutlass_profiler -j16

14、构建GEMM和卷积内核的子集（减少构建时间）

要严格编译一个内核或一小部分内核，可以使用带有通配符的逗号分隔内核名称列表来减少内核集。以下示例显示了为NVIDIA Ampere和图灵架构构建一个或一个内核子集：

构建子集Tensor Core GEMM内核

要编译具有FP32累积和针对NVIDIA Ampere和图灵架构的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...
=============================Problem ID: 1Provider: CUTLASSOperationKind: gemmOperation: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8Status: SuccessVerification: ONDisposition: Passedreference_device: PassedcuBLAS: PassedArguments: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:column --B=f16:row --C=f32:column --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=1024Bytes: 118489088  bytesFLOPs: 115992428544  flopsRuntime: 1.55948  msMemory: 70.7616 GiB/sMath: 74378.8 GFLOP/s=============================
...

构建一个CUDA Core GEMM内核

要编译一个针对NVIDIA Ampere和图灵架构的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=============================Problem ID: 1Provider: CUTLASSOperationKind: gemmOperation: cutlass_simt_sgemm_128x128_8x2_nn_align1Status: SuccessVerification: ONDisposition: PassedcuBLAS: PassedArguments: --m=3456 --n=4096 --k=4096 --A=f32:column --B=f32:column --C=f32:column --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=1024Bytes: 180355072  bytesFLOPs: 115992428544  flopsRuntime: 6.73655  msMemory: 24.934 GiB/sMath: 17218.4 GFLOP/s=============================

构建张量核心卷积内核的子集

要编译实现前向传播（fprop）的Tensor核心卷积内核子集，使用FP32累加和针对NVIDIA Ampere和图灵架构的FP16输入，请使用以下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...
=============================Problem ID: 1Provider: CUTLASSOperationKind: conv2dOperation: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwcStatus: SuccessVerification: ONDisposition: Passedreference_device: PassedArguments: --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 --Activation=f16:nhwc --Filter=f16:nhwc --Output=f32:nhwc  \--conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1  \--eq_gemm_provider=none --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=1024Bytes: 1130659840  bytesFLOPs: 118482796544  flopsRuntime: 0.711496  msMemory: 1479.99 GiB/sMath: 166526 GFLOP/s=============================
...

构建一个卷积CUDA内核

要编译和运行一个实现前向传播（fprop）的CUDA Core卷积内核，该内核具有F32累加和针对NVIDIA Ampere和图灵架构的FP32输入，请使用以下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=3=============================Problem ID: 1Provider: CUTLASSOperationKind: conv2dOperation: cutlass_simt_sfprop_optimized_128x128_8x2_nhwcStatus: SuccessVerification: ONDisposition: Passedreference_device: PassedArguments: --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 --Activation=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc  \--conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1  \--eq_gemm_provider=none --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=1024Bytes: 2055798784  bytesFLOPs: 118482796544  flopsRuntime: 7.34266  msMemory: 260.752 GiB/sMath: 16136.2 GFLOP/s=============================

15、有关编译CUTLASS内核和 CUTLASS Profiler的更多详细信息

- GEMM CMake示例
- 隐式GEMM卷积CMake示例
此处描述了有关CUTLASS Profiler的更多详细信息。

关于

CUTLASS由NVIDIA Corporation作为开源软件发布 3条款“新”BSD许可证。

贡献者

CUTLASS开发者和贡献者的官方列表可在此处获得：贡献者。

版权所有

  Redistribution and use in source and binary forms, with or withoutmodification, are permitted provided that the following conditions are met:1. Redistributions of source code must retain the above copyright notice, thislist 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 documentationand/or other materials provided with the distribution.3. Neither the name of the copyright holder nor the names of itscontributors may be used to endorse or promote products derived fromthis 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, THEIMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE AREDISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLEFOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIALDAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS ORSERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVERCAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USEOF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.