2016-05-14 17 views
0

CUDAの最小支配的な問題を解決しています。すべてのスレッドは、いくつかの局所的な結果を発見し、私は最高のものを見つける必要があります。グローバル結果(dev_bestConfigおよびdev_bestValue)には__device__の変数が使用されています。CUDAスレッドの集計結果

__device__ configType dev_bestConfig = 0; 
__device__ int dev_bestValue = INT_MAX; 

__device__ void findMinimalDominantSet(int count, const int *matrix, Lock &lock) 
{ 
    // here is some algorithm that finds local bestValue and bestConfig 

    // set device variables 
    if (bestValue < dev_bestValue) 
    { 
     dev_bestValue = bestValue; 
     dev_bestConfig = bestConfig; 
    } 
} 

私はより多くのスレッドが、私はこのクリティカルセクションを使用ので、同時にメモリアクセスするので、これは動作しないことを知っている:

// set device variables 
    bool isSet = false; 
    do 
    { 
     if (isSet = atomicCAS(lock.mutex, 0, 1) == 0) 
     { 
      // critical section goes here 
      if (bestValue < dev_bestValue) 
      { 
       dev_bestValue = bestValue; 
       dev_bestConfig = bestConfig; 
      } 
     } 
     if (isSet) 
     { 
      *lock.mutex = 0; 
     } 
    } while (!isSet); 

私はこのような何かをする必要がありますこれは実際には期待通りに動作しますが、は本当に遅いです。たとえば、このクリティカルセクションがない場合は、0.1秒がかかります。このクリティカルセクションでは1.8秒かかります。

もっと速くするにはどうすればよいですか?

+2

標準の並列リダクションを使用します。パラレルリダクションでは、最小のbestValueとそれを生成したスレッドのIDが見つかります。削減が完了したら、このIDを使って 'bestConfig'を取得するだけです。あるいは 'bestValue'と' bestConfig'で直接パラレルリダクションを実行することもできます。基本的なパラレルリダクションのチュートリアルは[こちら](https://docs.nvidia.com/cuda/samples/6_Advanced/reduction/doc/reduction.pdf)です。 –

+0

あなたの問題についての詳細を教えてください。実行可能なコードを提供して、それが十分に速くないことを実証することができますか?配列の最小値を見つけようとしていますか?パラレルリダクションはそれを可能にします。 – kangshiyin

+0

ローカルの結果を配列に保存してから、最良の値を検索する方が速く、クリティカルセクションが必要ないので、問題が解決します。もともと共有配列を使用しないようにしたいと思っていました。 –

答えて

1

実際には、クリティカルセクションを避け、最後にロックしました。私は局所的な結果を配列に保存してから、最良のものを探しました。検索は、順次または並行して行うことができます。