第6章流与并发 - 重叠计算与传输（Streams and Concurrency）

一、导读

1.1 本章学习目标

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

理解 CUDA 流（Stream）的概念和作用

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

学会使用多个流实现计算与数据传输的重叠

理解默认流的行为及其对并发的影响

掌握异步内存拷贝和异步内核启动的使用

学会使用 CUDA 事件进行精确的性能测量

理解 GPU 并发执行的硬件限制和调度机制

能够设计和实现高效的流水线并行程序

掌握多 GPU 编程的基本方法

了解 CUDA Graph 的作用和使用场景

1.2 为什么本章如此重要

在前面的章节中，我们主要关注单个内核的执行和优化。然而，在实际应用中，程序往往包含多个阶段：数据加载、预处理、计算、后处理、结果输出等。如果这些阶段串行执行，GPU 可能在某些阶段处于空闲状态，造成资源浪费。

CUDA 流提供了一种机制，允许不同的操作并发执行：

计算与传输重叠：在内核执行的同时进行数据传输
多个内核并发：在支持 Hyper-Q 的架构上，多个内核可以同时执行
流水线并行：将任务分解为多个阶段，形成流水线

掌握流的用法可以显著提高程序的吞吐量，特别是在处理大规模数据或实时应用时。此外，对于多 GPU 系统，流是协调多个设备工作的关键工具。

1.3 核心问题引导

在深入学习之前，请思考以下问题：

什么是 CUDA 流？它与 CPU 线程有什么区别？

默认流有什么特殊行为？为什么它可能限制并发？

如何实现计算与数据传输的重叠？

什么条件下多个内核可以并发执行？

CUDA 事件与 CPU 计时器相比有什么优势？

什么是 Hyper-Q？哪些架构支持它？

如何设计一个高效的流水线程序？

多 GPU 编程面临哪些挑战？

二、核心概念详解

6.1 CUDA 流基础

什么是流？

CUDA 流是一个操作序列，这些操作按照提交的顺序在设备上执行。流内的操作保证按序执行，但不同流之间的操作可以并发执行。

操作的类型

内核启动（Kernel Launch）
内存拷贝（Memory Copy）
内存设置（Memory Set）
事件记录（Event Record）

默认流（Default Stream）

每个 CUDA 上下文都有一个默认流（也称为流 0）。不使用显式流的操作都在默认流中执行。

cpp

// 这些操作都在默认流中执行
kernel<<<grid, block>>>(args);  // 隐式使用默认流
cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);  // 同步拷贝

默认流的特殊性

默认流是一个同步流：

在默认流中的操作会等待所有其他流完成
其他流中的操作会等待默认流中的前序操作完成

这种行为在 CUDA 7.0 之后可以通过编译选项 --default-stream per-thread 改变，使每个主机线程拥有独立的默认流。

6.2 创建和管理流

创建流

cpp

cudaStream_t stream;
cudaStreamCreate(&stream);

使用流

cpp

// 异步内存拷贝
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);

// 在流中启动内核
kernel<<<grid, block, 0, stream>>>(args);

销毁流

cpp

cudaStreamDestroy(stream);

流的完整生命周期示例

cpp

#include <cuda_runtime.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() {
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    float *h_data, *d_data;
    cudaMallocHost(&h_data, size);  // 页锁定内存，用于异步传输
    cudaMalloc((void **)&d_data, size);
    
    // 异步拷贝到设备
    cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);
    
    // 在流中启动内核
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    simpleKernel<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(d_data, n);
    
    // 异步拷贝回主机
    cudaMemcpyAsync(h_data, d_data, size, cudaMemcpyDeviceToHost, stream);
    
    // 等待流完成
    cudaStreamSynchronize(stream);
    
    // 清理
    cudaFreeHost(h_data);
    cudaFree(d_data);
    cudaStreamDestroy(stream);
    
    return 0;
}

6.3 异步操作

异步内存拷贝

cudaMemcpyAsync 是非阻塞的，它在发起传输后立即返回，允许主机继续执行其他操作。

cpp

// 同步拷贝：阻塞直到完成
cudaMemcpy(dst, src, size, kind);

// 异步拷贝：立即返回，在指定流中执行
cudaMemcpyAsync(dst, src, size, kind, stream);

重要要求

异步拷贝需要使用页锁定内存（Pinned Memory）：

cpp

// 分配页锁定内存
float *h_data;
cudaMallocHost(&h_data, size);

// 或者
float *h_data = (float *)malloc(size);
cudaHostRegister(h_data, size, cudaHostRegisterDefault);

