在 for 循环中重复调用内核的 CUDA 程序中的性能下降答案

【问题标题】：Performance hit in CUDA program that calls kernel repeatedly within a for loop在 for 循环中重复调用内核的 CUDA 程序中的性能下降
【发布时间】：2013-02-27 00:46:48
【问题描述】：

我有一个 CUDA 程序，它在 for 循环中重复调用内核。这代码使用前一个中计算的值计算矩阵的所有行直到整个矩阵完成。这基本上是一种动态规划算法。下面的代码并行填充许多单独矩阵的 (i,j) 条目内核。

for(i = 1; i <=xdim; i++){

  for(j = 1; j <= ydim; j++){ 

    start3time = clock();
    assign5<<<BLOCKS, THREADS>>>(Z, i, j, x, y, z)
    end3time = clock(); 
    diff = static_cast<double>(end3time-start3time)/(CLOCKS_PER_SEC / 1000); 
    printf("Time for i=%d j=%d is %f\n", i, j, diff); 
  }

}

内核assign5很简单

__global__ void assign5(float* Z, int i, int j, int x, int y, int z) {

  int id = threadIdx.x + blockIdx.x * blockDim.x;

  char ch = database[j + id];

  Z[i+id] = (Z[x+id] + Z[y+id] + Z[z+id])*dev_matrix[i][index[ch - 'A']];

  }

}

我的问题是，当我运行这个程序时，每个 i 和 j 的时间大部分都是 0 时间，但有时是 10 毫秒。所以输出看起来像

Time for i=0 j=0 is 0
Time for i=0 j=1 is 0
.
.
Time for i=15 j=21 is 10
Time for i=15 j=22 is 0
.

我不明白为什么会这样。我没有看到线程竞争条件。如果我添加

if(i % 20 == 0) cudaThreadSynchronize();

在第一个循环之后，i 和 j 的时间大多为 0。但是时间 for sync 有时是 10 甚至 20。看起来 CUDA 正在执行许多操作以低成本，然后为以后的费用收取很多费用。任何帮助将不胜感激。

【问题讨论】：

肯定只是因为时间源的精度太低而造成混叠？
我忘了补充变量 Z 和数据库是设备全局数组和 dev_matrix 以及索引和设备常量数组。内存访问被合并。
结果不一致，因为您的时钟分辨率约为 10 毫秒，就像 talonmies 所说的那样。最好只测量整个循环的时间。如果您需要高分辨率的每次迭代测量，那么this answer 应该会有所帮助。
我现在使用上面指示的计时器以微秒为单位测量了时间。每次早期迭代（所有 i
不，它在 Linux 上。我相信我可能错误地测量了运行时间。另一个论坛上的某个人说我应该使用 cudaEventRecord 和 cudaEventSynchronize。我现在得到了一致的结果。

标签： for-loop cuda gpu dynamic-programming

【解决方案1】：

我认为您对 CUDA 中的内核调用在主机上的实际作用存在误解。内核调用是非阻塞的，并且只添加到设备的队列中。如果您在内核调用之前和之后测量时间，那么差异与内核调用花费的时间无关（它将测量将内核调用添加到队列中所花费的时间）。

您应该在每个内核调用之后和测量 end3time 之前添加一个 cudaThreadSynchronize()。 cudaThreadSynchronize() 阻塞并在队列中的所有内核都完成工作时返回。

这就是为什么

if(i % 20 == 0) cudaThreadSynchronize();

在您的测量结果中出现峰值。

【讨论】：

谢谢，现在说得通了。此外，cudaMemcpy 似乎会阻塞，直到所有线程都完成。
是的，调用主机上的 cudaMemcpy 块，直到所有线程都完成。这是标准的 memcopy 功能，因为这是您大部分时间想要的。还有 cudaMemcpyAsync()，它是非阻塞的变种。
作为未来读者的参考：cudaThreadSynchronize 已被弃用，而应使用 cudaDeviceSynchronize。