1

最近、プログラミングで単純な概念に遭遇しましたが、それを cuda に実装しようとすると行き詰まりました。何千もの要素があり、それらの間で最も近いペアを見つけたいとします。私atomicMINはグローバルメモリで使用します（縮小したくないと仮定します）ので、各スレッドによって計算された距離がグローバル変数に格納されている距離よりも小さい場合、atomicCAS はそれをより小さい値に置き換えます。たとえば、私はグローバル変数を持っていますfloat gbl_min_dist

これを行うには、次のコードを使用します。

__device__ inline float atomicMin(float *addr, float value){
    float old = *addr, assumed;
    if( old <= value ) return old;
    do{
        assumed = old;
        old = atomicCAS((unsigned int*)addr, __float_as_int(assumed), __float_as_int(value));
    }while( old!=assumed );
    return old;
}

ここで、互いに接近していて、が古い最小距離をそれらの 2 つのポイントによって計算された距離に正常に置き換えた2 つのポイントのインデックスを保存したいとします。atomicMINつまり、グローバル変数で距離が正常にスワップされた場合にのみ、現在距離が小さい2つのポイントのインデックスのみを保存したいということです 

typedef struct {float gbl_min_dist, 
                unsigned int point1,
                unsigned int point2;} global_closest_points;

したがって、ここで、スレッドが を実行するときに、atomicMIN比較対象のスレッドによって提案された値が でスワップされるgbl_min_dist場合、p1、p2 もスレッドからの値とスワップする必要があります。がスワップされていない場合、gbl_min_distポイントを保存したくありません。これにより、間違ったポイントが得られますが、最小距離は正しくなります。

atomicCASスワップが行われたかどうかを確認する戻り値はありますか?

内でこれを実装する方法についてのアイデアはありatomicMINますか?

前もって感謝します

4

3 に答える 3

1

クリティカル セクションを構築して、最小値と対応するポイント インデックスをアトミックに更新できます。atomicCAS()次のリンクは、とを使用して CS を構築する方法の例を示していatomicExch()ます。

https://github.com/ArchaeaSoftware/cudahandbook/blob/master/memory/spinlockReduction.cu

一方、アトミック最小操作を並列削減アルゴリズムに置き換えることをお勧めします。これにより、パフォーマンスが向上する場合があります。

于 2013-10-14T15:03:29.760 に答える
1
  1. クリティカル セクションを使用して、各スレッドが更新中にデータに排他的にアクセスできるようにすることができます。
  2. あなたgbl_min_distは 32 ビット値であるため、p1とを 1 つの 32 ビット値に絞り込む方法を見つけられる場合は、ここでp2提供したカスタム アトミックの回答のようなアプローチを使用できます。

atomicCAS最初のスワップを作成したかどうかを単純に使用して、追加のコードを更新するように調整した場合でも、スレッドの更新間でデータが同期しなくなる競合状態が発生する可能性があると思いますp1。p2

于 2013-10-14T15:05:22.090 に答える