在CUDA核函数中初始化数据仍出现GPU页错误的原因排查

我是CUDA/C++新手，正在学习Unified Memory（统一内存），参考了英伟达官方入门教程。为减少数据迁移开销，我采用教程中核函数初始化数据的示例代码，但删除末尾的CPU结果检查循环后，通过NVPROF性能分析发现仍存在GPU页错误，对此感到困惑，想了解其中原因。

示例代码 #

#include <iostream>
#include <math.h>

// initialize arrays on device
__global__ void init(int n, float *x, float *y) {
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }
}

// CUDA kernel to add elements of two arrays
__global__ void add(int n, float *x, float *y){
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = index; i < n; i += stride){
        y[i] = x[i] + y[i];
    }
}

int main(void)
{
    int N = 1<<20;
    float *x, *y;

    // Allocate Unified Memory -- accessible from CPU or GPU
    cudaMallocManaged(&x, N*sizeof(float));
    cudaMallocManaged(&y, N*sizeof(float));

    // Launch kernel on 1M elements on the GPU
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;

    init<<<numBlocks, blockSize>>>(N, x, y);
    add<<<numBlocks, blockSize>>>(N, x, y);

    // Wait for GPU to finish before accessing on host
    cudaDeviceSynchronize();

    // Free memory
    cudaFree(x);
    cudaFree(y);

    return 0;
}

NVPROF性能分析结果 #

==4242== NVPROF is profiling process 4242, command: /content/src/add_unifmem_initonkernel
==4242== Profiling application: /content/src/add_unifmem_initonkernel
==4242== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   96.00%  1.4178ms         1  1.4178ms  1.4178ms  1.4178ms  init(int, float*, float*)
                    4.00%  59.070us         1  59.070us  59.070us  59.070us  add(int, float*, float*)
      API calls:   99.21%  263.47ms         2  131.74ms  54.879us  263.42ms  cudaMallocManaged
                    0.54%  1.4273ms         1  1.4273ms  1.4273ms  1.4273ms  cudaDeviceSynchronize
                    0.15%  401.83us         2  200.91us  197.33us  204.49us  cudaFree
                    0.05%  120.55us       101  1.1930us     139ns  50.860us  cuDeviceGetAttribute
                    0.04%  96.692us         2  48.346us  40.043us  56.649us  cudaLaunchKernel
                    0.01%  28.565us         1  28.565us  28.565us  28.565us  cuDeviceGetName
                    0.00%  6.9460us         1  6.9460us  6.9460us  6.9460us  cuDeviceGetPCIBusId
                    0.00%  2.0890us         3     696ns     225ns  1.5490us  cuDeviceGetCount
                    0.00%  1.0370us         2     518ns     314ns     723ns  cuDeviceGet
                    0.00%     502ns         1     502ns     502ns     502ns  cuDeviceTotalMem
                    0.00%     500ns         1     500ns     500ns     500ns  cuModuleGetLoadingMode
                    0.00%     230ns         1     230ns     230ns     230ns  cuDeviceGetUuid

==4242== Unified Memory profiling result:
Device "Tesla T4 (0)"
   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
      13         -         -         -           -  1.695805ms  Gpu page fault groups

按照教程说明，删除结果检查循环后不应出现GPU页错误，但实际仍存在，请问我忽略了什么？

解答 #

你遇到的GPU页错误，本质是统一内存的初始页分配机制导致的：

• 用cudaMallocManaged分配的内存，默认是在CPU端首次触达时分配物理页，但你的代码里是GPU先访问这些内存。当GPU首次访问未分配物理页的统一内存区域时，会触发页错误，此时CUDA runtime会在GPU端分配对应的物理页，再映射到虚拟地址空间，这个过程就会被NVPROF统计为GPU页错误。

• 教程里说删除CPU检查循环后不应出现页错误，可能是指避免了CPU触发的页迁移（也就是把GPU的数据迁回CPU），但GPU首次访问时的初始页分配带来的页错误，依然会存在——这是统一内存的正常初始化流程，不是数据迁移的开销，和你理解的"页错误"可能不是同一个概念。

• 如果你想消除这些初始页错误，可以在GPU访问前，用cudaMemPrefetchAsync把统一内存预取到GPU：

  // 在启动核函数前添加：
  cudaMemPrefetchAsync(x, N*sizeof(float), 0);
  cudaMemPrefetchAsync(y, N*sizeof(float), 0);
  

  这样runtime会提前在GPU端分配并映射好物理页，GPU访问时就不会触发页错误了。

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