📚 课程目录第2章显存分配与释放

2. 显存分配与释放

显存分配，说白了就是给GPU腾地方放数据。很多新手一上来就调kernel，结果发现数据根本传不过去——问题就出在分配这步。我个人习惯把显存分配分成四个层次：基础分配、统一内存、页锁定内存、内存池。咱们一层层说。

2.1 cudaMalloc / cudaFree：最基础的显存分配

cudaMalloc 是CUDA里最直接的显存分配函数。它的用法跟 malloc 几乎一样，只不过分配的是GPU显存。

float *d_data;
cudaMalloc((void**)&d_data, N * sizeof(float));
// ... 使用 d_data ...
cudaFree(d_data);

嗯，这里要注意：cudaMalloc 返回的是设备指针，你不能在CPU代码里直接解引用它。我在项目中遇到过有人拿 d_data[0] 直接读数据，结果段错误——这坑我踩过。

常见陷阱：

忘记检查返回值：cudaMalloc 可能返回 cudaErrorMemoryAllocation
分配0字节：行为未定义，不同驱动表现不同
忘记 cudaFree：显存泄漏，跑几次就OOM

cudaFree 释放显存。但有个细节：它是同步操作，会等所有正在使用这块显存的kernel执行完才释放。我曾经在调试时发现 cudaFree 卡住，就是因为前面有个死循环kernel没结束。

2.2 cudaMallocManaged：统一内存，省心但别滥用

cudaMallocManaged 是CUDA 6引入的“统一内存”。它分配的内存，CPU和GPU都能直接访问。说白了就是CUDA帮你自动做数据迁移。

float *managed_data;
cudaMallocManaged(&managed_data, N * sizeof(float));
// CPU可以读写
managed_data[0] = 1.0f;
// GPU也能直接访问
my_kernel<<<grid, block>>>(managed_data);
cudaDeviceSynchronize();
cudaFree(managed_data);

听起来很美好对吧？但我在项目中用过一段时间后，发现它有个致命问题：性能不可控。CUDA驱动会在page fault时自动迁移数据，但迁移时机和粒度你控制不了。如果你的访问模式是“CPU写一次，GPU读一万次”，那统一内存会反复迁移，性能惨不忍睹。

我的建议：

原型验证阶段可以用 cudaMallocManaged 快速迭代
性能敏感场景，还是手动管理显存
如果非要用，配合 cudaMemPrefetchAsync 手动指定数据位置

2.3 cudaHostAlloc：页锁定内存，加速数据传输

默认的 malloc 分配的是可分页内存。GPU访问可分页内存时，需要先拷贝到页锁定内存（pinned memory），这个过程是隐式的，但很慢。

cudaHostAlloc 直接分配页锁定内存，省去了中间拷贝。我做过测试：用 cudaHostAlloc 配合 cudaMemcpy，带宽能提升2-3倍。

float *host_pinned;
cudaHostAlloc(&host_pinned, N * sizeof(float), cudaHostAllocDefault);
// 使用
cudaMemcpy(d_data, host_pinned, N * sizeof(float), cudaMemcpyHostToDevice);
// 释放
cudaFreeHost(host_pinned);

  关键参数：
  
    标志位作用我什么时候用
cudaHostAllocDefault普通页锁定大多数场景
cudaHostAllocWriteCombined写合并内存CPU只写、GPU只读时
cudaHostAllocPortable跨设备可用多GPU场景
cudaHostAllocMapped映射到GPU地址空间零拷贝场景

标志位	作用	我什么时候用
`cudaHostAllocDefault`	普通页锁定	大多数场景
`cudaHostAllocWriteCombined`	写合并内存	CPU只写、GPU只读时
`cudaHostAllocPortable`	跨设备可用	多GPU场景
`cudaHostAllocMapped`	映射到GPU地址空间	零拷贝场景

但页锁定内存有个代价：它会占用系统物理内存，不能被换出。分配太多会导致系统内存紧张。我曾经在一台32GB内存的机器上分配了28GB页锁定内存，结果系统直接卡死——嗯，这个教训很深刻。

2.4 内存池技术：告别频繁分配

频繁调用 cudaMalloc 和 cudaFree 是有开销的。每次分配都要和驱动交互，还要做地址映射。如果你的程序需要反复分配释放同样大小的显存，内存池是更好的选择。

CUDA 11.2 引入了 cudaMemPool 系列API。你可以创建一个内存池，预先分配一大块显存，然后从池子里“借”和“还”。

cudaMemPool_t pool;
cudaMemPoolCreate(&pool, 0);

// 从池中分配
float *d_pool_data;
cudaMallocAsync(&d_pool_data, N * sizeof(float), pool);

// 释放回池中
cudaFreeAsync(d_pool_data, pool);

// 销毁池
cudaMemPoolDestroy(pool);

我做过对比：在一个需要反复分配128MB缓冲区的推理服务里，用内存池后分配延迟从平均200μs降到了5μs以下。为什么？因为 cudaMallocAsync 只是从池里取一块已映射好的内存，不需要和驱动做完整协商。

内存池最佳实践：

为不同大小的分配创建不同的池
设置合理的初始大小，避免频繁扩容
使用 cudaMemPoolSetAttribute 控制池的行为
多流场景下，池可以跨流复用内存

2.5 四种分配方式对比

我把这四种方式的核心差异整理成了一张表，方便你快速选择：

分配方式	CPU可访问	GPU可访问	分配开销	适用场景
`cudaMalloc`	否	是	中	标准GPU数据
`cudaMallocManaged`	是	是	低（首次）	原型、小数据量
`cudaHostAlloc`	是	是（需映射）	高	频繁Host-Device传输
内存池	否	是	极低（复用后）	频繁分配释放

2.6 本章知识体系

下面这张图展示了四种分配方式的关系和选择路径：

你想想看，实际项目中90%的性能问题都出在数据搬运上。选对了分配方式，就等于赢在了起跑线上。我个人习惯是：原型用 cudaMallocManaged，上线前根据profile结果换成手动管理。内存池则是“真香”技术——一旦用上就回不去了。