第13章 CUDA流与异步 - 重叠计算与通信 (CUDA Streams and Asynchrony - Overlapping Computation and Communication)

一、导读

13.1 本章学习目标

通过本章的学习，读者将能够：

理解CUDA流（Stream）的概念及其在异步执行中的作用

掌握创建、管理和同步CUDA流的方法

深入理解计算与数据传输重叠的原理和实现技巧

学会使用多个流实现并发kernel执行

掌握CUDA事件（Event）用于精确计时和同步

了解默认流与自定义流的交互行为，包括隐式同步点

能够设计高效的流水线策略，最大化GPU利用率

了解现代CUDA特性如CUDA Graphs对异步执行的改进

理解多GPU系统中的流使用策略

掌握异步编程中的常见问题和调试技巧

理解流优先级及其对调度的影响

掌握异步内存传输的最佳实践

13.2 为什么本章如此重要

在现代GPU应用中，性能优化不仅仅关乎单个kernel的执行速度，更关乎如何有效地组织和调度整个工作流程。CUDA流和异步机制是实现这一目标的核心工具。

首先，GPU拥有强大的计算能力，但如果数据不能及时供应，计算单元就会闲置等待。传统的主机-设备同步模式（启动kernel，等待完成，传输结果）导致大量的空闲时间。通过使用CUDA流，我们可以将数据传输和计算重叠，使得GPU在等待数据的同时执行其他任务。

其次，现代GPU具有多个执行引擎：复制引擎（用于数据传输）、多个SM阵列、Tensor Core等。这些引擎可以并行工作，但需要程序员显式地使用多个流来利用这种并行性。单个流只能串行执行操作，无法充分利用硬件的并发能力。

第三，异步执行对于延迟敏感的应用至关重要。在实时渲染、视频处理和高频交易中，减少端到端延迟比提高吞吐量更重要。通过将大任务分解为小任务并使用多个流并发执行，可以显著降低延迟。

在实际项目中，流的正确使用可以带来显著的性能提升。例如，在深度学习推理中，可以将模型的不同层分配到不同的流，实现流水线执行；在多GPU系统中，可以使用流来重叠通信和计算；在视频处理中，可以使用流来处理多个帧的并发编码。

然而，异步编程也带来了复杂性：竞态条件、死锁、资源竞争等问题可能出现。理解流的语义和同步机制对于编写正确的异步代码至关重要。

异步编程的演进：

CUDA从早期版本就支持流和异步操作，但随着架构的演进，其功能不断增强：

Fermi：引入了真正的并发kernel执行
Kepler：增加了多个复制引擎，改善了并发能力
Maxwell：进一步优化了复制引擎
Volta：引入了独立线程调度，改善了流的行为
CUDA 10+：引入了CUDA Graphs，减少了启动开销
Ampere/Hopper：增加了更多的复制引擎和更好的流调度

13.3 核心问题引导

什么是CUDA流？它与线程有什么关系？

默认流和自定义流有什么区别？何时应该使用自定义流？

如何实现kernel执行与数据传输的重叠？需要什么条件？

多个流中的操作如何排序？哪些操作会隐式同步？

CUDA事件的作用是什么？如何使用事件进行精确计时？

什么是页锁定内存（pinned memory）？为什么它对异步传输很重要？

CUDA Graphs如何改进传统的流管理？它的优势是什么？

在多GPU系统中，如何使用流来重叠跨GPU通信和本地计算？

流优先级如何影响调度？

如何调试异步代码中的问题？

二、核心概念详解

13.1 CUDA流的基本概念

流的定义：

CUDA流是一个操作序列，这些操作按照提交的顺序执行。操作包括kernel启动、内存传输、内存设置等。同一个流中的操作保证按序执行，但不同流中的操作可能并发执行。

默认流（Default Stream）：

每个CUDA上下文都有一个默认流（stream 0）。如果不显式指定流，所有操作都在默认流中执行。默认流有一个特殊属性：它会与所有其他流隐式同步。

cpp

// Using default stream
kernel<<<gridSize, blockSize>>>(args);  // Implicitly uses stream 0
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);  // Also stream 0

自定义流：

程序员可以创建多个自定义流，每个流独立于其他流执行操作。

