ネストされたループ内の配列の2次元の累積合計 - CUDA実装？

Question

私は削減を使ってCUDAでこの操作を実行する方法を考えていましたが、それを達成する方法について少し迷っています。 Cコードは以下の通りです。注意すべき重要な部分 - 変数precalculatedValueは、とループイテレータに依存します。また、変数ngoは、mのすべての値に固有のものではありません。 メートル = 4,5,6,7,8、私は場合に、ループイテレータのサイズが含まれている等NGO = 2を有することができる一方、メートル = 0,1,2は、NGO = 1を持っているかもしれませんより良い実装提案を提供するのに役立ちます。前場合

// 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は、内側ループ変数のみの関数であった場合、私はユニークな配列インデックスの値を保存し、事後に平行還元（推力）でそれらを添加しました。しかし、この場合、私は困惑しています。mの値は、ngoの値に一意にマップされません。したがって、私は、このコードを効率化する（または実行可能にする）方法を、削減を使用する方法は見ていません。どんなアイディアも大歓迎です。

Answer 1

だけ刺す...

私はあなたのループを転置することに役立つかもしれないと思われます。

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を作り、nよりもngo（mの関数）でより急速に変化するので、私はこれをしなかった理由です。注：私はprecalculatedValueを（m、n）の関数にしました。あなたはそれがそうであると言いました。これは擬似コードをより明確にします。

次に、ホストに外側のループを残して、内部ループ用のカーネルを作成することから始めます（64,480方向の並列処理は、最新のGPUを満たすのに十分です）。

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

キー：簡単に開始し、パフォーマンスを測定し、最適化の方法を決めます。

この計算は類似した課題を持つヒストグラム計算と非常によく似ていることに注意してください.GPUヒストグラムがGoogleにどのように実装されているかを確認することをおすすめします。

ひとつのアイデアは、thrust::sort_by_key()を使用して（rInd[m]、m）のペアをソートし、次に（rInd重複が一緒にグループ化されますので）し、あなたがそれらを反復し、衝突することなく、自分の削減を行うことができます。これはヒストグラムを行う1つの方法です。thrust::reduce_by_key()でこれを行うこともできます。