Tesla C2075 上的 cudaMemcpy 太慢了答案

【问题标题】：cudaMemcpy is too slow on Tesla C2075Tesla C2075 上的 cudaMemcpy 太慢了
【发布时间】：2013-01-09 13:36:16
【问题描述】：

我目前正在使用具有 2 个支持 cuda 的 GPU 的服务器：Quadro 400 和 Tesla C2075。我做了一个简单的向量加法测试程序。我的问题是，虽然 Tesla C2075 GPU 应该比 Quadro 400 更强大，但完成这项工作需要更多时间。我发现 cudaMemcpy 占用了大部分执行时间，并且在更强大的 gpu 上运行更慢。这是来源：

void get_matrix(float* arr1,float* arr2,int N1,int N2)
{
  int Nx,Ny;
  int n_blocks,n_threads;
  int dev=0; // 1
  float time;
  size_t size;
  clock_t start,end;
  cudaSetDevice(dev);
  cudaDeviceProp deviceProp;
  start = clock();
  cudaGetDeviceProperties(&deviceProp, dev);
  Nx=N1;
  Ny=N2;
  n_threads=256;
  n_blocks=(Nx*Ny+n_threads-1)/n_threads;
  size=Nx*Ny*sizeof(float);
  cudaMalloc((void**)&d_A,size);
  cudaMalloc((void**)&d_B,size);
  cudaMemcpy(d_A, arr1, size, cudaMemcpyHostToDevice);
  cudaMemcpy(d_B, arr2, size, cudaMemcpyHostToDevice);
  vector_add<<<n_blocks,n_threads>>>(d_A,d_B,size);
  cudaMemcpy(arr1, d_A, size, cudaMemcpyDeviceToHost);
  printf("Running device %s \n",deviceProp.name);
  end = clock();
  time=float(end-start)/float(CLOCKS_PER_SEC);
  printf("time = %e\n",time);
}

int main()
{
int const nx = 20000,ny = nx;
static float a[nx*ny],b[nx*ny];
for(int i=0;i<nx;i++)
  {
  for(int j=0;j<ny;j++)
  {
    a[j+ny*i]=j+10*i;
    b[j+ny*i]=-(j+10*i);
  }
}
get_matrix(a,b,nx,ny);
return 0;
}

输出是：

Running device Quadro 400
time = 1.100000e-01

Running device Tesla C2075
time = 1.050000e+00

我的问题是：

是否应该根据要使用的 GPU 修改代码？
代码中指定的块数、每个块的线程数与 GPU 上可用的多处理器数、每个多处理器的内核数之间是否存在任何联系？

我正在运行 Linux Open Suse 11.2。源码使用nvcc编译器（4.2版）编译。

感谢您的帮助！

【问题讨论】：

您没有正确测量时间。您测量的cudaMemcpy 时间实际上是内核执行时间（启动是非阻塞调用）。在内核启动之后和cudaMemcpy 调用之前添加cudaDeviceSynchronize() 并观察复制调用“加速”。
您还应该检查每个 API 调用的返回状态（请参阅this question 了解更多信息）。两个 GPU 之间的“性能”差异很可能是由于代码没有在较慢的 GPU 上运行，但是因为您没有任何错误检查，所以您根本不会注意到它。
只是添加到@talonmies 错误检查注释：您的网格尺寸 (n_blocks) 是 (20000*20000 + 255)/256 = 1562500。但是 grid.x 的最大尺寸是 65535你提到的设备。代码没有按照您的想法执行。请进行错误检查。

标签： cuda gpu nvidia tesla

【解决方案1】：

尝试调用get_matrix(a,b,nx,ny) 两次，获取第二次计时结果。第一次调用 CUDA API 将创建 cuda 上下文。这通常需要很长时间。

关于如何确定块大小和网格大小，请参考 CUDA C 最佳实践指南中的this section。

【讨论】：