cpp

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// Launch kernels in different streams
kernel1<<<gridSize, blockSize, 0, stream1>>>(args1);
kernel2<<<gridSize, blockSize, 0, stream2>>>(args2);

// Operations in stream1 and stream2 may execute concurrently

流的语义：

流内有序：同一个流中的操作严格按照提交顺序执行。

流间无序：不同流中的操作没有固定的执行顺序，可能并发执行。

隐式同步：默认流与所有其他流之间存在隐式同步点。

流与线程的关系：

流是GPU执行的操作队列，与CPU线程是不同的概念。多个CPU线程可以向同一个流提交操作，但通常建议每个流由一个CPU线程管理以避免竞争。

13.2 流的创建与管理

创建流：

cpp

cudaStream_t stream;
cudaStreamCreate(&stream);           // Default flags
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);  // No sync with default stream

cudaStreamNonBlocking标志创建的流不会与默认流隐式同步，这对于最大化并发性很有用。

销毁流：

cpp

cudaStreamDestroy(stream);

销毁流时，之前提交到该流的所有操作必须已经完成，否则函数会阻塞等待。

查询流状态：

cpp

cudaError_t err = cudaStreamQuery(stream);
if (err == cudaSuccess) {
    // All operations in this stream have completed
} else if (err == cudaErrorNotReady) {
    // Stream is still busy
}

同步流：

cpp

cudaStreamSynchronize(stream);  // Block host until stream completes

13.3 异步内存传输

同步 vs 异步传输：

cudaMemcpy：同步传输，阻塞主机直到传输完成
cudaMemcpyAsync：异步传输，立即返回，传输在后台进行

cpp

// Synchronous copy (blocks host)
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);

// Asynchronous copy (returns immediately)
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

页锁定内存（Pinned Memory）：

异步传输要求主机内存是页锁定的（pinned），否则CUDA会回退到同步传输。

cpp

float* h_data;
cudaMallocHost((void**)&h_data, size);  // Allocate pinned memory
// ... use h_data ...
cudaFreeHost(h_data);

页锁定内存不能被操作系统交换到磁盘，因此DMA引擎可以直接访问，实现高速异步传输。但过多的页锁定内存可能影响系统整体性能，应谨慎使用。

双向异步传输：

使用两个流可以实现H2D和D2H传输的重叠：

cpp

cudaStream_t stream_h2d, stream_d2h;
cudaStreamCreate(&stream_h2d);
cudaStreamCreate(&stream_d2h);

cudaMemcpyAsync(d_input, h_input, size, cudaMemcpyHostToDevice, stream_h2d);
cudaMemcpyAsync(h_output, d_output, size, cudaMemcpyDeviceToHost, stream_d2h);

// Both transfers may overlap

异步内存设置：

cpp

cudaMemsetAsync(d_data, 0, size, stream);

13.4 计算与通信重叠

基本原理：

GPU通常有多个执行引擎：

复制引擎（Copy Engine）：负责H2D和D2H传输
SM阵列：负责kernel执行

这些引擎可以并行工作。通过将数据传输放在一个流中，kernel执行放在另一个流中，可以实现重叠。

简单重叠示例：

cpp

cudaStream_t stream_compute, stream_transfer;
cudaStreamCreate(&stream_compute);
cudaStreamCreate(&stream_transfer);

// Transfer next batch while computing current batch
for (int i = 0; i < numBatches; i++) {
    // Transfer data for batch i+1
    if (i + 1 < numBatches) {
        cudaMemcpyAsync(d_data[i+1], h_data[i+1], batchSize, 
                       cudaMemcpyHostToDevice, stream_transfer);
    }
    
    // Compute on batch i
    kernel<<<gridSize, blockSize, 0, stream_compute>>>(d_data[i]);
    
    // Ensure transfer completes before compute on same data
    cudaStreamSynchronize(stream_transfer);
}

流水线策略：

将数据分成多个块，形成流水线：

Time →
Stream 1: [Transfer Block 1] [Compute Block 1] [Transfer Block 3] [Compute Block 3]
Stream 2:                    [Transfer Block 2] [Compute Block 2] [Transfer Block 4] ...

通过交错安排传输和计算，可以隐藏传输延迟。

多阶段流水线：

cpp

