4. IGX内存管理与数据搬运优化

在IGX这类边缘AI平台上，GPU算力往往不是瓶颈，真正拖慢推理吞吐的常常是数据在CPU内存、系统内存与GPU显存之间的往返搬运。本节围绕统一内存（Unified Memory）、零拷贝（Zero-Copy）、锁页内存（Pinned Memory）、异步传输（Async Copy）四条主线，说明它们各自解决什么问题、代价是什么，以及如何用带宽指标定位瓶颈。

一、先建立内存层级与搬运路径的直觉

IGX上的内存大致分为三层：CPU侧的可分页主机内存（Pageable）、锁页主机内存（Pinned/Page-Locked），以及GPU侧的设备显存（Device Memory）。可分页内存由操作系统按需换页，CUDA无法直接对其做DMA，必须先拷贝到一块临时的锁页缓冲区，再传给GPU——这就是"隐式同步+双拷贝"的来源。锁页内存则被固定在物理内存中，GPU的DMA引擎可以直接读取，因此是高效传输的前提。

统一内存是另一条思路：用 cudaMallocManaged 分配一块CPU和GPU都能寻址的地址空间，由驱动按需在两端迁移页面。它让代码更简洁，但页面迁移（page fault）会带来不可预测的延迟，在实时性要求高的推理场景中需要谨慎使用。

二、锁页内存与异步传输：把拷贝藏进计算里

锁页内存的分配方式很简单：cudaMallocHost(&h_ptr, bytes) 或 cudaHostAlloc。它的代价是占用不可换出的物理内存，分配过多会拖累整机。实践中通常只对"每帧都要传"的输入张量、输出缓冲区使用锁页内存，而不是把所有主机内存都锁住。

异步传输的关键是 cudaMemcpyAsync 配合非默认流（non-default stream）。默认流上的拷贝会与计算隐式同步，无法重叠；只有把拷贝和kernel放进同一个非默认流，或用多个流分块流水，才能让DMA引擎与SM并行工作。典型的分块流水写法如下：

const int N = 1 << 22;          // 4M floats
const int CHUNK = 1 << 20;      // 每块 1M floats
const int nStreams = 4;

float *h_a, *d_a;
cudaMallocHost(&h_a, N * sizeof(float));   // 锁页
cudaMalloc(&d_a, N * sizeof(float));

cudaStream_t s[nStreams];
for (int i = 0; i < nStreams; ++i) cudaStreamCreate(&s[i]);

for (int i = 0; i < N / CHUNK; ++i) {
    int off = i * CHUNK;
    cudaMemcpyAsync(d_a + off, h_a + off, CHUNK * sizeof(float),
                    cudaMemcpyHostToDevice, s[i % nStreams]);
    kernel<<<(CHUNK + 255) / 256, 256, 0, s[i % nStreams]>>>(d_a + off, CHUNK);
}

这段代码把4MB数据切成4块、用4条流轮转，H2D拷贝与kernel在不同流上交错执行。判断是否真正重叠，不能只看代码，要用 nsys profile 或 ncu 看时间线上拷贝条与计算条是否并行；如果它们仍然首尾相接，多半是流用错（用了默认流）或数据量太小、启动开销占主导。

三、零拷贝：让GPU直接读主机内存

零拷贝（Zero-Copy）通过 cudaHostAlloc(..., cudaHostAllocMapped) 分配映射锁页内存，再用 cudaHostGetDevicePointer 拿到设备侧指针，kernel就能通过PCIe/NVLink直接访问主机内存，省去显式H2D拷贝。它适合"数据只被访问一次、且访问量不大"的场景，例如小批量推理的输入特征、稀疏查表。但要注意：每次访问都要走互连总线，延迟远高于显存，若kernel反复读写同一块数据，零拷贝反而比"先拷进显存再算"更慢。

四、带宽瓶颈分析与实测方法

定位瓶颈的第一步是算清"理论下界"。以PCIe Gen4 x16为例，单向理论带宽约32 GB/s，实测有效带宽通常在20～26 GB/s。若一帧输入是8 MB，H2D理论耗时约 8MB / 25GB/s ≈ 0.32 ms；如果实测拷贝耗时是2 ms，说明有效带宽只有4 GB/s，问题多半出在可分页内存的隐式暂存、拷贝粒度太小或未用异步。

实测建议用 cudaEvent 计时，并区分三种配置：可分页+同步、锁页+同步、锁页+异步多流。下面是一个可直接套用的计时骨架：

cudaEvent_t t0, t1;
cudaEventCreate(&t0); cudaEventCreate(&t1);
cudaEventRecord(t0, s[0]);
for (int i = 0; i < N / CHUNK; ++i) {
    int off = i * CHUNK;
    cudaMemcpyAsync(d_a + off, h_a + off, CHUNK * sizeof(float),
                    cudaMemcpyHostToDevice, s[i % nStreams]);
}
cudaEventRecord(t1, s[0]);
cudaEventSynchronize(t1);
float ms; cudaEventElapsedTime(&ms, t0, t1);
printf("effective BW = %.2f GB/s\n",
       (N * sizeof(float)) / (ms * 1e-3) / 1e9);

把三种配置的 effective BW 列成表对比，就能直观看到锁页与异步带来的收益。经验上，从可分页同步切到锁页异步，有效带宽常能提升2～4倍；若提升不明显，应检查是否真的用了非默认流、块大小是否过小（建议每块≥256 KB）、以及是否触发了页错误。

最后提醒：统一内存虽然方便，但在IGX上做实时推理时，建议对延迟敏感的张量显式使用 cudaMemPrefetchAsync 预取到设备，或干脆改用锁页+异步的显式路径，避免运行中突发页迁移造成尾延迟抖动。带宽优化的本质，是让每一次字节搬运都发生在正确的内存类型上，并且尽可能被计算掩盖。