异步内存设置

cpp

cudaMemsetAsync(d_data, 0, size, stream);

6.4 计算与传输重叠

基本原理

现代 GPU 具有独立的拷贝引擎和计算引擎，可以同时进行数据传输和内核执行。

时间线示意：
Stream 0: [Copy H2D] [Kernel 1] [Kernel 2] [Copy D2H]
Stream 1:           [Copy H2D] [Kernel 1] [Kernel 2] [Copy D2H]

重叠效果：
        [Copy H2D][Kernel 1 ][Kernel 2 ][Copy D2H]
                  [Copy H2D][Kernel 1 ][Kernel 2 ][Copy D2H]
                ←─────── 重叠区域 ──────→

实现条件

使用多个流

使用异步操作

使用页锁定内存

数据分块处理

示例：双缓冲流水线

cpp

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

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

#define NUM_CHUNKS 4
#define CHUNK_SIZE (1 << 18)  // 每个块的大小

int main() {
    int totalSize = CHUNK_SIZE * NUM_CHUNKS;
    size_t bytes = totalSize * sizeof(float);
    size_t chunkBytes = CHUNK_SIZE * sizeof(float);
    
    // 分配页锁定内存
    float *h_data;
    cudaMallocHost(&h_data, bytes);
    
    // 分配设备内存
    float *d_data;
    cudaMalloc((void **)&d_data, chunkBytes);
    
    // 初始化数据
    for (int i = 0; i < totalSize; i++) {
        h_data[i] = (float)i;
    }
    
    // 创建流
    cudaStream_t streams[NUM_CHUNKS];
    for (int i = 0; i < NUM_CHUNKS; i++) {
        cudaStreamCreate(&streams[i]);
    }
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (CHUNK_SIZE + threadsPerBlock - 1) / threadsPerBlock;
    
    // 流水线执行
    for (int chunk = 0; chunk < NUM_CHUNKS; chunk++) {
        float *chunkPtr = h_data + chunk * CHUNK_SIZE;
        
        // 异步拷贝到设备
        cudaMemcpyAsync(d_data, chunkPtr, chunkBytes, 
                       cudaMemcpyHostToDevice, streams[chunk]);
        
        // 在流中执行内核
        processKernel<<<blocksPerGrid, threadsPerBlock, 0, streams[chunk]>>>(
            d_data, CHUNK_SIZE);
        
        // 异步拷贝回主机
        cudaMemcpyAsync(chunkPtr, d_data, chunkBytes, 
                       cudaMemcpyDeviceToHost, streams[chunk]);
    }
    
    // 等待所有流完成
    for (int i = 0; i < NUM_CHUNKS; i++) {
        cudaStreamSynchronize(streams[i]);
    }
    
    // 清理
    for (int i = 0; i < NUM_CHUNKS; i++) {
        cudaStreamDestroy(streams[i]);
    }
    cudaFreeHost(h_data);
    cudaFree(d_data);
    
    return 0;
}

6.5 CUDA 事件

事件的作用

CUDA 事件用于：

精确测量 GPU 操作的时间
同步不同流之间的操作
检测操作的完成状态

创建和销毁事件

cpp

cudaEvent_t event;
cudaEventCreate(&event);
cudaEventDestroy(event);

记录事件

cpp

cudaEventRecord(event, stream);

等待事件

cpp

// 主机等待事件
cudaEventSynchronize(event);

// 流等待事件
cudaStreamWaitEvent(stream, event, 0);

测量时间

cpp

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

cudaEventRecord(start, stream);
// ... 一些操作 ...
cudaEventRecord(stop, stream);

cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);

完整示例

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] = sin(data[idx]) + cos(data[idx]);
        }
    }
}

int main() {
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    float *d_data;
    cudaMalloc((void **)&d_data, size);
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    // 记录开始事件
    cudaEventRecord(start, stream);
    
    // 执行内核
    timedKernel<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(d_data, n);
    
    // 记录结束事件
    cudaEventRecord(stop, stream);
    
    // 等待完成
    cudaEventSynchronize(stop);
    
    // 计算时间
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("Kernel execution time: %.3f ms\n", milliseconds);
    
    cudaFree(d_data);
    cudaStreamDestroy(stream);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    
    return 0;
}

6.6 并发执行的硬件支持

Hyper-Q

从 Kepler 架构开始，NVIDIA GPU 支持 Hyper-Q 技术，允许多个硬件工作队列同时存在。这使得不同流中的操作可以真正并发执行。