// Three-stage pipeline: H2D -> Compute -> D2H
cudaStream_t streams[3];
for (int i = 0; i < 3; i++) {
    cudaStreamCreate(&streams[i]);
}

for (int phase = 0; phase < numPhases; phase++) {
    int idx = phase % 3;
    
    // H2D transfer
    if (phase < numPhases - 2) {
        cudaMemcpyAsync(d_data[(idx+2)%3], h_data[(idx+2)%3], size,
                       cudaMemcpyHostToDevice, streams[idx]);
    }
    
    // Compute
    if (phase >= 1) {
        kernel<<<gridSize, blockSize, 0, streams[(idx+2)%3]>>>(d_data[(idx+2)%3]);
    }
    
    // D2H transfer
    if (phase >= 2) {
        cudaMemcpyAsync(h_result[(idx+1)%3], d_result[(idx+1)%3], size,
                       cudaMemcpyDeviceToHost, streams[(idx+1)%3]);
    }
}

13.5 CUDA事件

事件的定义：

CUDA事件是时间线上的标记点，可以用于：

精确计时

流间同步

查询操作完成状态

创建和销毁事件：

cpp

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

// ... use events ...

cudaEventDestroy(start);
cudaEventDestroy(stop);

记录事件：

cpp

cudaEventRecord(start, stream);
// ... some operations ...
cudaEventRecord(stop, stream);

等待事件：

cpp

// Make stream2 wait for event in stream1
cudaStreamWaitEvent(stream2, event, 0);

计时：

cpp

cudaEventRecord(start, stream);
kernel<<<gridSize, blockSize, 0, stream>>>(args);
cudaEventRecord(stop, stream);

cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("Kernel time: %.3f ms\n", milliseconds);

cudaEventElapsedTime提供高精度的计时，分辨率可达微秒级。

事件查询：

cpp

cudaError_t err = cudaEventQuery(event);
if (err == cudaSuccess) {
    // Event has been recorded (operation completed)
}

13.6 流的优先级

CUDA支持流优先级，允许高优先级的流抢占低优先级流的资源。

cpp

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

cudaStream_t highPriorityStream;
cudaStreamCreateWithPriority(&highPriorityStream, cudaStreamDefault, priority_high);

高优先级流的操作可能被优先调度，但这不保证严格的优先级顺序，仅作为提示。

优先级应用场景：

实时渲染：高优先级流处理当前帧，低优先级流处理背景效果
交互式应用：用户交互响应使用高优先级，后台计算使用低优先级

13.7 CUDA Graphs

CUDA Graphs是CUDA 10引入的新特性，允许预先录制一系列操作，然后一次性提交执行。

传统方式的问题：

每次kernel启动和内存传输都有CPU端的开销。对于小kernel，这些开销可能占主导。

Graphs的优势：

减少启动开销：Graph一次性提交，减少了多次API调用的开销

更好的调度：运行时可以看到整个依赖图，做出更优的调度决策

可重复执行：录制的Graph可以多次执行，无需重新构建

捕获依赖：自动捕获操作间的依赖关系

使用Graphs：

cpp

cudaGraph_t graph;
cudaGraphExec_t instance;

// Begin capturing
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

// Record operations
kernel1<<<gridSize, blockSize, 0, stream>>>(args1);
kernel2<<<gridSize, blockSize, 0, stream>>>(args2);

// End capture
cudaStreamEndCapture(stream, &graph);

// Instantiate and launch
cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
cudaGraphLaunch(instance, stream);
cudaStreamSynchronize(stream);

// Cleanup
cudaGraphDestroy(graph);
cudaGraphExecDestroy(instance);

更新Graph：

cpp

// Update parameters without recreating the graph
cudaGraphExecUpdate(instance, graph, NULL, &updateResult);

13.8 多GPU系统中的流

在多GPU系统中，流可以用于重叠跨GPU通信和本地计算。

cpp

// GPU 0
cudaSetDevice(0);
cudaStream_t stream0;
cudaStreamCreate(&stream0);

// GPU 1
cudaSetDevice(1);
cudaStream_t stream1;
cudaStreamCreate(&stream1);

// Overlap communication and computation
cudaSetDevice(0);
cudaMemcpyPeerAsync(d_data_1, 1, d_result_0, 0, size, stream0);  // Send to GPU 1

