CUTLASS 的定位
CUTLASS 是 NVIDIA 提供的开源 GPU 计算项目，主要提供实现高性能矩阵乘法及相关计算的抽象与组件。它将线程与数据的划分、内存访问、矩阵乘加和输出处理组织为可组合的实现，使开发者能够按数据类型、矩阵形状和目标 GPU 配置或扩展 Kernel。项目的主要内容包括 CUDA C++ 模板库、CuTe 核心抽象，以及用 Python 语法编写 CUDA Kernel 的 CUTLASS Python。官方项目总览同时介绍了这些编写方式和相关示例。
CUDA 提供 NVIDIA GPU 的并行计算平台与编程模型。CUTLASS 基于这一平台表达并执行计算：外层代码准备参数并启动 Kernel，设备上的线程完成数据访问与运算。因此，理解 CUTLASS 的前提是掌握线程块、共享内存和 Kernel 启动等基本概念；相关说明见本站的 CUDA 编程基础。
CUTLASS 的主要使用方式是复用计算组件构建自定义 Kernel。开发者可以从完整的 GEMM 实现开始调整配置，也可以使用其中的布局、搬运和计算组件编写自己的算法。这种组织方式适合需要改变数据布局、计算调度或融合操作的场景。
GEMM 分块与数据流
GEMM（General Matrix Multiplication，通用矩阵乘法）的常见计算形式为 D = αAB + βC。其中 A 的形状为 M × K，B 为 K × N，C 和 D 为 M × N，α、β 为标量。输出元素 Dij 需要对 K 维的乘积求和，再与 Cij 合并。M、N 决定输出规模，K 决定每个输出元素的归约长度。
直接为每个输出元素独立读取整行 A 和整列 B，会重复访问许多输入元素。分块（tiling）将计算划分为较小的矩阵区域，使参与同一输出区域计算的线程共享输入数据，并让数据在共享内存和寄存器中被多次使用。分块同时决定并行任务的数量、数据复用程度和片上存储需求。
一个输出块的计算沿 K 维分阶段进行：每次读取 A 的一部分列和 B 的一部分行，计算它们的乘积并累加到该输出块。以 M = N = 256、K = 128 为例，若输出块大小为 64 × 64，则共有 16 个输出块；若每次处理 32 个 K 维元素，每个输出块需要完成 4 次部分乘积累加。这些数字仅用于说明分块关系。
Mainloop（主循环）负责沿 K 维反复搬运输入块并执行乘加。MMA（Matrix Multiply-Accumulate，矩阵乘加）表示矩阵乘积与已有结果的累加；CUTLASS 可以利用 Tensor Core 的矩阵指令，也提供由常规线程算术实现的计算路径。数据搬运和 MMA 必须使用相匹配的线程划分与数据布局。
Epilogue（收尾阶段）负责将累积结果转换为最终输出，完成 α、β 缩放以及输出存储。合适的 epilogue 还可以融合偏置、激活函数或类型转换。例如，将 Y = ReLU(αAB + βC + bias) 的后处理与 GEMM 写回结合，可以减少单独后处理所需的中间结果访问；其中 ReLU 按元素取 max(0, x)。具体可融合的运算由所选实现和接口决定。
流水线（pipeline）协调不同阶段的搬运与计算。在支持相应异步操作的实现中，当前数据块用于计算时，后续数据块可以同时进入可用缓冲区；同步机制保证计算只使用已到达的数据，并在读完后释放缓冲区。增大 tile 或增加缓冲阶段会占用更多共享内存、寄存器等资源，配置时需要兼顾数据复用、并行任务数量和资源容量。
CuTe 与 CUTLASS 的主要组件
CuTe 是 CUTLASS 中用于描述和操作多维线程与数据布局的核心库。它用统一的布局表达支持数据分块、线程分配、搬运与矩阵运算，CUTLASS 3.x 的 C++ GEMM 组件广泛使用这些抽象。该关系见 CUTLASS 3.x Design。
CuTe 的 Layout（布局）描述从逻辑坐标到索引的映射。Shape 给出各维大小，Stride 给出坐标变化对应的索引步长。例如，一个紧密存放的 2 × 3 行主序数组可以使用 Shape = (2, 3)、Stride = (3, 1)，坐标 (i, j) 的索引为 3i + j。同一套布局运算还可用于描述线程及其负责的数据，从而明确分块后的元素分配。定义及布局组合规则见 CuTe Layouts。
Tensor（张量）将布局与数据访问对象结合起来。访问某个逻辑坐标时，布局计算对应索引，再由指针或迭代器访问数据。这样，算法可以按多维坐标表达访问，而数据的物理组织由布局描述。相关接口见 CuTe Tensors。
CUTLASS 3.x 的 C++ GEMM API 将完整计算分成以下层次。布局抽象支撑内层的数据划分，mainloop 和 epilogue 则在外层被组合为可启动的 Kernel。
| 层次 | 主要职责 | 典型接口 |
|---|---|---|
| Device | 提供 Host 侧的参数检查、初始化与 Kernel 调用接口 | GemmUniversalAdapter |
| Kernel | 组合主循环与收尾阶段，并组织完整计算 | GemmUniversal |
| Collective | 协调协作线程的数据搬运、计算或输出处理 | CollectiveMma、collective epilogue |
| Tiled MMA / Copy | 将基本操作按布局扩展到多个线程和数据元素 | TiledMma、TiledCopy |
| Atom | 封装基本搬运或矩阵乘加操作及其线程、数据要求 | Copy_Atom、Mma_Atom |
按照 GEMM API 文档，构建完整 C++ GEMM 的典型关系是：配置 collective mainloop 和 epilogue，将二者组合成 Kernel 类型，再用 Device 层适配器包装调用。Device 是这套 API 的层次名称，其调用接口位于 Host 侧。
CUTLASS Python提供用 Python 语法编写 CUDA Kernel 的开发方式，其中 CuTe DSL（Domain-Specific Language，领域专用语言）沿用布局、张量、Copy Atom 和 MMA Atom 等概念。作者仍需表达线程与数据的组织方式，Kernel 代码通过 CUTLASS Python 编译栈和 CUDA 工具链生成设备代码。它与 C++ 路径共享核心概念，具体编写与编译接口见 CUTLASS Python Overview。
CUTLASS 与 CUDA、Triton 的关系
CUDA 是平台与编程模型，CUDA C++ 是直接编写 Kernel 的常见方式。CUTLASS 提供面向 CUDA 计算的实现组件和编写接口。Triton则提供基于 Python 的并行编程语言与编译器，以张量块表达一个 program 实例处理的数据和运算。下面按同一 GEMM 的实现任务比较这些方式，Triton 部分以标准 triton.language 接口为范围。
| 编写方式 | 主要表达对象 | 开发者配置或实现的内容 |
|---|---|---|
| CUDA C++ | 线程、线程块及其内存访问与运算 | 显式编写索引、共享数据、同步和计算逻辑 |
| CUTLASS C++ | 可组合的 GEMM 组件与策略 | 选择类型、布局、tile、主循环和 epilogue，并按需扩展组件 |
| CuTe DSL | Python 语法中的布局、张量、搬运和计算操作 | 显式描述线程与数据划分、存储使用和流水线 |
| Triton | program 负责的张量块及其运算 | 选择块大小、访问方式和计算逻辑，调优编译配置 |
对于一个输出矩阵块，CUDA C++ 作者可以直接编写协作线程的加载、同步和乘加逻辑；CUTLASS C++ 作者可以选择已实现的组件并组合 mainloop 与 epilogue；CuTe DSL 作者用 Python 语法明确布局和数据移动；Triton 作者通常加载 A、B 的张量块，以块级运算累积结果，由编译器安排更多底层布局与资源使用。Triton 矩阵乘法教程给出了这一块级计算过程，Triton 官方编程模型说明解释了标准 Triton 中由编译器处理的布局、数据移动等细节。
这些方式的差异主要体现在复用边界与控制粒度。使用完整 CUTLASS GEMM 时，开发者主要选择组件和策略；直接使用 CuTe 编写 Kernel 时，则需要参与更多布局与数据流设计。因此，评估学习成本和开发工作量时，应明确采用的接口层次以及需要定制的部分。
CUTLASS 的适用场景
CUTLASS 适合以矩阵计算为核心、需要定制实现的场景，例如对特定矩阵形状选择分块方案，对输入、输出和累积数据类型分别配置，或将后处理融合进 GEMM。其组件也可用于相关的自定义算法；实际支持范围取决于具体接口、目标架构和计算形式。
对于标准矩阵乘法，已有的 cuBLAS / cuBLASLt 接口通常可以直接满足调用需求。cuBLASLt 提供矩阵布局、计算类型和算法选择等配置，部分后处理也可通过其 epilogue 功能完成。需要开发自定义 Kernel 时，再根据现有接口能覆盖的操作和所需控制粒度选择 CUTLASS、CuTe DSL 或 Triton。cuBLASLt 的配置范围见 官方 API 文档。
比较实现性能时，需要固定目标 GPU、矩阵形状、数据类型、布局和融合操作，验证结果误差后再计时；使用即时编译的实现时，还应区分编译开销与后续执行时间。CUTLASS 提供 CUTLASS Profiler 用于运行、验证和测量候选算子。本文的讨论范围为编程关系与计算机制。
结论与参考资料
CUTLASS 将矩阵计算中的分块、布局、搬运、乘加和输出处理表达为可组合的实现。GEMM 数据流说明了这些组件的职责，CuTe 则统一描述其中的坐标与数据划分。理解这两层关系后，可以根据定制需求选择完整 GEMM 接口或更细的 Kernel 编写方式。
- CUTLASS 项目总览：项目组成、示例与文档入口。
- Efficient GEMM in CUDA：分块、数据复用和流水线机制。
- CUTLASS 3.x GEMM API：组件层次与 Kernel 组合关系。
- CuTe Layouts 与 CUTLASS Python Overview：布局抽象与 Python 编写接口。
资料核对日期：2026 年 10 月 10 日。两张图均依据官方文档绘制，用于说明计算阶段与编写方式。