通过在内核中初始化数据而出现“GPU 页面错误”

Question

我是 CUDA/C++ 新手，正在研究统一内存。我找到了关于这个主题的this介绍。但是，我对其中一个示例有疑问。

为了减轻迁移开销，有一个在内核中初始化数据的示例：

#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;
}

在我之前放置的链接中，据说对于这种情况“仍然存在设备到主机页面错误，但这是由于程序末尾检查CPU结果的循环所致。”。但是，我已经删除了最后的循环，对此的分析是

==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 页面错误，但如果我正确的话，这种情况应该不会发生。

我在这里缺少什么？

Answer 1

您的

init

内核仍然遇到页面错误。您可以通过注意

init

内核（约 1400 微秒）和

add

内核（约 60 微秒）的持续时间之间的巨大时间差异来获得额外的线索。

原因是至少在两种情况下可能会发生页面错误，并且与同一核心问题相关：代码所触及的页面不存在于设备内存中。也许这种情况的典型情况是数据物理上存在于其他一些处理器上并且需要迁移。在这种情况下，页面错误的目的是触发迁移，当发生这种情况时，

nvprof

通常会报告与错误相关的附加数据，例如迁移的数据量、块的大小、迁移的数量等。

但是您的报告中缺少所有这些。这是第二种线索，表明这些页面错误的起源和目的略有不同。基本思想是一些分配器是“所谓的”“惰性分配器”。分配器创造了数据存在的可能性，但实际上并不为其分配内存位置。在这方面，cudaMallocManaged是一个惰性分配器。内存分配将在“第一次触摸”时发生。在你的例子中，第一次接触发生在

init

内核中。如果您想让这种效果完全消失，您将需要在某处实际实例化数据。在典型的编程用法中，您通常会通过在“某处”初始化数据来完成此操作，因此，如果您在

init

内核中执行此操作，您将遇到此类 GPU 页面错误，如果您在主机代码中执行此操作，然后允许当数据迁移到 GPU 时，您会在 CPU 代码中遇到此类页面错误。

对于此处的特定程序，可以采取的消除此影响的一种方法是在第一个 (

init

) 内核调用之前插入以下内容：

cudaMemPrefetchAsync(x, N*sizeof(float), 0);
cudaMemPrefetchAsync(y, N*sizeof(float), 0);

根据我对您的代码的测试，这应该会导致 

init

内核的持续时间下降到与

add

内核的持续时间大致相当，并且它应该消除所有页面错误，据我所知。我无法复制您的确切情况，因为我目前无法访问可以运行

nvprof

的 GPU，该 GPU 也具有有效的需求分页功能，但我看到的内容与您使用

nsys profile --stats=true ... 看到的类似

您可以在

此在线培训系列

的第 6 单元中获得有关托管内存行为的更多讨论。