为什么在使用 cudaMallocManaged 时 NVIDIA Pascal GPU 在 运行 CUDA 内核上运行缓慢
Why is NVIDIA Pascal GPUs slow on running CUDA Kernels when using cudaMallocManaged
我正在测试新的 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;
编辑：这是 运行ning NVIDIA Visual Profiling 的结果。
对于旧的 Maxwell Titan，内存传输时间约为 205 毫秒，内核启动时间约为 268 微秒。
对于 Pascal Titan，内存传输时间约为 202 毫秒，内核启动时间约为长得离谱的 8343 微秒，这让我相信出了什么问题。
我通过将 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 毫秒，内核启动时间约为 264 微秒。
对于 Pascal Titan，双向内存传输时间约为 5 毫秒，内核启动时间约为 194 微秒，这实际上导致了我希望看到的性能提升...
为什么在使用 cudaMallocManaged 时 Pascal GPU 在 运行ning CUDA 内核上这么慢？如果我必须将所有使用 cudaMallocManaged 的现有代码还原为 cudaMalloc，那将是一种讽刺。这个实验也说明了使用cudaMallocManaged的内存传输时间比使用cudaMalloc慢了很多，也感觉哪里不对。如果使用它导致 运行 时间变慢，即使代码更简单，这应该是不可接受的，因为使用 CUDA 而不是普通 C++ 的全部目的是加快速度。我做错了什么，为什么我观察到这种结果？
我可以在 1060 和 1080 上用三个程序重现这一点。例如，我使用带有程序传递函数的体积渲染，它在 960 上几乎是实时交互的，但在 1080 上只是一个轻微的显示。所有数据都存储在只读纹理中，只有我的传递函数在托管内存中。与我的其他代码不同，体积渲染运行特别慢，这是因为与我的其他代码不同，我的传递函数从内核传递到其他设备方法。
我相信这不仅仅是用 cudaMallocManaged 数据调用内核。我的经验是内核或设备方法的每次调用都有这种行为并且效果加起来。此外，体积渲染的基础部分是提供的不带托管内存的 CudaSample，它按预期在 Maxwell 和 Pascal GPU（1080、1060,980Ti、980,960）上运行。
我昨天才发现这个错误，因为我们把所有的研究系统都改成了 pascal。在接下来的几天里，我将在 980 和 1080 之间分析我的软件。我还不确定我是否应该在 NVIDIA 开发者专区报告错误。
在带有 Pascal GPU 的 CUDA 8 下，统一内存 (UM) 机制下的托管内存数据迁移通常与以前的架构不同，您正在体验这种影响。 （另请参阅末尾关于 windows 的 CUDA 9 更新行为的注释。）
对于以前的架构（例如 Maxwell），特定内核调用使用的托管分配将在内核启动时立即全部迁移，就像您调用 cudaMemcpy 自己移动数据一样。
使用 CUDA 8 和 Pascal GPU，数据迁移通过请求分页发生。在内核启动时，默认情况下，没有数据显式迁移到设备 (*)。当 GPU 设备代码试图访问不在 GPU 内存中的特定页面中的数据时，将发生页面错误。此页面错误的最终结果是：
- 导致 GPU 内核代码（访问页面的一个或多个线程）停止（直到第 2 步完成）
- 导致那页内存从 CPU 迁移到 GPU
此过程将根据需要重复，因为 GPU 代码会触及不同的数据页。除了实际移动数据所花费的时间之外，上述步骤 2 中涉及的操作序列还涉及一些 延迟 处理页面错误。由于此过程将一次移动一页数据，因此与使用 cudaMemcpy 或通过导致所有数据在内核启动（无论是否需要，也不管内核代码何时真正需要它）。
两种方法各有利弊，我不想争论优劣或各种意见或观点。按需分页过程为 Pascal GPU 启用了许多重要的特性和功能。
但是，这个特定的代码示例没有任何好处。这是预料之中的，因此为了使行为与之前的（例如 maxwell）保持一致，建议使用 behavior/performance 是在内核启动之前调用 cudaMemPrefetchAsync()。
您将使用 CUDA 流语义强制此调用在内核启动之前完成（如果内核启动未指定流，您可以将流参数 NULL 传递给 select默认流）。我相信这个函数调用的其他参数是不言自明的。
在你的内核调用之前调用这个函数，覆盖有问题的数据，你应该不会在 Pascal 案例中观察到任何页面错误，并且配置文件行为应该类似于 Maxwell 案例。
正如我在评论中提到的，如果您创建了一个依次涉及两个内核调用的测试用例，您会发现即使在 Pascal 案例中，第二次调用也几乎全速运行，因为所有数据已经通过第一次内核执行迁移到GPU端。因此，不应将此预取功能的使用视为强制或自动，而应谨慎使用。在某些情况下，GPU 可能能够在某种程度上隐藏页面错误的延迟，并且显然不需要预取已经驻留在 GPU 上的数据。
请注意，上面第 1 步中提到的 "stall" 可能具有误导性。内存访问本身不会触发停顿。但是，如果请求的数据实际上是某个操作所需要的，例如a multiply，那么 warp 将在 multiply 操作处停止，直到必要的数据可用。那么，相关的一点是以这种方式从主机到设备的数据需求分页只是另一个 "latency" GPU 可能隐藏在它的延迟隐藏架构中，如果有足够的其他可用 "work" 参加。
作为补充说明，在 CUDA 9 中，pascal 及更高版本的请求分页机制仅在 linux 上可用；先前在 CUDA 8 中宣传的对 Windows 的支持已被删除。参见 here。在 windows 上，即使对于 Pascal 设备及更高版本，从 CUDA 9 开始，UM 制度与 maxwell 和之前的设备相同；数据在内核启动时整体迁移到 GPU。
(*) 这里的假设是主机上的数据是 "resident"，即在托管分配调用之后已经 "touched" 或在 CPU 代码中初始化。托管分配本身会创建与设备关联的数据页，当 CPU 代码 "touches" 这些页面时，CUDA 运行时将请求分页必要的页面以驻留在主机内存中，以便 CPU 可以使用它们。如果您执行分配但从未 "touch" CPU 代码中的数据（可能是一种奇怪的情况），那么当内核运行时它实际上已经 "resident" 在设备内存中，并且观察到行为会有所不同。但对于这个特定的 example/question.
而言，情况并非如此
this 博客文章中提供了更多信息。
这是 NVIDIA 在 Windows 系统上的一个 BUG 出现在 PASCAL 架构上。
几天前我就知道了，但不能写在这里，因为我正在度假，没有网络连接。
详情见评论：https://devblogs.nvidia.com/parallelforall/unified-memory-cuda-beginners/
NVIDIA 的 Mark Harris 确认了这个错误。它应该用 CUDA 9 更正。他还告诉它应该与 Microsoft 沟通以帮助解决问题。但是直到现在我还没有找到合适的 Microsoft Bug Report Page。
我正在测试新的 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;
编辑：这是 运行ning NVIDIA Visual Profiling 的结果。
对于旧的 Maxwell Titan，内存传输时间约为 205 毫秒，内核启动时间约为 268 微秒。
对于 Pascal Titan，内存传输时间约为 202 毫秒，内核启动时间约为长得离谱的 8343 微秒，这让我相信出了什么问题。
我通过将 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 毫秒，内核启动时间约为 264 微秒。
对于 Pascal Titan，双向内存传输时间约为 5 毫秒，内核启动时间约为 194 微秒，这实际上导致了我希望看到的性能提升...
为什么在使用 cudaMallocManaged 时 Pascal GPU 在 运行ning CUDA 内核上这么慢？如果我必须将所有使用 cudaMallocManaged 的现有代码还原为 cudaMalloc，那将是一种讽刺。这个实验也说明了使用cudaMallocManaged的内存传输时间比使用cudaMalloc慢了很多，也感觉哪里不对。如果使用它导致 运行 时间变慢，即使代码更简单，这应该是不可接受的，因为使用 CUDA 而不是普通 C++ 的全部目的是加快速度。我做错了什么，为什么我观察到这种结果？
我可以在 1060 和 1080 上用三个程序重现这一点。例如，我使用带有程序传递函数的体积渲染，它在 960 上几乎是实时交互的，但在 1080 上只是一个轻微的显示。所有数据都存储在只读纹理中，只有我的传递函数在托管内存中。与我的其他代码不同，体积渲染运行特别慢，这是因为与我的其他代码不同，我的传递函数从内核传递到其他设备方法。
我相信这不仅仅是用 cudaMallocManaged 数据调用内核。我的经验是内核或设备方法的每次调用都有这种行为并且效果加起来。此外，体积渲染的基础部分是提供的不带托管内存的 CudaSample，它按预期在 Maxwell 和 Pascal GPU（1080、1060,980Ti、980,960）上运行。
我昨天才发现这个错误，因为我们把所有的研究系统都改成了 pascal。在接下来的几天里，我将在 980 和 1080 之间分析我的软件。我还不确定我是否应该在 NVIDIA 开发者专区报告错误。
在带有 Pascal GPU 的 CUDA 8 下，统一内存 (UM) 机制下的托管内存数据迁移通常与以前的架构不同，您正在体验这种影响。 （另请参阅末尾关于 windows 的 CUDA 9 更新行为的注释。）
对于以前的架构（例如 Maxwell），特定内核调用使用的托管分配将在内核启动时立即全部迁移，就像您调用 cudaMemcpy 自己移动数据一样。
使用 CUDA 8 和 Pascal GPU，数据迁移通过请求分页发生。在内核启动时，默认情况下，没有数据显式迁移到设备 (*)。当 GPU 设备代码试图访问不在 GPU 内存中的特定页面中的数据时，将发生页面错误。此页面错误的最终结果是：
- 导致 GPU 内核代码（访问页面的一个或多个线程）停止（直到第 2 步完成）
- 导致那页内存从 CPU 迁移到 GPU
此过程将根据需要重复，因为 GPU 代码会触及不同的数据页。除了实际移动数据所花费的时间之外，上述步骤 2 中涉及的操作序列还涉及一些 延迟 处理页面错误。由于此过程将一次移动一页数据，因此与使用 cudaMemcpy 或通过导致所有数据在内核启动（无论是否需要，也不管内核代码何时真正需要它）。
两种方法各有利弊，我不想争论优劣或各种意见或观点。按需分页过程为 Pascal GPU 启用了许多重要的特性和功能。
但是，这个特定的代码示例没有任何好处。这是预料之中的，因此为了使行为与之前的（例如 maxwell）保持一致，建议使用 behavior/performance 是在内核启动之前调用 cudaMemPrefetchAsync()。
您将使用 CUDA 流语义强制此调用在内核启动之前完成（如果内核启动未指定流，您可以将流参数 NULL 传递给 select默认流）。我相信这个函数调用的其他参数是不言自明的。
在你的内核调用之前调用这个函数，覆盖有问题的数据，你应该不会在 Pascal 案例中观察到任何页面错误，并且配置文件行为应该类似于 Maxwell 案例。
正如我在评论中提到的，如果您创建了一个依次涉及两个内核调用的测试用例，您会发现即使在 Pascal 案例中，第二次调用也几乎全速运行，因为所有数据已经通过第一次内核执行迁移到GPU端。因此，不应将此预取功能的使用视为强制或自动，而应谨慎使用。在某些情况下，GPU 可能能够在某种程度上隐藏页面错误的延迟，并且显然不需要预取已经驻留在 GPU 上的数据。
请注意，上面第 1 步中提到的 "stall" 可能具有误导性。内存访问本身不会触发停顿。但是，如果请求的数据实际上是某个操作所需要的，例如a multiply，那么 warp 将在 multiply 操作处停止，直到必要的数据可用。那么，相关的一点是以这种方式从主机到设备的数据需求分页只是另一个 "latency" GPU 可能隐藏在它的延迟隐藏架构中，如果有足够的其他可用 "work" 参加。
作为补充说明，在 CUDA 9 中，pascal 及更高版本的请求分页机制仅在 linux 上可用；先前在 CUDA 8 中宣传的对 Windows 的支持已被删除。参见 here。在 windows 上，即使对于 Pascal 设备及更高版本，从 CUDA 9 开始，UM 制度与 maxwell 和之前的设备相同；数据在内核启动时整体迁移到 GPU。
(*) 这里的假设是主机上的数据是 "resident"，即在托管分配调用之后已经 "touched" 或在 CPU 代码中初始化。托管分配本身会创建与设备关联的数据页，当 CPU 代码 "touches" 这些页面时，CUDA 运行时将请求分页必要的页面以驻留在主机内存中，以便 CPU 可以使用它们。如果您执行分配但从未 "touch" CPU 代码中的数据（可能是一种奇怪的情况），那么当内核运行时它实际上已经 "resident" 在设备内存中，并且观察到行为会有所不同。但对于这个特定的 example/question.
而言，情况并非如此this 博客文章中提供了更多信息。
这是 NVIDIA 在 Windows 系统上的一个 BUG 出现在 PASCAL 架构上。
几天前我就知道了，但不能写在这里，因为我正在度假，没有网络连接。
详情见评论：https://devblogs.nvidia.com/parallelforall/unified-memory-cuda-beginners/ NVIDIA 的 Mark Harris 确认了这个错误。它应该用 CUDA 9 更正。他还告诉它应该与 Microsoft 沟通以帮助解决问题。但是直到现在我还没有找到合适的 Microsoft Bug Report Page。