企业官网建设流程全解析

CUTLASS Blackwell Blockwise/Groupwise GEMM 完全指南：FP8 分块缩放与分组 GEMM 的配置、剖析与性能调优

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

导读

本文围绕 CUTLASS 在 NVIDIA Blackwell（SM100）架构上提供的 Blockwise（分块缩放）与 Groupwise（分组缩放）GEMM 以及 Grouped GEMM（分组批 GEMM）展开，讲解如何通过"按累加器类型的软件缩放"在可配置粒度上控制数值精度，特别适用于量化神经网络中对张量不同区域采用不同缩放需求（per-block/per-group 缩放）的场景。读完本文，你将掌握缩放因子张量 SFA/SFB 的数学语义与 CuTe 布局表示、Sm100BlockwiseScaleConfig的配置方法、跨框架（如 PyTorch）张量格式转换要点、利用 CUTLASS Profiler 自动调优选择最优 kernel 的完整命令行流程，以及 kernel 命名规范和 MMA 维度相关的性能调优技巧。

本指南的主体基于 examples/81_blackwell_gemm_blockwise/README.md，并辅以 include/cutlass/detail/blockwise_scale_layout.hpp、include/cutlass/gemm/collective/sm100_mma_warpspecialized_blockwise_scaling.hpp 等源码与 examples/81_blackwell_gemm_blockwise 目录下的四个示例程序进行纵深佐证。

背景：为什么需要 Blockwise/Groupwise 缩放

标准的 GEMM 计算为 $D = \alpha A B + \beta C$。在量化推理场景中，若将 A、B 量化到 FP8（如e4m3），一个全局缩放因子往往无法同时兼顾张量不同区域动态范围的差异。Blockwise/Groupwise GEMM 的核心思想是引入两个缩放因子张量SFA 与 SFB，把计算改写为：

$$D = \alpha \ (\text{SFA} * A) \ (\text{SFB} * B) + \beta C$$

其中*表示逐元素相乘。缩放因子以可配置的粒度作用于 A 和 B，使得张量的不同区域可以拥有各自独立的缩放，从而实现更细粒度的数值精度控制。CUTLASS 通过**软件方式（按累加器类型）**完成这一缩放：缩放因子以f32（累加器类型）表示，而 A、B 以e4m3等低精度类型参与矩阵乘，缩放发生在计算过程中。

从源码结构看，这类 kernel 的实现分布在 include/cutlass/gemm/collective/sm100_mma_warpspecialized_blockwise_scaling.hpp（单 CTA 的 Warp Specialized 主循环）与 include/cutlass/gemm/collective/sm100_mma_array_warpspecialized_blockwise_scaling.hpp（指针数组/Grouped 形式）中，对应 include/cutlass/gemm/dispatch_policy.hpp 中定义的调度策略：

• KernelScheduleSm100Blockwise：Blockwise 调度的基类；
• KernelTmaWarpSpecializedBlockwise1SmSm100/KernelTmaWarpSpecializedBlockwise2SmSm100：1SM 与 2SM 的 TMA + Warp Specialized 变体；
• KernelScheduleSm100PtrArrayBlockwise与KernelPtrArrayTmaWarpSpecializedBlockwise1SmSm100：服务于 Grouped GEMM 的指针数组（PtrArray）形式。

缩放因子张量 SFA 与 SFB

缩放因子张量由两个粒度参数完全确定：

• SFA：作用于 A。在由scale granularity m与scale granularity k定义的分块内广播同一个缩放值。这两个粒度参数也常被称为scale vector m与scale vector k。
• SFB：作用于 B。在由scale granularity n与scale granularity k定义的分块内广播同一个缩放值。同理也称为scale vector n与scale vector k。

也就是说，SFA 的尺寸为 $(M / \text{scale granularity M}) \times (K / \text{scale granularity K})$，SFB 的尺寸为 $(N / \text{scale granularity N}) \times (K / \text{scale granularity K})$（均不计 batch 维）。

CuTe 布局表示

这两个张量在 CuTe 中可以分别表示为如下 Layout：

