企业官网建设流程全解析

做算子优化这些年，我一直有个体会：真正卡住团队进度的，往往不是算法本身，而是底层那套矩阵计算的基础设施不够趁手。上个月调一个融合了残差和 LayerNorm 的线性层，性能怎么都追不平 cuBLAS 的高峰，最后翻出团队早期用 CATLASS 这套矩阵计算模板库做的老底子重新梳理，才意识到问题根本不是算子写得不对，而是数据流从设计之初就没有按模板提供的分块思路来。这篇文章就借这次复盘，把 CATLASS 模板库的核心设计、GEMM 数据流拆解，以及混合精度优化里那些文档不会明说的细节，一次性讲透。

CATLASS 是一套面向 CUDA 平台的高性能矩阵计算模板库，核心思路是把 GEMM（通用矩阵乘法）这类算子彻底模板化——从分块大小、数据布局、指令派发到流水线策略全部变成编译期参数。它解决的痛点是：直接写 CUDA kernel 虽然灵活，但每次新业务场景都要从头优化一遍，重复造轮子成本极高；而直接用 cuBLAS 又很难做算子融合和定制。CATLASS 正好卡在中间，既能通过拆解好的 Tile 迭代器、Warp 级指令封装拿到接近手工调优的性能，又能让上层业务以组合模板的方式快速搭建新算子。适合谁看？正在做深度学习推理引擎、算子库研发，或者想搞懂高性能 GEMM 内部到底发生了什么的人。

1. CATLASS 模板库定位与整体设计思路

1.1 模板库到底解决的是什么问题

做高性能计算的人都会遇到一个尴尬：写一个能跑的 GEMM kernel 并不难，难的是让它跑得跟 cuBLAS 一样快。cuBLAS 之所以快，是因为它对每一种 shape、每一种数据类型、每一种架构都做了针对性调优，而这份调优经验被固化成了庞大的代码库。CATLASS 这类模板库的思路，就是把这套经验中"与具体硬件相关"的部分抽象成模板参数，让使用者通过组合模板来生成高性能 kernel。

模板参数化的核心收益是把"计算逻辑"和"数据布局"解耦。比如在 CATLASS 里，一个矩阵乘法的模板参数会包含：数据类型、矩阵布局（行主序/列主序）、Tile 大小、每个线程负责的元素数、是否启用双缓冲、是否使用 WMMA 张量核心指令等。这些参数全部在编译期确定，编译器可以做完整的常量展开和循环展开，最终生成的代码接近手工极致调优的版本，同时又不失灵活性。

举个例子，要给一份 FP16 的矩阵乘法适配 Ampere 架构张量核心，你只需要改一行模板参数，把指令策略从 SIMT 换成 WMMA；要给一个 128x128 的 Tile 改成 256x128，也只是一个常量。这种"改参数而不是改代码"的体验，对算子库团队来说极其重要——你从零写一个高性能 kernel 可能要两周，调模板一天就能搞定。

1.2 模板粒度：粗了不够灵活，细了编译爆炸

CATLASS 设计上最值得琢磨的一点，是模板粒度的取舍。模板参数如果过细，每一个循环层次都暴露给用户，确实灵活，但代价是模板瞬间变得极其冗长，编译时间暴涨，报错信息人类几乎无法阅读；如果过粗，性能和定制能力又受限。CATLASS 的折中方案是采用"三层映射"结构：线程块级负责 Tile 调度和 Shared Memory 切分，Warp 级负责子 Tile 的排布，线程级负责逐元素乘加和向量化加载。

这种粒度划分背后是有道理的。以 CUDA 编程模型来看，对性能影响最大的三个决策分别是：数据如何从全局内存搬入共享内存（涉及访问模式、向量化）、数据如何从共享内存搬到寄存器（涉及 bank conflict、寄存器重用）、计算用哪种指令（SIMT 指令、FMA 融合乘加，还是 Tensor Core 的矩阵指令）。CATLASS 把这三点拆成了不同层级的模板模块，互相之间可以独立替换，这是我用过之后觉得它比早期很多"一个大类搞定一切"的实现要高明的地方。

粒度选择还有一个实践层面的考量——可读性。团队做模板库不是写完一锤子买卖，后续要长期演进和维护。CATLASS 的模板分层让新成员能按"线程块级→Warp 级→线程级"的顺序逐步深入，而不是一上来就要理解所有宏展开。这个设计对整个开源社区的高性能计算项目都有参考价值。

