4

基数ソートを行うカーネル内から専用スキャン関数を呼び出したい。しかし、排他的スキャンは、その作業を行うためにスレッドの半分しか必要としません。

排他的スキャンアルゴリズムには、いくつかの__syncthreads（）が必要です。私が最初に次のようなステートメントを持っている場合

if（threadIdx.x> NTHREADS / 2）return;

これらのスレッドは、許可されていない排他的スキャン同期スレッドには参加しません。この問題を回避する方法はありますか？__syncthread（）で囲まれた排他的スキャンの呼び出しがあります。

4

2 に答える 2

4

このようなものはうまくいくはずです（早期リターンを使用しないでください）：

__syncthreads(); // at entry to exclusive scan region
// begin exclusive scan function
if (threadIdx.x < NTHREADS/2) {
  // do first phase of exclusive scan up to first syncthreads
  }
__syncthreads(); // first syncthreads in exclusive scan function
if (threadIdx.x < NTHREADS/2) {
  // do second phase of exclusive scan up to second syncthreads
  }
__syncthreads(); // second syncthreads in exclusive scan function
(... etc.)
__syncthreads(); // at exit from exclusive scan region

__syncthreads() これはやや面倒ですが、私が知っている唯一の使用方法に関する法律の文言を遵守する方法です。また、指定した方法でコードをそのままにしておくこともできます。作業を行っていないスレッドは早期にリターン/終了します。それはうまくいくかもしれません、おそらくうまくいくでしょう。ただし、将来のアーキテクチャや新しいツールチェーンで機能するという保証はありません。

于 2013-03-18T18:25:55.307 に答える
2

別の方法を指摘すると、 に
相当するインライン アセンブリを使用することもできます__syncthreads()。これにより、計算機能 2.0 以降で使用できる参加スレッドの数にオプションの引数を使用できます。このようなものが動作するはずです:

#define __syncthreads_active(active_threads) asm volatile("bar.sync 0, %0;" :: "r"(active_threads));

if(threadIdx.x >= NTHREADS/2) return;

int active_warps = (NTHREADS/2 + warpSize) / warpSize;
int active_threads = active_warps * warpSize; // hopefully the compiler will optimize this to a simple active_threads = (NTHREADS/2 + warpSize) & ~32

__syncthreads_active(active_threads);
// do some work...
__syncthreads_active(active_threads);
// do some more work...
__syncthreads_active(active_threads);

免責事項: ブラウザで記述され、完全にテストされていません!

しかし、それが苦労する価値があるかどうかは別の問題です.

于 2013-03-19T03:17:39.007 に答える