先说结论Blackwell 这套新架构出来之后最让我关注的并不是纸面上翻倍的 FP8 算力反而是看起来有点“传统”的 BF16 GEMM。原因很简单大模型训练、微调、推理甚至 YOLOv8 这类目标检测任务底层跑得最勤快的算子就是 GEMM通用矩阵乘法而 BF16 又是当前训练和推理里用得最广泛的数值格式之一。把 Blackwell、BF16、GEMM 三个词拼在一起基本就拼出了现役 AI 计算最核心的一条主线。这篇文章是一个持续更新的项目手记我会把在 Blackwell 上写和调 BF16 GEMM 时积累的实现细节、架构理解、工具链用法和踩坑记录都沉淀在这里。内容不追求一次讲完而是按“硬件原理 → 软件栈 → 调优方法 → 现场问题”的顺序不断补全。如果你正在做 GPU 算子开发、大模型微调、推理服务优化或者单纯想搞清楚 NVIDIA 新架构到底改了什么这篇内容应该能帮你少走很多弯路。1. 先把这几个词拆开Blackwell、BF16、GEMM 分别解决什么问题1.1 GEMM所有大模型算力的基本盘GEMM 全称 General Matrix Multiplication公式很简单就是 C Alpha * A * B Beta * C。但就是这个看似简单的矩阵乘撑起了深度学习中绝大部分的计算量。以一个大模型的一次训练迭代为例无论是 Transformer 里的 QKV 投影、Attention 分数计算还是 MLP 层的两个线性变换本质都是 GEMM。更直观点说前向传播加反向传播GEMM 的耗时占比经常能到 80% 以上。所以在 GPU 上搞 AI 计算本质上就是在和 GEMM 打交道。你写卷积也好写 attention 也好最终都会被拆分或者映射成若干个 GEMM 子问题。这也是为什么每次新 GPU 架构发布NVIDIA 总是最先秀 GEMM 的 FP16/BF16/FP8 算力数字这个指标的提升幅度直接决定了新卡在 AI 工作负载上的战斗力。1.2 BF16数值格式里的大模型答案BF16 全称 bfloat16它是一种 16 位浮点格式分配方式是 1 位符号、8 位指数、7 位尾数。很多人第一次看到这个格式会觉得奇怪尾数只有 7 位精度不是被砍得很惨吗但放在深度学习场景里这个设计恰恰是优点。FP16 的问题是尾数 10 位、指数只有 5 位数值范围只有大约 6x10^-8 到 65504。训练过程中梯度稍微大一点或者学习率调得激进一点非常容易上溢成 inf所以需要用 loss scaling 这类技巧去续命。BF16 则把指数位加到了 8 位和 FP32 完全一致数值范围直接对齐 FP32。这意味着你可以把 FP32 模型几乎无脑转成 BF16基本不用担心数值范围爆炸的问题——这在大模型训练里是一个巨大的工程简化。而尾数只有 7 位这件事在日常使用中并没有想象中那么致命因为 GEMM 的累加过程可以使用 FP32 累加器。也就是说两个 BF16 矩阵相乘中间的乘积和累加都在 FP32 下完成最后再把结果转回 BF16。这样既享受了 BF16 低显存占用、高吞吐的好处又保住了累加精度。大模型训练实践中混合精度方案里 BF16 已经是绝对的主流选择。1.3 Blackwell为了喂饱这些 GEMM 而生的新一代 GPU 架构Blackwell 是 NVIDIA 在 Hopper 架构之后的下一代 GPU 架构。如果只看名字你可能觉得它只是数字变大了但站在 GEMM 实现的角度看这一代改动非常关键。Hopper 已经引入了专门为矩阵运算设计的 Tensor Core 指令集与异步内存搬运机制Blackwell 则是在这条路上继续加码Tensor Core 的能力更强了显存带宽和容量上去了芯片互联也更快了。对写算子的人来说最关心的其实不是纸面 TFLOPs 涨了多少而是这些算力能不能通过实际的 GEMM 内核喂饱。Blackwell 在硬件上做了很多配合性改进比如共享内存容量、异步拷贝、寄存器资源、线程调度粒度等等这些细节直接决定了你可以用多大的 tile 去切分矩阵、能塞下多少层流水线。这一章后面我会详细展开。2. 硬件层面怎么变Tensor Core、显存与指令进化的底层逻辑2.1 Tensor Core 的演进从 AI 代偿到通用矩阵单元Tensor Core 最早在 Volta/Turing 上出现时给人的感觉是一个偏科生只在特定形状的矩阵乘法上有加速效果。到 Ampere 的第三代已经支持了比较灵活的 FP16/BF16 混合精度矩阵乘。到了 Hopper 的第四代Tensor Core 进一步强化并且引入了新的矩阵指令抽象让开发者可以在 warp 级别直接发起大规模的矩阵乘操作而不是像以前那样在底层指令里一个个排列组合。Blackwell 上的 Tensor Core 属于最新一代。从公开文档和工具链信息来看它的核心思路是继续把矩阵乘单元当作机器里的一等公民来对待而不是把它藏在一个普通指令后面。从我个人的理解角度把它类比成 CPU 里的 SIMD 单元更合适当年 CPU 从标量进化到 SIMD就是让一条指令能同时处理多个数据Tensor Core 做的事情类似只不过粒度更大一条指令直接算一个子矩阵。Blackwell 把这条路的颗粒度和并发能力又推高了一截。这里有个容易误会的点Tensor Core 并不是用更大的 FP32 单元去跑矩阵乘而是用专门设计的硬件阵列去算。所以你不能拿普通 CUDA 核的 FMA 吞吐去估算 GEMM 峰值必须看 Tensor Core 的矩阵乘吞吐。调优 BF16 GEMM 的时候目标就是尽量让数据流源源不断地送进 Tensor Core不要让它空转。2.2 显存与共享内存大容量和异步搬运是两条硬腿GEMM 的性能瓶颈通常不在计算本身而在数据搬运。BF16 虽然只有 2 字节但大模型里的 GEMM 动辄是几千乘几千的矩阵数据量仍然非常可观。Blackwell 这一代把显存带宽又往上抬了一大截同时 HBM 容量也做得更大了。这解决的是能不能装下和能不能喂饱的宏观问题。更细节的是片上共享内存。一个 SM流式多处理器内部的共享内存容量越大你能缓存的 tile 就越大循环里重复加载数据的次数就越少。Blackwell 在共享内存容量上的规格比 Hopper 继续提升这意味着你可以把更大的 BLOCK_K 切块塞进去减少全局内存访问次数。GEMM 调优里经典的分块tiling策略很大程度上就是被共享内存容量卡出来的容量大了调优空间自然更大。另外Blackwell 上异步数据搬运机制也进一步强化了。TMA张量内存加速器这种异步拷贝引擎可以把数据从全局内存搬到共享内存甚至寄存器并且不占用线程的算力。想象一下一个工厂里如果工人既要搬货又要加工零件效率肯定低现在有了专门的传送带工人只需要把加工好的零件送走、把传送带上的新零件拿过来加工流水线效率自然上去。2.3 指令集与异步流水线数据不等人计算器才不会被浪费GEMM 的极致性能拼到最后就是流水线编排。一个成熟的 BF16 GEMM 内核通常会把 K 维度的循环切成很多小段每一段都执行同样的动作把 A 和 B 的下一块数据从全局内存搬到共享内存再从共享内存搬到寄存器然后丢给 Tensor Core 执行矩阵乘同时累加结果。理想状态下这几个步骤应该像流水线一样重叠起来——计算当前这一段的时候下一段的数据已经在搬运路上了。这个思路说起来容易做起来需要硬件配合。Hopper 时代引入的 mbarrier、异步拷贝等机制就是为了让软件能精确控制这个流水线。Blackwell 在指令层面继续强化了这种异步能力让开发者可以更高效地用少量线程去驱动整个数据流而不是让所有线程都阻塞在同步点上。我在写 GEMM 内核时的体会是硬件的异步能力越强软件编排的灵活性越大但调试难度也越高。以前卡在一个 barrier 上线程等一等就完事了现在的流水线里你脑子里必须时刻清楚哪块数据在全局内存、哪块在共享内存、哪块在寄存器、Tensor Core 正在算哪一块。好在主流框架已经把这一层封装得不错并不是所有人都要直接面对这些底层次细节。3. 软件栈怎么选从快速验证到极致优化的三条路径3.1 路径一cuBLAS / cuBLASLt黑盒但可靠如果你只是想尽快把 Blackwell 的 BF16 GEMM 跑起来或者要的是一个通用的高性能矩阵乘实现那直接调 cuBLAS 或者 cuBLASLt 是最省事的选择。NVIDIA 在发布新硬件的时候一定会同步把 cuBLAS 的 GEMM 内核优化到位尤其是 BF16/F16 这种主流形态。绝大多数场景下你不需要自己写内核调库就够了。cuBLASLt 比老 cuBLAS 更灵活它支持你通过矩阵布局、算法 ID、权重等参数去搜索最优实现还提供了对分段 GEMM这类特殊需求的支持。实际工程里我建议优先用 cuBLASLt 的 matmul 接口因为它的 API 设计更贴近 GEMM 的各种变体需求也方便你做 benchmark 对比。不过调库的缺点也很明显你无法控制内部实现细节。当你的矩阵形状特殊、比如 M 很小、N 很大或者 K 特别深但矩阵很瘦的时候cuBLAS 默认选到的内核不一定最优。这时候你可能会需要自己写针对性的内核或者用更底层的库去定制。3.2 路径二Triton兼顾开发效率与性能Triton 是现在非常流行的一种 GPU 编程语言特别适合快速实现深度学习算子。它的核心思想是让你用 Python 语法写出类似 CUDA kernel 的逻辑但你不必手动管理线程块、线程束的很多细节。你只需要告诉它我要把矩阵切成多大的块、每个块怎么计算Triton 会自动帮你生成 CUDA 代码并调度到硬件上。下面是一个用 Triton 写 BF16 GEMM 的最小骨架import triton import triton.language as tl triton.jit def bf16_gemm_kernel( a_ptr, b_ptr, c_ptr, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr, ): pid_m tl.program_id(0) pid_n tl.program_id(1) offs_m pid_m * BLOCK_M tl.arange(0, BLOCK_M) offs_n pid_n * BLOCK_N tl.arange(0, BLOCK_N) offs_k tl.arange(0, BLOCK_K) a_ptrs a_ptr offs_m[:, None] * stride_am offs_k[None, :] * stride_ak b_ptrs b_ptr offs_k[:, None] * stride_bk offs_n[None, :] * stride_bn acc tl.zeros((BLOCK_M, BLOCK_N), dtypetl.float32) for k in range(0, K, BLOCK_K): a tl.load(a_ptrs) b tl.load(b_ptrs) acc tl.dot(a, b, acc, out_dtypetl.float32) a_ptrs BLOCK_K * stride_ak b_ptrs BLOCK_K * stride_bk c_ptrs c_ptr offs_m[:, None] * stride_cm offs_n[None, :] * stride_cn tl.store(c_ptrs, acc.to(tl.bfloat16))Triton 的好处是代码简短、可读性强而且当你切换到 Blackwell 这类新架构时Triton 编译器通常会在第一时间适配新硬件的指令集。对于快速验证 BF16 GEMM 的性能上限、测试不同 tile 尺寸组合Triton 是我最常用的上手工具。但 Triton 也不是万能的。它的自动调度对一些非标准矩阵布局、自定义数据流要求并不灵活而且排错时你很难看到底层到底生成了什么指令。做产品级长期优化时最终往往还是要落到 CUTLASS。3.3 路径三CUTLASS把 GEMM 从算子变成模板CUTLASS 是 NVIDIA 开源的 CUDA 模板库它的目标是把高性能 GEMM 的各个组成部分比如全局内存加载、共享内存存储、寄存器分块、Tensor Core 指令调用、epilogue 计算全部模板化。你可以通过拼接不同的模板参数组合出符合自己需求的 GEMM 内核。用 CUTLASS 写 BF16 GEMM 的典型代码大致长这样using Gemm cutlass::gemm::device::Gemm cutlass::bfloat16_t, // A 元素类型 cutlass::layout::RowMajor, // A 布局 cutlass::bfloat16_t, // B 元素类型 cutlass::layout::ColumnMajor, // B 布局 cutlass::bfloat16_t, // C 元素类型 cutlass::layout::RowMajor, // C 布局 float, // 累加器类型 cutlass::arch::OpClassTensorOp, // 使用 Tensor Core cutlass::arch::Sm100, // 对应 Blackwell 的架构标识 cutlass::gemm::GemmShape128, 128, 64, // 线程块级别的 tile 形状 cutlass::gemm::GemmShape64, 64, 64, // warp 级别的 tile 形状 cutlass::gemm::GemmShape16, 8, 32 // 指令级别的 tile 形状 ;注意上面的Sm100这部分具体到不同 Blackwell 型号可能会不同建议直接查你手上那份 CUDA Toolkit 的cudaDeviceProp输出决定。实际使用中CUTLASS 的模板参数非常多不看一遍官方文档直接上手确实会懵。但它的价值在于底下每一层都是透明的你可以精确控制每个阶段的算法也可以在里面插入自己的 epilogue 逻辑比如做 GELU、残差连接、量化缩放等等。如果是正式项目里需要做生产级 GEMM 优化我的建议是先跑 cuBLAS 拿一个性能基线再用 CUTLASS 的 09b 等示例去复现这个基线最后才根据 profile 结果去改模板参数或者定制数据流。千万不要一上来就从零手写。4. 从零调好一个 BF16 GEMM 内核核心参数与优化顺序4.1 经典 GEMM 的分块复用逻辑前面提到 GEMM 的本质是数据搬运那怎么减少搬运呢答案是分块复用。假设我们算一个 4096x4096 乘 4096x4096 的矩阵乘如果不做任何分块每一个输出元素都要遍历 A 的一整行和 B 的一整列总共需要访问的数据量非常大这对显存带宽是毁灭性的。分块之后我们把输出矩阵切成一个个小 tile比如 128x128然后每个线程块负责计算一个 tile。为了算出这个 128x128 的 tile需要遍历 K 维每次加载 A 的一个 128xBK 的块和 B 的一个 BKx128 的块。如果 BK 取 32 或 64这些小块能很好地放入共享内存并且被同一个线程块内的多次计算反复使用。全局内存的访问量就从原来每个输出元素都要读一次降到了每个输出 tile 读若干次复用率大大提升。这个思路对任何架构都适用但 Blackwell 在硬件层面把数据复用的层级和效率都抬高了共享内存更大寄存器更多异步搬运更顺滑。所以你在 Blackwell 上可以放心用更大的 tile让数据在片上多转几圈。4.2 关键维度参数如何选择BLOCK_M / BLOCK_N / BLOCK_K调 GEMM 时最绕不开的参数就是这三个。BLOCK_M 和 BLOCK_N 决定了一个线程块一次算多大的输出区域BLOCK_K 决定了 K 维循环里一次加载多深的输入。它们的选择直接影响到寄存器压力、共享内存占用、warp 数量和流水线深度。一般来说BLOCK_M 和 BLOCK_N 越大数据复用率越高但寄存器压力和共享内存占用也越大太大会导致占用率下降甚至编译失败。BLOCK_K 越大全局内存访问次数越少但需要更大的共享内存来缓存数据同时也会影响流水线更新的速度。以 Blackwell 级别的硬件和 128x128 输出 tile 为例BLOCK_K 通常选 32 或 64 比较常见。比较稳妥的方法是用 Triton 或者 CUTLASS 做一个参数扫描脚本把常见的组合都跑一遍。别觉得这是笨办法实际调优里理论推算只能帮你确定一个大概范围最终还是要靠实测数据拍板。以下是一个我常用的 Triton 参数扫描思路固定 BLOCK_M 和 BLOCK_N从小到大调整 BLOCK_K观察算力利用率。固定 BLOCK_K把 BLOCK_M / BLOCK_N 从 64 调到 256。对每个组合记录 FLOPS、耗时、共享内存占用画成表格对比。实测中我遇到最多的现象是BLOCK_M 调到 256 时性能有提升但提升幅度往往不如 BLOCK_K 从 32 调到 64 明显。原因是 K 维的循环深度直接决定了流水线能够隐藏多少内存延迟。4.3 算力打不满时先查什么性能分析的优先级很多人调 GEMM 上来就纠结 API 参数但如果你连瓶颈是计算还是带宽都没搞清楚调参数就是瞎忙。我在实际项目里的排查顺序是这样的第一步先用 Nsight Compute 跑一次 profile看内核的关键指标。重点看 ComputeSMThroughput 和 Memory Throughput 谁先到 100%。如果计算吞吐很高但显存吞吐很低说明是计算瓶颈可以继续压缩指令开销如果显存吞吐已经拉满但计算还有空余说明这是一个带宽受限的 GEMM再调计算 tile 也没用得从数据复用和缓存机制下手。第二步看 occupancy占用率和 stall 原因。Nsight Compute 会告诉你 warp 卡在什么阶段是等待全局内存、共享内存还是等待计算依赖。这一步能帮你定位是不是流水线编排有问题。第三步检查尾效应。如果矩阵的 M、N 不是 tile 尺寸的整数倍边界上的线程块会非常浪费。典型做法是让 kernel 支持predicate off即边界外不计算或者把矩阵 pad 一下让尺寸对齐。5. 真实场景里的 BF16 GEMM训练、推理与目标检测中的不同形态5.1 训练场景GEMM 形状决定了反向传播的性能上限训练和推理最大的区别是训练要同时算前向和反向而且反向传播里会频繁出现矩阵转置形式的 GEMM。比如线性层权重 W 的梯度计算本质上是用 Loss 对输出的梯度去乘输入矩阵的转置。这种情况下矩阵的布局RowMajor 还是 ColumnMajor对 GEMM 性能影响非常大。如果你用的是 PyTorch并且想确认自己的 BF16 GEMM 有没有走对路径可以留意一下这几个配置torch.backends.cuda.matmul.allow_tf32 False # 训练时大多数情况下应该关闭 TF32使用 BF16 torch.backends.cudnn.allow_tf32 True在 BF16 训练时PyTorch 会把 Linear 和 Conv 底层映射到 cuBLAS 的 BF16 GEMM。如果你的显存带宽足够瓶颈一般还是在 GEMM 计算本身。训练场景里 M 通常等于 batch size 乘序列长度很大N 和 K 对应模型宽度相对均衡。这种大 M、大 N、大 K的形态最适合用大 tile 去跑性能也最容易打满。5.2 推理场景小批量、KV Cache 与低时延下的取舍推理场景和高性能训练的场景差别很大。服务线上请求时batch size 经常很小比如 1 到 16。此时 GEMM 的形状变成瘦高型M 很小N 和 K 很大。这种形状下数据复用率天然低因为每个 thread block 分到的输出 tile 很少。这时为了优化推理业界经常会把多个请求拼在一起增大 batch或者用 continuous batching 技术动态拼 batch。在算子层面就需要针对小 M 的 GEMM 做特殊处理不用裁一个大 tile而是用更小的 BLOCK_M比如 16 或 32同时尽量利用 Tensor Core 的指令形状去凑效率。另外推理时如果做了 KV CacheAttention 部分会频繁出现批量 GEMM也就是 GEMV矩阵乘向量或者更小规模的 GEMM这些操作的瓶颈往往在显存带宽而不是算力。这种情况你用再多的 Tensor Core 也没用得靠优化数据布局、减少 KV Cache 读取量来解决。5.3 YOLOv8 这类检测任务少步推理、多路并发时的 GEMM 表现YOLOv8 这类目标检测模型虽然是以卷积为主但卷积在 GPU 上通常会被转化为 GEMM 来实现尤其是 1x1 卷积本质上就是一个矩阵乘。实际部署时大家常用 TensorRT 做优化把多个层的 GEMM 融合在一起减少内核启动开销和显存读写。以图像分辨率 768p、8 步推理这样的场景为例整个推理过程其实包含大量小 GEMM 和 GEMV。如果你用 BF16 精度要特别关注前处理里图像归一化会不会把数值范围压缩到 BF16 表达不好的区域检测头部分的卷积虽然有 GEMM 性质但通道数往往不大容易变成带宽受限。我在部署这类模型时发现单靠换 BF16 不一定能带来预期的加速更关键的是把算子融合做好、降低每层的数据搬运量。6. 常见问题与排查技巧实录6.1 算力利用率只有 40%初查清单这是我最经常被问到的问题没有之一。写好的 BF16 GEMM 内核跑起来Nsight Compute 一看计算吞吐只有 40%感觉自己白努力了。出现这种情况先按下面的清单过一遍矩阵尺寸是不是太小小 GEMM 本身就不适合打满算力这时不要盯着利用率看看端到端时延。有没有开异步流水线如果代码里每个循环都同步等待流水线完全没重叠利用率肯定上不去。共享内存是不是有 bank conflict这会让本来并行的内存访问退化成串行吞吐瞬间掉一截。是不是被边界判断拖慢了大量的 if-else 会导致部分 warp 空转。寄存器溢出有没有如果局部变量太多寄存器不够用会 spill 到本地内存那个代价非常大。6.2 共享内存 bank conflict 与尾效应共享内存是按 bank 组织的比如常见的是 32 个 bank每个 bank 在一个周期内只能服务一次访问。如果多个线程同时访问同一个 bank 但地址不同就会发生冲突硬件只能串行处理。在 GEMM 里共享内存布局如果没对齐很容易踩到冲突。排查方法是看 Nsight Compute 的 Shared Memory 报告里面会明确标出 bank conflict 的次数。修复手段一般是调整共享内存数组的 padding即在每行末尾多分配几个元素让不同线程的访问落到不同 bank 上。别小看这一步在一些场景里一个 padding 就能让性能提升 20% 以上。尾效应则是指矩阵尺寸不能被 tile 整除时多余部分计算浪费或者被跳过。处理办法是让 kernel 在边界处做掩码判断或者在上层把矩阵 padding 到对齐尺寸。实测时我倾向于 padding因为判断分支的开销在 GEMM 这种计算密集循环里非常不值得。6.3 BF16 精度是不是一定够什么时候要退回 FP32 累加BF16 的尾数只有 7 位在训练中通常配合 FP32 累加效果稳定。但一些对精度特别敏感的场景比如梯度很小、模型非常大、或者训练后期学习率很小的时候BF16 的组织误差可能会被放大。这时候你至少应该确认你用的是 FP32 累加而不是 BF16 累加。在 PyTorch 里torch.matmul默认使用 FP32 累加这点不用太担心。但如果你自己做自定义 kernel非常容易把累加器也声明成 BF16那就出大问题了。累加器必须是 FP32这是 BF16 GEMM 实现的底线。如果是推理部署量化到 BF16 之后再跑一层校准、看一下输出分布和 FP32 baseline 的差异是判断精度是否可接受的最快方法。6.4 工具速查nsys、ncu、nvidia-smi 怎么配着用调试 GEMM 时我日常三件套是 nvidia-smi、Nsight Systems 和 Nsight Compute分工很明确nvidia-smi dmon -i 0 -s mpucv -d 1这条命令可以实时监控 GPU 的利用率、显存、温度、功耗。跑长训练任务时我会一直开着它确认 GPU 有没有被喂饱。如果利用率频繁跌到 0说明 CPU 侧数据加载成了瓶颈。nsys profile --tracecuda,osrt -o gemm_profile python test_gemm.pyNsight Systems 用于看整个程序的执行时间线能帮你发现显式的 GPU 空闲和内核启动间隙。如果两个 GEMM 内核之间有明显的大段空白说明上层调度有问题或者数据还没有准备好。ncu --set full --kernel-name regex:gemm python test_gemm.pyNsight Compute 是内核级别剖析的利器看的就是单个 kernel 内部的计算与访存行为。所有关于 utilization、occupancy、stall、bank conflict 的数据都从这里来。调优时必须用它光靠 nvidia-smi 的利用率数字根本不够。这三者的关系可以理解成先拿 nvidia-smi 看面再用 Nsight Systems 看线最后拿 Nsight Compute 看点从宏观到微观一层层排除。7. 一点实操心得我踩过很多次坑之后发现BF16 GEMM 在 Blackwell 上的优化路径并不是越复杂越好。很多场景下直接用 cuBLAS 就能拿到非常强的性能。真正需要手写内核的往往是那些形状特殊、需要融合自定义算子、或者你有非常明确的极致性能目标的时候。这时候再考虑 CUTLASS 或者 Triton 也不迟。还有一个心得是调 GEMM 一定要做对照实验不要只看一次跑完的时间。我习惯在每个参数组合上都跑至少三遍取中位数或者最小值并且用相同的数据和 warmup 来保证公平。否则数据本身的波动可能会导致你做出错误的判断。另外Blackwell 的硬件能力虽然强但新架构的驱动和工具链迭代还需要时间。我建议在正式做大规模优化之前先确认你的 CUDA Toolkit、驱动、PyTorch 版本是否已经支持目标型号。很多时候排查了半天性能问题最后发现是工具链没跟上那就非常尴尬了。这个项目本身还在持续更新后续我会继续补充 CUTLASS 底层指令排布、更复杂的 tile 调度、以及在真实大模型训练和推理中的实测案例。也欢迎你在底下留言交流你遇到的具体 GEMM 性能问题我会挑有代表性的案例放到后面的更新里。