CUDA统一内存设备到主机预取效果异常？三类技术疑问

测试环境 #

• NVIDIA Quadro P1000（笔记本）
• Ubuntu 18.04
• CUDA 11.8

测试代码（统一内存实现） #

constexpr uint32_t N{512};
constexpr uint32_t DATA_SIZE{sizeof(float) * N * N};

__managed__ float ma[N * N];
__managed__ float mb[N * N];
__managed__ float mc[N * N];

__global__ void kernel()
{
  for (uint32_t i{0}; i < N * N; ++i)
  {
    mc[i] = ma[i] + mb[i];
  }
}

int main(int argc, char *[])
{
  for (uint32_t i{0}; i < N * N; ++i)
  {
    ma[i] = 1.0f;
    mb[i] = 2.0f;
  }

  int deviceId{};
  gpuErrchk(cudaGetDevice(&deviceId));
  gpuErrchk(cudaMemPrefetchAsync(ma, DATA_SIZE, deviceId, nullptr));
  gpuErrchk(cudaMemPrefetchAsync(mb, DATA_SIZE, deviceId, nullptr));

  kernel<<<1, 1>>>();
  gpuErrchk(cudaPeekAtLastError());

  gpuErrchk(cudaMemPrefetchAsync(mc, DATA_SIZE, cudaCpuDeviceId, nullptr));

  gpuErrchk(cudaDeviceSynchronize());

  float result{0.0f};
  for (uint32_t i{0}; i < N * N; ++i)
  {
    result += mc[i];
  }

  return static_cast<int>(result);
}

性能分析结果 #

开启预取后的nvprof结果 #

==29300== Unified Memory profiling result:
Device "Quadro P1000 (0)"
   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
       2  1.0000MB  1.0000MB  1.0000MB  2.000000MB  164.9620us  Host To Device
      20  153.60KB  4.0000KB  1.0000MB  3.000000MB  266.0500us  Device To Host
      19         -         -         -           -  551.9440us  Gpu page fault groups
Total CPU Page faults: 9

• Host To Device（HtoD）的2次传输符合ma、mb各1MB的预取操作，但Device To Host（DtoH）存在异常：预取效果不明显，总传输量达3MB远超请求的1MB。

注释预取后的nvprof结果 #

==30051== Unified Memory profiling result:
Device "Quadro P1000 (0)"
   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
      20  102.40KB  4.0000KB  508.00KB  2.000000MB  189.9230us  Host To Device
      29  105.93KB  4.0000KB  512.00KB  3.000000MB  278.4960us  Device To Host
      24         -         -         -           -  1.311533ms  Gpu page fault groups
Total CPU Page faults: 14

• 预取仅将HtoD迁移次数从20降至2、DtoH从29降至20，效果有限。

cudaMalloc实现的性能对比 #

Type  Time(%)      Time     Calls       Avg       Min       Max  Name
            0.00%  164.80us         2  82.401us  82.209us  82.593us  [CUDA memcpy HtoD]
            0.00%  81.665us         1  81.665us  81.665us  81.665us  [CUDA memcpy DtoH]

• 统一内存性能与cudaMalloc差距明显。

技术疑问 #

1. 为何向主机预取内存几乎无明显效果（迁移次数29 vs 20）？
2. 为何传输到主机的数据量超出请求的1MB，达到3MB？
3. 为何即使开启预取，统一内存性能仍比cudaMalloc分配的设备内存慢一个数量级？

1. 主机预取效果不明显的原因 #

• 编译器优化干扰：开启-O3优化后，CPU端对mc的遍历求和操作可能被重排或优化，导致预取mc到主机的操作未完全覆盖后续访问的内存范围，触发额外按需迁移。
• 页粒度迁移特性：统一内存以4KB页为单位迁移，预取虽标记了整个mc数组，但如果后续CPU访问的页未被预取完全（或预取页被系统调度影响），仍会触发页故障和额外迁移。
• 异步预取同步时机：cudaMemPrefetchAsync是异步操作，即使后续调用cudaDeviceSynchronize，也可能存在预取未完成CPU就开始访问mc的情况，导致部分页仍需按需迁移。

2. DtoH传输量达3MB的原因 #

• Pascal架构限制：Quadro P1000属于Pascal架构，其统一内存仅支持基于页故障的按需迁移，无法精准只迁移mc的页。当CPU访问统一内存时，系统可能连带迁移ma、mb的页（即使设备未修改它们），总传输量达到三个数组的总和3MB。
• 编译器优化触发额外访问：-O3优化下，CPU对mc的求和操作可能被优化为向量访问，间接触发对ma、mb的隐性访问，导致这些数组的页被迁移回主机。
• 页迁移批量特性：系统处理页故障时，会批量迁移相邻页，即使仅需mc的部分页，也可能连带迁移ma、mb的相邻页，导致总传输量超标。

3. 统一内存性能落后于cudaMalloc的原因 #

• 页故障额外开销：即使开启预取，统一内存仍存在页故障处理开销（硬件中断、操作系统调度、页表更新等），而cudaMalloc的memcpy是批量DMA传输，无此类额外开销。
• 架构特性差异：Pascal架构的统一内存基于页迁移实现，远不如Volta及以后架构的细粒度统一内存高效；显式memcpy是直接的批量数据传输，效率更高。
• 预取局限性：预取仅能减少按需迁移次数，但无法消除统一内存本身的页管理开销，而显式memcpy的开销相对固定且更低。
• 单线程kernel放大差距：测试中kernel为单线程执行，本身耗时较长，导致统一内存的页迁移开销占比更高，与显式memcpy的差距被进一步放大。

内容的提问来源于stack exchange，提问作者nikitablack