性能优化：流水线并行

CH-18-02-035·约 19 分钟·AI硬件与算力·更新 2026-08-11

假设你开了一家生意火爆的茶馆。客人点茶后，学徒负责烧水，师傅负责泡茶，小二负责端上桌。如果坚持“一壶茶从头盯到尾”的规矩——烧水时师傅在旁干等，泡茶时学徒闲着——明明两个人都没停过，客人却抱怨太慢。问题不在谁偷懒，而在工序之间互相等待。流水线并行，就是打破这种等待的艺术：学徒烧第二壶水的同时，师傅正在泡第一壶茶。两个人各忙各的，谁也不等谁。GPU 程序也面临同样的困境——把数据拷进显存、GPU 做计算、结果拷回内存，三个阶段串行执行时，GPU 大量时间在空转。这篇文章用“茶馆流水线”的视角，带你理解 CUDA 流、双缓冲、计算通信重叠和多级流水线调度。学完你会明白：高手写的 CUDA 程序看起来很“绕”，却跑得比谁都快，是因为他们把等待藏进了别人的工序里。

岗位关联

一句话价值：掌握流水线并行是 GPU 性能优化的分水岭，它能让训练与推理的吞吐翻倍，也是高性能计算岗位面试的高频考点。

可视化演示
流水线原理：一壶茶从头盯到尾
CUDA流基础：多雇几个学徒
双缓冲技术：两把水壶轮着用
计算通信重叠：边烧水边泡茶
多级流水线：拆出更多工序
优化避坑：老师傅的心法
AI 应用场景
Python 演示
常见误区
面试题

可视化演示

三步展示流水线并行：阶段执行（warmup预热→steady稳态→cooldown冷却）、调度策略对比（GPipe全前向全反向 vs 1F1B交替前向反向）、双缓冲重叠（compute与load交替执行隐藏延迟）。

本节配有交互动效（进入视口后自动加载，约 3KB）。

流水线阶段执行 warmup:各stage依次启动前向填充流水线 steady:各stage并行处理不同微批次（稳态） cooldown:反向传播逐步排空流水线 bubble_ratio ∝ (stages-1) / (stages+micro_batches-1) warmup→steady→cooldown三阶段调度。

流水线原理：一壶茶从头盯到尾

茶馆刚开张时，阿茶立下规矩：一壶茶从烧水到端上桌，由一个人全程盯完。学徒烧水时师傅候着，水开了师傅冲泡时学徒又闲着。客人络绎不绝，可每壶茶都要等足三分钟。两个伙计都喊累，效率却低得吓人——每道工序都在等上一道工序结束，谁也没有真正“并行”过。

GPU 程序里也有同样的一壶茶。一次完整的 GPU 计算拆成三步：数据从内存搬到显存（H2D 搬运）、GPU 流处理器做计算（kernel）、结果拷回内存（D2H 回传）。三步串行时，搬运期间 GPU 计算单元完全空转。别以为搬运很快——PCIe 总线带宽远低于显存带宽，数据量大时搬运耗时和 kernel 计算在同一量级，串行方案等于把 GPU 一半的生命浪费在等待上。

流水线的核心思想是把操作拆成独立阶段，让不同阶段在时间上重叠：搬运第 2 批数据的同时，GPU 正在计算第 1 批。设每批搬运耗时 t_copy、计算耗时 t_compute，n 批串行总时间为 n(t_copy + t_compute)，流水线总时间约等于 n×max(t_copy, t_compute) + min(t_copy, t_compute)。当 n 足够大，加速比趋近两阶段耗时之和除以较大者，两阶段等耗时时节加速比逼近 2。

让学徒烧水与师傅泡茶同时进行后，第一壶茶依然要等三分钟，但从第二壶开始，每隔一分钟就有一壶新茶出锅。等待并没有消失，它只是被悄悄藏进了别人的工序里。GPU 流水线同理：单批延迟没变，单位时间处理的批数——吞吐量——却大幅提升。这就是流水线最迷人的地方。

S = \frac{T_{\text{串行}}}{T_{\text{流水线}}} = \frac{n (t_{\text{copy}} + t_{\text{compute}})}{n \times \max(t_{\text{copy}}, t_{\text{compute}}) + \min(t_{\text{copy}}, t_{\text{compute}})}