架构	硬件队列数量
Fermi	1
Kepler	32
Maxwell	32
Pascal	32
Volta	32
Ampere	32+

并发内核执行的条件

GPU 必须有足够的资源（SM、寄存器、共享内存）同时运行多个内核

内核不能占用全部资源（否则无法并发）

使用非默认流

内核之间没有依赖关系

查询设备属性

cpp

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, deviceId);
printf("Concurrent Kernels: %s\n", 
       prop.concurrentKernels ? "Supported" : "Not Supported");
printf("Async Engine Count: %d\n", prop.asyncEngineCount);

6.7 流优先级

CUDA 允许为流设置优先级，高优先级的流中的操作会被优先调度。

cpp

int priorityHigh, priorityLow;
cudaDeviceGetStreamPriorityRange(&priorityLow, &priorityHigh);

cudaStream_t highPriorityStream;
cudaStreamCreateWithPriority(&highPriorityStream, cudaStreamNonBlocking, priorityHigh);

cudaStream_t lowPriorityStream;
cudaStreamCreateWithPriority(&lowPriorityStream, cudaStreamNonBlocking, priorityLow);

6.8 CUDA Graph

CUDA Graph 允许将一系列操作记录为图，然后重复执行，减少启动开销。

cpp

// 创建图
cudaGraph_t graph;
cudaGraphExec_t instance;

// 开始捕获
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

// 记录操作
cudaMemcpyAsync(..., stream);
kernel<<<..., stream>>>(...);
cudaMemcpyAsync(..., stream);

// 结束捕获
cudaStreamEndCapture(stream, &graph);

// 实例化图
cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);

// 执行图（可重复）
cudaGraphLaunch(instance, stream);
cudaStreamSynchronize(stream);

// 清理
cudaGraphDestroy(graph);
cudaGraphExecDestroy(instance);

三、程序员的视角

6.1 在实际项目中的应用

视频处理流水线

帧读取 → 预处理 → GPU 处理 → 后处理 → 帧输出
   ↓         ↓          ↓          ↓         ↓
Stream 0  Stream 1   Stream 2   Stream 3  Stream 4

深度学习推理

批次数据分批传输
多层网络的不同阶段并发
多请求并行处理

科学模拟

边界条件更新与内部计算重叠
多物理场耦合的并发求解

6.2 常见陷阱与误区

误区一：更多流总是更好

事实：过多的流会增加管理开销，且受硬件资源限制。通常 4-8 个流就足以实现良好的重叠。

误区二：异步操作不需要同步

事实：异步操作只是不阻塞主机，但仍需要在适当的时候同步以确保数据就绪。

陷阱：忘记使用页锁定内存

普通 malloc 分配的内存不能用于异步传输，会导致退化为同步操作。

陷阱：默认流的隐式同步

在默认流和其他流之间切换可能导致意外的同步点。

6.3 与其他层级的联系

第4章：异步传输需要页锁定内存
第5章：流内操作保持顺序，跨流操作需要同步
第8章：流是实现高性能流水线的关键

四、代码实战

6.1 多流性能对比

cpp

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

__global__ void computeKernel(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        for (int i = 0; i < 50; i++) {
            data[idx] = sqrt(data[idx] * data[idx] + 1.0f);
        }
    }
}

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

#define NUM_STREAMS 4
#define CHUNK_SIZE (1 << 19)

void testSingleStream(int numChunks) {
    size_t totalSize = CHUNK_SIZE * numChunks * sizeof(float);
    size_t chunkBytes = CHUNK_SIZE * sizeof(float);
    
    float *h_data;
    CUDA_CHECK(cudaMallocHost(&h_data, totalSize));
    
    float *d_data;
    CUDA_CHECK(cudaMalloc((void **)&d_data, chunkBytes));
    
    cudaStream_t stream;
    CUDA_CHECK(cudaStreamCreate(&stream));
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (CHUNK_SIZE + threadsPerBlock - 1) / threadsPerBlock;
    
    CUDA_CHECK(cudaEventRecord(start, stream));
    
    for (int i = 0; i < numChunks; i++) {
        float *chunkPtr = h_data + i * CHUNK_SIZE;
        CUDA_CHECK(cudaMemcpyAsync(d_data, chunkPtr, chunkBytes, 
                   cudaMemcpyHostToDevice, stream));
        computeKernel<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(d_data, CHUNK_SIZE);
        CUDA_CHECK(cudaMemcpyAsync(chunkPtr, d_data, chunkBytes, 
                   cudaMemcpyDeviceToHost, stream));
    }
    
    CUDA_CHECK(cudaEventRecord(stop, stream));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float milliseconds = 0;
    CUDA_CHECK(cudaEventElapsedTime(&milliseconds, start, stop));
    printf("Single stream time (%d chunks): %.3f ms\n", numChunks, milliseconds);
    
    CUDA_CHECK(cudaFreeHost(h_data));
    CUDA_CHECK(cudaFree(d_data));
    CUDA_CHECK(cudaStreamDestroy(stream));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
}

