CUDA内存管理的终极秘密

目录

前言

内存分配和释放

内存传输

带宽与吞吐量

固定内存

零拷贝内存

统一虚拟寻址

统一内存寻址UM（Unified Memory）

总结

前言

前面我们不断的使用的开辟内存的函数，这在c语言当中是非常常见的，对于c这类语言，我们程序员需要手动管理内存，但是对于python这类语言，是自动帮我们管理内存的，但是学会内存管理对开发来说更友好

cuda为什么不像python那样自动管理内存？自动管理器gc无法自动的安排你的数据，没有足够的信息做出最优的决策，CUDA的设计哲学是"把硬件能力完整暴露给程序员，让程序员根据具体算法做出最优决策"。

本章的学习是

• 分配和释放设备内存
• 在主机和设备之间搬运数据

即使以上我们之前粗略的讲过，但是本章将详细的讲解，结合之前的内存模型概述

内存分配和释放

// 分配
cudaError_t cudaMalloc(void **devPtr, size_t size);

// 使用示例
int *d_data = nullptr;
CHECK(cudaMalloc((void**)&d_data, sizeof(int) * n));

// 正确做法：检查返回值
cudaError_t err = cudaMalloc((void**)&d_data, sizeof(int)*n);
if(err != cudaSuccess){
    printf("cudaMalloc失败：%s\n", cudaGetErrorString(err));
    // 处理错误
}

这在之前我们是经常使用的，注意这个是void**，因为这个函数要修改devPtr本身，所以要传入地址，指针的地址那就需要二级指针接收，否则你传递值本身是会拷贝，无法进行修改的

如果出错则返回错误码，这个cudaMalloc是一个同步函数，一定要注意检查返回值，我们之前使用的是一个宏函数

//给分配的显存初始化value
cudaError_t cudaMemset(void *devPtr, int value, size_t count);

这个函数是一个同步函数，cpu端会等待这个函数执行完毕再往下异步执行

cudaError_t cudaFree(void * devPtr)

释放gpu内存，注意要和前面分配的指针一致，和cpu一样，不能重复的释放，完成之后才能cudaFree，设备内存的分配和释放非常影响性能，所以，尽量重复利用！

内存传输

注意主机线程不能直接访问设备内存，设备线程也不能直接访问主机内存（CUDA 提供了一种叫做零拷贝内存的机制。通过 cudaHostAlloc 分配的主机端内存，可以被映射到 GPU 的地址空间。这样，GPU 核函数就可以在核函数内部直接访问这块主机内存。）

也就是不能主机端直接*ptr或ptr[i]这样，你只能通过专门的api搬运数据函数，通过DMA把数据搬运过去

主机端 (CPU 内存)                       设备端 (GPU 显存)
┌──────────────┐                      ┌──────────────┐
│ h_a (0x3000) │ ←── PCIe 总线 ──→    │ d_a (0x7000) │
│              │      DMA 搬运        │              │
│ CPU 能直接   │                      │ CPU 不能直接  │
│ 读写 h_a     │                      │ 读写 d_a     │
└──────────────┘                      └──────────────┘

cudaMemcpy 的作用：让 DMA 沿着 PCIe 总线，
把 h_a 的数据搬到 d_a，或者反过来。
CPU 只负责下命令，不负责亲自搬运。

流程：
cpu下达指令，gpu命令处理器接收，下达给gpu的DMA
DMA接收到源地址和目标地址，然后进行搬运，全程不需要cpu和cuda
完成之后通过中断提醒CPU

搬运内存

cudaError_t cudaMemcpy
(void *dst, const void *src, size_t count, cudaMemcpyKind kind);

• dst：目的端指针（数据要搬到哪里去）。

• src：源端指针（数据从哪里来）。

• count：要搬运的 字节数（不是元素个数！）。

• kind：指定搬运方向，是一个 cudaMemcpyKind 枚举值。

• 返回值：cudaError_t，成功返回 cudaSuccess，否则返回错误码。

枚举值	数据流向
cudaMemcpyHostToDevice	主机内存 → 设备显存
cudaMemcpyDeviceToHost	设备显存 → 主机内存
cudaMemcpyDeviceToDevice	设备显存内的两个区域间
cudaMemcpyHostToHost	主机内存内部（极少使用，可直接用 memcpy）

这也是一个同步函数，注意检查返回值

带宽与吞吐量

带宽（Bandwidth）：理论上单位时间内能传输的最大数据量，是硬件规格书上写死的物理上限，不随实际使用情况变化。（RTX 4060显存带宽：272 GB/s）

吞吐量（Throughput）：实际单位时间内传输的数据量，取决于你的访问模式、合并效率、cache命中率等，永远小于等于带宽。