串行是“两个人轮流干活、总有一个在等”，流水线是“两人各干各的、互相掩护”。批次足够多时，加速比由两阶段是否均衡决定，最理想是翻倍。

串行总时间n 批数据，每批都先搬运再计算，时间依次累加。
$T_{\text{串行}} = n \times (t_{\text{copy}} + t_{\text{compute}})$
流水线总时间第一批搬运完全暴露，之后每批只增加两阶段中较慢者的耗时。
$T_{\text{流水线}} = n \times \max(t_{\text{copy}}, t_{\text{compute}}) + \min(t_{\text{copy}}, t_{\text{compute}})$
加速比极限两者相除，当批次数 n 趋于无穷时，常数项影响消失。
$S \to 1 + \frac{\min(t_{\text{copy}}, t_{\text{compute}})}{\max(t_{\text{copy}}, t_{\text{compute}})} \le 2$

CUDA流基础：多雇几个学徒

流水线的第一步是“让两个人各干各的”。阿茶于是多雇了几个学徒：每个学徒拎一把壶，排队烧水、递水、传话，互不干扰。师傅手里那把壶见底了，另一个学徒的水正好递上来——不需要任何人喊“该你了”。

在 CUDA 里，“学徒”就是 CUDA 流（stream）。流本质是一条独立的 GPU 命令队列：你把搬运命令放进流 A、计算命令放进流 B，GPU 的调度器就会尽量让复制引擎（copy engine）与计算核心（SM）同时忙碌。需要特别小心默认流（default stream）——它有隐式同步特性，会与所有其他流互相等待。想让流水线转起来，必须用 cudaStreamCreate 创建非默认流，把异步操作提交进去。

多流并发的写法很直接：cudaMemcpyAsync(……, streamA) 提交搬运，kernel<<<grid, block, 0, streamB>>>(……) 提交计算。数据有依赖时，用 cudaEventRecord 记录事件、cudaStreamWaitEvent 让计算流等待搬运流。切记：流之间没有依赖才有并发，有依赖必须显式表达先后，否则会出现读到半截数据的竞态 bug。M 条流把 N 批任务均分，理想情况下总耗时降到单流的 1/M。

多雇几个学徒后，茶馆的瓶颈从“没人烧水”变成“灶台不够用”，这就是资源竞争的代价。CUDA 流也一样：流是组织并发的方式，不是免费的加速器，数量要结合 GPU 的计算单元、复制引擎和任务粒度来定。而下一节的主角“双缓冲”，正是让每个学徒手里的水壶时刻都是满的。

S_M = \frac{T_{\text{单流}}}{T_{\text{多流}}} = M

单流是一个学徒包揽所有水壶；多流是 M 个学徒各管几把壶，任务均分后同时开工，理想加速比等于流的数量。同步开销和资源竞争会让实际值略低。

单流时间所有命令在一条队列排队执行，时间是各批之和。
$T_{\text{单流}} = N \times (t_{\text{copy}} + t_{\text{compute}})$
多流时间N 批任务均分给 M 条流，每条流只处理 N/M 批，流间并行。
$T_{\text{多流}} = \frac{N}{M} \times (t_{\text{copy}} + t_{\text{compute}})$
理想加速比两者相除得 M，流数越多理论越快，但调度开销会稀释收益。
$S_M = M$

双缓冲技术：两把水壶轮着用

流水线跑起来后，阿茶又发现新问题：学徒烧好一壶水交给师傅，师傅泡完，学徒才能烧下一壶——两人还是得互相等。阿茶一拍脑袋：买两把壶！学徒先烧满壶 A，师傅泡茶时，学徒立刻烧壶 B；师傅用完 A，B 的水正好开了，学徒转头又把 A 续满。两把壶轮流上岗，师傅手边永远有开水。

这就是双缓冲（double buffering）：在显存里分配两份缓冲，一份供当前 kernel 计算读取，另一份同时由 DMA 引擎写入下一批数据。计算和搬运操作不同的存储区域，互不踩踏。实现上申请两块显存 d_buf[2]，第 i 批数据搬运到 d_buf[i % 2]，kernel 计算 d_buf[(i+1) % 2]——批号每增加一批，两个缓冲的角色就翻转一次，和两把水壶轮流上岗一模一样。

