CUDA縮減和樣本：數據競賽？

我是CUDA新手，目前我正在調查總和縮減樣本，這與我的最終目標相關。CUDA縮減和樣本：數據競賽？

提供的文檔描述瞭如何對內核進行優化以快速減少塊之間的大型數組。 reduction_kernel.cu中的主機函數reduce使用模板在編譯時優化各種內核。

template <class T> 
void reduce(int size, int threads, int blocks, 
      int whichKernel, T *d_idata, T *d_odata) 
{ 
    // 
    // Long list with switch statement to have all optimized functions at compile-time 
    // 

    // amongst which (for instance): 
    case 32: 
     reduce5<T, 32><<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size); 
     break;

編輯：內核reduce5填補了d_odata與d_idata部分款項。更具體地說，它將g_idata的元素與指數2*blockSize*blockIdx.x加到2*blockSize*(blockIdx.x + 1)（不包括）並且將結果存儲在g_odata[blockIdx.x]中。（EDIT-END）

總和是通過跨塊減少直到剩下一個塊而獲得的。主機代碼用於通過在縮小陣列上重複啓動內核來跨「層次」同步內核。代碼中reduction.cpp相關位：

template <class T> 
T benchmarkReduce(int n, numThreads, numBlocks, /* more args */, 
        T *h_odata, T *d_idata, T *d_odata) { 

    // first kernel launch 
    reduce<T>(n, numThreads, numBlocks, whichKernel, d_idata, d_odata); 

    // repeated kernel launches 
    int s=numBlocks; 
    int kernel = whichKernel; 

    while (s > cpuFinalThreshold) 
    { 
     int threads = 0, blocks = 0; 
     getNumBlocksAndThreads(kernel, s, maxBlocks, maxThreads, blocks, threads); 

     reduce<T>(s, threads, blocks, kernel, d_odata, d_odata); 

     if (kernel < 3) 
      s = (s + threads - 1)/threads; 
     else 
      s = (s + (threads*2-1))/(threads*2); 
    } 
}

我很高興與第一內核調用，存儲的d_idata的部分和在d_odata。我擔心第二次內核啓動（在while循環內）：即，內核將讀取和寫入到d_odata，這可能導致數據競爭。例如，第二個塊可以在第一個塊讀取其原始值之前將其部分總和寫入d_odata[1];這是第一個塊的部分總和所需要的。

我錯過了一個細節？

來源

2015-06-29 Mouz

它取決於您還沒有提供的縮減內核的實現。沒有數據競爭就可以進行就地減少。 – Jez

是的，這是可能的，但它不會在示例代碼中發生。例如'reduce5'（由'reduce'封裝）讀取'd_idata [blockIdx.x *（blockSize * 2）+ threadIdx.x]'和'd_idata [blockIdx.x *（blockSize * 2）+ threadIdx。 x + blockSize]'並存儲在'd_odata [blockIdx.x]'中。 – Mouz

是那個真實的代碼？我看不到內核啓動。 –

這已在CUDA 8.0軟件包中修復。 CUDA 8.0即將推出。

來源

2016-05-22 01:44:25

CUDA縮減和樣本：數據競賽？

回答

相關問題