1.3 核心抽象：Tile、Iterator 与布局

CATLASS 里我最常用到的是三个核心抽象：Tile、Tile Iterator 和 Layout。Tile 是一块放进共享内存或寄存器中的子矩阵；Tile Iterator 负责按预定模式从全局内存/共享内存搬运数据；Layout 则描述矩阵元素在内存中的排列方式，比如行主序下第 m 行第 n 列的元素在 offset 为 m * ldm + n 的位置。

这三者的关系可以这么理解：Tile 是"容器"，Iterator 是"搬运机器人"，Layout 是"仓库货架分布图"。GEMM 的执行流程，就是一个一个 Tile 被 Iterator 按 Layout 指引搬到共享内存和寄存器，完成乘累加，再把结果 Tile 写回全局内存。由于这三者都是模板参数，你可以给同一个计算核心配上完全不同的 Layout 和 Iterator。

比如处理非连续内存的稀疏场景，你可以写一个自定义 Iterator，而主计算循环一行都不用动。这种可替换性一度帮我省过很多事：当时要给一个按列分块的业务矩阵写 GEMM，naive 做法是把列分块重排成连续内存（额外一次数据拷贝），但我直接写了个列访问 Iterator，计算循环原封不动，性能提升明显且省了一次全局内存读写。

2. GEMM 数据流拆解：从全局内存到寄存器的三层搬移

2.1 分块计算（Tiling）是 GEMM 性能的基石

GEMM 的计算公式是 C = A * B + C，核心操作是乘累加。一个 1024x1024 的矩阵乘，如果不做任何分块，A 的每一行要跟 B 的每一列做内积，意味着 A 的元素会被反复从全局内存读取 1024 次（对每个输出列都要取一次）。全局内存带宽是稀缺资源，这种重复读取必然把性能压到内存带宽的天花板之下。

分块计算的思路是：把输出矩阵 C 切成若干个小的 Tile（比如 128x128），每个线程块只负责这一个 Tile 的计算。计算这个 Tile 只需要 A 的一个水平条带（128 行）和 B 的一个垂直条带（128 列），这两块数据在计算过程中会反复被汇编代码使用，因此可以提前搬入共享内存或寄存器中做重复利用。这样一来，全局内存的访问次数从"每个输出元素读一次源数据"降为"每个 Tile 读一次源数据"，降幅大约为 Tile 尺寸的量级。

CATLASS 在 Tiling 上的一个重要设计是 Tile 尺寸的选择。128x128 是一个经验上很均衡的尺寸：共享内存占用合理（FP16 下两个矩阵各 32KB，加上主矩阵和额外开销正好压住 48KB~100KB 的范围），线程块内并行度高（可容纳 256 个线程），寄存器利用率高。我用过超大 Tile（256x256），性能并不一定更好，因为单个线程块要算的数据多了，双缓冲流水线深度会被共享内存容量压住，反而导致流水线气泡。

2.2 三层搬移：全局内存 → 共享内存 → 寄存器 → 计算

CATLASS 的 GEMM 内层循环，本质上是一条三级流水线。第一级把 A 和 B 的 Tile 从全局内存用向量化加载（如 float4，即连续读 16 字节）搬入共享内存；第二级把共享内存中的数据搬到线程各自的寄存器中，并按乘累加的因子分配好每个线程需要计算的子区域；第三级就是寄存器里的 FMAs（融合乘加）。

这里面每一级的代价差距非常大。全局内存访问延迟一般几百个周期，共享内存延迟约为几十个周期，寄存器则零延迟。所以优化的核心思想是：尽量把同一块数据在寄存器中反复使用，避免频繁访问共享内存，更不要频繁访问全局内存。

CATLASS 把这一层逻辑做得非常细。在寄存器层级，每个线程通常负责一个 8x8 或者 16x8 的微 Tile。以 FP16 利用 WMMA 张量核心为例，每个线程通过一条mma.sync指令可以完成 16x16x16 的矩阵乘累加，即一次指令做 4096 次乘累加。如果走普通 SIMT 的 FMA，一次指令只能做一次乘累加。这也是混合精度性能和 FP32 拉开差距的本质原因之一——除了数值格式省了带宽，指令本身的吞吐也完全不同。

2.3 共享内存 Bank Conflict 规避与数据布局

