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

2023-11-29 11:33:28

我正在测试新的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博客文章.

码农公寓

相关文章