双缓冲的收益可以算一笔时间账：串行总时间 n(t_copy + t_compute)，双缓冲后是 n×max + min，两者之差是 (n-1)×min。也就是说从第二批起，每一批都省下了较快阶段的那段等待。代价也很直观：显存占用翻倍。显存紧张时可以用“分块 + 单缓冲”折中，但吞吐会下降；绝大多数场景下，双缓冲是隐藏传输延迟性价比最高的手段。

两把水壶让茶馆真正实现“烧水不断档”。但阿茶很快意识到，水壶轮换只是基础，真正的功夫在于让学徒和师傅的节奏严丝合缝——这就是计算与通信重叠：不仅是“有水可用”，而是烧水和泡茶同时发出声音。

T_{\text{节省}} = T_{\text{串行}} - T_{\text{双缓冲}} = (n-1) \times \min(t_{\text{copy}}, t_{\text{compute}})

省下的不是“慢的环节”而是“快的环节”——慢的环节本来就要等，快的环节才是被白白浪费掉的时间。

串行总时间n 批数据搬运与计算依次完成。
$T_{\text{串行}} = n(t_{\text{copy}} + t_{\text{compute}})$
双缓冲总时间第一批搬运暴露在外，之后每批只增加较慢阶段耗时。
$T_{\text{双缓冲}} = n \times \max(t_{\text{copy}}, t_{\text{compute}}) + \min(t_{\text{copy}}, t_{\text{compute}})$
节省时间两者相减，n-1 表示从第二批起每批获益，min 表示省下的是较短阶段。
$T_{\text{节省}} = (n-1) \times \min(t_{\text{copy}}, t_{\text{compute}})$

计算通信重叠：边烧水边泡茶

茶馆里最动人的时刻，是学徒往壶 B 灌水时，师傅正提着壶 A 冲泡。烧水声和斟茶声同时响起，两个人都在忙，谁也不等谁。“重叠”的要义不在于等待消失了，而在于没有人在空转。

在 CUDA 里，“通信”指主机与设备之间的数据搬运（cudaMemcpyAsync），“计算”指 kernel。重叠的关键是用异步接口：cudaMemcpyAsync 提交后立刻返回，CPU 不等它完成；kernel 启动同样异步。把搬运放进流 1、kernel 放进流 2，GPU 的复制引擎和计算核心就能同时忙碌。前提是数据必须放在固定内存（pinned memory）里，否则异步拷贝会退化为同步。

真正的难点是表达依赖关系：第 i 批数据搬运完，第 i 批 kernel 才能开始；第 i 批 kernel 算完，第 i 批结果才能回传。做法是：搬运流用 cudaEventRecord 打点，计算流用 cudaStreamWaitEvent 等待事件。依赖表达错乱会出现读半截数据的竞态；全部丢进默认流则隐式同步强制排队，重叠彻底失效。检查重叠是否真实，要用 nsys 或 ncu 看 timeline：搬运条和 kernel 条在时间轴上有没有真正交叠。

烧水声和斟茶声同时响起，客人体验到的是“茶永远端得出来”。但从工程角度看，重叠不是自动发生的，而是靠精确的“事件”编排出来的。就像学徒和师傅之间不用喊话，靠的是水壶位置的默契——CUDA 里，这份默契就是流和事件。

\text{隐藏率} = \frac{\min(t_{\text{copy}}, t_{\text{compute}})}{t_{\text{copy}} + t_{\text{compute}}}

隐藏率衡量“被计算掩护掉的搬运时间”占总搬运时间的比例。两阶段等耗时隐藏率 50%；计算远慢于搬运时，搬运几乎被完全隐藏，隐藏率趋近 100%。

可重叠时间两阶段中较短者的耗时，在它结束前较长者仍在执行，这段时间可以被藏起来。
$\Delta t = \min(t_{\text{copy}}, t_{\text{compute}})$
串行总时间两阶段时间的总和。
$t_{\text{总}} = t_{\text{copy}} + t_{\text{compute}}$
隐藏率可重叠时间占总时间的比例，越高说明重叠越充分、利用率越好。
$\text{隐藏率} = \frac{\Delta t}{t_{\text{总}}}$

