为啥在使用 cudaMallocManaged 时 NVIDIA Pascal GPU 运行 CUDA 内核速度很慢

Posted 2023-04-15

【中文标题】为啥在使用 cudaMallocManaged 时 NVIDIA Pascal GPU 运行 CUDA 内核速度很慢

【英文标题】：Why is NVIDIA Pascal GPUs slow on running CUDA Kernels when using cudaMallocManaged为什么在使用 cudaMallocManaged 时 NVIDIA Pascal GPU 运行 CUDA 内核速度很慢

【发布时间】：2017-02-08 12:20:12

【问题描述】：

我正在测试新的 CUDA 8 和 Pascal Titan X GPU，并希望我的代码能够加快速度，但由于某种原因它最终变慢了。我在 Ubuntu 16.04 上。

这是可以重现结果的最少代码：

CUDASample.cuh

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

CUDASample.cu

__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 来进一步隔离问题，并进行了一些分析并观察了一些有趣的结果。

CUDASample.cu

__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 ms，内核启动时间约为 194 us，这实际上导致了我希望看到的性能提升......

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

【问题讨论】：

1.向量相加并不是一个特别有趣的 GPU 速度测试。 2. 不可能准确地知道你在计时什么，或者如何计时。 3. 在 any GPU 上，在 4096 个元素上的向量添加内核不可能花费 ~70ms。 70us 更合理。这是一个很小的问题，几乎可以肯定您测量的是某种开销，而不是实际的 GPU 计算性能

将问题大小增加到 100M 个元素。修改您的代码以连续两次调用您的内核。然后使用nvprof 运行您的代码。内核的第二次调用应该在较新的 Titan X 上运行得更快。

@RobertCrovella 我用 nvvp 的结果进行了更新。请看一下。谢谢！

你好，你有什么驱动版本？

@harrism 最新367.44

【参考方案1】：

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

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

使用 CUDA 8 和 Pascal GPU，数据迁移通过按需分页发生。在内核启动时，默认情况下，没有数据显式迁移到设备 (*)。当 GPU 设备代码尝试访问不在 GPU 内存中的特定页面中的数据时，将发生页面错误。此页面错误的最终结果是：

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

此过程将根据需要重复，因为 GPU 代码会触及不同的数据页面。除了实际移动数据所花费的时间之外，在处理页面错误时，上述第 2 步中涉及的操作序列还涉及一些延迟。由于此过程将一次移动一页数据，因此它可能比使用cudaMemcpy 或通过导致在内核启动时移动所有数据的pre-Pascal UM 安排一次移动所有数据的效率要低得多（无论是否需要，也不管内核代码何时真正需要它）。

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

但是，这个特定的代码示例并没有好处。这是意料之中的，因此建议使用 cudaMemPrefetchAsync() 调用在内核启动之前使行为与之前的（例如 maxwell）行为/性能保持一致。

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

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

正如我在 cmets 中提到的，如果您创建了一个包含两个顺序内核调用的测试案例，您会观察到即使在 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 博客文章中提供了更多信息。

【讨论】：

使用 cudaMemPrefetchAsync 为我解决了这个问题。帕斯卡和麦克斯韦之间的区别也得到了很好的解释！

@user3667089 不开玩笑。这个答案很精彩。我怀疑任何人都可以更简洁有效地解释它。

【参考方案2】：

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

我相信这不仅仅是使用 cudaMallocManaged 数据调用内核。我的经验是内核或设备方法的每次调用都有这种行为，并且效果加起来。此外，体积渲染的基础部分是提供的不带托管内存的 CudaSample，它可以在 Maxwell 和帕斯卡 GPU (1080、1060,980Ti,980,960) 上按预期运行。

我昨天才发现这个错误，因为我们将所有研究系统都更改为 pascal。我将在接下来的几天内在 980 和 1080 上分析我的软件。我还不确定是否应该在 NVIDIA 开发者专区报告错误。

【讨论】：

如果这不是 NVIDIA 的错误，请告诉我您是否设法解决此问题。我仍然坚持使用 maxwell GPU，因为我不想将所有代码从 cudaMallocManaged 更改为 cudaMalloc

【参考方案3】：

这是 NVIDIA 在 Windows 系统上的一个 BUG，发生在 PASCAL 架构上。

几天前我就知道了，但是因为我正在度假而没有互联网连接，所以无法在这里写下来。

详情见：https://devblogs.nvidia.com/parallelforall/unified-memory-cuda-beginners/ 来自 NVIDIA 的 Mark Harris 确认了这个错误。应该用 CUDA 9 纠正它。他还告诉它应该传达给 Microsoft 以帮助解决问题。但是直到现在我还没有找到合适的 Microsoft 错误报告页面。

【讨论】：

你混淆了两种不同的东西。您所指的错误是 Windows WDDM 上托管内存的实现问题。它与 OP 报告的问题不同，并且 OP 已经确认接受的答案修复了观察到的问题。事实上，CUDA 9 实际上已经放弃了对 windows 请求分页的支持，并且 UM 的 windows 行为恢复到 pre-pascal 机制（无请求分页），即使对于 pascal 设备及更高版本也是如此。见here。