┌─────────────────────────────────────────────────────────────────┐
│                        显卡 (Graphics Card)                     │
│  ┌───────────────────────────────────────────────────────────┐ │
│  │                  GPU 核心芯片 (片上 SRAM)                  │ │
│  │  ┌─────────────────────────────────────────────────────┐  │ │
│  │  │  SM (流多处理器) × 24                              │  │ │
│  │  │  ┌─────────────────┐  ┌─────────────────┐          │  │ │
│  │  │  │  寄存器文件      │  │  共享内存 / L1   │          │  │ │
│  │  │  │  (256 KB/SM)    │  │  (约 128 KB/SM)  │          │  │ │
│  │  │  │  ● 线程私有      │  │  ● Block内共享   │          │  │ │
│  │  │  │  ● 速度最快      │  │  ● 速度极快      │          │  │ │
│  │  │  └─────────────────┘  └─────────────────┘          │  │ │
│  │  │  ┌─────────────────────────────────────────────┐    │  │ │
│  │  │  │  CUDA 核心 (128/SM)  │  Tensor Core (4/SM) │    │  │ │
│  │  │  │  Warp Scheduler ×4   │  SFU, LD/ST 等      │    │  │ │
│  │  │  └─────────────────────────────────────────────┘    │  │ │
│  │  └─────────────────────────────────────────────────────┘  │ │
│  │                          │                                  │ │
│  │  ┌───────────────────────┴─────────────────────────────┐  │ │
│  │  │              L2 缓存 (24 MB, 所有 SM 共享)          │  │ │
│  │  └───────────────────────┬─────────────────────────────┘  │ │
│  │                          │                                  │ │
│  │  ┌───────────────────────┴─────────────────────────────┐  │ │
│  │  │             显存控制器 (Memory Controller)          │  │ │
│  │  │             与 L2 通过内部高速总线连接              │  │ │
│  │  └───────────────────────┬─────────────────────────────┘  │ │
│  └──────────────────────────┼────────────────────────────────┘ │
│                             │                                   │
│          ┌──────────────────┼──────────────────┐               │
│          │  128 位物理数据线 (位宽 = 128 bit)  │               │
│          └──────────────────┼──────────────────┘               │
│                             │                                   │
│  ┌──────────────────────────┴──────────────────────────────┐  │
│  │                显存颗粒 (片外 DRAM)                      │  │
│  │                总容量：8 GB GDDR6                        │  │
│  │  ┌─────────────────────────────────────────────────────┐ │  │
│  │  │               全局内存 (Global Memory)              │ │  │
│  │  │  - 通过 cudaMalloc 分配                             │ │  │
│  │  │  - 所有线程/SM 可读写                                │ │  │
│  │  │  - 包含本地内存、常量内存数据区、纹理内存数据区     │ │  │
│  │  └─────────────────────────────────────────────────────┘ │  │
│  └──────────────────────────────────────────────────────────┘  │
│                                                                 │
│  ┌──────────────────────────────────────────────────────────┐  │
│  │                   PCIe 接口 (连接 CPU)                   │  │
│  └──────────────────────────────────────────────────────────┘  │
└─────────────────────────────────────────────────────────────────┘

片上就是我们说的GPU，然后片外就是显存颗粒，片上和片外是通过位宽来传递数据的，之间还要通过显存控制器

CPU和GPU之间的传输是通过PCIe总线传输的，我的4060是8通道，带宽是16GB/s左右，这和显存的带宽相比简直差远了，

以上的图示标错了，PCIe是16GB/s左右，每个人的硬件不一样

DDR 内存与早期的 SDRAM 不同，它能在时钟信号的上升沿和下降沿各传输一次数据。这意味着，一个时钟周期可以传 2 次数据。

之前我们计算的频率全都是总的了，不用再x2了，已经把x2算进去了

这里的意思就是大致了解一下显存上面的是使用了DDR技术

结论

• 实际的瓶颈应该在于最慢的PCIe，才16GB/s，那么我们的解决措施就是减少数据搬运次数，不要来回搬运，能在GPU上计算的就一直在GPU上计算
• 在GPU内部，显存带宽是最大的瓶颈，那么我们GPU内部优化的时候应该要提高显存<----->L2缓存的利用率
• 应该把频繁的访问的数据放到共享内存当中，用一次全局显存读取换N次共享内存读取
• 本地内存是在片外，所以我们要防止寄存器溢出，否则你还要去访问全局内存

固定内存（pinned）

在讲Linux系统编程的时候，我们已经详细的了解了内存的分页管理，系统是把物理内存和虚拟地址空间都切分成一页一页的管理，也就是通俗的4KB，这样分的好处在于让程序员眼中是连续的地址，而不是直接操控物理内存，cpu通过页表进行映射，内部有个MMU帮忙翻译地址