共享内存位于芯片内部，吞吐远高于全局内存，但它有 32 个 bank，每个 bank 在单周期内只能返回一个 4 字节数据。如果同一周期内多个线程访问的地址映射到了同一个 bank（称为 bank conflict），这些访问会被串行化，吞吐直接除以冲突次数。一个常见的 2-way bank conflict 就会让共享内存读取性能掉一半，这个损耗在 GEMM 这种高频循环里会被放大成肉眼可见的性能滑坡。

CATLASS 处理 bank conflict 的经典手法是 padding，也就是在共享内存的 Tile 每一行末尾多塞几个元素，把每行的实际 stride 从 64 字节调成 68 字节或类似值。这样原本按对角线访问模式的线程，地址被错开，避免多个线程同时命中同一 bank。我当初第一次跑自己的 GEMM kernel 时，性能比 cuBLAS 低约三成，用 Nsight Compute 看到 LDS（加载共享内存）指令的 bank conflict 统计高达 12%，就是因为没有做 padding。

换一个角度说，bank conflict 的排查往往也是高性能算子优化中收益最直接的几个手段之一。你们如果发现自己写的 GEMM 或者 Implicit GEMM 卷积算子，性能怎么调都上不去，第一步不是怀疑指令选型，而是打开 profiler 看内存访问那一档是不是有红黄指标。

2.4 双缓冲与流水线：用异步把延迟藏起来

数据搬移和计算天然存在依赖：要等数据到齐才能算，算完才能搬下一批。如果串行执行，"等待数据"的时间就得靠发虚。CATLASS 的解法是流水线化，典型实现是双缓冲：共享内存里开两份缓冲区，一份用于当前计算，一份用于预取下一批数据。计算当前 Tile 的同时，通过cp.async指令异步把下一个 Tile 的数据从全局内存搬入另一份缓冲区，两件事并行发生，等待延迟被完全隐藏。

Ampere 架构之后，cp.async指令可以直接让数据绕开寄存器从全局内存异步拷贝到共享内存，拷完自动触发完成机制，不需要显式再走一次 LDS。用这套机制时，流水线 stage 数的选择非常关键。stage=2 是双缓冲，stage=3 是三级流水。stage 越多，共享内存占用越大，但抵抗延迟波动的能力越强。CATLASS 在部分 kernel 中默认 stage=4，因为它按最坏情况延迟来设计，而对很多业务场景来说 stage=2 就够。

实际使用中我建议从 stage=2 开始往上加，观察实际吞吐变化。有的场景共享内存容量已经被 Tile 占得差不多，硬塞 stage=4 会导致 occupancy（占用率）下降，反而不如 stage=2 来的稳。混合精度下感受会更明显——低精度 Tile 本身占的内存少，流水线深度提高带来的收益比 FP32 大得多。

3. 混合精度优化的核心实践

3.1 混合精度为什么能快：算力与带宽双线并进

混合精度（Mixed Precision）在训练和推理中的收益，来自两个层面。第一是算力层面：现代 GPU 的 Tensor Core 对 FP16/BF16 的矩阵乘累加吞吐，通常可以达到 FP32 的数倍以上。以 A100 为例，FP32 FMA 的标称算力约 19.5 TFLOPS，而 Tensor Core FP16 可以到 312 TFLOPS，差距接近 16 倍。第二是带宽层面：FP16 每个元素只占 2 字节，是 FP32 的一半，全局内存和共享内存的单位时间搬运元素数量在同样带宽下直接翻倍。对于矩阵乘这类访存和计算并重的算子，两个层面的收益可叠加。

值得注意的是，"算力差 16 倍"并不等于"端到端快 16 倍"。GEMM 如果是访存受限型（矩阵比较小、无法充分复用数据），主要的瓶颈在搬运，Tensor Core 高算力完全使不上劲。比如一批 64x64 的小 GEMM，性能差异可能只有 1.5 倍。混合精度优化的正确思路，是先判断算子是 compute-bound 还是 memory-bound，再决定是否值得上 Tensor Core 路径。我在实践中遇到不少团队直接把所有算子改成 FP16，结果收益很有限，反而还引入了精度风险，就是忽略了这一步判据。

3.2 FP16 / BF16 / TF32 三种格式的选型对比

混合精度里最常碰到的三种低精度格式，各有各的使用场景。

