使用cudaMallocManaged时，为什么NVIDIA Pascal GPU在运行CUDA内核时会变慢

Question

我正在测试新的CUDA 8以及Pascal Titan X GPU，并期待我的代码加速，但由于某种原因它最终会变慢。我在Ubuntu 16.04上。

以下是可以重现结果的最小代码：

CUDASample.cuh

class CUDASample{
 public:
  void AddOneToVector(std::vector<int> &in);
};

CUDA sample.粗

__global__ static void CUDAKernelAddOneToVector(int *data)
{
  const int x  = blockIdx.x * blockDim.x + threadIdx.x;
  const int y  = blockIdx.y * blockDim.y + threadIdx.y;
  const int mx = gridDim.x * blockDim.x;

  data[y * mx + x] = data[y * mx + x] + 1.0f;
}

void CUDASample::AddOneToVector(std::vector<int> &in){
  int *data;
  cudaMallocManaged(reinterpret_cast<void **>(&data),
                    in.size() * sizeof(int),
                    cudaMemAttachGlobal);

  for (std::size_t i = 0; i < in.size(); i++){
    data[i] = in.at(i);
  }

  dim3 blks(in.size()/(16*32),1);
  dim3 threads(32, 16);

  CUDAKernelAddOneToVector<<<blks, threads>>>(data);

  cudaDeviceSynchronize();

  for (std::size_t i = 0; i < in.size(); i++){
    in.at(i) = data[i];
  }

  cudaFree(data);
}

Main.cpp的

std::vector<int> v;

for (int i = 0; i < 8192000; i++){
  v.push_back(i);
}

CUDASample cudasample;

cudasample.AddOneToVector(v);

唯一的区别是NVCC标志，对于Pascal Titan X来说是：

-gencode arch=compute_61,code=sm_61-std=c++11;

对于旧Maxwell Titan X来说：

-gencode arch=compute_52,code=sm_52-std=c++11;

编辑：以下是运行NVIDIA Visual Profiling的结果。

对于旧的Maxwell Titan，内存传输的时间约为205 ms，内核启动时间约为268 us。

对于Pascal Titan来说，内存传输的时间大约是202毫秒，而内核的启动时间大约为8343 us，这让我觉得有些不对劲。

我进一步通过将cudaMallocManaged替换为好的旧cudaMalloc并进行一些分析并观察一些有趣的结果来解决问题。

CUDA sample.粗

__global__ static void CUDAKernelAddOneToVector(int *data)
{
  const int x  = blockIdx.x * blockDim.x + threadIdx.x;
  const int y  = blockIdx.y * blockDim.y + threadIdx.y;
  const int mx = gridDim.x * blockDim.x;

  data[y * mx + x] = data[y * mx + x] + 1.0f;
}

void CUDASample::AddOneToVector(std::vector<int> &in){
  int *data;
  cudaMalloc(reinterpret_cast<void **>(&data), in.size() * sizeof(int));
  cudaMemcpy(reinterpret_cast<void*>(data),reinterpret_cast<void*>(in.data()), 
             in.size() * sizeof(int), cudaMemcpyHostToDevice);

  dim3 blks(in.size()/(16*32),1);
  dim3 threads(32, 16);

  CUDAKernelAddOneToVector<<<blks, threads>>>(data);

  cudaDeviceSynchronize();

  cudaMemcpy(reinterpret_cast<void*>(in.data()),reinterpret_cast<void*>(data), 
             in.size() * sizeof(int), cudaMemcpyDeviceToHost);

  cudaFree(data);
}

对于旧的Maxwell Titan，内存传输的时间大约为5 ms，内核启动时间约为264 us。

对于Pascal Titan来说，内存传输的时间大约是5毫秒，内核启动时间约为194 us，这实际上导致性能提升，我希望看到...

使用cudaMallocManaged时，为什么Pascal GPU在运行CUDA内核时这么慢？如果我必须将使用cudaMallocManaged的所有现有代码恢复为cudaMalloc，那将是一个讽刺。这个实验还表明，使用cudaMallocManaged的内存传输时间比使用cudaMalloc慢很多，这也感觉有些不对劲。如果使用这会导致运行缓慢，即使代码更容易，这也是不可接受的，因为使用CUDA而不是普通C ++的整个目的是加快速度。我做错了什么，为什么我会观察这种结果？

Answer 1

在使用Pascal GPU的CUDA 8下，统一内存（UM）系统下的托管内存数据迁移通常会与以前的体系结构不同，并且您正在体验这种影响。 （另请参阅最后关于Windows的CUDA 9更新行为的说明。）

使用以前的体系结构（例如Maxwell），特定内核调用使用的托管分配将在内核启动时一次性迁移，就像您调用cudaMemcpy自己移动数据一样。

