先用库，再决定是否写 Kernel

需求
├─ 高层框架已有高质量算子？         → 使用并 profile
├─ CUDA-X/CCCL/CUTLASS 已覆盖？      → 使用库，配置 stream/layout/workspace
├─ 能通过组合 + Graph 达到目标？      → 组合并捕获
├─ 需要融合/特殊布局/特殊语义？       → 自定义 kernel
└─ 需要架构极限或新指令研究？         → CUTLASS/CuTe/PTX + 专属 dispatch

成熟库通常已经处理多代架构、边界 shape、数值细节和算法选择。手写版本应先与库建立 correctness 和性能基线。

CUDA 生态地图

CUB、Thrust 与 libcu++ 统一在 CCCL（CUDA C++ Core Libraries）生态下。CUB 的 warp/block/device 分层尤其适合自定义 kernel 内部复用。

Library 集成检查表

• handle 是否绑定到正确 stream；
• 输入/输出 layout、leading dimension、stride 是否准确；
• host pointer mode 还是 device pointer mode；
• workspace 的大小、对齐、生命周期；
• 算法是否要求特定对齐、dtype、CC；
• 是否允许 TF32/降精度/非确定性算法；
• capture 到 CUDA Graph 时该 API 是否支持；
• handle 是否在线程间安全共享；
• 错误状态是否检查；
• benchmark 是否包含 heuristic/autotune 的一次性成本。

多 GPU 的三层问题

算法层：数据并行 / 张量并行 / 流水线并行 / 任务并行
通信层：P2P copy / collective / host staging / 网络
拓扑层：PCIe switch / NVLink / NVSwitch / NUMA / NIC affinity

多 GPU 不是把 device_count 乘上去。通信量、同步点、拓扑和负载均衡决定可扩展性。

显式设备管理

int count = 0;
cudaGetDeviceCount(&count);

for (int d = 0; d < count; ++d) {
  cudaSetDevice(d);
  cudaDeviceProp p{};
  cudaGetDeviceProperties(&p, d);
  // 每个 device 分别创建 stream、event、allocation、library handle。
}

CUDA 对象通常与创建它的 device/context 有关联。多 host thread 管多 GPU 时，在线程开头显式 cudaSetDevice，不要依赖之前调用遗留的 current device。

Peer-to-Peer 与拓扑

int can = 0;
cudaDeviceCanAccessPeer(&can, dst, src);
if (can) {
  cudaSetDevice(dst);
  cudaDeviceEnablePeerAccess(src, 0);
}

随后可使用 peer copy 或在受支持条件下访问 peer memory。P2P 可用不等于一定走最高速 NVLink，也不等于任意访问模式都高效；使用 nvidia-smi topo -m、NVML 和 Systems 验证实际拓扑与路径。

若 P2P 不可用，Runtime 可能需要 host staging，带宽和同步行为会不同。跨 NUMA CPU socket 的 pinned buffer 也应考虑内存亲和性。

Collective 与通信复杂度

训练中常见：

NCCL 会按拓扑和消息规模选择算法/协议。优化重点不是自己重写 ring，而是：

• 减少通信字节（分片、压缩、低精度）；
• 将大通信拆分并与计算重叠；
• 避免过细 collective；
• 正确设置 rank-device 映射与拓扑；
• 用 Systems/NCCL tracing 找关键路径；
• 保证所有 rank 以一致顺序调用 collective，避免死锁。

扩展效率

强扩展：总问题规模固定，GPU 增加；通信和串行部分最终主导。Amdahl 近似：

$$S(p)=\frac{1}{s+(1-s)/p}$$

弱扩展：每 GPU 工作量固定，总问题规模随 GPU 增长；更关注通信增长和负载均衡。报告多 GPU 性能时至少给出：单卡基线、全局/每卡 batch、吞吐、延迟、通信占比和 scaling efficiency。

生产级 C++ 资源管理

使用 RAII 封装 device memory、stream、event 和 library handle，避免异常/早退泄漏：

class CudaStream {
 public:
  CudaStream() { CUDA_CHECK(cudaStreamCreateWithFlags(&s_, cudaStreamNonBlocking)); }
  ~CudaStream() { if (s_) cudaStreamDestroy(s_); }
  CudaStream(const CudaStream&) = delete;
  CudaStream& operator=(const CudaStream&) = delete;
  CudaStream(CudaStream&& other) noexcept : s_(other.s_) { other.s_ = nullptr; }
  cudaStream_t get() const { return s_; }
 private:
  cudaStream_t s_{};
};

析构函数不能安全抛异常；异步工作与资源析构的先后必须由所有权协议保证。高频 allocation 使用内存池，长期 buffer 使用显式生命周期，避免隐式 device-wide 同步。

推荐工程结构

cuda-project/
├─ CMakeLists.txt
├─ include/
│  ├─ cuda_check.hpp
│  └─ ops.hpp
├─ src/
│  ├─ ops.cu              # dispatch + kernels
│  ├─ reference.cpp       # CPU reference
│  └─ main.cpp
├─ tests/
│  ├─ correctness.cpp
│  └─ sanitizer_cases.cpp
├─ benchmarks/
│  └─ bench_ops.cpp
└─ scripts/
   ├─ profile_systems.*
   └─ profile_compute.*

每个算子至少有：

1. 清晰的 shape/dtype/layout/aliasing 合同；
2. CPU 或可信库 reference；
3. 边界和随机 correctness test；
4. 独立 benchmark；
5. profile 入口；
6. 支持架构与 fallback；
7. 数值精度和确定性说明。

架构 Dispatch

不要在业务各处散落 if (cc >= ...)。集中注册：

op()
├─ capability / dtype / shape 检查
├─ Blackwell/Hopper specialized path（若编译且设备支持）
├─ Ampere+ async-copy path
├─ portable baseline CUDA path
└─ unsupported → 明确错误或 CPU/library fallback

编译时 __CUDA_ARCH__ 只在 device 编译阶段定义；host dispatch 应读取 device attributes。a/f 架构目标的兼容范围更窄，应保留 baseline code object 或独立二进制策略。

高级能力地图

高级特性不是学习清单上的勋章。只有当基础 profile 指向它解决的问题时再引入。

四个递进项目

项目 1：图像处理流水线

• RGB→Gray、Gaussian blur、Sobel；
• global-only 与 shared halo 对照；
• pinned memory + streams 重叠批量图像；
• CUDA Graph replay；
• 指标：吞吐、端到端延迟、有效带宽。

项目 2：并行 primitives

• reduce、scan、histogram、compact；
• 自写版对照 CUB；
• 测均匀/热点/边界输入；
• Sanitizer 与浮点误差报告。

项目 3：GEMM 与 fused epilogue

• naive → shared tile → register tile；
• 对照 cuBLAS；
• 用 cuBLASLt 实现 bias + activation；
• 解释 Roofline 位置和资源限制。

项目 4：多 GPU mini-system

• 每卡独立 stream 与 buffer；
• P2P 或 NCCL AllReduce；
• 计算/通信 overlap；
• Systems 时间线标注 NVTX；
• 报告 1/2/4/8 卡 scaling efficiency。

12 周学习计划

速查：遇到问题先去哪一章

官方参考总表

• CUDA Documentation Portal
• CUDA Programming Guide
• CUDA C++ Best Practices Guide
• CUDA Runtime API
• CUDA Driver API
• NVCC Compiler Driver
• CUDA Core Compute Libraries
• CUDA Libraries
• NCCL Documentation
• CUTLASS Documentation
• CUDA Samples