Cuda编程：流(Streams)和事件(Events)

最新推荐文章于 2025-05-16 13:26:14 发布

倔强老吕

最新推荐文章于 2025-05-16 13:26:14 发布

阅读量406

点赞数 11

分类专栏： Cuda编程文章标签： c++

本文链接：https://blog.csdn.net/WaitingHereFor/article/details/148001391

版权

Cuda编程专栏收录该内容

8 篇文章

订阅专栏

1. 流(Streams)基础

1.1 什么是CUDA流

CUDA流是一系列按顺序执行的命令序列，不同流中的命令可以并行执行。流提供了以下关键功能：

并发执行：不同流中的操作可以同时进行
执行顺序控制：同一流中的操作按提交顺序执行
资源隔离：不同流的操作可以独立管理

1.2 流的创建与销毁

cudaStream_t stream;
cudaStreamCreate(&stream);  // 创建流
// ... 使用流 ...
cudaStreamDestroy(stream);  // 销毁流

1.3 默认流

CUDA有一个特殊的"默认流"(也称为空流或流0)：

所有未指定流的操作都在默认流中执行
默认流是同步的，会阻塞主机直到操作完成

2. 流的使用模式

2.1 基本使用示例

// 创建两个流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 在stream1中执行
kernel1<<<blocks, threads, 0, stream1>>>(...);
cudaMemcpyAsync(dev1, host1, size, cudaMemcpyHostToDevice, stream1);

// 在stream2中执行
kernel2<<<blocks, threads, 0, stream2>>>(...);
cudaMemcpyAsync(dev2, host2, size, cudaMemcpyHostToDevice, stream2);

// 等待两个流完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

2.2 计算与数据传输重叠

cudaStream_t computeStream, dataStream;
cudaStreamCreate(&computeStream);
cudaStreamCreate(&dataStream);

// 异步传输数据到设备
cudaMemcpyAsync(dev_data, host_data, size, cudaMemcpyHostToDevice, dataStream);

// 在计算流中执行核函数
kernel<<<blocks, threads, 0, computeStream>>>(dev_data);

// 异步将结果传回主机
cudaMemcpyAsync(host_result, dev_result, size, cudaMemcpyDeviceToHost, dataStream);

3. 事件(Events)基础

3.1 什么是CUDA事件

CUDA事件是流中的标记点，用于：

同步流的执行
测量时间间隔
建立流间的依赖关系

3.2 事件的创建与销毁

cudaEvent_t event;
cudaEventCreate(&event);  // 创建事件
// ... 使用事件 ...
cudaEventDestroy(event);  // 销毁事件

4. 事件的使用模式

4.1 计时示例

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, 0);  // 记录开始时间
kernel<<<blocks, threads>>>(...);
cudaEventRecord(stop, 0);  // 记录结束时间

cudaEventSynchronize(stop);  // 等待事件完成
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);  // 计算时间差

printf("Kernel time: %.2f ms\n", elapsedTime);

4.2 流间同步

cudaEvent_t event;
cudaEventCreate(&event);

// 在stream1中执行操作并记录事件
kernel1<<<..., stream1>>>(...);
cudaEventRecord(event, stream1);

// stream2等待event完成后再执行
cudaStreamWaitEvent(stream2, event, 0);
kernel2<<<..., stream2>>>(...);

5. 高级流技术

5.1 优先级流

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);

5.2 回调函数

void CUDART_CB myCallback(cudaStream_t stream, cudaError_t status, void *data) {
    printf("Callback executed in stream %p\n", stream);
}

// 在流中插入回调
kernel<<<..., stream>>>(...);
cudaStreamAddCallback(stream, myCallback, nullptr, 0);

6. 总结

最佳实践

流数量：通常创建2-4个流即可获得良好并发性，监控：使用nvidia-smi查看GPU利用率
非阻塞流：使用cudaStreamNonBlocking创建非阻塞流
资源竞争：避免过多流竞争有限资源
错误检查：检查流操作返回值
默认流同步：注意默认流会与其他流同步

常见问题与调试

隐式同步点：某些操作(如设备内存分配)会导致隐式同步，解决方案：预分配内存，使用非阻塞流
流优先级：高优先级流不保证绝对优先，只是调度提示
事件计时：确保事件已完成再查询时间，不准方案：使用CUDA Profiler获取更精确计时
多设备流：每个设备有自己的流和事件空间

通过合理使用流和事件，可以显著提高CUDA程序的并行度和整体性能。

7. 流和事件性能优化技巧

1. 计算与数据传输重叠优化

1.1 双缓冲技术(Double Buffering)

