CUDA：为 sm_20 显示错误的 lmem 统计信息？

Question

当指定sm_20 GPU 架构时，使用选项 --ptxas-options=-v 编译的 CUDA 内核似乎显示错误的 lmem（本地内存） 统计信息。对于 sm_10 / sm_11 / sm_12 / sm_13 架构，同样可以提供有意义的 lmem 统计信息。

有人能澄清一下 sm_20 lmem 统计数据是否需要以不同的方式阅读，或者它们是完全错误的吗？

这是内核：

__global__ void fooKernel( int* dResult )
{
        const int num = 1000;
        int val[num]; 

        for ( int i = 0; i < num; ++i )
        val[i] = i * i; 

        int result = 0; 

        for ( int i = 0; i < num; ++i )
        result += val[i]; 

        *dResult = result;

        return;
}

--ptxas-options=-v 和sm_20 举报：

1>ptxas info    : Compiling entry function '_Z9fooKernelPi' for 'sm_20'
1>ptxas info    : Used 5 registers, 4+0 bytes lmem, 36 bytes cmem[0]

--ptxas-options=-v 和sm_10 / sm_11 / sm_12 / sm_13 举报：

1>ptxas info    : Compiling entry function '_Z9fooKernelPi' for 'sm_10'
1>ptxas info    : Used 3 registers, 4000+0 bytes lmem, 4+16 bytes smem, 4 bytes cmem[1]

sm_20 报告 4 字节 的 lmem，如果您看到内核中使用了 4x1000 字节数组，这根本不可能。较旧的 GPU 架构报告正确的 4000 字节 lmem 统计信息。

已在 CUDA 3.2 上进行了尝试。我参考了 NVCC 手册 (v3.2) 的 Printing Code Generation Statistics 部分，但它无助于解释这种异常情况。

Answer 1

编译器是正确的。通过巧妙的优化，数组不需要存储。您所做的实际上是计算result += i * i，而不会将临时数据存储到val。

查看生成的 ptx 代码不会显示 sm_10 与 sm_20 的任何差异。使用 decuda 反编译生成的 cubin 将显示优化。

顺便说一句：尽量避免使用本地内存！它和全局内存一样慢。