CUDA流与事件：流的概念、创建与同步

核心原理：流是GPU任务的执行序列

CUDA流（Stream）是GPU上按顺序执行的一组操作序列。从硬件视角看，GPU包含多个独立的执行引擎（如计算引擎、拷贝引擎、张量核心等），流的作用就是将这些操作按依赖关系组织起来，提交给GPU调度器执行。理解流的关键在于：同一个流中的操作严格按提交顺序执行，不同流中的操作可以并发执行。

默认情况下，所有CUDA操作（内核启动、内存拷贝、事件记录）都在默认流（Default Stream，流0）中执行。默认流具有隐式同步特性——它与其他所有流之间存在同步屏障，这会阻止并发执行。因此，要实现真正的多流并发，必须创建显式流（Non-default Stream）。

流与硬件执行引擎的关系：现代GPU（如A100、H100）拥有独立的DMA拷贝引擎和计算引擎。当操作位于不同流时，GPU调度器可以将拷贝操作分配给拷贝引擎、将计算操作分配给计算引擎，从而实现拷贝与计算的重叠。但同一流内的操作必须等待前一个操作完成，这是流的基本约束。

上图展示了两个流与硬件引擎的映射关系。Stream 1中的H2D拷贝、Kernel A、D2H拷贝、Kernel B按顺序执行；Stream 2中的Kernel C独立运行。由于拷贝引擎和计算引擎相互独立，Stream 1的H2D拷贝可以与Stream 2的Kernel C同时执行，实现计算与传输的重叠。

流的创建与同步方法

CUDA提供完整的流管理API，包括创建、销毁、同步和查询。核心函数如下：

// 创建流
cudaStream_t stream;
cudaStreamCreate(&stream);                    // 普通流
// 或创建带优先级/标志的流
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
cudaStreamCreateWithPriority(&stream, cudaStreamNonBlocking, priority);

// 异步操作提交到流
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);
kernel<<<grid, block, 0, stream>>>(args...);

// 流同步
cudaStreamSynchronize(stream);                // 阻塞主机直到流中所有操作完成
cudaStreamQuery(stream);                      // 非阻塞查询，返回cudaSuccess或cudaErrorNotReady

// 销毁流
cudaStreamDestroy(stream);

关键标志说明：cudaStreamNonBlocking标志使该流不与默认流隐式同步，这是多流并发的前提。默认创建的流（cudaStreamCreate）实际上也带有非阻塞标志，但为了代码清晰，建议显式指定。

事件：精确计时与流间同步

CUDA事件（Event）是流中的标记点，用于两个核心目的：测量操作耗时和实现流间同步。事件本身不执行计算，只记录GPU执行到该点的时间戳。

API	功能	典型使用场景
`cudaEventCreate`	创建事件对象	初始化计时或同步点
`cudaEventRecord`	将事件记录到指定流	标记操作开始/结束
`cudaEventSynchronize`	阻塞主机直到事件完成	等待GPU执行到事件点
`cudaEventElapsedTime`	计算两个事件间毫秒数	测量内核/拷贝耗时
`cudaStreamWaitEvent`	让流等待事件完成	跨流依赖控制

事件计时是性能分析的基础工具。相比CPU计时器（如clock()），事件计时直接在GPU时间线上打点，不受CPU调度和PCIe传输延迟影响，精度更高。

// 事件计时示例
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream);
kernel<<<grid, block, 0, stream>>>(args...);
cudaEventRecord(stop, stream);

cudaEventSynchronize(stop);  // 等待GPU完成
float ms = 0.0f;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel 耗时: %.3f ms\n", ms);

案例：多流并发流水线

以典型的数据处理流水线为例：输入数据分块传输到GPU、逐块执行计算、结果回传。使用多流可以将三个阶段重叠，大幅提升吞吐。

流水线设计：将数据分为N个块，创建M个流（通常M=2~4）。每个流负责一个数据块的完整处理链：H2D拷贝 → Kernel计算 → D2H拷贝。由于不同流的操作可以重叠，当流0执行Kernel时，流1可以同时执行H2D拷贝，流2可以执行D2H拷贝。

// 多流流水线核心代码
const int NUM_STREAMS = 4;
cudaStream_t streams[NUM_STREAMS];
for (int i = 0; i < NUM_STREAMS; ++i) {
    cudaStreamCreateWithFlags(&streams[i], cudaStreamNonBlocking);
}

for (int i = 0; i < num_chunks; ++i) {
    int s = i % NUM_STREAMS;
    int offset = i * chunk_size;

    // 1. H2D 拷贝
    cudaMemcpyAsync(d_data + offset, h_data + offset,
                    chunk_bytes, cudaMemcpyHostToDevice, streams[s]);

    // 2. 计算内核
    process_kernel<<<grid, block, 0, streams[s]>>>(d_data + offset, chunk_size);

    // 3. D2H 拷贝
    cudaMemcpyAsync(h_result + offset, d_data + offset,
                    chunk_bytes, cudaMemcpyDeviceToHost, streams[s]);
}

// 等待所有流完成
for (int i = 0; i < NUM_STREAMS; ++i) {
    cudaStreamSynchronize(streams[i]);
    cudaStreamDestroy(streams[i]);
}

该模式的关键在于：每个流内部保持H2D→Kernel→D2H的顺序依赖，但不同流之间完全独立。GPU调度器会自动将可并行的操作分配到空闲引擎上。

误区与复盘

常见误区一：忽略默认流的隐式同步。如果使用默认流（流0）执行任何操作，它会阻塞所有其他流。例如，在循环中调用cudaMemcpy（同步版本）会强制GPU等待，破坏并发。必须使用cudaMemcpyAsync并指定非默认流。

常见误区二：过度创建流。创建大量流（如100个）并不会带来线性加速。GPU硬件并发能力有限（通常2~4个引擎），过多流反而增加调度开销。建议流数量与数据块大小匹配，一般2~8个流即可。

常见误区三：事件同步位置错误。cudaEventSynchronize会阻塞CPU，如果在内核启动后立即调用，就退化为同步执行。正确做法是：先记录事件，完成所有异步提交后再同步等待。

复盘建议：使用nsys profile或ncu工具查看GPU时间线，确认不同流的操作是否真正重叠。如果时间线显示串行执行，检查是否存在隐式同步（如默认流使用、cudaDeviceSynchronize调用、内存分配未使用cudaMallocAsync等）。

流和事件是CUDA并发编程的基石。掌握流的创建、同步和事件计时，是后续实现多流并发、多CUDA流推理加速的前提。实际开发中，建议先用事件计时验证单流性能，再逐步引入多流，通过性能分析工具确认重叠效果，避免盲目增加流数量。