cudaMemcpyAsync必须配非默认stream才能真正异步，因默认stream会隐式同步整个设备；且host内存必须页锁定（如cudaMallocHost），否则退化为同步拷贝。
cudaMemcpyAsync 为什么必须配 stream？
因为 cudaMemcpyAsync 本身不保证异步——它只是“允许”异步，真正触发并发执行的前提是绑定非默认 stream。用默认 stream（即 0）调用时，行为等价于同步拷贝，还会隐式同步整个 device，彻底废掉流水线。实际中常见错误是只改函数名没配 stream，结果 GPU 利用率卡在 30% 以下。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 显式创建 stream：
cudaStream_t stream; cudaStreamCreate(&stream); - 所有配套操作（kernel 启动、
cudaMemcpyAsync）都传同一个 stream 句柄 - 避免跨 stream 依赖未加同步；若需顺序保障，用
cudaStreamSynchronize(stream)或事件（cudaEventRecord+cudaStreamWaitEvent）
host 端内存必须页锁定（pinned）才能异步？
是的。cudaMemcpyAsync 对 host 内存有硬性要求：只有页锁定内存（pinned memory）才能被 DMA 引擎直接访问。普通 malloc 或 new 分配的内存会触发隐式同步，退化为同步拷贝，且可能引发 cudaErrorInvalidValue 错误。
实操建议：
立即学习“C++免费学习笔记（深入）”；
- 用
cudaMallocHost(&h_ptr, size)分配 host 内存（自动页锁定 + write-combined 优化） - 或用
cudaHostAlloc(&h_ptr, size, cudaHostAllocWriteCombined)控制属性 - 释放必须用
cudaFreeHost(h_ptr)，不能用free - 注意：页锁定内存占用物理 RAM，不宜长期大量分配
拷贝方向选错会导致 silent failure？
不会报错，但数据写到错误位置，后续 kernel 读到脏数据或全零。CUDA 不校验地址合法性，cudaMemcpyAsync 的 kind 参数一旦设错（比如把 cudaMemcpyHostToDevice 写成 cudaMemcpyDeviceToHost），GPU 就往 host 指针里瞎写，极难调试。
关键点：
- 方向由第 4 个参数决定，必须与指针类型严格匹配：
h_ptr是 host 地址，d_ptr是 device 地址 - 常见组合只有三种：
cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice - 别依赖宏定义缩写（如
cudaMemcpyHTD），手敲全称更安全
为什么 memcpy 异步了，整体耗时却没降？
大概率是 host 端计算没跟上，或者 kernel 启动/拷贝没重叠。异步拷贝只是移除了 CPU 等待时间，不代表 GPU 工作变快；若 host 侧处理太慢（比如预处理逻辑卡在 CPU），device 始终空转，吞吐还是上不去。
排查方向：
- 用
nvprof或 Nsight Compute 看 timeline：确认 H2D / kernel / D2H 是否有重叠 - 检查 host 计算是否阻塞在锁、IO、低效算法上
- 考虑用多个 stream 轮询喂数据（double buffering），例如 stream A 拷贝 batch 1 时，stream B 启动 batch 0 的 kernel
真正高效的异步拷贝，从来不是单靠一个 API 调用，而是 host 计算节奏、内存布局、stream 划分三者咬合的结果。最容易被忽略的是 pinned 内存的生命周期管理——分配一次、反复复用，比每次 cudaMallocHost/cudaFreeHost 更关键。