cudaSetDevice(1);
kernel_on_gpu1<<<gridSize, blockSize, 0, stream1>>>(d_data_1);  // Compute while receiving

使用NVLink等高速互联，跨GPU传输可以与计算有效重叠。

多GPU流水线：

GPU 0: [Compute] [Send to GPU 1] [Receive from GPU 1] [Compute]
GPU 1:            [Receive]       [Compute]          [Send back]

13.9 流回调

CUDA支持在流中注册回调函数，当流执行到该点时调用。

cpp

void CUDART_CB myCallback(cudaStream_t stream, cudaError_t status, void* userData) {
    printf("Stream reached callback point\n");
}

cudaStreamAddCallback(stream, myCallback, NULL, 0);

注意：回调函数在CPU上执行，不应包含耗时操作。

三、程序员的视角

3.1 在实际项目中的应用

深度学习推理：将模型的不同层分配到不同的流，实现流水线执行。同时处理多个请求时，每个请求可以使用独立的流。

视频处理：每个视频帧的处理可以使用独立的流，实现多帧并发编码/解码。

科学模拟：在多物理场耦合模拟中，不同的物理过程可以使用不同的流，实现部分并发。

数据库加速：查询处理的各个阶段（扫描、过滤、聚合）可以使用不同的流，形成流水线。

实时渲染：几何处理、光照计算、后处理可以使用不同的流，减少帧时间。

3.2 常见陷阱与误区

忘记页锁定内存：使用普通malloc分配的主机内存进行异步传输，导致静默回退到同步传输。

隐式同步陷阱：在默认流和自定义流之间交替使用，导致意外的同步点，破坏并发性。

资源竞争：多个流同时访问同一全局内存区域，导致数据竞争。需要适当的同步。

过度并行化：创建过多流可能导致资源碎片化和调度开销增加。应根据硬件能力选择合适的流数量。

事件泄漏：创建事件后忘记销毁，导致资源泄漏。

回调阻塞：在回调函数中执行耗时操作，阻塞流的执行。

3.3 与其他层级的联系

与内存层次的联系：异步传输依赖于DMA引擎和页锁定内存。理解这一机制有助于优化数据传输。

与同步机制的联系：事件和流同步是异步编程的基础。正确使用这些原语对于避免竞态条件至关重要。

与硬件的联系：不同GPU架构的复制引擎数量和带宽不同，影响重叠效果。应针对目标硬件调优。

与CUDA Graphs的联系：Graphs可以看作是一种高级的流管理方式，将多个操作打包为一个单元。

四、代码实战

4.1 基本的流使用示例

cpp

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void simpleKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] *= 2.0f;
    }
}

int main() {
    int N = 1000000;
    size_t size = N * sizeof(float);
    
    // Allocate pinned host memory
    float *h_data1, *h_data2;
    cudaMallocHost((void**)&h_data1, size);
    cudaMallocHost((void**)&h_data2, size);
    
    for (int i = 0; i < N; i++) {
        h_data1[i] = (float)i;
        h_data2[i] = (float)(i * 2);
    }
    
    // Allocate device memory
    float *d_data1, *d_data2;
    cudaMalloc((void**)&d_data1, size);
    cudaMalloc((void**)&d_data2, size);
    
    // Create streams
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);
    
    // Launch operations in parallel streams
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    
    // Stream 1: transfer and compute
    cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
    simpleKernel<<<gridSize, blockSize, 0, stream1>>>(d_data1, N);
    cudaMemcpyAsync(h_data1, d_data1, size, cudaMemcpyDeviceToHost, stream1);
    
    // Stream 2: transfer and compute
    cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);
    simpleKernel<<<gridSize, blockSize, 0, stream2>>>(d_data2, N);
    cudaMemcpyAsync(h_data2, d_data2, size, cudaMemcpyDeviceToHost, stream2);
    
    // Wait for both streams to complete
    cudaStreamSynchronize(stream1);
    cudaStreamSynchronize(stream2);
    
    printf("Stream processing completed.\n");
    
    // Cleanup
    cudaFreeHost(h_data1);
    cudaFreeHost(h_data2);
    cudaFree(d_data1);
    cudaFree(d_data2);
    cudaStreamDestroy(stream1);
    cudaStreamDestroy(stream2);
    
    return 0;
}

