CUDA 算子性能优化：从访存合并到占用率的实战路径
CUDA 优化最容易踩的坑是直接套“高占用率”或“多用共享内存”这样的口诀。正确顺序应该是先测量瓶颈，再改变一个变量，并用基准测试确认收益。
先判断是带宽受限还是计算受限
向量加法、类型转换等算子每个元素只做少量运算，通常受显存带宽限制；矩阵乘、卷积和复杂归一化则可能受 Tensor Core 或 FP32 吞吐限制。用有效带宽衡量前者：
有效带宽 = 读写总字节数 / Kernel 时间如果有效带宽已经接近设备实测峰值，继续减少几条指令不会有明显收益；应转向减少数据搬运或做算子融合。
让全局访存合并
相邻线程访问相邻地址，才能合并成少量内存事务。对行主序张量，连续维通常映射到 threadIdx.x。跨步访问、按列读取或不规则 gather 会降低吞吐。必要时先转置、重排布局，或者为常见 shape 提供专门 Kernel。
共享内存与 Bank Conflict
共享内存适合缓存一个 Tile 内会重复使用的数据，例如矩阵乘中的 A、B 子块。它不是“更快的显存替代品”：加载一次、复用多次才值得。多个线程访问同一 Bank 的不同地址会产生 bank conflict，可以通过 padding 或改变访问顺序缓解。
寄存器、占用率和 Block 大小
寄存器用量过高会降低一个 SM 同时驻留的 Warp 数，甚至溢出到 local memory。占用率只是隐藏延迟的手段，不是目标本身；一个使用较少 Warp、但每个线程计算复用充分的 Kernel 可能更快。用 Nsight Compute 观察 register/thread、achieved occupancy、memory throughput 和 eligible warps per cycle，再调 128/256/512 等 Block 配置。
融合通常比微调更值
bias + activation、scale + mask 等连续算子如果各自读写显存，会产生多次带宽开销。将它们融合到一个 Kernel，往往比单独优化每个 Kernel 更有效。但融合会增加寄存器压力和编译变体，应该只针对真实工作负载的 shape 做基准。
可复现的优化循环
固定输入和 warmup 次数，报告 p50/p95 延迟；每次只改一处；同时检查数值误差和不同 shape。性能优化的终点不是某个指标最大，而是在精度、吞吐、显存和维护成本之间找到稳定的 Pareto 点。
