6. 主机侧：拷贝、流、事件与多卡

kernel 写得再好，数据和它之间还隔着一条 PCIe。这一节讲主机侧的账：一次拷贝有多贵、怎么让拷贝和计算重叠、怎么正确地计时。很多“GPU 没提速”的案例，问题根本不在 kernel。

6.1 分离的内存

显存和主机内存是两套物理内存。设备指针在主机上解引用会崩，主机指针也不能直接在 kernel 里用（除非用了统一内存或 zero-copy）。完整骨架长这样：

host.cu

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

#define CK(x) do { cudaError_t e_ = (x); if (e_ != cudaSuccess) { \
    fprintf(stderr, "%s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(e_)); \
    std::exit(1); } } while (0)

int main() {
    const size_t bytes = 64ull << 20;              // 64 MiB
    float *h = nullptr, *d = nullptr;

    CK(cudaHostAlloc((void**)&h, bytes, cudaHostAllocDefault));  // pinned
    CK(cudaMalloc((void**)&d, bytes));

    CK(cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice));
    kernel<<<(bytes / 4 + 255) / 256, 256>>>(d, bytes / 4);
    CK(cudaGetLastError());
    CK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost));

    CK(cudaFree(d));
    CK(cudaFreeHost(h));                            // pinned 用 cudaFreeHost
}

cudaMalloc 是同步的，而且可能很慢：驱动要建立页表映射，几 MB 的分配经常是几百微秒。不要在热路径里反复分配；用一次分配、重复使用，或者用 cudaMallocAsync 的池（CUDA 11.2+）。

UVA（统一虚拟地址）让主机指针和设备指针落在同一个地址空间，可以用 cudaPointerGetAttributes 判断一个指针属于谁。它不改变物理上的分离，只是省掉了“这个指针是哪边的”这类手工记账。

6.2 为什么必须用 pinned memory

普通 malloc 的页可能被换出，DMA 引擎不能直接拿它的物理地址。驱动只能先把数据拷到一块临时的锁定页，再发起传输——于是你的“异步”拷贝退化成同步，带宽也打折。用 cudaHostAlloc 或 cudaHostRegister 锁住页之后，DMA 可以直接读写：

// 错误：pageable 内存 + Async，实际是同步的
cudaMemcpyAsync(d, h_pageable, bytes, cudaMemcpyHostToDevice, s);

// 正确：pinned 内存，DMA 真正异步
cudaHostAlloc((void**)&h, bytes, cudaHostAllocDefault);
cudaMemcpyAsync(d, h, bytes, cudaMemcpyHostToDevice, s);

代价是锁住的页不能被换出、也不能被 OS 随意挪动，锁太多会拖累整个系统。通常只锁“正在流式传输”的那几块缓冲区。PCIe 4.0 x16 的单向带宽约 32 GB/s，5.0 x16 约 64 GB/s；H2D 和 D2H 可以同时进行，总线和双向计数时要注意别重复算。

6.3 流：把拷贝和计算叠起来

stream 是一个有序队列：同一个流里的操作按顺序执行，不同流之间可以并发。把数据切成几块，每块用独立的流做“拷贝进 → 计算 → 拷贝出”，就能让 DMA 和 SM 同时工作：

streams.cu

const int nStreams = 4;
cudaStream_t s[nStreams];
for (int i = 0; i < nStreams; ++i) CK(cudaStreamCreate(&s[i]));

size_t n = bytes / sizeof(float);
size_t chunk = n / nStreams;
for (int i = 0; i < nStreams; ++i) {
    size_t off = i * chunk;
    size_t cnt = (i == nStreams - 1) ? (n - off) : chunk;   // 最后一块收尾
    CK(cudaMemcpyAsync(d + off, h + off, cnt * sizeof(float),
                       cudaMemcpyHostToDevice, s[i]));
    kernel<<<(cnt + 255) / 256, 256, 0, s[i]>>>(d + off, cnt);
    CK(cudaMemcpyAsync(h + off, d + off, cnt * sizeof(float),
                       cudaMemcpyDeviceToHost, s[i]));
}
for (int i = 0; i < nStreams; ++i) CK(cudaStreamSynchronize(s[i]));
for (int i = 0; i < nStreams; ++i) CK(cudaStreamDestroy(s[i]));

三个前提同时满足才有重叠：内存是 pinned 的、传输和计算在不同的流、它们访问的缓冲区不冲突。切块大小也有讲究：太小（< 1 MB）被 launch 开销吃掉，太大（> 几十 MB）末尾的尾巴盖不住。经验值 1–16 MB 一块。

6.4 用事件计时

主机计时器会被调度抖动污染，而且 cudaDeviceSynchronize() 之前的代码并不代表 GPU 真的开始/结束。用 CUDA event 在流里打点：

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

CK(cudaEventRecord(start, s));
for (int i = 0; i < iters; ++i) kernel<<<blocks, 256, 0, s>>>(d, n);
CK(cudaEventRecord(stop, s));
CK(cudaEventSynchronize(stop));

float ms = 0.0f;
CK(cudaEventElapsedTime(&ms, start, stop));
printf("%.3f ms/iter\n", ms / iters);

注意循环里不要放同步，否则测到的是“每次 launch + 同步”的和，不是吞吐。事件本身也有约 1–2 µs 的开销，测 5 µs 的 kernel 时要多跑几次取平均。

6.5 统一内存与 zero-copy

cudaMallocManaged 返回一块“主机和设备都能访问”的指针，第一次在设备上访问时按页迁移。它省掉了显式拷贝，但：

• 首次访问触发页错误，延迟是微秒级，可能比拷贝还慢；
• 主机和设备来回访问会反复迁移（抖动）；
• 要性能就得手动 cudaMemPrefetchAsync 预取，等于把拷贝又写回来了。

zero-copy（映射的 pinned 内存）让 GPU 直接通过 PCIe 读主机内存，适合“只读一次、用完就丢”的数据；每次访问都要过 PCIe，延迟高、带宽低。

GH200 这类用 NVLink-C2C 连接 CPU/GPU 的机器让边界变模糊，但编程模型仍然是分离的：缓存一致性解决的是正确性，不是带宽。

6.6 多卡

多卡的基本操作是“每张卡一个上下文 + 显式指定设备”：

int can = 0;
CK(cudaDeviceCanAccessPeer(&can, 1, 0));    // 1 能否直接访问 0 的显存
if (can) {
    CK(cudaSetDevice(1));
    CK(cudaDeviceEnablePeerAccess(0, 0));
}

同一台机器上，NVLink 的卡间带宽约 900 GB/s（H100），PCIe 只有 32–64 GB/s。多卡训练/推理的集合通信（all-reduce、all-gather）用 NCCL，不要手写；它会在 NVLink / PCIe / IB 之间自动选路径。

6.7 和 CPU 的对照

6.8 容易写错的地方