我正在测试新的 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 的结果。
For the old Maxwell Titan, the time for memory transfer is around 205 ms, and the kernel launch is around 268 us.
For the Pascal Titan, the time for memory transfer is around 202 ms, and the kernel launch is around an insanely long 8343 us, which makes me believe something is wrong.
我通过将 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);
}
For the old Maxwell Titan, the time for memory transfer is around 5 ms both ways, and the kernel launch is around 264 us.
For the Pascal Titan, the time for memory transfer is around 5 ms both ways, and the kernel launch is around 194 us, which actually results in the performance increase I am hoping to see...
为什么使用 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 中涉及的操作顺序涉及一些latency除了实际移动数据所花费的时间之外，还需要处理页面错误。由于此过程一次将数据移动一页，因此它可能比一次移动所有数据的效率低得多，或者使用cudaMemcpy或者通过 Pascal 之前的 UM 安排，导致所有数据在内核启动时移动（无论是否需要，也无论内核代码何时实际需要它）。
两种方法都有其优点和缺点，我不想争论优点或不同的意见或观点。请求调页过程为 Pascal GPU 提供了许多重要的特性和功能。
然而，这个特定的代码示例并没有受益。这是预料之中的，因此建议使用在内核启动之前使用cudaMemPrefetchAsync() call.
您可以使用 CUDA 流语义强制此调用在内核启动之前完成（如果内核启动未指定流，您可以为流参数传递 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博客文章。