0

次のカーネルがあります。 

__global__ void myKernel(int k, int inc, int width, int* d_Xco, int* d_Xnum, bool* 
        Xvalid, float* d_X)
    {

        int i, k1;  
        i = threadIdx.x + blockIdx.x * blockDim.x;
        //k1 = threadIdx.y + blockIdx.y * blockDim.y;

        if( (i < k)  ){
           for(k1 = 0; k1 < inc; k1++){

             int mul = (d_X[i*inc + k1] >= 2e2);
             d_X[i*inc + k1] *= (float)(!mul);
             d_Xco[i*width + k1] = k*mul;
             d_Xnum[i] += mul;
             d_Xvalid[i*inc + k1] = (!mul) ; 

            }
         }// of if

 }

これは次のように呼び出されます:

  int bx = (int)(k/32)+1;
  int by = (int)(inc/32)+1;

  dim3 b(bDim, 1);
  dim3 t(tDim, 1);
  cmyKernel<< b, t >>>( k, inc, width, d_Xco, d_Xnum, d_Xvalid, d_X );

  cudaThreadSynchronize();

kは約 9000 で、inc約 5000 ですので、ブロック数を超えていないと確信しています。myKernelが次元で 1thread/1block で呼び出された場合y、カーネルは正常に動作しているように見えますが、たとえば、次元のスレッドとブロックの数yを 10 に変更すると、カーネル内で実際に使用していない場合でも、間違った出力が得られます。のスレッドとブロックy。for()理想的には、私は使用を取り除きたいですk = threadIdx.y + blockIdx.y * blockDim.y

4

2 に答える 2

2

y次元= 10でカーネルを起動すると、それらを使用しています。スレッド識別子を使用していないという理由だけthreadIdx.yでblockIdx.y、スレッドが起動されていないという意味ではありません。y 次元 = 10 でカーネルを起動すると、i = 0 で 10 個のスレッド、i = 1 で 10 個のスレッドなどになります。

簡単にするために 2x2 スレッドを起動するとします。スレッドが (0,0)(0,1) (1,0) (1,1) になります。あなたのコードでは、2 つのスレッド (0,0) と (0,1) の i 変数は 0 ですが、threadIdx.y違います。これは、両方のスレッドが同じ i 変数のコードを評価し、競合状態が発生することを意味します。

反復間の依存関係を解決する必要があります (d_Xnum[i] += mul)。これを行う 1 つの方法は、atomicAdd(..). のコメントを外し、k1ループを に置き換えてif(k1 < inc)、atomicAdd. これにより、正しい動作が得られるはずです。

于 2012-10-02T14:52:30.007 に答える
1

コメントで既に述べたように、現在のソリューションは、それぞれの作業をメモリの同じスペースに適用する複数のスレッドを起動しています。これは、すべて同じ値の threadIdx.x を持ち、threadIdx.y の値が異なる複数のスレッドが生成されるためです。これは、複数のスレッドが同時に同じメモリ空間を読み書きすることを意味します。これには多くの潜在的な問題があります。ここに簡単な説明があります。

これを回避するには、いくつかの手順を実行できます。たとえば、同期データ アクセスを使用できます (これにより、他のスレッドがデータ アクセスを完了するのをスレッドが待機するため、大幅な速度低下が発生します)。各スレッドで 1 つのセル要素を処理する場合は、for ループを削除し、代わりに以前のように k1 を使用する必要がありますが、メモリの読み取りと書き込みを慎重に考慮する必要があります。他は別スレで！

ここでの核となる理解は、スレッド間の操作のシーケンスに依存することは決してできないということです!

データ アクセスに関しては、すべてのデータ構造をグリッドとして考えると役立ちます。各スレッドは、独自の座標内のデータにのみアクセスして変更する必要があります。たとえば、threadIdx.x == 3 のスレッドの場合は (3,2) です。および threadIdx.y == 2. このようにして、スレッドの動作と潜在的な競合状態を簡単に視覚化できます。これを使用する最も簡単な方法は、出力データの要素ごとに 1 つのグリッド エントリを作成することです。したがって、9000x5000 要素のマトリックスがある場合、最初にその量のスレッドを生成し、そこから最適化することができます。もちろん、これにより GPU はすべてのユニットで数回実行する必要がありますが、これは良い出発点です。

オスロ大学には、とりわけこのトピックに関する大学院レベルのコースがあります。これらのスライド は、理解を深めるために非常に関連性が高いと思われるかもしれません。特に、スレッドのバッチ処理、グリッド、およびブロックに関するセクションを参照してください。

于 2012-10-02T13:16:00.000 に答える