Streams 与 events：用依赖组织异步工作

Stream 是有序工作序列

把复制输入、计算、复制输出连续写在 host 程序里，不代表三步已在设备完成。异步接口允许函数先返回，实际工作随后推进。Stream 表达有序命令序列；同一 stream 中的工作按接口规定的顺序执行。它不是一个固定 GPU 核，也不意味着每个 stream 独占一套硬件。

如果计算读取复制产生的数据，两者之间存在依赖。这条链无论放进多少 stream 都不能被数学上消除。增加 stream 给没有依赖的工作留下重叠机会，是否重叠还取决于复制条件、设备能力、kernel 资源占用和调度。用了两个 stream，并不自动得到两倍吞吐。

跨序列必须画出一条边

设 S0 复制输入后运行 producer，S1 运行 consumer。仅在 host 上先提交 producer、再提交 consumer，不能代替独立 stream 之间的执行依赖。可以在 producer 后记录 event，让 S1 等待此 event，再提交 consumer。Event 标记某个命令位置的进展，不是自动保护所有内存的 mutex。

// HIP API顺序示意，未在SDK/设备运行。
// producer 已在 s0 排队；各资源已创建，各调用须检查错误。
hipEventRecord(produced, s0);
hipStreamWaitEvent(s1, produced, 0);
// 随后把 consumer 提交到 s1。

记录必须发生在 producer 之后。把 event 放在 producer 前面，会得到形式上存在、语义上错误的依赖：消费者等到了某个事件，但所需数据还没有产生。调试时要检查 event 所代表的具体工作位置，而不只是搜索程序里有没有 wait。

用事件图理解，不用睡眠猜测

例一设上传结束于逻辑时刻二，producer耗时三，consumer耗时一。即使S1从时刻零空闲，consumer仍应从五开始，在六结束。这些数字是教学单位，不是毫秒或设备周期。模型计算依赖约束，不能模拟真实调度。

例二有两批独立数据，每批复制两单位、计算三单位。完全串行需要十；若明确假设复制和计算能使用独立资源流水，第二次复制可在第一次计算时进行，总长八。差异来自资源与依赖假设，真实设备是否符合必须看trace，不能把模型数字套到GPU上。

Buffer 的生命随异步工作延长

Host buffer交给异步复制后马上改写，可能破坏尚未读取的数据。输出地址交给下一次工作，也需确认上一轮消费者结束。双缓冲不是仅分配两块数组：每块数组都需经历可写、生产中、可读、消费中、可复用。复用条件由事件或同步表达，不能靠固定睡眠保证。

异步host/device复制要实现预期重叠，通常还需符合API要求的host memory条件。普通分页内存可能引入额外处理或同步，应查对应版本文档。初学时先用明确同步确保正确，再缩小同步范围寻找重叠，比一开始删除所有等待更容易定位问题。

先给计时范围命名

CPU只围住异步launch，常测到提交开销。围住提交加完成等待，才可能覆盖端到端时间；event计时则要明确两个标记在哪条时间线上，是否包含依赖等待。优化前后应固定输入、预热、资源、次数和范围。回答“多stream为什么没加速”时，先查依赖和瓶颈，再讨论更多并发。

输入、边界、状态变化

// Original CPU teaching model; not a device simulator or benchmark.
#include <iostream>
#include <algorithm>
#include <stdexcept>

void require(bool ok) { if (!ok) throw std::runtime_error("model check failed"); }

int main() {
 const int producer_end=2+3;
 const int consumer_end=std::max(0,producer_end)+1;
 require(consumer_end==6);
 std::cout<<"producer_end="<<producer_end<<" consumer_end="<<consumer_end<<'\n';
}

clang++ -std=c++20 -Wall -Wextra -Wpedantic -Werror -pthread 34-a.cpp -o example && ./example

producer_end=5 consumer_end=6

// Original CPU teaching model; not a device simulator or benchmark.
#include <iostream>
#include <algorithm>
#include <stdexcept>

void require(bool ok) { if (!ok) throw std::runtime_error("model check failed"); }

int main() {
 const int first_compute=2+3, second_copy=2+2;
 const int pipelined=std::max(first_compute,second_copy)+3;
 const int serial=2*(2+3);
 require(serial==10 && pipelined==8);
 std::cout<<"serial="<<serial<<" pipelined="<<pipelined<<'\n';
}

clang++ -std=c++20 -Wall -Wextra -Wpedantic -Werror -pthread 34-b.cpp -o example && ./example

serial=10 pipelined=8