cudaStream_t 本身不自动隐藏通信，需 CPU 主动发起非阻塞操作并合理同步；关键前提是使用 cudaMallocHost 分配 pinned memory，否则 cudaMemcpyAsync 退化为同步；多 GPU 下 stream 必须与 device 绑定，且 pinned memory 需 cudaFreeHost 配对释放。
cudaStream_t 如何与 CPU 线程协同实现重叠
关键不是“能不能”，而是“怎么让 CPU 不卡 GPU”。cudaStream_t 本身不自动隐藏通信，它只是提供异步执行的上下文；真正起作用的是 CPU 主动发起非阻塞操作 + 合理插入同步点。常见错误是调用 cudaMemcpy 时没指定流（默认 0），结果所有传输串行排队，GPU 空转等数据。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 所有
cudaMemcpyAsync必须传入显式创建的cudaStream_t，且该流需与后续 kernel 启动使用同一 stream - CPU 端用
std::thread或std::async预加载下一批数据到 pinned memory（用cudaMallocHost分配），避免cudaMemcpyAsync内部隐式拷贝 - 避免在 stream 中混用同步 API（如
cudaMemcpy、cudaDeviceSynchronize），它们会阻塞整个流，破坏重叠 - 用
cudaEventRecord+cudaEventQuery替代轮询cudaStreamQuery，减少 CPU 开销
为什么 pinned memory 是通信隐藏的前提
没有 pinned memory，cudaMemcpyAsync 实际退化为同步行为——驱动会在后台偷偷做 host-to-host 拷贝再发 DMA，完全失去异步性。这不是 bug，是设计限制。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 用
cudaMallocHost(&h_src, size)分配 host 端内存，而非new或malloc - 若需多线程写入同一 pinned buffer，加锁粒度要细（例如按 buffer 分片加锁），避免 CPU 成为瓶颈
- pinned memory 有限（通常几百 MB），不能无节制分配；可用
cudaMemGetInfo监控剩余量 - Linux 下需确保用户有
lock memory权限（ulimit -l unlimited），否则cudaMallocHost返回cudaErrorMemoryAllocation
如何验证重叠是否真实发生
靠肉眼观察 nvvp 或 nsys 时间线不够——CPU 和 GPU 时间轴不同步，容易误判。必须用事件打点测实际重叠宽度。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 在 CPU 准备数据完成时记录 event A，在 kernel launch 后记录 event B，两者时间差就是 CPU 等待 GPU 的空闲窗口
- 用
cudaEventElapsedTime测 A→B，若接近 0 表明 CPU 已经“追上”GPU，重叠充分；若 > kernel 执行时间，说明 CPU 落后太多 - 注意：
cudaEventRecord本身有微秒级开销，高频打点会污染结果，单次 pipeline 建议只打 2–3 对事件 - 对比基线：关掉 pinned memory / 全用默认 stream，看
nsys profile --trace=cuda,nvtx中 memcpy 和 kernel 是否出现明显 gap
多 GPU 场景下 stream 与 device 切换的坑
cudaStream_t 是 device-local 的。在 multi-GPU 异构系统中，cudaSetDevice 必须在 stream 创建和使用前调用，且不能跨 device 复用 stream 句柄——否则行为未定义，常见表现为 kernel 不执行或 cudaErrorInvalidValue。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 每个 GPU 对应独立的
std::vector<cudastream_t></cudastream_t>，配合cudaSetDevice(i)初始化 - CPU 数据分发逻辑需显式绑定 device ID，例如用
std::vector<void> d_buffers</void>存每卡的 device pointer - 避免在 stream callback（
cudaLaunchHostFunc）里调用cudaSetDevice，callback 运行在 driver 内部线程，device 上下文可能错乱 - 统一用
cudaGetLastError()检查每个 stream 操作，别依赖最后一条报错——错误状态是 per-device 的
delete 释放，必须配对调用 cudaFreeHost；漏掉会导致内存泄漏且不报错，只在下次 cudaMallocHost 时失败。 