多级流水线：拆出更多工序

生意越做越大，阿茶把泡茶也拆开了：温杯的只管温杯，投茶的只管投茶，注水的只管注水，出汤的只管出汤。每个人动作都极简，流水线从“烧水→泡茶→端茶”三级变成六级的细长队形，每一级都在为不同客人的茶忙碌。

GPU 流水线同样可以多级化。除了搬运与计算重叠，还可以把 kernel 拆成预处理、核心计算、后处理多个子 kernel，并把数据切成多个 chunk。这样 chunk 1 的后处理、chunk 2 的核心计算、chunk 3 的预处理可以同时进行。多级流水线的节奏由最慢的一级决定——就像注水最慢时，整条线都得按它的节拍走。

多级流水线调度有三个要点：第一，阶段划分要均衡，最慢级是瓶颈，优化它收益最大；第二，级数不是越多越好，每级都有启动和同步开销，流水线“填满”需要更多批次，小批量任务可能不赚反亏；第三，跨级依赖用事件链管理，每级完成时记录事件，下一级等待后再启动。吞吐极限是每个周期完成一批，周期 T = max_i(t_i)，所以优化目标是缩短最慢级耗时，而非平均用力。

六级流水线让茶馆的产出稳定如钟摆，但阿茶发现：客人太少时，六级工序反而成了负担——人没站满，节奏还乱了。多级流水线是给“稳定大批量”场景准备的利器，任务量小时，简单的串行反而更好。优化永远是场景的函数。

T_{\text{cycle}} = \max_{1 \le i \le m} t_i, \qquad \text{吞吐} = \frac{1}{T_{\text{cycle}}}

一条流水线的速度由最慢的工人决定。想提高整条线的产出，先找到那个“注水最慢”的环节去优化。

每级耗时t_i 表示第 i 级处理一批数据所需时间，m 是级数。
$t_i,\quad i = 1, 2, \dots, m$
流水线周期所有级同时忙碌时，每批产出的间隔等于最慢级的耗时。
$T_{\text{cycle}} = \max_{1 \le i \le m} t_i$
吞吐量单位时间产出的批数，等于周期的倒数，周期越短吞吐越高。
$\text{吞吐} = \frac{1}{T_{\text{cycle}}}$

优化避坑：老师傅的心法

经营半年后，阿茶总结出一套心法：水壶不是越多越好——三个人挤一个灶台反而添乱；水也不是一次烧得越多越好——壶太大烧水慢，客人等得更久。茶馆的快，靠的是各环节匹配，而不是某个环节的蛮力。

GPU 流水线优化也是同样的平衡术。先说流数量：流太多增加调度开销，流太少无法充分重叠，通常从 2~4 条流起步，用 profiler 实测后调整。再看分块大小：数据切成小块流水线才能转起来，但块太小则启动开销占比过高，经验上让每块处理时间落在数十微秒到毫秒量级。最容易被忽视的是同步频率：cudaDeviceSynchronize 是流水线的头号敌人，它一出现，全线重叠瞬间归零。

实战中还有几个高频技巧：用固定内存提升 cudaMemcpyAsync 的实际带宽；把 CPU 端的数据预处理也编排进流水线，让 CPU 与 GPU 同时忙碌；用 CUDA Graphs 把整条流水线的启动开销一次性打包；显存紧张时用二维异步拷贝或流式统一内存访问缓解压力。最后，一切优化都要用数据说话——跑完 nsys 看 timeline，如果搬运条和计算条之间还有明显空隙，那就是还没优化到位。

老茶客都知道，阿茶茶馆真正的秘诀不是某一把壶、某一个学徒，而是整套工序严丝合缝的配合。GPU 性能优化也一样：流水线并行不是某个 API 的开关，而是一种“让每一份硬件都不闲着”的思维方式。掌握了它，你就拿到了 GPU 性能世界的入门钥匙。

U = \frac{n \times t_{\text{compute}}}{n \times \max(t_{\text{copy}}, t_{\text{compute}}) + \min(t_{\text{copy}}, t_{\text{compute}})}

GPU 利用率 = 真正干活的时间 ÷ 总时间。流水线做得好，利用率逼近 1；做不好，利用率就是串行比例。

