2013-07-07 56 views
2

我正在使用OpenCL。我感興趣的是如何在下面的例子中執行工作項目。工作項目執行順序

我的10000與512的內核工作組大小的一維範圍是跟隨着:

__kernel void 
doStreaming() { 
    unsigned int id = get_global_id(0); 

    if (!isExecutable(id)) 
    return; 

/* do some work */ 
} 

這檢查它是否需要與以下ID或不會進行元素。

假設執行從512個大小的第一個工作組開始,其中20個被isExecutable拒絕。 GPU是否繼續執行其他20個元素而不等第492個元素?

沒有涉及任何障礙或其他同步技術。

回答

2

當一些工作項目遠離通常的分支*做一些工作* /時,他們可以通過從下一個波前(amd)或下一個warp(nvidia)獲取指令來使用管線佔用優勢,因爲當前的warp /波前工作項目正忙做其他事情。但是這會導致內存訪問序列化並清除工作組的訪問順序,從而降低性能。

避免發散的經紗/波前:如果你在if循環中做if語句,它真的很糟糕,所以你找到另一種方法更好。

如果工作組中的每個工作項都具有相同的分支,那麼它就可以。

如果每個工作項目每百次計算的分支數很少,那就沒問題。

嘗試爲所有工作項(難以理解的並行數據/算法)生成相同的條件以利用gpu帶來的功耗。

我知道擺脫最簡單的分支與計算案例的最佳方式是使用全局的yes-no數組。 0 =是,1 =否:總是計算,然後將結果與工作項的yes-no元素相乘。通常每個內核添加1個字節的元素內存訪問要比每個內核執行一個分支要好得多。實際上使對象長度爲2的冪在添加這個1字節後可能會更好。

1

是和否。以下闡述基於NVIDIA的文檔,但我懷疑它在ATI硬件上有什麼不同(儘管實際數字可能有所不同)。通常,工作組的線程是在所謂的warps中執行的,它們是工作組大小的子塊。在NVIDIA硬件上,每個工作組分成每個32個線程的變形。並且每個經紗都以鎖定步驟執行,因此完全平行(可能不是實時並行的,這意味着可能有16個並行線程,然後是16個線程,但在概念上它們完全平行) 。因此,如果這32個線程中只有一個執行了附加代碼,其他代碼將等待它。但所有其他經紗中的紗線都不會關心這一切。

所以是的,可能有線程會不必要地等待其他線程,但是發生的規模要小於整個工作組的大小(任何NVIDIA硬件上的32)。這就是爲什麼儘可能避免經內分支偏差的原因,並且這也是爲什麼只保證在單個經紗內工作的代碼不需要任何同步的原因。共享內存訪問(算法的通用優化)。