概述

CUTLASS 4.5.0

CUTLASS 4.5.0 - 2026年3月

CUTLASS是一套用于在CUDA中实现高性能矩阵乘法（GEMM）及相关计算的抽象库，覆盖了所有层次和规模。它结合了分层分解和数据移动策略，将这些“可移动部分”拆解为可重用、模块化的软件组件和抽象。

针对概念性并行化层次中的不同层级，可以通过自定义分块大小、数据类型及其他算法策略进行专门化和调优。这种灵活性使得它们能够更简便地作为自定义内核和应用程序中的构建模块使用。

自2017年以来，CUTLASS一直提供用于高性能线性代数的CUDA C++模板抽象，并广泛支持多种计算任务，包括混合精度计算、专用数据移动（异步复制）以及适用于FP64、FP32、TF32、FP16、BF16等的数据累加抽象；还支持通过张量核心指令实现的FP32模拟计算[参见GitHub仓库：NVIDIA/cutlass/tree/main/examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm]；此外，还支持8位浮点类型（e5m2和e4m3）、块尺度数据类型（如NVIDIA NVFP4及OCP标准MXFP4、MXFP6、MXFP8）、窄整数类型（4位和8位有符号与无符号整数）以及二进制1位数据类型（在架构原生支持的情况下），上述功能均能在NVIDIA的Volta、Turing、Ampere、Ada、Hopper和Blackwell架构上运行。

为了进一步丰富这一基于C++的内核编程抽象生态系统，CUTLASS 4引入了CUTLASS DSL。这是一种Python原生接口，允许开发者基于CUTLASS和CuTe的核心概念编写高性能CUDA内核，且不会带来任何性能损失。这不仅显著降低了学习曲线，还将编译时间缩短了多个数量级，实现了与深度学习框架的原生集成而无需编写胶水代码，同时提供了更加直观的元编程方式，无需深厚的C++专业知识。

总体而言，我们设想CUTLASS DSL将成为一系列领域特定语言（DSL）。随着4.0版本的发布，我们首先推出了CuTe DSL。这是一种低层次编程模型，完全兼容CuTe C++抽象——暴露了布局、张量、硬件原子等核心概念，并对硬件线程和数据层次结构拥有完全控制权。

CuTe DSL展示了针对NVIDIA Ampere、Hopper和Blackwell架构所实现的可编程高吞吐量张量核心的最佳矩阵乘法和其他线性代数操作。

我们相信，它将成为学生、研究人员和性能工程师不可或缺的工具——帮助他们降低GPU编程的学习难度，快速原型化内核设计，并将优化后的解决方案投入生产。

CuTe DSL目前处于公开测试阶段，预计将于2025年夏季末结束测试并正式发布。

欲快速入门，请参考：

• CUTLASS C++快速入门指南。
• CuTe DSL快速入门指南。

CUTLASS 4.5的新特性

CuTe DSL

• 错误修复与改进
  • 改进了源代码与性能分析/调试之间的关联性

CUTLASS C++

• 添加示例95，以支持绿色上下文SM分区
  • 允许在部分SM分配的流上启动GEMM。
• 修复了一些内核问题：
  • 修复了Blackwell SM100/SM120内核模板中l2_capacity=0的处理问题
  • 修复了CUTLASS的Clang构建问题
• 修复了一些性能分析工具的问题：
  • 为分块GEMM性能分析工具添加了缺失的参考内核
• 来自社区和CUTLASS团队的各种改进与修复。感谢所有提交PR的贡献者！
• 使用CUDA工具包13.2版本生成最优代码。

注意：已知CUTLASS 4.x版本在Windows平台上无法正常构建，无论使用哪个CUDA工具包。CUTLASS团队正在努力修复此问题。

有关所有过往版本及更新的详细信息，请参阅CHANGELOG。

性能

CUTLASS提供的原语非常高效。当用于构建设备级别的GEMM内核时，它们几乎可以达到理论峰值吞吐量的充分利用。下图显示了CUTLASS 3.8在NVIDIA Blackwell SM100架构GPU上运行时，针对不同输入输出数据类型的理论峰值利用率百分比。

