GPU Kernel Tips: From Warps to Tile Sizes

No, but they ask the same question: when the 32 threads of a warp access memory together, how are their addresses laid out?

不是一回事，但问的是同一个问题：一个 warp 的 32 个线程同时访存时，地址是怎么分布的。

Because a warp's 32 threads are the batch that issues memory requests together. That batch of addresses is what the hardware looks at.

因为一个 warp 的 32 个线程是同时发访存请求的那一批，硬件看的就是这一批地址长什么样。

Coalescing is about global memory. When 32 threads read 128 consecutive bytes, the reads merge into one transfer. When the addresses are scattered, say each thread one row apart, they split into as many as 32 transfers.

coalesce 说的是 global memory。32 个线程读连续的 128 字节，能合并成一次传输；地址要是散的，比如每个线程隔一行读一个，最多会拆成 32 次传输。

Those are about shared memory. It is split into 32 banks, each 4 bytes wide; the n-th 4-byte word lives in bank n % 32, and each bank can serve one address per cycle. If the 32 threads hit 32 different banks, the read finishes in one pass. If several threads hit different addresses in the same bank, they queue up, up to 32 times slower. If they all read the same address, that is a broadcast, which is fine.

说的是 shared memory。它分成 32 个 bank，每个 bank 4 字节宽，第 n 个 4 字节落在 bank n % 32，一个周期里每个 bank 只能服务一个地址。32 个线程落在 32 个不同的 bank，一次完成；多个线程读同一个 bank 的不同地址，就要排队，最坏慢 32 倍；大家都读同一个地址是广播，没问题。

From the top down. A grid is every block one kernel launch starts. A block, also called a CTA (cooperative thread array), has up to 1024 threads; the two names mean the same thing. A warp is a group of 32 threads cut out of a block. The SM is the hardware, and blocks are what get placed on SMs to run.

从大往小说。grid 是一次 kernel launch 起的所有 block。block 也叫 CTA（cooperative thread array），是同一个东西的两个名字，最多 1024 个线程。warp 是 block 里每 32 个线程切出来的一组。SM 是硬件，block 是被分配到 SM 上去跑的。

No. A block stays on one SM from start to finish. Think about what a block can do: its threads share one shared memory and can meet at __syncthreads(). Shared memory is on-chip storage inside an SM, and that barrier is implemented inside the SM too. Split the block across two SMs and neither would work.

不会。一个 block 从开始到结束都只待在一个 SM 上。想想 block 能做什么：它的线程共用一块 shared memory，还能在 __syncthreads() 那里会合。shared memory 是 SM 里的片上存储，这个 barrier 也在 SM 内部实现，block 拆到两个 SM 上，这两样就都用不了了。

Whenever an SM has room, the scheduler gives it the next block, and when one finishes, another takes its place. If resources allow, one SM holds several blocks at once.

哪个 SM 有空位，调度器就塞一个 block 过去，跑完一个再补一个。资源够的话，一个 SM 能同时放好几个 block。

By linear thread index: tx + ty * blockDim.x + …. A (32, 32) block is 32 warps, one per row ty. Grids and blocks are how you split up the work in code; the warp is the unit the hardware actually issues instructions for.

按线性编号切：tx + ty * blockDim.x + …。一个 (32, 32) 的 block 就是 32 个 warp，每一行 ty 正好一个。grid 和 block 是你写代码时的划分，warp 才是硬件真正一条条发指令的单位。

Only through global memory. That is why a reduction has each block write one partial sum, and the total takes either a second pass or an atomic.

只能走 global memory。所以做归约时，每个 block 先写出一个部分和，最后要么再跑一轮，要么用 atomic。

With the two things we just talked about: the shared memory everyone in the block shares, and __syncthreads(). Each block adds its 256 numbers pairwise in its own shared memory and syncs after every round:

__global__ void block_sum(const float* x, float* out, int n) {
    __shared__ float s[256];                 // this block's copy, on its SM
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;
    s[tid] = (i < n) ? x[i] : 0.f;
    __syncthreads();                         // waits for this block only

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) s[tid] += s[tid + stride];
        __syncthreads();
    }
    if (tid == 0) out[blockIdx.x] = s[0];    // one partial sum per block
}

就用刚才说的那两样：大家共用的 shared memory，和 __syncthreads()。每个 block 在自己的 shared memory 里把 256 个数两两相加，每加一轮同步一次：

__global__ void block_sum(const float* x, float* out, int n) {
    __shared__ float s[256];                 // 每个 block 各有一份，在它的 SM 上
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;
    s[tid] = (i < n) ? x[i] : 0.f;
    __syncthreads();                         // 只等本 block 的线程

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) s[tid] += s[tid + stride];
        __syncthreads();
    }
    if (tid == 0) out[blockIdx.x] = s[0];    // 每个 block 写一个部分和
}

