3. 图像数据在GPU上的内存布局与传输优化

在CUDA图像处理中，算法本身的并行度往往不是瓶颈，真正的瓶颈常常出现在“数据怎么放、怎么搬”上。一张1920×1080的8位灰度图约2MB，彩色图约6MB，若按每秒30帧处理，仅输入带宽就接近180MB/s，再加上输出与中间结果，PCIe与显存之间的往返很容易成为限制因素。本章围绕四类GPU内存的选型，以及cudaMemcpy2D、pinned memory与异步流如何协同隐藏传输延迟展开。

一、四类内存的定位与选型依据

GPU上的内存并非只有“显存”一种。全局内存容量大、可读写、所有线程可见，但访问延迟高，且未合并访问时带宽浪费严重；纹理内存本质是全局内存的只读缓存路径，通过纹理单元访问，具备空间局部性缓存和硬件插值，适合图像这类二维邻域访问；常量内存容量小（通常64KB），位于片外但带常量缓存，当同一warp内所有线程读取同一地址时广播效率极高，适合卷积核权重、颜色查找表等；共享内存位于SM内部，延迟极低，由线程块内所有线程共享，适合做分块（tiling）复用，把全局内存的重复读取转化为一次加载、多次使用。

选型的核心判断是“访问模式”而非“数据大小”。若每个像素只被读取一次且访问连续，用全局内存配合合并访问即可；若存在大量邻域重叠（如3×3、5×5卷积），应把图像分块载入共享内存；若访问带有随机性或需要边界处理，纹理内存的缓存与clamp寻址更省心；若参数被所有线程以相同方式读取，常量内存是最优解。实际工程中常常是组合使用：输入图像走纹理或共享内存，卷积核走常量内存，输出写回全局内存。

二、cudaMemcpy2D：处理带步长的图像行

图像在内存中通常不是完全紧凑的。例如从相机采集的一帧可能带有行对齐填充（pitch），或者我们只关心大图中的某个ROI区域。此时逐行调用cudaMemcpy会产生大量小传输，效率极低。cudaMemcpy2D正是为此设计：它一次调用即可完成“带源步长、带目标步长”的二维拷贝。

// 假设 src 为 pinned 主机内存，spitch 为源行字节数
// dst 为设备内存，dpitch 为设备行字节数
// 拷贝 width 字节宽、height 行的图像区域
cudaMemcpy2D(dst, dpitch,
             src, spitch,
             width, height,
             cudaMemcpyHostToDevice);

关键参数是pitch（行间距，单位字节）。若源图每行有效像素为W字节，但实际存储行距为spitch，则拷贝时以spitch为步长逐行取W字节。设备端同理。这样既避免了逐行拷贝的调用开销，又能正确处理对齐填充。对于ROI提取，只需调整src指针偏移与width/height即可。需要注意的是，pitch通常建议为256字节的倍数以获得最佳合并访问，cudaMallocPitch可自动完成这一对齐。

三、Pinned Memory：让传输真正跑满PCIe

默认的cudaMalloc分配的是可分页（pageable）主机内存。当从这种内存拷贝到设备时，驱动必须先将其复制到一块临时的pinned缓冲区，再DMA到显存，相当于多了一次内存拷贝，且无法与计算重叠。使用cudaHostAlloc或cudaMallocHost分配的pinned内存则锁定在物理内存中，DMA可直接访问，带宽通常能提升2～3倍，并且是异步传输的前提。

float *h_data;
cudaHostAlloc(&h_data, bytes, cudaHostAllocDefault);
// ... 使用 h_data ...
cudaFreeHost(h_data);

代价是pinned内存不可被操作系统换出，过量分配会拖慢系统。实践中常采用双缓冲或内存池：预先分配2～3块pinned缓冲区循环使用，而不是每帧临时分配。对于图像流水线，输入帧与输出帧各用一块pinned缓冲即可满足大多数场景。

四、异步流：用计算掩盖传输

即使使用了pinned内存，如果调用同步的cudaMemcpy，CPU仍会阻塞等待传输完成，GPU在此期间可能空闲。真正的优化是把传输和计算放入不同的CUDA流，让DMA引擎与SM并行工作。典型的三段式流水线是：流A传输第n+1帧的同时，流B计算第n帧，流C写回第n-1帧的结果。

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

for (int f = 0; f < frames; f++) {
    int cur = f % 2;
    cudaMemcpyAsync(d_in[cur], h_in[cur], bytes,
                    cudaMemcpyHostToDevice, s[cur]);
    kernel<<<grid, block, 0, s[cur]>>>(d_in[cur], d_out[cur]);
    cudaMemcpyAsync(h_out[cur], d_out[cur], bytes,
                    cudaMemcpyDeviceToHost, s[cur]);
}

要真正重叠，必须满足三个条件：主机内存是pinned的；使用cudaMemcpyAsync而非同步版本；不同流之间没有隐式同步。此外，若使用默认流（stream 0），它会与所有其他流串行，因此应显式创建非默认流。判断是否成功隐藏延迟，可用Nsight Systems观察时间线：理想情况下传输条与计算条应错开重叠，总耗时接近max(传输, 计算)而非两者之和。

五、一个可复用的参数参考

需要强调的是，表中收益为典型场景下的经验范围，实际数值取决于图像尺寸、kernel复杂度与PCIe版本。建议在自己的目标平台上用CUDA Events测量各阶段耗时，再决定是否值得引入共享内存分块或多流。内存布局与传输优化没有万能公式，只有与访问模式匹配的方案。