OpenCL：關於SIMT執行模型的基本問題

Question

「SIMT」體系結構的一些概念和設計對我而言仍然不清楚。

從我所看到和閱讀的內容來看，分歧的代碼路徑和if（）總是一個相當糟糕的主意，因爲許多線程可能會以鎖步方式執行。那究竟是什麼意思？什麼是這樣的：

kernel void foo(..., int flag) 
{ 
    if (flag) 
     DO_STUFF 
    else 
     DO_SOMETHING_ELSE 
} 

參數「標誌」是一切工作的單位和同一分支同樣是採取所有的工作單位。現在，GPU是否會執行所有的代碼，然後序列化所有的東西，並且基本上仍然採取未被採用的分支？或者它是否更聰明一些，只要執行了所採用的分支，只要所有線程都同意採用分支？這將永遠是這種情況。

I.e.序列化總是會發生還是僅在需要時纔會發生？對不起，這個愚蠢的問題。 ;）

Answer 1

不，總是不會發生。只有條件在本地工作組中的線程之間不一致時才執行這兩個分支，這意味着如果條件評估爲本地工作組中工作項目之間的不同值，則當前代GPU將執行兩個分支，但只有正確的分支會寫出數值併產生副作用。

因此，保持一致性對於GPU分支的性能至關重要。

Answer 2

對ati不太確定，但對於nvidia來說 - 它很聰明。 如果warp中的每個線程都採用相同的方式，則不會有序列化。

Answer 3

在你的例子中，flag對於所有的工作項都具有相同的值，所以一個好的編譯器會生成代碼，它將把所有的工作項都放在同一個方向上。

但考慮以下情況：

kernel void foo(..., int *buffer) 
{ 
    if (buffer[get_global_id(0)]) 
     DO_STUFF 
    else 
     DO_SOMETHING_ELSE 
} 

這不能保證所有的工作項目將採取相同的路徑，因此需要序列化或控制流的消除。