使用2個GPU同時調用cudaMalloc時的性能不佳

我有一個應用程序，可以在用戶系統中的GPU之間分配處理負載。基本上，每個GPU有一個CPU線程，當由主應用程序線程週期性地觸發時，它啓動一個GPU處理間隔。使用2個GPU同時調用cudaMalloc時的性能不佳

考慮以下圖像（使用NVIDIA的CUDA分析器工具生成），以GPU處理間隔爲例 - 此處應用程序使用單個GPU。

enter image description here

正如你所看到的GPU處理時間的很大一部分是由兩個分揀作業消耗，我使用這個（推力:: sort_by_key）的推力庫。另外，在啓動實際排序之前，它看起來像thrust :: sort_by_key在引擎蓋下調用一些cudaMallocs。

現在考慮同樣的處理間隔，其中應用已經普及的處理負荷超過兩個GPU：

enter image description here

在一個完美的世界裏，你所期望的2 GPU處理間隔正好一半的單GPU（因爲每個GPU都在做一半的工作）。正如您所看到的，這並非部分原因，因爲由於某種爭用問題，cudaMallocs在同時調用時（有時延長2-3倍）似乎需要更長的時間。我不明白爲什麼需要這樣做，因爲這兩個GPU的內存分配空間是完全獨立的，所以不應該有cudaMalloc上的全系統鎖定 - 每個GPU鎖定會更合理。

爲了證明我的假設，即同時發生的cudaMalloc調用問題，我創建了一個帶有兩個CPU線程（每個GPU）的可笑簡單程序，每個線程多次調用cudaMalloc。我第一次運行這個程序，使得單獨的線程不會在同一時間撥打cudaMalloc：

enter image description here

你看，它需要每〜175分配微秒。接下來，我跑的程序與線程同時調用cudaMalloc：

enter image description here

在這裏，每個呼叫拿了〜538微秒或比以前的情況下，長3倍！毋庸置疑，這大大減緩了我的應用程序的運行速度，而且這意味着只有2個以上的GPU纔會導致問題惡化。

我已經注意到在Linux和Windows上的這種行爲。在Linux上，我使用的是Nvidia驅動程序版本319.60，在Windows上我使用的是327.23版本。我正在使用CUDA工具包5.5。

可能的原因： 我在這些測試中使用GTX 690。這張卡基本上是2 680個像GPU一樣的單元。這是我運行過的唯一的「多GPU」設置，所以或許cudaMalloc問題與690的2個GPU之間的一些硬件依賴有關？

來源

2013-10-05 rmccabe3701

高性能代碼的常見建議是使malloc操作脫離任何性能循環。我意識到這並不是一個微不足道的問題，因爲你使用推力。有高性能的排序庫可以替代推力sort_by_key，這將允許您提前進行分配並將其重新用於排序操作。 [CUB]（http://nvlabs.github.io/cub/），[b40c]（http://code.google.com/p/back40computing/）和[MGPU]（http：//nvlabs.github .io/moderngpu /）都是可能的。 –

是的，我已經看過CUB和B40C（B40C網站說該項目已被棄用）。在我做清除推力的工作之前，我想看看圖書館之間的一些比較圖。你能指點我一些表演數字嗎？你推薦哪個圖書館？ ......似乎推力並不是非常高的性能，例如，我已經用我自己的定製內核交換了一堆推力:: reduce和reduce_by_key調用 - 這樣做將我的處理時間減半。不是開玩笑。 – rmccabe3701

推力實際上是基於b40c的一個特定變體（或曾經是）。對於等效的測試用例，我在b40c和MGPU之間的測試沒有太大差異。在我運行的一個測試中，我只對32位值的22位進行排序。 MGPU有一個撥號盤，我可以轉而只用22位分辨率，我觀察到這樣做的速度提高了40％。我沒有使用CUB多。如果你瀏覽這些鏈接，你可能會發現一些性能數據。例如，某些MGPU性能數據[here]（http://nvlabs.github.io/moderngpu/performance.html#performance） –

總結問題和給出一個可能的解決方案：

的cudaMalloc爭可能是由驅動程序級別的競爭（可能是由於需要切換設備上下文爲talonmies suggestsed）莖和一個能避免這種額外的延遲性能關鍵部分預先通過cudaMalloc-ing和臨時緩衝區。

它看起來像我可能需要重構我的代碼，以便我不調用任何排序例程，調用引擎蓋下的cudaMalloc（在我的情況下是thrust :: sort_by_key）。在這方面，CUB library 看起來很有希望。作爲獎勵，CUB還向用戶公開了CUDA流參數，這也可以提高性能。

查看CUB (CUDA UnBound) equivalent of thrust::gather瞭解從推力轉向CUB的一些細節。

UPDATE：

我退縮了有利於幼崽:: DeviceRadixSort :: SortPairs的推力:: sort_by_key呼叫。
這樣做我的每間隔處理時間減少了毫秒。此外，多GPU爭用問題已經解決 - 卸載至2個GPU幾乎可以將處理時間縮短50％，如預期的那樣。

來源

2013-10-06 16:00:21 rmccabe3701

如果你可以通過這個和你較舊的CUDA問題並接受一些你認爲適合的答案，這將是一件好事。它將它們從未得到答覆的列表中刪除（我們主動儘量保持儘可能短的內容），並且通過搜索可以使其他人更容易找到您是否可以這樣做。謝謝。 – talonmies

哎呀，對不起，我曾想過，當答案被投票時，它會被「接受」。我回去接受了一些對我的舊問題的回答。再一次，對不起，我對這個網站還是有些新鮮的。 – rmccabe3701

我將在前言中聲明一個免責聲明：我並不知道NVIDIA驅動程序的內部，所以這是有點投機的。

您看到的放緩只是由多個線程同時調用設備malloc引起的驅動程序級爭用。設備內存分配需要多個OS系統調用，驅動程序級別上下文切換也是如此。兩種操作都有一個不平凡的延遲時間。當兩個線程同時嘗試並分配內存時，您可能會看到額外的時間，這是由於在兩個設備上分配內存所需的系統調用序列中，從一個設備切換到另一個設備所引起的額外驅動程序延遲。

我能想到的幾個方法，你應該能夠減輕這一點：

你可以推的內存分配的系統調用的開銷降低到零通過編寫自己的自定義推力內存分配器的該設備在初始化過程中分配一塊內存。這將消除每個sort_by_key內的所有系統調用開銷，但編寫自己的用戶內存管理器的努力並非微不足道。另一方面，它會使您的推力代碼的其餘完好無損。
您可以切換到另一個排序庫並收回自己管理臨時內存的分配。如果在初始化階段執行所有分配，則在每個線程的使用期限內，一次內存分配的成本可以分攤到幾乎爲零。

在基於多GPU CUBLAS的線性代數代碼中，我編寫了這兩個想法，並編寫了一個獨立的用戶空間設備內存管理器，該管理器使用一次分配的設備內存池。我發現消除所有中間設備內存分配的開銷成本產生了有用的加速。您的使用案例可能受益於類似的策略。

來源

2013-10-05 10:22:48 talonmies

使用2個GPU同時調用cudaMalloc時的性能不佳

回答

相關問題