FP16（半精度）：1 位符号、5 位指数、10 位尾数。它能表示的最大值约 65504，超过就上溢到无穷；最小正常值约 6e-5，更小的数会逐步损失精度，一直到约 6e-8 的次正规数边界。FP16 的优势是 Tensor Core 支持最成熟、前后端软件栈适配最好；劣势是指数范围窄，训练时如果梯度或中间激活值超过 65504 或者跌到 1e-4 以下，都会出问题。

BF16（Brain Floating Point）：1 位符号、8 位指数、7 位尾数。指数范围和 FP32 完全一致，最大可表示约 3.4e38，最小正常值约 1e-38，基本不会出现上溢或下溢的问题。代价是尾数只有 7 位，相对精度低。BF16 在训练场景下因为不容易炸，又配合损失缩放，已经成为主流训练格式；但在推理场景，如果模型权重数值本身很小（比如某些蒸馏后的模型），BF16 的尾数不足容易导致精度明显下降。

TF32：本质是在 Tensor Core 计算中把 FP32 数据截断为 10 位尾数参与矩阵乘。它不用改内存里的数据格式，输入输出仍按 FP32 存储，因此兼容性最好，适合那些想换取 Tensor Core 加速又不愿改动数据通路的 FP32 模型。TF32 的精度约等于"劣化 FP32"，只保留大约 10 位十进制有效数字，对大多数深度学习模型在可接受范围内。

我的选型经验：训练首选 BF16 + 动态损失缩放，推理首选 FP16（如果模型精度敏感，再做逐层精度检测来决定哪些层退回 FP32）。TF32 适合基础设施层不想大动干戈、只是想白拿一部分算力的场景，但因为内存占用没有减半，带宽收益全无。

3.3 损失缩放（Loss Scaling）：混合精度不能绕过的一环

FP16 训练中一个经典问题是梯度下溢。深度学习训练时，梯度的数值往往远小于权重，早期层梯度降到 1e-6 甚至更小很常见，而 FP16 的最小正常值只有 6e-5，更小的梯度直接变成 0，参数就再也不更新了。解决办法是损失缩放：在反向传播前，把损失值乘以一个较大的系数（通常 1024 或 2048），梯度也会同比例放大，FP16 能容纳的数值范围自动覆盖原本小到会消失的梯度；反向传播完成后，更新参数前再把梯度除以同样的系数。

混合精度库一般在内部自动处理了损失缩放，但用 CATLASS 这类底层模板库做自定义算子时，这个环节要自己注意。我踩过一次坑：用 FP16 写一个自定义 GEMM 核反向传播，算出来的梯度在小数值区域明显偏小，最后定位到是没有对中间梯度做任何 scale，梯度在矩阵乘内部就截断成了 0。这里建议梯度的缩放因子遵循"先放大、算完、再缩放"的顺序，不要在乘累加内部反复乘除，因为每次缩放都会引入一次舍入误差。

损失缩放还有一个细节是动态缩放。固定缩放因子 1024 在部分模型里不够用（太大会让权重上溢，太小挡不住梯度下溢），更好的方案是动态检测：连续一段时间没有出现 inf/NaN，就适当调大缩放因子；一旦检测到 inf/NaN，立即调小并跳过本轮更新。PyTorch 的 GradScaler 就是这么做的，底层模板库使用者可以仿照这个逻辑自己实现，成本不高但效果明显。

3.4 混合精度在矩阵乘法中的落点：哪些算子能改

不是所有算子都适合切混合精度。性能层面，切低精度的收益来自算力翻倍和带宽减半，所以 compute-bound 的大矩阵乘收益最大，memory-bound 的 activation 类算子收益主要靠带宽。精度层面，几何型算子（做大量连续乘法和指数运算的）对精度变化敏感，规约型算子（求和、求均值）相对不敏感。

以 Transformer 为例，Attention 里的 QK 点积矩阵乘、以及 MLP 块里的两个大线性层，是混合精度的最大受益者。它们都是大矩阵乘，内存访问次数相对计算量占比低，Tensor Core 高吞吐正好发挥。而 Softmax、LayerNorm 这类逐元素算子，数值范围对精度影响大且本身不是计算密集型的，通常保持 FP32，或只在数据通路用 FP16 而将内部规约动态提升为 FP32。CATLASS 模板库因为把计算核心和规约算子分开设计，做这种"矩阵乘低精度、规约高精度"的组合非常方便。

