1

オプションでコンパイルされた CUDA カーネルは、 GPU アーキテクチャが指定されている場合、--ptxas-options=-v誤った lmem (ローカル メモリ)統計を表示しているようです。sm_20同じことで、アーキテクチャに関する意味のある lmem 統計が得られsm_10 / sm_11 / sm_12 / sm_13ます。

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) の「コード生成統計の印刷」セクションを参照しましたが、この異常の説明には役立ちません。

4

1 に答える 1

1

コンパイラは正しいです。巧妙な最適化により、配列を保存する必要はありません。あなたがしていることはresult += i * i、一時的なものを に保存することなく、基本的に計算することvalです。

生成された ptx コードを見ると、sm_10 と sm_20 の違いはわかりません。生成された cubin を decuda で逆コンパイルすると、最適化が明らかになります。

ところで：ローカルメモリを避けるようにしてください！グローバルメモリと同じくらい遅いです。

于 2011-02-24T12:39:28.293 に答える