for-loop - 在 for 循环中重复调用内核的 CUDA 程序的性能受到影响

Question

我有一个 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); 
  }

}

内核分配5很简单

__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。但是接下来的时间forsync 有时是 10 甚至 20。看起来 CUDA 正在执行许多操作成本低，然后对以后的产品收取高额费用。任何帮助将不胜感激。

Answer 1

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

您应该在每个内核调用之后以及测量 end3time 之前添加 cudaThreadSynchronize()。如果队列中的所有内核都完成了工作，cudaThreadSynchronize() 会阻塞并返回。

这就是原因

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

您的测量值出现峰值。