0

这是我的问题，为了加快我的项目，我想将内核内部生成的值保存到共享内存中，但是，我发现保存该值需要很长时间。如果我删除“THIS LINE”（参见下面的代码），即删除“THIS LINE”，保存该值非常快（加速 100 倍！）。

extern __shared__ int sh_try[];

__global__ void xxxKernel (...)
{
  float v, e0, e1;
  float t;
  int count(0);
  for (...)
  {
     v = fetchTexture();
     e0 = fetchTexture();
     e1 = fetchTexture();
     t = someDeviceFunction(v, e0, e1);
     if (t>0.0 && t < 1.0)  <========== <THIS LINE>
       count++;
  }
  sh_try[threadIdx.x] = count;
}

main()
{
  sth..
  START TIMING:

  xxxKernel<<<gridDim.x, BlockDim.x, BlockDim.x*sizeof(int)>>> (...);

  cudaDeviceSynchronize();

  END TIMING.
  sth...
}

为了解决这个问题，我简化了只将数据保存到共享内存中的代码。并停下来。据我所知，共享内存。是最有效的内存。除了注册，我想知道这种高延迟是正常的还是我做错了什么。请给我一些建议！！！先谢谢各位了！！！

特鲁迪

更新： 如果我用全局内存替换共享内存，它几乎需要相同的时间，没有“这条线”需要 33 毫秒，需要 297 毫秒。将数据保存到全局内存是正常的吗？与共享内存占用相同的时间。？这也是“编译器优化”的一部分吗？

我还检查了stackoverflow上的其他类似问题，即是否将数据保存到共享内存之间存在巨大的时间间隔，这可能是由编译器优化引起的，因为它对计算数据没有意义但不保存它们，所以编译器只是“删除”了那些毫无意义的代码。

我不确定我是否有同样的原因，因为这条线改变了游戏是一个假设 - “这条线”，当我评论它时，变量“计数”在每次迭代中都会增加，当我取消评论时，它会增加t 是有意义的。

有任何想法吗？请...

4

1 回答 1

1

通常，当由于相对较小的代码更改（例如在内核中添加或删除一行代码）导致较大的性能变化时，性能变化并不是由于该行代码的实际性能影响，而是由于由于编译器做出不同的优化决策，这可能导致内核中机器代码的大量添加或删除。

帮助确认这一点的相对简单的方法是查看生成的机器代码。例如，如果生成的机器代码的大小由于添加或删除单行源代码而发生显着变化，则可能是编译器做出了严重影响代码的优化决策。

尽管它不是机器代码，但出于这些目的，一个合理的代理是查看生成的 PTX 代码，这是编译器创建的中间代码。

-ptx您可以通过简单地将开关添加到您的编译命令来生成 ptx ：

nvcc -ptx mycode.cu

这将生成一个名为mycode.ptx您可以检查的文件。当然，如果您的常规编译命令需要额外的开关（例如 -I/path/to/include/files），那么此命令可能需要这些相同的开关。该nvcc 手册提供了更多关于代码生成选项的信息，并且有一个 PTX手册可以帮助您了解 PTX，但是您可能仅根据生成的 PTX 的大小（例如.ptx文件中的行数）。

于 2013-11-16T13:57:02.457 回答