3

削減を使用してCUDAでこの操作を実行する方法を考えていましたが、それを実行する方法について少し戸惑っています。Cコードは以下のとおりです。覚えておくべき重要な部分-変数precalculatedValueは両方のループイテレータに依存します。また、変数ngoはmのすべての値に固有ではありません...たとえば、 m =0,1,2はngo =1である可能性がありますが、 m =4,5,6,7,8はngo =2である可能性があります。より良い実装提案を提供するのに役立つ場合に備えて、ループイテレータのサイズを含めました。

// macro that translates 2D [i][j] array indices to 1D flattened array indices
#define idx(i,j,lda) ( (j) + ((i)*(lda)) )

int Nobs = 60480;
int NgS  = 1859;
int NgO  = 900;
// ngo goes from [1,900]

// rInd is an initialized (and filled earlier) as:
// rInd = new long int [Nobs];

for (m=0; m<Nobs; m++) {        
    ngo=rInd[m]-1;

    for (n=0; n<NgS; n++) {
            Aggregation[idx(n,ngo,NgO)] += precalculatedValue;
    }
}

以前のケースでは、precalculatedValueが内部ループ変数の関数のみであった場合、値を一意の配列インデックスに保存し、事後に並列削減（Thrust）を追加しました。しかし、この場合は困惑しました。mの値はngoの値に一意にマップされていません。したがって、このコードを効率的に（または実行可能に）して、削減を使用する方法がわかりません。どんなアイデアでも大歓迎です。

4

1 に答える 1

2

ただ刺す...

ループを入れ替えると役立つのではないかと思います。

for (n=0; n<NgS; n++) {
    for (m=0; m<Nobs; m++) {            
        ngo=rInd[m]-1;
        Aggregation[idx(n,ngo,NgO)] += precalculatedValue(m,n);
    }
}

私がこれを行った理由は、 （の関数）の方が、よりもidx急速に変化するためです。したがって、mを内側のループにすると、コヒーレンスが向上します。また、precalculatedValueを（m、n）の関数にしたことに注意してください。これは、疑似コードをより明確にするためです。ngomn

次に、外側のループをホストに残し、内側のループのカーネルを作成することから始めることができます（64,480ウェイの並列処理で、現在のほとんどのGPUを埋めることができます）。

内側のループでは、atomicAdd（）を使用して衝突を処理することから始めます。頻度が低い場合は、FermiGPUのパフォーマンスはそれほど悪くないはずです。いずれの場合も、この計算の算術強度は低いため、帯域幅に制限があります。したがって、これが機能したら、達成している帯域幅を測定し、GPUのピークと比較します。かなり離れている場合は、さらに最適化することを検討してください（おそらく、外側のループを並列化することによって-スレッドブロックごとに1回の反復を行い、共有メモリとスレッドの連携の最適化を使用して内側のループを実行します）。

重要なのは、簡単に始めてパフォーマンスを測定し、最適化する方法を決定することです。

この計算はヒストグラムの計算と非常によく似ていることに注意してください。ヒストグラムの計算にも同様の課題があります。そのため、GPUヒストグラムをグーグルで検索して、それらがどのように実装されているかを確認することをお勧めします。

1つのアイデアは、を使用して（rInd[m]、m）ペアを並べ替えthrust::sort_by_key()てから（rInd重複がグループ化されるため）、それらを繰り返し処理して、衝突することなく削減を行うことができます。（これはヒストグラムを作成する1つの方法です。）これは。を使用して行うこともできますthrust::reduce_by_key()。

于 2012-03-02T03:58:33.747 に答える