从GPU到CPU的复制速度比从CPU到GPU慢。

Question

我已经开始学习CUDA一段时间了，我遇到了以下问题：

请看下面的操作步骤： 复制GPU

int* B;
// ...
int *dev_B;    
//initialize B=0

cudaMalloc((void**)&dev_B, Nel*Nface*sizeof(int));
cudaMemcpy(dev_B, B, Nel*Nface*sizeof(int),cudaMemcpyHostToDevice);
//...

//Execute on GPU the following function which is supposed to fill in 
//the dev_B matrix with integers


findNeiborElem <<< Nblocks, Nthreads >>>(dev_B, dev_MSH, dev_Nel, dev_Npel, dev_Nface, dev_FC);

复制CPU

cudaMemcpy(B, dev_B, Nel*Nface*sizeof(int),cudaMemcpyDeviceToHost);

1. Copying array B to dev_B takes only a fraction of a second. However copying array dev_B back to B takes forever.

2. The findNeiborElem function involves a loop for each thread e.g. it looks like that

   __ global __ void findNeiborElem(int *dev_B, int *dev_MSH, int *dev_Nel, int *dev_Npel, int *dev_Nface, int *dev_FC){
   
       int tid=threadIdx.x + blockIdx.x * blockDim.x;
       while (tid<dev_Nel[0]){
           for (int j=1;j<=Nel;j++){
                // do some calculations
                B[ind(tid,1,Nel)]=j// j in most cases do no go all the way to the Nel reach
                break; 
           }
       tid += blockDim.x * gridDim.x; 
       }
   }
   

非常奇怪的是，将dev_B复制到B所需的时间与j索引的迭代次数成正比。

例如，如果Nel=5，则时间约为5秒。

当我增加Nel = 20时，时间大约为20秒。

我本来期望复制时间应该独立于内部迭代次数，只需要分配Matrix dev_B的值即可。

同时，我也期望从CPU复制相同的矩阵的时间与之相同。

您有任何想法出了什么问题吗？

- giorgk

Answer 1

与其使用clock()来测量时间，你应该使用事件：

使用事件，你可以像这样编写代码：

  cudaEvent_t start, stop;   // variables that holds 2 events 
  float time;                // Variable that will hold the time
  cudaEventCreate(&start);   // creating the event 1
  cudaEventCreate(&stop);    // creating the event 2
  cudaEventRecord(start, 0); // start measuring  the time

  // What you want to measure
  cudaMalloc((void**)&dev_B, Nel*Nface*sizeof(int));
  cudaMemcpy(dev_B, B, Nel*Nface*sizeof(int),cudaMemcpyHostToDevice);

  cudaEventRecord(stop, 0);                  // Stop time measuring
  cudaEventSynchronize(stop);               // Wait until the completion of all device 
                                            // work preceding the most recent call to cudaEventRecord()

  cudaEventElapsedTime(&time, start, stop); // Saving the time measured

编辑：附加信息：

“在内核启动完成之前，内核启动就将控制权交还给CPU线程。因此，您的计时构造既测量了内核执行时间，也测量了第二个memcpy。当在内核之后计时复制时，您的计时器代码立即被执行，但是cudaMemcpy在开始之前会等待内核完成。这也解释了为什么您对数据返回的计时测量似乎会根据内核循环迭代而变化。它也解释了为什么您的内核函数所花费的时间“可以忽略不计”。”由Robert Crovella提供。

- dreamcrash

---

你应该使用这种方法。在内核启动之后，控制权会在内核完成之前返回给CPU线程。因此，你的时间构造函数测量了内核执行时间以及第二个memcpy。当在内核之后计时复制时，你的计时器代码会立即执行，但是cudaMemcpy需要等待内核完成才能开始执行。这也解释了为什么你对数据返回的计时测量似乎会根据内核循环迭代而有所不同。它还解释了为什么你的内核函数所花费的时间“可以忽略不计”。 - Robert Crovella

Answer 2

关于你的第二个问题

 B[ind(tid,1,Nel)]=j// j in most cases do no go all the way to the Nel reach

在GPU上进行计算时，由于同步原因，每个完成工作的线程都不会执行任何计算，直到同一工作组中的所有线程都完成为止。

换句话说，执行此计算所需的时间将是最坏情况下的时间，即使大多数线程没有完全执行也无关紧要。

我不确定您的第一个问题，如何测量时间？我对cuda不太熟悉，但我认为在从CPU复制到GPU时，实现会缓冲您的数据，隐藏涉及的有效时间。

- sbabbi

---

感谢您的回答，但是让我澄清一下。我纯粹是指在复制过程中花费的时间。在函数“findNeiborElem”上花费的时间实际上是可忽略不计的。为了测量时间，我使用了以下代码：start = std::clock();cudaMemcpy(B, dev_B, NelNfacesizeof(int),cudaMemcpyDeviceToHost);duration = ( std::clock() - start ) / (double) CLOCKS_PER_SEC;std::cout<<"从GPU复制数据所需时间： " << duration << "秒\n"; - giorgk

这取决于你想要测量哪个时间点，记住大多数操作都是异步执行的。换句话说，除非过程发出同步点（这是cudaMemcpyDeviceToHost的情况），否则GPU还没有完成任何或仅完成了少量工作。在这里https://dev59.com/gE_Ta4cB1Zd3GeqPFf3s有更好的解释。 - sbabbi

直到CUDA出现之前，数据路径GPU->CPU的性能并不重要，因此它可能没有经过彻底优化（甚至AGP也做出了传输速度应该是非对称的有意决定）。 - Simon Richter