改善CUDA中的異步執行

我目前正在編寫一個程序，使用CUDA API在GPU上執行大型模擬。爲了加速性能，我試圖同時運行我的內核，然後再次將結果異步複製到主機內存中。代碼看起來大致是這樣的：改善CUDA中的異步執行

#define NSTREAMS 8 
#define BLOCKDIMX 16 
#define BLOCKDIMY 16 

void domainUpdate(float* domain_cpu,  // pointer to domain on host 
        float* domain_gpu,  // pointer to domain on device 
        const unsigned int dimX, 
        const unsigned int dimY, 
        const unsigned int dimZ) 
{ 
    dim3 blocks((dimX + BLOCKDIMX - 1)/BLOCKDIMX, (dimY + BLOCKDIMY - 1)/BLOCKDIMY); 
    dim3 threads(BLOCKDIMX, BLOCKDIMY); 

    for (unsigned int ii = 0; ii < NSTREAMS; ++ii) { 

     updateDomain3D<<<blocks,threads, 0, streams[ii]>>>(domain_gpu, 
                  dimX, 0, dimX - 1, // dimX, minX, maxX 
                  dimY, 0, dimY - 1, // dimY, minY, maxY 
                  dimZ, dimZ * ii/NSTREAMS, dimZ * (ii + 1)/NSTREAMS - 1); // dimZ, minZ, maxZ 

     unsigned int offset = dimX * dimY * dimZ * ii/NSTREAMS; 
     cudaMemcpyAsync(domain_cpu + offset , 
         domain_gpu+ offset , 
         sizeof(float) * dimX * dimY * dimZ/NSTREAMS, 
         cudaMemcpyDeviceToHost, streams[ii]); 
    } 

    cudaDeviceSynchronize(); 
}

總而言之，這只是一個簡單的循環，遍歷所有的流（8在這種情況下）和分工。這實際上是一個更快的交易（高達30％的性能增益），儘管可能比我希望的要少。我分析了Nvidia的計算視覺探查一個典型的週期，並執行如下：

CUDA API trace in the Compute Visual Profiler

正如可以在圖片中可以看出，內核確實重疊，但從來沒有超過兩個內核在運行同時。我對不同數量的流和模擬域的不同大小進行了相同的嘗試，但情況總是如此。

所以我的問題是：有沒有辦法鼓勵/強制GPU調度程序在同一時間運行兩件事？或者，這是否取決於代碼中無法表示的GPU設備的限制？我的系統規格爲：64位Windows 7和GeForce GTX 670圖形卡（即開普勒架構，計算能力3.0）。

來源

2013-04-25 Yellow

僅當GPU有剩餘資源才能運行第二個內核時，內核纔會重疊。一旦GPU完全加載，並行運行更多內核就沒有收益，因此驅動程序不會這樣做。

來源

2013-04-25 11:30:37 tera

即使有一個很小的內核，就像一些塊一樣，同一時間內也不會有超過兩個內核運行。所以GPU的物理尺寸不可能是整個故事，可以嗎？ – Yellow 2013-04-25 12:18:16

是的，它可以。什麼是「小內核」？多少塊？每塊有多少個線程？他們使用共享內存嗎？寄存器？除非你分析了內核的資源利用率，否則你不知道可以運行多少個內核。 Windows（GPU處於WDDM模式時）也可以通過批量GPU活動來影響併發性。 GPU不限於同時運行兩件事。 – 2013-04-25 13:14:51

這是一個很好的觀點，我沒有完全想到所有的共享內存和註冊要求，我不明白這會影響性能。我嘗試的一個'小'內核是例如8x8塊和16x16線程。其中，理論上適合GPU的負載更多，我會說。它每個線程使用33個寄存器，每塊大約2 kB共享內存。這很多嗎？ – Yellow 2013-04-26 13:19:05

改善CUDA中的異步執行

回答

相關問題