先记住一句话

先修 global-memory access 和重复 HBM traffic，再谈 occupancy；一个 100% occupancy、32 路 scattered transactions 的 kernel 仍会很慢。

correct baseline→profile traffic/stalls→coalesce + align→tile for reuse→control registers/occupancy→autotune real shapes
HBMcoalesced transactions
tile load
Shared memoryblock reuse
fragment
Registersthread reuse
compute
Outputcoalesced store
优化核心是让每个慢速字节服务更多 FLOPs，同时避免 bank conflict、spill 与同步开销。

1. Coalesced global access

warp 的 32 个 threads 发出一条 load/store instruction 时，硬件按地址覆盖的 memory segments 合并 transactions。相邻 threads 访问相邻 elements 通常最理想；stride/scatter 会请求更多 sectors、浪费带宽。

global load efficiency ≈ requested bytes / transferred bytes

2. Alignment 与 vectorized access

float2/float4 等 vector load 可减少 instructions 并提高 transaction utilization，但 pointer、stride、tensor offset 必须满足 alignment；misaligned slice 不能盲 cast。尾部仍需 masked/scalar path。

3. Shared-memory tiling

coalesced load global tile → shared memory
__syncthreads()
many threads reuse tile / reorder access
__syncthreads()
load next tile

tile 增加 reuse 和允许 transpose/reordering。现代架构可用 asynchronous copy 与 multi-stage pipeline overlap global→shared transfer 和 compute。

4. Bank conflicts

shared memory 分 banks；warp 同一 instruction 若多个不同地址映射同 bank，就被拆成多次服务（同地址 broadcast 例外）。经典 transpose tile [32][32] 按列读产生冲突，padding 成 [32][33] 可打散映射。

5. Register tiling

让每 thread 计算多个 outputs，可把 input/accumulator 留 registers，提高 reuse 与 instruction-level parallelism；但 registers/thread 上升会降低 resident warps，甚至 spill。编译器报告与 Nsight metrics 必须一起看。

6. Occupancy 的正确用法

  1. 先查 block size、registers、shared memory 谁限制 active blocks；
  2. 确认是否 latency-bound，更多 warps 是否有帮助；
  3. 尝试 launch bounds、tile/warps/stages，而非只压寄存器；
  4. 同时比较 achieved occupancy、stall reasons 与 runtime。

牺牲 reuse 来追 occupancy 往往得不偿失。

7. Warp primitives

shuffle 让 warp lanes 直接交换 registers，适合 reduction/scan，避免 shared memory + block barrier；ballot 汇总 predicates；需传正确 active mask，并理解 independent thread scheduling 后的显式同步要求。

8. Atomics 与 contention

atomic correctness 简单，但大量 threads 更新同一地址会 serialization。常见优化是 warp/block local aggregation，再做少量 global atomic；histogram 可 privatize shared bins。atomic throughput 与 dtype/architecture 相关。

9. Fusion 的边界

融合 elementwise/reduction epilogue 可省 traffic/launch；但融合巨大 expression 会提高 register pressure、compile time 与 code size。若主干是 cuBLAS GEMM，custom fused kernel 可能失去成熟 Tensor Core pipeline，整体反而慢。

10. Shape 与 autotuning

最佳 block/tile/num warps/stages 随 M/N/K、dtype、architecture 改变。生产实现通常有 heuristic/autotuner + cache；benchmark 要覆盖真实分布而非只调一个方阵。

11. 优化检查顺序

  1. correctness 与真实 baseline；
  2. timeline：launch/CPU/copy 还是 kernel？
  3. Roofline：compute or memory；
  4. coalescing/DRAM sectors/cache；
  5. reuse/shared-bank/register spill；
  6. occupancy/stalls/instruction mix；
  7. 不同 shapes、dtypes、warm/cold cache 与 end-to-end。

12. 四个 kernel 优化手算

例 1：coalescing

32-thread warp 连续读取 32 个 FP32，共 128 bytes，可由少量 aligned transactions 服务；stride=32 时地址跨约 4 KB，可能需要接近 32 个独立 sectors，requested bytes 相同但实际 traffic 大增。

例 2：tile reuse

16×16 GEMM tile，每个 K-step 从 global 读 A/B 各 256 elements，共 512；产生 16×16×16×2=8192 FLOPs，每个输入 element 在 block 内复用 16 次。

例 3：shared bank conflict

32 banks 下，warp threads 访问 shared[threadIdx×32] 时 index mod 32 全为 0，形成 32-way conflict；把 stride padding 成 33 后 bank indices 为 0…31，冲突消失。

例 4：atomic contention

1024 threads 都 atomicAdd 同一个 counter，要串行化 1024 updates。若先在 32 warps 内 reduction、每 warp 只 atomic 一次，global atomics 降到 1024/32=32 次。

常见误解：

shared memory 比 global 快，不代表多用就快。若数据只用一次，cooperative copy + barrier 是纯开销；shared 的价值来自 reuse 或把不合并访问重排成合并访问。

自测

1. padding 为什么能消除 transpose bank conflict？

它改变每一行起始地址对 bank 的模映射，使按列访问不再全部落同一 bank。

2. register spill 到哪里？

编译器所谓 local memory，物理上经 cache/device memory hierarchy，latency 远高于 registers。

3. 为什么最高 occupancy 不一定最快？

达到它可能需要更小 tile/更少 registers，降低数据复用或增加指令；且某些 kernel 在较低 occupancy 已隐藏足够 latency。

官方资料

CUDA Best Practices Guide详细说明 coalescing、shared-memory banks 与 matrix tiling；Nsight Compute用于确认 transactions、occupancy 和 stalls。