Yes. It covers this block's 256 threads, that is, 8 warps. Other blocks sit on other SMs, add up their own numbers, and write their partial sums back to global memory.

对，它只管本 block 的 256 个线程，也就是 8 个 warp。别的 block 在别的 SM 上，各加各的，最后把部分和写回 global memory。

Once stride drops below 32, if (tid < stride) leaves part of one warp working and the rest idle. That is called warp divergence. By then everything left fits in one warp, so __shfl_down_sync passes values directly between that warp's registers, with no shared memory and no __syncthreads().

The warp runs both sides: first the if, with the other threads masked off and idle, then the else. You pay for both branches. Different loop trip counts count too; the whole warp waits for its slowest thread.

两边都跑：先跑 if，另一半线程被 mask 掉空转；再跑 else。代价是两个分支的时间相加。循环次数不一样也算，整个 warp 要陪最慢的线程跑完。

Write short branches as cond ? a : b, which usually compiles to a select that every thread executes. For long ones, sort or group the data so the threads of one warp take the same path. Keep data-dependent branches out of hot loops; if one is unavoidable, compute both sides and select at the end.

短分支写成 cond ? a : b，一般会编译成所有线程都执行的条件选择指令。长的就按条件把数据排好或分组，让同一个 warp 里的线程走同一条路。热循环里别放依赖数据的分支，实在躲不开，就两边都算，最后 select。

Inside one SM. A warpgroup is 4 consecutive warps in the same block, 128 threads, with the first warp index aligned to a multiple of 4; in code it is usually just threadIdx.x / 128. A block lives on one SM, so a warpgroup does too.

在一个 SM 里面。warpgroup 是同一个 block 里连续的 4 个 warp，128 个线程，起始 warp 编号对齐到 4 的倍数，代码里一般就是 threadIdx.x / 128。block 只在一个 SM 上，warpgroup 当然也只在一个 SM 里。

The SM is split into 4 partitions, each with its own warp scheduler and Tensor Core. A block's warps generally go to the 4 partitions round-robin by index, so 4 consecutive warps land one per partition: a warpgroup spans all 4 partitions of the SM.

SM 分成 4 个分区，每个分区有自己的 warp 调度器和 Tensor Core。block 的 warp 一般按编号轮流分到这 4 个分区上，所以连续的 4 个 warp 正好每个分区一个：一个 warpgroup 横跨整个 SM 的 4 个分区。

Because Hopper's Tensor Core instruction, wgmma, has to be issued by 4 warps together; one warp can't issue it alone. The picture above shows why: the 4 warps sit in 4 partitions, and wgmma has the 4 partitions' Tensor Cores compute one block of the matrix together. Warpgroups exist for wgmma.

因为 Hopper 的 Tensor Core 指令 wgmma 要求 4 个 warp 一起发起，一个 warp 单独发不了。看上面那张图就明白了：4 个 warp 分在 4 个分区，wgmma 是让 4 个分区的 Tensor Core 一起算一块矩阵。warpgroup 就是为它引入的。

Several, each with its own job. One can be a producer that only moves data, and one or two others consumers that only run wgmma. They share the same shared memory, so they must be in the same block. More on the split later.

好几个，各干各的。比如一个当 producer 专门搬数据，另外一两个当 consumer 专门跑 wgmma。它们要共用同一块 shared memory，所以必须在同一个 block 里。怎么分工，后面细说。

Blackwell's MMA instruction is tcgen05, and a single thread can issue it. The accumulator also moves out of registers into dedicated Tensor Memory.

Blackwell 的 MMA 指令叫 tcgen05，由单个线程就能发起，accumulator 也从寄存器挪到了专门的 Tensor Memory 里。

Much less. You no longer need 4 warps lined up, the accumulator no longer takes registers, and the register pressure we'll get to later is much lighter on Blackwell.

没那么核心了。不用再凑齐 4 个 warp，accumulator 也不占寄存器，后面要讲的寄存器压力在 Blackwell 上也轻得多。

Clearly not. Keep two numbers apart. Resident warps: up to 64 can live on an SM at once on Hopper, each holding its own registers. Issuing warps: the SM has 4 partitions with a scheduler each, so about 4 warps issue an instruction every cycle.

显然不是。要分清两个数。驻留的 warp：Hopper 上一个 SM 同时能挂 64 个，寄存器各占各的。在发指令的 warp：SM 分成 4 个分区，每个分区一个 scheduler，所以每个周期大约 4 个 warp 在发指令。

