1

私は自分のゲーム プロジェクト (タワー ディフェンス) に取り組んでおり、共有メモリを使用して JCuda ですべてのクレーターとタワーの間の距離を計算しようとしています。タワーごとに、N スレッドで 1 ブロックを実行します。ここで、N はマップ上の生き物の数に等しくなります。特定のブロックのすべてのクレーターとその塔の間の距離を計算し、これまでに見つかった最小の距離をブロックの共有メモリに保存します。私の現在のコードは次のようになります。

extern "C"

__global__ void calcDistance(int** globalInputData, int size, int
critters, int** globalQueryData, int* globalOutputData) {

 //shared memory
 __shared__ float minimum[2];

 int x = threadIdx.x  + blockIdx.x * blockDim.x;
 int y = blockIdx.y;

 if (x < critters) {

   int distance = 0;
   //Calculate the distance between tower and criter
   for (int i = 0; i < size; i++) {
     int d = globalInputData[x][i] - globalQueryData[y][i];
     distance += d * d;
   }

   if (x == 0) {        
     minimum[0] = distance;
     minimum[1] = x;
   }

   __syncthreads();



   if (distance < minimum[0]) {
     minimum[0] = distance;
     minimum[1] = x;
   }

   __syncthreads();
   globalOutputData[y * 2]     = minimum[0];
   globalOutputData[y] = minimum[1];

 }


}

問題は、同じ入力を使用してコードを複数回再実行すると(各実行後にホストとデバイスの両方ですべてのメモリを解放します)、コードがブロック (タワー) 番号 > 27 に対して実行されるたびに異なる出力が得られることです。グローバルメモリを使用するようにコードを書き換えると、コードが実行されるたびに同じ結果が得られるため、共有メモリとそれを処理する方法に関係があると確信しています。何か案は？

4

1 に答える 1

1

ここで、そのカーネルにはメモリ競合の問題 (つまり、書き込み後の読み取りの正確性) があります。

   if (distance < minimum[0]) {
     minimum[0] = distance;
     minimum[1] = x;
   }

実行されると、ブロック内のすべてのスレッドが最小値の読み取りと書き込みを同時に試みます。ワープ内の複数のスレッドが同じ共有メモリ ロケーションに書き込もうとしたときに何が起こるかは保証されません。また、同じブロック内の他のワープが、書き込み先のメモリ ロケーションからロードするときにどのような値を読み取るかは保証されません。メモリアクセスはアトミックではなく、コードが実行しようとしているタイプのリダクション操作を確実に実行するロックやシリアライゼーションはありません。

同じ問題のより穏やかなバージョンが、カーネルの最後にあるグローバル メモリへの書き戻しに適用されます。

   __syncthreads();
   globalOutputData[y * 2]     = minimum[0];
   globalOutputData[y] = minimum[1];

書き込み前のバリアにより、最小値への書き込みは、"最終的な" (一貫性はありませんが) 値が最小値に格納される前に完了しますが、その後、ブロック内のすべてのスレッドが書き込みを実行します。

各スレッドに距離を計算させ、ブロック全体の距離値の最小値をグローバル メモリに書き出すことを意図している場合は、アトミック メモリ操作を使用する必要があります (共有メモリの場合、これは計算でサポートされています)。 1.2/1.3 および 2.x デバイスのみ)、または明示的な共有メモリ削減を記述します。その後、1 つのスレッドだけがグローバル メモリへの書き戻しを実行する必要があります。

最後に、カーネルがハングする原因となる可能性のある同期の正確性の問題も考えられます。__syncthreads()(PTX bar 命令にマップされます)は、ブロック内のすべてのスレッドが到着し、カーネルが続行する前に命令を実行することを要求します。この種の制御フローを持つ:

 if (x < critters) {
 ....
   __syncthreads();
 ....
 }

ブロック内の一部のスレッドがバリアを迂回して終了し、他のスレッドがバリアで待機している場合、カーネルがハングします。CUDA でのカーネルの実行の正確性を確保するために、__syncthreads() 呼び出しの周りで分岐分岐があってはなりません。

要約すると、現在のコードの少なくとも 3 つの問題について、振り出しに戻ります。

于 2011-04-16T16:29:37.933 に答える