CUDA：atomicAdd需要太多時間，序列化線程

我有一個內核，它進行一些比較並決定兩個對象是否發生碰撞。我想將碰撞對象的ID存儲到輸出緩衝區。我不想在輸出緩衝區中有空隙。我想將每次碰撞記錄到輸出緩衝區中的唯一索引。CUDA：atomicAdd需要太多時間，序列化線程

所以我在共享內存（局部總和）和全局內存（全局總和）中創建了一個原子變量。下面的代碼顯示了發現碰撞時共享變量的增量。現在我沒有在全局內存中增加原子變量的問題。

__global__ void mykernel(..., unsigned int *gColCnt) { 
    ... 

    __shared__ unsigned int sColCnt; 
    __shared__ unsigned int sIndex; 

    if (threadIdx.x == 0) { 
     sColCnt = 0; 
    } 

    __syncthreads(); 

    unsigned int index = 0; 
    if (colliding) 
     index = atomicAdd(&sColCnt, 1); //!!Time Consuming!! 

    __syncthreads(); 

    if (threadIdx.x == 0) 
     sIndex = atomicAdd(gColCnt, sColCnt); 

    __syncthreads(); 

    if (sColCnt + sIndex > outputSize) { //output buffer is not enough 
     //printf("Exceeds outputsize: %d + %d > %d\n", sColCnt, sIndex, outputSize); 
     return; 
    } 

    if (colliding) { 
     output[sIndex + index] = make_uint2(startId, toId); 
    } 
}

我的問題是，當許多線程嘗試增加原子變量時，它們會被序列化。在寫前綴和之前，我想問問是否有辦法有效地完成這項工作。

由於這一行，我的內核已用時間從13msec增加到44msec。

我發現了一個前綴總和示例代碼，但其引用的鏈接失敗，因爲NVIDIA的討論板已關閉。 https://stackoverflow.com/a/3836944/596547

編輯：我加入了我的代碼末尾太以上。事實上，我確實有一個層次結構。爲了觀察每條代碼行的影響，我設置了每個對象相互碰撞的場景，極端情況以及另一個極端情況，即幾乎沒有對象發生碰撞的情況。

最後，我將共享的原子變量添加到全局變量（gColCnt）中，以通知外部碰撞次數並查找正確的索引值。我想我必須以任何方式使用atomicAdd。

來源

2012-07-22 phoad

'atomicAdd'序列化的定義，所以當你預測的碰撞將是稀疏的，你應該只依賴於它。也許你可以重構計算以分層使用原子：首先，積累到每個線程塊中的__shared__變量。在後處理中（例如，在上面內核的第3個__syncthreads之後），可以將每個塊的衝突累加到全局內存中的單個變量中。 – 2012-07-22 06:17:43

其實我有層次結構。但是同一塊中的線程也會在atomicAdd上爲__shared__變量序列化，至少在第一種極端情況下，每個對象都會相互碰撞。 – phoad 2012-07-22 11:46:42

www.cuvilib.com/Reduction.pdf我找到了M.哈里斯的教程。我會盡量利用它。 – phoad 2012-07-22 12:17:08

考慮使用並行流壓縮算法，例如thrust::copy_if。

來源

2012-07-22 14:24:24

我想我已經想過爲什麼我不能調用thrust :: copy_if，但我現在無法弄清楚原因。我會嘗試內核中的小前綴和，然後重新思考並嘗試這個並通知。謝謝。 – phoad 2012-07-22 14:49:01

是的，我不知道會發現多少碰撞。所以輸出緩衝區的大小是未知的（可能太大）。我做了初始輸出緩衝區大小估計，並儘可能多地將衝突複製到緩衝區。如果需要更多緩衝區，我將緩衝區放大並再次調用內核。這是一個好方法嗎？ – phoad 2012-07-22 15:58:17

看起來像一個單線程可以找到零或一個碰撞？然後，您可以嘗試爲每個線程分配一個插槽，並在發現衝突時寫入該插槽。然後，看看用流式壓縮算法收集結果與找出衝突的內核相比多久，然後看看它是否值得追求更高級的解決方案。 – 2012-07-22 17:13:08

NVIDIA博客文章涉及：http://devblogs.nvidia.com/parallelforall/gpu-pro-tip-fast-histograms-using-shared-atomics-maxwell/

來源

2015-03-21 07:11:18 Tanguy

CUDA：atomicAdd需要太多時間，序列化線程

回答

相關問題