Waiting, for example on a memory request. When one warp stalls, its scheduler picks another ready warp on the next cycle. That is how a GPU hides latency.

在等，比如发了访存请求在等数据。某个 warp 卡住了，scheduler 下个周期就挑另一个就绪的 warp。GPU 就是这样藏延迟的。

Nothing is saved. A resident warp keeps its own slice of the register file the whole time, so switching just means issuing from a different warp. The cost shows up elsewhere: the register file has to hold every resident warp at once.

什么都不用存。驻留的 warp 一直占着自己那块寄存器堆，切换只是换一个 warp 发指令。代价在别处：寄存器堆要同时装下所有驻留的 warp。

Enough that some warp is always ready. Try it:

够让任何时候都有 warp 是就绪的。拖一下试试：

Occupancy is the fraction of the warp limit that is resident. Its only job is hiding latency; once latency is hidden, more does nothing.

occupancy 是驻留 warp 占上限的比例。它唯一的作用是藏延迟，延迟藏住了，再往上加就没用。

Plenty on Hopper. Instead of many warps, they hide latency inside one block with asynchronous copies and asynchronous math. When we get to FlashAttention running one block per SM, this is exactly that case.

Hopper 上很多。它们不靠 warp 多，而是靠一个 block 内部的异步搬运和异步计算来藏延迟。后面说到 FlashAttention 一个 SM 只放一个 block，就是这种情况。

All 4 warps do have to reach the same wgmma, with uniform control flow. But you don't insert a __syncthreads(), and how the hardware coordinates the 4 warps isn't documented.

4 个 warp 确实都要执行到同一条 wgmma，控制流也必须一致。但不需要你插 __syncthreads()，硬件内部怎么协调这 4 个 warp 也没有公开。

The key is that wgmma is asynchronous. wgmma.mma_async returns as soon as it is issued; the warps keep going, and the results land in the accumulators later. What actually stops you is the step where you choose to wait for the results.

关键在于 wgmma 是异步的。wgmma.mma_async 发出去就返回，warp 接着往下走，结果过一阵才落到 accumulator 里。真正会让你停下来的，是你主动去等结果的那一步。

Three instructions work together:

• wgmma.fence: before issuing, makes earlier writes to registers and shared memory visible to wgmma.
• wgmma.commit_group: bundles the wgmmas issued so far into a group.
• wgmma.wait_group N: blocks until at most N groups are still in flight.

三条指令配合：

• wgmma.fence：发射前用，保证之前对寄存器和 shared memory 的写入对 wgmma 可见。
• wgmma.commit_group：把已经发出去的几条 wgmma 打包成一组。
• wgmma.wait_group N：阻塞到只剩 N 组没完成。

Right. You must wait before reading the accumulators and before overwriting shared memory that still holds operands. Spend the time in between on other work, like waiting for the next tile or running softmax, and the cost is hidden. The worst thing you can do is wait right after issuing.

对。读 accumulator 之前、覆盖还装着操作数的 shared memory 之前，必须等。中间这段时间拿去干别的，比如等下一块数据、做 softmax，开销就藏住了。最亏的是发完马上就等。

It does. It is each SM's own on-chip SRAM: 256 KB per SM, shared between L1 and shared memory. Up to 228 KB can be shared memory, and one block can use up to 227 KB.

有关。它是每个 SM 自己的片上 SRAM：每个 SM 256 KB，L1 和 shared memory 共用这一块，最多 228 KB 能划给 shared memory，单个 block 最多用 227 KB。

To keep the Tensor Cores fed. The faster the Tensor Cores, the earlier data has to be brought in: tiles get bigger, and several buffers are kept so the next tile is already in place while the current one is computed. All of that lives in shared memory.

为了喂饱 Tensor Core。Tensor Core 越快，就越要提前把数据搬进来：tile 要更大，还要留好几份缓冲，让下一块数据在算当前这块时就已经到位。这些都要占 shared memory。

A block gets 48 KB by default, and static __shared__ arrays can't go past that. For more, use dynamic shared memory (extern __shared__) and raise the limit before launch with cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes).

每个 block 默认只能用 48 KB，静态声明的 __shared__ 数组超不过这个数。要更多就用动态 shared memory（extern __shared__），并且在 launch 之前用 cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes) 把上限调上去。

Reuse. The classic case is GEMM. In the naive version each thread computes one C[i][j]. How many times is each element of A read from HBM?

复用。最经典的是 GEMM。naive 版本里一个线程算一个 C[i][j]。你算算，A 的每个元素要从显存读几次？

N times, and each element of B M times. The tiled version gives each block a TS×TS tile of C. It walks along K, loads the matching tiles of A and B into shared memory, and every thread reads from there. HBM traffic drops to 1/TS.

