CUDA Streams 并发实战
基于 RTX 5090 (2 个 async copy engine) 和 A100-SXM4-80GB (3 个 async copy engine) 双平台实测。本文演示 CUDA Streams 如何实现 H2D + Kernel + D2H 重叠执行，并测量实际加速比。
1. 背景
1.1 Default Stream vs Explicit Streams
CUDA 的默认 stream（stream 0）是同步的——所有操作按提交顺序串行执行：H2D0 → Kernel0 → D2H0 → H2D1 → Kernel1 → D2H1。这意味着 GPU 的 compute engine 在数据传输时空闲，copy engine 在 kernel 执行时空闲——两个 engine 交替闲置，硬件利用率低下。
显式 stream（cudaStreamCreate）打破了这种串行约束。每个 stream 内部仍然保持顺序，但不同 stream 之间可以并发。当 stream 0 的 kernel 在执行时，stream 1 的 H2D 可以同时进行——GPU 的三个硬件引擎（2 个 copy engine + 1 个 compute engine）可以同时忙。
1.2 硬件并发能力
RTX 5090 有 2 个 async copy engine，理论最多可实现 2 路并发数据传输。加上 compute engine，最多 3 个操作可同时进行。A100-SXM4-80GB 有 3 个 async copy engine，理论上可实现 3 路并发数据传输（4 个操作同时进行）。但实际并发度都受限于：
- PCIe 带宽共享：多个 copy engine 共享同一 PCIe 链路（RTX 5090: Gen 5 ×16 ~63 GB/s；A100: Gen 4 ×16 ~32 GB/s 理论值）。如果多个 engine 同时做 H2D，带宽被均分。
- H2D 和 D2H 的方向冲突：H2D（写 GPU）和 D2H（读 GPU）共享 PCIe 双向带宽。同时做 H2D + D2H 时，双向带宽可能小于两个单向之和（nvbandwidth 实测：单向 56.3 GB/s，双向各 50.3 GB/s）。
- Kernel 执行时长：只有当 kernel 执行时间与数据传输时间大致匹配时，才能最大化重叠。如果 kernel 太短，数据传输追不上；如果 kernel 太长，copy engine 闲置。
1.3 测试策略
以下程序用 4 个 stream 各处理 256 MB 数据（含 H2D + kernel + D2H），对比串行和并发的总耗时。kernel 的循环次数 (K=1024) 经过调整，使单 stream 的计算时间与传输时间大致相当——这是最大化重叠的前提。
2. 对比测试程序
cat > stream_overlap.cu << 'EOF'
#include <cuda_runtime.h>
#include <stdio.h>
#define N (64 * 1024 * 1024) // 256 MB per buffer
#define K 1024 // dummy compute iterations
#define STREAMS 4
#define CHECK(cmd) do { \
cudaError_t e = cmd; \
if (e != cudaSuccess) { \
printf("Error: %s\n", cudaGetErrorString(e)); \
exit(1); \
} \
} while(0)
__global__ void dummy_kernel(float *d, int n, float s) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
float x = d[i];
for (int j = 0; j < K; j++)
x = x * s + 1.0f;
d[i] = x;
}
}
int main() {
float *h[STREAMS], *d[STREAMS];
cudaStream_t s[STREAMS];
for (int i = 0; i < STREAMS; i++) {
CHECK(cudaMallocHost(&h[i], N * sizeof(float)));
CHECK(cudaMalloc(&d[i], N * sizeof(float)));
CHECK(cudaStreamCreate(&s[i]));
}
cudaEvent_t start, stop;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
// === Test 1: Sequential (no streams) ===
CHECK(cudaEventRecord(start, 0));
for (int i = 0; i < STREAMS; i++) {
CHECK(cudaMemcpy(d[i], h[i], N * sizeof(float),
cudaMemcpyHostToDevice));
dummy_kernel<<<(N + 255) / 256, 256>>>(d[i], N, 1.0001f);
CHECK(cudaMemcpy(h[i], d[i], N * sizeof(float),
cudaMemcpyDeviceToHost));
}
CHECK(cudaDeviceSynchronize());
CHECK(cudaEventRecord(stop, 0));
CHECK(cudaEventSynchronize(stop));
float t_seq;
cudaEventElapsedTime(&t_seq, start, stop);
// === Test 2: Concurrent with streams ===
CHECK(cudaEventRecord(start, 0));
for (int i = 0; i < STREAMS; i++) {
CHECK(cudaMemcpyAsync(d[i], h[i], N * sizeof(float),
cudaMemcpyHostToDevice, s[i]));
dummy_kernel<<<(N + 255) / 256, 256, 0, s[i]>>>(d[i], N, 1.0001f);
CHECK(cudaMemcpyAsync(h[i], d[i], N * sizeof(float),
cudaMemcpyDeviceToHost, s[i]));
}
CHECK(cudaDeviceSynchronize());
CHECK(cudaEventRecord(stop, 0));
CHECK(cudaEventSynchronize(stop));
float t_async;
cudaEventElapsedTime(&t_async, start, stop);
float total = ((float)N * sizeof(float) * 3 * STREAMS)
/ (1024.0 * 1024.0 * 1024.0);
printf("Data: %.0f MB x %d streams = %.1f GB total\n",
(float)N * sizeof(float) / (1024 * 1024), STREAMS, total);
printf("Sequential: %7.2f ms (%.1f GB/s)\n",
t_seq, total / t_seq * 1000);
printf("Streamed: %7.2f ms (%.1f GB/s)\n",
t_async, total / t_async * 1000);
printf("Speedup: %.2fx\n", t_seq / t_async);
for (int i = 0; i < STREAMS; i++) {
CHECK(cudaFreeHost(h[i]));
CHECK(cudaFree(d[i]));
CHECK(cudaStreamDestroy(s[i]));
}
return 0;
}
EOF
nvcc -o stream_overlap stream_overlap.cu
./stream_overlap
3. 实测结果
RTX 5090 输出：
Data: 256 MB x 4 streams = 3.0 GB total
Sequential: 63.81 ms (47.0 GB/s)
Streamed: 27.01 ms (111.1 GB/s)
Speedup: 2.36x
3.1 为什么不是 3x
2 个 async copy engine + 1 个 compute engine 理论上可 3 路并发，但实际受限于：
- D2H 和 H2D 共享 copy engine：4 个 stream 的 8 次传输竞争 2 个 engine
- PCIe 半双工特性：同时 H2D + D2H 需要 PCIe 双向带宽（nvbandwidth 实测双向 ~50 GB/s vs 单向 ~56 GB/s）
- Kernel 执行时间不完美匹配：dummy kernel 时长可能不精确等于传输时间
4. A100 实测结果（官方 cuda-samples）
以下使用 NVIDIA 官方 cuda-samples 13.1 中的 simpleStreams、simpleHyperQ、simpleMultiCopy 在 A100-SXM4-80GB (CC 8.0, 3 copy engines) 上验证。
4.1 simpleStreams — 基础流并发
[simpleStreams] - GPU: NVIDIA A100-SXM4-80GB, 6912 Cores
> array_size = 16777216 (64.00 MB)
memcopy: 2.80 ms
kernel: 32.33 ms
non-streamed: 2.80 ms
4 streams: 2.57 ms
A100 的 4 stream 重叠比非重叠仅快 ~8%。这是因为 kernel 执行时间 (32.33 ms) 远大于 memcopy 时间 (2.80 ms)，kernel 成为瓶颈后数据拷贝的重叠空间被压缩。
4.2 simpleHyperQ — 32 路硬件队列并发
A100 支持 HyperQ——32 个独立的硬件工作队列，允许多个 stream 真正并行执行：
[simpleHyperQ] - GPU: NVIDIA A100-SXM4-80GB
Expected serial: 0.330s - 0.640s
Expected concurrent: ~0.020s
Measured: 0.027s
实测 0.027s，接近理论并发时间。32 路 kernel 在执行时间上几乎完全重叠——HyperQ (32 硬件工作队列) 是 A100 数据中心 GPU 的标志性能力。消费级 GPU 通常不支持或仅有受限的 HyperQ 功能。
4.3 simpleMultiCopy — 多 copy engine 带宽
[simpleMultiCopy] - GPU: NVIDIA A100-SXM4-80GB
Memcpy H2D: 2.72 ms (6.17 GB/s)
Memcpy D2H: 1.96 ms (8.54 GB/s)
Kernel: 0.004 ms
Avg. serialized: 1.638 ms
Avg. 4 streams: 1.682 ms
Speedup: -0.044 ms (negative!)
为什么重叠反而更慢？ Kernel 只有 0.004 ms——太轻量了。4 个 stream 的管理开销超过了重叠收益。这揭示了一个关键瓶颈：只有当 kernel 计算时间与传输时间相近时，重叠才有收益。本例中 kernel 比传输快 500 倍，GPU compute engine 在 99.8% 时间里空等数据。
关键推论：A100 有 3 个 copy engine 仍出现负加速比，说明 engine 数量不是瓶颈——kernel/传输比才是。始终用 Nsight Systems 确认重叠是否实际发生。
5. Nsight Systems 可视化 (概念)
在 Nsight Systems 中，stream 并发效果体现为时间线上各 stream 的操作互相交错：
Stream 0: [H2D ][Kernel ][D2H ]
Stream 1: [H2D ][Kernel ][D2H ]
Stream 2: [H2D ][Kernel ][D2H ]
Stream 3: [H2D ][Kernel ][D2H ]
Time: 0ms ──────────────────────────────── 27ms
没有 stream 的时序：
Default: [H2D0][K0][D2H0][H2D1][K1][D2H1][H2D2][K2][D2H2][H2D3][K3][D2H3]
Time: 0ms ──────────────────────────────────────────────────── 64ms
6. Stream 最佳实践
| 实践 | 说明 |
|---|---|
使用 cudaMemcpyAsync |
非阻塞拷贝是重叠的前提 |
| 每个 stream 独立 buffer | 避免 data hazard，每个 stream 有独立的 H2D/Kernel/D2H buffer |
pinned memory (cudaMallocHost) |
普通 pageable memory 无法异步 H2D |
| 先 issue 所有 work，再同步 | issue 阶段无阻塞，GPU 调度器自行安排并发 |
| Nsight Systems 验证 | 眼见为实——工具确认重叠是否真的发生 |