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

2. 显存分配与释放工程实践

CUDA 提供了多种主机端与设备端内存分配接口，它们的语义差异直接决定了数据能否被核函数访问、是否支持页锁定、是否可被统一虚拟地址管理，以及释放时应该调用哪个 API。工程中最常见的错误不是“不会分配”，而是“分配接口与释放接口不匹配”“忽略对齐要求导致向量化访存失效”“错误码未检查导致泄漏被掩盖”。本节围绕 cudaMalloc、cudaFree、cudaMallocManaged、cudaMallocHost、cudaHostAlloc 五个接口，梳理语义、对齐、错误处理与泄漏排查，并给出一套可复用的显存管理封装。

一、五个接口的语义差异

cudaMalloc(void** devPtr, size_t size) 在设备全局内存中分配线性内存，返回的指针只能在设备代码中解引用，主机端不能直接读写。它是最基础的设备内存分配方式，释放必须使用 cudaFree。

cudaMallocManaged(void** devPtr, size_t size, unsigned int flags) 分配统一内存（Unified Memory），返回的指针在主机和设备上都可访问，由 CUDA 运行时按需在 CPU 与 GPU 之间迁移页面。它简化了编程模型，但首次访问会触发缺页迁移，若访问模式跨页跳跃，性能可能显著低于显式拷贝。释放同样使用 cudaFree。

cudaMallocHost(void** ptr, size_t size) 分配页锁定（pinned）主机内存，主机可直接访问，且 DMA 拷贝带宽通常高于普通可分页内存。它不能被核函数直接解引用，必须通过 cudaMemcpy 传输。释放必须使用 cudaFreeHost，用 free 或 cudaFree 释放都是错误的。

cudaHostAlloc(void** ptr, size_t size, unsigned int flags) 是 cudaMallocHost 的增强版，通过 flags 控制映射、可移植、写合并等属性。常用标志包括 cudaHostAllocMapped（获得设备端可访问的映射指针，配合 cudaHostGetDevicePointer 使用）、cudaHostAllocPortable（多 GPU 上下文可见）、cudaHostAllocWriteCombined（写合并，适合只写不读的传输缓冲）。释放同样使用 cudaFreeHost。

二、对齐要求与向量化访存

cudaMalloc 返回的设备指针至少按 256 字节对齐，这通常足以满足 float4（16 字节）、double2（16 字节）等向量类型的对齐需求。但工程中真正的风险来自“偏移访问”：如果在一个已对齐的基址上做非对齐偏移，再强制转换为向量指针，就会破坏对齐假设。

例如，基址 d_ptr 按 256 字节对齐，若执行 float4* v = reinterpret_cast<float4*>(d_ptr + 1);，此时地址偏移 4 字节，不再是 16 的倍数。核函数中解引用 v[i] 会触发 misaligned address 运行时错误，报错信息通常为 an illegal memory access was encountered 或直接 misaligned address。排查时应检查所有向量化指针的转换点，确认偏移量是向量宽度的整数倍。

页锁定内存同样有对齐要求。cudaHostAlloc 返回的指针也满足系统页对齐，但若使用 cudaHostAllocWriteCombined，该内存不适合主机端读取，只适合作为设备写入的传输缓冲，否则读性能会明显下降。

三、错误处理：每次分配都必须检查

CUDA 运行时 API 返回 cudaError_t，但很多代码习惯性忽略返回值。显存分配失败最常见的原因是显存不足，返回 cudaErrorMemoryAllocation。如果不检查，后续核函数启动会以“非法内存访问”的形式报错，掩盖真正的根因。

推荐做法是封装一个检查宏，在每次分配、拷贝、释放后立即检查并打印文件名、行号与错误字符串：

#define CUDA_CHECK(call)                                                     \
  do {                                                                       \
    cudaError_t err = (call);                                                \
    if (err != cudaSuccess) {                                                \
      fprintf(stderr, "CUDA error %s:%d: %s\n", __FILE__, __LINE__,          \
              cudaGetErrorString(err));                                      \
      exit(EXIT_FAILURE);                                                    \
    }                                                                        \
  } while (0)

// 使用示例
float* d_data = nullptr;
CUDA_CHECK(cudaMalloc(&d_data, N * sizeof(float)));
CUDA_CHECK(cudaMemset(d_data, 0, N * sizeof(float)));
CUDA_CHECK(cudaFree(d_data));