#define TS 32
__global__ void gemm(const float* A, const float* B, float* C, int M, int N, int K) {
    __shared__ float sA[TS][TS], sB[TS][TS];
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TS + ty, col = blockIdx.x * TS + tx;
    float acc = 0.f;
    for (int t = 0; t < K; t += TS) {
        sA[ty][tx] = (row < M && t + tx < K) ? A[row * K + t + tx] : 0.f;
        sB[ty][tx] = (t + ty < K && col < N) ? B[(t + ty) * N + col] : 0.f;
        __syncthreads();                 // the whole tile is loaded
        for (int k = 0; k < TS; ++k) acc += sA[ty][k] * sB[k][tx];
        __syncthreads();                 // everyone is done before it is overwritten
    }
    if (row < M && col < N) C[row * N + col] = acc;
}

N 次，B 的每个元素读 M 次。tiled 版本让一个 block 负责 C 的一个 TS×TS 小块，沿 K 方向每次把 A、B 对应的小块搬进 shared memory，所有线程从里面读，显存读取量降到 1/TS。

#define TS 32
__global__ void gemm(const float* A, const float* B, float* C, int M, int N, int K) {
    __shared__ float sA[TS][TS], sB[TS][TS];
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TS + ty, col = blockIdx.x * TS + tx;
    float acc = 0.f;
    for (int t = 0; t < K; t += TS) {
        sA[ty][tx] = (row < M && t + tx < K) ? A[row * K + t + tx] : 0.f;
        sB[ty][tx] = (t + ty < K && col < N) ? B[(t + ty) * N + col] : 0.f;
        __syncthreads();                 // 等整块搬完
        for (int k = 0; k < TS; ++k) acc += sA[ty][k] * sB[k][tx];
        __syncthreads();                 // 等所有人算完再覆盖
    }
    if (row < M && col < N) C[row * N + col] = acc;
}

Yes. The first waits until the tile is loaded; without the second, the next iteration overwrites data other threads are still reading. Pad with zeros at the edges too, or any size that isn't a multiple of TS reads out of bounds. Another common mistake is swapping the indices: sA holds this block's rows × the current K slice, and sB holds the current K slice × this block's columns.

都要。第一个等整块搬完；少了第二个，下一轮会覆盖别的线程还在读的数据。边界还要补 0，不然尺寸不是 TS 的整数倍就会越界。还有一个常见错误是下标写反：sA 存的是“本 block 的行 × 当前 K 段”，sB 存的是“当前 K 段 × 本 block 的列”。

Check it yourself. tx is contiguous within a warp, so the loads from A and B hit consecutive addresses: coalesced. In the inner loop, sA[ty][k] is the same address for the whole warp, a broadcast, and sB[k][tx] spans 32 different banks. No conflicts.

你自己检查一下。warp 内 tx 连续，读 A、B 的地址都连续，是 coalesced。内层循环里 sA[ty][k] 整个 warp 读同一个地址，是广播；sB[k][tx] 落在 32 个不同的 bank。没有冲突。

Very far.

很远。

In production it's hard to beat cuBLAS, true. But every next step carries over to other kernels. Roughly by payoff: have each thread compute several outputs (say 4×4) and reuse values from registers; vectorized float4 loads; double buffering, prefetching the next tile while computing this one; and finally Tensor Cores.

生产里确实很难打过 cuBLAS。但往下优化的每一步，换个算子都用得上。大致按收益排：一个线程算多个输出（比如 4×4），值放寄存器里复用；float4 向量化加载；double buffering，算这块的同时预取下一块；最后上 Tensor Core。

Yes, plus TMA to move data asynchronously. Take double buffering all the way and you get the multi-stage ring buffer FlashAttention uses.

对，再加上用 TMA 异步搬数据。double buffering 走到头，就是 FlashAttention 里那种多级环形缓冲。

Not directly. Four limits decide it, and the tightest one wins: shared memory, registers (65,536 per SM), warps (64), and blocks (32). Move the sliders:

不能直接指定。它由四个上限决定，取最紧的那个：shared memory、寄存器（每个 SM 65,536 个）、warp 数（64）、block 数（32）。拖一下滑块：

Right, by adjusting block size, shared memory use, and register use. To see the actual result, ask the runtime: cudaOccupancyMaxActiveBlocksPerMultiprocessor.

对，靠调 block 大小、shared memory 用量和寄存器用量。想知道实际结果，问 runtime：cudaOccupancyMaxActiveBlocksPerMultiprocessor。

It depends on the kernel. An ordinary kernel wants several, since more warps hide more latency. Something like FlashAttention is the opposite: one block takes most of the shared memory, so only one fits per SM.

