sm_20显示错误的lmem统计数据？

Question

使用--ptxas-options=-v选项编译的CUDA内核似乎在指定 GPU体系结构时显示错误的sm_20 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报告了一个lmem的4字节，如果您看到内核中使用的4x1000字节数组，这是完全不可能的。旧的GPU体系结构报告正确的4000字节 lmem统计量。

这是在CUDA 3.2上进行的。我已经提到了打印代码生成统计信息部分( NVCC手动 (v3.2) )，但它无助于解释这种异常。

Answer 1

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

查看生成的ptx代码不会显示sm_10与sm_20之间的任何区别，用decuda对生成的cubins进行反编译将显示优化。

顺便说一句:尽量避免本地记忆！它和全球记忆一样慢。