2015-11-03 22 views
0

標題不能包含整個問題:我有一個內核正在執行流壓縮,之後它繼續使用較少的線程數。內核可以更改其塊大小嗎?

我知道一種避免執行未使用線程的方法:返回並執行第二個較小塊大小的內核。我想問的是,如果未使用的線程發散並結束(返回),並且假設它們完全對齊,我可以安全地假定它們不會浪費執行嗎?

除了在兩個連續的內核執行中進行拆分以外,是否有這種常見做法?

非常感謝!

回答

3

單位執行調度和資源調度內SM是32個線程的warp組。

在內核代碼中使用return以任何順序退出線程是完全合法的。然而,存在至少2點考慮:

  1. __syncthreads()在設備代碼的使用取決於具有參與該塊的每個線程。因此,如果線程遇到return語句,那麼該線程不可能參與將來的__syncthreads()語句,因此在一個或多個線程退役後使用__syncthreads()是非法的。

  2. 從執行效率的角度來看(並且還從一個資源調度的角度看,雖然這後一種概念是沒有很好的記載和較爲複雜證明),經編仍然會消耗執行(和其他)資源,直到在所有線程經紗已經退役了。

如果你能在經單位退休的線程,並且不需要的__syncthreads()使用,你應該能夠使GPU資源的相當有效的使用,甚至在退役一些經線threadblock。

爲了完整起見,線程塊的維度是在內核啓動時定義的,並且它們在其後的任何時刻都不能也不會改變。所有線程塊都有最終退休的線程。退休線程的概念不會改變線程塊的尺寸,這在我的用法中(並且與__syncthreads()的用法一致)。

雖然可能與您的問題沒有直接關係,但CUDA Dynamic Parallelism可能是另一種方法來允許線程塊「動態」管理動態變化的執行資源。但是對於給定的線程塊本身,上述所有註釋也適用於CDP案例。

+0

優秀的答案羅伯特,非常有用。我在計算能力3.2的Jetson上,不支持CUDA Dynamic Parallelism。我確實需要__syncthreads(),所以我決定使用實際作業壓縮所有線程,並在返回之前用其他__syncthreads()保留其餘部分。另一種方法是壓縮並重新啓動較小的內核,但我擔心啓動時間開銷。 –

相關問題