使用CUDA 8和Pascal GPU，数据迁移通过请求分页进行。在内核启动时，默认情况下，没有数据显式迁移到设备（*​​）。当GPU设备代码尝试访问未驻留在GPU内存中的特定页面中的数据时，将发生页面错误。此页面错误的净效果是：

1. 导致GPU内核代码（访问该页面的一个或多个线程）停止（直到第2步完成）
2. 导致该内存页面从CPU迁移到GPU

当GPU代码触及各种数据页面时，将根据需要重复此过程。除了实际移动数据所花费的时间之外，上述步骤2中涉及的操作序列还涉及处理页面错误时的一些延迟。由于此过程将一次一页地移动数据，因此使用cudaMemcpy或者通过pre-Pascal UM排列导致所有数据在内核启动时移动，可能显着低于一次移动所有数据的效率（无论是它是否需要，无论内核代码何时实际需要它。

这两种方法都有其优点和缺点，我不想就优点或各种意见或观点进行辩论。请求分页过程为Pascal GPU提供了许多重要的特性和功能。

然而，这个特定的代码示例没有受益。这是预料之中的，因此推荐使用与先前（例如maxwell）行为/性能一致的行为是在内核启动之前使用cudaMemPrefetchAsync()调用。

您将使用CUDA流语义在内核启动之前强制完成此调用（如果内核启动未指定流，则可以为stream参数传递NULL，以选择默认流）。我相信这个函数调用的其他参数是非常明显的。

在内核调用之前调用此函数，覆盖有问题的数据，您不应该在Pascal情况下观察到任何页面错误，并且配置文件行为应该与Maxwell情况类似。

正如我在评论中提到的，如果你创建了一个依次涉及两个内核调用的测试用例，你会发现第二个调用即使在Pascal情况下也会以大约全速运行，因为所有数据都已经被迁移了通过第一次内核执行到GPU端。因此，不应将此预取功能的使用视为强制或自动，但应谨慎使用。存在GPU可能能够在某种程度上隐藏页面错误的延迟的情况，并且显然已经驻留在GPU上的数据不需要被预取。

请注意，上面步骤1中提到的“失速”可能会产生误导。内存访问本身不会触发停顿。但是，如果操作实际需要所请求的数据，例如乘以，然后经线将在乘法运算时停止，直到必要的数据变为可用。一个相关的观点是，以这种方式从主机到设备的数据请求分页只是GPU可能隐藏在其延迟隐藏架构中的另一个“延迟”，如果有足够的其他可用“工作”参与的话至。

另外需要注意的是，在CUDA 9中，pascal及更高版本的请求 - 分页机制仅适用于Linux;以前对CUDA 8中宣传的Windows的支持已被删除。见here。在Windows上，即使对于Pascal设备及其他设备，从CUDA 9开始，UM机制与maxwell和之前的设备相同;在内核启动时，数据将迁移到GPU集群。

（*）这里的假设是数据在主机上“驻留”，即在管理的分配调用之后已经在CPU代码中“触摸”或初始化。托管分配本身创建与设备关联的数据页面，当CPU代码“触摸”这些页面时，CUDA运行时将要求页面驻留在主机内存中，以便CPU可以使用它们。如果执行分配但从不“触摸”CPU代码中的数据（可能是一种奇怪的情况），那么当内核运行时它实际上已经“驻留”在设备内存中，并且观察到的行为将是不同的。但对于这个特定的例子/问题，情况并非如此。

有关更多信息，请参阅this博客文章。

Answer 2

我可以在1060和1080上的三个程序中重现这个。例如我使用带有程序传递函数的voulme渲染，这在960上几乎是交互式实时，但在1080上是轻微的显示。所有数据都存储在只读纹理中，只有我的传输函数位于托管内存中。与我的其他代码不同，卷渲染运行速度特别慢，这与我的其他代码不同，我的传递函数从内核传递到其他设备方法。

我相信它不仅是使用cudaMallocManaged数据调用内核。我的expierence转到内核或设备方法的每次调用都有这种行为，效果加起来。此外，卷渲染的基础部分是提供的没有托管内存的CudaSample，它在Maxwell上运行，如pascal GPU（1080,1060,980Ti，980,960）。

我昨天发现了这个错误，因为我们将所有的oure reaserch系统改为pascal。我将在接下来的日子里以980的comapre配置我的软件到1080.我还不确定我是否应该报告NVIDIA开发人员区域中的错误。

Answer 3

它是Windows系统上NVIDIA的一个BUG，它与PASCAL架构一起发生。

我知道这几天了，但是不能在这里写，因为我在没有网络连接的假期。

有关详细信息，请参阅以下评论：https://devblogs.nvidia.com/parallelforall/unified-memory-cuda-beginners/来自NVIDIA的Mark Harris确认了Bug。它应该用CUDA 9来纠正。他还告诉它应该传达给微软以帮助解决问题。但到目前为止，我还没有找到合适的Microsoft Bug Report Page。