概述

CUTLASS 4.4.1

CUTLASS 4.4.1 - 2026年2月

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

针对概念上并行化层次结构的不同层级，其原语可以通过自定义的平铺大小、数据类型和其他算法策略进行专门化和调优。由此产生的灵活性简化了它们作为自定义内核和应用程序中构建块的使用。

自 2017 年以来，CUTLASS 一直为高性能线性代数提供 CUDA C++ 模板抽象，这些抽象为广泛的计算提供了广泛支持，包括混合精度计算、专门的数据移动（异步拷贝）以及针对 FP64、FP32、TF32、FP16、BF16、通过张量核心指令实现的 FP32 仿真、8位浮点类型（e5m2 和 e4m3）、块缩放数据类型（NVIDIA NVFP4 和 OCP 标准 MXFP4、MXFP6、MXFP8）、窄整数类型（4位和8位有符号及无符号整数）以及二进制1位数据类型（在架构允许原生支持此类数据类型的场景下），覆盖了 NVIDIA 的 Volta、Turing、Ampere、Ada、Hopper 和 Blackwell 架构。

在这个丰富的基于 C++ 的内核编程抽象生态系统之上，CUTLASS 4 新增了 CUTLASS DSLs。这是一套原生 Python 接口，用于基于核心 CUTLASS 和 CuTe 概念编写高性能 CUDA 内核，且不牺牲任何性能。这带来了更平缓的学习曲线、数量级更快的编译时间、与深度学习框架的原生集成（无需编写胶水代码），以及更直观的元编程体验，无需深厚的 C++ 专业知识。

总体而言，我们将 CUTLASS DSLs 设想为一个领域特定语言（DSL）家族。随着 4.0 版本的发布，我们推出了其中的第一个成员：CuTe DSL。这是一个低层级的编程模型，与 CuTe C++ 抽象完全一致——它暴露了布局（layouts）、张量（tensors）、硬件原子（hardware atoms）等核心概念，并提供对硬件线程和数据层次结构的完全控制。

CuTe DSL 展示了针对 NVIDIA Ampere、Hopper 和 Blackwell 架构实现的可编程、高吞吐量 张量核心 的优化矩阵乘法及其他线性代数运算。

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

CuTe DSL 目前处于公开测试阶段，预计将于 2025 年夏季结束前结束测试。

快速入门请参考：
- CUTLASS C++ 快速入门指南.
- CuTe DSL 快速入门指南.

CUTLASS 4.4 的新特性

CuTe DSL

• 新功能
• CuTe DSL 现在支持 CUDA 工具包 13.1！
  • 使用 cutlass/python/CuTeDSL/setup.sh --cu13 进行设置
  • 更多详情请参考 https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/quick_start.html
• CuTe DSL 现在支持 CTK 13.1 下的 GB300
  • 参考 SM103 批处理 3xFP4 块缩放 GEMM 内核 获取示例内核
• cute.experimental：在现有 CuTe DSL API 之上引入一个更高层级、可组合的层（并非独立的抽象），可以与现有的 Cute DSL 构建块混合使用。
  • 无片段编程模型：copy/dot API 直接接受内存引用（memrefs）而非描述符/片段。
  • 自动 TMA 描述符生成和更新插入。
  • 为 SIMT 拷贝自动向量化和谓词化。
  • 新的流水线抽象及便捷包装器。
  • 新的分区操作以简化分区逻辑。
  • 设备端 TMA 描述符的分配、初始化和管理。
  • 这些示例可在此处找到：https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/experimental
• 现在支持提前（Ahead of Time，AoT）编译！
  • 参考 https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/cute/export 下的文件获取使用示例
• JAX 支持 - 您现在可以将 CuTeDSL 与 JAX 一起使用
  • 参考 https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/jax 下的文件获取使用示例
• 在 DSL 中引入版本支持：
  • cutlass.__version__ 用于获取 DSL 版本的字符串表示
  • cutlass.CUDA_VERSION 用于获取一个版本类，以告知 DSL 使用的 CUDA 版本
• 新增 CopyDsmemStoreOp，用于将数据存储到分布式共享内存并执行显式同步。
• 分组 GEMM 示例现在支持仅设备端的问题形状。
• 允许在没有主机端问题形状的情况下进行网格划分（grid carve-out）。
• 用于加载+解包窄宽度类型的 Tma+LdMatrix 功能（参考 mixed_input_fmha_decode.py 获取使用示例）。

• 现在可以通过 Python Epilogue Fusion Configuration (EFC) 函数为持久化密集 GEMM 定制化尾融合（epilogue fusion），某种程度上类似于 CUTLASS C++ 的 EVT。它还提供了一个 PyTorch 评估器来比较结果。

• 更多编写峰值性能内核的示例

• SM103 批处理 3xFP4 块缩放 GEMM 内核
• 支持 int4 KV 的混合输入 FMHA 解码示例（int8 KV 在 4.3 版本中已支持）
• 引入了新的 acc_scale 分组混合输入 gemm 内核变体，为解码场景提供更好的性能。

• 所有 mixed_input_gemm 示例已移至单独的文件夹 mixed_input_gemm。通用工具函数也已提取到同一文件夹下的 mixed_input_host_utils.py 中。

• 错误修复与改进

• 修复了 if 语句两个分支都被执行的问题
• 修复了 cute.printf 使用 f-string 的问题
• 修复了标量张量的索引问题
• 修复了 Blackwell SM100 静态调度持久化密集块缩放 GEMM 中，在 cta_tile_n = 256 且启用累加器重叠优化时，小 K 参考检查错误的问题。