以下两张图展示了自CUTLASS 3.1以来，在NVIDIA H100（NVIDIA Hopper架构）上CUTLASS性能的持续提升。CUTLASS 3.5.1是使用CUDA 12.5u1工具包编译的。张量核心操作是通过CUDA的mma和wgmma指令实现的。

CuTe

CUTLASS 3.0引入了一个新的核心库CuTe，用于描述和操作线程与数据的张量。CuTe是一组C++ CUDA模板抽象，用于定义和操作线程与数据的分层多维布局。CuTe提供了Layout和Tensor对象，能够紧凑地封装数据的类型、形状、内存空间和布局，同时为用户自动完成复杂的索引计算。这样一来，程序员只需关注算法的逻辑描述，而繁琐的底层管理工作则由CuTe代劳。借助这些工具，我们可以快速设计、实现和修改所有的密集线性代数运算。

CuTe的核心抽象是分层多维布局，它可以与数据数组组合起来表示张量。这种布局表示能力强大，几乎可以涵盖实现高效密集线性代数所需的全部内容。布局还可以通过函数式组合进行合并与操作，我们在此基础上构建了一整套常用的操作，例如分块和划分。

从CUTLASS 3.0开始，其模板中的GEMM层次结构全面采用了CuTe。这大大简化了设计过程，并提高了代码的可组合性和可读性。更多关于CuTe的文档可以在其专用文档目录中找到。

兼容性

最低要求：

• 架构：Volta（计算能力 7.0）
• 编译器：必须支持至少 C++17
• CUDA 工具包版本：11.4

CUTLASS 需要一个支持 C++17 的主机编译器，并且在使用 CUDA 12.8 工具包 构建时表现最佳。它也兼容 CUDA 11.4、CUDA 11.5、CUDA 11.6、CUDA 11.7、CUDA 11.8，以及所有其他 CUDA 12.x 版本。

操作系统

我们已测试了以下环境。

操作系统	编译器
Ubuntu 18.04	GCC 7.5.0
Ubuntu 20.04	GCC 10.3.0
Ubuntu 22.04	GCC 11.2.0

注意：GCC 8.5.0 在折叠表达式和重载运算符方面存在已知的回归问题。建议使用 GCC 7.5.0 或（更优）GCC >= 9。

注意：已知 CUTLASS 3.x 在 Windows 平台上无法构建，无论使用哪个 CUDA 工具包。CUTLASS 团队正在努力修复此问题。

硬件

CUTLASS 可以在以下 NVIDIA GPU 上成功运行，并且预计在基于 Volta、Turing、Ampere、Ada 和 Hopper 架构的 NVIDIA GPU 上具有高效性能。

GPU	CUDA 计算能力	CUTLASS-3 所需的最低 CUDA 工具包版本
NVIDIA V100 Tensor Core GPU	7.0	11.4
NVIDIA TitanV	7.0	11.4
NVIDIA GeForce RTX 20x0 系列	7.5	11.4
NVIDIA T4	7.5	11.4
NVIDIA A100 Tensor Core GPU	8.0	11.4
NVIDIA A10	8.6	11.4
NVIDIA GeForce RTX 30x0 系列	8.6	11.4
NVIDIA GeForce RTX 40x0 系列	8.9	11.8
NVIDIA L40	8.9	11.8
NVIDIA H100 Tensor Core GPU	9.0	11.8
NVIDIA H200 Tensor Core GPU	9.0	11.8
NVIDIA B200 Tensor Core GPU	10.0	12.8
NVIDIA B300 Tensor Core GPU	10.3	13.0
NVIDIA DRIVE Thor	11.0	13.0
NVIDIA GeForce RTX 50x0 系列	12.0	12.8
NVIDIA DGX Spark	12.1	13.0

目标架构

一般来说，为某一目标架构生成的 PTX 代码可以在未来的架构上运行（即向前兼容）。然而，CUDA 12.0 引入了“架构加速特性”的概念，其 PTX 不具备向前兼容性保证。Hopper 和 Blackwell 的一些 PTX 指令属于这一类架构加速特性，因此需要指定 sm_90a 或 sm100a 作为目标架构（注意末尾的“a”）。有关此类及其他架构加速指令的详细信息，请参阅 CUDA 文档。

