加快CUDA原子計算的許多箱/幾個箱

Question

我想優化我在CUDA中的直方圖計算。它使我在相應的OpenMP CPU計算上有了很好的加速。但是，我懷疑（按照直覺）大多數像素都屬於幾個桶。出於論點的緣故，假設我們有256個像素落入讓我們說，兩個桶。

做到這一點最簡單的方法是它似乎是

1. 負載變量到共享內存中
   • 如果需要做量化的負載爲無符號字符等做。
2. 做一個原子加入共享內存
3. 做全球聚結寫。

像這樣：

__global__ void shmem_atomics_reducer(int *data, int *count){ 
    uint tid = blockIdx.x*blockDim.x + threadIdx.x; 

    __shared__ int block_reduced[NUM_THREADS_PER_BLOCK]; 
    block_reduced[threadIdx.x] = 0; 

    __syncthreads(); 

    atomicAdd(&block_reduced[data[tid]],1); 
    __syncthreads(); 

    for(int i=threadIdx.x; i<NUM_BINS; i+=NUM_BINS) 
    atomicAdd(&count[i],block_reduced[i]); 

} 

這個內核的性能下降（自然），當我們在32個箱減少箱櫃的數目，從大約45 GB/s到大約10 GB/s的1 bin。爭用和共享內存銀行衝突是理由。我不知道是否有任何方法可以以任何重要的方式去除這些計算中的任何一個。

我也一直在使用__ballot獲取warp結果，然後使用__popc（）來執行warp級別降低，從而使用_parallelforall博客中的另一個（美麗）想法來實現warp級別降低。

__global__ void ballot_popc_reducer(int *data, int *count){ 
    uint tid = blockIdx.x*blockDim.x + threadIdx.x; 
    uint warp_id = threadIdx.x >> 5; 

    //need lane_ids since we are going warp level 
    uint lane_id = threadIdx.x%32; 

    //for ballot 
    uint warp_set_bits=0; 

    //to store warp level sum 
    __shared__ uint warp_reduced_count[NUM_WARPS_PER_BLOCK]; 
    //shared data 
    __shared__ uint s_data[NUM_THREADS_PER_BLOCK]; 

//load shared data - could store to registers 
    s_data[threadIdx.x] = data[tid]; 

    __syncthreads(); 


//suspicious loop - I think we need more parallelism 
    for(int i=0; i<NUM_BINS; i++){ 
     warp_set_bits = __ballot(s_data[threadIdx.x]==i); 

     if(lane_id==0){ 
     warp_reduced_count[warp_id] = __popc(warp_set_bits); 
     } 

    __syncthreads(); 

     //do warp level reduce 
     //could use shfl, but it does not change the overall picture 
     if(warp_id==0){ 
     int t = threadIdx.x; 
     for(int j = NUM_WARPS_PER_BLOCK/2; j>0; j>>=1){ 
      if(t<j) warp_reduced_count[t] += warp_reduced_count[t+j]; 
      __syncthreads(); 
     } 
     }                                                                 



     __syncthreads(); 


     if(threadIdx.x==0){ 
     atomicAdd(&count[i],warp_reduced_count[0]); 
     } 

    }                                                            

    } 

這給了體面的數量（當然，這是沒有實際意義 - 峯值裝置MEM體重是133 GB /秒，事情似乎取決於啓動配置）爲單箱情況下（35-40 GB/s的1 bin，而使用原子的則爲10-15 GB/s），但是當我們增加bin的數量時，性能會急劇下降。當我們運行32個bin時，性能下降到大約5 GB/s。原因可能是因爲單線程遍歷所有的bin，要求並行化NUM_BINS循環。

我已經嘗試了幾種並行化NUM_BINS循環的方法，其中沒有一個似乎正常工作。例如，人們可以（非常不恰當地）操作內核爲每個bin創建一些塊。這似乎表現了相同的方式，可能是因爲我們將再次遭受試圖從全局內存中讀取的多個塊的爭用。此外，編程是笨重的。同樣，在y方向上並行化分箱也給出了類似的不令人鼓舞的結果。

我嘗試過的另一個想法是動態並行，爲每個bin啓動一個內核。這個速度很慢，可能是由於沒有真正的子核心計算工作和發射開銷。

最有前途的方法似乎是 - 從尼古拉斯·張伯倫article

使用含箱在共享內存每個線程，這表面上是對SHMEM使用非常沉重的這些所謂的私有化直方圖（只有我們Maxwell上每SM有48 kB）。

也許有人可以對此問題有所瞭解？我覺得應該改變算法，不要使用直方圖，要使用頻率較低的東西。否則，我想我們只是使用原子版本。

編輯：我的問題的上下文是計算概率密度函數用於模式分類。我們可以通過使用非參數方法（如Parzen Windows或Kernel Density Estimation）來計算近似直方圖（更準確地說是pdf）。然而，這並沒有克服維數問題，因爲我們需要對每個箱的所有數據點進行求和，當箱數變大時這變得很昂貴。看到這裏：Parzen

Answer 1

我面臨類似的chalanges與羣集工作，但在botton結束，最好的解決方案是使用掃描模式來分組處理。所以，我認爲它不適合你。既然你問了一些這方面的經驗，我會和你分享。

的問題

在你的第一個代碼，我想這與倉減少的數量低性能的交易是與扭曲失速，既然你對每個評價數據進行非常少的處理。當bin的數量增加時，該內核的處理和全局內存負載（數據信息）之間的關係也會增加。您可以使用Nsight的性能分析中的「問題效率」實驗輕鬆檢查。可能你得到的循環週期至少有一個可變的扭曲（Warp Issue Efficiency）。由於我無法將可變經紗的數量提高到接近95％的地方，所以我放棄了這種方法，因爲在某些情況下，情況會變得更糟（內存依賴性佔用了我處理週期的90％。

的洗牌和表決減少是非常有用的，如果倉的數量並不大。如果是大的，線程少量應爲每斌過濾器活躍。所以，你可以用一個結束很多代碼分歧，這對於並行處理來說是非常不希望的，你可以嘗試對分歧進行分組以便消除分支並且具有良好的控制流，所以整個warp/block呈現類似的處理，但是很多機會穿越街區。

一種可行的方案

我不知道在哪裏，但也有您的問題，各地，我看到很好的解決方案。你有沒有試過this one？

你也可以使用一個vectorized load並嘗試類似的東西，但我不知道有多少會是提高你的表現：

__global__ hist(int4 *data, int *count, int N, int rem, unsigned int init) { 

__shared__ unsigned int sBins[N_OF_BINS]; // you may want to declare this one dinamically 
int idx = blockIdx.x * blockDim.x + threadIdx.x; 
if (threadIdx.x < N_OF_BINS) sBins[threadIdx.x] = 0; 

for (int i = 0; i < N; i+= warpSize) { 
    atomicAdd(&sBins[data[i + init].w], 1); 
    atomicAdd(&sBins[data[i + init].x], 1); 
    atomicAdd(&sBins[data[i + init].y], 1); 
    atomicAdd(&sBins[data[i + init].z], 1); 
} 

//process remaining elements if the data is not multiple of 4 
// using recast and a additional control 
for (int i = 0; i < rem; i++) { 
    atomicAdd(&sBins[reinterpret_cast<int*>(data)[N * 4 + init + i]], 1); 
} 
//update your histogram data here 
}