void testMultiStream(int numChunks) {
    size_t totalSize = CHUNK_SIZE * numChunks * sizeof(float);
    size_t chunkBytes = CHUNK_SIZE * sizeof(float);
    
    float *h_data;
    CUDA_CHECK(cudaMallocHost(&h_data, totalSize));
    
    float *d_data;
    CUDA_CHECK(cudaMalloc((void **)&d_data, chunkBytes));
    
    cudaStream_t streams[NUM_STREAMS];
    for (int i = 0; i < NUM_STREAMS; i++) {
        CUDA_CHECK(cudaStreamCreate(&streams[i]));
    }
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (CHUNK_SIZE + threadsPerBlock - 1) / threadsPerBlock;
    
    CUDA_CHECK(cudaEventRecord(start, 0));
    
    for (int i = 0; i < numChunks; i++) {
        int streamIdx = i % NUM_STREAMS;
        float *chunkPtr = h_data + i * CHUNK_SIZE;
        
        CUDA_CHECK(cudaMemcpyAsync(d_data, chunkPtr, chunkBytes, 
                   cudaMemcpyHostToDevice, streams[streamIdx]));
        computeKernel<<<blocksPerGrid, threadsPerBlock, 0, streams[streamIdx]>>>(
            d_data, CHUNK_SIZE);
        CUDA_CHECK(cudaMemcpyAsync(chunkPtr, d_data, chunkBytes, 
                   cudaMemcpyDeviceToHost, streams[streamIdx]));
    }
    
    CUDA_CHECK(cudaEventRecord(stop, 0));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float milliseconds = 0;
    CUDA_CHECK(cudaEventElapsedTime(&milliseconds, start, stop));
    printf("Multi-stream time (%d chunks): %.3f ms\n", numChunks, milliseconds);
    
    for (int i = 0; i < NUM_STREAMS; i++) {
        CUDA_CHECK(cudaStreamDestroy(streams[i]));
    }
    CUDA_CHECK(cudaFreeHost(h_data));
    CUDA_CHECK(cudaFree(d_data));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
}

int main() {
    int numChunks = 8;
    
    printf("Testing with %d chunks of %d elements each:\n", numChunks, CHUNK_SIZE);
    
    testSingleStream(numChunks);
    testMultiStream(numChunks);
    
    return 0;
}

6.2 流依赖同步

cpp

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

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

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

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

int main() {
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    float *h_data;
    CUDA_CHECK(cudaMallocHost(&h_data, size));
    
    float *d_data;
    CUDA_CHECK(cudaMalloc((void **)&d_data, size));
    
    for (int i = 0; i < n; i++) {
        h_data[i] = (float)i;
    }
    
    cudaStream_t stream1, stream2;
    CUDA_CHECK(cudaStreamCreate(&stream1));
    CUDA_CHECK(cudaStreamCreate(&stream2));
    
    cudaEvent_t dependency;
    CUDA_CHECK(cudaEventCreate(&dependency));
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    // Stream 1: 拷贝 + Kernel A
    cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream1);
    kernelA<<<blocksPerGrid, threadsPerBlock, 0, stream1>>>(d_data, n);
    
    // 记录依赖事件
    CUDA_CHECK(cudaEventRecord(dependency, stream1));
    
    // Stream 2: 等待事件后执行 Kernel B
    CUDA_CHECK(cudaStreamWaitEvent(stream2, dependency, 0));
    kernelB<<<blocksPerGrid, threadsPerBlock, 0, stream2>>>(d_data, n);
    
    // 拷贝回主机
    cudaMemcpyAsync(h_data, d_data, size, cudaMemcpyDeviceToHost, stream2);
    
    // 等待完成
    CUDA_CHECK(cudaStreamSynchronize(stream2));
    
    printf("First 10 results: ");
    for (int i = 0; i < 10; i++) {
        printf("%.1f ", h_data[i]);
    }
    printf("\n");
    
    CUDA_CHECK(cudaFreeHost(h_data));
    CUDA_CHECK(cudaFree(d_data));
    CUDA_CHECK(cudaStreamDestroy(stream1));
    CUDA_CHECK(cudaStreamDestroy(stream2));
    CUDA_CHECK(cudaEventDestroy(dependency));
    
    return 0;
}

