1

大量の共有メモリを必要とする N 体問題に取り組んでいます。

基本的に、N独立したタスクがあり、それぞれが 4 つの double 変数、つまり 32 バイトを使用します。そして、1 つのタスクがスレッドによって実行されます。

迅速にするために、これらの変数には共有メモリを使用しています (レジスタはスレッドによっても使用されているため)。Nコンパイル時にはタスクの数がわからないため、共有メモリは動的に割り当てられます。

  • グリッドの次元と共有メモリNは、ブロック サイズに応じて計算されます。

    const size_t BLOCK_SIZE = 512;
const size_t GRID_SIZE = (N % BLOCK_SIZE) ? (int) N/BLOCK_SIZE : (int) N/BLOCK_SIZE +1;
const size_t SHARED_MEM_SIZE = BLOCK_SIZE * 4 * sizeof(double);

  • 次に、これら 3 つの変数を使用してカーネルが起動されます。

    kernel_function<<<GRID_SIZE, BLOCK_SIZE, SHARED_MEM_SIZE>>>(N, ...);

smallNの場合、これは正常に機能し、カーネルはエラーなしで実行されます。

しかし、exceed の場合N = 1500、カーネルの起動は失敗します (次のメッセージが複数回表示されます)。

========= Invalid __global__ write of size 8
=========
========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaLaunch.

私が理解している限り、これは割り当てられた共有メモリの境界外に書き込もうとしたためです。これは、カーネルで、グローバル メモリが共有メモリにコピーされている場合に発生します。

__global__ void kernel_function(const size_t N, double *pN, ...)
{
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;

    if(idx<N)
    {
        extern __shared__ double pN_shared[];
        for(int i=0; i < 4; i++)
        {
            pN_shared[4*idx + i] = pN[4*idx + i];
        }
        ...
    }
}

このエラーは、共有メモリN > 1500の全体量が48kBを超えた場合にのみ発生します( )。 この制限は、グリッドとブロック サイズに関係なく同じです。1500 * 4 * sizeof(double) = 1500 * 32 = 48000

CUDA の仕組みを正しく理解していれば、グリッドが使用する共有メモリの累積量は48kBに制限されません。これは、単一のスレッド ブロックで使用できる共有メモリの制限にすぎません。

共有メモリの累積量は、ストリーミング マルチプロセッサ間でグリッドがスケジュールされる方法にのみ影響するはずなので、このエラーは私には意味がありません (さらに、GPU デバイスには 15 個の SM が自由に使用できます)。

4

3 に答える 3

1

インデックス idx*4+0:3 で共有配列にアクセスしています。N > BLOCK_SIZE から始まるプログラムは正しくありません。幸いなことに、1500 までは動作するようです。ただし、cuda mem-check を使用すると、問題が指摘されるはずです。関連トピックでは、別の場所に静的に割り当てられた共有メモリが共有メモリを使用する可能性があることに注意してください。ポインターの値を出力すると、理解するのに役立ちます。

于 2016-05-20T15:59:25.603 に答える