企业官网建设流程全解析

做模型部署这几年，我越来越觉得一件事：推理引擎里真正值钱的算子，翻来覆去就那几个。GEMM、Softmax、各种归一化和激活，再加 Attention 里的组合算子，几乎把 profiling 报告占满了。这也是我把这套“CUDA 算子优化实战”的第一篇放在 Softmax 和 GEMM 上的原因——一个访存密集型的典型，一个算力密集型的代表，两者几乎覆盖了手写算子遇到的所有底层问题。这篇文章会直接给出能编译运行的 CUDA 代码，讲清楚每一步为什么这么做，再把我踩过的坑和排查思路整理出来。适合已经写过简单 CUDA kernel、想进一步理解优化套路的人。

1. 为什么第一个实战就选 Softmax 和 GEMM

1.1 它们在推理链路里到底有多高频

随便打开一个 Transformer 模型的解码过程看：每个 token 的生成都要过一遍 GEMM 相关的投影计算，同时 Attention 内部要对 KV 缓存做一次 Softmax。推理一张图、跑一段语音识别，也躲不开分类头的 Softmax 和骨干网络里大量 GEMM。也就是说，这两个算子的性能直接决定了最终端到端的延迟和吞吐。

很多同学一开始喜欢去卷那些看起来很酷的融合算子，比如 Flash Attention。但 Flash Attention 本质上就是把 GEMM、Softmax、掩码、rescale 这些基础操作揉在一起。你如果连单独的 Softmax 和 GEMM 都写不好，贸然上手融合算子，只会被一堆边界条件和同步问题绕晕。先把这两个基础算子调明白，后面再碰 Attention 优化就会顺很多。

还有一个现实原因：这两个算子在 profiling 里特别好分析。Softmax 的瓶颈基本在显存带宽，GEMM 的瓶颈基本在计算管线利用率。把这两个极端情况吃透，中间态的那些算子（比如各种 elementwise 融合）就没什么能难住你的了。

1.2 分清算力瓶颈和带宽瓶颈

优化 GPU kernel 之前，第一件事是判断它属于哪一类问题。我习惯用一个很粗的指标：计算量除以访存量。每一轮要做的浮点运算多、从显存搬的数据少，就是算力密集；反过来，搬数据的时间远大于算数的时间，就是访存密集。

Softmax 对每个元素基本就是一次指数、一次除法、几次加减法，但每个元素都要从显存读进来一次、写出去一次。指数运算虽然也不便宜，但在现代 GPU 上，跑 Softmax 时通常还是显存带宽先打满。GEMM 则完全相反，理论上它的计算访存比可以做到很高，问题是你能不能把数据复用起来。如果你没做好分块，A 矩阵的元素会被反复从低价缓存甚至显存里捞，计算单元反而在空转。

判断错了类型，优化方向就会完全跑偏。我见过有人对着一个明显的带宽瓶颈算子疯狂调计算指令顺序，结果毫无变化；也见过有人在 GEMM 里花大力气压共享内存流量，核心的寄存器累加却没做。所以动刀之前先花十分钟算清楚：这个形状下，到底谁在拖后腿。

1.3 两个算子背后是同一套方法论

虽然一个偏访存、一个偏算力，但优化思路高度一致：第一，尽量减少不必要的数据搬运；第二，让在片上的数据尽可能被多次使用；第三，把访存指令向量化，把计算指令排满。Softmax 的“在线计算”和 GEMM 的“分块累加”，说白了都是围绕这三条在做文章。

另外还有一个容易被忽略的点：验证手段要提前准备好。我每次改 kernel 都会先写一个 CPU 上的朴素实现做对照，再固定一组随机输入和一个可比的性能测量方式。没有这个底子，后面所有优化都是空中楼阁。写这篇之前的准备也一样，所有代码都在同一台机器、同一组数据下对比，避免拿别人的数字误导自己。

2. Softmax 优化：目标是把三次显存读压成一次

2.1 教科书实现为什么会把带宽白白烧掉

最朴素的 Softmax 写法是分三步：第一步扫描整行找最大值，第二步再扫描一次算每个元素减最大值后的指数和，第三步又扫描一次做归一化写出。也就是每个元素被从显存读了三次、写了一次。在带宽受限的算子上，这基本上是灾难。