• SFA Layout： $((\text{scale granularity M},\ M / \text{scale granularity M}),\ (\text{scale granularity K},\ K / \text{scale granularity K})) : ((0,\ int),\ (0,\ int))$
• SFB Layout： $((\text{scale granularity N},\ N / \text{scale granularity N}),\ (\text{scale granularity K},\ K / \text{scale granularity K})) : ((0,\ int),\ (0,\ int))$

其中内层形如(scale granularity, count)的 mode 对，配合 stride 中的0 元素步长（(0, int)），确保同一分块内的所有坐标都映射到缩放因子张量中的同一个元素——这正是"块内广播"的布局级实现。在 include/cutlass/detail/blockwise_scale_layout.hpp 中可以看到与之对应的类型定义：

using ShapeSFA = Shape<Shape<Int<SFVecSizeM>, int32_t>, Shape<Int<SFVecSizeK>, int32_t>, int32_t>; using StrideSFA = conditional_t<majorSFA == UMMA::Major::MN, Stride<Stride<_0,_1>, Stride<_0,int32_t>, int32_t>, Stride<Stride<_0,int32_t>, Stride<_0,_1>, int32_t>>;

配置：Sm100BlockwiseScaleConfig

为方便使用，Blockwise 与 Groupwise 的实现提供了配置类：

cutlass::detail::Sm100BlockwiseScaleConfig<ScaleGranularityM, ScaleGranularityN, ScaleGranularityK>

它负责推导 SFA/SFB 的 Layout并管理紧凑张量。默认情况下，该配置使所有张量以 M/N mode 为主序（major），但可以通过模板参数覆盖。例如：

cutlass::detail::Sm100BlockwiseScaleConfig<ScaleGranularityM, ScaleGranularityN, ScaleGranularityK, UMMA::Major::K, UMMA::Major::MN>

表示 SFA 在 K 维度上为主序，而 SFB 在 N 维度上为主序。模板定义的完整签名（见 include/cutlass/detail/blockwise_scale_layout.hpp）为：

template<int SFVecSizeM, int SFVecSizeN, int SFVecSizeK, UMMA::Major majorSFA = UMMA::Major::MN, UMMA::Major majorSFB = UMMA::Major::MN> struct Sm1xxBlockwiseScaleConfig { ... }; template<int SFVecSizeM, int SFVecSizeN, int SFVecSizeK, UMMA::Major majorSFA = UMMA::Major::MN, UMMA::Major majorSFB = UMMA::Major::MN> using Sm100BlockwiseScaleConfig = Sm1xxBlockwiseScaleConfig<SFVecSizeM, SFVecSizeN, SFVecSizeK, majorSFA, majorSFB>;

同一文件还提供了Sm120BlockwiseScaleConfig与Sm90BlockwiseScaleConfig（SM90 目前仅支持 MN 主序的 SFA/SFB），以及"平凡"（trivial）配置辅助函数sm100_trivial_blockwise_scale_config(MmaTileShape_MNK{})——它将缩放粒度直接取为 MMA tile 的各维大小。

该配置类提供了两组关键接口：

• deduce_layoutSFA()/deduce_layoutSFB()：返回仅含静态零步长信息的"原子" Layout，步长只遍历标量、不含零；
• smem_atom_layoutSFA()/smem_atom_layoutSFB()：根据 CTA tile 形状推导共享内存中的原子布局；
• tile_atom_to_shape_SFA(problem_shape)/tile_atom_to_shape_SFB(problem_shape)：将原子布局平铺（tile）到实际的动态问题尺寸（M、N、K、L）上，此时步长才包含零。

示例程序 81_blackwell_gemm_blockwise.cu 中即采用了 trivial 配置：

using ScaleConfig = decltype(cutlass::detail::sm100_trivial_blockwise_scale_config(MmaTileShape_MNK{})); using LayoutSFA = decltype(ScaleConfig::deduce_layoutSFA()); using LayoutSFB = decltype(ScaleConfig::deduce_layoutSFB());

而 81_blackwell_gemm_groupwise.cu 则展示了显式指定粒度的 Groupwise 用法：

constexpr int ScaleGranularityM = 1; constexpr int ScaleGranularityN = 128; constexpr int ScaleGranularityK = 128; using ScaleConfig = cutlass::detail::Sm100BlockwiseScaleConfig<ScaleGranularityM, ScaleGranularityN, ScaleGranularityK>;

