CUDAの削減を実装する際の問題

Question

私は配列の中で最大の要素を見つけるcudaのreduceを実装しようとしています。私は最小値を見つけるためにこのカーネルを使いました。それはうまく動作しますが、最大値を見つけようとするとうまくいきません。私はアルゴリズムを何度も繰り返してきたので、バグを見つけることができません。どんな助けでも本当に感謝しています。 （私はそこにprint文のコメントを解除する場合に加えて、私は....もかなり頭痛である、異なる出力を得る）

__global__ 
void findMaxAndMin(const float* const d_logLuminance, float* reduceCopy, int length, float* min_logLum, float* max_logLum){ 
    int idx = threadIdx.x + blockDim.x*blockIdx.x; 
    if(idx >= length){ 
     return; 
    } 
    reduceCopy[idx] = d_logLuminance[idx]; 
    __syncthreads(); 

    //do a reduction with max 

    for(int offset = 1;offset < length;offset = offset*2){ 
     if(idx % (offset*2) == 0){ 
      int compIdx = idx + offset; 
      if(compIdx < length){ 
       float newVal = a_max(reduceCopy[idx], reduceCopy[compIdx]); 
       if(idx == 0){ 
        //printf("val %f \n", newVal); 
       } 
       __syncthreads(); 
       reduceCopy[idx] = newVal; 
       __syncthreads(); 
      } 
     } 
     __syncthreads(); 
    } 
    __syncthreads(); 
    if(idx == 0){ 
     *max_logLum = reduceCopy[0]; 
    } 

} 

Answer 1

コードにはいくつかの問題があります。もしそれが最小限に働くなら、あなたは幸運でした。

• 複数のブロックを起動するとします（blockIdx.xを使用します）。あるブロックの結果を別のブロックで繰り返し使用すると、reduceCopy[compIdx]は別のブロックによって設定されることがあります。これに頼ることはできません。ブロックの実行順序を予測したり、ブロックを同期させることはできません。 __syncthreads()は、1つのブロック内でのみ動作するバリアです！

• if(idx >= length) returnは、すべてのスレッドが次の値に達するとは限らないため、危険です。__syncthreads

• __syncthreads()あなたは発散条件if(compIdx < length)にあります。

• a_maxは未定義です。常に最低限の実例を含めることを忘れないでください。私はその機能が何をすべきかを推測することができますが、おそらく別のバグが潜んでいますか？

理論的には並列削減についてよく理解しているようですが、実装はCUDA固有の動作のために失敗します。

具体的には、CUDAの並列リダクションを行う方法の例をお読みになることをお勧めします。