CPU-GPU 异构流水线
异构计算不是“把代码扔给 GPU”，而是让 CPU 与 GPU 各自承担合适阶段，并控制数据移动和依赖。
text
input -> CPU parse/control -> H2D -> GPU bulk compute -> D2H -> CPU output
^ |
+---- scheduling ----+CPU 适合复杂分支、系统调用、少量延迟敏感工作；GPU 适合规模大、规则、算术密度高的批处理。边界应按阶段特征与数据位置划分，而不是按代码行数平均切割。
1. 先算值不值得
端到端时间：
text
T_total = T_prepare + T_H2D + T_kernel + T_D2H + T_sync + T_finish即使 Kernel 比 CPU 版本快 20 倍，若传输和同步占大头，整体可能只快一点甚至更慢。尽可能让数据留在设备端，连续执行多个阶段后再取回。
2. 分块流水
大输入可切成 Chunk，并用多个 Stream 形成流水：
text
time --->
copy engine: H2D A | H2D B | D2H A | H2D C | D2H B
compute: kernel A | kernel B | kernel CChunk 太小会增加启动与管理开销，太大则降低重叠机会并占用更多内存。需要测量寻找折中。
3. Unified Memory 不是“没有数据移动”
统一内存提供统一地址空间和自动迁移能力，降低编程负担，但页面仍可能在 CPU/GPU 间迁移。访问模式不佳会引发缺页与抖动。它改变管理方式，并没有消除物理层的数据位置和带宽限制。
Pinned host memory 可提高并支持某些异步传输，但它占用不可分页资源，过量使用会影响系统，适合有边界地复用缓冲区。
4. 错误与同步
CUDA 工作常异步，错误可能在后续同步点才暴露。工程代码应分别检查启动错误和异步执行结果，并用事件表达跨 Stream 依赖，避免用全设备同步把所有并发都串行化。
5. Amdahl 定律
若不可加速比例为 s，其余部分加速 p 倍：
text
speedup = 1 / (s + (1-s)/p)当 s = 0.1，即使 GPU 部分无限快，整体加速也不超过 10 倍。这解释了为什么应优化端到端路径，而非只展示最快 Kernel。
深入原理与工程实践
前面的内容负责建立统一心智模型；下面把同一主题继续拆到执行过程、代码、性能代价与工程判断。
<!-- migrated-deep-dive:start -->
完整迁入：原 CPU-GPU 异构计算全文
异构计算——CPU 与 GPU 如何协同工作 / Heterogeneous Computing with CPUs and GPUs
📅 创建时间：2026-06-02 🏷️ 标签：#异构计算 #CPU #GPU #流水线 #异步执行 #PinnedMemory 📚 前置知识：[[04-cuda-programming-model]]（CUDA 编程） [[08-mpi-cluster-hpc]]（集群通信） 📚 相关知识：[[10-dl-training-optimization]]（训练优化实战）
先抓住直觉
CPU 与 GPU 不是竞争关系，而是一条流水线的不同工位。最差的情况是 CPU 准备数据时 GPU 闲着，GPU 计算时传输通道又闲着；协同优化就是让准备、搬运和计算尽量重叠。
- 必须理解：异步执行、流水线、Pinned Memory 为何能帮助传输。
- 用到再查：Stream API、Unified Memory 提示函数和双缓冲模板。
- 读完能回答：为什么 GPU 利用率低不一定是 GPU Kernel 的问题？
场景：CPU 和 GPU 谁才是瓶颈？
┌─────────────────────────────────────────────────────────────┐
│ │
│ 训练一个 epoch，耗时分析： │
│ │
│ GPU Forward/Backward： 10 秒（90%） │
│ CPU 数据加载： 1 秒（9%） │
│ CPU→GPU 数据传输： 0.1 秒（1%） │
│ │
│ 看起来 GPU 占据了大部分时间。 │
│ 但为什么 GPU 利用率只有 60%？ │
│ │
│ 原因：CPU 数据加载太慢，GPU 经常在等待数据！ │
│ │
│ 本章学习如何让 CPU 和 GPU 协同工作，最大化效率。 │
│ │
└─────────────────────────────────────────────────────────────┘第1节：异构系统的基本架构
CPU-GPU 数据流
┌─────────────────────────────────────────────────────────────┐
│ CPU-GPU 异构系统 │
├─────────────────────────────────────────────────────────────┤
│ │
│ CPU │
│ ┌─────────────────────────────────────────────────────┐ │
│ │ │ │
│ │ 数据加载 数据预处理 调度控制 日志记录 │ │
│ │ (SSD) (OpenMP) (CUDA API) │ │
│ │ │ │
│ └────────────────────────────┬────────────────────────┘ │
│ │ PCIe │
│ ↓ │
│ ┌────────────────────────────┐ │
│ │ GPU 显存 (HBM) │ │
│ │ ┌──────────────────────┐ │ │
│ │ │ Forward | Backward │ │ ← GPU 计算 │
│ │ │ Kernel | Kernel │ │ │
│ │ └──────────────────────┘ │ │
│ └────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────┘
关键问题：
1. 数据从 SSD → CPU → GPU，链路长
2. 如果 CPU 太慢，GPU 等待
3. 需要 CPU 和 GPU 并行工作异步执行原理
cpp
// ❌ 同步执行：CPU 等 GPU，浪费时间
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
kernel<<<...>>>(d_data); // GPU 执行时 CPU 空闲
cudaMemcpy(h_result, d_result, size, cudaMemcpyDeviceToHost);
// ================================================================
// ✅ 异步执行：CPU 和 GPU 并行工作
// ================================================================
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// CPU 做数据加载（stream1）
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream1);
// GPU 执行 kernel（stream1）
kernel<<<blocks, threads, 0, stream1>>>(d_data);
// CPU 同时做其他事情（不阻塞！）
preprocess_next_batch_cpu(); // CPU 预处理下一批
log_to_disk(); // CPU 写日志
// 最后才等待 GPU 完成
cudaStreamSynchronize(stream1);第2节：Pinned Memory——加速 CPU-GPU 传输
Pageable vs Pinned Memory
┌─────────────────────────────────────────────────────────────┐
│ Pinned Memory（页锁定内存） │
├─────────────────────────────────────────────────────────────┤
│ │
│ Pageable Memory（普通内存）： │
│ - 操作系统可以把这块内存换页到磁盘（swap） │
│ - CPU 访问可能触发缺页中断 │
│ - GPU 传输前需要先拷贝到临时缓冲区 │
│ → 传输慢 │
│ │
│ Pinned Memory（页锁定内存）： │
│ - 永远保持在物理内存，不会被换页 │
│ - GPU 可以直接传输，不需要临时拷贝 │
│ → 传输快 2-3 倍！但会占用系统内存 │
│ │
└─────────────────────────────────────────────────────────────┘Pinned Memory 代码
cpp
// 普通方式（pageable memory）
float* h_data = (float*)malloc(size); // pageable
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// ================================================================
// Pinned Memory 方式
// ================================================================
float* h_data_pinned;
cudaMallocHost(&h_data_pinned, size); // 分配页锁定内存
// 使用方式一样，但传输更快
cudaMemcpy(d_data, h_data_pinned, size, cudaMemcpyHostToDevice);
// 释放
cudaFreeHost(h_data_pinned);异步传输 + Pinned Memory
cpp
// 最佳实践：Pinned Memory + Async + Double Buffering
cudaStream_t stream;
cudaStreamCreate(&stream);
// 两套缓冲区，轮换使用
float *d_buf[2], *h_buf[2];
for (int i = 0; i < 2; i++) {
cudaMalloc(&d_buf[i], size);
cudaMallocHost(&h_buf[i], size); // pinned memory
}
int buf_idx = 0;
bool first_iteration = true;
// 主循环
while (has_more_data()) {
// CPU 加载下一批数据到当前缓冲区
load_data(h_buf[buf_idx], size);
// CPU→GPU 异步拷贝（数据加载完成后）
cudaMemcpyAsync(d_buf[buf_idx], h_buf[buf_idx],
size, cudaMemcpyHostToDevice, stream);
// 如果不是第一次，等待上一个 kernel 完成
if (!first_iteration) {
int prev_idx = 1 - buf_idx;
cudaStreamSynchronize(stream); // 等待前一批 GPU 计算完成
// 处理前一批的结果
process_results(d_buf[prev_idx]);
}
// 启动当前批次的 kernel（与下一次 CPU 加载并行！）
kernel<<<blocks, threads, 0, stream>>>(d_buf[buf_idx], size);
buf_idx = 1 - buf_idx; // 切换缓冲区
first_iteration = false;
}第3节：CUDA Streams 实现流水线
数据加载流水线
┌─────────────────────────────────────────────────────────────┐
│ 3-Stage 流水线 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Time → │
│ │
│ Stage 1: [Load GPU0][Load GPU1][Load GPU2][Load GPU3]...│
│ Stage 2: [Preprocess][Preprocess][Preprocess]... │
│ Stage 3: [GPU Fwd][GPU Fwd][GPU Fwd]... │
│ │
│ 理想情况：3 个阶段并行，满载运转 │
│ │
│ 如果数据加载太慢（Stage 2 瓶颈）： │
│ Stage 1: [Load][Load][Load][Load][Load][Load]... │
│ Stage 2: [Pre ][Pre ][Pre ][Pre ][Pre ][Pre ] │
│ Stage 3: [Fwd ][Fwd ][Fwd ][Fwd ] │
│ ─────────→ GPU 大量空闲等待 │
│ │
└─────────────────────────────────────────────────────────────┘多 Stream 实现流水线
cpp
// 3-Stream 流水线实现
cudaStream_t stream_load, stream_preprocess, stream_compute;
cudaStreamCreate(&stream_load);
cudaStreamCreate(&stream_preprocess);
cudaStreamCreate(&stream_compute);
// 使用事件控制流水线顺序
cudaEvent_t preprocess_done, compute_done;
int batch = 0;
while (true) {
// Stage 1: CPU 加载数据到 pinned memory
load_data_cpu(h_pinned[batch % 2]);
// 等待 preprocess 完成，才能开始下一次加载
if (batch > 0) {
cudaEventSynchronize(preprocess_done);
}
// 异步拷贝到 GPU
cudaMemcpyAsync(d_pinned[batch % 2], h_pinned[batch % 2],
size, cudaMemcpyHostToDevice, stream_load);
// Stage 2: GPU 预处理（resize, normalize 等）
preprocess<<<blocks, threads, 0, stream_preprocess>>>(
d_pinned[batch % 2], d_preprocessed[batch % 2]);
cudaEventRecord(preprocess_done, stream_preprocess);
// 等待 preprocess 完成，才能开始计算
cudaEventSynchronize(preprocess_done);
// Stage 3: GPU 训练
train<<<blocks, threads, 0, stream_compute>>>(
d_preprocessed[batch % 2], model_params);
cudaEventRecord(compute_done, stream_compute);
batch++;
}第4节：统一内存的异构使用
┌─────────────────────────────────────────────────────────────┐
│ Unified Memory 在异构场景的应用 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 场景：CPU 和 GPU 交替访问同一块数据 │
│ │
│ 例如： │
│ - CPU 加载数据 → GPU 处理 → CPU 读取结果 → CPU 写日志 │
│ - 循环很多次 │
│ │
│ 传统方式：每次都需要 cudaMemcpy │
│ Unified Memory：自动迁移，无需手动管理 │
│ │
└─────────────────────────────────────────────────────────────┘cpp
// Unified Memory 方式
float* data;
cudaMallocManaged(&data, size);
// CPU 处理
preprocess_cpu(data, size);
// GPU 处理（OS 自动把数据迁移到 GPU）
cudaMemcpy(data, size, cudaMemcpyHostToDevice); // 显式提示迁移
train<<<blocks, threads>>>(data, size);
cudaDeviceSynchronize();
// CPU 读取结果（OS 自动把数据迁移回 CPU）
read_results_cpu(data, size);
// 性能提示：预取到 GPU
cudaMemPrefetchAsync(data, size, 0); // deviceId=0 = GPU
train<<<blocks, threads>>>(data, size);
cudaDeviceSynchronize();
// 预取回 CPU
cudaMemPrefetchAsync(data, size, cudaCpuDeviceId);
read_results_cpu(data, size);第5节：CPU 和 GPU 的任务分配策略
┌─────────────────────────────────────────────────────────────┐
│ CPU vs GPU 任务分配 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 适合 CPU： │
│ - IO 密集型：磁盘读写、网络通信 │
│ - 分支密集型：if-else 多的逻辑 │
│ - 串行任务：数据预处理、格式转换 │
│ - 任务量小的任务：启动 GPU 开销不划算 │
│ │
│ 适合 GPU： │
│ - 计算密集型：矩阵乘法、卷积 │
│ - 数据并行：大量相同操作 │
│ - 算术强度高：计算量 >> 内存访问量 │
│ │
│ 实际训练中的分配： │
│ ┌─────────────────────────────────────────────────────┐ │
│ │ CPU: 数据加载 → Tokenize → Resize → Normalize │ │
│ │ ↓ (Pinned Memory + Async) │ │
│ │ GPU: Embedding → Transformer → Loss → Backward │ │
│ │ ↑ (CUDA Stream 并行) │ │
│ │ CPU: 日志记录 → Checkpoint → Metrics 计算 │ │
│ └─────────────────────────────────────────────────────┘ │
│ │
│ 核心原则：CPU 和 GPU 永远不要同时闲着！ │
│ │
└─────────────────────────────────────────────────────────────┘"AI 可查 vs 必须理解"清单
AI 可查：
✅ cudaStreamCreate/cudaEventCreate 的具体参数
✅ Pinned Memory 的具体使用限制
✅ Unified Memory prefetch 的具体 API
必须理解：
🔴 CPU-GPU 数据传输是瓶颈，需要异步化
🔴 Pinned Memory 比普通内存传输快 2-3 倍
🔴 Double Buffering：两套缓冲区轮换，CPU 和 GPU 并行工作
🔴 CUDA Stream 实现流水线：Load → Preprocess → Compute
🔴 任务分配：CPU 做 IO/预处理，GPU 做计算密集型任务学习状态：🟡 开始学习
完整迁入：原并行计算版异构计算
异构计算：CPU、GPU、NPU 如何协同工作 / Heterogeneous Computing with CPUs, GPUs, and NPUs
📅 创建时间：2026-07-20 🏷️ 标签：#异构计算 #CPUGPU #NPU #任务调度 📚 前置知识：[[07-parallel-patterns]] 📚 相关知识：[[/02-systems-and-performance/02-computer-architecture-and-hardware/09-heterogeneous-computing]] [[/02-systems-and-performance/02-computer-architecture-and-hardware/12-npu-landscape]]
1. 为什么需要异构系统
不存在对所有任务都最优的处理器：
- CPU 擅长控制流、串行逻辑和低延迟
- GPU 擅长规则的大规模数据并行
- NPU 擅长特定张量算子和低精度计算
- FPGA 擅长确定性数据流和定制硬件流水线
异构计算的目标，是让每类工作运行在最适合的设备上，同时控制数据迁移成本。
2. 一个典型异构流水线
text
CPU：读取文件、解析格式
↓
CPU：构建批次、数据预处理
↓ PCIe / 共享内存
GPU：矩阵计算或仿真核心
↓
CPU：业务判断、结果整理
↓
GPU：可视化或后处理如果每一步都频繁来回传输小数据，GPU 的计算收益可能被传输抵消。
3. 任务放置的四个问题
- 任务是否有足够并行度？
- 数据当前位于哪块内存？
- 迁移数据需要多长时间？
- 设备是否支持所需精度、操作和容量？
粗略决策：
text
收益 = 设备计算节省时间 - 数据迁移 - 启动与同步开销只有收益显著为正时，迁移到加速器才值得。
4. 数据管理模式
显式拷贝
程序明确分配 Host/Device 内存并传输。控制最清楚，优化空间最大，但代码复杂。
统一虚拟地址
CPU 和 GPU 使用统一地址空间，但物理数据仍可能需要迁移。
Unified Memory
运行时按需迁移页面，简化开发。访问模式不佳时会产生频繁 Page Fault，因此便利不等于没有成本。
零拷贝与共享内存
某些集成式架构让 CPU 和 GPU 共享物理内存，减少显式拷贝，但仍要考虑缓存一致性、带宽竞争和同步。
5. 计算与通信重叠
把数据分成多个批次：
text
时间 →
传输批次 0 | 传输批次 1 | 传输批次 2
计算批次 0 | 计算批次 1 | 计算批次 2
回传批次 0 | 回传批次 1实现重叠需要：
- 异步传输
- 多个 Stream 或队列
- 独立缓冲区
- 足够大的任务粒度
- 避免不必要的全局同步
6. 可移植编程模型
| 模型 | 特点 |
|---|---|
| CUDA | NVIDIA 生态成熟、控制力强 |
| HIP | 面向 AMD，并提供部分 CUDA 迁移能力 |
| SYCL | 基于现代 C++ 的跨设备模型 |
| OpenCL | 开放、设备广泛，开发体验较底层 |
| OpenMP Offload | 用指令扩展把部分循环卸载到设备 |
可移植性通常会牺牲部分特定硬件优化空间。工程上应先明确目标设备、生命周期和性能要求。
7. 调度与资源竞争
实际系统可能同时运行多个模型或求解任务，需要处理：
- 显存分配和碎片
- Stream 优先级
- 多进程共享 GPU
- CPU 数据线程抢占
- PCIe 和内存带宽竞争
- 设备故障与任务回退
异构系统优化不能只看单个 Kernel，还要观察端到端流水线。
核心总结
- 异构计算是任务放置和数据管理问题。
- 最快设备不一定带来最快端到端系统。
- 数据应尽量长期停留在使用它的设备附近。
- 批处理、异步执行和双缓冲可帮助重叠通信与计算。
- 可移植性、性能和开发成本需要共同权衡。
下一篇：[[09-distributed-parallelism]] <!-- migrated-deep-dive:end -->
面试速答
什么任务适合 GPU offload？ 并行度高、访问规则、计算量足以覆盖传输和启动成本，且能让数据在设备端复用的任务。
统一内存是否等于零拷贝？ 不是。它统一寻址并自动管理迁移，物理页面的数据移动与一致性成本仍存在。
自测
- Kernel 快 10 倍为何应用可能只快 1.2 倍？
- 为什么频繁
cudaDeviceSynchronize()会伤害性能？ - 分块是否越小越容易重叠、越好？
答案：非 Kernel 阶段占比高；它建立全局等待并破坏异步流水；不是，过小会放大启动和调度成本。