動態並行 - 啓動許多小內核很慢

我正在嘗試使用動態並行來改進我在CUDA中的算法。在我最初的CUDA解決方案中，每個線程計算每個塊通用的數字。我想要做的是首先啓動一個粗略的（或低分辨率）內核，其中線程只計算一次公共值（就像每個線程代表一個塊一樣）。然後每個線程創建一個1塊（16x16線程）的小網格，併爲其傳遞公共值啓動一個子內核。從理論上講，它應該更快，因爲它可以節省許多冗餘操作。但實際上，解決方案的工作非常緩慢，我不知道爲什麼。動態並行 - 啓動許多小內核很慢

這是代碼，非常簡單，只是想法。

__global__ coarse_kernel(parameters){ 
    int common_val = compute_common_val(); 
    dim3 dimblock(16, 16, 1); 
    dim3 dimgrid(1, 1, 1); 
    child_kernel <<< dimgrid, dimblock >>> (common_val, parameters); 

} 

__global__ child_kernel(int common_val, parameters){ 
    // use common value 
    do_computations(common_val, parameters); 
}

child_kernels的數量很多，每個線程一個，並且必須有大約400x400線程。據我所知，GPU應該並行處理所有這些內核，對吧？

或子內核以某種方式順序處理？

我的結果顯示，性能比我原來的解決方案慢10倍以上。

來源

2014-01-07 labotsirc

啓動內核，無論是父母還是子女都有成本。如果您的子內核沒有提取太多的並行性，並且對其非平行對象沒有太多好處，那麼您的微弱優勢可能會被子內核啓動開銷抵消。 – JackOLantern

我嘗試在公式中解釋我上面的含義。設置執行子內核的開銷，執行時間和執行相同代碼的時間，而不需要動態並行的幫助。使用動態並行的加速是'ts /（to + te）'。也許（但是這不能從你的代碼中得到）'te JackOLantern

啓動內核，無論是父母還是孩子，都有成本。如果您的子內核沒有提取太多的並行性，並且對其非平行對象沒有太多好處，那麼您的微弱優勢可能會被子內核啓動開銷抵消。

在公式中，讓to是執行一個孩子的內核，其te執行時間和ts無動態並行的幫助下，執行相同的代碼的時間開銷。使用動態並行的加速比爲ts/(to+te)。也許（但這不能從你的代碼中得到）te<ts但te,ts<<to，所以ts/(to+te)約爲(ts/to)<1，你觀察到減速而不是加速。

來源

2014-01-08 22:10:23 JackOLantern

Jack認爲當超過32個內核啓動時，$ to $實際上會爆炸嗎？ – labotsirc

我並不是說'要'實際上爆炸了。我只是解釋你的代碼可能發生了什麼，但是上面的代碼片段中的細節太少，無法提供明確的答案。如果您的子內核沒有太多要做的事情，那麼子內核啓動開銷可能會阻止您觀察到使用動態並行機制的任何改進。 – JackOLantern

它的好處。我會繼續研究這個，但我開始看到DP正常工作的必要條件。謝謝你的幫助。 – labotsirc

動態並行 - 啓動許多小內核很慢

回答

相關問題