CUTLASS 4.7 深度指南:CUDA 高性能线性代数模板库与 CuTe DSL 全景解析
【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass
CUTLASS 是 NVIDIA 开源的 CUDA C++ 模板库,为在 CUDA 的所有层级与规模下实现高性能矩阵乘法(GEMM)及相关计算提供了分层分解与数据搬移的模块化抽象;CUTLASS 4 又在此基础上加入了 Python 原生的CUTLASS DSL(CuTe DSL),在不损失性能的前提下大幅降低 GPU 内核开发门槛。本文以当前仓库 README.md 为骨架,结合 CMakeLists.txt、include/cutlass、include/cute、tools/profiler 等源码,系统讲解 CUTLASS 的能力边界、4.7 新特性、兼容性矩阵、构建流程、项目结构与 Profiler 实战用法。
完整的 CUDA GEMM 层次分解示意图
一、CUTLASS 是什么:为 CUDA 中所有层级与规模的 GEMM 而生的模板库
CUTLASS 是一组用于在 CUDA 内实现高性能矩阵乘法(GEMM)及相关计算的抽象集合,核心思路是层次化分解(hierarchical decomposition)与数据搬移(data movement)。它把 GEMM 中复杂的"活动部件"拆解为可复用、模块化的软件组件与抽象,程序员可以针对概念化并行层次的不同层级,通过自定义 tile 尺寸、数据类型与算法策略对原语进行专门化与调优,从而让这些抽象在自定义内核与应用中充当灵活的建筑积木。
自 2017 年以来,CUTLASS 一直为高性能线性代数提供 CUDA C++ 模板抽象,支持极其广泛的数值类型组合:
- 混合精度计算:FP64、FP32、TF32、FP16、BF16;
- 基于 Tensor Core 指令的 FP32 仿真(3xTF32 快速且精确的实现,参见 examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm);
- 8 位浮点类型:e5m2 与 e4m3;
- 块缩放数据类型:NVIDIA NVFP4 与 OCP 标准的 MXFP4、MXFP6、MXFP8;
- 窄整数类型:4 位与 8 位有符号/无符号整数;
- 二进制 1 位数据类型(在硬件原生支持的架构上);
- 覆盖 NVIDIAVolta、Turing、Ampere、Ada、Hopper、Blackwell全系列架构。
在丰富的 C++ 内核编程抽象生态之上,CUTLASS 4 加入了CUTLASS DSLs——基于核心 CUTLASS 与 CuTe 概念编写高性能 CUDA 内核的 Python 原生接口。其优势包括:更平滑的学习曲线、数量级更快的编译时间、无需胶水代码即可与深度学习框架原生集成、以及不需要深厚 C++ 功底即可完成的直观元编程。CUTLASS 4.0 发布了其中的第一个 DSL——CuTe DSL,它是一种与 CuTe C++ 抽象完全一致的低层编程模型,暴露了布局(layouts)、张量(tensors)、硬件原子(hardware atoms)等核心概念,并让开发者完全掌控硬件线程与数据层次。CuTe DSL 针对 NVIDIA Ampere、Hopper 与 Blackwell 架构可编程、高吞吐的 Tensor Core 演示了最优的矩阵乘法与其他线性代数算子。按 README 的说明,CuTe DSL 目前处于公开测试阶段(public beta),预计 2026 年夏季末正式毕业(graduating out of beta)。
二、CUTLASS 4.7 新特性速览
当前仓库对应CUTLASS 4.7.0(2026 年 8 月发布),4.7 的新特性分布在三个层面:
2.1 CuTe DSL:Primitives API、任务调度框架与更强的编译诊断
- Primitives API:在 CuTe 之下提供了更底层的抽象,允许通过 SIMT 进行 Tensor Core 编程,是 NVVM 操作之上的稳定、轻薄封装,适合在 CuTe 抽象拖慢开发速度的场景使用。Primitives 目前以 experimental 状态发布,会依据用户反馈演进;README 特别注明这是一套过渡性 API,直到可用的 CUDA Python 式解决方案出现为止。
- Task Scheduling 框架:为 warp-specialized 内核提供执行调度的静态分析。当检测到已知并发问题时,编译会直接停止(compilation stops when known concurrency issues are detected),同时还提供了可视化资源/任务依赖、分析内核调度结构的工具。
- 改进的编译诊断:寄存器溢出(register spills)与本地内存(local memory)使用可在编译期以源码行号形式报告;当使用 Primitives API 时,初步支持在编译期检测若干类 NVVM 同步与执行危害(synchronization and execution hazards);此前不含源码行号的编译器报错现在也有了更好的行号报告。
2.2 CUTLASS Operator API:自定义 epilogue 融合增强
- 新增对**标量归约(scalar reductions)**的支持;
- 新增能力:可为每个被加载/存储的 operand 单独指定数据搬移策略(data movement strategy)。
2.3 C++ 层改进与修复
- 在 examples/77_blackwell_fmha 中新增了面向FP8 的 2-kernel backward实现:新增一个带多精度(FP16/FP8)的反向融合多头注意力 benchmark,支持可配置的 batch/sequence/head 尺寸、变长与掩码选项,并内置正确性检查与运行时间/吞吐量报告。按 README 说明,在 Blackwell SM103 芯片上无掩码 FP8 场景下,2-kernel backward 相比 1-kernel 实现约有25% 的提升(这是项目自测数据,实际收益取决于具体硬件与问题规模)。
- NVRTC:支持 CUDA 12.6 及更新版本的结构化绑定头文件(structured bindings headers)。
- NumericArrayConverter:新增 fp32/fp16/bf16/e4m3/e5m2 → e2m1(FP4)的转换路径。
- 优化E2M1 → FP16 的 LUT 解码辅助函数
_e2m1_to_half_x2与_e2m1_to_half_x4(合并 mask 后再执行 prmt)。 - 问题修复:更新 streamk 启发式算法以优化混合 cluster 尺寸的内核;修复 CuTe tuple 算法中 integer-sequence get 的歧义;修复一个 TMA 创建驱动 bug(通过检查张量是否紧凑,保守判断首 128KiB 是否被映射,从而决定翻转或清零对应 bit)。
- 支持CUDA Toolkit 13.3下的最优代码生成(optimal code generation)。
注意:README 明确提示,CUTLASS 4.x 在 Windows 平台上的构建对所有 CUDA Toolkit 均存在问题,CUTLASS 团队正在修复中。完整的历史发布记录见 CHANGELOG.md。
三、性能表现
CUTLASS 原语非常高效:当用于构造 device-wide GEMM 内核时,它们能够接近理论峰值吞吐(peak theoretical throughput)的利用水平。README 给出的下图展示了 CUTLASS 3.8 在 NVIDIA Blackwell SM100 GPU 上、针对多种输入输出数据类型以理论峰值利用率百分比呈现的性能:
CUTLASS 3.8 在 Blackwell SM100 上的 GEMM 峰值性能(占理论峰值百分比)
另外两张图展示了自 CUTLASS 3.1 以来在 NVIDIA H100(Hopper 架构)上的持续性能改进。CUTLASS 3.5.1 使用 CUDA 12.5u1 Toolkit 编译,其 Tensor Core 操作基于 CUDA 的mma与wgmma指令实现:
CUTLASS 3.5.1 在 H100 上的 GEMM 峰值性能 CUTLASS 3.5.1 在 H100 上的 FP8 GEMM 峰值性能
四、CuTe 核心库:Layout 与 Tensor 的模板抽象
CUTLASS 3.0 引入了新的核心库CuTe,用于描述和操作线程与数据的张量。CuTe 是一组 C++ CUDA 模板抽象,用于定义和操作线程与数据的分层多维布局(hierarchically multidimensional layouts)。它提供Layout与Tensor对象,将数据的类型、形状、内存空间与布局紧凑地打包在一起,同时替用户完成复杂的索引计算,让程序员专注于算法的逻辑描述,而由 CuTe 负责机械化的簿记工作。
CuTe 的核心抽象——分层多维布局——可以与数据数组组合来表示张量,其布局表示能力足以覆盖高效稠密线性代数所需的几乎所有场景;布局之间还可以通过**函数复合(functional composition)**进行组合与操作,并在此基础上构建起 tiling、partitioning 等大量常见操作。
在仓库中,CuTe 的实现位于 include/cute,其内部组织为:
- layout.hpp、tensor.hpp:核心类型
Shape、Stride、Layout、Tensor及关联操作; - atom/mma_atom.hpp:
cute::Mma_Atom与cute::TiledMma; - atom/copy_atom.hpp:
cute::Copy_Atom与cute::TiledCopy; - atom 下的
*sm*.hpp:各架构(arch)特有的 copy 与 math 操作元信息; - arch:copy 与 math 指令的极简 PTX 封装结构体;
- algorithm:copy、gemm 及
cute::tuple操作等核心操作定义。
CUTLASS 3.0 及后续版本在其模板中全面采用 CuTe 贯穿 GEMM 层次,大大简化了设计并提升了代码的可组合性与可读性。
五、兼容性矩阵:架构、编译器与操作系统
5.1 最低要求
- 架构:Volta(compute capability 7.0)
- 编译器:至少支持 C++17
- CUDA Toolkit 版本:11.4
CUTLASS 要求 C++17 宿主编译器,并且使用CUDA 12.8 Toolkit构建时表现最佳;同时兼容 CUDA 11.4、11.5、11.6、11.7、11.8 及所有 12.x 版本。在仓库根目录 CMakeLists.txt 中可看到对应的构建约束:cmake_minimum_required(VERSION 3.19 FATAL_ERROR),并在检测到 CUDA 版本低于 11.4 时打印告警(CUTLASS ${CUTLASS_VERSION} requires CUDA 11.4 or higher)。
5.2 测试过的操作系统
| 操作系统 | 编译器 |
|---|---|
| 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 在 fold expressions 与重载运算符方面存在已知回归,建议使用 GCC 7.5.0 或(更推荐)GCC ≥ 9。
5.3 硬件支持
CUTLASS 在下列 NVIDIA GPU 上运行成功,并预期在 Volta、Turing、Ampere、Ada、Hopper 架构 GPU 上保持高效:
| GPU | CUDA Compute Capability | 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 |
六、目标架构与 architecture-accelerated features
一般来说,为某目标架构生成的 PTX 代码可以前向兼容地运行在未来的架构上。但 CUDA 12.0 引入了"架构加速特性(architecture-accelerated features)",这类特性的 PTX不保证前向兼容。若干 Hopper 与 Blackwell 的 PTX 指令即属于此类,因此需要使用sm_90a或sm100a目标架构(注意结尾的 "a")。
目标架构通过 CMake 变量CUTLASS_NVCC_ARCHS传递给 CUTLASS。为了在 Hopper GH100 上获得最大性能,必须以90a为目标架构构建;如果用户误用 SM90(缺少 "a")目标构建使用了 SM90a 特性(如 Hopper Tensor Core 指令)的内核,无论是 CUDA Toolkit 12 还是 11.8,该内核预计都会以运行时错误失败:
cmake .. -DCUTLASS_NVCC_ARCHS="90a"或
cmake .. -DCUTLASS_NVCC_ARCHS="100a"注意:数据中心产品使用的 NVIDIA Blackwell SM100 架构,其 compute capability 与 NVIDIA Blackwell GeForce RTX 50 系列 GPU 底层的 SM120 架构不同。因此,使用sm100a架构条件特性编译的 Blackwell SM100 内核无法兼容 RTX 50 系列 GPU。
从仓库 CMakeLists.txt 的实现看,CUTLASS_NVCC_ARCHS的默认值会根据 CUDA Toolkit 版本动态生成受支持的架构列表(例如 75、80、86、87 / 70、72 / 89、90 / 90a / 100、100a、120、120a、121、121a,以及 101a、110a、100f、120f、121f、103a、103f 等),并可通过CUTLASS_NVCC_ARCHS_ENABLED控制实际编译哪些架构的代码;传入不支持的架构会打印Using unsupported or deprecated compute capabilities ...告警。
七、构建 CUTLASS:从 header-only 到单元测试
CUTLASS 是一个header-only 模板库,被其他项目使用时无需构建——客户端应用只需将仓库的include/目录加入 include 路径即可(参见 include/cutlass 与 include/cute)。
CUTLASS 的单元测试、示例与工具可以通过 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 构建内核;为缩短编译时间,可通过修改CUTLASS_NVCC_ARCHS指定目标架构:
$ mkdir build && cd build $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # compiles for NVIDIA's Ampere Architecture在build/目录中,通过构建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.在支持的平台上所有测试都应通过(具体测试数量会随时间变化)。对应的单元测试源码位于仓库的 test 目录,基于 Google Test 实现,覆盖 Core API 组件的基础用法与 CUTLASS GEMM 计算的完整测试。
八、项目结构:模板库、工具与示例
CUTLASS 以 header-only 库为核心,配以 Utilities、Tools、Examples 与单元测试。以下是 README.md 归纳的主要组件:
8.1 CUTLASS 模板库(include/)
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这些目录在当前仓库中均真实存在,例如 include/cutlass/gemm、include/cutlass/epilogue、include/cute/atom 等,读者可对照查阅具体头文件。
8.2 CUTLASS SDK Examples 与 Tools
- CUTLASS SDK examples:位于 examples,将 CUTLASS 模板应用到具体的基础计算中,如 FP32 仿真 GEMM(examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm)与 Blackwell FMHA(examples/77_blackwell_fmha)。
- tools/目录:
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.当前仓库中 tools/profiler(含 src/ 与 include/)、tools/library 等均存在;此外 CUTLASS 4 还新增了 cutlass_compiler(含 CuTe DSL 编译器 cutlass_compiler/cute_ir)与 python(Python 侧 DSL 与绑定)等目录,构成了完整的 C++ / Python 双栈项目形态。
九、性能剖析实战:CUTLASS Profiler
tools/profiler 目录包含一个命令行工具,用于启动各类 GEMM 内核进行性能剖析,构建方式如下:
$ make cutlass_profiler -j169.1 构建全部 GEMM 与卷积内核(编译时间长)
默认情况下,每种数据类型、数学指令与布局只实例化一种 tile 尺寸。若想实例化全部,需要在空的 build/ 目录中运行 CMake 时设置以下环境变量。注意:这会产生数万个内核和极长的构建时间,同时二进制体积巨大,部分平台上链接器可能因此构建库失败。因此强烈建议只生成内核的一个子集(见 9.2 节):
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ make cutlass_profiler -j16从 CMakeLists.txt 的实现看,CUTLASS_LIBRARY_KERNELS是"逗号分隔的内核名过滤器列表":未指定时仅启用最大的 tile 尺寸;指定字符串all时启用全部内核。还可以用CUTLASS_LIBRARY_IGNORE_KERNELS(仅在设置CUTLASS_LIBRARY_KERNELS时生效)与CUTLASS_LIBRARY_EXCLUDE_KERNELS(始终生效)进一步裁剪内核集合。
9.2 构建 GEMM 与卷积内核子集(编译时间短)
编译严格一个或一小撮内核时,可以使用带通配符的逗号分隔内核名列表来缩小集合。下面按 README 给出四类典型场景。
场景一:Tensor Core GEMM 子集
为 Ampere 与 Turing 架构编译 FP32 累加、FP16 输入的 Tensor Core GEMM 内核子集:
$ 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 ============================= ...该输出展示了 CUTLASS Profiler 的关键字段:Operation为具体的内核名(如cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8),其中编码了数学指令(s1688)、tile 尺寸(256x128x32)、stage 数(2)、矩阵转置(nt)与对齐(align8)等信息;Arguments列出全部可调参数,包括cta_m/cta_n/cta_k(CTA tile)、warps_m/warps_n/warps_k(warp 划分)、inst_m/inst_n/inst_k(指令形状)、stages(流水线级数)、min_cc/max_cc(适用 compute capability 范围);Memory(GiB/s)与Math(GFLOP/s)则是剖析的核心性能指标。
场景二:单个 CUDA Core GEMM 内核
为 Ampere 与 Turing 架构编译单个 SGEMM 内核:
$ 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 卷积内核子集
为 Ampere 与 Turing 架构编译实现前向传播(fprop)、FP32 累加、FP16 输入的 Tensor Core 卷积内核子集:
$ 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 内核
为 Ampere 与 Turing 架构编译并运行一个实现 fprop、FP32 累加、FP32 输入的 CUDA Core 卷积内核:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ make cutlass_profiler -j16剖析单个卷积 CUDA 内核:
$ ./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: CUTLASS OperationKind: conv2d Operation: cutlass_simt_sfprop_optimized_128x128_8x2_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=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 =============================卷积剖析中的参数语义:--conv_kind=fprop指定前向卷积;--n/--h/--w/--c/--k分别为 batch、高、宽、输入通道、输出通道;--r/--s为卷积核尺寸;--p/--q为输出尺寸;--pad_h/--pad_w/--stride_h/--stride_w/--dilation_h/--dilation_w控制填充、步长与膨胀;--iterator_algorithm=optimized指定隐式 GEMM 迭代器算法;--conv_mode=cross指定相关(cross-correlation)卷积模式。
9.3 编译内核与 Profiler 的更多细节
- GEMM 与隐式 GEMM 卷积的更多选择性编译 CMake 示例,可参见仓库中的各 examples 子目录对应的 CMakeLists.txt;
- 关于 CUTLASS Profiler 的完整说明,可查阅 tools/profiler 目录下的源码与头文件。
十、许可证、贡献者与版权
CUTLASS 由 NVIDIA Corporation 以开源软件形式发布,采用3-clause "New" BSD 许可证(完整文本见 LICENSE.txt);官方开发者与贡献者列表见 CONTRIBUTORS.md;版权为 Copyright (c) 2017 - 2026 NVIDIA CORPORATION & AFFILIATES,SPDX-License-Identifier: BSD-3-Clause。
小结:CUTLASS 4.7 一方面延续了其作为"CUDA 中所有层级与规模的高性能线性代数模板库"的定位,通过 include/cutlass 与 include/cute 提供从架构指令(mma/wgmma)、CuTe Layout/Tensor 到 GEMM/卷积线程块与 warp 级实现、再到 Profiler 与 Instance Library 的全栈能力;另一方面通过 CUTLASS DSL 家族(CuTe DSL 及 cutlass_compiler/cute_ir)把这套能力开放给 Python 生态。对希望深入 CUDA 内核优化的开发者而言,从 README.md 的构建与 Profiler 章节入手,对照 CMakeLists.txt 中的CUTLASS_NVCC_ARCHS、CUTLASS_LIBRARY_KERNELS等配置,即可快速搭建起一套"选定架构 → 裁剪内核集合 → 编译 → 剖析调优"的高效工作流。
【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考