CUTLASS 4.7 深度指南：CUDA 高性能线性代数模板库与 CuTe DSL 全景解析

CUTLASS 4.7 深度指南CUDA 高性能线性代数模板库与 CuTe DSL 全景解析【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlassCUTLASS 是 NVIDIA 开源的 CUDA C 模板库为在 CUDA 的所有层级与规模下实现高性能矩阵乘法GEMM及相关计算提供了分层分解与数据搬移的模块化抽象CUTLASS 4 又在此基础上加入了 Python 原生的CUTLASS DSLCuTe 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_gemm8 位浮点类型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.02026 年 8 月发布4.7 的新特性分布在三个层面2.1 CuTe DSLPrimitives 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 → e2m1FP4的转换路径。优化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 H100Hopper 架构上的持续性能改进。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.hppcute::Mma_Atom与cute::TiledMmaatom/copy_atom.hppcute::Copy_Atom与cute::TiledCopyatom 下的*sm*.hpp各架构arch特有的 copy 与 math 操作元信息archcopy 与 math 指令的极简 PTX 封装结构体algorithmcopy、gemm 及cute::tuple操作等核心操作定义。CUTLASS 3.0 及后续版本在其模板中全面采用 CuTe 贯穿 GEMM 层次大大简化了设计并提升了代码的可组合性与可读性。五、兼容性矩阵架构、编译器与操作系统5.1 最低要求架构Voltacompute capability 7.0编译器至少支持 C17CUDA Toolkit 版本11.4CUTLASS 要求 C17 宿主编译器并且使用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.0Ubuntu 20.04GCC 10.3.0Ubuntu 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 ToolkitNVIDIA V100 Tensor Core GPU7.011.4NVIDIA TitanV7.011.4NVIDIA GeForce RTX 20x0 series7.511.4NVIDIA T47.511.4NVIDIA A100 Tensor Core GPU8.011.4NVIDIA A108.611.4NVIDIA GeForce RTX 30x0 series8.611.4NVIDIA GeForce RTX 40x0 series8.911.8NVIDIA L408.911.8NVIDIA H100 Tensor Core GPU9.011.8NVIDIA H200 Tensor Core GPU9.011.8NVIDIA B200 Tensor Core GPU10.012.8NVIDIA B300 Tensor Core GPU10.313.0NVIDIA DRIVE Thor11.013.0NVIDIA GeForce RTX 50x0 series12.012.8NVIDIA DGX Spark12.113.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_ARCHS90a或cmake .. -DCUTLASS_NVCC_ARCHS100a注意数据中心产品使用的 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_ARCHS80 # compiles for NVIDIAs 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 builds 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 与 ToolsCUTLASS SDK examples位于 examples将 CUTLASS 模板应用到具体的基础计算中如 FP32 仿真 GEMMexamples/27_ampere_3xtf32_fast_accurate_tensorop_gemm与 Blackwell FMHAexamples/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与 pythonPython 侧 DSL 与绑定等目录构成了完整的 C / Python 双栈项目形态。九、性能剖析实战CUTLASS Profilertools/profiler 目录包含一个命令行工具用于启动各类 GEMM 内核进行性能剖析构建方式如下$ make cutlass_profiler -j169.1 构建全部 GEMM 与卷积内核编译时间长默认情况下每种数据类型、数学指令与布局只实例化一种 tile 尺寸。若想实例化全部需要在空的 build/ 目录中运行 CMake 时设置以下环境变量。注意这会产生数万个内核和极长的构建时间同时二进制体积巨大部分平台上链接器可能因此构建库失败。因此强烈建议只生成内核的一个子集见 9.2 节$ cmake .. -DCUTLASS_NVCC_ARCHS90a -DCUTLASS_LIBRARY_KERNELSall ... $ 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_ARCHS75;80 -DCUTLASS_LIBRARY_KERNELScutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ make cutlass_profiler -j16对其中的 Tensor Core GEMM 内核子集进行剖析的示例命令与输出./tools/profiler/cutlass_profiler --kernelscutlass_tensorop_s*gemm_f16_*_nt_align8 --m3456 --n4096 --k4096 ... 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_kinduniversal --m3456 --n4096 --k4096 --Af16:column --Bf16:row --Cf32:column --alpha1 \ --beta0 --split_k_slices1 --batch_count1 --op_classtensorop --accumf32 --cta_m256 --cta_n128 \ --cta_k32 --stages2 --warps_m4 --warps_n2 --warps_k1 --inst_m16 --inst_n8 --inst_k8 --min_cc75 \ --max_cc1024 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_kCTA tile、warps_m/warps_n/warps_kwarp 划分、inst_m/inst_n/inst_k指令形状、stages流水线级数、min_cc/max_cc适用 compute capability 范围MemoryGiB/s与MathGFLOP/s则是剖析的核心性能指标。场景二单个 CUDA Core GEMM 内核为 Ampere 与 Turing 架构编译单个 SGEMM 内核$ cmake .. -DCUTLASS_NVCC_ARCHS75;80 -DCUTLASS_LIBRARY_KERNELScutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ make cutlass_profiler -j16剖析单个 SGEMM CUDA 内核$ ./tools/profiler/cutlass_profiler --kernelssgemm --m3456 --n4096 --k4096 Problem ID: 1 Provider: CUTLASS OperationKind: gemm Operation: cutlass_simt_sgemm_128x128_8x2_nn_align1 Status: Success Verification: ON Disposition: Passed cuBLAS: Passed Arguments: --m3456 --n4096 --k4096 --Af32:column --Bf32:column --Cf32:column --alpha1 --beta0 --split_k_slices1 \ --batch_count1 --op_classsimt --accumf32 --cta_m128 --cta_n128 --cta_k8 --stages2 --warps_m4 \ --warps_n2 --warps_k1 --inst_m1 --inst_n1 --inst_k1 --min_cc50 --max_cc1024 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_ARCHS75;80 -DCUTLASS_LIBRARY_KERNELScutlass_tensorop_s*fprop_optimized_f16 ... $ make cutlass_profiler -j16剖析 Tensor Core 卷积内核子集$ ./tools/profiler/cutlass_profiler --kernelscutlass_tensorop_s*fprop_optimized_f16 --n8 --h224 --w224 --c128 --k128 --r3 --s3 ... 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_kindfprop --n8 --h224 --w224 --c128 --k128 --r3 --s3 --p224 --q224 --pad_h1 --pad_w1 \ --stride_h1 --stride_w1 --dilation_h1 --dilation_w1 --Activationf16:nhwc --Filterf16:nhwc --Outputf32:nhwc \ --conv_modecross --iterator_algorithmoptimized --alpha1 --beta0 --split_k_modeserial --split_k_slices1 \ --eq_gemm_providernone --op_classtensorop --accumf32 --cta_m128 --cta_n128 --cta_k32 --stages5 \ --warps_m2 --warps_n2 --warps_k1 --inst_m16 --inst_n8 --inst_k16 --min_cc80 --max_cc1024 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_ARCHS75;80 -DCUTLASS_LIBRARY_KERNELScutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ make cutlass_profiler -j16剖析单个卷积 CUDA 内核$ ./tools/profiler/cutlass_profiler --kernelscutlass_simt_sfprop_optimized_128x128_8x2_nhwc --n8 --h224 --w224 --c128 --k128 --r3 --s3 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_kindfprop --n8 --h224 --w224 --c128 --k128 --r3 --s3 --p224 --q224 --pad_h1 --pad_w1 \ --stride_h1 --stride_w1 --dilation_h1 --dilation_w1 --Activationf32:nhwc --Filterf32:nhwc --Outputf32:nhwc \ --conv_modecross --iterator_algorithmoptimized --alpha1 --beta0 --split_k_modeserial --split_k_slices1 \ --eq_gemm_providernone --op_classsimt --accumf32 --cta_m128 --cta_n128 --cta_k8 --stages2 --warps_m4 \ --warps_n2 --warps_k1 --inst_m1 --inst_n1 --inst_k1 --min_cc50 --max_cc1024 Bytes: 2055798784 bytes FLOPs: 118482796544 flops Runtime: 7.34266 ms Memory: 260.752 GiB/s Math: 16136.2 GFLOP/s 卷积剖析中的参数语义--conv_kindfprop指定前向卷积--n/--h/--w/--c/--k分别为 batch、高、宽、输入通道、输出通道--r/--s为卷积核尺寸--p/--q为输出尺寸--pad_h/--pad_w/--stride_h/--stride_w/--dilation_h/--dilation_w控制填充、步长与膨胀--iterator_algorithmoptimized指定隐式 GEMM 迭代器算法--conv_modecross指定相关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 AFFILIATESSPDX-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），仅供参考