目标架构信息通过 CMake 标志 CUTLASS_NVCC_ARCHS 传递给 CUTLASS。为了在 Hopper GH100 上实现最佳性能，用户需要将目标架构设置为 90a 来构建 CUTLASS。如果用户不小心构建了一个使用 SM90a 特性（例如 Hopper Tensor Core 指令）的内核，并使用 SM90 目标架构（缺少“a”），无论使用 CUDA 工具包 12 还是 11.8，该内核都可能会因运行时错误而失败。

cmake .. -DCUTLASS_NVCC_ARCHS="90a"

或者

cmake .. -DCUTLASS_NVCC_ARCHS="100a"

注意：数据中心产品中使用的 NVIDIA Blackwell SM100 架构与支撑 NVIDIA Blackwell GeForce RTX 50 系列 GPU 的架构（SM120）具有不同的计算能力。因此，为 Blackwell SM100 架构编译并包含架构条件特性的内核（使用 sm100a）与 RTX 50 系列 GPU 不兼容。

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

文档

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 中的 2D 和 3D 卷积
• 代码组织 - 描述了 CUTLASS 项目的组织结构和内容
• 术语 - 解释了代码中使用的术语
• 编程指南 - 提供了编写高效现代 CUDA C++ 的指导原则
• 基础类型 - 描述了 CUTLASS 中用于表示数值和数组的基本 C++ 类
• 布局 - 描述了矩阵和张量在内存中的布局
• 分块迭代器 - 描述了用于在内存中遍历矩阵分块的 C++ 概念
• CUTLASS 性能分析工具 - 基于命令行的性能分析应用程序
• CUTLASS 实用工具 - 提供了额外的模板，以促进快速开发
• 依赖内核启动 - 描述了 Hopper 中的一项新特性，允许在同一流中重叠执行依赖内核，以及它在 CUTLASS 中的应用。

资源

我们还在2018年GPU技术大会上发表的演讲中描述了高效GEMM的结构： GPU技术大会2018。

• CUTLASS：CUDA中各层次和规模的密集线性代数软件原语
• 开发CUDA内核以将Tensor Cores推至NVIDIA A100的极限
• 使用CUTLASS中的Tensor Cores加速卷积
• 通过提高CUTLASS中Tensor Cores的利用率来加速反向数据梯度计算
• CUTLASS：Python API、增强功能以及NVIDIA Hopper

构建CUTLASS

CUTLASS是一个仅包含头文件的模板库，无需构建即可被其他项目使用。客户端应用程序应在其包含路径中指定CUTLASS的include/目录。

CUTLASS的单元测试、示例和实用工具可以通过CMake进行构建。CMake的最低版本在快速入门指南中给出。请确保CUDACXX环境变量指向您系统上安装的CUDA Toolkit中的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               # 为NVIDIA的Ampere架构编译

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

单元测试被组织成几个二进制文件，分别对应CUTLASS的顶级命名空间，可以通过make的-j命令行参数并行执行。

$ make test_unit -j
...
...
...
[----------] 全局测试环境清理
[==========] 共有57个测试用例中的946个测试已运行。（总耗时10812毫秒）
[  PASSED  ] 946个测试。

在支持的平台上，所有测试都应通过，尽管测试的具体数量可能会随时间变化。

项目结构

CUTLASS由一个仅包含头文件的库以及实用工具、工具、示例和单元测试组成。Doxygen文档提供了CUTLASS项目中定义的文件、类和模板概念的完整列表。

关于源代码组织的详细说明可以在CUTLASS文档中找到，但以下几个主要组件在此简要概述。

CUTLASS模板库

