先泼一盆冷水醒醒脑：在 ROCm 上，hipMemcpyAsync这个函数名里的 Async，是最容易让人产生错觉的三个字母。不少人一看名字就以为“异步嘛，调用完马上返回，GPU自己搞定”，结果要么拿到脏数据，要么发现耗时跟同步拷贝一模一样，要么多stream开了一堆，GPU利用率还是上不去。我这两年在 AMD GPU 上调训练和推理任务，遇到的大部分“玄学性能问题”，追到最后都落在 MemcpyAsync 和任务调度这一层上。
这篇文章不打算给你堆 API 文档，我想从实际踩过的坑出发，把hipMemcpyAsync的语义、stream 的调度模型、event 的跨流同步，以及怎么让拷贝和计算真正重叠起来，完整拆一遍。适合刚接触 ROCm 的迁移用户，也适合已经在跑大模型但被“搬运速度”卡住的老手。读完你能少走很多弯路。
1. 先说 Async 这个单词，别急着高兴
1.1 hipMemcpyAsync 到底是哪里异步
先说最基本的概念。hipMemcpyAsync是 HIP 里做主机与设备之间数据搬运的核心函数，它和同步版hipMemcpy表面差别就是 Async 这个后缀，但背后的执行模型完全不一样。
同步版hipMemcpy很直白：函数返回的那一刻，数据拷贝要么已经完成，要么已经报错。调用线程在这段时间里就是等着，什么也干不了。这个语义适合小数据量、边界同步、或者程序里无所谓性能的关键路径。
hipMemcpyAsync就狡猾多了。它只保证“任务提交”这个动作是异步的，也就是说，函数调用返回时，搬运任务已经排队进了你指定的 stream，但 GPU 上那个 DMA 引擎可能还没开始搬，甚至任务还没被 GPU 前端调度器取走。真正的数据搬运发生在提交之后的某个时刻，具体什么时候完成，由 stream 里的队列顺序和硬件调度决定，不由 host 线程决定。
打个比方，同步拷贝是打电话等对方签收，异步拷贝是发快递单，你只管下单，对方什么时候收到货，取决于物流调度。很多人把这个逻辑弄反，以为异步拷贝等于“调用完数据就绪”，结果在跨 stream 场景里，kernel 跑得比拷贝还快，读到的还是一块没写完的显存，程序跑出来结果不对还不知道错在哪。
所以第一个要记住的心智模型是：hipMemcpyAsync是给 GPU 排队干活，不是让 host 线程什么都不用管。如果你要保证数据可见，必须通过 stream、event 或同步 API 来建立“拷贝已完成”的边界。这个问题不解决，后面全白谈。
1.2 内存是不是 pinned，差了一个数量级
异步拷贝能不能真正生效，第一个硬件前提是：主机端内存必须是 pinned memory，也就是常说的固定内存、锁页内存。
为什么强调这个？GPU 的 DMA 引擎只能访问物理页固定、不会被操作系统换到磁盘的内存。普通malloc得到的 pageable 内存，物理页随时可能被换走，DMA 引擎没法直接碰。要让 GPU 拷贝这种内存，运行时必须先在内部找一块固定的 staging buffer，把数据从 pageable 内存倒腾到 staging buffer，再启动 DMA。这一倒腾，问题就来了。
第一，路径变长了。一次 H2D 拷贝变成“主存到 staging”和“staging 到显存”两段，带宽开销翻倍，延迟也变高。第二，运行时为了保证 staging buffer 的安全，往往会在某个环节插入隐式同步，让你的 Async 徒有其名。实测下来，用 pageable 内存调用hipMemcpyAsync，性能经常跟同步拷贝差不多，甚至更差。
所以在热路径上，凡是高频执行的 H2D/D2H 拷贝，都建议用hipHostMalloc分配主机缓冲，或者对已有内存调用hipHostRegister把它注册成 pinned，用完了再hipHostUnregister。这俩函数本质都是告诉操作系统：这些页面锁住，不许换页。
这里有一个很多人忽略的坑：hipHostMalloc别滥用。它会把物理页锁住，锁太多会导致系统内存碎片化，甚至换页压力变大。我见过有人把模型全部参数都做成 pinned，结果机器卡成幻灯片。正确做法是只锁经常参与拷贝的传输缓冲，比如每个 batch 的输入输出、特征张量，而不是把整个数据集都塞进去。
整理成一张对照表，方便你判断该用哪个：
| 调用方式 | 返回时机 | 数据何时可用 | 适用场景 |
|---|---|---|---|
hipMemcpy | 拷贝完成后 | 返回时已就绪 | 小数据、程序边界、首次初始化 |
hipMemcpyAsync+ pinned memory | 提交后立即返回 | 需要 stream/event 同步确认 | 大块数据搬运、拷贝计算重叠 |
hipMemcpyAsync+ pageable memory | 看运行时，可能内部阻塞 | 容易出现隐式同步 | 不推荐用于热路径 |
一句话总结：想让 Async 名副其实，先检查你的主机内存是不是 pinned。这是性价比最高的一步，也是绝大多数人第一步就走错的地方。
2. Stream 是任务调度的最小单元
2.1 不要把 Stream 理解成线程
很多从线程模型迁移过来的人，第一个反应是把 stream 当成“GPU 线程”。这是个很自然的误解，但必须纠正。stream 本质是一条先进先出的任务提交队列，队列里的任务按提交顺序排队执行，不会乱序。之所以这么设计，是为了让你在同一 stream 里用天然顺序表达依赖。
举个例子，你在同一个 stream 里先提交一个hipMemcpyAsync，再提交一个 kernel，那么 GPU 会先执行拷贝，拷贝结束才启动 kernel。不需要额外加任何同步，顺序天然安全。这是 stream 最省心的用法，也是新手最容易理解的部分。
但注意，stream 不是线程，并不会因为你有两个 stream，GPU 就有两个内核同时跑。stream 只是描述了“哪一堆任务之间存在先后关系”。硬件上，命令处理器会把不同 stream 的任务分发给对应的执行引擎，比如拷贝任务给 DMA 引擎，计算任务给计算引擎。如果多个 stream 里都是计算任务，它们会共享计算引擎，以任务切片的形式交错执行，而不是严格各占一个计算核心。所以“多 stream = 多核并行”这个想法，在绝大多数情况下是不成立的。
另一个常见误区是疯狂创建 stream。有人一算，拷贝一个 stream、前向一个 stream、优化器一个 stream，再搞几个数据加载 stream，一口气建十几个。实际收益往往没有，反而把调度复杂化，还增加了事件等待和队列占用。硬件上的执行队列是有限的资源，stream 太多只会增加提交阶段的开销，甚至因为互相等待而把流水线堵死。我自己的经验是：能用一条 stream 表达顺序，就用一条；真正需要并行的是“拷贝”和“计算”这种不同类型的任务，才考虑拆开。
2.2 默认流是并发毒药
所有没指定 stream 的 HIP 调用，都会落到默认流上，也就是 stream 0。这个默认流有个非常坑的特性：它是 legacy 默认流，会跟其他所有流发生隐式同步。
具体表现是：向默认流提交任务之前，运行时需要等所有其他流的工作完成；反过来，其他流要等默认流上的任务完成之后才能继续。这意味着，只要你的代码里混用了默认流和非默认流，并发能力就会被一个看不见的全局同步点给掐断。很多人的代码表面上开了好几个 stream，结果发现 GPU 利用率和纯串行差不多，原因就在这里。
所以，真正想做多 stream 并发，我强烈建议用hipStreamCreateWithFlags显式创建非阻塞流，flag 传hipStreamNonBlocking。这样创建的流不跟默认流绑定那一套隐式同步，调度的自由度大很多。示例写法：
hipStream_t stream; hipStreamCreateWithFlags(&stream, hipStreamNonBlocking);还有一件事要提醒：hipDeviceSynchronize少用。它会把整个设备的所有 stream 全部同步一次，虽然是万金油，但每次调用都相当于全局屏障，把好不容易做出来的重叠全部打散。只在真正需要“所有工作完成”的边界上调用它，其他情况下优先用更细粒度的hipStreamSynchronize或hipEventSynchronize。
2.3 Event 是跨流同步的钥匙
不同 stream 之间要表达“谁等谁”，靠的是 event。event 像一个里程碑，记录某个 stream 的执行进度。你可以在 stream A 里插一个 event，然后让 stream B 等待这个 event，stream B 后续的任务就会在 event 到达之后才继续。
典型调用链路是四步：
hipEventCreate创建一个事件对象；hipEventRecord(event, streamA)把事件排进 streamA 的队列；hipStreamWaitEvent(streamB, event, 0)让 streamB 等待 event；hipEventSynchronize(event)让 host 线程等待这个里程碑。
这套机制是 MemcpyAsync 场景里最核心的粘合剂。跨 stream 时，A stream 的拷贝还没完成，B stream 的 kernel 就不能启动；一旦 event 到达，B stream 自动继续，中间不需要 host 介入。
我在实际代码里的习惯是：只要两个 stream 之间有数据依赖，就显式画一条 event 边。宁可多用几个 event，也不要靠“时间差不多”来赌。GPU 不会像 CPU 那样自动做数据依赖检查，它只认队列顺序和事件关系。你把依赖关系忘了，它就真的不管，该并行并行，该出错出错。
event 还有个很实用的附加价值：计时。GPU 侧操作如果用std::chrono测，测到的往往只是 host 提交耗时，不是设备执行耗时。用事件测才贴近真实：
hipEvent_t start, stop; hipEventCreate(&start); hipEventCreate(&stop); hipEventRecord(start, stream); // ... 提交任务 ... hipEventRecord(stop, stream); hipEventSynchronize(stop); float ms = 0.f; hipEventElapsedTime(&ms, start, stop);这个时间戳是 GPU 端记录的，比 host 侧计时靠谱得多。后面做重叠优化，基本都靠它来判断效果。
3. 从一块 GPU 内部看任务到底怎么跑
3.1 提交、门铃和队列
要理解任务调度，不能光看 API 层。Host 调用hipMemcpyAsync之后，实际发生的事可以简化为三步。
第一步，runtime 把这次拷贝包装成一个任务描述，写到对应 stream 的提交队列里。这个队列在内存里是一块环形缓冲，由 runtime 维护。第二步，写完队列之后，host 通过门铃机制通知 GPU：有新任务了。第三步，GPU 端命令处理器看到消息，把任务取出来，按任务类型分发给对应的执行引擎，比如拷贝走 DMA 引擎，kernel 走计算引擎。
所以你看，API 调用返回快，本质是“提交快”，不是“执行快”。任务可能还在队列里等着，连 GPU 前端都还没拿到。这个模型能解释很多诡异现象，比如为什么调用完立刻查数据还是旧的，为什么 event 还没到 host 就不该读 buffer。
还有一点很多人想不到：提交队列不是无限深的。如果 host 疯狂提交任务，GPU 消费不过来，队列满了之后，API 会在提交阶段等待队列空出，这时候异步调用照样阻塞。所以“Async 永远不阻塞”是错的。压力测试中如果发现某个 MemcpyAsync 调用耗时突然飙高，先怀疑是不是任务提交速度超过了 GPU 消费速度，队列水位满了。
从这个模型还能推出一个重要结论：异步拷贝的源缓冲区，在拷贝完成之前千万不能改。你把任务提交了，host 又回头把源数据改了，DMA 读到一半就会是新旧混合的数据，结果完全不可控。这个坑我在 4.2 节会再强调一遍。
3.2 为什么拷贝和计算可以同时发生
很多人觉得，同一个 GPU 上一会儿拷贝一会儿计算，怎么可能并行？但硬件设计恰恰是支持并行的。GPU 内部有独立的 DMA 拷贝引擎和计算引擎，hipMemcpyAsync走 DMA 引擎，kernel 走计算引擎，两者本身是可以同时干活的。这也是“拷贝和计算重叠”能带来收益的硬件基础。
不过，这里有个隐藏瓶颈：两个引擎共享显存带宽、内存控制器和 L2 缓存。如果一块拷贝任务已经把带宽吃满，计算 kernel 需要的显存访问就会被排队拖慢，整体耗时可能不是 max(拷贝, 计算)，而是接近两者之和。所以“重叠必然加速”是错的。
我的经验判断是这样的：大块 H2D 拷贝搭配计算密集型的 kernel，重叠收益最明显。因为计算密集型 kernel 对显存带宽的压力相对小，两者可以在不同维度同时推进。反过来，小拷贝搭配访存密集型的 kernel，重叠意义不大，还引入 stream 管理和 event 同步的开销，不如老老实实串行。
补一句关于“多个 stream 并行”的正确理解：GPU 前端调度器可以同时处理多个 stream，但执行引擎是共享的。多个 stream 里都是 kernel，它们在计算引擎上是交错执行，不是严格各占一块。调度的优先级、切分粒度由硬件决定，代码层控制不了太多。你唯一能做的是：把依赖关系理清楚，把隐式同步拆掉，剩下的交给调度器。
3.3 版本和环境带来的调度差异
现在要说一件很现实的事：同样的 HIP 代码，在 ROCm 不同版本、不同发行版、不同芯片上的表现，可能差别很大。
版本差异会直接影响hipMemcpyAsync是不是真的“异步”。比如某些旧版本对 pageable 内存的 staging 处理更激进，H2D 异步很容易退化成同步；新版本改了 staging 路径，行为又变了。更麻烦的是默认流的语义，不同版本对 legacy 默认流跟其他流的隐式同步执行力度不太一样。如果从 CUDA 迁到 ROCm，千万别假设默认流行为和 CUDA 完全对齐，关键代码里显式创建流和事件是底线。
新发行版的问题更隐蔽。像 Debian 13 这类系统，glibc 和系统库版本都很新，而 ROCm 官方包有时候跟不上，装好后容易出现运行时和系统依赖不匹配。表现出来的症状很暧昧：简单 tensor 搬运能跑，一旦涉及多 stream 并发或者大批量 MemcpyAsync，就开始卡顿、报错、性能倒挂。最近社区里经常看到“gx1031 rocm 哪个版本支持 pytorch”这类问题，背后往往不是 PyTorch 本身，而是运行时版本和芯片、发行版的匹配问题。
我的排查建议是：换版本、换系统、换芯片之后，别急着直接上大框架。先用一个极简 HIP 程序验证 MemcpyAsync 和多 stream 调度链路通不通，再跑 PyTorch。基础链路如果都不稳，换什么层都是把问题换一种形式露出来。确认版本用rocminfo、hipcc --version，再看芯片型号，最后才回去查业务代码。
4. 实战：把拷贝和计算真正重叠起来
4.1 一份可以直接改的代码骨架
理论说完了，上实操。下面这份代码展示一个最典型的重叠模型：一个 stream 做 H2D 拷贝，另一个 stream 等拷贝完成之后做计算，两个操作在引擎和调度层面尽量拉开距离。
#include <hip/hip_runtime.h> #include <cstdio> #include <numeric> __global__ void scale_kernel(float* data, float scale, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) data[i] *= scale; } int main() { const int n = 64 * 1024 * 1024; const size_t bytes = n * sizeof(float); // 1. 关键：主机侧必须用 pinned memory，热路径不能省 float* h_src = nullptr; hipHostMalloc(&h_src, bytes, hipHostMallocDefault); for (int i = 0; i < n; ++i) h_src[i] = 1.0f; float* d_data = nullptr; hipMalloc(&d_data, bytes); // 2. 创建非阻塞流，避免默认流的隐式同步 hipStream_t copy_stream, compute_stream; hipStreamCreateWithFlags(©_stream, hipStreamNonBlocking); hipStreamCreateWithFlags(&compute_stream, hipStreamNonBlocking); // 3. 事件用于跨流同步 hipEvent_t copy_done; hipEventCreate(©_done); // 4. 拷贝任务提交到 copy_stream hipMemcpyAsync(d_data, h_src, bytes, hipMemcpyHostToDevice, copy_stream); // 5. 在 copy_stream 上记录事件 hipEventRecord(copy_done, copy_stream); // 6. compute_stream 等待拷贝完成事件 hipStreamWaitEvent(compute_stream, copy_done, 0); // 7. 计算 kernel 提交到 compute_stream hipLaunchKernelGGL(scale_kernel, dim3((n + 255) / 256), dim3(256), 0, compute_stream, d_data, 0.5f, n); // 8. 只需要同步 compute_stream，拷贝依赖的完成状态已经包含在内 hipStreamSynchronize(compute_stream); // 9. 结果拷回 host 验证，注意用 pinned 接收缓冲 float* h_check = nullptr; hipHostMalloc(&h_check, bytes); hipMemcpy(h_check, d_data, bytes, hipMemcpyDeviceToHost); printf("h_check[0]=%f\n", h_check[0]); hipHostFree(h_src); hipHostFree(h_check); hipFree(d_data); hipEventDestroy(copy_done); hipStreamDestroy(copy_stream); hipStreamDestroy(compute_stream); return 0; }这段代码的骨架可以套用到绝大多数实际场景。你把copy_stream理解成数据搬运专线，compute_stream理解成计算专线，它们之间用事件搭了一座桥。拷贝没完成，计算这边绝不越线。
4.2 实测：重叠前后的差别
用上面的代码跑一个参考测试，我拿到的典型数据大概是这样的：
| 写法 | 参考耗时 | 说明 |
|---|---|---|
| 同步拷贝，再执行 kernel | 约 3.5 ms | 两段完全串行 |
| MemcpyAsync + event + kernel 分两流 | 约 2.6 ms | 拷贝与计算有一定重叠 |
| MemcpyAsync 但 host 内存非 pinned | 约 3.4 ms | staging 拷贝 + 隐式同步，几乎没收益 |
| 跨流但漏掉 event 依赖 | 结果错误 | 依赖关系缺失，运行顺序不可控 |
注意，这个表不是让你背数字，而是让你建立正确预期。重叠能不能达到理想中的 max(copy, compute)，取决于 kernel 的访存特征和拷贝大小。我上面已经说过，如果 kernel 访存很重，总耗时可能更接近两者之和而不是 max，因为带宽被抢了。
还有个特别容易踩的坑，我要单独拎出来说。异步拷贝提交之后，host 侧如果立刻往h_src里写新数据，DMA 引擎还没搬完，数据就会新旧混合。比如你在一个循环里重复用同一个 pinned buffer，做完一次拷贝就马上填充下一批数据，忘了等事件完成，结果就是模型训练数据偶尔出错，还特别难复现。解决的办法也不复杂：要么拷贝完成前不碰源 buffer，要么准备两块 pinned buffer 交替使用，让上一轮拷贝和下一轮数据填充错开。
4.3 从 PyTorch 看 MemcpyAsync 调度
如果你不是直接写 HIP，而是用 PyTorch 跑模型，同样会遇到这一层调度问题，只是被封装在框架后面。
PyTorch 的 ROCm 后端在很多地方会调用hipMemcpyAsync。最常见的两个入口：一个是tensor.to(device, non_blocking=True)，另一个是DataLoader的pin_memory=True。前一个对应 H2D 搬运，后一个是在 CPU 后台把样本放进 pinned buffer，再用异步拷贝把数据搬到 GPU 上，和当前 step 的计算重叠。
很多人说“我开了 non_blocking 怎么还是慢”，这时候先检查两件事：第一，源张量是不是在 CPU 内存上，且是否来自 pinned buffer；第二，目标 stream 是不是和当前计算流在同一个依赖链上。如果 CPU 侧是普通 pageable 内存，non_blocking=True很多时候不会真正异步，因为底层还是走了 staging 路径。如果 DataLoader 不开pin_memory=True，H2D 搬运大概率也是同步的，GPU 在前几个 step 会频繁出现利用率回落。
多卡场景还要注意 D2D 拷贝。tensor.to('cuda:1')这种跨卡搬运，底层可能走hipMemcpyPeerToPeerAsync或者事件同步机制，发送流和接收流的关系如果理不清，就会出现莫名其妙的卡顿或者数据滞后。我的建议是：多卡同步、广播、梯度规约这些操作尽量依赖框架自带的 DDP 或通讯封装，不要自己手写跨卡拷贝和 event 链，水太深。
回到热词里那个“gx1031 rocm 哪个版本支持 pytorch”的问题。这类问题我见的太多了，它本质上是版本配对问题，不只是 PyTorch 和 ROCm 版本号对上那么简单，还牵涉运行时调度行为是否正常。装上之后先跑一个简单的张量搬运加多 stream 重叠 demo，如果基础链路都是通的，再上大模型。Python 里调半天参数，最后发现是 HIP 运行时调度有问题，那才是真的浪费时间。
5. 常见故障排查实录
5.1 症状对照表与排查顺序
预告了这么多，最后把常见故障和排查顺序整理成一张速查表：
| 症状 | 最常见原因 | 排查步骤 |
|---|---|---|
| 拷贝后不等待，kernel 结果偶尔错 | 跨 stream 依赖缺失，或用了 pageable 内存 | 检查是否 pinned，确认 event 依赖是否建立 |
hipMemcpyAsync耗时和同步版一样 | 主机内存不是 pinned，运行时走了 staging | 换成hipHostMalloc再测 |
| 多 stream 开了，性能反而下降 | 默认流隐式同步、stream 过多、频繁hipDeviceSynchronize | 用hipStreamNonBlocking，减少 stream 数，收敛同步点 |
| 主机改源 buffer 后数据变脏 | 异步拷贝未完成就写入源内存 | 等 event 后再复用 buffer，或双缓冲切换 |
| 错误在很远的地方才暴露 | 异步 API 错误被延迟到同步点返回 | 每个逻辑阶段用hipGetLastError检查一次 |
| 新发行版或新芯片上 PyTorch 搬运卡顿 | ROCm 运行时和系统依赖不匹配 | 用极简 HIP demo 验证运行时链路，再查框架层 |
排查顺序我一般固定成五步，不乱跳：
- 先确认版本。运行时、驱动、PyTorch wheel 三者是否匹配，这是性价比最高的一步。
- 跑最小 HIP 案例。看
hipMemcpyAsync的异步和重叠是否正常。 - 检查内存类型。普通
malloc分配的内存别指望真正异步。 - 梳理 stream 依赖。把每个拷贝、kernel、event 画成依赖链，看有没有环，有没有漏边。
- 找隐式同步。在代码里全局搜
hipDeviceSynchronize、hipMemcpy、默认流调用，看是不是它们在偷偷打断并行。
5.2 代码里应该养成的四个习惯
排查经验久了，我慢慢把一些原则固化成了自己的编码习惯，这里分享给你参考。
第一，显式创建 stream，并且明确指定hipStreamNonBlocking。不要让默认流的隐式同步悄悄决定你的性能。代码里哪怕只有一个 stream，也值得显式创建，至少避免以后扩展时踩坑。
第二，一个 stream 内放无依赖操作，跨 stream 依赖全部通过 event 表达。也就是说，同一个 stream 内只要顺序对就行；跨 stream 必须让依赖关系在 code 里看得见。
第三，把同步收敛到少数几个点。逻辑阶段结束时同步一次，不要在循环体里频繁 sync。频繁同步等于反复拆墙，重叠再好也白搭。
第四，关键路径上不要复用 pageable 内存做拷贝。分配一次 pinned buffer 反复使用，或者用双缓冲切换。分配和释放hipHostMalloc的开销不小，频繁调用来回折腾，性能很难看。
这四条看着简单，真能严格执行，能避开 90% 的 MemcpyAsync 调度问题。
5.3 工具建议与一点体会
排查这类问题，工具不需要多复杂。我最常用的三个：第一个是rocprof或omnitrace，看任务时间线里 kernel 和拷贝引擎的占用情况；第二个是rocm-smi，看显存带宽和引擎利用率，辅助判断是不是带宽撞顶；第三个也是最笨但最有效的——event 打点。在关键位置用 event 记录时间，比任何花哨工具都直观，能快速定位到底卡在拷贝、kernel 还是同步等待上。
如果是换了新版本或者新系统，我建议在正式压测之前，先跑一遍 4.1 里的代码骨架，确认三件事：MemcpyAsync 确实异步、两个 stream 能重叠、数据依赖没有错乱。三件事全过，再上大框架。这套流程我帮人排查过很多次，每次都能快速缩小范围，省下大量对着编译日志干瞪眼的时间。
我自己后来把这事变成了固定 checklist：先看内存，再看 stream，再看 event，最后才怀疑编译器和驱动。大部分 MemcpyAsync 的妖问题，都倒在前两步。Async 这个名字很容易让人觉得自己“天生快人一步”，但实际上它只是把决策权还给了你——什么时候完成、什么时候等待、要不要重叠，都需要你明确告诉 GPU。把这条链想清楚，MemcpyAsync 在调度里就是最好用的一根橡皮筋。
