弯刀 3.5 - 2024 年 4 月
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.5 是 CUTLASS 的更新，添加了：
- 通过 WGMMA + TMA im2col针对 Hopper SM90A 的隐式 GEMM 卷积。
- 使用 CuTe 在 CUTLASS 3.x 中进行本机实现，镜像与GEMM 相同的设计层次结构。
- 以与等级无关的方式支持 1D、2D 和 3D 卷积。
- 支持Fprop、Dgrad和Wgrad算法。
- CUTLASS 分析器支持通过 3.x API 实现的 2D 和 3D 卷积。
- 注意：这是测试版。 CUTLASS 的进一步更新将包括重大性能改进、功能启用以及在 3.7 版本发布之前可能对 API 进行的重大更改。欢迎您对设计提出反馈！
- 通过 2.x API支持Ada (SM89) FP8 张量核心。需要 CUDA 12.4 或更高版本。
- CuTe 和 CUTLASS 3.x 中的
安培聚集/散射卷积示例。
- 展示如何使用 CUTLASS 3.x 和 CuTe 编写和优化自定义内核，以及将卷积实现为 GETT 专业化的一般策略。
- 实施粗粒度稀疏聚集/分散内核，在安培级张量核心上实现峰值性能。
- CUTLASS 2.x 中添加了 32x 和 16x 切片尺寸，以提高窄高和宽短矩阵的性能。
- 更新了MMA 原子
cute::Tensor<>的 CuTe 文档，以及经过彻底修改的CuTe GEMM 教程系列。 - CuTe 的扩展以支持L2 预取和TMA 存储+缩减。
- 删除了一些 CUTLASS 2.x API 头文件的 C++11 要求。所有 CUTLASS 文件现在都需要 C++17。
- 修复以大大减少构建警告。
- 来自社区的更新和错误修复（谢谢！）
最低要求：
- 建筑：沃尔特
- 编译器：必须至少支持 C++17
- CUDA 工具包版本：11.4
从 CUTLASS 3.0 开始，CUTLASS 删除了对以下内容的支持：
- Maxwell 和 Pascal GPU 架构
- 乌班图16.04
- CUDA 10.2
- C++ 语言版本低于 17。
有关版本和更新的详细列表，请参阅变更日志。
CUTLASS 原语非常高效。当用于构建设备范围的 GEMM 内核时，它们在标量 GEMM 计算方面表现出与 cuBLAS 相当的峰值性能。上图显示了NVIDIA H100（NVIDIA Hopper 架构）、NVIDIA L40（NVIDIA Ada 架构）、NVIDIA A100（NVIDIA Ampere 架构）
和NVIDIA A40 （NVIDIA Ampere 架构）上大矩阵维度下 CUTLASS 相对于 cuBLAS 的性能。 CUTLASS 3.0 是使用CUDA 12.0 工具包编译的。 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 CUTLASS_NVCC_ARCHS。为了最大限度地提高 Hopper GH100 的性能，用户需要构建 CUTLASS 作为90a目标架构。如果用户意外地使用 SM90 目标（注意缺少“a”）以及 CTK 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 实用程序- 用于促进快速开发的附加模板
我们还在2018 年 GPU 技术大会上的演讲中描述了高效 GEMM 的结构 。
- CUTLASS：CUDA 内所有级别和规模的密集线性代数的软件基元
- 开发 CUDA 内核将 Tensor Core 推向 NVIDIA A100 的绝对极限
- 在 CUTLASS 中使用张量核心加速卷积
- 通过增加 CUTLASS 中张量核心的利用率来加速后向数据梯度
- CUTLASS：Python API、增强功能和 NVIDIA Hopper
CUTLASS 是一个仅包含头文件的模板库，不需要构建即可供其他项目使用。客户端应用程序应include/在其包含路径中以 CUTLASS 目录为目标。
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 的体系结构
CUTLASS_NVCC_ARCHS。
$ mkdir build && cd build$ cmake .. -DCUTLASS_NVCC_ARCHS=80 # compiles for NVIDIA's Ampere Architecture
从build/目录中，通过使用 make 构建目标来编译并运行 CUTLASS 单元测试test_unit。
单元测试被组织为几个二进制文件，镜像 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.
所有测试都应在支持的平台上通过，但测试的确切数量可能会随着时间的推移而变化。
CUTLASS 与实用程序、工具、示例和单元测试一起被安排为仅包含头文件的库。 Doxygen 文档提供了 CUTLASS 项目中定义的文件、类和模板概念的完整列表。
源代码组织的详细说明可以在 CUTLASS 文档中找到，但下面总结了几个主要组件。
include/ # client applications should target this directory in their build's include pathscutlass/ # 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 operationscute/ # 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
默认情况下，只会为每种数据类型、数学指令和布局实例化一个图块大小。要实例化所有内容，请在从空目录运行 CMake 时设置以下环境变量build/。请注意，这会导致数以万计的内核和较长的构建时间。这还会导致二进制大小过大，并且在某些平台上链接器无法构建库。因此，强烈建议仅生成内核的子集，如下面小节所示。
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ 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=4096Problem ID: 1
Provider: CUTLASSOperationKind: gemm Operation: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Status: Success Verification: ON Disposition: Passedreference_device: Passed cuBLAS: Passed
Arguments: --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=1024 Bytes: 118489088 bytes FLOPs: 115992428544 flops Runtime: 1.55948 ms Memory: 70.7616 GiB/s Math: 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============================= Problem ID: 1
Provider: CUTLASSOperationKind: gemm Operation: cutlass_simt_sgemm_128x128_8x2_nn_align1
Status: Success Verification: ON Disposition: Passed cuBLAS: Passed Arguments: --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=1024 Bytes: 180355072 bytes FLOPs: 115992428544 flops Runtime: 6.73655 ms Memory: 24.934 GiB/s Math: 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=3Problem ID: 1
Provider: CUTLASSOperationKind: conv2d Operation: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Status: Success Verification: ON Disposition: Passedreference_device: Passed
Arguments: --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=1024 Bytes: 1130659840 bytes FLOPs: 118482796544 flops Runtime: 0.711496 ms Memory: 1479.99 GiB/s Math: 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============================= Problem ID: 1
Provider: CUTLASSOperationKind: conv2d Operation: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Status: Success Verification: ON Disposition: Passedreference_device: Passed
Arguments: --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=1024 Bytes: 2055798784 bytes FLOPs: 118482796544 flops Runtime: 7.34266 ms Memory: 260.752 GiB/s Math: 16136.2 GFLOP/s=============================
- 请点击以下链接，获取有关选择性编译 CUTLASS 内核的更多 CMake 示例：
- 有关 CUTLASS Profiler 的更多详细信息请参见此处。
CUTLASS 是 NVIDIA 公司根据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:
Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.
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.
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.