6.3 多 GPU 简单示例

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

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

int main() {
    int deviceCount;
    CUDA_CHECK(cudaGetDeviceCount(&deviceCount));
    
    if (deviceCount < 2) {
        printf("Need at least 2 GPUs for this example.\n");
        return 0;
    }
    
    printf("Found %d GPUs\n", deviceCount);
    
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    // 在每个 GPU 上执行
    for (int dev = 0; dev < deviceCount && dev < 2; dev++) {
        CUDA_CHECK(cudaSetDevice(dev));
        
        float *h_data = (float *)malloc(size);
        float *d_data;
        
        for (int i = 0; i < n; i++) {
            h_data[i] = (float)i;
        }
        
        CUDA_CHECK(cudaMalloc((void **)&d_data, size));
        CUDA_CHECK(cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice));
        
        int threadsPerBlock = 256;
        int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
        
        cudaEvent_t start, stop;
        CUDA_CHECK(cudaEventCreate(&start));
        CUDA_CHECK(cudaEventCreate(&stop));
        
        CUDA_CHECK(cudaEventRecord(start));
        simpleKernel<<<blocksPerGrid, threadsPerBlock>>>(d_data, n);
        CUDA_CHECK(cudaEventRecord(stop));
        CUDA_CHECK(cudaEventSynchronize(stop));
        
        float milliseconds = 0;
        CUDA_CHECK(cudaEventElapsedTime(&milliseconds, start, stop));
        
        cudaDeviceProp prop;
        CUDA_CHECK(cudaGetDeviceProperties(&prop, dev));
        printf("GPU %d (%s): %.3f ms\n", dev, prop.name, milliseconds);
        
        CUDA_CHECK(cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost));
        
        CUDA_CHECK(cudaFree(d_data));
        CUDA_CHECK(cudaEventDestroy(start));
        CUDA_CHECK(cudaEventDestroy(stop));
        free(h_data);
    }
    
    return 0;
}

五、深入思考

5.1 历史背景

早期的 CUDA 版本只支持简单的串行执行模型。随着应用需求的增长，NVIDIA 逐步引入了流、并发内核、Hyper-Q 等特性，使 GPU 能够更好地处理复杂的工作负载。

5.2 前沿方向

CUDA Graph：减少内核启动开销
Multi-Instance GPU (MIG)：将单个 GPU 划分为多个独立实例
NVLink Switch：多 GPU 之间的高速互连

5.3 延伸阅读

NVIDIA CUDA C Programming Guide - Streams 章节
CUDA Concurrent Execution 白皮书
NVIDIA Developer Blog 上的流相关文章

六、实验笔记

实验 1：测量流重叠的效果

目标：验证计算与传输重叠带来的性能提升。

步骤：

实现单流和多流版本

在不同数据规模下测试

使用 nsight 可视化执行时间线

分析重叠程度

实验 2：探索并发内核的限制

目标：了解什么条件下内核可以并发执行。

步骤：

编写资源占用不同的内核

测试不同组合的并发执行情况

分析资源占用与并发度的关系

七、本章小结

本章介绍了 CUDA 流与并发编程的核心概念。我们学习了如何创建和管理流、实现计算与传输的重叠、使用事件进行同步和计时。

关键要点包括：

流是操作的有序序列，不同流可以并发执行
异步操作需要页锁定内存才能有效工作
CUDA 事件提供精确的 GPU 端计时和同步机制
Hyper-Q 技术支持多个硬件工作队列
流优先级可以控制调度顺序

在下一章中，我们将探讨 CUDA 库的使用，学习如何利用现成的高效库来简化开发。

八、习题

解释 CUDA 流的概念及其与 CPU 线程的区别。

默认流有什么特殊行为？如何改变这种行为？

为什么异步内存拷贝需要页锁定内存？

列举实现计算与传输重叠的必要条件。

CUDA 事件有哪些用途？

什么是 Hyper-Q？它如何改善并发性能？

编写代码演示如何使用流优先级。

解释 CUDA Graph 的优势和使用场景。

如何查询设备支持的并发特性？

设计一个多 GPU 的数据并行处理方案。

第6章 流与并发 - 重叠计算与传输（Streams and Concurrency）