CUTLASS 编程示例全解析从基础 SGEMM 到 Blackwell Tensor Core 的实战学习路线【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlassexamples/是 CUTLASS 官方随仓库一起维护的编程示例集Programming Examples覆盖从单精度 GEMM、混合精度 Tensor Core 运算、卷积隐式 GEMM、算子融合到 Hopper/Blackwell 新一代内核的全部核心用法。本文以 examples/README.md 为骨架逐类剖析每个示例的定位并结合 00_basic_gemm 的源码、examples/CMakeLists.txt 的构建机制、CuTe 教程 与 Python 接口示例让你既能按索引快速找到目标示例也能理解示例背后 CUTLASS 模板库的组织方式最终具备看懂示例 → 修改示例 → 编写自定义内核的完整能力。重要提示原文档原意保留这些示例仅用于演示 CUTLASS 功能并未针对性能基准测试进行优化NOT for Benchmarking。若需准确的性能测量请改用 CUTLASS Profiler推荐或者在 Profiler 无法覆盖时手动对示例做自动调优。示例索引总览一份按能力分层的路线图examples/README.md 本质上是 CUTLASS 官方示例的目录索引按编号00 起组织了上百个独立示例每个示例都聚焦一个明确的功能点。综合仓库中实际的目录结构与 examples/CMakeLists.txt 中注册的构建目标可以把这些示例按能力分层归纳如下。第一层入门与基础组件00_basic_gemm以单精度FP32输入输出启动一个基础 GEMM是最简可运行内核。01_cutlass_utilities演示 CUTLASS Utilities 如何分配与初始化张量这些工具类在 tools/util 中维护。02_dump_reg_shmem调试工具用于打印寄存器register与共享内存shared memory中的内容。03_visualize_layout可视化 CUTLASS 中所有布局layout函数的工具。04_tile_iterator演示对内存中数据块tile进行迭代的迭代器。19_tensorop_canonical/20_simt_canonical分别使用 Tensor Core 与 SIMTCUDA Core实现Canonical规范GEMM是最贴近教材式数据流的内核写法适合初学者对照阅读。第二层GEMM 形态变体05_batched_gemmCUTLASS 的批处理 strided GEMM同一内核处理多个矩阵对。06_splitK_gemmSplit-K 并行归约内核把 K 维度切分给多个线程块再做累加归约。24_gemm_grouped一批问题规模各不相同的 GEMM 操作Grouped GEMM。31_basic_syrk对称 Rank-K 更新Symmetric Rank-K update。32_basic_trmm三角矩阵-矩阵乘法Triangular Matrix-Matrix multiply。33_ampere_3xtf32_tensorop_symm使用 FP32 模拟3xTF32的对称矩阵乘法。38_syr2k_grouped一批问题规模各异的 SYR2K 操作。39_gemm_permute批处理 GEMM 输出以重排张量permuted reshaped tensors形式写出。23_ampere_gemm_operand_reduction_fusion在计算 GEMM 时沿 K 维度对其中一个操作数做归约并融合进内核。第三层Tensor Core 与数据精度07_volta_tensorop_gemm使用 Volta Tensor Core 的混合精度 GEMM。08_turing_tensorop_gemm使用 Turing Tensor Core 的整数 GEMM。14_ampere_tf32_tensorop_gemmFP32 输入隐式转换为 TF32 的 GEMM。15_ampere_sparse_tensorop_gemm使用 Ampere Sparse Tensor Core2:4 结构化稀疏。18_ampere_fp64_tensorop_affine2_gemmAffine-2 GEMM。21_quaternion_gemm四元数QuaternionGEMM 计算。27_ampere_3xtf32_fast_accurate_tensorop_gemm用三次 TF32 操作模拟快速且精确的 SGEMMFP32 精度模拟的经典方案。28_ampere_3xtf32_fast_accurate_tensorop_fprop用 TF32 操作模拟快速精确的 FP32 卷积fprop。29_ampere_3xtf32_fast_accurate_tensorop_complex_gemm用 TF32 操作模拟快速精确的复数 CGEMM。35_gemm_softmax使用 Ampere Tensor Core 的混合精度 GEMM 与 Softmax 融合。58_ada_fp8_gemm通过 CUTLASS 2.x API 使用 Ada 架构 FP8 Tensor Core 的 GEMM 内核。54_hopper_fp8_warp_specialized_gemm实例化并运行一个 Hopper FP8 GEMM 内核。55_hopper_mixed_dtype_gemm使用 CUTLASS 3.x API 的 A/B 不同数据类型mixed dtypeGEMM并带有融合反量化fused dequantization。第四层卷积隐式 GEMM09_turing_tensorop_conv2dfprop使用 Turing Tensor Core 的整数隐式 GEMM 卷积前向传播 fprop。16_ampere_tensorop_conv2dfpropNHWC 布局张量上的前向卷积。17_fprop_per_channel_bias卷积与逐通道偏置per channel bias及 ReLU 融合。22_quaternion_conv四元数卷积。25_ampere_fprop_mainloop_fusion将激活函数的逐通道 scalebiasReLU 融合进 fprop 主循环mainloop。26_ampere_wgrad_mainloop_fusion将逐通道 scalebiasReLU 融合进 wgrad 主循环。30_wgrad_split_k结合 Split-K 计算卷积对权重的梯度wgrad。34_transposed_conv2d使用 CUTLASS conv2d Dgrad 内核计算二维转置卷积又称反卷积 deconvolution。42_ampere_tensorop_group_conv使用 Tensor Core 运行分组卷积group convolution。46_depthwise_simt_conv2dfprop使用 SIMT 指令的深度可分离depthwise二维卷积。第五层算子融合与内存模式12_gemm_bias_reluGEMM 融合偏置bias与 ReLU这是 epilogue 融合的入门示例。13_two_tensor_op_fusion在一个内核中融合两个 GEMM 或两个卷积。36_gather_scatter_fusion把 GEMM 前的 gather 与 GEMM 后的 scatter 融合进同一个 GEMM 内核。37_gemm_layernorm_gemm_fusion把 gemm→layernorm→gemm 三段融合进一个内核。45_dual_gemm把共享同一左输入矩阵的两个 GEMM 融合进一个内核。52_hopper_gather_scatter_fusionHopper 架构下把 gather 与 scatter 融合进同一内核。53_hopper_gemm_permuteHopper 架构下把张量置换permutation操作与 GEMM 内核融合。61_hopper_gemm_with_topk_and_softmax带 Top-K 与 Softmax epilogue 融合的 Hopper GEMM 内核。第六层复数与注意力10_planar_complex平面复数planar complexGEMM 内核。11_planar_complex_array支持批次各自独立问题规模的平面复数内核。41_fused_multi_head_attentionREADME 中写作 41_multi_head_attention非固定序列长度输入的自注意力示例目录中同时包含前向fused_multihead_attention_fixed_seqlen.cu、fused_multihead_attention_variable_seqlen.cu与反向fused_multi_head_attention_backward.cu实现。44_fused_multi_head_attentionREADME 中写作 44_fused_multi_head_attention使用共享内存的融合多头注意力支持固定与可变序列长度。第七层稀疏计算43_ell_block_sparse_gemmBlock-Ell 稀疏 GEMM。80_blackwell_geforce_sparse_gemm面向 GeForce RTX 50 系列SM120CUDA Core 的稀疏 MMA 内核。83_blackwell_sparse_gemmBlackwell SM100 稀疏 GEMM 内核。84_blackwell_narrow_precision_sparse_gemmBlackwell Block Scaled SM100 稀疏 GEMM 内核。第八层并行分解与调度Stream-K47_ampere_gemm_universal_streamk对比 GEMM 线程块的 Stream-K 并行分解与经典数据并行classic>#include cutlass/gemm/device/gemm.h // 通用 GEMM 计算模板类 cudaError_t CutlassSgemmNN(int M, int N, int K, float alpha, float const *A, int lda, float const *B, int ldb, float beta, float *C, int ldc) { using ColumnMajor cutlass::layout::ColumnMajor; // 用 5 个模板参数定义 CUTLASS GEMM 类型A/B/C 的数据类型与布局 using CutlassGemm cutlass::gemm::device::Gemm float, // A 的数据类型 ColumnMajor, // A 的布局 float, // B 的数据类型 ColumnMajor, // B 的布局 float, // C 的数据类型 ColumnMajor; // C 的布局 CutlassGemm gemm_operator; // 构造参数对象问题维度、A/B/C 的张量引用指针leading dimension、D 矩阵、alpha/beta 标量 CutlassGemm::Arguments args( {M, N, K}, // GEMM 问题维度 {A, lda}, // 源矩阵 A 的张量引用 {B, ldb}, // 源矩阵 B 的张量引用 {C, ldc}, // 源矩阵 C 的张量引用 {C, ldc}, // 目标矩阵 D可与源 C 不同内存 {alpha, beta}); // Epilogue 中使用的标量 cutlass::Status status gemm_operator(args); // 启动内核 return (status cutlass::Status::kSuccess) ? cudaSuccess : cudaErrorUnknown; }这段代码体现了一个贯穿 CUTLASS 的经典设计模式把主机端可构造的 GEMM 参数指针、步长、标量等打包成参数对象按值传给内核好处是参数结构化、可组合同时把内核入口的初始化开销降到最低。模板参数只指定了数据类型与布局线程块 tile 尺寸默认 128x128x8、warp 划分、流水线阶段数等调优参数均由 default_gemm_configuration.h 提供合理的默认值完整的设备端 API 接口可见 gemm.h。main函数中可以看到该示例的命令行用法源码注释与解析逻辑均给出00_basic_gemm M N K alpha beta默认问题规模为128 128 128alpha、beta默认取1与0即标准C A * B。程序随后会做两件事来验证正确性其一用InitializeMatrix_kernel填充三张矩阵A 用种子 0、B 用种子 17、C 用种子 101其二用朴素参考 GEMM 内核ReferenceGemm_kernel在设备端逐元素计算C alpha * A * B beta * C再把 CUTLASS 结果与参考结果拷回主机逐位比较一致则打印Passed.。这种CUTLASS 内核 朴素参考内核 逐位校验的验证模式几乎贯穿所有 CUTLASS 示例与单元测试是学习每个示例时值得留意的通用套路。示例的构建与运行CMake 机制与实操命令CUTLASS 是 header-only 模板库示例、单元测试与工具通过 CMake 构建。examples/目录自身的 CMakeLists.txt 定义了统一的示例注册机制公共源码目录CUTLASS_EXAMPLES_COMMON_SOURCE_DIR指向 examples/common其中的 helper.h 等头文件被所有示例共享00_basic_gemm顶部即#include helper.h自定义函数cutlass_example_add_executable包装了cutlass_add_executable为每个示例链接CUTLASS、cutlass_tools_util_includes并依据CUTLASS_ENABLE_CUBLAS可选链接 cuBLAS所有示例目录通过foreachadd_subdirectory批量注册同时挂到cutlass_examples与test_examples两个自定义目标上后者负责把每个示例注册为可执行测试。构建示例的典型流程前提CUDACXX指向 CUDA Toolkit 的 nvcc$ export CUDACXX${CUDA_INSTALL_PATH}/bin/nvcc $ mkdir build cd build $ cmake .. -DCUTLASS_NVCC_ARCHS80 # 只编译 Ampere 架构显著缩短编译时间 $ make 00_basic_gemm -j # 只构建单个示例 # 或者构建全部示例 $ make cutlass_examples -j运行示例同样简单$ ./examples/00_basic_gemm/00_basic_gemm 128 128 128 1 0 Passed.如需获取准确的性能数据请构建并运行 CUTLASS Profilermake cutlass_profiler -j16它支持用--kernels通配符精确筛选内核、用--m/--n/--k指定问题规模并输出运行时、显存带宽GiB/s与算力GFLOP/s等指标详细用法见仓库根 README.md 的 Performance Profiling 一节。CuTe 编程示例脱离 CUTLASS 直接学习布局与内核编写examples/README.md 明确指出不依赖 CUTLASS、直接展示 CuTe 特性的示例位于 examples/cute。CuTe 是 CUTLASS 3.0 引入的核心库用于描述和操作线程与数据的分层多维布局提供Layout与Tensor对象。examples/cute/tutorial 下的教程式示例按硬件代际递进基础 SGEMM 教学sgemm_1.cu最简、sgemm_2.cu、sgemm_sm70.cu、sgemm_sm80.cu逐步引入架构特性数据搬运tiled_copy.cu分块拷贝、tiled_copy_if.cu带谓词的分块拷贝Hopperhopper/wgmma_sm90.cuwarpgroup 级 MMA 指令、hopper/wgmma_tma_sm90.cu结合 TMA 异步拷贝Blackwellblackwell/下的01_mma_sm100.cu、02_mma_tma_sm100.cu、03_mma_tma_multicast_sm100.cu、04_mma_tma_2sm_sm100.cu、05_mma_tma_epi_sm100.cu从裸 MMA 逐步走向 TMA、多播multicast、双 SM 与 epilogue 组合。此外CuTe 核心的布局Layout与布局代数layout algebra在 test/unit/cute/core 中拥有自己的单元测试例如composition.cpp、complement.cpp、coalesce.cpp、inverse_left.cpp、logical_divide.cpp、compact_xmajor.cpp等。这些测试本身就是很好的 CuTe 使用范例——每个文件都针对一个布局代数运算复合、补集、合并、逆、逻辑划分等给出了可运行的调用代码与断言是比教程更细粒度的学习材料。Python 接口示例从 C 到 Python 的桥梁CUTLASS 同时提供 Python 接口相关示例位于 examples/python与 python/ 目录中的 Python 包配套。从 python/README.md 可以确认仓库维护着cutlass_cppgenCUTLASS Python 接口可让用户在 Python 中编译并运行 CUTLASS 内核与cutlass_library用于枚举和生成 CUTLASS 内核 C 代码的工具两个包。其典型用法是把操作声明成高层对象后直接运行import cutlass import numpy as np plan cutlass.op.Gemm(elementnp.float16, layoutcutlass.LayoutType.RowMajor) A, B, C, D [np.ones((1024, 1024), dtypenp.float16) for i in range(4)] plan.run(A, B, C, D)需要注意该接口的定位它优先易用性会自动为操作选择合理默认配置但不以选取最优内核配置为目标追求极限性能应使用 CUTLASS Profiler 手动调优也不是快速的容器或 Python-to-CUDA JIT 引擎部署高性能内核时应直接使用接口产出的 C 代码或利用其 emitter 自动生成框架扩展如 PyTorch CUDA 扩展。examples/python 目录下可以看到三个子区域CuTeDSL/CuTe 领域专用语言示例对应 CUTLASS 4 新增的 Python 原生 DSL、advanced_compiler_control/编译器高级控制与deprecated/已弃用示例按需取用即可。README 中对应的入门示例是40_cutlass_py。如何高效利用这份示例集结合上述分层推荐的学习路径是先跑通 00_basic_gemm 理解模板实例化 参数对象 内核启动三步范式再用 12_gemm_bias_relu 与 35_gemm_softmax 理解 epilogue 融合接着对照 19_tensorop_canonical、20_simt_canonical 与 47_ampere_gemm_universal_streamk 理解不同并行分解最后按目标硬件进入 Hopper48_hopper_warp_specialized_gemm与 Blackwell70_blackwell_gemm系列。若你关心的是卷积则从 16_ampere_tensorop_conv2dfprop 与 34_transposed_conv2d 入手若关心推理融合36_gather_scatter_fusion、37_gemm_layernorm_gemm_fusion、41_fused_multi_head_attention 是更贴近真实业务的参考。需要再次强调的是示例是用来理解 CUTLASS 功能与 API 语义的不是用来跑分的。在发表任何性能结论之前务必改用 CUTLASS Profiler 进行基准测量或在 Profiler 不可用时对示例进行手动自动调优否则得到的数字既不准确也缺乏可比性。【免费下载链接】cutlassCUDA Templates and Python DSLs for High-Performance Linear Algebra项目地址: https://gitcode.com/GitHub_Trending/cu/cutlass创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考