2. 显存分配机制剖析

在GPU程序里，“申请一块显存”看似只是一行API调用，实际却牵涉驱动、运行时、操作系统页表与设备页表的多层协作。理解这些层次，才能解释为什么同样大小的分配，第一次调用可能耗时数百微秒，而后续调用可能只需几微秒；也才能理解为什么高频分配/释放会让训练循环明显变慢。本节从API语义出发，拆解分配开销来源，再引出内存池化思想，并给出一个可运行的最小分配器实现与对比实验方案。

一、三类核心分配API的语义差异

cudaMalloc(void** devPtr, size_t size) 是最基础的设备端分配：它在当前设备上申请一段线性地址空间，返回的指针只能在设备代码或 cudaMemcpy 系列函数中使用，主机端不能直接解引用。其分配粒度受驱动影响，通常按页对齐（常见为2MB大页或64KB页），因此申请1字节也会占用一个页粒度。释放必须用 cudaFree，且指针必须与分配时完全一致。

cudaMallocManaged(void** devPtr, size_t size, unsigned int flags) 分配的是统一内存（Unified Memory），同一指针在主机与设备端均可访问，由驱动在缺页时按需迁移数据。它极大简化了编程模型，但首次访问会触发页错误与迁移，延迟不可预测；在无 cudaMemAdvise 预取的情况下，频繁跨端访问可能比显式拷贝更慢。flags 常用 cudaMemAttachGlobal（默认，所有流可见）或 cudaMemAttachHost（需显式attach到流）。

此外还有 cudaMallocPitch（按行对齐分配二维数组，适合纹理/矩阵访问）、cudaMallocAsync/cudaFreeAsync（流序分配，CUDA 11.2+，天然支持异步与池化）以及 cudaHostAlloc（锁页主机内存，用于零拷贝或异步拷贝）。选择哪一类，取决于数据生命周期、访问模式与是否需要跨端共享。

二、分配开销到底花在哪里

一次 cudaMalloc 的耗时并非恒定，主要来自四部分：

• 驱动与运行时状态切换：用户态运行时需进入内核态与驱动通信，涉及上下文切换与锁竞争，这是固定成本，通常在几微秒量级。
• 虚拟地址空间管理：驱动需在设备的虚拟地址空间中寻找空闲区间、建立映射，若地址空间碎片化，搜索成本上升。
• 物理页分配与页表更新：真正把物理显存页挂到页表上，大块分配时需批量更新页表，成本随size增长。
• 隐式同步：传统 cudaMalloc 是同步操作，会阻塞主机线程，且在某些驱动版本下会与正在执行的流产生隐式同步，打断流水线。

一个常被忽视的事实是：释放（cudaFree）同样昂贵，它需要同步设备、回收页表、归还物理页。因此在“分配—使用—释放”高频循环的代码中，开销往往不是单次分配，而是分配与释放的叠加，以及由此引入的同步等待。

三、内存池化思想

池化的核心是“预分配大块、内部切分、复用而不归还”：启动时一次性向驱动申请较大显存（如256MB），之后所有小分配都从池内切出，释放时只标记为空闲而不调用 cudaFree。这样把昂贵的驱动交互从“每次分配”降为“池扩容时一次”，把同步阻塞变为纯用户态指针运算。

池化需要解决三个问题：切分策略（首次适应、最佳适应或按2的幂分级）、碎片控制（相邻空闲块合并）、对齐与线程安全（多线程下加锁或用线程本地池）。CUDA自带的 cudaMallocAsync 配合 cudaMemPool 就是官方池化实现，可通过 cudaDeviceSetLimit(cudaLimitMemPoolDeallocDelay, ...) 等参数调节回收行为。

四、一个最小可用分配器的实现

下面给出一个简化版池化分配器，采用“首次适应 + 空闲链表 + 相邻合并”策略，块头记录大小与空闲标志，所有块按地址顺序串成链表。为便于教学，省略了多线程锁与设备端校验，实际工程中应加互斥并处理对齐。

struct Block {
    size_t size;        // 块大小（含头部）
    bool   free;        // 是否空闲
    Block* next;        // 地址序下一块
};

class SimplePool {
    char*  base_ = nullptr;   // 池起始地址
    size_t cap_  = 0;         // 池容量
    Block* head_ = nullptr;   // 首块
public:
    bool init(size_t bytes) {
        if (cudaMalloc(&base_, bytes) != cudaSuccess) return false;
        cap_  = bytes;
        head_ = reinterpret_cast<Block*>(base_);
        head_->size = bytes;
        head_->free = true;
        head_->next = nullptr;
        return true;
    }
    void* alloc(size_t n) {
        size_t need = (n + sizeof(Block) + 15) & ~size_t(15); // 16B对齐
        for (Block* b = head_; b; b = b->next) {
            if (!b->free || b->size < need) continue;
            if (b->size - need >= sizeof(Block) + 16) {       // 可切分
                Block* rest = reinterpret_cast<Block*>(
                    reinterpret_cast<char*>(b) + need);
                rest->size = b->size - need;
                rest->free = true;
                rest->next = b->next;
                b->size = need;
                b->next = rest;
            }
            b->free = false;
            return reinterpret_cast<char*>(b) + sizeof(Block);
        }
        return nullptr;   // 池耗尽，实际应触发扩容
    }
    void dealloc(void* p) {
        if (!p) return;
        Block* b = reinterpret_cast<Block*>(
            reinterpret_cast<char*>(p) - sizeof(Block));
        b->free = true;
        // 相邻空闲块合并
        for (Block* cur = head_; cur; cur = cur->next) {
            while (cur->free && cur->next && cur->next->free) {
                cur->size += cur->next->size;
                cur->next  = cur->next->next;
            }
        }
    }
};

关键点在于：alloc 返回的是块头之后的地址，dealloc 通过指针回退 sizeof(Block) 找回块头；切分阈值保证剩余空间至少能容纳一个块头加最小数据，避免产生无法使用的碎片。合并循环在每次释放后线性扫描，块数较少时可接受，块数多时应改为按地址排序的伙伴系统或分级空闲链表。

五、性能对比实验设计

设计一个对照实验：固定总分配次数 N=100000，每次申请 4KB，使用后立即释放，分别走“直接 cudaMalloc/cudaFree”与“SimplePool”两条路径，用 cudaEvent 记录主机侧总耗时（池化路径因无设备同步，也可用 std::chrono）。为避免编译器优化掉分配，需对返回指针写入一个字节并读回校验。

表中数值为教学示例量级，实际结果随GPU型号、驱动版本与请求尺寸变化，应以本机实测为准。实验时应额外记录“池扩容次数”，若扩容次数远小于N，说明池容量设置合理；若频繁扩容，则需增大初始池或引入分级池。另一个值得观察的指标是碎片率：在混合尺寸请求下，统计最大可分配块与总空闲量之比，比值越低说明碎片越严重，此时应调整切分策略或引入按2的幂分级。