4.2 使用事件进行精确计时

cpp

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void timedKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        for (int i = 0; i < 100; i++) {
            data[idx] = sqrtf(data[idx] * data[idx] + 1.0f);
        }
    }
}

int main() {
    int N = 10000000;
    size_t size = N * sizeof(float);
    
    float *d_data;
    cudaMalloc((void**)&d_data, size);
    
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    
    // Record start event
    cudaEventRecord(start, stream);
    
    // Launch kernel
    timedKernel<<<gridSize, blockSize, 0, stream>>>(d_data, N);
    
    // Record stop event
    cudaEventRecord(stop, stream);
    
    // Wait for completion
    cudaEventSynchronize(stop);
    
    // Calculate elapsed time
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    
    printf("Kernel execution time: %.3f ms\n", milliseconds);
    
    // Cleanup
    cudaFree(d_data);
    cudaStreamDestroy(stream);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    
    return 0;
}

4.3 流水线数据传输与计算

cpp

#include <cuda_runtime.h>
#include <stdio.h>

#define NUM_CHUNKS 4

__global__ void processChunk(float* data, int chunkSize) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < chunkSize) {
        data[idx] *= 2.0f;
    }
}

int main() {
    int totalSize = 10000000;
    int chunkSize = totalSize / NUM_CHUNKS;
    size_t chunkBytes = chunkSize * sizeof(float);
    
    // Allocate pinned host memory for each chunk
    float *h_data[NUM_CHUNKS];
    float *d_data[NUM_CHUNKS];
    
    for (int i = 0; i < NUM_CHUNKS; i++) {
        cudaMallocHost((void**)&h_data[i], chunkBytes);
        cudaMalloc((void**)&d_data[i], chunkBytes);
        
        for (int j = 0; j < chunkSize; j++) {
            h_data[i][j] = (float)(i * chunkSize + j);
        }
    }
    
    // Create streams for each chunk
    cudaStream_t streams[NUM_CHUNKS];
    for (int i = 0; i < NUM_CHUNKS; i++) {
        cudaStreamCreate(&streams[i]);
    }
    
    int blockSize = 256;
    int gridSize = (chunkSize + blockSize - 1) / blockSize;
    
    // Pipeline: overlap transfer and compute
    for (int i = 0; i < NUM_CHUNKS; i++) {
        // Transfer chunk i to device
        cudaMemcpyAsync(d_data[i], h_data[i], chunkBytes, 
                       cudaMemcpyHostToDevice, streams[i]);
        
        // Process chunk i
        processChunk<<<gridSize, blockSize, 0, streams[i]>>>(d_data[i], chunkSize);
        
        // Transfer result back
        cudaMemcpyAsync(h_data[i], d_data[i], chunkBytes,
                       cudaMemcpyDeviceToHost, streams[i]);
    }
    
    // Wait for all streams
    for (int i = 0; i < NUM_CHUNKS; i++) {
        cudaStreamSynchronize(streams[i]);
    }
    
    printf("Pipeline processing completed.\n");
    
    // Cleanup
    for (int i = 0; i < NUM_CHUNKS; i++) {
        cudaFreeHost(h_data[i]);
        cudaFree(d_data[i]);
        cudaStreamDestroy(streams[i]);
    }
    
    return 0;
}

4.4 CUDA Graphs示例

cpp

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void graphKernel1(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] *= 2.0f;
    }
}

__global__ void graphKernel2(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] += 1.0f;
    }
}

int main() {
    int N = 1000000;
    size_t size = N * sizeof(float);
    
    float *d_data;
    cudaMalloc((void**)&d_data, size);
    
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    
    cudaGraph_t graph;
    cudaGraphExec_t instance;
    
    // Begin capturing
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
    
    // Record operations
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    
    graphKernel1<<<gridSize, blockSize, 0, stream>>>(d_data, N);
    graphKernel2<<<gridSize, blockSize, 0, stream>>>(d_data, N);
    
    // End capture
    cudaStreamEndCapture(stream, &graph);
    
    // Instantiate
    cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
    
    // Launch multiple times
    for (int i = 0; i < 100; i++) {
        cudaGraphLaunch(instance, stream);
    }
    
    cudaStreamSynchronize(stream);
    
    printf("Graph executed 100 times.\n");
    
    // Cleanup
    cudaGraphDestroy(graph);
    cudaGraphExecDestroy(instance);
    cudaStreamDestroy(stream);
    cudaFree(d_data);
    
    return 0;
}