CPU和GPU的数据传输是通过PCIe总线传输的，上面有个DMA控制器，也就是cpu发出指令，它不需要cpu本身自己去拿数据传输数据，只要发给DMA即可，DMA就会自动的帮我们传输数据，好了之后在通知CPU即可，这样就能异步的发挥性能，而不需要等待

但是DMA需要真实的物理地址，即使你在cpu端编写代码的时候是连续的，但这个是虚拟地址，它需要映射到真实的物理内存地址，DMA需要的是真实的内存地址，如果这个内存地址是零散分布在四处，而不是连续的时候，数据的搬运就很耗时，并且如果正在传输时os进行换出，把物理页暂时的换到磁盘上，那数据的传输就会出现问题

当数据不连续的时候，必须CPU介入，先把数据拼凑成连续的，DMA才能工作

• 不连续问题：DMA 需要真实的物理地址，而可分页内存的物理地址是分散的，导致 DMA 无法一次高效完成传输。

• “页”被换出问题：操作系统可能把一块数据所在的页临时换到硬盘上，如果此时 DMA 正在传输这块数据，就会出错或丢失数据。

基于以上的问题，提出来固定内存的解决方案

如果没有一开始分配固定内存就是左边，右边是一开始就分配固定内存

很明显如果少了拼凑那效率绝对又会提升一个档次

• 普通内存 + 拼凑：你分配的是普通可分页内存，物理上不连续。每次你要把它搬到 GPU 时，CUDA 驱动必须临时在背后做三件事：① 分配一块连续的固定内存作为临时缓冲区；② CPU 把你的数据从分散的普通内存拷贝到这个临时缓冲区；③ DMA 才从这个临时缓冲区搬到 GPU。这三步是每次传输时实时发生的，CPU 拷贝耗时巨大。

• 固定内存：一开始就用 cudaHostAlloc 分配了一块物理连续、不会被换出的内存。数据直接放进去。当需要搬到 GPU 时，DMA 可以直接从这块内存搬运，零 CPU 拼凑开销。这三步在分配时就已经完成了一次性投资。

固定内存作为一种稀缺资源，过度使用可能导致系统内存不足，所以通常只对需要频繁传输的关键数据使用。

API函数

cudaError_t cudaHostAlloc(void** pHost, size_t size, unsigned int flags);

• pHost：指向主机指针的指针，函数会把分配好的地址写进去。

• size：要分配的字节数。

• flags：控制内存行为的标志，常用选项包括：

  • cudaHostAllocDefault：最基础的固定内存，行为和上面理解的一样。

  • cudaHostAllocMapped：零拷贝内存标志。这会让分配的这块主机内存能被映射到 GPU 的地址空间，让你在核函数里直接访问它。

  • cudaHostAllocPortable：让这块内存在所有 CUDA 设备上都可用（多卡环境）。

  • cudaHostAllocWriteCombined：优化写操作，提升主机到设备的传输速度，但会降低 CPU 本身的读取速度。

cudaError_t cudaFreeHost(void* ptr);

#include <cuda_runtime.h>
#include <stdio.h>
#include <cstring>
#include "../freshman.hpp"   

