为什么以下程序的相同袖口代码需要不同的时间？

Question

我在 cufft (cuda 9) (Nvidia 1080) 中运行了以下代码。所有执行的代码都是相同的。但是，执行时间（在代码下方）变化很大。任何人都可以描述如何始终获得最低时间以及这种行为背后的原因吗？

int NX 2048
int BATCH 96

cufftHandle plan;
cufftHandle rev_plan;
cufftDoubleReal *idata;
cufftDoubleComplex *odata;

int BLOCKSIZE  = 1024;
int gridSize = (NX * BATCH)/BLOCKSIZE;

cufftPlan1d(&plan, NX, CUFFT_D2Z, BATCH);
cufftPlan1d(&rev_plan, NX, CUFFT_Z2D, BATCH);


cudaMalloc((void **) &idata, sizeof(cufftDoubleReal) * NX * BATCH);
cudaMalloc((void **) &odata, sizeof(cufftDoubleComplex) * (NX / 2 + 1) * BATCH);
inputData << < gridSize, BLOCKSIZE >> > (idata, NX * BATCH);

double sT = omp_get_wtime();
for (int i = 0; i < 500; ++i) {
    cufftExecD2Z(plan, idata, odata);
    cufftExecZ2D(plan, odata, idata);
}
printf("Time taken: %f\n", omp_get_wtime() - sT);

sT = omp_get_wtime();
for (int i = 0; i < 500; ++i) {
    cufftExecD2Z(plan, idata, odata);
    cufftExecZ2D(plan, odata, idata);
}
printf("Time taken: %f\n", omp_get_wtime() - sT);

sT = omp_get_wtime();
for (int i = 0; i < 500; ++i) {
    cufftExecD2Z(plan, idata, odata);
    cufftExecZ2D(plan, odata, idata);
}
printf("Time taken: %f\n", omp_get_wtime() - sT);

sT = omp_get_wtime();
for (int i = 0; i < 500; ++i) {
    cufftExecD2Z(plan, idata, odata);
    cufftExecZ2D(plan, odata, idata);
}
printf("Time taken: %f\n", omp_get_wtime() - sT);

cudaFree(idata);
cudaFree(odata);

所用时间：0.004334 所用时间：0.022906 所用时间：0.027820 耗时：0.027786

Answer 1

对 cufft 例程的调用可以是异步的

这意味着调用可能在工作完成之前返回。

这只能在一定限度内发生。有一个异步启动队列。填满队列后，队列中的新插槽仅在调度队列项时打开。这意味着启动过程不再是异步的。

这会扭曲您的计时结果。

要“修复”此问题，请在每个计时区域结束之前添加一个 cudaDeviceSynchronize(); 调用（即紧接在每个 printf 语句之前）。这将大大平衡结果。这会强制所有 GPU 工作在您完成计时测量之前完成。

$ cat t37.cu
#include <cufft.h>
#include <omp.h>
#include <cuda_runtime_api.h>
#include <cstdio>

int main(){

  const int NX = 2048;
  const int BATCH = 96;

  cufftHandle plan;
  cufftHandle rev_plan;
  cufftDoubleReal *idata;
  cufftDoubleComplex *odata;

  //int BLOCKSIZE  = 1024;
  //int gridSize = (NX * BATCH)/BLOCKSIZE;

  cufftPlan1d(&plan, NX, CUFFT_D2Z, BATCH);
  cufftPlan1d(&rev_plan, NX, CUFFT_Z2D, BATCH);


  cudaMalloc((void **) &idata, sizeof(cufftDoubleReal) * NX * BATCH);
  cudaMalloc((void **) &odata, sizeof(cufftDoubleComplex) * (NX / 2 + 1) * BATCH);
  //inputData << < gridSize, BLOCKSIZE >> > (idata, NX * BATCH);

  double sT = omp_get_wtime();
  for (int i = 0; i < 500; ++i) {
            cufftExecD2Z(plan, idata, odata);
            cufftExecZ2D(plan, odata, idata);
  }
  #ifdef FIX
  cudaDeviceSynchronize();
  #endif
  printf("Time taken: %f\n", omp_get_wtime() - sT);

  sT = omp_get_wtime();
  for (int i = 0; i < 500; ++i) {
            cufftExecD2Z(plan, idata, odata);
            cufftExecZ2D(plan, odata, idata);
  }
  #ifdef FIX
  cudaDeviceSynchronize();
  #endif
  printf("Time taken: %f\n", omp_get_wtime() - sT);

  sT = omp_get_wtime();
  for (int i = 0; i < 500; ++i) {
            cufftExecD2Z(plan, idata, odata);
            cufftExecZ2D(plan, odata, idata);
  }
  #ifdef FIX
  cudaDeviceSynchronize();
  #endif
  printf("Time taken: %f\n", omp_get_wtime() - sT);

  sT = omp_get_wtime();
  for (int i = 0; i < 500; ++i) {
            cufftExecD2Z(plan, idata, odata);
            cufftExecZ2D(plan, odata, idata);
  }
  #ifdef FIX
  cudaDeviceSynchronize();
  #endif
  printf("Time taken: %f\n", omp_get_wtime() - sT);

  cudaFree(idata);
  cudaFree(odata);
}
$ nvcc -o t37 t37.cu -lcufft -lgomp
$ ./t37
Time taken: 0.007373
Time taken: 0.185308
Time taken: 0.196998
Time taken: 0.196857
$ nvcc -o t37 t37.cu -lcufft -lgomp -DFIX
$ ./t37
Time taken: 0.197076
Time taken: 0.196994
Time taken: 0.196937
Time taken: 0.196916
$

有人可能会问，“为什么没有调用cudaDeviceSynchronize() 的总时间明显低于使用它的总时间？”这基本上是由于相同的原因。异步启动队列充满了待处理的工作，但程序在队列中的所有工作启动之前终止（没有最终的cudaDeviceSynchronize()）。在每种情况下，这都会导致总执行时间之间的明显差异。通过只添加最后一个cudaDeviceSynchronize() 调用，可以观察到这种效果。