【发布时间】：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
标签： c++ cuda gpgpu nvidia nvcc