看 kernel。普通 kernel 希望多放几个，warp 多了才好藏延迟。FlashAttention 这种正好反过来：一个 block 就吃掉一大半 shared memory，一个 SM 只放一个。

First, what is the bottleneck on Hopper?

先问你，Hopper 上的瓶颈是什么？

And each SM has a fixed number of them. If the consumers in this one block issue wgmma back to back, the Tensor Cores are already full; more blocks would only fight over the same units. MFU and HFU both measure whether the Tensor Cores are working. cutlass's Hopper GEMM runs one block per SM and still gets close to peak.

每个 SM 上的 Tensor Core 数量是固定的。只要这一个 block 的 consumer 能一条接一条地发 wgmma，Tensor Core 就已经满了，多放几个 block 也只是抢同一批 Tensor Core。MFU 也好、HFU 也好，看的都是 Tensor Core 有没有在干活。cutlass 的 Hopper GEMM 就是一个 SM 一个 block，照样接近峰值。

When the Tensor Cores go idle, usually for one of three reasons: waiting for data, the consumer busy with something else, or prologue and epilogue overhead. Compare three designs:

Tensor Core 空转的时候。一般是三种情况：等数据、consumer 在忙别的、头尾开销。对比三种设计：

Make shared memory a multi-stage ring buffer, say 4 stages, so the producer loads several tiles ahead. That is where most of FA3's 200 KB goes.

把 shared memory 做成多级环形缓冲，比如 4 级，让 producer 提前把后面几块搬进来。FA3 那 200 KB 大部分就花在这里。

In attention, softmax follows QKᵀ, and the Tensor Cores sit idle while it runs. FA3 has two consumer warpgroups take turns: while one runs softmax, the other runs its GEMM.

attention 里算完 QKᵀ 要做 softmax，这段时间 Tensor Core 闲着。FA3 让两个 consumer warpgroup 轮流来：一个做 softmax 时，另一个在跑 GEMM。

Right direction, but not quite the same. They solve the same problem: in a matrix multiply the same number gets multiplied many times, so you build a dedicated multiply-add array where data is read in once and reused inside the array, instead of going back to the registers for every multiplication.

方向对，但不完全一样。它们要解决的是同一个问题：矩阵乘里同一个数要被乘很多次，那就做一块专门的乘加阵列，让数据读进来一次、在阵列里被反复用，而不是每做一次乘法都去寄存器里取一趟。

Start with scale. A TPU's matrix unit is one big systolic array, for example 128 × 128 multiply-add units, and a chip has only a few of them. A GPU spreads many small units over its SMs: each Hopper SM has 4 Tensor Cores, one per partition, 528 in an H100.

先看规模。TPU 的矩阵单元是一整块很大的脉动阵列，比如 128 × 128 个乘加单元，一颗芯片上就那么几块。GPU 是把很多个小的分散在各个 SM 里：Hopper 每个 SM 有 4 个 Tensor Core，每个分区一个，H100 一共 528 个。

Work back from a number we'll get to later: each SM does 4096 FLOPs per cycle, that is, 2048 multiply-adds. Split over 4 Tensor Cores, that's 512 multiply-adds per Tensor Core per cycle. Much smaller than a TPU's whole array.

可以拿后面会算到的那个数倒推：每个 SM 每周期 4096 FLOPs，也就是 2048 个乘加，分到 4 个 Tensor Core 上，每个每周期 512 个乘加。比 TPU 那一整块小得多。

NVIDIA hasn't published its internals. It is generally thought to be closer to a set of dot-product units: one instruction comes in and finishes the multiply-adds on a small matrix in one go, rather than streaming data through hundreds of stages. So rather than a scaled-down TPU systolic array, think of it as a fixed-shape small matrix multiplier in each partition.

NVIDIA 没有公开内部结构。一般认为它更像一组点积单元：一条指令进来，在一小块矩阵上把乘加一口气做完，而不是让数据在几百级的阵列里流过去。所以与其说它是缩小版的 TPU 脉动阵列，不如说是每个分区里一块固定形状的小矩阵乘法器。

They are fed differently. A TPU's array is filled in whole blocks laid out by the compiler, and matrix dimensions are best rounded up to multiples of 128. A GPU's Tensor Cores are fed instruction by instruction by warps through mma / wgmma, and you move the data yourself from HBM into shared memory and on into the Tensor Cores. Working out N ≥ 64 later is working out whether that feeding path is fast enough.

喂法不一样。TPU 那块阵列由编译器排好数据整块灌进去，矩阵维度最好凑成 128 的倍数。GPU 的 Tensor Core 是 warp 用 mma / wgmma 指令一条条喂的，数据得你自己从显存搬到 shared memory，再送进 Tensor Core。后面算 N ≥ 64，算的就是这条喂料的路够不够快。