粒度约束（编译期检查）

从 include/cutlass/gemm/collective/sm100_mma_warpspecialized_blockwise_scaling.hpp 的实现可以看出，缩放粒度并非任意取值，主循环在编译期通过static_assert施加如下约束：

• ScaleGranularityM 必须整除 TileShape 的 M 维，且不大于 TileShape 的 M 维；
• ScaleGranularityN 必须整除 TileShape 的 N 维，且不大于 TileShape 的 N 维；
• ScaleGranularityK 必须整除 TileShape 的 K 维，且不大于 TileShape 的 K 维；
• ScaleGranularityK 必须能被 MMA 指令的 K 维（MMA_K）整除——缩放粒度不能小于单个 MMA 原子指令在 K 方向覆盖的范围，这保证了缩放可以按 MMA 指令为单位对齐执行。

与其他框架（如 PyTorch）的集成

从 Torch 等框架迁移时，SFA 的形状通常为 $(M / \text{ScaleGranularityM},\ K / \text{ScaleGranularityK})$，而 SFB 的形状为 $(K / \text{ScaleGranularityK},\ N / \text{ScaleGranularityN})$。接入 CUTLASS 时需要注意：

1. 对 SFB 与 B 执行转置，使其符合 CUTLASS 的规范 CuTe 布局形式，保证K 始终是第二个 mode；
2. 利用张量自身的 strides 判断每个张量是 MN 主序还是 K 主序，据此直接构造布局，或使用上述便捷包装类（Sm100BlockwiseScaleConfig等）完成布局推导。

换句话说，只要保证 SFA 的 mode 顺序为 (M, K)、SFB 的 mode 顺序为 (K, N)，并把每个 mode 内的粒度信息正确映射到(granularity, count)的 CuTe mode 对中，即可无缝对接。

Kernel 选择与 Profiling

要为自己的 workload 确定性能最优的 Blockwise/Groupwise GEMM 或 Grouped GEMM kernel，官方推荐使用 CUTLASS Profiler。

编译期裁剪 kernel 数量

所有使用f32缩放、e4m3或运行时f8类型的 Blockwise/Groupwise GEMM 及 Group GEMM，都可以通过在 CMake 配置时传入 kernel 子集来启用/裁剪：

-DCUTLASS_LIBRARY_KERNELS="cutlass3x*f32xe4m3_*f32xe4m3*,cutlass3x*f32xf8_*f32xf8*"

进一步地，可以通过指定 SFA 与 SFB 的缩放粒度来减少生成的 kernel 数量，例如：

-DCUTLASS_LIBRARY_KERNELS="cutlass3x*1x128f32xe4m3_*128x128f32xe4m3*"

使用 Profiler 自动调优

使用 profiler 最简单的方式是传入m、n、k以及scale_vec_size_m、scale_vec_size_n、scale_vec_size_k。这三个scale_vec_size_*参数对应 profiler 源码 tools/profiler/src/blockwise_gemm_operation_profiler.cu 中注册的命令行参数（"Scale vector size in GEMM M/N/K dimension"）。

加上enable-best-kernel-for-fixed-shape后，profiler 会对每个 kernel 执行自动调优，寻找最佳的 rasterization 顺序、swizzle 与 cluster 尺寸。通过operation标志传入blockwiseGemm或GroupedGemm，可以决定剖析哪一组操作。

例如，下面的命令剖析所有支持"scale granularity m = 1、scale granularity n = 128、scale granularity k = 128"的已编译 kernel，在 8192x8192x8192 问题规模上的性能：

cutlass_profiler --operation=blockwiseGemm \ --enable-best-kernel-for-fixed-shape \ --m=8192 --n=8192 --k=8192 \ --scale_vec_size_m=1 --scale_vec_size_n=128 --scale_vec_size_k=128 \ --verification-enabled=false

Kernel 命名规范

