高性能算子

CUDA 与 GPU 架构

1. CUDA 的 grid、block、warp、thread 如何组织和调度？
2. GPU 的寄存器、Shared Memory、L1/L2 Cache、Global Memory、HBM 各有什么特点？
3. CUDA Core、Tensor Core、MMA/WMMA 分别适合什么计算？
4. SIMT、Warp Divergence、Warp Shuffle 的原理和性能影响是什么？
5. Occupancy 是什么？受哪些资源因素影响？Occupancy 越高是否一定越快？
6. Global Memory 合并访问的条件是什么？如何判断访存是否对齐？
7. Shared Memory Bank Conflict 是什么？常见规避方法有哪些？
8. 为什么 float4 等向量化加载可能提升性能？有哪些对齐前提？
9. H2D、D2H、Kernel 之间如何实现异步和重叠？
10. 一个 Block 是否可能在多个 SM 间迁移执行？
11. CUDA stream、event、同步和 memory fence 分别解决什么问题？
12. 如何用 Nsight Systems、Nsight Compute、PyTorch Profiler 定位 GPU 瓶颈？

算子优化与 Kernel

1. 如何判断算子是 Compute-bound、Memory-bound、Launch-bound 还是同步受限？
2. Roofline 模型中的算术强度如何计算和使用？
3. GEMM 的 tiling、寄存器复用、Shared Memory、Warp-level 分块如何设计？
4. 为什么沿 K 维度分块能提高 GEMM 的计算访存比？
5. 双缓冲、cp.async、TMA 和流水线如何隐藏访存延迟？
6. 算子融合为什么能加速？什么情况下融合反而变慢？
7. Online Softmax 如何保证数值稳定？
8. FlashAttention V1 与 V2 的核心改进分别是什么？
9. Reduce 算子如何做线程内、Warp 内和 Block 内分层归约？
10. 什么时候可以使用 atomicAdd？如何降低原子操作冲突？
11. 多 shape、不同 dtype 和不同硬件如何做 Kernel dispatch？
12. 如何做算子正确性、数值误差和性能回归？