CUDA主机到设备传输比设备到主机传输快速。

Question

我正在开发一款简单的cuda程序，其中我发现90%的时间都花在了一个语句上，即从设备到主机的cudamemcpy。该程序在600-700微秒内从主机传输了约2MB的数据，并在10毫秒内从设备复制回了4MB的数据。我的程序总共需要13毫秒的时间。我的问题是，为什么两个复制从设备到主机和从主机到设备之间存在不对称性？这是因为CUDA开发者认为复制回来的数据通常比较小吗？我的第二个问题是，是否有任何方法可以规避这个问题。

我使用的是Fermi GTX560图形卡，具有343个核心和1GB内存。

- Dipendra Kumar Mishra

---

这很可能是时间上的假象，而不是真实情况。内核启动是异步的，因此很有可能设备和主机之间10毫秒的传输时间包括了内核执行时间。 - talonmies

我不这么认为。我正在使用rdtsc，它是一个硬件计数器，并且我已经在cudaMemcpy(...);上方和下方放置了两个计数器标记。此外，为了防止噪音进入系统，我一遍又一遍地重复实验。内核启动是异步的，但我没有使用cudaMemcpyAsync，并且它不能在内核调用结束之前执行。 - Dipendra Kumar Mishra

尝试在设备到主机复制之前放置一个cudaDeviceSynchronize()函数调用。我预测cudaMemcpy函数的测量时间将会大大缩短。 - talonmies

谢谢，它有效地减少了为1毫秒。有趣的是，我曾因速度原因注释掉了cudaDeviceSynchronize :(。再次感谢。 - Dipendra Kumar Mishra

@talonmies 那个方法神奇地起作用了 - 有没有解释为什么会这样？我成功将 memCpy 的时间从 15 毫秒减少到小于 1 毫秒。 - Gokul

Answer 1

CUDA函数的定时与CPU略有不同。首先，请确保您在应用程序开始时调用CUDA函数时不要考虑CUDA初始化成本，否则它可能会在您开始计时时进行初始化。

int main (int argc, char **argv) {
   cudaFree(0);
   ....//cuda is initialized..

}

使用Cutil计时器的方法如下：

unsigned int timer;
cutCreateTimer(&timer);
cutStartTimer(timer);

//your code, to assess elapsed time..

cutStopTimer(timer);
printf("Elapsed: %.3f\n", cutGetTimerValue(timer));
cutDeleteTimer(timer);

现在，在这些预备步骤之后，让我们来看看问题。当调用内核时，CPU 部分仅在调用被传递到 GPU 时会停顿。GPU 将继续执行，而 CPU 也将继续执行。如果您调用 cudaThreadSynchronize(..)，那么 CPU 将停顿，直到 GPU 完成当前调用。cudaMemCopy 操作也需要 GPU 完成其执行，因为请求应由内核填充的值。

kernel<<<numBlocks, threadPerBlock>>>(...);

cudaError_t err = cudaThreadSynchronize();
if (cudaSuccess != err) {
  fprintf(stderr, "cudaCheckError() failed at %s:%i : %s.\n", __FILE__, __LINE__, cudaGetErrorString( err ) );
  exit(1);
}

//now the kernel is complete..
cutStopTimer(timer);

在调用停止定时器函数之前，请先进行同步。如果在内核调用后放置了内存复制，则内存复制的经过时间将包括内核执行的一部分。因此，memCopy操作可以放置在计时操作之后。

还有一些性能分析器计数器可用于评估内核的某些部分。 如何为CUDA内核分析全局内存事务数量？ 如何分析和优化CUDA内核？

- phoad

---

cudaThreadSynchronize()已被弃用，应使用cudaDeviceSynchronize()代替。来源1，来源2 - Ewa