前几课已经发现，很多 CUDA 程序并不是慢在 GPU kernel，而是慢在 H2D 和 D2H 数据搬运。
因此，第四课的重点从“怎么写 kernel”转向“怎么让数据更快地进出 GPU”。

普通主机内存属于 Pageable Memory（可分页内存）。使用这种内存进行 cudaMemcpy 时，数据往往需要先经过一个临时的锁页缓冲区，再通过 DMA 传输到 GPU，过程相对多了一次中转。

而 Pinned Memory 是通过 cudaMallocHost 申请的固定页内存，它不会被操作系统换出，GPU 可以更直接地进行 DMA 传输，因此能够显著提升 Host 与 Device 之间的数据传输带宽。

• DMA 全称是 Direct Memory Access，直接内存访问。
• DMA 要求数据源地址在传输过程中稳定不变，否则搬到一半，操作系统把这块内存换走了，DMA 就懵了。

写代码进行测试（代码见文末）

实验结果：

从实测数据看，Pinned Memory 的提升非常明显。

在 64 MB 数据规模下，Pageable Memory 的 H2D 带宽约为 4.745 GB/s，D2H 带宽约为 5.255 GB/s；而 Pinned Memory 的 H2D 带宽提升到 12.307 GB/s，D2H 提升到 13.112 GB/s。这意味着 H2D 大约提升 2.59 倍，D2H 大约提升 2.50 倍。

在 256 MB 数据规模下，结果同样稳定。Pageable Memory 的 H2D 带宽约为 4.711 GB/s，D2H 约为 4.942 GB/s；Pinned Memory 的 H2D 达到 12.356 GB/s，D2H 达到 13.136 GB/s。对应提升约为：H2D 2.62 倍，D2H 2.66 倍。这说明 Pinned Memory 并不是偶然加速，而是在不同数据规模下都能稳定提升传输效率。

工程结论很明确：如果 CUDA 程序中存在大规模数据传输，或者频繁进行 CPU↔GPU 拷贝，应优先考虑使用 Pinned Memory。 它可以显著降低数据传输成本，从而改善端到端性能。

不过，Pinned Memory 也不是越多越好。因为它会锁定物理内存，不能被操作系统轻易换出。如果申请过多，可能影响系统整体内存调度。因此更合理的做法是：在大数据传输、批量推理、流式处理、异步拷贝等关键路径上使用 Pinned Memory，而不是把所有主机内存都改成 Pinned Memory。

将第四课总结为一句话：

Pinned Memory 通过减少 Host 侧中转，让 CPU↔GPU 数据传输从约 5 GB/s 提升到约 13 GB/s，是优化 CUDA 端到端性能的重要一步。

附代码

jupyter cuda环境

# 查看当前 Colab/Notebook 分配到的 GPU 信息
!nvidia-smi

# 查看 nvcc 编译器版本，确认 CUDA 编译工具链可用
!nvcc --version

# 安装 nvcc4jupyter，让 notebook 可以直接运行 CUDA C/C++ 代码单元
!pip install nvcc4jupyter

# 加载 nvcc4jupyter 扩展，之后可以使用 %%cuda 魔法命令
%load_ext nvcc4jupyter

编写cuda代码

%%writefile cuda_pin_memory.cu
#include <cuda_runtime.h>
#include <iostream>
#include <chrono>

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            std::cerr << "CUDA Error: " << cudaGetErrorString(err) << std::endl; \
            exit(1); \
        } \
    } while (0)

float measure_bandwidth(void* h_ptr, void* d_ptr, size_t bytes, cudaMemcpyKind kind) {
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    CUDA_CHECK(cudaEventRecord(start));
    if (kind == cudaMemcpyHostToDevice) {
        CUDA_CHECK(cudaMemcpy(d_ptr, h_ptr, bytes, kind));
    } else if (kind == cudaMemcpyDeviceToHost) {
        CUDA_CHECK(cudaMemcpy(h_ptr, d_ptr, bytes, kind));
    } else {
        std::cerr << "Unsupported cudaMemcpyKind" << std::endl;
        exit(1);
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));

    float ms = 0;
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));

    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));

    return ms;
}

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

    std::cout << "Data size: " << bytes / 1024.0 / 1024.0 << " MB\n";

    // ---------------- Pageable ----------------
    float* h_pageable = (float*)malloc(bytes);
    float* d_mem;
    CUDA_CHECK(cudaMalloc(&d_mem, bytes));

    for (size_t i = 0; i < N; ++i) h_pageable[i] = i;

    float h2d_page_ms = measure_bandwidth(h_pageable, d_mem, bytes, cudaMemcpyHostToDevice);
    float d2h_page_ms = measure_bandwidth(h_pageable, d_mem, bytes, cudaMemcpyDeviceToHost);

    // ---------------- Pinned ----------------
    float* h_pinned;
    CUDA_CHECK(cudaMallocHost(&h_pinned, bytes));

    for (size_t i = 0; i < N; ++i) h_pinned[i] = i;

    float h2d_pin_ms = measure_bandwidth(h_pinned, d_mem, bytes, cudaMemcpyHostToDevice);
    float d2h_pin_ms = measure_bandwidth(h_pinned, d_mem, bytes, cudaMemcpyDeviceToHost);

    auto gbps = [&](float ms) {
        return bytes / (ms / 1000.0) / 1e9;
    };

    std::cout << "\n===== Pageable Memory =====\n";
    std::cout << "H2D: " << h2d_page_ms << " ms, " << gbps(h2d_page_ms) << " GB/s\n";
    std::cout << "D2H: " << d2h_page_ms << " ms, " << gbps(d2h_page_ms) << " GB/s\n";

    std::cout << "\n===== Pinned Memory =====\n";
    std::cout << "H2D: " << h2d_pin_ms << " ms, " << gbps(h2d_pin_ms) << " GB/s\n";
    std::cout << "D2H: " << d2h_pin_ms << " ms, " << gbps(d2h_pin_ms) << " GB/s\n";

    CUDA_CHECK(cudaFree(d_mem));
    free(h_pageable);
    CUDA_CHECK(cudaFreeHost(h_pinned));

    return 0;
}