using namespace std;
void sumArrays(float * a,float * b,float * res,const int size)
{
  for(int i=0;i<size;i+=4)
  {
    res[i]=a[i]+b[i];
    res[i+1]=a[i+1]+b[i+1];
    res[i+2]=a[i+2]+b[i+2];
    res[i+3]=a[i+3]+b[i+3];
  }
}
__global__ void sumArraysGPU(float*a,float*b,float*res)
{
  int i=blockIdx.x*blockDim.x+threadIdx.x;
  res[i]=a[i]+b[i];
}
int main() {
    int dev = 0;
    cudaSetDevice(dev);

    // 数据量：64 MB
    int nElem = 1 << 24;          // 16M 个 float
    int nByte = nElem * sizeof(float);
    printf("数据量: %d 个 float, %.2f MB\n", nElem, nByte / (1024.0 * 1024.0));

    // ========== 分配设备内存 ==========
    float *d_a, *d_b, *d_res;
    CHECK(cudaMalloc(&d_a, nByte));
    CHECK(cudaMalloc(&d_b, nByte));
    CHECK(cudaMalloc(&d_res, nByte));

    // ========== 分配主机端普通可分页内存 ==========
    float *h_a_pageable = (float*)malloc(nByte);
    float *h_b_pageable = (float*)malloc(nByte);
    float *h_res_pageable = (float*)malloc(nByte);

    // ========== 分配主机端固定内存 ==========
    float *h_a_pinned, *h_b_pinned;
    CHECK(cudaHostAlloc(&h_a_pinned, nByte, cudaHostAllocDefault));
    CHECK(cudaHostAlloc(&h_b_pinned, nByte, cudaHostAllocDefault));

    // ========== 初始化数据 ==========
    initialData(h_a_pageable, nElem);
    initialData(h_b_pageable, nElem);
    // 将同样的数据复制到固定内存，保证公平对比
    memcpy(h_a_pinned, h_a_pageable, nByte);
    memcpy(h_b_pinned, h_b_pageable, nByte);

    // ========== 预热：消除首次调用 API 的初始化开销 ==========
    // 做一次 H→D 传输（普通内存 + 固定内存）和一次核函数调用
    sumArraysGPU<<<dim3((nElem + 1023) / 1024), dim3(1024)>>>(d_a, d_b, d_res);
    cudaMemcpy(d_a, h_a_pageable, nByte, cudaMemcpyHostToDevice);
    cudaMemcpy(d_a, h_a_pinned, nByte, cudaMemcpyHostToDevice);
    cudaDeviceSynchronize();

    // 创建计时事件
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    float time_pageable, time_pinned;

    // ========== 测试普通可分页内存 H→D 传输 ==========
    cudaEventRecord(start);
    cudaMemcpy(d_a, h_a_pageable, nByte, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b_pageable, nByte, cudaMemcpyHostToDevice);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time_pageable, start, stop);

    // ========== 测试固定内存 H→D 传输 ==========
    cudaEventRecord(start);
    cudaMemcpy(d_a, h_a_pinned, nByte, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b_pinned, nByte, cudaMemcpyHostToDevice);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time_pinned, start, stop);

    // ========== 输出结果 ==========
    printf("普通内存 H→D 传输耗时: %.4f ms\n", time_pageable);
    printf("固定内存 H→D 传输耗时: %.4f ms\n", time_pinned);
    printf("加速比: %.2fx\n", time_pageable / time_pinned);

    // ========== 释放资源 ==========
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_res);
    free(h_a_pageable); free(h_b_pageable); free(h_res_pageable);
    cudaFreeHost(h_a_pinned); cudaFreeHost(h_b_pinned);
    cudaEventDestroy(start); cudaEventDestroy(stop);

    return 0;
}

用nsys测，不要用ncu测，nsys是测宏观的，我们需要看host和device的传输耗时，而ncu测的是细微的某个核函数的内部，而不是宏观总体的

 nsys profile --stats=true ./main

占大头的是开辟内存，并且是cudaHostAlloc，说明固定内存的开辟消耗非常大

总的来说固定内存的耗时是相较于普通内存小的，但是开辟的时候所需要的时间多余普通内存，核心就在于还需要向os请求锁定物理页面，这个本来就比普通的malloc耗时，但是总的效益还是有的，说明在传输上面节省了很多时间

这里计算得到

• 8.20GB/s
• 10.53GB/s

所以固定内存适合需要频繁、大量数据传输，且对带宽利用率要求极高的场景。它通过牺牲操作系统对内存的灵活管理，换取 DMA 的极致传输效率。

1. 适合固定内存的典型场景

场景	为什么适合	典型应用
大数据块反复传输	数据量越大，节省的 CPU 拼凑开销越显著。固定内存让 DMA 直接连续搬运，避免每次传输都临时分配缓冲区。	高清视频流处理、大型矩阵运算、批量图像预处理。
频繁的 H↔D 数据交换	普通内存每次 cudaMemcpy 都要驱动层介入拼凑。固定内存一次锁定，多次传输，摊薄了锁定开销。	迭代算法中每次循环都要把中间结果传回 CPU 检查或记录。
异步传输与流重叠	固定内存是 cudaMemcpyAsync 的硬性要求。普通可分页内存无法与核函数计算并行执行。	高性能计算中让数据拷贝和核函数执行重叠，隐藏传输延迟。（后面详讲）
实时或低延迟系统	普通内存的 CPU 拼凑耗时不可控，可能引起抖动。固定内存的传输时间更稳定。	自动驾驶中摄像头帧的实时上传、医疗影像实时重建。

2. 不适合固定内存的情况

情况	原因
小数据量（< 1MB）	固定内存的 cudaHostAlloc 分配成本可能比传输本身还高。我第一次实验用 64KB 时加速比只有 0.76x ，有时候跑出比1小有时候比1大，小数据不占优势。
仅使用一次的数据	分配固定内存的额外开销无法被多次传输摊薄。直接用普通内存 + 一次 cudaMemcpy 更省事。
内存资源紧张	固定内存会减少操作系统可管理的物理内存。大量分配会导致系统整体性能下降，甚至触发内存不足。
数据传输本身不是瓶颈	如果程序主要时间都花在核函数计算上，传输占比很小，优化传输带来的整体收益有限。应该先用 nsys 确认传输确实是瓶颈。

零拷贝内存

之前提到过，主机不能访问设备内存，设备不能访问主机内存，即不可能通过解引用去访问

但是后来出现了一个例外-----零拷贝内存（一种特殊的固定内存）
 