include/                     # 客户端应用程序应在构建的包含路径中指定此目录

  cutlass/                   # 线性代数子程序和求解器的CUDA模板——仅包含头文件

    arch/                    # 直接暴露架构特性（包括指令级GEMM）

    conv/                    # 针对卷积优化的代码

    epilogue/                # 针对GEMM/卷积后处理阶段优化的代码

    gemm/                    # 针对通用矩阵乘法计算优化的代码

    layout/                  # 矩阵、张量及其他数学对象在内存中的布局定义

    platform/                # 具备CUDA功能的标准库组件

    reduction/               # 带宽受限的归约内核，不符合“gemm”模型

    thread/                  # 可在CUDA线程内执行的SIMT代码

    transform/               # 针对布局、类型和域转换优化的代码

    *                        # 核心词汇类型、容器和基本数值运算

  cute/                      # CuTe布局、布局代数、MMA/Copy原子、分块MMA/Copy

    algorithm/               # 对核心操作的定义，如复制、GEMM以及对cute::tuple的操作

    arch/                    # 复制和数学指令的裸PTX包装结构体

    atom/                    # 与arch/运算符相关联或基于arch/运算符构建的元信息

      mma_atom.hpp           # cute::Mma_Atom和cute::TiledMma

      copy_atom.hpp          # cute::Copy_Atom和cute::TiledCopy

      *sm*.hpp               # 针对复制和数学操作的特定架构元信息

    *                        # 核心库类型，如Shape、Stride、Layout、Tensor及其相关操作

CUTLASS SDK示例

CUTLASS SDK示例应用CUTLASS模板来实现基本计算。

工具

tools/
  library/                   # CUTLASS实例库——包含所有受支持的CUTLASS模板的实例化版本
    include/
      cutlass/
        library/

  profiler/                  # CUTLASS性能分析工具         ——用于在CUTLASS库中执行操作的命令行实用程序

  util/                      # CUTLASS实用工具        ——包含大量帮助类，用于
    include/                 #                            管理设备内存中的张量、GEMM的参考实现、
      cutlass/               #                            张量的随机初始化以及输入输出。
        util/

测试

test/unit/目录包含使用Google Test实现的单元测试，展示了核心API组件的基本用法以及对CUTLASS GEMM计算的全面测试。

构建和运行单元测试的说明在快速入门指南中有所介绍。

性能剖析

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

$ make cutlass_profiler -j16

构建所有 GEMM 和卷积核（耗时极长）

默认情况下，对于每种数据类型、数学指令和布局，仅实例化一个分块大小。 若要实例化所有分块大小，请在从空的 build/ 目录运行 CMake 时设置以下环境变量。 请注意，这将生成 数以万计 的内核，并导致漫长的构建时间。 此外，还会使二进制文件体积庞大，在某些平台上可能导致链接器在构建库时失败。 因此，强烈建议仅生成一部分内核，如下一小节所示。

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

构建部分 GEMM 和卷积核（缩短构建时间）

若要严格编译单个内核或少量内核，可以使用逗号分隔的内核名称列表，并配合通配符来缩小内核集合。以下示例展示了为 NVIDIA Ampere 和 Turing 架构精确编译单个或部分内核的方法：

构建部分 Tensor Core GEMM 内核

若要编译针对 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

          状态：成功
    验证：开启
     处理结果：通过

reference_device：通过
          cuBLAS：通过

       参数：--gemm_kind=universal --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

           字节数：118,489,088 字节
           浮点运算次数：115,992,428,544 次浮点运算

         运行时间：1.55948 毫秒
          内存带宽：70.7616 GiB/s

            计算强度：74,378.8 GFLOP/s



=============================
...

构建单个 CUDA Core GEMM 内核

若要编译针对 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

           字节数：180,355,072 字节
           浮点运算次数：115,992,428,544 次浮点运算

         运行时间：6.73655 毫秒
          内存带宽：24.934 GiB/s

            计算强度：17,218.4 GFLOP/s

=============================

构建部分 Tensor Core 卷积核

若要编译针对 NVIDIA Ampere 和 Turing 架构、采用 FP32 累加且 FP16 输入、实现前向传播（fprop）的部分 Tensor Core 卷积核，请使用以下 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
   操作类型：2D 卷积
       操作：cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc

          状态：成功
    验证：开启
     处理结果：通过

reference_device：通过

       参数：--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=串行 --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

           字节数：1,130,659,840 字节
           浮点运算次数：118,482,796,544 次浮点运算

         运行时间：0.711496 毫秒
          内存带宽：1,479.99 GiB/s

            计算强度：166,526 GFLOP/s

=============================
...

构建一个卷积 CUDA 内核

要编译并运行一个使用 F32 累加和 FP32 输入、针对 NVIDIA Ampere 和 Turing 架构实现前向传播（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


