CUDA共享內存廣播和__syncthreads行爲

我遇到一個奇怪的問題，以及至少對我來說，它看起來很奇怪，我希望有人也許能夠擺脫它的一些光。我有一個CUDA內核，它依靠共享內存來實現快速本地訪問。據我所知，如果半warp內的所有線程訪問相同的共享內存bank，則該值將被廣播給warp中的線程。另外，從多個經紗到同一家銀行的訪問不會導致銀行衝突，他們只會被序列化。牢記這一點，我已經創建了一個小內核來測試（在我的原始內核遇到問題後）。這裏的片段：CUDA共享內存廣播和__syncthreads行爲

#define NUM_VALUES 16 
#define NUM_LOOPS 1024 

__global__ void shared_memory_test(float *output) 
{ 
    // Create some shared memory 
    __shared__ int dm_delays[NUM_VALUES]; 

    // Loop over NUM_LOOPS 
    float accumulator = 0; 
    for(unsigned c = 0; c < NUM_LOOPS; c++) 
    { 
     // Force shared memory update 
     for(int d = threadIdx.x; d < NUM_VALUES; d++) 
      dm_delays[d] = c * d; 

     // __syncthreads(); 
     for(int d = 0; d < NUM_VALUES; d++) 
      accumulator += dm_delays[d]; 
} 

    // Store accumulated value to global memory 
    for(unsigned d = 0; d < NUM_VALUES; d++) 
     output[d] = accumulator; 
}

我已經16塊上運行這個（半經，不是非常有效，但它只是用於測試目的）。所有線程都應該尋址同一個共享內存組，所以不應該有衝突。但是，情況正好相反。我在Visual Studio 2010上使用Parallel Nsight進行此測試。

對我來說更加神祕的事實是，如果我在外循環中取消對__syncthreads調用的註釋，那麼銀行衝突的數量會急劇增加。

只是一些號碼給你的想法（這是含有一個塊具有16個線程，所以一個單一的半warp，NUM_VALUES = 16，NUM_LOOPS = 1024的網格）：

而不__syncthreads： 4組衝突
與__syncthreads：4096組衝突

我在GTX 670上運行此，設定在compute_capability 3.0

預先感謝您

更新：有人指出，沒有__syncthreads，由於dm_delays的值永遠不會改變，因此外部循環中的NUM_LOOPS讀數正在被編譯器優化掉。現在，在這兩種情況下，我都會遇到4,096次銀行衝突，這對於共享內存的廣播行爲仍然不太合適。

來源

2013-03-25 lessju

一些修改：訪問相同*地址*被廣播（假定一個較新的設備），但訪問相同*銀行*（假設地址不是跨線程相同）被串行化，並且該序列是*定義*銀行衝突的影響。所以說訪問共享內存位置導致序列化，但不是銀行衝突沒有多大意義。 – 2013-03-25 15:41:09

是的你是對的。在代碼段的訪問都應該橫跨threadblock相同的共享存儲器地址（所有的線程應被訪問dm_delays [d]在同一時間，並且存在分配給dm_delays只有16個浮點值） – lessju 2013-03-25 15:53:39

由於dm_delays的值不會改變，如果__syncthreads不存在，這可能是編譯器將1024次讀取優化爲共享內存的情況。在那裏使用__syncthreads，它可能會認爲該值可能會被另一個線程改變，所以它會一遍又一遍地讀取該值。

來源

2013-03-25 15:35:51

有效點。我已經更新了外層循環，以便在每次迭代中更改dm_delays的值。現在我得到了不計「__syncthreads」的4,096次銀行衝突。 – lessju 2013-03-25 15:42:59

我不認爲它可以優化讀取，但它可以減少內部循環的單次讀取，然後將該值乘以NUM_LOOPS（或將16個共享內存值中的每一個帶入寄存器，相同的效果）。可能這就是答案中的含義。您還應該能夠通過在共享內存變量「dm_delays」前面使用'volatile'關鍵字來消除這種影響。所以如果你對此感到滿意的話，你可以接受這個答案。 – 2013-03-25 15:53:56

是的，忘了所有關於易變的，再次。 Howvever我仍然不確定爲什麼我在每次迭代時訪問相同的共享內存位置時出現單個半變形的銀行衝突（如果我強制內部循環使用特定的共享內存位置用於所有迭代，說第二個位置，我仍然會發生衝突，除非我使用第0個位置，在這種情況下我不會）。這是對原始問題的部分回答。 – lessju 2013-03-25 16:11:28

CUDA共享內存廣播和__syncthreads行爲

回答

相關問題