Exactly, and that's the main line again.

对，这就回到主线了。

When the main loop is short. Each tile has to set up the pipeline at the start and write results back to HBM at the end; with a long loop along K that is negligible, with a short one it shows. MoE is the classic case: each expert gets few tokens, so its matrices are small, sometimes with fewer tiles than SMs.

主循环短的时候。每个 tile 开头要建流水线，结尾要把结果写回显存，K 方向循环很长时可以忽略，循环短就很显眼。MoE 是典型：每个 expert 分到的 token 少，矩阵小，甚至 tile 数比 SM 数还少。

Write the grouped GEMM as a persistent kernel: launch as many blocks as SMs, keep each block resident, and let it loop, pulling tiles from a scheduler. That saves the repeated launches and pipeline setups.

把 grouped GEMM 写成 persistent 的：block 数等于 SM 数，每个 block 常驻，循环着从调度器领 tile，省掉反复 launch 和建流水线的开销。

The experts share K and N, but each has its own M. Take a prefix sum over their tile counts; whichever range a global tile index falls into names its expert, found with a binary search. This comes back when we extend FlashFFN to MoE at the end.

各个 expert 的 K、N 相同，M 不同。对各 expert 的 tile 数做前缀和，全局 tile 编号落在哪个区间就属于哪个 expert，二分查一下就行。最后讲 FlashFFN 拓展到 MoE 时，这个问题还会回来。

First see how far measured TFLOPS is from peak. Then open Nsight Compute: check tensor pipe utilization and what the warps are stalled on (the stall reasons). Stalling at a barrier while waiting for data looks completely different from stalling on other math.

先看实测 TFLOPS 离峰值多远。再打开 Nsight Compute，看 tensor pipe 的利用率和 warp 卡在什么上（stall reason）。卡在等数据的 barrier 上，和卡在算别的东西上，是两种完全不同的样子。

Why 3:1?

为什么是 3:1？

Then tell me how three warps issue one wgmma together.

那你说，三个 warp 怎么一起发一条 wgmma？

Right, the units are mixed up: the split is counted in warpgroups. A common setup is 1 producer warpgroup plus 2 consumer warpgroups, as in cutlass's Hopper GEMM and mostly in FA3's forward pass too. That's 12 warps and 384 threads, or 4:8 counted in warps.

对，单位搞混了，分工是按 warpgroup 算的。常见配置是 1 个 producer warpgroup 加 2 个 consumer warpgroup：cutlass 的 Hopper GEMM 是这样，FA3 的 forward 也以这个为主。一共 12 个 warp、384 个线程，按 warp 数算是 4:8。

The wgmma shape is m64 × nN × k16 for half precision: M is fixed at 64, and N goes from 8 to 256. Each consumer owns 64 rows of the output tile, so together M is 128. Both read the same B from shared memory and keep their accumulators in their own registers.

wgmma 的形状是 m64 × nN × k16（半精度）：M 固定 64，N 从 8 到 256。每个 consumer 管输出 tile 的 64 行，拼起来 M 就是 128。两边读 shared memory 里同一份 B，accumulator 各放在自己的寄存器里。

The producer uses TMA to copy data from HBM into the ring buffer in shared memory. TMA copies are asynchronous, and only one thread in the whole producer warpgroup actually issues them.

producer 用 TMA 把数据从显存搬进 shared memory 的环形缓冲。TMA 是异步拷贝，整个 producer warpgroup 里真正发 TMA 指令的只有一个线程。

Through barriers in shared memory (mbarriers). Each stage of the buffer has two. When the producer finishes loading a tile, the stage's "full" barrier completes and the consumers waiting on it go ahead; when the consumers are done with a stage, they arrive on its "empty" barrier, and only then can the producer overwrite it. Each side runs on its own and they only meet at the barriers.

靠 shared memory 里的 barrier（mbarrier）。每一级缓冲配两个：producer 搬完一块，“满”的 barrier 到位，consumer 在上面等；consumer 用完一级，到“空”的 barrier 上报到，producer 才能覆盖这一级。两边各跑各的，只在 barrier 上对一下。

Which is why cutlass and FA3 both start by rebalancing registers with setmaxnreg: the producer drops to a few dozen per thread, and the consumers take the freed budget and go above two hundred. Consumers keep their accumulators in registers, so that is exactly where the space is needed.

所以 cutlass 和 FA3 开头都会用 setmaxnreg 重新分配寄存器：producer 降到每线程几十个，省下的额度给 consumer 升到两百多个。consumer 要在寄存器里放 accumulator，正缺这些。

Lay out the shapes first. Take the first GEMM, S = Q·Kᵀ.

