CUDAカーネルでの三つ組の発見

Question

私はp.ntp試験粒子を有し、i番目の粒子はデカルト座標tp.rh[i].x,tp.rh[i].y,tp.rh[i].zを有する。このセットの中で私はCLUSTERSを見つける必要があります。それは、hill2 (tp.D_rel < hill2)未満のi番目の粒子に近い粒子を探していることを意味します。そのようなメンバーの数はN_convに格納されます。

データセットを経由するこのサイクルfor (int i = 0; i < p.ntp; i++)を使用します。各i番目のパーティクルについては、セット内の他のメンバとの平方された距離は、tp.D_rel[idx]です。次に私は最初のスレッド(idx == 0)を使用して、自分の状態を満たす症例の数を見つけます。最後に、1より多くの肯定的なケース（N_conv > 1）がある場合、私は可能なクラスターを形成するすべての粒子を一緒に書く必要があります（トリプレット、...）。

私のコードは、i < blockDim.xの場合にのみうまくいきます。どうして？一般的な方法は、どのようにデータセット内のクラスターを見つけるが、トリプレットなどを書くだけですか？

注：私はいくつかのケースが2回見つかることを知っています。

__global__ void check_conv_system(double t, struct s_tp tp, struct s_mp mp, struct s_param p, double *time_step) 
{ 
const uint bid = blockIdx.y * gridDim.x + blockIdx.x; 
const uint tid = threadIdx.x; 
const uint idx = bid * blockDim.x + tid; 
double hill2 = 1.0e+6; 

__shared__ double D[200]; 
__shared__ int ID1[200]; 
__shared__ int ID2[200]; 

if (idx >= p.ntp) return; 
int N_conv; 

for (int i = 0; i < p.ntp; i++) 
{ 

    tp.D_rel[idx] = (double)((tp.rh[i].x - tp.rh[idx].x)*(tp.rh[i].x - tp.rh[idx].x) + 
          (tp.rh[i].y - tp.rh[idx].y)*(tp.rh[i].y - tp.rh[idx].y) + 
          (tp.rh[i].z - tp.rh[idx].z)*(tp.rh[i].z - tp.rh[idx].z)); 
    __syncthreads(); 
    N_conv = 0; 

    if (idx == 0) 
    { 

     for (int n = 0; n < p.ntp; n++) { 
     if ((tp.D_rel[n] < hill2) && (i != n)) { 

     N_conv = N_conv + 1; 
     D[N_conv] = tp.D_rel[n]; 
     ID1[N_conv] = i; 
     ID2[N_conv] = n; 

     } 
     } 

     if (N_conv > 0) { 
     for(int k = 1; k < N_conv; k++) { 
      printf("%lf %lf %d %d \n",t/365.2422, D[k], ID1[k], ID2[k]); 
      } 
     } 

     } //end idx == 0 

    } //end for cycle for i   

} 

Answer 1

RobertCrovellaが言及しているように、MCVの例がないと、わかりにくいです。

しかし、tp.D_del配列はidxインデックスを持つに書き込まれているようだ、とリードバックを__syncthreads()た後、フルレンジインデックスnで。 __syncthreads()への呼び出しは、ブロック内での同期を実行するだけで、デバイス全体では実行されません。その結果、一部のスレッド/ブロックはまだ計算されていないデータにアクセスするため、障害が発生します。

ブロックで計算された値が互いに依存しないように、コードを確認する必要があります。