1

私はCUDAの初心者です。ここにあるのは、2 つのスレッドで実行されるカーネルです。すべてのスレッドは、結果を共有変数に保存する必要があります。3 つすべてが終了した後、結果はsum12 になるはずですが、6 になりました。

__global__ void kernel (..)
{
    int i=blockDim.x*blockIdx.x+threadIdx.x;

    __shared__ double sum;

        ...

    if(i==0)
        sum=0.0;
    __syncthreads();

    if(i<=1)
        sum+= 2.0*3.0;
    __syncthreads();

    //sum should be 12 here, but I get 6. Why?
}

によって呼び出された

test<<<1,2>>>(..);
4

1 に答える 1

9

コードにメモリ競合があります。これ：

sum+= 2.0*3.0;

複数のスレッドが同時に合計に累積される可能性があります。あなたの例では、両方のスレッドが同時に同じアドレスにロードおよびストアしようとしました。これは CUDA では未定義の動作です。

この問題を回避する通常の方法は、アルゴリズムの再設計です。複数のスレッドが同じメモリ位置に書き込むことは避けてください。メモリ競合なしで共有メモリ配列から合計を累積するために使用できる、非常に広く説明されている共有メモリ削減手法があります。

あるいは、メモリアクセスをシリアル化するために使用できるアトミックメモリアクセスプリミティブがあります。あなたの例は倍精度浮動小数点です。これには、組み込みのアトミック加算関数がないと確信しています。プログラミング ガイドには、倍精度のユーザー空間アトミック加算の例が含まれています。64 ビット共有メモリのアトミック操作は計算機能 2.x および 3.x デバイスでのみサポートされているため、ハードウェアによっては、共有メモリ変数で使用できる場合と使用できない場合があります。いずれにせよ、メモリ アクセスをシリアル化するとパフォーマンスが大幅に低下するため、アトミック メモリ操作は慎重に使用する必要があります。

于 2013-05-28T06:54:59.697 に答える