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

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); 
  }

}

内核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.但是同步的时间有时是10甚至20.看起来CUDA以低成本执行许多操作然后为后来的那些收费很多.任何帮助,将不胜感激.

Answer 1

我认为您对CUDA中的内核调用实际上对主机的影响存在误解.内核调用是非阻塞的,只会添加到设备的队列中.如果你在内核调用之前和之后测量时间,那么差异与内核调用花费的时间无关(它将测量将内核调用添加到队列所花费的时间).

您应该在每次内核调用之后和测量end3time之前添加cudaThreadSynchronize().如果队列中的所有内核都完成了工作,cudaThreadSynchronize()将阻塞并返回.

这就是为什么

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

在你的测量中产生了尖峰.