注意 cudaGetErrorString 返回的是最近一次错误的描述，若在检查前又调用了其他 CUDA API，错误状态可能被覆盖。因此检查必须紧跟目标调用，不要延迟到函数末尾统一检查。

四、资源泄漏排查

显存泄漏的典型表现是：程序反复执行分配-释放循环后，cudaMemGetInfo 报告的可用显存持续下降。排查手段有三类。

第一，使用 cudaMemGetInfo(&freeMem, &totalMem) 在关键节点打印可用显存，观察是否单调递减。第二，使用 compute-sanitizer --tool memcheck 运行程序，它可以报告非法访问与部分泄漏信息。第三，使用 cuda-memcheck --leak-check full（旧版本）或 compute-sanitizer --leak-check full，在程序退出时列出未释放的分配记录，包括分配大小与调用栈。

工程上更可靠的做法是从源头避免泄漏：所有分配都通过 RAII 封装管理，析构时自动释放；禁止在循环体内裸调用 cudaMalloc 而不配对释放；对异常路径也要保证释放，例如使用 std::unique_ptr 配合自定义 deleter。

五、可复用的显存管理封装

下面给出一个轻量封装，覆盖设备内存、统一内存与页锁定内存三类资源，统一错误检查与自动释放。它使用模板与自定义 deleter，避免手写 cudaFree 遗漏。

template <typename T>
struct DeviceDeleter {
  void operator()(T* p) const { if (p) CUDA_CHECK(cudaFree(p)); }
};

template <typename T>
struct HostDeleter {
  void operator()(T* p) const { if (p) CUDA_CHECK(cudaFreeHost(p)); }
};

template <typename T>
using DevicePtr = std::unique_ptr<T, DeviceDeleter<T>>;

template <typename T>
using HostPtr = std::unique_ptr<T, HostDeleter<T>>;

template <typename T>
DevicePtr<T> makeDevice(size_t n) {
  T* p = nullptr;
  CUDA_CHECK(cudaMalloc(&p, n * sizeof(T)));
  return DevicePtr<T>(p);
}

template <typename T>
DevicePtr<T> makeManaged(size_t n) {
  T* p = nullptr;
  CUDA_CHECK(cudaMallocManaged(&p, n * sizeof(T)));
  return DevicePtr<T>(p);
}

template <typename T>
HostPtr<T> makePinned(size_t n, unsigned int flags = cudaHostAllocDefault) {
  T* p = nullptr;
  CUDA_CHECK(cudaHostAlloc(&p, n * sizeof(T), flags));
  return HostPtr<T>(p);
}

使用方式如下，资源在作用域结束时自动释放，即使中途抛出异常也不会泄漏：

{
  auto d_in  = makeDevice<float>(N);
  auto d_out = makeDevice<float>(N);
  auto h_buf = makePinned<float>(N, cudaHostAllocWriteCombined);

  CUDA_CHECK(cudaMemcpy(d_in.get(), h_buf.get(), N * sizeof(float),
                        cudaMemcpyHostToDevice));
  kernel<<<blocks, threads>>>(d_in.get(), d_out.get(), N);
  CUDA_CHECK(cudaGetLastError());
  CUDA_CHECK(cudaDeviceSynchronize());
}  // d_in / d_out / h_buf 在此自动释放

封装的关键点有三条：一是分配与释放接口严格配对，设备内存与统一内存走 cudaFree，页锁定内存走 cudaFreeHost；二是所有 CUDA 调用都经过 CUDA_CHECK，错误在发生点暴露；三是用 RAII 保证异常安全，避免在 return 或 throw 路径上漏掉释放。

不要用 free() 释放 cudaMallocHost 或 cudaHostAlloc 返回的指针，也不要用 cudaFree 释放页锁定内存。混用会导致未定义行为，常见表现是程序在退出时崩溃或堆损坏。统一内存虽然主机可访问，但它的释放接口仍是 cudaFree，不是 free。

最后给出一个排查清单，便于在出现分配失败或泄漏时逐项核对：确认每个 cudaMalloc 都有对应的 cudaFree；确认 cudaMallocHost/cudaHostAlloc 都走 cudaFreeHost；确认向量化指针的偏移是向量宽度的整数倍；确认每次分配后立即检查 cudaError_t；确认循环体内没有未释放的临时分配；确认异常路径也有释放逻辑。把这六项做成代码审查项，可以消除绝大多数显存管理类故障。