NVIDIA/cutlass

 # 概述 # CUTLASS 4.4.1 _CUTLASS 4.4.1 - 2026年2月_ CUTLASS 是一个用于在 CUDA 中实现高性能矩阵乘法 (GEMM) 及相关计算的抽象集合，涵盖了各个层级和规模。它包含了分层分解和数据移动的策略。CUTLASS 将这些“动态部分”分解为可重用、模块化的软件组件和抽象。 概念并行化层级中不同层级的原语可以通过自定义平铺尺寸、数据类型和其他算法策略进行专门化和调整。由此产生的灵活性简化了它们作为自定义内核和应用程序中构建块的使用。 自 2017 年以来，CUTLASS 一直为高性能线性代数提供 CUDA C++ 模板抽象，这些抽象广泛支持多种计算，包括混合精度计算、专门的数据移动和乘累加 抽象，支持 FP64、FP32、TF32、FP16、BF16、 [通过 Tensor Core 指令进行 FP32 仿真](https://github.com/NVIDIA/cutlass/tree/main/examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm)、 8b 浮点类型 (e5m2 和 e4m3)、 块缩放数据类型 (NVIDIA NVFP4 和 OCP 标准 MXFP4, MXFP6, MXFP8)、 窄整数类型 (4 位和 8 位有符号及无符号整数)、 以及二进制 1b 数据类型 (在架构允许原生支持此类数据类型的情况下)，覆盖 NVIDIA 的 Volta, Turing, Ampere, Ada, Hopper, 和 Blackwell 架构。 在这个丰富的基于 C++ 的内核编程抽象生态系统之上，CUTLASS 4 增加了 CUTLASS DSL。这些是原生的 Python 接口，用于基于核心 CUTLASS 和 CuTe 概念编写高性能 CUDA 内核，且没有任何性能损失。这使得学习曲线更加平滑，编译速度提高了数个数量级，无需编写粘合代码即可与 DL 框架原生集成，以及更直观的元编程，不再需要深厚的 C++ 专业知识。 总的来说，我们将 CUTLASS DSL 视为一个领域特定语言 家族。随着 4.0 版本的发布，我们发布了其中的第一个：CuTe DSL。这是一个底层编程模型，与 CuTe C++ 抽象完全一致 —— 暴露了布局、张量、硬件原子 以及对硬件线程和数据层级的完全控制等核心概念。 CuTe DSL 展示了针对 NVIDIA Ampere, Hopper, 和 Blackwell 架构实现的可编程、高吞吐量 _Tensor Cores_ 的最佳矩阵乘法和其他线性代数运算。 我们相信它将成为学生、研究人员和性能工程师不可或缺的工具 —— 它降低了 GPU 编程的学习曲线，能够快速原型化内核设计，并将优化后的解决方案投入生产环境。 CuTe DSL 目前处于公开测试阶段，并将于 2025 年夏季结束前正式发布。 为了快速入门 - 请参考： - [CUTLASS C++ 快速入门指南](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/quickstart.html)。 - [CuTe DSL 快速入门指南](https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/quick_start.html)。 # CUTLASS 4.4 更新内容 ## CuTe DSL * 新功能 - CuTe DSL 现已支持 CUDA toolkit 13.1！ + 使用 cutlass/python/CuTeDSL/setup.sh --cu13 进行设置 + 详情请参阅 https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/quick_start.html - 搭配 CTK 13.1，CuTe DSL 现已支持 GB300 + 请参阅 [SM103 batched 3xFP4 blockscaled GEMM kernel](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/blackwell/sm103_dense_blockscaled_gemm_persistent.py) 查看示例内核 - cute.experimental：在现有 CuTe DSL API 之上引入了一个更高层次、可组合的层（不是单独的抽象），可以与现有的 Cute DSL 构建块混合使用。 + 无 Fragment 编程模型：copy/dot API 直接接受 memrefs 而不是 descriptors/fragments。 + 自动 TMA descriptor 生成和更新插入。 + SIMT copy 的自动向量化和断言。 + 带有便捷包装器的新 pipeline 抽象 + 新的 Partition 操作以简化分区逻辑。 + 设备端 TMA descriptor 分配、初始化和管理 + 这些示例可在 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 以通过显式同步将数据存储到分布式共享内存。 - Grouped GEMM 示例现在支持仅限设备端的问题形状。 - 我们允许在主机端没有问题形状的情况下进行 grid carve-out。 - Tma+LdMatrix 功能用于加载+解包窄宽度类型（示例用法请参考 mixed_input_fmha_decode.py）。 - 现在可以通过 Python Epilogue Fusion Configuration (EFC) 函数为持久密集 GEMM 定制 epilogue fusion，这在某种程度上类似于 CUTLASS C++ EVT。它还提供了一个 PyTorch 评估器来比较结果。 * 更多编写峰值性能内核的示例 - [SM103 batched 3xFP4 blockscaled GEMM kernel](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/blackwell/sm103_dense_blockscaled_gemm_persistent.py) - 混合输入 FMHA decode 示例，支持 int4 KV (int8 KV 在 4.3 中支持) - 引入了新的 acc_scale grouped mixed input gemm 内核变体，以为解码情况提供更好的性能。 - 所有 mixed_input_gemm 示例已移至单独的文件夹 `mixed_input_gemm`。通用实用函数也被提取到同一文件夹下的 mixed_input_host_utils.py 中。 * 错误修复和改进 - 修复了 if 的两个分支都被执行的问题 - 修复了 `cute.printf` 的 f-string 问题 - 修复了标量张量的索引问题 - 修复了在 [Blackwell SM100 persistent dense blockscaled GEMM with static scheduling](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/blackwell/dense_blockscaled_gemm_persistent.py) 中，带有重叠累加器优化的 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 copy traits 已更新以忠实于无排列的 PTX。排列变体重命名为 LdMatrix16x8x8bOp。 - Grouped GEMM 示例接受参数 --host_problem_shape_available。如果提供了该参数，grid 将基于主机端的问题形状进行 carve out，否则，我们将启动最大可能的 SMs。 - hardware_info.get_max_active_cluster 支持传入特定的 stream 进行查询。适用于基于 green context 的 SM 分区。 - async bulk copy 示例中的 group_bulk_copy_modes 现已弃用，请直接使用 group_modes。 - 弃用 nvvm wrapper 使用 nvvm enum，改用 str。 - cute.arch.calc_packed_f32x2_op 默认启用 ftz 改为默认禁用 ftz - 在带有 CTK 13.1 的 CuTe DSL 中，cutlass.cute.arch 中的以下 API 现在需要字符串字面量而不是 enum 作为参数： + 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 toolkit 为特定内核调整的内部编译器设置，以获得更好的 GPU 内核代码。有关如何创建这些控制文件的更多详细信息和文档将在未来的 CUDA toolkit 版本中提供。注意：高级编译器控制文件不适用于未为其调整的内核。没有兼容性保证，控制文件不适用于不同版本的 CUDA toolkit。 ## CUTLASS C++ * 为 Blackwell 低延迟生成阶段 GQA 内核添加了 [示例 93](https://github.com/NVIDIA/cutlass/tree/main/examples/93_blackwell_low_latency_gqa/)。 - 带 cluster reduction 的 Flash Decoding。 - 内核设计详情请查看 [Readme](https://github.com/NVIDIA/cutlass/tree/main/examples/93_blackwell_low_latency_gqa/readme.md)。 * 在 [示例 112](https://github.com/NVIDIA/cutlass/tree/main/examples/112_blackwell_ssd) 中添加了 Blackwell SM100 State Space Decomposition (SSD) 内核。 * 在 [示例 111](https://github.com/NVIDIA/cutlass/tree/main/examples/111_hopper_ssd) 中添加了 Hopper SM90 State Space Decomposition (SSD) 内核。 * 添加了 Hopper e2m1 到 fp32 的优化转换以及 e2m1 * TF32 tensor core GEMM。 - 启用带有 TF32 支持的 [示例 55](https://github.com/NVIDIA/cutlass/tree/main/examples/55_hopper_mixed_dtype_gemm) * 添加了 [示例 94](https://github.com/NVIDIA/cutlass/tree/main/examples/94_ada_fp8_blockwise/)，用于 Ada FP8xFP8 -> BF16 GEMM，在 MMA 循环中对输入矩阵进行块级反量化，并进行 FP32 累加。 * 添加了对块缩放张量 的任意应用程序提供步幅的支持。 - 用户和应用程序现在必须在所有情况下传递有效的块缩放步幅，即使张量是 packed 的。 * 支持 CUDA 13.1 的 4x blockscale 公共 ptx。 * 允许 `AuxTmaParams` 中使用非静态 `TmaGbasis`。 - attention 内核中的某些情况可能需要非静态 `tma_gbasis`。 - 放宽了对 `AuxTmaParams` 的 `TmaGbasis` 参数的限制，允许用户手动构造动态 gbasis。 * 修复了一些内核问题： - 修复了 MSVC 预处理问题。 - 修复了 GEMV 内核中的自赋值问题。 - 修复了 TMA descriptor 错误，即 CUDA 驱动程序未正确设置 OOB 地址生成模式。 - 修复了 Blackwell SM120 pingpong 内核中 clc scheduler 的内存栅栏 问题。 - 修复了 Blackwell SM120 缩放因子中缺失的 SMEM 对齐。 - 修复了 grouped gemm 的 PDL 问题。 - 修复了 sm100 implicit gemm 内核中 canimplement 的除以零问题。 - 修复了 Grouped GEMM 的 cluster swizzle。 + 将主机端 swizzle 启发式算法移至设备端。 + 根据问题形状和最大 swizzle 大小按组应用 swizzle。 + 改进了示例和单元测试。 * 修复了一些 profiler 问题： - 修复了 nvfp4 grouped GEMM 内核的核心转储问题。 - 修复了不一致的 GEMM 验证逻辑。 - 重新设计了针对不同类型的 grouped gemm 验证逻辑。 - 修复了使用 nvMatmulHeuristics 时的 API 破坏性变更。 * 修复了 `media/docs` 下的一些失效链接。 注意：已知 CUTLASS 4.x 版本在所有 CUDA toolkits 的 Windows 平台上无法构建。 CUTLASS 团队正在致力于修复。 **有关所有过去版本和更新的详情，请参阅 [更新日志](https://docs.nvidia.com/cutlass/latest/CHANGELOG.html)。** # 性能 CUTLASS 原语非常高效。当用于构建设备级 GEMM 内核时，它们展现出接近峰值理论吞吐量的最佳利用率。下图展示了 CUTLASS 3.8 在 NVIDIA Blackwell SM100 架构 GPU 上运行时，在各种输入和输出数据类型下的性能（以理论峰值利用率的百分比表示）。  下图展示了自 CUTLASS 3.1 以来，在 [NVIDIA H100](https://www.nvidia.com/en-us/data-center/h100/) (NVIDIA Hopper 架构) 上 CUTLASS 性能的持续改进。 CUTLASS 3.5.1 使用 [CUDA 12.5u1 Toolkit](https://developer.nvidia.com/cuda-downloads) 编译。 Tensor Core 操作使用 CUDA 的 [mma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma) 和 [wgmma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions) 指令实现。   # CuTe CUTLASS 3.0 引入了一个新的核心库 CuTe，用于描述和操作线程和数据张量。 CuTe 是一组 C++ 模板抽象，用于定义和操作线程和数据的分层多维布局。 CuTe 提供了 `Layout` 和 `Tensor` 对象，它们紧凑地封装了数据的类型、形状、内存空间和布局，同时为用户执行复杂的索引工作。 这让程序员可以专注于算法的逻辑描述，而 CuTe 为他们处理机械的簿记工作。借助这些工具，我们可以快速设计、实现和修改所有密集线性代数运算。 CuTe 的核心抽象是分层多维布局，它可以与数据数组组合来表示张量。 布局的表示形式非常强大，足以表示我们实现高效密集线性代数所需的一切。 布局还可以通过函数组合进行组合和操作，在此基础上我们构建了大量常见操作，例如 tiling 和 partitioning。 CUTLASS 3.0 及更高版本在其模板的整个 GEMM 层级中采用了 CuTe。 这大大简化了设计，并提高了代码的可组合性和可读性。 关于 CuTe 的更多文档可以在其[专用文档目录](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/00_quickstart.html)中找到。 # 兼容性 最低要求： - 架构：Volta (计算能力 7.0) - 编译器：必须至少支持 C++17 - CUDA Toolkit 版本：11.4 CUTLASS 需要 C++17 主机编译器，并且在使用 [**CUDA 12.8 Toolkit**](https://developer.nvidia.com/cuda-downloads) 构建时性能最佳。 它还兼容 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 版本在所有 CUDA toolkits 的 Windows 平台上无法构建。 CUTLASS 团队正在致力于修复。 ## 硬件 CUTLASS 在以下 NVIDIA GPU 上成功运行，并且预期在基于 Volta, Turing, Ampere, Ada, 和 Hopper 架构的 NVIDIA GPU 上高效运行。 |**GPU**|**CUDA 计算能力**|**CUTLASS-3 要求的最低 CUDA Toolkit**| |---|---|---| |NVIDIA V100 Tensor Core GPU |7.0|11.4| |NVIDIA TitanV |7.0|11.4| |NVIDIA GeForce RTX 20x0 series |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 series |8.6|11.4| |NVIDIA GeForce RTX 40x0 series |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 series |12.0|12.8| |NVIDIA DGX Spark |12.1|13.0| ## 目标架构 通常，为一个目标架构生成的 PTX 代码可以在未来的架构上运行（即，它是向前兼容的）。 但是，CUDA 12.0 引入了“架构加速特性”的概念，其 PTX 不具备向前兼容性保证。 几个 Hopper 和 Blackwell PTX 指令属于这类架构加速特性，因此需要 `sm_90a` 或 `sm100a` 目标架构（注意附加的“a”）。有关此点和其他架构加速指令的更多详细信息，请参阅 [CUDA 文档](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#feature-availability)。 目标架构信息通过 cmake 标志 `CUTLASS_NVCC_ARCHS` 传递给 CUTLASS。为了在 Hopper GH100 上最大化性能，用户需要使用 `90a` 作为目标架构来构建 CUTLASS。 如果用户意外地构建了一个使用 SM90a 特性（例如 Hopper Tensor Core 指令）的内核，却使用了 SM90 目标（注意缺少“a”），无论是使用 CUDA Toolkit 12 还是 11.8，该内核预计会因运行时错误而失败。 ``` cmake .. -DCUTLASS_NVCC_ARCHS="90a" ``` 或者 ``` cmake .. -DCUTLASS_NVCC_ARCHS="100a" ``` 注意：数据中心产品中使用的 NVIDIA Blackwell SM100 架构与支撑 NVIDIA Blackwell GeForce RTX 50 系列 GPU (SM120) 的架构具有不同的计算能力。因此，为带有架构条件特性（使用 `sm100a`）的 Blackwell SM100 架构编译的内核与 RTX 50 系列 GPU 不兼容。 请参阅[功能文档](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/functionality.html)以了解哪些内核需要哪些目标架构的详细信息。 # 文档 CUTLASS 在以下文档和随附的 [Doxygen 文档](https://nvidia.github.io/cutlass)中进行了描述。 - [快速入门指南](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/quickstart.html) - 构建和运行 CUTLASS 的基础知识 - [功能](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/functionality.html) - 总结了 CUTLASS 中可用的功能 - [CUDA 中的高效 GEMM](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/efficient_gemm.html) - 描述了如何在 CUDA 中高效实现 GEMM 内核 - [CUTLASS 3.x 设计](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cutlass_3x_design.html) - 描述了 CUTLASS 3.x 设计、其优势以及 CuTe 如何使我们能够编写更具可组合性的组件 - [GEMM API 3.x](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/gemm_api_3x.html) - 描述了 CUTLASS 3.x GEMM 模型和 C++ 模板概念 - [GEMM API 2.x](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/gemm_api.html) - 描述了 CUTLASS 2.x GEMM 模型和 C++ 模板概念 - [Implicit GEMM 卷积](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/implicit_gemm_convolution.html) - 描述了 CUTLASS 中的 2-D 和 3-D 卷积 - [代码组织](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/code_organization.html) - 描述了 CUTLASS 项目的组织和内容 - [术语](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/terminology.html) - 描述了代码中使用的术语 - [编程指南](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/programming_guidelines.html) - 编写高效现代 CUDA C++ 的指南 - [基础类型](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/fundamental_types.html) - 描述了 CUTLASS 中用于表示数值量和数组的基本 C++ 类 - [布局](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/layout.html) - 描述了矩阵和张量在内存中的布局 - [Tile 迭代器](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/tile_iterator_concept.html) - 描述了用于迭代内存中矩阵 tile 的 C++ 概念 - [CUTLASS Profiler](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/profiler.html) - 命令行驱动的性能分析应用程序 - [CUTLASS 实用工具](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/utilities.html) - 用于促进快速开发的附加模板 - [依赖内核启动](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/dependent_kernel_launch.html) - 描述了 Hopper 中的一个新特性，该特性允许在同一 stream 中重叠依赖 内核，以及它在 CUTLASS 中是如何使用的。 # 资源 我们还在 [GPU Technology Conference 2018](http://on-demand.gputechconf.com/gtc/2018/presentation/s8854-cutlass-software-primitives-for-dense-linear-algebra-at-all-levels-and-scales-within-cuda.pdf) 的演讲中描述了高效 GEMM 的结构。 - [CUTLASS: CUDA 各层级和规模密集线性代数的软件原语](https://www.nvidia.com/en-us/on-demand/session/gtcsiliconvalley2018-s8854/) - [开发 CUDA 内核以在 NVIDIA A100 上将 Tensor Cores 推向极限](https://www.nvidia.com/en-us/on-demand/session/gtcsj20-s21745/) - [在 CUTLASS 中利用 Tensor Cores 加速卷积](https://www.nvidia.com/en-us/on-demand/session/gtcspring21-s31883/) - [在 CUTLASS 中通过提高 Tensor Core 利用率加速反向数据梯度](https://www.nvidia.com/en-us/on-demand/session/gtcspring22-s41996/) - [CUTLASS: Python API, 增强功能与 NVIDIA Hopper](https://www.nvidia.com/en-us/on-demand/session/gtcfall22-a41131/) # 构建 CUTLASS CUTLASS 是一个仅头文件 的模板库，不需要构建即可被其他项目使用。客户端应用程序应在其包含路径中定位 CUTLASS 的 `include/` 目录。 CUTLASS 单元测试、示例和实用工具可以使用 CMake 构建。 CMake 的最低版本在 [快速入门指南](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/quickstart.html) 中给出。 确保 `CUDACXX` 环境变量指向您系统上安装的 CUDA Toolkit 中的 NVCC。 ``` $ export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc ``` 在 CUTLASS 项目中创建一个 build 目录，然后运行 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 # compiles for NVIDIA's Ampere Architecture ``` 从 `build/` 目录中，通过使用 make 构建目标 `test_unit` 来编译并运行 CUTLASS 单元测试。 单元测试组织为多个二进制文件，镜像了 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 文档](https://nvidia.github.io/cutlass) 提供了 CUTLASS 项目中定义的文件、类和模板概念的完整列表。 关于源代码组织的详细说明可以在 [CUTLASS 文档](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/code_organization.html) 中找到，但下面总结了几个主要组件。 ## 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 SDK 示例](https://github.com/NVIDIA/cutlass/tree/main/examples) 应用 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/ # managing tensors in device memory, reference cutlass/ # implementations for GEMM, random initialization util/ # of tensors, and I/O. ``` ### 测试 `test/unit/` 目录包含使用 Google Test 实现的单元测试，演示了 Core API 组件的基本用法，并对 CUTLASS GEMM 计算进行了完整测试。 构建和运行单元测试的说明在 [快速入门指南](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/quickstart.html) 中有描述。 # 性能分析 `tools/profiler/` 目录包含一个命令行实用工具，用于启动每个 GEMM 内核。 它可以按如下方式构建： ``` $ make cutlass_profiler -j16 ``` ## 构建所有 GEMM 和卷积内核（_较长_ 的构建时间） 默认情况下，每种数据类型、数学指令和布局只实例化一种 tile 大小。 要实例化所有，请在从空 `build/` 目录运行 CMake 时设置以下环境变量。 注意，这会导致 *数万个* 内核和漫长的构建时间。 这也会导致二进制文件体积过大，并且在某些平台上链接器在构建库时会失败。 因此，强烈建议只生成一部分内核，如下面子节所示。 ``` $ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ make cutlass_profiler -j16 ``` ## 构建一部分 GEMM 和卷积内核（_减少_ 构建时间） 要严格编译一个内核或一小组内核，可以使用带有通配符的内核名称的逗号分隔列表来减少内核集。以下示例展示了为 NVIDIA Ampere 和 Turing 架构构建恰好一个或一部分内核： ### 构建一部分 Tensor Core GEMM 内核 要编译具有 FP32 累加和 FP16 输入、针对 NVIDIA Ampere 和 Turing 架构的一部分 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 ... ============================= Problem ID: 1 Provider: CUTLASS OperationKind: gemm Operation: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8 Status: Success Verification: ON Disposition: Passed reference_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 ============================= ... ``` ### 构建一个 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 ============================= Problem ID: 1 Provider: CUTLASS OperationKind: 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 ============================= ``` ### 构建一部分 Tensor Core 卷积内核 要编译具有 FP32 累加和 FP16 输入、针对 NVIDIA Ampere 和 Turing 架构的实现前向传播 的一部分 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 ... ============================= Problem ID: 1 Provider: CUTLASS OperationKind: conv2d Operation: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc Status: Success Verification: ON Disposition: Passed reference_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 ============================= ... ``` ### 构建一个卷积 CUDA 内核 要编译并运行一个具有 F32 累加和 FP32 输入、针对 NVIDIA Ampere 和 Turing 架构的实现前向传播 的 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 卷积内核的示例命令行：

标签：Ampere, Bash脚本, BF16, Blackwell, CUDA, CUDA内核, CUTLASS, C++模板, FP16, FP64, GEMM, Hopper, HPC, JIT编译, Python DSL, SIMD, Tensor Core, TF32, 代码生成, 凭据扫描, 半精度计算, 并行计算, 库开发, 底层优化, 张量核心, 微架构, 数学库, 整数运算, 深度学习, 混合精度, 渗透测试工具, 矩阵乘法, 科学计算, 算子开发, 线性代数, 计算机视觉, 逆向工具, 高性能计算