cudaError_t cudaHostAlloc(void** pHost, size_t size, unsigned int flags);

• pHost：指向主机指针的指针，函数会把分配好的地址写进去。

• size：要分配的字节数。

• flags：控制内存行为的标志，常用选项包括：

  • cudaHostAllocDefault：最基础的固定内存，行为和上面理解的一样。

  • cudaHostAllocMapped：零拷贝内存标志。这会让分配的这块主机内存能被映射到 GPU 的地址空间，让你在核函数里直接访问它。

  • cudaHostAllocPortable：让这块内存在所有 CUDA 设备上都可用（多卡环境）。

  • cudaHostAllocWriteCombined：优化写操作，提升主机到设备的传输速度，但会降低 CPU 本身的读取速度。

我们可以通过选项去开辟一块零拷贝内存，注意这块内存是主机端的内存，包括上面说的固定内存也是主机端的内存

1. 分配：用 cudaHostAlloc 并指定 cudaHostAllocMapped 标志，在主机端分配一块固定内存。

2. 映射：驱动会为这块内存创建一个到 GPU 地址空间的映射。

3. 获取设备指针：在主机端调用 cudaHostGetDevicePointer，得到这块内存在 GPU 地址空间的合法指针（例如 d_mapped）。

4. 核函数直接访问：把这个 d_mapped 传给核函数，GPU 就能像访问普通显存一样读写它了。

cudaError_t cudaHostGetDevicePointer(void ** pDevice,void * pHost,unsigned flags);

• pDevice：一个输出参数，用来存放映射后的 GPU 设备指针。调用成功后，你就能用这个指针在核函数里直接读写主机内存。这是一个 void**，意味着函数需要间接写入你的指针变量。

• pHost：输入参数，指向一块已经分配好的零拷贝固定内存的主机指针。这块内存必须是通过 cudaHostAlloc 搭配 cudaHostAllocMapped 标志分配的，否则调用会失败。

• flags：保留参数，目前必须设为 0。在多卡系统中，可以通过指定非零的 flags 来获取针对特定 GPU 的设备指针

• 返回值：cudaSuccess 表示成功，否则返回错误码。

float *h_data;
cudaHostAlloc(&h_data, size, cudaHostAllocMapped); // 1. 分配映射的固定内存

float *d_mapped;
cudaHostGetDevicePointer(&d_mapped, h_data, 0);     // 2. 获取映射后的设备指针

kernel<<<grid, block>>>(d_mapped); // 3. 在核函数里直接使用

// 核函数内部可以像正常设备指针一样读写 d_mapped
__global__ void kernel(float *ptr) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    ptr[tid] = ...;  // 每次访问都实时经过 PCIe 总线
}

实际在传输数据的过程当中，你没有显式的使用cudaMemcpy，但是你把地址映射过去了，本质是由核函数访存指令隐式的触发数据传输，所以每次使用都会触发PCIe总线传输

cudaMemcpy是一次传输到显存，开销大，你以后都能直接用，后续不需要传输了，但是零拷贝内存是你每次使用都会传输，那么缺点就很明显，PCIe的带宽很小，会造成速度慢延迟高等缺点，而且GPU多线程环境无法合并成一次大事务，只能产生多次小数据传输，对于零拷贝内存，这些地址都在主机内存里。LSU的合并逻辑无法跨过PCIe总线去把分散在主机内存里的零散地址打包。每次访问，GPU只能向主机内存发起一个独立的、小粒度的读取请求，等数据传回来后，再处理下一个。

有时候假设你有64MB的数据，核函数只是会访问某些零碎的数据，但是你不知道访问哪些，你无法用cudaMemcpy去一次传输，如果你知道核函数要用哪些，当然一次传输的效率肯定更高，但是针对不知道，这个决策必须在 GPU 内部，根据运行时计算结果才能做出。你可以使用零拷贝，就不用传输64MB的数据，让核函数访存去隐式的触发即可，所以这个是能够提高PCIe的吞吐率，能够把传输用在有效数据上，而不是大量的无效数据

使用零拷贝的场景

• 当设备内存不够时，使用主机内存
• 无法预测核函数需要访问哪些数据
• 小数据多次访问（不要变成多次的cudaMemcpy调用）

实验
我们在固定内存的核函数上面做实验，我们需要比较一下在针对特定场景，即无法预测核函数需要访问哪些数据时，零拷贝内存是否占优势，分为密集型（全部访问）和稀疏型（只访问特定的数据）

#include <cuda_runtime.h>
#include <stdio.h>
#include <cstdlib>
#include <ctime>

#define N (1 << 24)  // 16M 个 float = 64 MB
#define ACCESS_RATIO 0.01f  // 稀疏访问比例 1%

// 核函数：密集访问（遍历全部数据）
__global__ void denseAccess(float *data, float *output, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        output[tid] = data[tid] * 2.0f;
    }
}