总计算时间n 批 kernel 的计算时间总和，这是真正“产出”的时间。
$T_{\text{计算}} = n \times t_{\text{compute}}$
流水线总时间与双缓冲总时间公式相同，是整个执行从开始到结束的时长。
$T_{\text{总}} = n \times \max(t_{\text{copy}}, t_{\text{compute}}) + \min(t_{\text{copy}}, t_{\text{compute}})$
利用率两者之比，目标是通过重叠让分母无限接近分子。
$U = \frac{T_{\text{计算}}}{T_{\text{总}}}$

AI 应用场景

大模型训练中的数据预取流水线

阶段：数据准备与加载阶段

解决的问题：深度学习训练时，GPU 每算完一个 mini-batch 就要等待 CPU 准备下一个 batch 的数据。传统 DataLoader 串行执行“读磁盘→预处理→拷进显存”，GPU 利用率常跌破 60%。用流水线并行后，CPU 在 GPU 计算当前 batch 的同时，提前完成下一个 batch 的读取和异步拷贝，GPU 几乎不再空等。

具体案例：某 LLM 微调任务中，训练脚本用固定内存 + 双缓冲 + 两个 CUDA 流，一个流负责 H2D 搬运下一批样本，另一个流执行前向反向 kernel。改造后单卡吞吐从 4.2 万 tokens/秒提升到 7.8 万 tokens/秒，GPU 利用率从 55% 升到 92%。

推理服务的多流批处理

阶段：模型推理部署阶段

解决的问题：在线推理服务通常按请求批次处理，如果每个请求都串行执行“预处理→推理→后处理→回传”，GPU 在等待响应时大量闲置。用多级流水线把预处理、推理、回传放进不同 CUDA 流，让多个请求批次在不同阶段重叠执行。

具体案例：Triton Inference Server 部署 Stable Diffusion 时，通过动态批处理器与多流调度，将图像生成的端到端吞吐提升 1.8 倍，同时保证单请求延迟不显著恶化；视频分析管线把解码、检测、编码拆成多级 kernel，用事件链调度实现多路视频并行处理。

高性能计算中的通信计算重叠

阶段：大规模分布式计算阶段

解决的问题：多卡训练中，梯度同步（AllReduce）通信常常让 GPU 计算停顿。通过把通信拆成小块并与下一轮前向计算放入不同流，实现通信与计算的深度重叠。

具体案例：某推荐系统模型使用 8 卡训练，梯度切分成 4 个 chunk，每个 chunk 的 AllReduce 用独立流异步发起，同时 GPU 继续计算后续 chunk 的前向。最终训练吞吐提升约 40%，与 NVIDIA 官方文档中“通信计算重叠”的最佳实践一致。

Python 演示

python38 行

"""模拟双缓冲流水线：对比串行调度与流水线调度的总耗时."""import time  def run_serial(batches, copy_time, compute_time):    """串行模式: 每批先搬运后计算, 时间为各环节之和."""    return batches * (copy_time + compute_time)  def run_pipeline(batches, copy_time, compute_time):    """双缓冲流水线: 总时间 = n*max + min."""    return batches * max(copy_time, compute_time) + min(copy_time, compute_time)  def show_roles(batches):    """打印两把水壶(缓冲)的角色翻转过程."""    for i in range(batches):        buf_a = "搬运" if i % 2 == 0 else "计算"        buf_b = "计算" if i % 2 == 0 else "搬运"        print(f"第{i+1:2d}批 -> 缓冲A:{buf_a}  缓冲B:{buf_b}")  def main():    batches = 10    copy_time = 2.0    compute_time = 3.0    serial = run_serial(batches, copy_time, compute_time)    pipe = run_pipeline(batches, copy_time, compute_time)    print(f"批数={batches} 搬运={copy_time} 计算={compute_time}")    print(f"串行耗时  : {serial:6.1f}")    print(f"流水线耗时: {pipe:6.1f}")    print(f"加速比    : {serial / pipe:5.2f}")    print(f"GPU利用率 : {batches * compute_time / pipe:5.2%}")    show_roles(batches)  if __name__ == "__main__":    main()

目的：用纯 Python 模拟 CUDA 双缓冲流水线的时序模型，直观展示串行与流水线的耗时差异，以及缓冲角色按批号翻转的规律。