先把形状摆出来。看第一个 GEMM，S = Q·Kᵀ。

Where does each of the three live?

这三块各放在哪？

Right. S stays in registers and is never written back to shared memory, so it adds no memory traffic. What about the 128 in the middle?

对。S 一直留在寄存器里，不写回 shared memory，所以它不带来额外的访存。中间那个 128 呢？

No. One wgmma has the shape m64 nN k16, so it takes only 16 along K at a time. The 128 is cut into 8 slices and accumulated into S over 8 iterations. So we only need to do the math for a single wgmma.

不是。wgmma 一条指令的形状是 m64 nN k16，K 方向每次只吃 16。128 要拆成 8 段，循环 8 次累加进 S。所以只要算清楚一条 wgmma 的账就够了。

You count. A multiply-add is two operations.

你来数。乘加算两次运算。

Now turn that into time. What we want is how much one SM computes per cycle.

再换成时间。我们要的是一个 SM 一个周期能算多少。

…and you get 4096 FLOPs, that is, 2048 multiply-adds per SM per cycle.

4096 FLOPs，也就是每个 SM 每周期 2048 个乘加。

989.4 TFLOPS is the dense bf16 peak of an H100 SXM (the sparse figure is twice that), 1830 MHz is its boost clock, and the SXM part has 132 SMs. Multiply back: 4096 × 132 × 1.83e9 ≈ 989.4e12.

989.4 TFLOPS 是 H100 SXM 的 dense bf16 峰值（带稀疏的数字是它的两倍），1830 MHz 是 boost 频率，SXM 版有 132 个 SM。乘回去：4096 × 132 × 1.83e9 ≈ 989.4e12。

Under full load it often runs lower. Compute per cycle stays the same, but HBM bandwidth doesn't follow the SM clock, so a slower clock actually gives each cycle a few more HBM bytes. The order of magnitude doesn't change.

满载时常常更低。每周期的算力不变，但 HBM 带宽不随 SM 频率变，频率低了，每个周期能分到的 HBM 字节反而多一点。量级不变。

Where do those bytes come from, and at what bandwidth?

这些字节从哪来，带宽是多少？

The number is right; let me tighten the wording. Shared memory and L1 are the same on-chip storage, and this isn't coalescing in the global-memory sense: shared memory has 32 banks, each 4 bytes wide, so without bank conflicts it delivers exactly 128 bytes per cycle.

数是对的，说法再准一点：shared memory 和 L1 是同一块片上存储，这里不是 global memory 的合并访问，而是 shared memory 有 32 个 bank、每个 bank 4 字节宽，没有 bank conflict 的时候一个周期正好 128 字节。

Right. Notice the 16: every wgmma re-reads A, a fixed cost.

对。注意那个 16：每发一条 wgmma，都要把 A 重新读一遍，这是固定成本。

It assumes a warp reads all 32 banks in one cycle. The real bandwidth of the path wgmma reads its operands through isn't published, so the N we're about to get is an order-of-magnitude answer, not an exact one.

它是按一个 warp 一个周期读满 32 个 bank 算的。wgmma 读操作数那条通路的真实带宽没有公开，所以等会儿算出来的 N 是量级上的结论，不是精确值。

Not necessarily. wgmma can also take the A operand straight from registers, which takes those 16 cycles off shared memory. FA3 feeds P to its second GEMM (P·V) this way: after softmax, P is already in registers.

不一定。wgmma 的 A 操作数也可以直接从寄存器读，那 16 个周期就不压在 shared memory 上了。FA3 的第二个 GEMM（P·V）就是这么喂 P 的：softmax 算完，P 本来就在寄存器里。

Write it out.

写出来。

Right:

compute time ≥ memory time
0.5 N ≥ 16 + 0.25 N
N ≥ 64

So your 64 sits exactly at break-even, with no margin at all. Drag N:

对：

计算时间 ≥ 访存时间
0.5 N ≥ 16 + 0.25 N
N ≥ 64

所以你的 64 正好踩在临界点上，一点余量都没有。拖一下 N：

Compute over memory is 1.00 at N = 64, 1.33 at 128, 1.60 at 256, and tops out at 2. wgmma pins M at 64, so the fixed 16 cycles get spread thinner as N grows, but never disappear.

计算和访存之比：N = 64 是 1.00，128 是 1.33，256 是 1.60，最多到 2。wgmma 把 M 固定在 64，那 16 个周期的固定成本被 N 摊得越来越薄，但永远摊不掉。

That is where N's lower bound comes from; its upper bound depends on registers. One level further up there's HBM. Both are follow-ups you can open.

N 的下限就是这么来的，上限要看寄存器。再往上一层还有 HBM，这两个可以接着问。