// 核函数：稀疏访问（只访问随机分布的 1% 数据）
__global__ void sparseAccess(float *data, float *output, int n, int *indices) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < int(n * ACCESS_RATIO)) {
        int idx = indices[i];
        output[idx] = data[idx] * 2.0f;
    }
}

int main() {
    int nBytes = N * sizeof(float);
    int sparseCount = int(N * ACCESS_RATIO);

    // ========== 分配设备内存 ==========
    float *d_data, *d_output;
    cudaMalloc(&d_data, nBytes);
    cudaMalloc(&d_output, nBytes);

    // ========== 分配主机端普通内存 ==========
    float *h_data_pageable = (float*)malloc(nBytes);
    float *h_output_pageable = (float*)malloc(nBytes);
    for (int i = 0; i < N; i++) h_data_pageable[i] = float(i);

    // ========== 分配零拷贝内存 ==========
    float *h_data_mapped;
    cudaHostAlloc(&h_data_mapped, nBytes, cudaHostAllocMapped);
    memcpy(h_data_mapped, h_data_pageable, nBytes);
    float *d_data_mapped;
    cudaHostGetDevicePointer(&d_data_mapped, h_data_mapped, 0);

    // ========== 生成随机访问索引（稀疏访问用） ==========
    int *h_indices = (int*)malloc(sparseCount * sizeof(int));
    int *d_indices;
    cudaMalloc(&d_indices, sparseCount * sizeof(int));
    srand(time(0));
    for (int i = 0; i < sparseCount; i++) h_indices[i] = rand() % N;
    cudaMemcpy(d_indices, h_indices, sparseCount * sizeof(int), cudaMemcpyHostToDevice);

    // ========== 预热 ==========
    denseAccess<<<(N+255)/256, 256>>>(d_data, d_output, N);
    sparseAccess<<<(sparseCount+255)/256, 256>>>(d_data, d_output, N, d_indices);
    cudaDeviceSynchronize();

    // ========== 计时工具 ==========
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    float timeA, timeB, timeC, timeD;

    dim3 denseGrid((N+255)/256), denseBlock(256);
    dim3 sparseGrid((sparseCount+255)/256), sparseBlock(256);

    // --- 实验 A：显式拷贝 + 密集访问 ---
    cudaMemcpy(d_data, h_data_pageable, nBytes, cudaMemcpyHostToDevice);
    cudaEventRecord(start);
    denseAccess<<<denseGrid, denseBlock>>>(d_data, d_output, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&timeA, start, stop);

    // --- 实验 B：零拷贝 + 密集访问 ---
    cudaEventRecord(start);
    denseAccess<<<denseGrid, denseBlock>>>(d_data_mapped, d_output, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&timeB, start, stop);

    // --- 实验 C：显式拷贝 + 稀疏访问 ---
    cudaMemcpy(d_data, h_data_pageable, nBytes, cudaMemcpyHostToDevice);
    cudaEventRecord(start);
    sparseAccess<<<sparseGrid, sparseBlock>>>(d_data, d_output, N, d_indices);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&timeC, start, stop);

    // --- 实验 D：零拷贝 + 稀疏访问 ---
    cudaEventRecord(start);
    sparseAccess<<<sparseGrid, sparseBlock>>>(d_data_mapped, d_output, N, d_indices);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&timeD, start, stop);

    // ========== 输出结果 ==========
    printf("=== 实验结果 (数据量: 64 MB) ===\n");
    printf("A. 显式拷贝 + 密集访问: %.4f ms\n", timeA);
    printf("B. 零拷贝   + 密集访问: %.4f ms\n", timeB);
    printf("C. 显式拷贝 + 稀疏访问 (1%%): %.4f ms\n", timeC);
    printf("D. 零拷贝   + 稀疏访问 (1%%): %.4f ms\n", timeD);
    printf("密集访问加速比 (A/B): %.2fx\n", timeA / timeB);
    printf("稀疏访问加速比 (C/D): %.2fx\n", timeC / timeD);

    // ========== 释放资源 ==========
    cudaFree(d_data); cudaFree(d_output); cudaFree(d_indices);
    free(h_data_pageable); free(h_output_pageable); free(h_indices);
    cudaFreeHost(h_data_mapped);
    cudaEventDestroy(start); cudaEventDestroy(stop);
    return 0;
}

这个的访问密度是百分之0.01

这个的访问密度是百分之1

基于以上，针对特定的场景，稀疏访问的时候零拷贝更加，但是注意如果你是多次的访问数据，这些数据每次都要经过PCIe总线，带宽太慢了相比于显存，所以如果你多次的访问，肯定会存在一个临界值，一次加载到显存优于零拷贝内存，所以我们要针对不同的场景去使用

统一虚拟寻址