return batches * (copy_time + compute_time)串行模式把每批的搬运与计算依次相加。
return batches * max(copy_time, compute_time) + min(copy_time, compute_time)流水线总时间：每批只增加较慢环节，再加上初始暴露的较快环节。
print(f"加速比 : {serial / pipe:5.2f}")输出加速比，验证两阶段等耗时极限接近 2 的结论。
print(f"GPU利用率 : {batches * compute_time / pipe:5.2%}")计算时间占总时间比例，代表 GPU 真正干活的时间占比。
buf_a = "搬运" if i % 2 == 0 else "计算"核心翻转逻辑：批号奇偶决定缓冲 A/B 各自承担的角色。

常见误区

误区听说流水线能加速，于是把代码里的拷贝和 kernel 随便扔进几个流，觉得马上就能快一倍。跑完发现不但没快，反而变慢了，十分困惑。

正解流水线加速有前提：搬运耗时必须与计算耗时在同一个量级，才有足够的“空隙”可隐藏；流之间的依赖必须正确表达，同步点会抹掉全部重叠；小批量任务来不及填满流水线，启动开销反而占大头。就像茶馆里客人只有一两位时，多雇学徒只会增加成本。

为何常见：直觉上“并行 = 快”，忽略了流水线需要足够的批量规模来摊薄启动成本，也忽略了同步对并发结构的破坏力。

误区流的数量越多，GPU 就越忙，速度就越快。有人一口气创建 32 个流，结果性能反而下降，还占用大量显存。

正解流是命令队列，不是线程。GPU 的计算单元和复制引擎数量有限，流太多只会增加调度开销和资源排队。通常从 2~4 条流起步，结合 profiler 实测调整——就像灶台只有两个火眼，雇十个学徒也是轮流用。

为何常见：把 CPU 多线程“线程越多并行度越高”的直觉直接搬到了 GPU 流上，忽略了硬件并发上限。

误区显存很贵，双缓冲要分配两份数据空间，觉得太浪费，宁愿用单缓冲省着点，结果 GPU 老是等数据。

正解双缓冲是隐藏传输延迟最直接的手段，省下的等待时间远比显存开销值钱。显存紧张时可用“分块 + 单缓冲”折中，但代价是吞吐下降。就像茶馆买两把壶看似多花钱，但翻台率提高赚回的钱远多于壶钱。

为何常见：只看到内存成本，没有把时间成本算进总账；双缓冲的本质是用空间换时间，评估时要算投入产出比。

误区看到 cudaMemcpyAsync 名字里有 Async，就觉得搬运和 kernel 已经自动并行了，再也没管过同步。

正解异步只是“不阻塞 CPU”，不代表“与 kernel 并发”。如果拷贝和 kernel 都在默认流，隐式同步会强制它们排队；即使分到不同流，也要用事件显式表达依赖，否则要么读到脏数据，要么悄悄退化成串行。用 nsys 看时间线才能确认重叠是否真的发生。

为何常见：API 名字里的 Async 给人“万事大吉”的错觉，实际上并发需要流的隔离和事件编排共同保证。

面试题

Q1. 流水线并行能提升性能的根本原因是什么？

单选题 · easy

A.提高了 GPU 核心的时钟频率
B.让不同阶段在时间上重叠执行← 正确答案
C.减少了每批数据的计算总量
D.将数据搬运压缩到零延迟

答案：B。流水线不降低单批延迟，也不改变计算量，而是通过让搬运与计算重叠，把 GPU 的空闲时间利用起来。提高频率、减少计算量、把搬运压到零都不现实。就像茶馆提速靠学徒和师傅同时忙，而不是让师傅手速翻倍。

Q2. CUDA 流（stream）的本质是什么？

单选题 · mid

A.一条独立有序的命令队列← 正确答案
B.GPU 上的一个物理计算核心
C.一块专用的显存缓冲区域
D.CPU 上的一个后台线程池

答案：A。流是 GPU 命令队列，是组织并发的一种逻辑结构，不是物理核心，也不是内存区域，更不是 CPU 线程池。把流理解成“学徒的工位和排班表”，真正干活的是 GPU 里的硬件单元。

Q3. 双缓冲中两份缓冲的角色如何切换？

