CUDA線程/線程塊之間的通信

我正試圖將幾個任務「映射」到CUDA GPU。有n個任務需要處理。（看僞代碼）CUDA線程/線程塊之間的通信

malloc an boolean array flag[n] and initialize it as false. 
for each work-group in parallel do 
    while there are still unfinished tasks do 
     Do something; 
     for a few j_1, j_2, .. j_m (j_i<k) do 
      Wait until task j_i is finished; [ while(flag[j_i]) ; ] 
      Do Something; 
     end for 
     Do something; 
     Mark task k finished; [ flag[k] = true; ] 
    end while 
end for

由於某些原因，我將不得不使用不同線程塊中的線程。

問題是如何實現等到任務j_i完成;和標記任務k完成; CUDA中的。我的實現是使用布爾數組作爲標誌。然後在任務完成後設置標誌，並讀取標誌以檢查任務是否完成。

但它只適用於小案例，一個大案例，GPU崩潰的原因不明。有沒有更好的方法在CUDA中實現Wait和Mark。

這基本上是CUDA上的線程間通信問題。

來源

2012-09-12 wwnigel

你將不得不使用原子操作這..這花費了很多。 – angainor

我認爲你不需要在CUDA中實現。每一件事都可以在CPU上實現。您正在等待任務完成，然後隨機執行另一個任務。如果你想在CUDA中實現，你不需要等待所有的標誌爲真。你最初知道所有的標誌都是假的。因此，只需對所有線程並行執行Do something並將該標誌更改爲true。

如果要在CUDA中實現，請採用int flag並在完成Do something後繼續添加1，以便您可以在執行Do something之前和之後知道標誌的更改。

如果我的問題不對，請評論。我會盡力改進答案。

來源

2012-09-14 18:43:50 Fr34K

使用__syncthreads()直接在線程塊內同步。然而，線程塊之間的同步更棘手 - 編程模型方法是分成兩個內核。

如果你仔細想想，這是有道理的。執行模型（針對CUDA和OpenCL）適用於處理單元上執行的大量模塊，但沒有提及什麼時候執行。這意味着一些塊會執行，但其他塊不會（他們會等待）。所以如果你有一個__syncblocks()，那麼你會冒死鎖的風險，因爲那些已經執行的將會停止，但那些沒有執行的將永遠不會到達屏障。

您可以在塊之間共享信息（例如使用全局內存和原子），但不能全局同步。

根據你想要做的事情，經常會有另一種解決或解決問題的方法。

來源

2012-09-12 20:54:07 Tom

對不起，我想我沒有把問題弄清楚。我想要做的是**請求一個線程等待來自另一個線程**的數據。例如，線程＃2需要線程＃1中的結果。這意味着線程＃2的一部分必須在線程＃1完成後啓動。
所以我想知道是否有像「sem_post」和「sem_wait」，通過2線程之間的按摩功能？或者我將不得不編寫（線程＃1）和讀取（線程＃2）全局內存，儘管它效率很低？ @Mark Ebersole – wwnigel

在一個塊中，您可以使用共享內存來進行線程間的通信（使用__syncthreads（）），在不同塊的線程之間進行通信，請參閱我的答案：由於可能導致死鎖，因此無法同步塊。 – Tom

我明白了。仍然在這個例子中，如果線程＃2需要來自線程＃1的數據。（線程＃2檢查「標誌」全局內存，線程＃1在完成時寫入「標誌」全局內存）。然後，線程＃2將總是檢查內存並佔用總線，因此線程＃1不能成功地將「標誌」寫入內存。如果它們在同一個線程塊中，我們可以使用__syncthreads（）來確保寫入完成，然後啓動線程＃2的其餘部分？ – wwnigel

由於線程塊可以按任意順序進行調度，並且在它們之間沒有簡單的同步或通信方法，所以您要求的操作並不容易。從CUDA編程指南：

對於並行工作負載，在算法，其中，因爲一些線程需要以互相共享數據同步並行壞了點，有兩種情況：要麼這些線程屬於到同一塊，在這種情況下，他們應該使用__syncthreads（）並通過同一內核調用中的共享內存共享數據，或者它們屬於不同的塊，在這種情況下，它們必須使用兩個獨立的內核調用通過全局內存共享數據，用於寫入和從全局內存中讀取。第二種情況不太理想，因爲它增加了額外內核調用和全局內存流量的開銷。因此，通過將算法映射到CUDA編程模型，以儘可能在單個線程塊內執行需要線程間通信的計算的方式，可以最小化其發生。

所以，如果你不能滿足你在一個線程塊中需要的所有通信，你將需要多個內核調用才能完成你想要的。

我不認爲與OpenCL有任何區別，但我也不能在OpenCL中工作。

來源

2012-09-12 20:58:10

@Tom在4分鐘內擊敗了我...... :) –

對不起，我想我沒有把問題弄清楚。我試圖做的是讓一個線程等待來自另一個線程的數據。例如，線程＃2需要線程＃1中的結果。這意味着線程＃2的一部分必須在線程＃1完成後啓動。
所以我想知道是否有像「sem_post」和「sem_wait」，通過2線程之間的按摩功能？或者我將不得不編寫（線程＃1）和讀取（線程＃2）全局內存，儘管它效率很低？ – wwnigel

謝謝你的回答。我重新描述了問題描述，現在看起來很清楚。你可以再次看看這個問題嗎？ – wwnigel

這種問題最好由一個稍微不同的方法解決：

不分配固定的任務，你的線程，迫使你的線程等待，直到他們的任務變得可用（這是不可能的，因爲CUDA線程不能阻塞）。

而是保留一個可用任務列表（使用原子操作），讓每個線程從該列表中獲取任務。

這仍然是棘手的實施和獲得正確的案件，但至少是可能的。

來源

2012-09-12 22:59:25 tera

這樣，我需要保留一個可用任務的列表，但在這個應用程序中它非常複雜。是否有可能在CUDA中實現「等待」和「標記」？ – wwnigel

不，線程無法等待任何超出其自身塊的內容。 – tera

您可以通過保留兩個單獨的可用任務列表來簡化任務列表的簿記工作 - 其中一個任務從中被處理，另一個任務在準備就緒時添加。一旦處理完第一個列表中的所有任務，就交換列表並啓動一個新的內核。這會導致輕微的性能損失，但是使列表管理非常簡單。 – tera

CUDA線程/線程塊之間的通信

回答

相關問題