4. CUDA流与并发执行

在图像预处理这类“数据搬运多、计算密度中等、任务可切分”的场景中，GPU 利用率低往往不是因为核函数写得慢，而是因为默认流把拷贝和计算串行化了。CUDA 流（Stream）是解决这一问题的核心机制：它把 H2D 拷贝、核函数执行、D2H 拷贝组织成一条条可独立调度的任务队列，让不同队列之间能够重叠执行。本节围绕“流—事件—异步拷贝—核函数重叠”这条主线，说明如何把单流串行预处理改造成多流流水线。

一、默认流为什么慢：串行时间线

不使用流时，所有操作都提交到默认流（legacy default stream），同一流内严格顺序执行。以一批 8 张 1920×1080 的 RGB 图像做“去噪 + 缩放 + 归一化”预处理为例，单张图像的处理链路是：H2D 拷贝约 6.2 MB、核函数执行、D2H 拷贝约 2.1 MB。若逐张串行，时间线近似为：

默认流: [H2D#0][Kernel#0][D2H#0][H2D#1][Kernel#1][D2H#1] ...
        拷贝引擎空闲时，SM 也在等数据；核函数执行时，拷贝引擎又空闲。

问题在于：H2D/D2H 由拷贝引擎（Copy Engine）执行，核函数由 SM 执行，二者本是不同硬件单元，却因为同处一条流而被强制串行。只要把它们拆到不同流，拷贝与计算就能在时间上重叠。

二、流、事件与异步拷贝的配合关系

CUDA 流是一个按提交顺序执行的操作队列，不同流之间没有顺序约束，可并发。事件（Event）则是流内的“时间戳/同步点”，用于跨流等待或测量耗时。异步拷贝（cudaMemcpyAsync）必须配合页锁定内存（cudaMallocHost 或 cudaHostAlloc）才能真正异步：普通可分页内存在拷贝时可能被驱动同步化，导致流失去重叠能力。

下面这张图给出本节的核心结构：主机端页锁定内存、设备端缓冲区、拷贝引擎与 SM 之间的数据流，以及事件如何把“拷贝完成”这一状态传递给后续核函数。

三、把预处理拆成可重叠的三段

要让重叠真正发生，需要把每张图像的处理拆成三个可独立提交的操作：H2D 拷贝、核函数、D2H 拷贝。核函数内部完成去噪、缩放和归一化，例如把 1920×1080 缩放到 640×640，并做 (x/255 - mean)/std 归一化。关键约束是：同一张图像的核函数必须等它的 H2D 拷贝完成，D2H 必须等核函数完成；不同图像之间则没有依赖。

实现上，为每张图像分配独立的设备缓冲，或使用双缓冲轮转。下面给出双流、双缓冲的核心代码骨架：

const int N = 8;
const int STREAMS = 2;
cudaStream_t stream[STREAMS];
cudaEvent_t  ev_h2d[STREAMS], ev_kernel[STREAMS];

for (int s = 0; s < STREAMS; ++s) {
    cudaStreamCreate(&stream[s]);
    cudaEventCreateWithFlags(&ev_h2d[s], cudaEventDisableTiming);
    cudaEventCreateWithFlags(&ev_kernel[s], cudaEventDisableTiming);
}

// 双缓冲：每个流使用自己的 d_in / d_out
float *d_in[STREAMS], *d_out[STREAMS];
for (int s = 0; s < STREAMS; ++s) {
    cudaMalloc(&d_in[s],  inBytes);
    cudaMalloc(&d_out[s], outBytes);
}

for (int i = 0; i < N; ++i) {
    int s = i % STREAMS;
    // 1) 异步 H2D，提交到 stream[s]
    cudaMemcpyAsync(d_in[s], h_in[i], inBytes,
                    cudaMemcpyHostToDevice, stream[s]);
    cudaEventRecord(ev_h2d[s], stream[s]);

    // 2) 核函数等待本流 H2D 完成（同流本已有序，事件用于跨流场景）
    cudaStreamWaitEvent(stream[s], ev_h2d[s], 0);
    preprocessKernel<<<grid, block, 0, stream[s]>>>(d_in[s], d_out[s]);
    cudaEventRecord(ev_kernel[s], stream[s]);

    // 3) 异步 D2H
    cudaStreamWaitEvent(stream[s], ev_kernel[s], 0);
    cudaMemcpyAsync(h_out[i], d_out[s], outBytes,
                    cudaMemcpyDeviceToHost, stream[s]);
}

cudaDeviceSynchronize();

注意：同一流内操作本身有序，因此 cudaStreamWaitEvent 在同流场景下是冗余的；它真正的价值在于跨流依赖，例如“流 1 的核函数要等流 0 的拷贝完成”。上面的写法保留了事件，是为了让依赖关系显式化，便于后续扩展成更复杂的流水线。

四、用事件测量重叠效果

判断流水线是否真的重叠，不能只看总耗时，还要看拷贝与计算是否在时间上交错。用事件在关键位置打点即可：

cudaEvent_t t0, t1;
cudaEventCreate(&t0); cudaEventCreate(&t1);
cudaEventRecord(t0);
// ... 提交所有流任务 ...
cudaEventRecord(t1);
cudaEventSynchronize(t1);
float ms = 0;
cudaEventElapsedTime(&ms, t0, t1);
printf("pipeline total: %.3f ms\n", ms);

更细的观察方式是给每个流分别记录起止事件，再用 cudaEventElapsedTime 得到各流耗时。若双流总耗时接近单流的一半，说明重叠生效；若几乎不变，常见原因是：使用了可分页内存导致 cudaMemcpyAsync 被同步化、两个流共用同一块设备缓冲产生写后读冲突、或核函数本身占满 SM 使两个流无法并行。此时应检查内存分配方式、缓冲是否按流隔离，以及核函数的资源占用（寄存器、共享内存、block 数）。

五、流水线参数与调优要点

流数量不是越多越好。流数过少，拷贝引擎和 SM 容易互相等待；流数过多，则上下文切换和缓冲占用上升。对图像预处理这类任务，通常 2～4 条流即可覆盖拷贝与计算的重叠窗口。判断依据是单张图像的 H2D 时间、核函数时间、D2H 时间三者的比例：若核函数时间明显长于拷贝时间，增加流数收益有限；若拷贝时间占主导，增加流数能更充分地填满拷贝引擎。

把单流串行改成多流流水线后，图像预处理任务的时间线会从“拷贝—计算—拷贝”依次排列，变成多条流交错推进。理解流、事件、异步拷贝三者的分工，是后续把预处理与推理、后处理进一步流水化的基础。