=============================
  问题 ID：1

        提供者：CUTLASS
   操作类型：conv2d
       操作：cutlass_simt_sfprop_optimized_128x128_8x2_nhwc

          状态：成功
    验证：开启
     处理结果：通过

reference_device：通过

       参数：--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

           字节数：2055798784 字节
           浮点运算次数：118482796544 次浮点运算

         运行时间：7.34266 毫秒
          内存带宽：260.752 GiB/秒

            计算性能：16136.2 GFLOPS/秒


=============================

编译 CUTLASS 内核和 CUTLASS Profiler 的更多细节

• 请参阅以下链接，获取有关选择性编译 CUTLASS 内核的更多 CMake 示例：
  • GEMM CMake 示例
  • 隐式 GEMM 卷积 CMake 示例
• 关于 CUTLASS Profiler 的更多详细信息请见此处。

关于

CUTLASS 由 NVIDIA 公司以开源软件的形式发布，采用 三条款“新”BSD 许可证。

贡献者

CUTLASS 的官方开发者和贡献者列表可在以下位置找到：CONTRIBUTORS。

版权

版权所有 © 2017 - 2026 NVIDIA CORPORATION & AFFILIATES。保留所有权利。 SPDX 许可证标识：BSD-3-Clause

  重新分发和使用源代码及二进制文件，无论是否修改，均在满足以下条件时被允许：

  1. 源代码的再分发必须保留上述版权声明、本条件列表以及以下免责声明。

  2. 二进制形式的再分发必须在随附的文档或其他材料中复制上述版权声明、本条件列表以及以下免责声明。

  3. 未经事先书面许可，不得使用版权持有者的名称或其贡献者的名称来认可或推广由此软件衍生的产品。

  本软件按“原样”提供给版权所有者和贡献者，不提供任何明示或暗示的保证，包括但不限于适销性和特定用途适用性的暗示保证。在任何情况下，版权所有者或贡献者均不对直接、间接、偶然、特殊、示范性或后果性损害承担责任（包括但不限于替代品或服务的采购、使用损失、数据丢失、利润损失或业务中断）。即使已被告知发生此类损害的可能性，亦不承担任何责任。

CUTLASS 快速上手指南

CUTLASS (CUDA Templates for Linear Algebra Subroutines) 是 NVIDIA 开源的高性能矩阵乘法（GEMM）及相关计算的 C++ 模板库。它支持从 Volta 到 Blackwell 的全系列 NVIDIA GPU 架构，并提供 C++ 模板抽象和全新的 Python CuTe DSL 接口。

环境准备

在开始之前，请确保您的开发环境满足以下最低要求：

系统要求

• 操作系统: 推荐 Ubuntu 18.04/20.04/22.04。
  • 注意: 目前 Windows 平台构建已知存在问题，建议使用 Linux。
• 编译器: 支持 C++17 标准的宿主编译器。
  • 推荐：GCC 9.0 或更高版本（GCC 7.5.0 可用，但避免使用 GCC 8.5.0）。
• CUDA Toolkit:
  • 最低版本：11.4
  • 推荐版本：CUDA 12.8 (以获得最佳性能和 Blackwell 架构支持)。

硬件支持

CUTLASS 支持以下架构的 NVIDIA GPU：

• Volta (V100), Turing (T4, RTX 20xx), Ampere (A100, RTX 30xx), Ada (RTX 40xx), Hopper (H100), Blackwell (B200, RTX 50xx)。
• 重要提示: 针对 Hopper (SM90) 和 Blackwell (SM100/SM120) 架构的新特性（如异步 Warp 组矩阵指令），编译时必须指定带 a 后缀的目标架构（例如 90a, 100a），否则运行时会报错。

前置依赖

• CMake: 3.18 或更高版本。
• Git: 用于克隆代码库。
• (可选) Python: 如果您计划使用新的 CuTe DSL (Python 接口)，需安装 Python 3.8+ 及相关依赖。

---

安装步骤

1. 克隆仓库

git clone https://github.com/NVIDIA/cutlass.git
cd cutlass

2. 创建构建目录

mkdir build && cd build

3. 配置构建 (CMake)

根据您的目标 GPU 架构选择配置命令。

场景 A：通用构建 (兼容大多数旧架构)

