CUTLASS默认模板实例化（如仅传四个类型参数）采用128x128x8分块等通用配置，虽能正确计算，但在小矩阵、非对齐shape或需bias/relu融合时性能常降30%～50%，因未显式控制分块、流水线深度、epilogue等关键变量。
直接用 cutlass::gemm::device::Gemm 模板能跑，但性能不达标
默认模板实例化（比如只传四个类型参数）会走预设的“通用路径”，在 Ampere 或 Hopper 上大概率用的是 128x128x8 分块 + 默认共享内存布局 + 无 Split-K。它能正确计算，但面对小矩阵、非对齐 shape、或需要融合 bias/relu 的场景，吞吐常掉 30%～50%。关键不是写不出来，而是没显式控制分块、流水线深度、epilogue 类型这些变量，等于把性能交给编译器猜。
Gemm 模板的五个必填维度：数据类型、布局、分块、流水线、epilogue
真正定制化的起点不是改 kernel，而是补全模板参数列表。CUTLASS 3.x 把 GEMM 拆成五层抽象，每一层都对应一个模板参数槽位：
-
ElementA/ElementB/ElementC：不只是float，还要考虑cutlass::half_t、cutlass::bfloat16_t，甚至int8_t+ scale； -
LayoutA/LayoutB/LayoutC：列主序（ColumnMajor）是 cuBLAS 兼容模式，但行主序（RowMajor）在某些访存 pattern 下更省 bank conflict； -
ThreadblockShape：不能只依赖默认值，比如cutlass::gemm::GemmShape对 M=1024/N=2048/K=512 更友好； -
WarpShape：必须和ThreadblockShape整除，且要匹配 Tensor Core 的 MMA 指令粒度（如32x32x16对应 FP16 MMA）； -
EpilogueOutputOp：这是融合算子的关键入口，cutlass::epilogue::thread::LinearCombination支持 alpha/beta，但若要加 relu，得换cutlass::epilogue::thread::ReLu并确保 C 数据类型可被 inplace 修改。
编译时选型错误会导致 nvcc 静默退化到 fallback kernel
CUTLASS 不报错，但生成的 kernel 性能断崖下跌，常见于三类配置冲突：
- 指定
cutlass::arch::OpClassTensorOp却在 SM 75（Turing）上编译——Tensor Core 在 Turing 上仅支持 INT8/FP16，FP32 会 fallback 到 SIMT； -
ElementA是cutlass::half_t，但没配cutlass::arch::OpClassTensorOp和对应ArchTag（如cutlass::arch::Sm80），结果用 SIMT 指令模拟 half 计算； - 自定义
Swizzle布局（如cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle）但没同步调整SharedMemoryAlignment，导致 bank conflict 实际未消除。
验证方法：用 cutlass_profiler 加 --kernels=all 看实际 dispatch 的 kernel 名，含 tensorop 字样的才是走 Tensor Core 路径。
调试时最容易忽略的是 Arguments 构造顺序和 stride 含义
模板类型定好后，运行时传参出错几乎 100% 导致结果错或越界，尤其注意：
-
TensorRef的第二个参数是 leading dimension，不是 total size。例如 A 是MxK列主序，lda == M，不是K； -
alpha/beta必须传入cutlass::NumericConverter包装后的值，否则混合精度下可能被截断（如 FP16 alpha 传 float 常量 2.0f，实际进 kernel 变成 2.0h）； - Split-K 场景下，
split_k_slices必须是 2 的幂且 ≤ K，否则cutlass::Status::kErrorInvalidProblem不抛异常，而是静默降级为单 slice。
建议在构造 Arguments 后立刻调用 gemm_operator.can_implement(args) 做合法性检查，比等 kernel launch 失败再排查快得多。
C++免费学习笔记（深入）：立即使用
在学习笔记中，你将探索 C++ 的入门与实战技巧！