// 分配双缓冲
float *d_buffer[2];
cudaMalloc(&d_buffer[0], size);
cudaMalloc(&d_buffer[1], size);

cudaStream_t computeStream, dataStream;
cudaStreamCreate(&computeStream);
cudaStreamCreate(&dataStream);

for(int i = 0; i < iterations; i++) {
    // 当前计算缓冲和下一数据缓冲
    int curr = i % 2;
    int next = (i + 1) % 2;
    
    // 异步传输下一批数据
    cudaMemcpyAsync(d_buffer[next], h_data + next*chunk, 
                   chunk_size, cudaMemcpyHostToDevice, dataStream);
    
    // 计算当前批数据
    kernel<<<blocks, threads, 0, computeStream>>>(d_buffer[curr]);
    
    // 确保计算完成后再交换
    cudaStreamSynchronize(computeStream);
}

优化效果：完全重叠计算和数据传输，最大化设备利用率

2. 流优先级精细控制

2.1 多优先级流配置

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

// 创建不同优先级的流
cudaStream_t high_pri_stream, low_pri_stream;
cudaStreamCreateWithPriority(&high_pri_stream, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&low_pri_stream, cudaStreamNonBlocking, priority_low);

// 关键任务用高优先级流
kernel_critical<<<..., high_pri_stream>>>(...);

// 后台任务用低优先级流
kernel_background<<<..., low_pri_stream>>>(...);

优化效果：确保关键任务优先获得计算资源

3. 事件驱动的精细同步

3.1 最小化同步点

cudaEvent_t kernel1_done, kernel2_done;
cudaEventCreate(&kernel1_done);
cudaEventCreate(&kernel2_done);

// 流1执行序列
kernel1<<<..., stream1>>>(...);
cudaEventRecord(kernel1_done, stream1);

// 流2等待流1的kernel1完成
cudaStreamWaitEvent(stream2, kernel1_done, 0);
kernel2<<<..., stream2>>>(...);
cudaEventRecord(kernel2_done, stream2);

// 主机只需要等待最后一个事件
cudaEventSynchronize(kernel2_done);

优化效果：减少不必要的全局同步，提高并发性

4. 流池(Stream Pool)模式

4.1 流池实现

class StreamPool {
public:
    StreamPool(size_t pool_size) {
        streams.resize(pool_size);
        for(auto& stream : streams) {
            cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
        }
    }
    
    cudaStream_t getNextStream() {
        return streams[counter++ % streams.size()];
    }
    
    ~StreamPool() {
        for(auto& stream : streams) {
            cudaStreamDestroy(stream);
        }
    }

private:
    std::vector<cudaStream_t> streams;
    std::atomic<size_t> counter{0};
};

// 使用示例
StreamPool pool(4);
auto stream = pool.getNextStream();
kernel<<<..., stream>>>(...);

优化效果：避免频繁创建销毁流，减少开销

5. 事件重用技术

5.1 事件池实现

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// 重用事件进行多次计时
for(int i = 0; i < 10; i++) {
    cudaEventRecord(start, 0);
    kernel<<<...>>>(...);
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    
    float time;
    cudaEventElapsedTime(&time, start, stop);
    printf("Iteration %d: %.2f ms\n", i, time);
}

优化效果：避免重复创建销毁事件的开销

6. 高级优化技巧

6.1 流回调精细控制

void CUDART_CB postKernelCallback(cudaStream_t stream, cudaError_t status, void* userData) {
    // 回调中启动下一个核函数
    kernel_next<<<..., stream>>>(...);
}

// 主程序
kernel_first<<<..., stream>>>(...);
cudaStreamAddCallback(stream, postKernelCallback, nullptr, 0);

优化效果：实现核函数间的精细流水线

6.2 多设备流协同

cudaStream_t stream[2];
cudaEvent_t event[2];

for(int dev = 0; dev < 2; dev++) {
    cudaSetDevice(dev);
    cudaStreamCreate(&stream[dev]);
    cudaEventCreate(&event[dev]);
}

// 设备0执行并记录事件
cudaSetDevice(0);
kernel<<<..., stream[0]>>>(...);
cudaEventRecord(event[0], stream[0]);

// 设备1等待设备0完成
cudaSetDevice(1);
cudaStreamWaitEvent(stream[1], event[0], 0);
kernel<<<..., stream[1]>>>(...);

优化效果：实现多GPU间的精细任务调度

7. 性能分析技巧

Nsight Systems：分析流和事件的执行时间线
```
nsys profile --trace=cuda ./your_program
```
Nsight Compute：分析核函数在流中的执行细节
关键指标：
- 流间的空闲时间
- 核函数启动延迟
- 数据传输与计算的重叠程度