使用 CUDA 转置:在博客中查询程序



我在下面的链接中给出了博客作为矩阵转置的示例,因为下面的链接显示了如何使用 3 种方法进行矩阵转置,朴素、凝聚和 Nobankconflict 凝聚

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

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

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循环呢?即使仅使用其中一个进行了测试,它们都给出了有效的输出,但时间不同,

希望我的问题清楚,如果需要注意,请告诉我,谢谢

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

第一个关键点是函数通常运行得如此之快,以至于您无法从定时函数中获得对其运行时的精确估计:因此,需要在循环中多次运行函数,以便获得更好的精度。

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

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

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

最新更新