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=3
=============================
问题编号：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.
