2. 流与异步API核心机制

在 CUDA 编程中，流（Stream）是异步执行的基本单位。一个流本质上是一条按顺序执行的操作队列：提交到同一流中的核函数、内存拷贝等操作，会严格按照提交顺序依次执行；而不同流之间的操作则可以并发或交错执行。理解 cudaStream_t 的创建与销毁、cudaMemcpyAsync 的异步语义、核函数启动的异步行为，以及默认流与显式流的差异，是掌握多流并发优化的第一步。

一、cudaStream_t 的创建与销毁

cudaStream_t 是一个不透明句柄，代表一条流。创建流使用 cudaStreamCreate，销毁使用 cudaStreamDestroy。基本用法如下：

cudaStream_t stream;
cudaError_t err = cudaStreamCreate(&stream);
if (err != cudaSuccess) {
    // 处理创建失败
}
// ... 在该流上提交操作 ...
cudaStreamDestroy(stream);   // 销毁流

需要注意几个关键点。第一，cudaStreamCreate 创建的是阻塞流（blocking stream），它与传统默认流之间存在隐式同步关系；若希望流与默认流完全独立，应使用 cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) 创建非阻塞流。第二，cudaStreamDestroy 是异步的：它会立即返回，但流中尚未完成的工作仍会继续执行，直到全部完成才真正释放资源。因此，销毁流之前通常需要先同步，例如 cudaStreamSynchronize(stream) 或 cudaDeviceSynchronize()。第三，流句柄在销毁后不应再被使用，否则行为未定义。

二、cudaMemcpyAsync 的异步语义

cudaMemcpyAsync 是异步拷贝接口，其函数原型为：

cudaError_t cudaMemcpyAsync(void* dst, const void* src,
                            size_t count,
                            cudaMemcpyKind kind,
                            cudaStream_t stream = 0);

与同步的 cudaMemcpy 不同，cudaMemcpyAsync 在提交拷贝请求后立即返回，实际的数据传输在指定流中异步进行。它有几个重要约束：

• 主机内存必须是页锁定内存（pinned memory）。如果传入的是普通可分页内存，cudaMemcpyAsync 可能退化为同步行为，甚至返回错误。因此，异步拷贝前通常先用 cudaMallocHost 或 cudaHostAlloc 分配页锁定内存。
• 拷贝方向由 kind 指定，如 cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice。
• 拷贝与同流中的核函数按序执行。例如先提交 H2D 拷贝，再提交核函数，核函数会等待拷贝完成后再开始。

一个典型的多流流水线写法如下：

float *h_a, *d_a;
cudaMallocHost(&h_a, N * sizeof(float));   // 页锁定内存
cudaMalloc(&d_a, N * sizeof(float));

cudaStream_t s;
cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);

cudaMemcpyAsync(d_a, h_a, N * sizeof(float),
                cudaMemcpyHostToDevice, s);
kernel<<<grid, block, 0, s>>>(d_a, N);
cudaMemcpyAsync(h_a, d_a, N * sizeof(float),
                cudaMemcpyDeviceToHost, s);

cudaStreamSynchronize(s);   // 等待该流全部完成
cudaStreamDestroy(s);

三、核函数启动的异步语义

核函数启动（<<<grid, block, sharedMem, stream>>>）本身是异步的：主机线程提交启动命令后立即继续执行后续代码，不会等待核函数在设备上执行完毕。核函数的执行由设备端调度，主机与设备通过流来协调顺序。

这意味着，如果主机在核函数启动后立刻读取设备内存结果，可能读到尚未完成的数据。要获得正确结果，必须显式同步，常用方式有：

• cudaDeviceSynchronize()：同步整个设备，等待所有流完成。
• cudaStreamSynchronize(stream)：只等待指定流完成。
• cudaEventSynchronize(event)：等待某个事件，适合精细控制。

此外，核函数启动的错误（如参数非法、资源不足）不会立即在主机端抛出，而是延迟到下一次同步或错误检查时才报告。因此，在调试阶段建议在关键启动后调用 cudaGetLastError() 或 cudaPeekAtLastError() 检查。

四、默认流与显式流的行为差异

默认流（也称空流，stream 0）具有特殊语义。在传统默认流模式下，默认流中的操作与其它阻塞流之间存在隐式同步：默认流会等待之前所有阻塞流中的操作完成，之后提交到阻塞流的操作也会等待默认流完成。这种隐式同步会破坏并发性，因此多流优化中通常使用非阻塞流来避免。

显式流（用户创建的流）则没有这种隐式同步关系，不同显式流之间可以真正并发执行。但要注意，同一流内仍然严格按序执行。下表对比了默认流与显式流的关键差异：

下面这张图展示了多流场景下，主机提交操作、流内顺序执行以及流间并发的关系。图中以两条非阻塞流为例，分别执行“H2D 拷贝 → 核函数 → D2H 拷贝”的流水线，两条流之间可以重叠执行。

五、实践中的关键判断

在实际工程中，判断是否真正实现了异步并发，可以从以下几点入手。第一，确认主机内存是否为页锁定内存，否则 cudaMemcpyAsync 可能退化为同步。第二，确认流是否为非阻塞流，避免默认流带来的隐式同步。第三，使用 CUDA 事件（cudaEventRecord / cudaEventElapsedTime）测量各流的时间线，观察传输与计算是否重叠。第四，注意资源竞争：如果多个流同时提交大量核函数，可能因 SM 占用、寄存器或共享内存不足而无法真正并发，此时需要调整网格大小或流数量。

总结来说，cudaStream_t 是并发调度的载体，cudaMemcpyAsync 和核函数启动是异步提交的两种主要操作，默认流与显式流的核心差异在于隐式同步。掌握这些机制后，才能进一步设计多流流水线，让数据传输与计算真正重叠，从而提升整体吞吐。