Count the registers first. S is [64, N] in fp32, spread over a warpgroup's 128 threads. How many per thread?

先数数寄存器。S 是 [64, N] 的 fp32，摊到一个 warpgroup 的 128 个线程上，每线程几个？

Plus the output accumulator O, [64, 128] in fp32: another 64 per thread. That's 192 already, against a cap of 255, and P, the softmax statistics, and addresses still need room.

还有输出累加器 O，[64, 128] 的 fp32，每线程又是 64 个。光这两个就 192 了，上限才 255，还要放 P、softmax 的统计量和地址。

Exactly. Shared memory sets N's lower bound; registers set its upper bound.

对。N 的下限由 shared memory 定，上限由寄存器定。

A KV tile is K and V, each [N, 128] in bf16. You count.

一个 KV tile 是 K、V 各一块 [N, 128] 的 bf16。你来数。

It does. At this level N doesn't matter; M does, the number of Q rows a block owns. Each byte of KV buys M FLOPs. The balance point is 989.4 ÷ 3.35 ≈ 295, so with M = 64 and all KV read from HBM, you're more than 4× short.

约掉了。这一层 N 不起作用，起作用的是 M，也就是一个 block 管多少行 Q：每字节 KV 换来 M 个 FLOPs。平衡点是 989.4 ÷ 3.35 ≈ 295，M = 64 时如果 KV 全从 HBM 读，就差了 4 倍多。

Because not every byte really comes from HBM. Many Q blocks of the same head read the same KV, so later blocks mostly hit L2. In this model, M = 64 needs an L2 hit rate of about 78%. Here it is at 80%:

因为不是每个字节都真的从 HBM 来。同一个 head 的很多 Q block 读的是同一批 KV，后来的 block 大多命中 L2。按这个模型，M = 64 时 L2 命中率要到 78% 左右。下面是 80% 的情况：

Why?

为什么？

Right. The effective arithmetic intensity falls back toward M, and the full memory-bound gap shows.

对。有效算术强度掉回 M 附近，memory-bound 的差距会完整地暴露出来。

Make M bigger by sharing the gather among more queries. For example, the query heads of one GQA group can share the same selected KV and be computed together, which multiplies M by the number of heads in the group.

让更多 query 共享同一次 gather，把 M 做大。比如同一个 GQA 组里的几个 query head 共享选出来的同一批 KV，一起算，相当于 M 乘上组内的 head 数。

Close to it: one token, or the few query heads of a GQA group. Decode attention is memory-bound no matter what.

差不多：一个 token，或者 GQA 组里的几个 query head。decode attention 怎么都是 memory-bound。

The goal changes from Tensor Core utilization to saturating HBM bandwidth. For example, split the KV along the sequence so more SMs read it in parallel (split-KV).

目标从 Tensor Core 利用率变成把 HBM 带宽跑满。比如把 KV 沿序列切开，让更多 SM 一起读（split-KV）。

Tell me about your FlashFFN paper. How did it start?

说说你那篇 FlashFFN 吧。起因是什么？

Why not just split it into heads?

直接拆成多头不就行了？

Two bottlenecks, then: exploding memory and slower compute.

两个瓶颈：显存爆炸，计算变慢。

Then use a flash kernel and never write the intermediate to HBM. Problem solved?

那用 Flash 算子，中间结果不落显存，不就解决了？

And after splitting into heads?

那拆成多头以后呢？

So each makes the other possible: heads make the flash kernel fit, and the flash kernel keeps heads from blowing up memory.

所以是互相成全：多头让 Flash 算子放得下，Flash 算子让多头不爆显存。

Why is the speedup so small?

速度为什么只快这么一点？

Do the N and M from the previous part apply to your kernel?

上一部分算的 N 和 M，在你这里还适用吗？

Can this idea carry over to MoE?

这个思路能用到 MoE 上吗？

So each head's experts are isolated.

这样每个头的 expert 是隔离的。

Sounds reasonable for the model. Have you thought about the cost in the kernel?

模型上听着合理。kernel 上你想过代价吗？

Routing moves from per token to per (token, head), so the rows each expert receives are cut into 1/16 of what they were. Matrices get smaller, M can hardly reach wgmma's 64, the main loop is short, and prologue and epilogue overhead takes a large share: exactly the kind of idle we talked about earlier.

路由的粒度从 token 变成了 (token, head)，每个 expert 分到的行数被切碎成原来的 1/16。矩阵更小，M 很难凑满 wgmma 的 64，主循环短，头尾开销占比就高，正是前面说的那种空转。

Yes, and the model-side gains have to be weighed against that cost.

对，而且模型上的收益要和这个代价放在一起算。