Blockwise 与 Groupwise kernel 的命名引入了新模式：对于每对张量缩放组合，采用<scale_granularity_m 或 scale_granularity_n>x<scale_granularity_k><累加器类型>x<被缩放张量类型>。以cutlass3x_sm100_tensorop_gemm_64x128f32xe4m3_1x128f32xe4m3_f32_f16_f16_64x128x128_1x1x1_0_nnn_align16_1sm为例，各段含义依次为：

另外值得注意：如果不需要beta * C缩放，C 可以为void（即不提供 C 张量）。

性能技巧与调优建议

MMA 维度选择

在 Blackwell 与 Hopper 的 Tensor Core 上，最小的MMA_M维是 64，但某些指令的MMA_N维可以小到 8。因此对于 M 较小的 problem size，应当考虑改为计算：

$$D^T = \alpha B^T A^T + \beta C^T$$

交换 A、B 并转置后，原本较小的 M 变成了 N 维，配合小的MMA_N可以更高效地分块（tiling），避免无谓的多余计算。

布局交换（Layout Swapping）

使用 profiler 优化时，可以交换m与n输入，并相应调整布局以反映这种交换与转置。例如，若原始布局为 row-major A、column-major B、row-major D，则交换张量后可以运行这样的 kernel：

• 左手矩阵（原 B 转置后）为 row-major；
• 右手矩阵（原 A 转置后）为 column-major；
• 输出（原 D 转置后）为 column-major。

使用 Blockwise/Groupwise GEMM 时，做上述优化必须同步交换缩放向量的大小：例如原本 scale granularity M = 1、scale granularity N = 128，交换后应运行 scale granularity M = 128、scale granularity N = 1 的 kernel。

该技巧同样体现在源码注释中：81_blackwell_gemm_groupwise.cu 指出，当一个 tile 内有多个缩放因子（如 M 方向每个 tile 有 128 个 scale）时，实现会尽可能限制在 16B 对齐以内（即 M 方向至少有 16B 的 scale）；此时可执行的最小 M 为 16。对于更小的 M，可以通过交换 A、B 并转置 A、B、C 与 scale 来规避，因为 $B^T A^T = C^T$。

参考示例与运行方式

本目录（examples/81_blackwell_gemm_blockwise）共包含四个示例：

这些示例统一以ElementA = ElementB = cutlass::float_e4m3_t、ElementAccumulator = float为数据布局（见 81_blackwell_gemm_blockwise.cu），且要求编译架构包含100a、CUDA 12+、GPU compute capability 为 100a（示例代码在cudaGetDeviceProperties中检查props.major == 10 && props.minor == 0）。构建注册见 CMakeLists.txt，其中CUTLASS_NVCC_ARCHS需匹配100a，并通过cutlass_example_add_executable注册四个可执行目标。

命令行可选项（以 blockwise 示例为例）：--m、--n、--k、--l（batch 维）、--alpha、--beta、--iterations、--skip-verification。例如：

./81_blackwell_gemm_blockwise --m=1024 --n=512 --k=1024 --alpha=2 --beta=0.707

示例在verify()阶段使用 include/cutlass/util/reference/host/gett.hpp 提供的Gemm3x参考实现（含GettBlockScalingMainloopParams）对 kernel 输出做逐元素校验，并统计平均运行时间与 GFLOPS（2 * m * n * k次浮点运算）。

小结

Blockwise/Groupwise GEMM 是 CUTLASS 面向 Blackwell SM100 提供的高精度量化 GEMM 方案：通过 SFA/SFB 两个缩放因子张量在 M/N/K 三个维度上以可配置粒度做软件缩放，结合Sm100BlockwiseScaleConfig的布局推导、Profiler 的自动调优与命名规范的快速定位，开发者可以为自己的量化模型快速找到最优 kernel 配置。性能调优时牢记两条主线：小 M 时交换 A/B 并转置计算 $D^T$，以及做布局交换时同步交换缩放粒度，即可在 Blackwell Tensor Core 上充分发挥 tcgen05 MMA 的吞吐能力。

如需深入，可继续阅读：media/docs/cpp/profiler.md（Profiler 完整使用手册）、include/cutlass/detail/blockwise_scale_layout.hpp（缩放布局推导源码）以及 include/cutlass/gemm/dispatch_policy.hpp（Blockwise 调度策略定义）。

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