cmake .. -DCUTLASS_NVCC_ARCHS="70;75;80;86"

场景 B：针对 Hopper (H100/H200) 优化

必须使用 90a 以启用架构加速特性。

cmake .. -DCUTLASS_NVCC_ARCHS="90a"

场景 C：针对 Blackwell (B200/RTX 50xx) 优化

数据中心卡 (SM100) 和消费级卡 (SM120) 需区分：

# 针对 B200/H200 等数据中心卡
cmake .. -DCUTLASS_NVCC_ARCHS="100a"

# 针对 RTX 5090/5080 等消费级卡
cmake .. -DCUTLASS_NVCC_ARCHS="120"

提示: 您可以组合多个架构，例如 -DCUTLASS_NVCC_ARCHS="80;90a;100a"。

4. 编译

利用多核加速编译过程（将 <jobs> 替换为您的 CPU 核心数，如 8 或 16）：

make -j<jobs>

(注：如果您仅需使用头文件或 Python DSL，通常无需完整编译所有示例，但构建测试用例可验证环境正确性。)

---

基本使用

CUTLASS 提供两种主要使用方式：传统的 C++ 模板 API 和新推出的 CuTe DSL (Python)。

方式一：C++ 模板 API (高性能生产环境)

CUTLASS 主要通过 C++ 模板实例化来生成内核。以下是一个最简单的调用现有 GEMM 操作的逻辑示例。

代码示例 (simple_gemm.cu):

#include "cutlass/gemm/device/gemm.h"

int main() {
    // 定义 GEMM 操作类型
    using Gemm = cutlass::gemm::device::Gemm<
        float,                              // ElementA
        cutlass::layout::RowMajor,          // LayoutA
        float,                              // ElementB
        cutlass::layout::RowMajor,          // LayoutB
        float,                              // ElementC
        cutlass::layout::RowMajor,          // LayoutC
        float                               // ElementAccumulator
    >;

    // 定义问题规模 (M, N, K)
    Gemm::Arguments args{
        {512, 512, 512},                    // Problem Size
        {nullptr, 512},                     // Ptr A & Leading Dim
        {nullptr, 512},                     // Ptr B & Leading Dim
        {nullptr, 512},                     // Ptr C & Leading Dim
        {nullptr, 512},                     // Ptr D & Leading Dim
        {1.0f, 0.0f}                        // Epilogue: alpha, beta
    };

    // 实例化并检查支持性
    Gemm gemm_op;
    auto status = gemm_op.can_implement(args);
    if (status != cutlass::Status::kSuccess) {
        return -1;
    }

    // 启动计算 (实际使用时需分配显存并传入指针)
    status = gemm_op(args);
    
    return (status == cutlass::Status::kSuccess) ? 0 : -1;
}

编译命令:

nvcc -std=c++17 -I../include simple_gemm.cu -o simple_gemm

方式二：CuTe DSL (Python 快速原型)

CUTLASS 4.x 引入了 Python 原生接口，无需编写 C++ 胶水代码即可定义高性能内核。

前置准备: 确保已安装相关 Python 依赖（通常在 python/ 目录下有 setup.py 或 requirements.txt）。

代码示例 (simple_gemm.py):

import cute
import cutlass

# 定义布局 (Layout) 和张量 (Tensor) 概念
# 这里演示如何声明一个简单的矩阵乘法逻辑结构
def run_gemm():
    # 定义矩阵维度
    M, N, K = 128, 128, 64
    
    # 使用 CuTe DSL 定义操作 (伪代码示意，具体 API 随版本迭代可能微调)
    # 详细用法请参考官方 CuTe DSL Quick Start
    print(f"Initializing GEMM for {M}x{N}x{K} using CuTe DSL...")
    
    # 在实际环境中，这里会调用 cutlass.dsl 下的具体算子构建器
    # 例如: kernel = cutlass.dsl.gemm(...)
    
    print("Kernel definition complete. Ready for compilation/execution.")

if __name__ == "__main__":
    run_gemm()

进阶学习:

• C++ 详细文档：CUTLASS C++ Quick Start
• Python DSL 教程：CuTe DSL Quick Start
• 示例代码库：查看 examples/ 目录，其中包含针对不同架构（如 95_blackwell_gemm）的具体实现。