实践上还有个经验：混合精度代码里尽量保持"计算时低精度、累加时高精度"的模式。加法树用 FP32 累加器，乘法和乘累加用 FP16，这样既吃到了 Tensor Core 算力，又把累加误差控制在 FP32 水平。这也是 CATLASS 混精度 kernel 里一个被写死的模板参数——如果你想改成全 FP16 累加，不是不能改，但绝大多数情况都不会更好。

4. 混合精度踩坑实录与排查方法

4.1 精度问题不一定是舍入，先查指数范围溢出

很多团队在混合精度上遇到的第一个问题不是性能不达标，而是模型直接产出 NaN。常规思路会怀疑舍入误差、累加顺序、FMA 融合，但排查了一圈下来，FP16 的上溢才是最大嫌疑。FP16 最大只到 65504，如果权重矩阵的初始值带几个偏差大的异常点，或者激活值过了一个带大 scale 的层，几乎必炸。

排查方法是逐层打印中间激活值和权重的绝对最大值、绝对值最小值、绝对值均值，对比 FP32 基线的同一层数值。如果 FP16 下某一层的 max 超过 65504，或者 min 掉到 1e-4 以下，基本可以锁定是范围问题而不是精度问题。解决手段有三条：一是把该层数据先做 affine 归一化再运算，二是在矩阵乘之前手动对输入做缩放，三是干脆将该层退回 FP32。这三条路我都在实际项目里用过，第一条路的性能损失最小，因为归一化通常可以融合进前面的算子。

4.2 性能问题：bank conflict、occupancy 和指令混合

混合精度 kernel 性能不如预期，通常不是 Tensor Core 没发力，而是旁边的数据通路拖了后腿。我排查得最多的三个方向是：共享内存 bank conflict、寄存器溢出（spill）和线程块占用率（occupancy）过低。这三个问题有一个共同特点——不会导致结果错误，所以更容易被忽略。

排查思路是打开 Nsight Compute，先看"Memory Workload Analysis"里的 LDS Bank Conflict 指标，超过 5% 就值得优化。再看 "Occupancy" 那一栏，如果实际值低于理论值，大概率是共享内存或寄存器资源超了。最后看一眼 warp state 里的待命周期（Stall），如果大量周期是等待共享内存或全局内存数据回来，说明流水线深度不足，需要增加 stage。

4.3 一个典型案例：错误配置导致性能反而下降

之前帮一个团队看一个 Transformer 推理引擎，他们把 QKV 的 GEMM 全部切成 FP16，FP32 基线大约 220 TFLOPS 的利用率，切完反而只有 120 TFLOPS，甚至不如基线。逐个参数排查后发现了两个问题叠加：第一，Tile 大小从 128x128 改成了 64x64，因为共享内存里给双缓冲留的缓冲区被 FP16 数据塞得更大了，看起来"应该刚好能塞下"，但忘记了 CATLASS 还需要给共享内存预留一行 padding 和 stage 缓冲区，结果 occupancy 掉了一半。第二，warp 数量从 8 个减到 4 个后，寄存器循环展开量不够，全局内存加载频率变高，流水线没法藏住延迟。

修正方案很简单：恢复 128x128 Tile，把双缓冲 stage 从 4 降到 2，给 padding 留出空间。改完性能直接回到接近 FP32 基线的 1.8 倍。这给团队的教训是：混合精度不只是换数据类型，它改变了一整条数据通路的资源占用面貌，所有依赖资源的设计（Tile、stage、寄存器数）都要重新平衡。

4.4 性能排查工具箱：关键工具与基本使用流程

我调 CATLASS kernel 时，一般会按下面这个流程走：

用ncu --set full跑一次 Nsight Compute，重点看 SOL（Speed of Light）里的 Compute 和 Memory Throughput。如果两者都没超过 60%，先假设 kernel 被延迟限制了，看 warp stall 原因再往前推。如果 Compute 接近 90%、Memory 接近 90%，那就是优化得很好的状态。接下来看 instructions 里的 LDS/UDS 数量、bank conflict 计数、global load/store 效率。最后用ncu --metrics launch__occupancy_limit_shared_mem,launch__occupancy_limit_registers确认资源瓶颈。

