2

私の状況: ワープ内の各スレッドは、完全に独立した個別のデータ配列で動作します。すべてのスレッドは、データ配列をループします。ループの反復回数はスレッドごとに異なります。（これには費用がかかります、私は知っています）。

for ループ内で、各スレッドは 3 つの float を計算した後、最大値を保存する必要があります。for ループの後、ワープ内のスレッドは、ワープ内の「隣接スレッド」のみによって計算された最大値をチェックすることによって「通信」します(パリティによって決定されます)。

質問:

  1. 乗算を行って「最大」操作の条件を回避すると、ワープの発散が回避されますよね? (以下のコード例を参照)
  2. (1.) で述べた余分な乗算操作は価値がありますよね? - つまり、あらゆる種類のワープ発散よりもはるかに高速です。
  3. ワープ発散 (すべてのスレッドに対して 1 セットの命令) を引き起こす同じメカニズムは、for ループの最後で暗黙的な「スレッド バリア」(ワープ用) として利用できます (「#pragma非 GPU コンピューティングでの omp for" ステートメント)。したがって、あるスレッドが別のスレッドによって保存された値をチェックする前に、for ループの後でワープの「syncthreads」呼び出しを行う必要はありません。(これは、「シンスレッド」が「GPU 全体」、つまりインターワープとインター MP のみを対象としているためですよね?)

コード例:

__shared__ int N_per_data;  // loaded from host
__shared__ float ** data;  //loaded from host
data = new float*[num_threads_in_warp];
for (int j = 0; j < num_threads_in_warp; ++j)
     data[j] = new float[N_per_data[j]];

// the values of jagged matrix "data" are loaded from host.


__shared__  float **max_data = new float*[num_threads_in_warp];
for (int j = 0; j < num_threads_in_warp; ++j)
     max_data[j] = new float[N_per_data[j]];

for (uint j = 0; j <  N_per_data[threadIdx.x]; ++j)
{
   const float a = f(data[threadIdx.x][j]);
   const float b = g(data[threadIdx.x][j]);
   const float c = h(data[threadIdx.x][j]);

  const int cond_a = (a > b)  &&  (a > c);
  const int cond_b = (b > a)  && (b > c);
  const int cond_c = (c > a)  && (c > b);

  // avoid if-statements.  question (1) and (2)
  max_data[threadIdx.x][j] =   conda_a * a  +  cond_b * b  +  cond_c * c; 
}



 // Question (3):
// No "syncthreads"  necessary in next line:

// access data of your mate at some magic positions (assume it exists):
float my_neighbors_max_at_7 = max_data[threadIdx.x + pow(-1,(threadIdx.x % 2) == 1) ][7];

アルゴリズムを GPU に実装する前に、アルゴリズムのあらゆる側面を調査して、実装作業に見合うだけの価値があることを確認しています。だから我慢してください..

4

1 に答える 1

2
  1. はい
  2. 私の推測では NO です。if を使用して他のバージョンをどのように記述するかによって異なります。
    コンパイラはおそらく述語を使用して、不要な書き込みをマスクします。この場合、実際のスレッドの分岐はなく、実行されたがマスクされた書き込み命令がいくつかあるだけです。
    コンパイラに魔法をかけてもらい、両方のバージョンの逆コンパイルされたコードを比較して、どちらがより良い解決策であるかを判断する必要があります。
    符号付き整数 d = a > b の最大値を計算する特定のケースでは? a : b は 1 つの PTX ISA 命令 max.s32 に変換されるため、実際には複雑にする必要はありません。最大値を一時変数に計算し、無条件書き込みを 1 回行うだけです。
  3. はい。ただし、synthreads バリアはブロック内バリアであり、ブロック間バリアではなく、もちろんインター MP バリアでもありません。
于 2013-02-19T19:24:06.347 に答える