c++ - 使用 CUDA 转置 : Inquiry on program in Blog

Question

我在下面的博客链接中给出了矩阵转置的示例，下面的链接显示了如何使用 3 种方法进行转置矩阵，naive、Coalesced 和 Nobankconflict Coalesced

https://github.com/parallel-forall/code-samples/blob/master/series/cuda-cpp/transpose/transpose.cu

在 Main() 中，当调用内核代码时，所有 3 个方法都以类似的方式调用它，如下代码部分(取自 main 函数主机):

cudaMemset(d_tdata, 0, mem_size);
// warmup
transposeNoBankConflicts << <dimGrid, dimBlock >> >(d_tdata, d_idata);
cudaEventRecord(startEvent, 0);
for (int i = 0; i < NUM_REPS; i++)
    transposeNoBankConflicts << <dimGrid, dimBlock >> >(d_tdata, d_idata);
cudaEventRecord(stopEvent, 0);
cudaEventSynchronize(stopEvent);
cudaEventElapsedTime(&ms, startEvent, stopEvent);
cudaMemcpy(h_tdata, d_tdata, mem_size, cudaMemcpyDeviceToHost);

我什至看了网上的代码解释，它是矩阵转置的一个很好的引用

http://devblogs.nvidia.com/parallelforall/efficient-matrix-transpose-cuda-cc/

但是有一部分没有解释:

为什么如您所见，内核被调用了两次:1\一次，正如评论所说的热身2\秒下for循环最多100次(NUM_REPS)这是一个初始化为100的#define值，

那么为什么不调用一次呢？为什么两次和第二次 100 次循环？尽管我只使用其中一个进行了测试，但它们都提供了有效输出，但时间不同，

希望我的问题很清楚，如果有什么需要注意的地方请告诉我，谢谢

Answer 1

这与矩阵转置无关:这些是准确计时代码块的一些基础知识。

第一个关键点是函数通常运行得如此之快，以至于您无法从计时函数中准确估计它们的运行时间:因此，需要在循环中多次运行该函数，以便获得更好的精度.

(而且你必须注意你实际上是在为你想要的时间计时；有时优化器很聪明，你尝试的最简单的事情实际上不会为你想要的时间计时；例如，优化器可能会找到一种有效混合在一起的方法代码的结尾和开头，或者它可能会注意到您没有使用循环中前 99 次的输出，因此它不会打扰运行它们。使用 nvcc< 可能不会受到这种影响，因为它几乎肯定会将内核编译成一个不透明的函数调用)

第二个关键点是第一次迭代通常会比后面的迭代慢，原因有很多，下面列出了一些。因此，为了获得准确的计时，您通常不希望在计时中包含第一次迭代。

• 也许您的数据通常在缓存中，但第一次迭代时还没有，所以第一次会出现很多缓存未命中。 (相反，如果您的代码预期运行时数据不在缓存中并且已经在缓存中会影响计时，您应该在迭代之间做一些事情来污染缓存)
• 在 CPU 上，第一次访问内存区域会产生页面错误，这可能会非常慢。我不确定这种效果是否会出现在 GPU 上
• GPU 上未充分利用的计算单元可能会切换到省电模式。我不熟悉细节，但关键是第一次(或者可能是前几次)GPU 可能被降频，或者许多功能单元可能被关闭，并且“预热”GPU 将恢复它的全部能力。