cuda - ワープ投票機能について

Question

CUDAプログラミングガイドでは、ワープ投票機能「_all」、「_ any」、「__ballot」の概念を紹介しました。

私の質問は、どのアプリケーションがこれら3つの機能を使用するのかということです。

Answer 1

のプロトタイプ__ballotは次のとおりです

unsigned int __ballot(int predicate);

predicateがゼロ以外の場合、 thビットが設定され__ballotた値を返します。ここで、はスレッドインデックスです。NN

atomicOrおよびと組み合わせて__popc、真の述語を持つ各ワープのスレッド数を累積するために使用できます。

確かに、のプロトタイプatomicOrは

int atomicOr(int* address, int val);

が指す値をatomicOr読み取り、を使用addressしてビットOR演算を実行しval、値をに書きaddress戻し、古い値を戻りパラメーターとして返します。

反対側で__popcは、-bit32パラメータで設定されたビット数を返します。

したがって、指示

volatile __shared__ u32 warp_shared_ballot[MAX_WARPS_PER_BLOCK];

const u32 warp_sum = threadIdx.x >> 5;

atomicOr(&warp_shared_ballot[warp_num],__ballot(data[tid]>threshold));

atomicAdd(&block_shared_accumulate,__popc(warp_shared_ballot[warp_num]));

述語が真であるスレッドの数をカウントするために使用できます。

詳細については、Shane Cook、CUDAプログラミング、MorganKaufmannを参照してください。

Answer 2

__ballotは、CUDAヒストグラムおよびCUDA NPPライブラリで使用され、ビットマスクをすばやく生成し、それを組み込みと組み合わせて__popc、ブール削減を非常に効率的に実装します。

__allの__any導入前に削減に使用されましたが__ballot、他の用途は考えられません。

Answer 3

__ballot APIを使用するアルゴリズムの例として、DM HughesetAlによるカーネル内ストリーム圧縮について説明します。これは、ストリーム圧縮のプレフィックス合計部分で、述語を通過した要素の数を（ワープごとに）カウントするために使用されます。

ここに紙。In-kストリーム圧縮

Answer 4

CUDAは、NVIDIAのアーキテクチャが効率的にサポートするワープ全体のブロードキャストおよびリダクション操作をいくつか提供します。たとえば、__ ballot（predicate）命令は、ワープのすべてのアクティブなスレッドの述語を評価し、ワープのN番目のスレッドとN番目のスレッドの述語がゼロ以外と評価された場合にのみN番目のビットが設定された整数を返します。アクティブです[参照：GPUアーキテクチャの柔軟なソフトウェアプロファイリング]。