五、深入思考

5.1 历史背景

CUDA流从早期版本就存在，但随着GPU架构的演进，其功能不断增强。Fermi架构引入了真正的并发kernel执行；Kepler增加了多个复制引擎；Volta引入了独立线程调度，改善了流的行为。

CUDA Graphs是近年来的重要创新，解决了小kernel启动开销大的问题。它代表了CUDA编程模型的演进方向：从细粒度的API调用向更高层次的抽象过渡。

5.2 前沿方向

异步算法：越来越多的算法被设计为原生异步的，充分利用GPU的并发能力。

多流优化：自动化工具正在出现，可以帮助程序员优化流的使用和调度。

与统一内存的集成：统一内存的页面迁移可以与流结合，实现更灵活的数据管理。

Graph优化：编译器可以自动优化Graph的执行顺序，减少依赖冲突。

量子-经典混合：在未来的量子-经典混合系统中，流可能用于协调量子处理器和经典GPU之间的任务。

5.3 延伸阅读

NVIDIA CUDA C Programming Guide - Streams and Events

"How to Overlap Data Transfers in CUDA" - NVIDIA Developer Blog

CUDA Graphs Documentation

Vasily Volkov, "Better Performance at Lower Occupancy"

"Efficient Multi-GPU Computing with CUDA Streams" - Research papers

六、实验笔记

实验1：单流vs多流性能对比

目的：量化多流带来的并发收益。

步骤：

使用单流处理4个数据块

使用4个流并发处理同样的数据块

比较总执行时间

预期：如果传输和计算可以完全重叠，多流应显著快于单流。

实验2：事件计时精度验证

目的：验证CUDA事件的计时精度。

步骤：

使用事件测量已知时间的kernel

与CPU高精度计时器对比

分析误差来源

预期：事件计时应与CPU计时一致，误差在微秒级别。

实验3：Graphs vs 传统API性能对比

目的：比较CUDA Graphs与传统API调用的性能差异。

步骤：

实现传统方式的kernel启动

使用Graphs录制同样的操作

比较重复执行时的开销

预期：对于小kernel，Graphs应显著减少启动开销。

七、本章小结

CUDA流是实现异步执行和并发操作的核心机制。通过使用多个流，程序员可以重叠计算和数据传输，最大化GPU利用率。

默认流提供了简单的编程模型，但与所有其他流隐式同步，限制了并发性。自定义流允许真正的并发执行，但需要仔细管理同步。

CUDA事件提供了精确计时和流间同步的能力。结合流和事件，可以构建复杂的流水线策略。

CUDA Graphs代表了异步执行的新方向，通过预先录制操作序列，减少了启动开销，提高了调度效率。

在实际应用中，流的正确使用可以带来显著的性能提升，特别是在数据传输频繁或多任务并发的场景中。

八、习题

流的基本操作：创建3个流，分别在每个流中启动一个kernel，然后等待所有流完成。

隐式同步分析：以下代码中有哪些隐式同步点？

cpp

kernel1<<<...>>>();  // Default stream
kernel2<<<..., stream1>>>();
kernel3<<<...>>>();  // Default stream

异步传输条件：为什么异步传输需要页锁定内存？如果使用普通malloc会发生什么？

流水线设计：设计一个三阶段流水线（传输→计算→传回），使用4个流交错执行。

事件同步：使用cudaStreamWaitEvent使stream2等待stream1中的某个事件。

优先级测试：创建高优先级和低优先级流，观察调度行为是否有差异。

Graph录制：录制一个包含两个kernel和一个内存传输的Graph，然后重复执行100次。

多GPU流：在双GPU系统中，使用流重叠跨GPU传输和本地计算。

性能建模：估算使用N个流的理论最大加速比。考虑硬件限制。

综合实现：实现一个完整的流水线视频处理框架，使用流并发处理多个帧。

回调应用：在流中注册回调函数，在kernel完成后执行某些操作。

错误处理：如何在异步操作中检测和报告错误？