UVA（Unified Virtual Addressing，统一虚拟寻址）

让 CPU 和 GPU 共享同一个虚拟地址空间，从而通过指针的数值本身，就能自动判断这块内存属于谁。

注意，以上所讲的都是虚拟地址，不是真实的物理地址

float *d_a;
cudaMalloc(&d_a, 128);

1. 驱动接收请求：cudaMalloc 是主机端运行时库函数，它把请求转发给 NVIDIA 内核驱动。

2. 驱动在虚拟地址空间中划拨地址：驱动检查当前 GPU（比如 GPU 0）的虚拟地址空间预留区域，找到一段未使用的、足够大的虚拟地址范围（比如从 0x7F_0000_0000 到 0x7F_0000_007F），这就是即将分配给 d_a 的虚拟地址。

3. 驱动向 GPU 显存管理器请求物理显存：驱动通过 PCIe 总线，向 GPU 内部的显存管理器发送命令，要求分配 128 字节的物理显存。GPU 的显存管理器从空闲显存池中划出 128 字节，得到物理显存地址（比如 0x1000）。

4. 驱动建立映射：驱动在自己的页表管理结构中，将步骤 2 的虚拟地址 0x7F_0000_0000 到 0x7F_0000_007F，映射到步骤 3 的物理显存地址 0x1000。

5. 驱动回填虚拟地址：驱动通过 &d_a 这个地址，把虚拟地址 0x7F_0000_0000（而不是物理地址 0x1000）写入你的 d_a 变量。你的 d_a 现在就是 0x7F_0000_0000。

6. 你使用这个指针时：当你用 d_a[0] 在主机端做 cudaMemcpy 时，驱动会根据 0x7F_0000_0000 这个虚拟地址，查自己的页表，找到物理地址 0x1000，然后发起正确的 PCIe 传输。当核函数内部使用这个虚拟地址时，GPU 内部的 MMU 也会查页表，找到物理地址 0x1000，然后访问显存。

在统一虚拟寻址出现之前：

• 设备指针和主机指针是不能混用的，主机的内存有自己的32/64位虚拟地址，每一块GPU都有自己的显存地址
• 核函数不能使用主机指针，必须手动的标记哪个是主机指针_h，哪个是设备指针_d
• 多GPU的场景下，必须通过CPU进行中转，GPU1不能识别到GPU2的显存地址，必须要经过主机内存中转，浪费PCIe带宽，增加延迟
• 零拷贝内存的使用非常麻烦，一块物理空间要两个地址进行标识

支持出现的前提：

• 64位虚拟地址的出现，32位仅仅支持4GB，但是64位达到了256TB，完全能够把多余的空间映射显存
• 第二代Fermi架构的硬件支持，GPU的MMU能够自动识别某个地址是主机还是哪块GPU，DMA 引擎可以直接使用统一地址发起传输，不需要驱动手动做地址转换

有了UVA出现之后，对于零拷贝内存，压根就不需要两个地址了

设想一下，如果没有UVA，虽然你是零拷贝内存，物理页被锁定了，但是host和device内部都有各自的页表，都有各自的虚拟地址映射物理地址，没有UVA时，你访问零拷贝地址，device需要从设备的虚拟地址映射到物理地址，发现物理地址属于主机零拷贝内存时，通过DMA传输到设备，此时就能用了

因为UVA的出现，每个设备的虚拟地址是不同的，如果没有UVA大家的虚拟地址可能相同，意味着设备的虚拟和主机的虚拟可能相同，那就不能混用，因为你不知道访问的是设备还是主机，一旦统一了，那就可以共用了，这个逻辑大家可以想一想

• 零拷贝内存是页锁的，这个是设备能够访问主机内存的前提，DMA可以直接搬运
• UVA是解决统一编址的问题，大家共用一套虚拟地址
• 实际的物理地址本身就是不同的，以上所讲的都是虚拟地址，大家注意辨别

统一内存寻址UM（Unified Memory）

统一虚拟寻址是更底层的基础，而统一内存是建立在此之上的高级特性。简单来说，统一虚拟寻址只是把 CPU 和 GPU 的虚拟地址空间合并了，但数据的物理位置依然泾渭分明。统一内存则更进一步，自动在 CPU 和 GPU 之间动态迁移数据，让你无需手动拷贝。UVA 只解决了「地址统一、指针通用」的问题，不会主动移动数据；UM 在此基础上，增加了缺页驱动的自动数据迁移能力。

• 无 UVA 零拷贝：数据永远在主机内存，两个指针，手动转换，GPU 远程 PCIe 访问
• 有 UVA 零拷贝：数据永远在主机内存，一个指针，直接用，GPU 远程 PCIe 访问
• 统一内存 UM：一个指针，数据自动搬到访问方的本地内存，GPU 访问时数据在显存，CPU 访问时数据在内存

