CUDA：--ptxas-options=-v 共享内存和 cudaFuncAttributes.sharedSizeBytes 不匹配 [关闭]答案

【问题标题】：CUDA: --ptxas-options=-v shared memory and cudaFuncAttributes.sharedSizeBytes do not match [closed]CUDA：--ptxas-options=-v 共享内存和 cudaFuncAttributes.sharedSizeBytes 不匹配 [关闭]
【发布时间】：2014-09-02 11:47:09
【问题描述】：

我试图在cudaFuncAttributes 结构中使用reqRegs 和sharedSizeBytes 在运行时动态优化内核的块大小。

我当前的实现从nvcc --ptxas-options=-v 中浏览标准输出文本以发现内核的寄存器和共享内存使用情况。这种方法有点笨拙，并且取决于来自--ptxas-options=-v 的输出文本的确切格式，可能会在没有警告的情况下更改。

我的问题是我看到--ptxas-options=-v 输出中报告的“smem”共享内存值与cudaFuncAttributes 结构中的sharedSizeBytes 之间存在差异，这让我担心共享内存估计我一直使用到现在是错误的，或者 sharedSizeBytes 变量不可靠，这意味着我不能将它用于运行时块大小优化。这是nvcc --ptxas-options=-v 对一个这样的内核的输出...

ptxas info    : Used 14 registers, 2088 bytes smem, 48 bytes cmem[1]

... 与运行时 cudaFuncAttributes.sharedSizeBytes = 296 的值相比，对于完全相同的内核。有人知道这里会发生什么吗？

这是另一个使用不同内核的示例：

ptxas info    : Used 18 registers, 2132 bytes smem, 48 bytes cmem[1]

其中cudaFuncAttributes.sharedSizeBytes = 340 在运行时。

谢谢。

【问题讨论】：

一些优化可能使用了未使用的共享内存，你编译的是哪个cc？
我写了a simple test，但无法观察到这种差异。您能否提供一个完整的测试用例，希望像我的一样简单，证明这种差异？

标签： c++ c cuda ptxas

【解决方案1】：

感谢 Robert 和 Marco 的回复。他们帮助我排除了一些情况。

事实证明，报告的共享内存使用量不匹配是由于第一次测试编译后使用的共享内存量（由--ptxas-options=-v 报告）与修改后的最终程序使用的共享内存量不同块大小（由cudaFuncAttributes.sharedSizeBytes 报告）。（为清楚起见进行编辑）

共享内存的差异是由共享内存数组分配依赖于块大小引起的；例如：

__shared__ float myArray[BLOCK_SIZE];

上面的语句在块大小为 256 的程序中使用不同数量的共享内存，而不是使用优化块大小为 192 编译的相同源代码。现在看起来很明显，但需要注意的是优化 CUDA 代码生成。

【讨论】：

关于您的答案，您显示的一行代码是静态分配，这意味着必须在编译时知道大小（BLOCK_SIZE）。编译时输出和运行时输出之间也不会有差异。（您还没有展示一个示例来证明差异。没有。）您可能遇到的唯一差异是扫描内核的错误（即旧的、陈旧的、以前的）编译时输出问题。
这是我的意思，先生。对nvcc 的第一次调用使用的BLOCK_SIZE 值与第二次nvcc 调用使用的值不同。因此，第一个程序使用的共享内存量与第二个程序使用的共享内存量不同。