即使你把三步并进一个 kernel，用 block 内部归约省掉中间 kernel 的启动开销，只要数据还留在全局内存里，每一轮扫描就仍然是一次完整的显存读取。显存带宽是死的，读三遍就是三倍的耗时。所以优化 Softmax 的核心目标非常明确：把对全局内存的遍历压到一次，读完的数据尽量留在寄存器或者共享内存里。

这里有个很多人问的问题：为什么不能直接把原始值存到寄存器，等归约出最大值和总和之后再做指数？当然可以，这正是寄存器驻留方案的做法。前提是每一行分给单个线程的元素数量要少到寄存器装得下。行长度比较短（比如 1024 个 float）时完全没问题，行很长（比如几千上万）时就得换共享内存方案。

2.2 在线 Softmax 的公式推导与防溢出设计

在线 Softmax 的核心思想是：不等到扫描完整行，就一边扫描一边维护“当前见过的最大值”和“当前见过的指数和”。假设已经处理了前 t 个数，手里有当前最大值 m_t 和当前和 s_t。现在来了一个新块，它的局部最大值是 m_new，组合后的全局最大值是：

m' = max(m_t, m_new)

原来的和 s_t 是基于 m_t 算的，现在基准变成了 m'，所以要把旧和做一个缩放：

s' = s_t * exp(m_t - m') + Σ exp(x - m')

对每个新元素都这样更新下去，等整行扫描完，s 就是真正的归一化分母。这样做的好处有两个：一是每个元素只从显存读一次，配合寄存器驻留就直接省掉后面两轮扫描；二是全程基准都跟当前最大值走，不会出现 exp 的参数太大导致溢出。

如果你选择把原始值存下来、等归约出最终最大值后再统一算指数，那实际上不需要这个缩放公式，直接在最后用最终最大值算一次指数就行。这也是寄存器驻留方案为什么简单：原始值在手，随时可以按最终最大值重算。在线缩放主要用在没法驻留原始数据、又想一遍扫完的场景（比如长行的共享内存方案，或者流式数据）。

2.3 一个可跑的向量化 Softmax Kernel

下面这个版本是我在项目里最常用的，每个 block 处理一行，行内按 float4 向量化访问，每个线程把分到的数据存进寄存器，全程只读一次全局内存：

#include <cuda_runtime.h> #include <float.h> __device__ __forceinline__ float warpReduceMax(float v) { for (int off = 16; off > 0; off >>= 1) v = fmaxf(v, __shfl_xor_sync(0xffffffff, v, off)); return v; } __device__ __forceinline__ float warpReduceSum(float v) { for (int off = 16; off > 0; off >>= 1) v += __shfl_xor_sync(0xffffffff, v, off); return v; } // 每个 block 处理一行；cols4 是 float4 数量 template<int VPT> __global__ void softmaxRowOnlineVec4(const float4* __restrict__ in, float4* __restrict__ out, int cols4) { const int row = blockIdx.x; const int tid = threadIdx.x; const int nt = blockDim.x; const float4* rowIn = in + (size_t)row * cols4; float4* rowOut = out + (size_t)row * cols4; float4 x[VPT]; float e[VPT * 4]; int idx[VPT]; float lmax = -FLT_MAX; for (int k = 0; k < VPT; ++k) { idx[k] = k * nt + tid; if (idx[k] < cols4) { x[k] = rowIn[idx[k]]; lmax = fmaxf(lmax, fmaxf(fmaxf(x[k].x, x[k].y), fmaxf(x[k].z, x[k].w))); } } float wmax = warpReduceMax(lmax); __shared__ float s_warp[32]; const int lane = tid & 31; const int wid = tid >> 5; if (lane == 0) s_warp[wid] = wmax; __syncthreads(); const int nwarps = nt >> 5; if (wid == 0) { float v = (lane < nwarps) ? s_warp[lane] : -FLT_MAX; v = warpReduceMax(v); if (lane == 0) s_warp[0] = v; } __syncthreads(); const float gmax = s_warp[0]; __syncthreads(); // 必须等所有线程读完 gmax，才能复用 s_warp float lsum = 0.0f; for (int k = 0; k < VPT; ++k) { if (idx[k] < cols4) { e[k*4+0] = __expf(x[k].x - gmax); e[k*4+1] = __expf(x[k].y - gmax); e[k*4+2] = __expf(x[k].z - gmax); e[k*4+3] = __expf(x[k].w - gmax); lsum += (e[k*4+0] + e[k*4+1] + e[k*4+2] + e[k*4+3]); } } float wsum = warpReduceSum(lsum); if (lane == 0) s_warp[wid] = wsum; __syncthreads(); if (wid == 0) { float v = (lane < nwarps) ? s_warp[lane] : 0.0f; v = warpReduceSum(v); if (lane == 0) s_warp[0] = v; } __syncthreads(); const float invSum = 1.0f / s_warp[0]; for (int k = 0; k < VPT; ++k) { if (idx[k] < cols4) { float4 o; o.x = e[k*4+0] * invSum; o.y = e[k*4+1] * invSum; o.z = e[k*4+2] * invSum; o.w = e[k*4+3] * invSum; rowOut[idx[k]] = o; } } }

几个要点说一下。第一，float4 一次搬 16 字节，配合每个 warp 内部连续线程访问连续地址，访存效率能接近峰值带宽。第二，指数结果也被存进了寄存器数组 e，这样后面写回的时候不用重新算一遍 exp，省掉的计算量在长行上非常可观。第三，块内归约先走 warp shuffle，再由每个 warp 的 lane 0 写共享内存，最后让第一个 warp 做终极归约。这种两级归约比全走共享内存快很多。

Shared memory 数组 s_warp 在这里被复用了两次，一次放最大值、一次放求和。第一次读完 gmax 之后必须插一个 __syncthreads()，否则有的线程可能还在读 s_warp[0]，别的线程已经把它改写成求和结果了，这是典型的数据竞争现场。真要偷懒也可以开两个共享变量，但做这种底层优化，还是养成“每次复用共享内存前想清楚同步边界”的习惯比较好。

2.4 长行怎么办：共享内存版本的取舍

上面这个寄存器驻留版本，VPT 越大，每个线程要持有的 float4 和指数结果就越多。比如一行 4096 个 float，256 线程，每个线程 4 个 float4，光数据就是 64 个寄存器，加上指数结果直接奔着 100 多寄存器去了，寄存器溢出到局部内存反而更慢。

这种情况我一般把数据先进共享内存：整个 block 合作把一行读进shared，然后基于共享内存算最大值和求和，最后再从共享内存读出来做归一化写回。全局内存的读取仍然只有一次，共享内存虽然访问了两次，但它的带宽比显存高得多，整体还是划算。

共享内存版本的边界条件是行的最大长度。一个 block 的共享内存默认上限一般只有几十 KB，一行 8192 个 float 就要 32KB，再加别的开销可能就超了。真遇到超长行，可以一个 block 处理半行、多 block 协作做跨 block 归约，但那套复杂度就上来了。我在实际部署里遇到的大部分 Softmax 都在 1024 到 4096 这个量级，寄存器驻留加共享内存这两板斧已经够用。

2.5 精度与特殊情况的处理

默认我用 __expf，它比 expf 快不少，精度大概低一两个 ulp，推理场景完全够。如果做训练或者对数值极其敏感，换回 expf 也不影响结构。还有一个常用技巧是用 exp2 的硬件指令代替 exp：exp(x - m) 可以写成 exp2((x - m) * log2e)，因为 GPU 上 exp2 是原生快速指令。实测很多卡上这样写更快，代价是精度再掉一点点。

需要特别小心的是全掩码或者空行的情况。如果整行都被 mask 成 -inf，最大值归约出来还是 -FLT_MAX，指数全变成 0，总和也是 0，最后除法直接得到 NaN。我在代码里一般会让调用方保证至少有一个有效元素，或者对这种行单独走一个常数输出路径。

3. GEMM 优化：分块、寄存器累加与共享内存的三角关系

3.1 三层循环为什么不行

很多第一次写 GEMM 的人上来就是最朴素的 i、j、k 三层循环：C[i][j] 累加 A[i][k] * B[k][j]。这个写法在 GPU 上会被打得很惨，原因不在计算本身，而在访存复用。每个输出元素都要把 A 的一整行和 B 的一整列拉进来，A 矩阵的每个元素实际上被重复读了 N 次，B 的每个元素被重复读了 M 次。

当然 L1 和 L2 缓存会帮你兜掉一部分重复读取，但这个写法完全依赖缓存命中的运气，无法保证数据在片上停留的时间足够长。真正的做法是人肉管理数据复用：让一个线程块先把一小块 A 和一小块 B 搬进共享内存，然后这个小块里所有的输出元素都从共享内存取数。共享内存再往寄存器搬，寄存器里的数据会被算很多次。这样数据搬运量就从“每个输出元素都要访存”降到了“每个输出分块只访存一次”。

这套思路还有一个副产品：共享内存天然起到了同步屏障的作用，一个 block 内所有线程看到的数据是一致的，不需要处理复杂的跨线程数据依赖。

3.2 分块 Kernel 的线程映射与代码结构

下面是我常用的一个教学级 SGEMM 结构：每个 block 算 128x128 的输出分块，256 个线程，每个线程负责一个 8x8 的输出子块。内层 K 每次推进 8 个元素：

#define BM 128 #define BN 128 #define BK 8 #define TM 8 #define TN 8 #define THREADS 256 // C = A * B，行主序，假设 M、N 都能被分块大小整除 __global__ void sgemmTileKernel(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, int M, int N, int K) { const int bx = blockIdx.x, by = blockIdx.y; const int tid = threadIdx.x; const int cRow = tid / 16; // 线程负责的 8x8 输出块行号 const int cCol = tid % 16; // 线程负责的 8x8 输出块列号 __shared__ float As[BM][BK + 1]; __shared__ float Bs[BK][BN + 1]; float acc[TM][TN] = {0.0f}; for (int k0 = 0; k0 < K; k0 += BK) { for (int i = tid; i < BM * BK; i += THREADS) { int r = i / BK; int c = i % BK; As[r][c] = A[(by * BM + r) * K + (k0 + c)]; } for (int i = tid; i < BK * BN; i += THREADS) { int r = i / BN; int c = i % BN; Bs[r][c] = B[(k0 + r) * N + (bx * BN + c)]; } __syncthreads(); float regA[TM]; float regB[TN]; #pragma unroll for (int k = 0; k < BK; ++k) { #pragma unroll for (int i = 0; i < TM; ++i) regA[i] = As[cRow * TM + i][k]; #pragma unroll for (int j = 0; j < TN; ++j) regB[j] = Bs[k][cCol * TN + j]; #pragma unroll for (int i = 0; i < TM; ++i) #pragma unroll for (int j = 0; j < TN; ++j) acc[i][j] += regA[i] * regB[j]; } __syncthreads(); } for (int i = 0; i < TM; ++i) for (int j = 0; j < TN; ++j) C[(by * BM + cRow * TM + i) * N + (bx * BN + cCol * TN + j)] = acc[i][j]; }

这个代码结构上有三个层次。第一层是全局内存到共享内存，cooperative load 让 256 个线程一起把 A 和 B 的分块搬进去。第二层是共享内存到寄存器，每个线程按自己负责的输出位置取对应的 A 行片段和 B 列片段。第三层是寄存器里的 8x8 累加，这 64 个 FMA 完全不需要访存。三层之间用 __syncthreads() 隔开，保证共享内存不会被过早覆盖。

这里线程映射的重点是：两个连续的线程（tid 差 1）对应的是同一个输出块行、相邻的输出块列，所以它们在算 B 时访问的是 B 共享内存同一行里相邻的 8 列。这个映射直接决定了访存的合并度，也决定了下文要说的 bank conflict 情况，改动之前一定要想清楚。

3.3 共享内存 bank conflict：怎么看、怎么规避

共享内存的硬件结构是 32 个 bank，每 bank 4 字节，一个 warp 里不同线程访问同一个 bank 的不同地址就会产生冲突，冲突会让这次访问被拆成多次串行执行。这是新手最容易忽略的性能杀手，因为代码完全正确、结果完全正确，就是慢。

上面这个 kernel 里，A 的共享内存访问基本是广播模式（同一 warp 的线程读相同地址），冲突很小。B 的访问则有问题：连续线程访问 Bs[k][cCol*8 + j]，地址之间隔了 8 个 float，也就是 32 字节。按 bank 编号算一下，8 的倍数对 32 取模只有 0、8、16、24 四种结果，16 个线程挤在 4 个 bank 上，形成 4 路冲突。这也是为什么很多人照这个模板写出来，跑起来带宽和性能都差一截。

规避手段分几个等级。最简单的给共享内存行加 padding，能解决的是跨行访问时的冲突，解决不了同一行内 8 步长带来的这种规律冲突。真正工程上用的方案是把 B 的存储做 swizzle：把列索引跟行索引做一个异或重排，让同一 warp 的线程访问时分散到不同 bank。另一个方案是改变线程映射，让 warp 沿线程行方向展开，但这样 A 那边又会出类似的冲突。再往上就是 CUTLASS 那套复杂的向量化加 swizzle 体系，不是一篇文章能讲完的。

我在实战里的态度是：先用简单 kernel 把正确性和整体性能跑通，再用 profiler 的 bank conflict 指标确认问题严重程度，最后决定要不要上 swizzle。很多形状下，改一版向量化浮点加载就能把冲突压下去不少，不需要一开始就把最复杂的方案搬上来。

3.4 让内存流水线跑起来：双缓冲与异步拷贝

上面代码的主循环里，全局内存加载和计算是串行交替的：先加载、同步、计算、同步，再加载下一块。加载下一块数据的时间被完全暴露在关键路径上，计算单元在等待。解决办法是双缓冲：准备两块共享内存，计算当前块的同时，用异步拷贝指令把下一块的数据从全局内存搬到另一块共享内存。

现代 CUDA 提供了 cp.async 这类异步拷贝能力，数据从全局内存到共享内存不需要经过寄存器，也不占用线程的计算时间。配合 mbarrier 做到达机制，可以让加载和计算的流水真正重叠起来。改造之后，主循环会变成：异步发起下一块加载，同步等待当前块数据到达，计算当前块。如果块数足够多，加载时间基本被计算时间完全盖住。

这个流水线改造对性能的提升在小矩阵上尤其明显，因为小矩阵本身计算时间短，加载时间占比高。不过它也会显著增加代码复杂度，而且依赖编译器版本和卡的功能支持。我的建议是：第一版先把双缓冲的手写版本跑通，理解同步关系，再考虑上异步拷贝。直接一步到位上 cp.async，出了问题很难排查，因为你分不清是数据竞争还是硬件特性没用好。

3.5 参数选择与下一步：Tensor Core 是另一个世界

分块参数不是越大越好。BM、BN 决定共享内存占用和每块的并行度，TM、TN 决定每个线程的寄存器压力。寄存器太多会导致 occupancy 下降，太少又无法充分掩盖访存延迟。我一般从 128x128 配 8x8 开始调，再根据 profiler 里的 occupancy、spill 指标微调。BK 一般取 8 或 16，太小访存指令占比高，太大共享内存装不下。

还有向量化加载的问题：全局内存加载如果能用 float4，一次搬 16 字节，会让共享内存加载和最终写回的效率都上一个台阶。要求是 K 方向的地址对齐，矩阵有 padding 或者 K 是 4 的倍数时可以直接上。

至于 Tensor Core，那是另一套编程模型，走的是 mma 指令或者 wmma API，对数据排布的要求完全不一样。GEMM 的叉乘形状也变了：普通 CUDA 核走 FMA 累加，Tensor Core 走的是 16x8x8 这类硬件矩阵指令，性能上限能再翻好几倍。这个我会放到这一系列的后几篇单独讲，先把手写 SGEMM 这套基本功练扎实。

4. 实测环节：性能数字怎么读、怎么验证正确性

4.1 固定测量方法再谈优化

我每次对比性能都固定用同一套流程：先写一个 CPU 朴素实现做正确性基准，再写一个 CUDA 朴素版本做性能基线，然后才是优化版。计时用 CUDA event，而不是 CPU 的 clock 函数，因为 kernel 是异步执行的，CPU 端时间不可靠。

cudaEvent_t t0, t1; cudaEventCreate(&t0); cudaEventCreate(&t1); cudaEventRecord(t0); kernel<<<grid, block>>>(args...); cudaEventRecord(t1); cudaEventSynchronize(); float ms = 0.0f; cudaEventElapsedTime(&ms, t0, t1);

注意要加 warmup。GPU 第一次跑 kernel 时有很多初始化开销，直接计时会把第一轮的数据混进去，数字会虚高或虚低。我习惯先跑三到五轮丢弃，再取后面几轮的稳定值。

Softmax 的性能用有效带宽衡量：读一遍加写一遍的总字节数除以耗时。GEMM 用 FLOPS 衡量：2 * M * N * K 除以耗时。两个指标都要标注输入形状，因为同一个 kernel 在不同形状下的表现差异巨大。我这边随手记录过一组数据（消费级 8GB 显存卡，fp32，仅供参考）：Softmax 朴素三遍读版本大概只有峰值带宽的一半左右，向量化在线版本能到 85% 以上；GEMM 三重循环版本不到 1 TFLOPs，分块加寄存器累加能到几个 TFLOPs，跟厂商闭源库比还有明显差距，但思路已经跑通了。

4.2 编译参数与代码检查

编译的时候一定要开 -O3，另外强烈建议加 -Xptxas -v 看寄存器使用情况。它会输出每个 kernel 用了多少寄存器、有没有发生 local memory spill。spill 一旦出现，性能会断崖式下跌，这是最常见的“代码看着没问题但就是慢”的元凶。看到寄存器超了，优先缩小 TM 或 TN。

nvcc -O3 -arch=sm_80 -Xptxas -v softmax_bench.cu -o softmax_bench

正确性验证我习惯用随机输入加边界值：全零、全负、包含一个很大的正数、包含一个很大的负数、行内全相等。Softmax 对这几个 case 很敏感，一个不注意就会出 NaN。GEMM 用几组不同形状的矩阵跟 CPU 结果做对比，误差阈值 fp32 一般放 1e-4 到 1e-5 之间，如果开了 __expf 之类快速指令，阈值要放宽一点。

4.3 profiler 里重点盯这几个指标

跑完一遍用 profiler 看数据，我最关心的指标随着算子类型变化。Softmax 先看 DRAM 总吞吐，理想状态下应该接近显存峰值；再看 L1/L2 命中率和 warp stall 原因，如果大量 stall 在访存等待上，说明向量化没做够。GEMM 先看 SM 活跃度和指令发射率，再看共享内存 bank conflict 的总次数，以及寄存器溢出量。

还有一个容易被忽略的点：grid 和 block 的配置要跟数据形状匹配。Softmax 如果行数很少，每个 block 只处理一行，可能只调度了几个 block，GPU 大部分计算单元在睡觉。这种情况要么让一个 block 处理多行，要么调整 block 内线程数。GEMM 同理，分块后 block 总数要能填满整张卡，否则算力永远上不去。

5. 常见问题与排查技巧实录

5.1 共享内存数据竞争

表现是结果时对时错，或者用同一份数据跑两次结果不一样。原因基本是缺了 __syncthreads()，或者共享内存数组被复用之前没有做同步。我排查这类问题最笨但最有效的办法是：在可疑位置先无脑加满 __syncthreads()，如果问题消失，再一个个删，直到找到真正的边界。虽然粗暴，但比盯着代码猜快得多。

5.2 为什么实测带宽只有一半

Softmax 向量化版本如果实测带宽上不去，先检查输入地址对齐。float4 要求 16 字节对齐，如果行首地址因为矩阵宽度原因不对齐，向量化加载会失败，编译器可能退化成多个标量加载。再检查索引类型，用 int 的时候行号乘行宽很容易溢出，一旦溢出访存就乱套了，速度自然上不去。这里我用 (size_t)row * cols4 是有意为之，老手也会在这里翻车。

5.3 bank conflict 数字到底在哪里看

profiler 的共享内存相关指标里会直接给 bank conflict 计数。如果没有明显冲突，这个数字应该是 0 或者很小。一旦看到几十万上百万次冲突，就要回到映射关系上重新想。我在前面的 kernel 里分析过 B 访问的 4 路冲突，实测中这个问题的性能损失常常在 20% 到 50% 之间，值得认真处理。

5.4 寄存器溢出导致变慢

代码逻辑没错，编译也过了，profiler 显示 occupancy 极低，多半是寄存器溢出了。用 -Xptxas -v 看一眼 local memory 的大小，非零就说明 spill 了。解决办法是降低每个线程的工作量，把 TM、TN 从 8x8 降到 4x8，或者减小 VPT。也可以调 launch bounds 限制最大寄存器数，但更推荐先降工作量，别一开始就上这种偏方。

5.5 我测出来没别人快，是不是写错了

别急着怀疑自己。厂商闭源库的 GEMM 是几十年积累的成果，包含 Tensor Core、自动调参、多级缓存优化、针对不同形状的特化路径。我手写 SGEMM 在第一篇能达到它的三分之一甚至五分之一都是正常水平。跟自己比才是正确的目标：上一版 500 GFLOPS，这一版 2 TFLOPS，这个进步是真实的、可复现的。

最后再分享一个我个人的经验：学 CUDA 算子优化，不要一开始就追 Tensor Core 和最高性能，先把 Softmax、GEMM 这两个最基础的算子，用最朴素到逐步优化的完整流程走一遍。这个过程能让你把向量化、分块、归约、共享内存同步、profiler 分析这些基本功全部过一遍，而且每一步的瓶颈都特别清晰。我后来写融合算子遇到的很多问题，追根溯源都能回到这些基础操作上。把这两个算子的手感练出来，后面碰到再复杂的 kernel 也不会慌。