UM 的核心机制是页故障（Page Fault）+ 按需迁移，全程由驱动和硬件自动完成，开发者完全感知不到：

1. 分配阶段

调用 cudaMallocManaged 分配一块 UM 内存，返回一个全局唯一的虚拟地址（这一步必须依赖 UVA 保证地址通用）。 此时这块内存的物理页落在主机内存里，GPU 页表中对应的条目被标记为「无效」，GPU 还不能直接访问。

2. GPU 访问触发迁移

核函数启动后，SM 访问这个 UM 地址时，GPU MMU 查页表发现该页无效，触发 GPU 缺页异常：

• 触发缺页的 Warp 暂时挂起等待

• 驱动接管，把对应的主机物理页临时锁定，通过 DMA 把数据整块拷贝到 GPU 显存

• 更新 GPU 页表，把虚拟地址映射到本地显存物理地址，标记为有效

• 恢复 Warp 执行，后续访问就和普通显存一样快

3. CPU 再次访问触发回迁

如果之后 CPU 再读写这块内存，CPU 侧会触发缺页：

• 驱动把数据从 GPU 显存拷回主机内存

• 更新页表，让 CPU 可以正常访问

• GPU 侧对应的显存页重新标记为无效，下次 GPU 再访问会再次触发迁移

你可能会想，这个效率也太低了吧，数据搬来搬去，为什么不直接存两份呢？？？

那你的思维又回到之前的cudaMemcpy

设想一下，有一份数据10GB，但是你的显存只有8GB，你如何cudaMalloc10GB出来呢？

只能在host开辟10GB，然后使用UM，你使用到的数据，硬件自动给你转换，这一步对你来说是透明的，只不过效率很慢（但是你第一次加载进去之后，后面的访问就可以享受显存的带宽了，不要频繁的在主机和设备之间来回访问，这样会出现乒乓球效应）

当数据是稀疏的时候也可以用，这个和零拷贝不同的是，零拷贝的数据是在主机的，而这个是会真实的搬运到设备，你完全可以把热点数据加载到设备当中，后续可以想用设备的高速带宽，只用一次PCIe即可

它同时解决了两个关键痛点：超大规模数据和稀疏/随机访问。

API

cudaError_t cudaMallocManaged
(void **devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal);


float *a;
cudaMallocManaged(&a, 1024 * sizeof(float));  // 分配统一内存

// CPU 可以直接读写
for (int i = 0; i < 1024; i++) a[i] = i;

// GPU 也可以直接读写，无需 cudaMemcpy
kernel<<<grid, block>>>(a);
cudaDeviceSynchronize();  // 等 GPU 完成，CPU 再读

// CPU 再次访问时，驱动自动把数据从显存迁回
printf("%f\n", a[0]);  

• devPtr：输出参数，指向分配好的统一内存指针。这个指针在 CPU 和 GPU 端都可以直接使用。

• size：要分配的字节数。

• flags：可选标志，默认 cudaMemAttachGlobal（全局可访问）。

为避免首次访问时的缺页中断延迟，可以用 cudaMemPrefetchAsync 提前把数据搬到目标设备：

// 在 CPU 初始化后，提前把数据迁移到 GPU，避免核函数首次访问时中断
cudaMemPrefetchAsync(a, size, 0);  // 0 表示当前设备

可以用 cudaMemAdvise 给驱动提供访问模式的建议：

// 告诉驱动：这块内存主要在 GPU 端读，偶尔在 CPU 端写
cudaMemAdvise(a, size, cudaMemAdviseSetPreferredLocation, cudaCpuDeviceId);
cudaMemAdvise(a, size, cudaMemAdviseSetAccessedBy, 0);  // 0 表示当前 GPU

释放

cudaFree(a);

可以看出UM是在UVA的基础上的高级用法，只有UVA的支持，UM才得以实现，因为你是同一个地址

总结

固定内存：消除 cudaMemcpy 时 CPU 拼凑数据的开销，让 DMA 一次性连续搬运，加速显式传输。

零拷贝内存：让 GPU 直接访问主机内存，省掉显式 cudaMemcpy。适用于无法预知访问模式或显存不足的场景，但每次访问都走 PCIe，速度较慢。

UVA (统一虚拟寻址)：让 CPU 和 GPU 共享同一个 64 位虚拟地址空间，指针值全局唯一，驱动可根据地址本身判断数据归属。这是所有现代 CUDA 特性的地基。

UM (统一内存)：建立在 UVA 之上，让数据在 CPU 和 GPU 之间按需自动迁移，只加载实际被访问的页面，对用户透明。适用于超大规模数据或访问模式不可预测的场景，但需警惕“乒乓效应”。

核心原则：根据数据量、访问模式、显存容量和性能要求，选择最合适的内存管理方式。