• 修复了在 aarch64 架构上使用 tvm-ffi 时的段错误问题

• API 变更

• 弃用 blackwell_helpers.py 中的 get_num_tmem_alloc_cols。请改用 tmem_allocator.py 中的版本。
• 弃用 SM100_TMEM_CAPACITY_COLUMNS 和 SM100_TMEM_MIN_ALLOC_COLUMNS。
• LdMatrix16x16x8bOp 和 StMatrix16x8x8bOp 现在在调用 __init__ 时需要显式指定 transpose=True，以避免数据转置的歧义。
• LdMatrix16x16x8bOp 的拷贝特性已更新，以忠实于 PTX 指令，不进行置换。置换变体重命名为 LdMatrix16x8x8bOp。
• 分组 GEMM 示例接受参数 --host_problem_shape_available。如果提供了该参数，则基于主机端问题形状进行网格划分；否则，我们将启动尽可能多的 SM。
• hardware_info.get_max_active_cluster 支持传入特定流进行查询。适用于基于绿色上下文（green context）的 SM 分区。
• 异步批量拷贝示例中的 group_bulk_copy_modes 现已弃用，请直接使用 group_modes。
• 弃用 nvvm 包装器使用 nvvm 枚举，改用字符串。
• cute.arch.calc_packed_f32x2_op 默认启用 ftz 改为默认禁用 ftz。

• 在使用 CTK 13.1 的 CuTe DSL 中，cutlass.cute.arch 中的以下 API 现在要求使用字符串字面量作为参数，而不是枚举：

  • fence_proxy
  • fence_view_async_tmem_op
  • calc_packed_f32x2_op
  • warp_redux_sync
  • atomic_add
  • atomic_and
  • atomic_or
  • atomic_xor
  • atomic_max
  • atomic_min
  • atomic_exch
  • atomic_cas
  • store
  • load

• 为混合输入 gemm 示例使用“高级控制文件”以获得更好性能。

• 高级控制文件是 CUDA 编译器的一个实验性功能。该控制文件包含针对特定内核和特定版本 CUDA 工具包调优的内部编译器设置，以获得更好的 GPU 内核代码。关于如何创建这些控制文件的更多细节和文档将在未来的 CUDA 工具包版本中提供。注意：高级编译器控制文件不适用于未经调优的内核。没有兼容性保证，且控制文件无法用于不同版本的 CUDA 工具包。

CUTLASS C++

• 新增 示例 93 用于 Blackwell 低延迟生成阶段 GQA 内核。
  • 包含集群归约的 Flash Decoding。
  • 内核设计详情请查看 Readme。
• 在 示例 112 中添加 Blackwell SM100 状态空间分解（SSD）内核。
• 在 示例 111 中添加 Hopper SM90 状态空间分解（SSD）内核。
• 添加 Hopper e2m1 到 fp32 的优化转换以及 e2m1 * TF32 张量核心 GEMM。
  • 为 示例 55 启用 TF32 支持
• 新增 示例 94 用于 Ada FP8xFP8 -> BF16 GEMM，在 MMA 循环中对输入矩阵进行块式反量化，并使用 FP32 累加。
• 为块缩放张量添加对任意应用程序提供的步幅的支持。
  • 用户和应用程序现在在所有情况下都必须传递有效的块缩放步幅，即使张量是打包的。
• 支持 CUDA 13.1 的 4x 块缩放公共 ptx。
• 允许 AuxTmaParams 中使用非静态的 TmaGbasis。
  • 注意力内核中的某些情况可能需要非静态的 tma_gbasis。
  • 放宽了对 AuxTmaParams 的 TmaGbasis 参数的限制，允许用户手动构造动态的 gbasis。
• 修复一些内核问题：
  • 修复 MSVC 预处理问题。
  • 修复 GEMV 内核中的自赋值问题。
  • 修复 TMA 描述符的一个 bug，该 bug 导致 CUDA 驱动程序未能正确设置 OOB 地址生成模式。
  • 修复 Blackwell SM120 pingpong 内核中 clc 调度器的内存栅栏问题。
  • 修复 Blackwell SM120 缩放因子中缺失的 SMEM 对齐问题。
  • 修复分组 gemm 的一个 PDL 问题。
  • 修复 sm100 隐式 gemm 内核的 canimplement 中的除零问题。
  • 修复分组 GEMM 的集群切换（cluster swizzle）问题。
    • 将主机端的切换启发式方法移至设备端。
    • 根据问题形状和最大切换大小对每个组应用切换。
    • 改进示例和单元测试。
• 修复一些性能分析器问题：
  • 修复 nvfp4 分组 GEMM 内核的核心转储问题。
  • 修复不一致的 GEMM 验证逻辑。
  • 为不同类型重新设计分组 gemm 验证逻辑。
  • 修复使用 nvMatmulHeuristics 时的 API 破坏性变更问题。
• 修复 media/docs 下的一些失效链接。

注意：已知 CUTLASS 4.x 版本在所有 CUDA 工具包下的 Windows 平台上构建失败。
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 版本兼容。

操作系统

我们已测试以下环境。

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

注意：已知 CUTLASS 3.x 版本在所有 CUDA 工具包下的 Windows 平台上构建失败。
CUTLASS 团队正在修复此问题。

硬件

CUTLASS 已在以下 NVIDIA GPU 上成功运行，并且预计在基于 Volta、Turing、Ampere、Ada 和