CUTLASS 4.7 深度指南:CUDA 高性能线性代数模板库与 CuTe DSL 全景解析
2026/9/15 20:56:31 网站建设 项目流程

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 的mmawgmma指令实现:

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)。它提供LayoutTensor对象,将数据的类型、形状、内存空间与布局紧凑地打包在一起,同时替用户完成复杂的索引计算,让程序员专注于算法的逻辑描述,而由 CuTe 负责机械化的簿记工作。

CuTe 的核心抽象——分层多维布局——可以与数据数组组合来表示张量,其布局表示能力足以覆盖高效稠密线性代数所需的几乎所有场景;布局之间还可以通过**函数复合(functional composition)**进行组合与操作,并在此基础上构建起 tiling、partitioning 等大量常见操作。

在仓库中,CuTe 的实现位于 include/cute,其内部组织为:

  • layout.hpp、tensor.hpp:核心类型ShapeStrideLayoutTensor及关联操作;
  • atom/mma_atom.hpp:cute::Mma_Atomcute::TiledMma
  • atom/copy_atom.hpp:cute::Copy_Atomcute::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.04GCC 7.5.0
Ubuntu 20.04GCC 10.3.0
Ubuntu 22.04GCC 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 上保持高效:

GPUCUDA Compute CapabilityCUTLASS-3 要求的最低 CUDA Toolkit
NVIDIA V100 Tensor Core GPU7.011.4
NVIDIA TitanV7.011.4
NVIDIA GeForce RTX 20x0 series7.511.4
NVIDIA T47.511.4
NVIDIA A100 Tensor Core GPU8.011.4
NVIDIA A108.611.4
NVIDIA GeForce RTX 30x0 series8.611.4
NVIDIA GeForce RTX 40x0 series8.911.8
NVIDIA L408.911.8
NVIDIA H100 Tensor Core GPU9.011.8
NVIDIA H200 Tensor Core GPU9.011.8
NVIDIA B200 Tensor Core GPU10.012.8
NVIDIA B300 Tensor Core GPU10.313.0
NVIDIA DRIVE Thor11.013.0
NVIDIA GeForce RTX 50x0 series12.012.8
NVIDIA DGX Spark12.113.0

六、目标架构与 architecture-accelerated features

一般来说,为某目标架构生成的 PTX 代码可以前向兼容地运行在未来的架构上。但 CUDA 12.0 引入了"架构加速特性(architecture-accelerated features)",这类特性的 PTX不保证前向兼容。若干 Hopper 与 Blackwell 的 PTX 指令即属于此类,因此需要使用sm_90asm100a目标架构(注意结尾的 "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 -j16

9.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_ARCHSCUTLASS_LIBRARY_KERNELS等配置,即可快速搭建起一套"选定架构 → 裁剪内核集合 → 编译 → 剖析调优"的高效工作流。

【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass

创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考

需要专业的网站建设服务?

联系我们获取免费的网站建设咨询和方案报价,让我们帮助您实现业务目标

立即咨询