数值精度排查方面，写一个小工具脚本：FP32 基线算一遍各层输出，混合精度算一遍，统计逐元素的绝对误差、相对误差、最大误差位置。通常相对误差在 1e-2 以内就可以接受，如果超过 0.1，需要定位是哪一层开始放大的。这个方法虽然土，但在追踪"精度在哪个算子被毁掉"时非常有效。

5. 从 GEMM 到更多算子：模板库能力的延伸应用

5.1 卷积的隐式 GEMM（Implicit GEMM）化

CATLASS 的 GEMM 核心不止能算常规矩阵乘，也是卷积算子的基础。卷积本质上是一个四维循环，可以转换成矩阵乘形式：把输入特征图按卷积窗口展开成矩阵（im2col），把卷积核展开成另一个矩阵，两者相乘再重组。naive 的 im2col 需要额外拷贝数据，内存开销大得吓人（展开后数据量通常膨胀几十倍），而隐式 GEMM 的思路是不真正展开数据，而是在 GEMM 的 Tile Iterator 里直接根据卷积映射关系去取数据。

这个思路在 CATLASS 里落地得很自然。Tile Iterator 本来就是"按布局搬运数据"的抽象，你给它定义好"第 m 行对应输入特征图的哪个位置、第 n 列对应卷积核的哪个位置"，它会自动在 GEMM 主循环里完成取数。这样做的好处是主循环计算逻辑完全复用 GEMM 的优化积累——Tiling、双缓冲、Tensor Core 指令全都是现成的。

5.2 注意力机制算子的融合实践

注意力机制里的 QK^T 和 Score·V 本质上也是矩阵乘，定制的空间在于中间要穿插 Softmax，并且要避免把整个 Score 矩阵写回全局内存。若完全按通用 GEMM 做，QK^T 算完把 NxN 的 Score 矩阵写回全局内存，再从全局内存读回来做 Softmax、再参与 PV 矩阵乘，访存开销非常大。CATLASS 的做法是把两段矩阵乘和中间的 Softmax 融合成一个大的 kernel，中间的 Score 留在寄存器或共享内存，只借道 CPU 或小块内存做规约。

这个方向最有名的实现是 FlashAttention，它的核心贡献之一就是用分块思想避免大 Score 矩阵落回全局内存。CATLASS 用户完全可以照着这个思路实现一个定制化的 Attention kernel，Tile 迭代器照用 GEMM 的，中间替换成融合了 Softmax 的 epilogue（后处理）。这种"保留 GEMM 主循环、替换 epilogue"的自定义模式，是模板库最有价值的部分。

5.3 自定义算子的复用模式与边界

CATLASS 这类模板库给了使用者很大的自由，但边界也要心里有数。适合复用的场景是：算子主体是矩阵乘或点积类计算，差异点在于取数方式和后处理逻辑。不适合复用的场景是：算子数据依赖关系非常非线性，比如复杂的稀疏路径、动态控制流极多的场景。这类算子用模板库强行套模板，最后只能越套越复杂，不如直接写 NVCC kernel。

我给团队定的经验准则是：如果算子中有超过 80% 的工作本质是"计算一个输出 Tile = 若干输入 Tile 的乘累加"，就用 CATLASS 改；否则直接手写。手写时可以借用 CATLASS 里的分块、双缓冲思路，但不需要强行上模板。很多团队项目失败，是因为过早地想把所有算子都用模板库统一了，结果收益没看到，只看到编译时间和排查成本的上升。

说到调试，CATLASS 模式下的模板代码报错信息向来不友好。一个通用技巧是按"C++ 模板类实例化展开"的方式去读报错：先把缺失的类型、常量推断出来，再反查是哪个模板参数类型不匹配。并且确保编译时打开-G调试模式和-lineinfo选项，这样反汇编和性能分析的定位会准确得多。

关于这套模板库的实际使用，我个人的体会是：它确实不是最快让人出活的那条路，初期学习成本比直接调 cuBLAS 高不少，但一旦把 Tile 迭代器、流水线、双缓冲、混合精度这些概念吃透，后续写任何高性能算子都会上一个台阶。最后再分享一个调试小技巧：当 kernel 性能不符合预期时，先别急着搜代码，把 Nsight Compute 的 SOL 页面截图，对照 cuBLAS 同规格算子的 SOL 图找差异区间，通常一眼就能看出是计算瓶颈还是访存瓶颈，省掉大量盲目尝试的时间。这套方法论，比背任何库的 API 都值钱。