多GPU任务调度必须显式绑定设备，每个worker线程启动时须调用cudaSetDevice(gpu_id)并校验，数据分发优先用cudaMemcpyPeer配合cudaDeviceEnablePeerAccess，跨设备同步依赖cudaEventRecord与cudaStreamWaitEvent，队列需按device隔离。
多 GPU 任务调度必须显式绑定设备，cudaSetDevice 不是可选的
CUDA 默认只在当前 host 线程绑定的 device 上执行 kernel，哪怕你调用 cudaMalloc 分配了多个 GPU 的显存，不切换设备就发 kernel，所有计算都会挤在默认卡（通常是 ID 0）上。常见错误是“开了多个线程但算力没分摊”，根本原因就是漏了 cudaSetDevice。每个 worker 线程启动前必须先调用一次，且不能和别的线程交叉——CUDA context 是线程局部的，跨线程复用会触发 cudaErrorInvalidValue 或静默失败。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 用
std::thread启动 worker 时，把 GPU ID 作为参数传入，在线程入口第一行调用cudaSetDevice(gpu_id) - 避免在 worker 中反复调用
cudaSetDevice；初始化后保持绑定，直到线程退出 - 用
cudaGetLastError()紧跟cudaSetDevice检查是否成功，ID 超出范围时返回cudaErrorInvalidDevice
数据分发要绕过 host 内存瓶颈，优先用 cudaMemcpyPeer
如果每个 GPU 都从 host 拷贝自己那份输入，带宽会被 PCIe 总线锁死，尤其在 4 卡以上场景，实际吞吐常不到理论值的 1/3。正确做法是：只让一个 GPU（比如 ID 0）从 host 加载完整数据，再用 P2P 拷贝分发到其他卡。前提是开启 cudaDeviceEnablePeerAccess，否则 cudaMemcpyPeer 会返回 cudaErrorPeerAccessUnsupported。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 初始化阶段对每对 GPU 调用
cudaDeviceEnablePeerAccess（注意方向：A→B 和 B→A 是两个独立操作） - 分发时用
cudaMemcpyPeer(dst_ptr, dst_device, src_ptr, src_device, size)，src_device 和 dst_device 必须是不同 GPU ID - 禁用 unified memory（
cudaMallocManaged）做跨卡共享——它底层仍走 PCIe，且 page fault 同步开销极大
同步不能只靠 cudaDeviceSynchronize，得用 stream + event 跨设备协调
多 GPU 任务若依赖全局顺序（比如 A 卡输出是 B 卡输入），仅在每个卡上调用 cudaDeviceSynchronize 无法保证时序：A 卡 finish 不代表 B 卡已收到数据。必须用 cudaEventRecord + cudaStreamWaitEvent 构建跨设备依赖链。典型错误是直接在 device 0 的 stream 里 wait device 1 的 event，这会报 cudaErrorInvalidResourceHandle——event 必须在目标 device 上创建。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 每个 GPU 创建自己的
cudaEvent_t，并在对应 device 上 record - 用
cudaStreamWaitEvent时，stream 和 event 必须属于同一 device；跨设备等待需中转：device 0 record event → device 1 的 stream wait 这个 event（CUDA 支持跨 device event wait，但 event 必须用cudaEventDisableTiming创建） - 避免混合使用
cudaDeviceSynchronize和 event：前者阻塞整个 device，后者只阻塞指定 stream，粒度更细
任务队列设计要隔离 device context，别让 std::queue 成为性能热点
用单个 std::queue 统一派发任务给多个 GPU worker，容易因 mutex 锁争用拖慢吞吐，尤其当任务粒度小（如每次 kernel 计算 1ms）。更糟的是，如果 queue 存的是 raw pointer，而分配内存的 device 和消费它的 device 不同，cudaMemcpy 会静默退化成 host 中转。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 每个 GPU 对应一个无锁队列（如
moodycamel::ConcurrentQueue），producer 按数据所在 device 把任务推到对应队列 - 任务结构体里显式存
device_id字段，并在 worker 中校验，防止误投 - GPU 内存指针（
void*）只在对应 device 的 worker 中解引用；跨 device 传递指针前，确保已用cudaMemcpyPeer完成数据迁移
真正难的不是写通多卡，而是让数据流、事件流、控制流三者严格对齐。任何一个环节漏掉 device scope 检查，都可能在 32 卡集群上跑三天才发现结果错了一位。
