嵌套循環中數組上的二維累積和 - CUDA實現？

我一直在考慮如何使用約簡在CUDA上執行此操作，但是我對如何實現它有點不知所措。 C代碼如下。需要記住的重要部分 - 變量預計算值取決於兩個循環迭代器。另外，變量ngo對於每個值m ... 米 = 0,1,2可能具有NGO = 1，而米 = 4,5,6,7,8可以有NGO = 2，等等。我已經包括循環迭代器的尺寸的情況下它有助於提供更好的實施建議。嵌套循環中數組上的二維累積和 - CUDA實現？

// 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的值。因此，我沒有看到使這種代碼有效（甚至可行）使用減少的方法。任何想法都是最受歡迎的。

來源

2012-03-02 vaca_loca

只是刺...

我懷疑你的換位循環可能的幫助。

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變化更迅速地與ngo（的m功能）比用n，所以使m內環提高的一致性。注意我也做了precalculatedValue函數（m，n），因爲你說這是 - 這使得僞代碼更清晰。

然後，你可以留下主機上外環，並作出內核爲內環開始（64480路並行足夠裝滿最新的GPU）。

在內部循環中，首先使用atomicAdd（）來處理衝突。如果它們不經常發生，費米GPU的性能應該不會太差。無論如何，由於該計算的算術強度較低，您將受到帶寬限制。因此，一旦這個工作正常，測量您正在實現的帶寬，並與GPU的峯值進行比較。如果你離開，然後考慮進一步優化（也許通過並行化外部循環 - 每個線程塊一次迭代，並使用一些共享內存和線程協作優化來完成內部循環）。

關鍵：從簡單開始，測量性能，然後決定如何優化。

注意，這種計算方法看起來非常相似，直方圖計算，它具有類似的挑戰，所以你可能想谷歌的GPU直方圖，看看它們是如何實現的。

一個想法是排序（rInd[m]，m）使用thrust::sort_by_key()，然後對（因爲rInd副本將被組合在一起），你可以在它們之間迭代，做你的削減沒有碰撞。（這是做直方圖的一種方法。）你甚至可以用thrust::reduce_by_key()來做到這一點。

來源

2012-03-02 03:58:33 harrism

+1：「從簡單開始，測量性能，然後決定如何優化。」 - 一般的好建議，特別是像GPU這樣的東西。 – 2012-03-02 12:18:48

謝謝您的詳細解答！作爲第一步，我確定原始C算法輸出與循環移位相同。有各種計算都依賴於兩個迭代器，並且我沒有意識到我可以在不影響輸出正確性的情況下做這個簡單的更改。我會從這裏開始，如果atomicAdd給我一個合適的表現。 – 2012-03-02 15:45:05

但是我確實忘記了一個問題。 precalculatedValue（a，b）類型是雙重的。（實際上它是一個具有2個double值的結構，所以操作重複兩次），我認爲這會造成性能瓶頸。 – 2012-03-02 16:40:46

嵌套循環中數組上的二維累積和 - CUDA實現？

回答

相關問題