3、多流并发实战：多流并发矩阵乘法、多流并发数据拷贝、流优先级与调度

好，咱们进入第三章。多流并发，这名字听着挺唬人，对吧？其实说白了，就是让GPU同时干好几件事。你想想看，一个GPU里有那么多计算单元，如果只跑一个任务，那多浪费啊。我刚开始接触CUDA时，也觉得一个核函数跑完再跑下一个天经地义。直到有一次，我在项目里做视频处理，发现GPU利用率才30%多，剩下的时间全在等数据拷贝。嗯，从那时候起，我才真正开始研究多流。

3.1 多流并发矩阵乘法

先来个最经典的场景：矩阵乘法。假设你有N个矩阵对要相乘，最笨的办法就是循环，一次算一对。但这样，GPU的利用率其实很低。为什么？因为每次启动核函数都有开销，而且数据搬运和计算是串行的。

我个人习惯的做法是，把大矩阵切成小块，分到不同的流里去算。每个流独立管理自己的核函数执行、数据拷贝。这样，流A在计算时，流B可能正在拷数据，流C可能正在准备下一批。说白了，就是流水线作业。

来看代码：

cudaStream_t streams[NUM_STREAMS];
for (int i = 0; i < NUM_STREAMS; i++) {
    cudaStreamCreate(&streams[i]);
}

// 每个流处理一块数据
for (int i = 0; i < NUM_STREAMS; i++) {
    int offset = i * chunk_size;
    // 异步拷贝到设备
    cudaMemcpyAsync(d_A + offset, h_A + offset, 
                    chunk_size * sizeof(float), 
                    cudaMemcpyHostToDevice, streams[i]);
    cudaMemcpyAsync(d_B + offset, h_B + offset, 
                    chunk_size * sizeof(float), 
                    cudaMemcpyHostToDevice, streams[i]);
    
    // 在同一个流里启动核函数
    matmul_kernel<<<grid, block, 0, streams[i]>>>(
        d_A + offset, d_B + offset, d_C + offset, chunk_size);
    
    // 异步拷回结果
    cudaMemcpyAsync(h_C + offset, d_C + offset,
                    chunk_size * sizeof(float),
                    cudaMemcpyDeviceToHost, streams[i]);
}

cudaDeviceSynchronize();

这里的关键点，就是每个流里的操作都是异步的。你启动完一个流的拷贝，不用等它结束，就可以去启动下一个流的拷贝。GPU的硬件调度器会自动安排这些操作，尽量让计算单元和拷贝引擎都忙起来。

核心要点：多流并发的本质，是让GPU的不同硬件单元（计算单元、拷贝引擎）同时工作。不是让一个核函数跑得更快，而是让整个系统的吞吐量更高。

3.2 多流并发数据拷贝

说到数据拷贝，这里有个坑。我曾经在一个项目里，天真地以为只要用了多个流，数据拷贝就能自动并行。结果一测，发现还是串行的。为什么？因为GPU的拷贝引擎只有一个！

嗯，这里要注意。对于大多数消费级GPU（比如RTX 3090），确实只有一个拷贝引擎。你开再多流，数据拷贝也是排队执行的。那多流并发数据拷贝还有什么意义？

有意义。因为你可以让计算和拷贝重叠。比如流1在计算时，流2可以同时拷数据。这样，计算单元和拷贝引擎就能同时工作。我建议你养成一个习惯：把数据准备、计算、结果回传，分别放到不同的流里，或者至少让它们错开时间。

来看一个更实际的例子：

// 流0：专门负责数据拷贝
// 流1：专门负责计算
cudaStream_t stream_copy, stream_compute;
cudaStreamCreate(&stream_copy);
cudaStreamCreate(&stream_compute);

for (int i = 0; i < num_batches; i++) {
    // 在拷贝流里准备下一批数据
    cudaMemcpyAsync(d_input[i], h_input[i], 
                    size, cudaMemcpyHostToDevice, stream_copy);
    
    // 在计算流里处理上一批数据
    // 注意：这里需要等拷贝流完成
    cudaEventRecord(event, stream_copy);
    cudaStreamWaitEvent(stream_compute, event, 0);
    
    process_kernel<<<grid, block, 0, stream_compute>>>(d_input[i], d_output[i]);
    
    // 把结果拷回去
    cudaMemcpyAsync(h_output[i], d_output[i],
                    size, cudaMemcpyDeviceToHost, stream_copy);
}

小技巧：如果你用的是支持双拷贝引擎的GPU（比如某些专业卡或服务器卡），可以同时进行双向拷贝。这时候多流并发的效果会更明显。可以用 cudaDeviceGetAttribute 查一下你的设备支持几个拷贝引擎。

3.3 流优先级与调度

最后聊聊流优先级。这个功能，说实话，我平时用得不多。但在某些实时性要求高的场景里，它确实能救命。

CUDA支持给流设置优先级。高优先级的流，它的核函数会优先被调度执行。低优先级的流，就得等一等。我记得有一次做实时渲染，需要保证渲染帧的延迟不能超过16毫秒。但后台还有个数据预处理的任务，很占资源。这时候，我就把渲染任务放到高优先级流里，预处理放到低优先级流里。效果立竿见影。

设置优先级很简单：

int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);

cudaStream_t stream_high, stream_low;
cudaStreamCreateWithPriority(&stream_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&stream_low, cudaStreamNonBlocking, priority_low);

// 高优先级流：实时任务
render_kernel<<<grid, block, 0, stream_high>>>(...);

// 低优先级流：后台任务
preprocess_kernel<<<grid, block, 0, stream_low>>>(...);

注意：流优先级只影响核函数的调度顺序，不影响数据拷贝。而且，如果高优先级流一直有任务，低优先级流可能会被饿死。我曾经就遇到过这个问题，后台任务一直得不到执行。解决办法是，给低优先级流也留一些时间片，或者用事件同步来强制切换。

关于调度，还有一点要提：默认情况下，同一个流里的操作是严格顺序执行的。但不同流之间的操作，顺序是不确定的。你没法保证流A的核函数一定在流B之前执行完。所以，如果你需要跨流同步，就得用事件（cudaEvent）或者cudaStreamWaitEvent。

我个人习惯是，尽量让每个流独立处理一块完整的数据。这样就不需要跨流同步了，省心。如果实在需要同步，那就用事件，别用cudaDeviceSynchronize，那个太粗暴了，会把所有流都停下来。

好了，这一章的内容就这些。多流并发的核心思想，就是让GPU的各个硬件单元都忙起来。你只要记住三点：

  • 计算和拷贝要重叠
  • 每个流尽量独立
  • 优先级是最后的手段，别滥用

下一章，我们会聊更高级的话题：动态并行和CUDA图。到时候你会发现，多流并发只是开始，真正的大招还在后面。

← 上一章 📖 返回目录 下一章 →