单选题 · easy

A.一份固定搬运一份固定计算
B.按批号奇偶交替使用缓冲区← 正确答案
C.根据显存剩余量动态分配
D.计算结束后销毁再重新创建

答案：B。双缓冲中两份缓冲地位对等，按批号奇偶轮换角色——偶数批用缓冲 A 搬运、缓冲 B 计算，奇数批互换。固定分工会让一半缓冲闲置；动态分配和销毁重建都有额外开销，都不是双缓冲的做法。

Q4. 影响计算与通信重叠效果的因素有哪些？

多选题 · mid

A.搬运时间与计算时间的相对大小← 正确答案
B.是否使用独立的非默认流← 正确答案
C.是否调用了全局同步点← 正确答案
D.GPU 风扇的散热效果好坏

答案：ABC。重叠效果取决于搬运与计算的耗时比例（决定隐藏率上限）、是否用非默认流（决定能否并发）、以及是否有同步点（决定重叠是否被打断）。风扇散热是硬件环境问题，与任务级流水线是否重叠无关。

Q5. cudaMemcpyAsync 需要配合什么才能真正实现异步？

单选题 · mid

A.默认流配合同步等待函数
B.非默认流和固定内存配合← 正确答案
C.多块 GPU 卡的交错使用
D.将 kernel 改为单线程执行

答案：B。cudaMemcpyAsync 的异步行为依赖两个条件：放入非默认流以避开隐式同步；数据放在固定内存（pinned memory），DMA 才能高效异步访问。默认流会让操作退化为同步；多卡交错和单线程内核与异步拷贝无关。

Q6. 两阶段耗时相等时，流水线加速比的极限接近多少？

单选题 · hard

A.1 倍
B.1.5 倍
C.2 倍← 正确答案
D.n 倍（批次数）

答案：C。两阶段等耗时（t_copy = t_compute）时，流水线每批只增加一个阶段耗时，总时间趋近 n×t，而串行是 2n×t，加速比极限为 2。因为最慢环节决定节奏，加速比不可能无限增长，更不可能等于批次数。

Q7. 多级流水线调度需要遵守哪些原则？

多选题 · hard

A.各级耗时尽量均衡以减小瓶颈← 正确答案
B.依赖关系用事件显式表达← 正确答案
C.阶段拆得越细收益一定越大
D.小批量任务也适合多级流水线

答案：AB。多级流水线要求各级均衡（瓶颈决定周期）、依赖用事件链显式表达。阶段拆得过细会产生启动与同步开销，小批量任务来不及填满流水线反而变慢，所以 C、D 错误。就像茶馆客人太少时，六级工序只会添乱。

Q8. “气泡”（bubble）在流水线中指什么？

单选题 · easy

A.显存中未被使用的空闲区域
B.流水线中某级空闲等待的时间← 正确答案
C.GPU 温度过高时的降频现象
D.数据拷贝时产生的错误帧

答案：B。“气泡”指流水线中某级无事可做的空闲时间，像茶水断档时师傅干等。它不是显存空洞、不是降频、更不是错误帧。消除气泡正是流水线优化的核心目标——让每一级都始终有活干。

Q9. 哪些操作适合放入不同 CUDA 流以重叠执行？

多选题 · mid

A.主机到设备的异步数据搬运← 正确答案
B.相互独立的不同 kernel 计算← 正确答案
C.对同一缓冲区的读写操作
D.设备到主机的异步结果回传← 正确答案

答案：ABD。搬运、回传和相互独立的 kernel 之间没有数据依赖，可以放入不同流并发执行。对同一缓冲区的读写是数据竞争，必须先同步再操作，绝不能并发。就像两个学徒不能同时往一个壶里灌水和倒水。

Q10. 判断流水线是否真正生效的可靠手段是什么？

单选题 · hard

A.查看代码中流对象的数量
B.用剖析工具观察时间线重叠← 正确答案
C.简单比较程序编译的快慢
D.检查内核函数命名的前缀

答案：B。判断重叠是否真实的唯一可靠手段是看剖析工具（如 nsys）的 timeline：搬运条和 kernel 条在时间轴上是否同时存在。数流数量、比编译时间、看函数名都是间接猜测，无法反映实际调度结果。