c – 使用cudaMallocManaged时,为什么NVIDIA Pascal GPU在运行CUDA内核时运行缓慢

2019-09-23 19:14:30 阅读：1624 来源： 互联网

我正在测试新的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毫秒,而内核的启动时间约为194微秒,这实际上导致性能提升,我希望看到……

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

解决方法:

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

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

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

>导致GPU内核代码(访问该页面的一个或多个线程)停止(直到第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博客文章.

标签：c,cuda,gpgpu,nvidia,nvcc
来源： https://codeday.me/bug/20190923/1815686.html

本站声明： 1. iCode9 技术分享网（下文简称本站）提供的所有内容，仅供技术学习、探讨和分享；
2. 关于本站的所有留言、评论、转载及引用，纯属内容发起人的个人观点，与本站观点和立场无关；
3. 关于本站的所有言论和文字，纯属内容发起人的个人观点，与本站观点和立场无关；
4. 本站文章均是网友提供，不完全保证技术分享内容的完整性、准确性、时效性、风险性和版权归属；如您发现该文章侵犯了您的权益，可联系我们第一时间进行删除；
5. 本站为非盈利性的个人网站，所有内容不会用来进行牟利，也不会利用任何形式的广告来间接获益，纯粹是为了广大技术爱好者提供技术内容和技术思想的分享性交流网站。

ICode9

c – 使用cudaMallocManaged时,为什么NVIDIA Pascal GPU在运行CUDA内核时运行缓慢