3. 设备内存与主机内存的数据搬运

在 GPU 计算中，真正决定端到端性能的往往不是核函数本身，而是主机（Host）与设备（Device）之间的数据搬运。PCIe 链路的双向带宽通常只有几十 GB/s，而 GPU 内部显存带宽可达数百 GB/s 甚至 TB/s 级，因此一次“拷贝—计算—拷回”的向量加法，很可能大部分时间都花在搬运上。本节围绕 cudaMemcpy、cudaMemcpyAsync、cudaMemcpy2D/3D 三类接口，讲清同步与异步语义、页锁定内存的作用，并完成一个可复现的向量加法端到端实验。

三类拷贝接口的定位

cudaMemcpy 是最基础的同步拷贝：调用返回时，数据已经完成搬运。它的方向由 cudaMemcpyKind 指定，常用 cudaMemcpyHostToDevice（H2D）、cudaMemcpyDeviceToHost（D2H）、cudaMemcpyDeviceToDevice（D2D）。同步语义的代价是：在拷贝期间主机线程被阻塞，无法提交后续工作，也无法与核函数执行重叠。

cudaMemcpyAsync 是异步拷贝：调用立即返回，真正的搬运在指定流（stream）中排队执行。它必须配合页锁定内存（pinned memory）才能获得真正的异步与高带宽；若源或目标是普通可分页内存，驱动往往需要先做一次隐式暂存拷贝，异步性会被削弱。异步拷贝的完成由流顺序或事件（event）保证，而不是由函数返回保证。

cudaMemcpy2D 与 cudaMemcpy3D 用于带步长（pitch）的二维、三维数据。它们的核心参数是 dpitch/spitch（目标/源每行字节数，含填充）与 width/height/depth（实际有效数据范围）。当矩阵行有对齐填充、或需要搬运图像/体数据的子区域时，必须用它们，否则逐行调用 cudaMemcpy 会带来大量小拷贝开销。

同步与异步语义的关键区别

同步拷贝的“完成”由函数返回保证；异步拷贝的“完成”必须由流或事件保证。常见错误是调用 cudaMemcpyAsync 后立刻在主机读取目标内存，此时数据可能尚未到达。正确做法是：若后续操作在同一流中，流内顺序天然保证；若主机需要读取结果，则调用 cudaStreamSynchronize(stream) 或 cudaEventSynchronize(event)。此外，异步拷贝要求主机内存在拷贝期间保持有效且不被换页，这正是页锁定内存的意义。

页锁定内存对带宽的影响

普通 malloc 分配的内存可分页，驱动在 DMA 前需将其锁定并可能复制到内部暂存区，导致带宽下降、异步性丢失。使用 cudaMallocHost 或 cudaHostAlloc 分配的页锁定内存可被 GPU 直接 DMA 访问，通常能显著提升 H2D/D2H 带宽，并使 cudaMemcpyAsync 真正异步。代价是页锁定内存占用物理页、分配较慢、过量使用会拖累主机系统，因此应只对参与传输的缓冲区使用，并在结束时 cudaFreeHost 释放。

向量加法端到端实验

下面给出一个可编译运行的示例。设 N = 1<<24（约 1677 万个 float，单缓冲区约 64 MB），计算 h_C = h_A + h_B。为对比，分别用可分页内存 + cudaMemcpy，以及页锁定内存 + cudaMemcpyAsync 两条路径，并用 CUDA 事件测量 H2D、核函数、D2H 三段耗时。

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

__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

int main() {
    const int N = 1 << 24;
    const size_t bytes = N * sizeof(float);

    // 页锁定主机内存
    float *h_A, *h_B, *h_C;
    cudaMallocHost(&h_A, bytes);
    cudaMallocHost(&h_B, bytes);
    cudaMallocHost(&h_C, bytes);
    for (int i = 0; i < N; ++i) { h_A[i] = 1.0f; h_B[i] = 2.0f; }

    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    cudaStream_t stream;
    cudaStreamCreate(&stream);
    cudaEvent_t e0, e1, e2, e3;
    for (auto e : {&e0, &e1, &e2, &e3}) cudaEventCreate(e);

    cudaEventRecord(e0, stream);
    cudaMemcpyAsync(d_A, h_A, bytes, cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(d_B, h_B, bytes, cudaMemcpyHostToDevice, stream);
    cudaEventRecord(e1, stream);

    int threads = 256, blocks = (N + threads - 1) / threads;
    vecAdd<<<blocks, threads, 0, stream>>>(d_A, d_B, d_C, N);
    cudaEventRecord(e2, stream);

    cudaMemcpyAsync(h_C, d_C, bytes, cudaMemcpyDeviceToHost, stream);
    cudaEventRecord(e3, stream);
    cudaEventSynchronize(e3);

    float tH2D, tK, tD2H;
    cudaEventElapsedTime(&tH2D, e0, e1);
    cudaEventElapsedTime(&tK,   e1, e2);
    cudaEventElapsedTime(&tD2H, e2, e3);
    printf("H2D %.3f ms, Kernel %.3f ms, D2H %.3f ms\n", tH2D, tK, tD2H);
    printf("H2D BW %.1f GB/s, D2H BW %.1f GB/s\n",
           2.0 * bytes / tH2D / 1e6, bytes / tD2H / 1e6);

    // 校验
    bool ok = true;
    for (int i = 0; i < N; ++i) if (h_C[i] != 3.0f) { ok = false; break; }
    printf("check: %s\n", ok ? "PASS" : "FAIL");

    cudaStreamDestroy(stream);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    cudaFreeHost(h_A); cudaFreeHost(h_B); cudaFreeHost(h_C);
    return 0;
}

编译命令：nvcc -O2 -arch=sm_80 vecadd.cu -o vecadd。运行后观察三段耗时：在典型 PCIe 4.0 x16 平台上，H2D 与 D2H 各约 64 MB，若页锁定路径测得 H2D 约 0.6～0.9 ms、D2H 约 0.6～0.9 ms，则对应带宽约 12～20 GB/s；核函数处理 64 MB 输入通常远小于 0.1 ms。若把页锁定换成 malloc，H2D/D2H 耗时往往明显上升，带宽下降，这就是页锁定内存的价值所在。

2D/3D 拷贝的适用场景

当数据是带填充的矩阵或体数据时，cudaMemcpy2D 的 spitch/dpitch 表示每行实际占用的字节数（含填充），width 是有效列字节数，height 是行数。例如源矩阵每行 256 B、有效 200 B、共 64 行，则一次 cudaMemcpy2D(dst, 256, src, 256, 200, 64, cudaMemcpyHostToDevice) 即可完成，避免 64 次小拷贝。三维场景（如体素、多通道图像）用 cudaMemcpy3D，通过 cudaPitchedPtr 与 cudaExtent 描述 pitch、width、height、depth。它们同样有异步版本 cudaMemcpy2DAsync/cudaMemcpy3DAsync，可放入流中与计算重叠。

结果判读与优化方向

实验闭环的判断依据有三点：一是校验 h_C[i] == 3.0f 全部通过，说明搬运与计算正确；二是三段耗时中 H2D/D2H 占比是否远大于核函数，若是，则瓶颈在传输；三是页锁定路径相对可分页路径的带宽提升是否明显。若带宽仍偏低，可检查是否误用可分页内存、是否未使用异步、是否传输了不必要的填充字节。进一步优化可考虑多流分块传输与计算重叠、使用 cudaMemcpy2DAsync 减少小拷贝，以及只回传真正需要的子区域。