CUDA，在共享内存中存储数据时性能低下答案

【问题标题】：CUDA, low performance in storing data in shared memroyCUDA，在共享内存中存储数据时性能低下
【发布时间】：2013-03-27 13:03:00
【问题描述】：

这是我的问题，为了加快我的项目，我想将内核内部生成的值保存到共享内存中，但是，我发现保存该值需要很长时间。如果我删除“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...
}

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

特鲁迪

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

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

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

有什么想法吗？请...

【问题讨论】：

当您在更改单行代码时看到这样的巨大速度差异时，很可能是因为编译器能够优化出一大块代码。由于您的内核仅将数据存储在共享内存中，因此它没有做任何有用的事情。编译器可以检测到这一点，并从本质上将其替换为空内核。您可以通过使用nvcc -ptx mycode.cu 查看这两种情况下的代码输出来看到差异。
使用“@name”通知评论者。 ptx 文件以某种方式可读。要检查的主要内容是函数的主体。它应该以.entry _Z6xxxKernelILi2EEvPj() { 开头。之后，类似汇编器的代码中的主体如下。
@stuhlo，感谢您的回复。也许我应该首先弄清楚如何使用 nvcc -ptx 进行编译。我收到一个错误“找不到 cutil_inline.h”有什么想法吗？
@Robert Crovella，感谢您的回复。对不起，我是 CUDA 的新手，请问如何使用 nvcc-ptx mycode.cu 检查代码输出？
是的，即使使用全局操作，添加或删除 THIS LINE 也可以让编译器摆脱代码片段。例如，没有这行代码，就不需要调用someDeviceFunction，因为计算值t 对代码的行为没有影响。因此编译器可以优化该调用。比较ptx，只要使用有行和不行的代码生成ptx，比较文件长度或指令总数的差异。

标签： cuda shared-memory

【解决方案1】：

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

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

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

您只需将-ptx 开关添加到您的编译命令即可生成ptx：

nvcc -ptx mycode.cu

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

【讨论】：