c++ - CUDA:--ptxas-options=-v 共享内存和 cudaFuncAttributes.sharedSizeBytes 不匹配

Question

已关闭。此问题需要 debugging details 。目前不接受答案。

---

编辑问题以包含 desired behavior, a specific problem or error, and the shortest code necessary to reproduce the problem 。这将有助于其他人回答问题。

已关闭 8 年前。

Improve this question

我正在尝试使用reqRegs和sharedSizeBytes在 cudaFuncAttributes struct 在运行时动态优化内核的 block 大小。

我当前的实现略读 nvcc --ptxas-options=-v 中的标准输出文本发现内核的寄存器和共享内存使用情况。此方法有点 hacky，并且依赖于 --ptxas-options=-v 输出文本的确切格式。 ，这可能会在没有警告的情况下发生变化。

我的问题是，我发现 --ptxas-options=-v 中报告的“smem”共享内存值之间存在差异。输出，和sharedSizeBytes在 cudaFuncAttributes struct，这让我担心我到目前为止一直使用的共享内存估计是错误的，或者 sharedSizeBytes变量不可靠，这意味着我不能将它用于运行时 block 大小优化。这是 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。

谢谢。

Answer 1

感谢罗伯特和马可的回复。他们帮助我排除了一些情况。

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

共享内存差异是由于共享内存数组分配依赖于 block 大小造成的；例如:

__shared__ float myArray[BLOCK_SIZE];

上面的语句在 block 大小为 256 的程序中使用的共享内存量与使用优化 block 大小 192 编译的相同源代码不同。现在看来很明显，但在使用时需要注意一些事情优化的 CUDA 代码生成。