。
如果直接用 blockIdx.x 当 tile 号：GPU 不保证 block 按 blockIdx 顺序调度，完全可能 block 100 已经在跑、block 5 还没被调度（SM 满了）。block 100 会一直自旋等 block 5，而 block 5 要等 block 100 退出才有 SM 可用 → 死锁。
用 atomicAdd(&counter, 1) 动态领号后，tile 号的顺序 == block 实际开始执行的顺序，所以任何 tile t 等待的 tile < t 一定都已经启动了，只是可能没算完——这是可以等到的。
另外还依赖 Volta 之后的独立线程调度：一个线程自旋时，同 warp 的其他线程能继续前进到 barrier。这道题能答明白，基本就能证明你理解 GPU 的执行模型。

`},{q:`scan 在推理里到底哪儿用得上？`,a:`

非常多，而且都在关键路径上：
• MoE 的 token 路由：算出每个 expert 收到多少 token、每个 token 该写到哪个位置 → 直方图 + exclusive scan + scatter；
• 变长 batch（continuous batching）：从每个序列的长度算出 cu_seqlens（累积长度），FlashAttention 的变长接口就吃这个；
• radix sort / TopK：每一位的 scatter 位置来自 scan；
• 稀疏 / 结构化 attention：构建非零块的索引；
• 投机解码：接受的 token 数是前缀性质，用 scan 算出各序列的写入偏移。
所以它虽然不像 GEMM 那样直接占时间，却是调度类 kernel 的通用零件。

`},{q:`__shfl_up_sync 里那个 if (lane >= off) 不写会怎样？`,a:`

会算错。__shfl_up_sync(mask, v, off) 对于 lane < off 的线程没有合法的源 lane，CUDA 规定此时返回该线程自己的值（而不是 0）。如果无条件相加，前 off 个 lane 就会把自己的值加了两遍。
所以固定写法是 float t = __shfl_up_sync(...); if (lane >= off) v += t;——注意 shfl 本身必须在 if 外面执行（所有 lane 都要参与），只有加法在 if 里面。
这个