1

私はかなり長い間、単純なスキャンを機能させようとしてきました。小さな問題の場合、出力は正しいですが、大きな出力の場合、正しい結果が得られることがあります。Apple の OpenCL の例を確認しましたが、基本的に同じことを行っています (atm を無視している銀行の競合を除いて)。最初のフェーズのコードは次のとおりです。

__kernel void
scan_init(__global int * input,
          __global int * sums)
{
  int gid = get_global_id(0);
  int lid = get_local_id(0);
  int chunk_size = get_local_size(0)*2;

  int chunk = gid/chunk_size;
  int offset = chunk*chunk_size;

  reduction(input, offset);

  // store sums
  if(lid==0)
  {
    sums[chunk] = input[(chunk+1)*chunk_size-1];
  }

  downsweep(input, offset);
}

そしてリダクション関数自体：

void reduction(__global int * input,
      int offset)
{
 int stride = 1;
 int grp_size = get_local_size(0);
 int lid = get_local_id(0);

 for(int d = grp_size; d > 0; d>>=1)
 {
   barrier(CLK_GLOBAL_MEM_FENCE);

   if(lid < d)
   {
     int ai = stride*(2*lid+1)-1+offset;
     int bi = stride*(2*lid+2)-1+offset;
     input[bi] += input[ai];
   }

   stride *= 2;
  }
}

2 番目のフェーズでは、パーシャルの合計を使用して各要素の合計を作成します。

void downsweep(__global int * input,
        const unsigned int offset)
{
  int grp_size = get_local_size(0);
  int lid = get_local_id(0);
  int stride = grp_size*2;

  for(int d = 1; d <= grp_size; d *=2)
  {
    barrier(CLK_GLOBAL_MEM_FENCE);

    stride >>=1;

    if(lid+1 < d)
    {
      int src = 2*(lid + 1)*stride-1+offset;
      int dest = src + stride;
      input[dest]+=input[src];
    }
  }
}

入力は、ローカル作業サイズの倍数のサイズにパディングされます。各ワーク グループは、そのサイズの 2 倍のチャンクをスキャンできます。結果を確認するために使用する sums 配列に各チャンクの合計を保存します。以下は、1 の配列の入力サイズ 4000 の出力です。

Chunk size: 1024
Chunks: 4
Scan global size: 4096
Local work size: 512
Sum size: 4
0:1024  1:1120  2:2904  3:928

ただし、期待される結果は

0:1024  1:1024  2:1024  3:928

コードをもう一度実行すると、次のようになります。

0:1056  1:5376  2:1024  3:928 
0:1024  1:1088  2:1280  3:992 
0:5944  1:11156 2:3662  3:1900  
0:7872  1:1056  2:2111  3:1248

カーネルへの呼び出しは次のとおりです。

clEnqueueNDRangeKernel(cl_ctx->queue, scan_init, 1, NULL, &scan_global_size, &local_work_size, 0, NULL, NULL);

グローバル サイズは 4096 で、ローカル サイズは 512 です。ローカル ワーク グループのサイズを 64 に制限すると、出力は次のようになります。

0:128  1:128  2:128  3:288  4:128  5:128  6:192  7:192  
8:192  9:254  10:128  11:256  12:128  13:360  14:128  15:128  
16:128  17:128  18:128  19:288  20:128  21:128  22:128  23:128  
24:192  25:128  26:128  27:192  28:128  29:128  30:128  31:32

入力サイズを 512 と任意のチャンク サイズに変更すると、すべてがうまく機能します。

最後に、入力サイズ 513 とグループ サイズ 256 を使用すると (つまり、それぞれ 512 要素を持つ 2 つのチャンクがあり、2 番目のチャンクには最初の要素のみが 1 に設定されている)、最初のフェーズの結果は次のようになります。

0:1  1:2  2:1  3:6  4:1  5:2  6:1  7:14  
8:1  9:2  10:1  11:6  12:1  13:2  14:1  15:28  
16:1  17:2  18:1  19:6  20:1  21:2  22:1  23:14  
24:1  25:2  26:1  27:6  28:1  29:2  30:1  31:56  
32:1  33:2  34:1  35:6  36:1  37:2  38:1  39:14  
40:1  41:2  42:1  43:6  44:1  45:2  46:1  47:28  
48:1  49:2  50:1  51:6  52:1  53:2  54:1  55:14  
56:1  57:2  58:1  59:6  60:1  61:2  62:1  63:148

あるべき場所:

0:1  1:2  2:1  3:4  4:1  5:2  6:1  7:8  
8:1  9:2  10:1  11:4  12:1  13:2  14:1  15:16  
16:1  17:2  18:1  19:4  20:1  21:2  22:1  23:8  
24:1  25:2  26:1  27:4  28:1  29:2  30:1  31:32  
32:1  33:2  34:1  35:4  36:1  37:2  38:1  39:8  
40:1  41:2  42:1  43:4  44:1  45:2  46:1  47:16  
48:1  49:2  50:1  51:4  52:1  53:2  54:1  55:8  
56:1  57:2  58:1  59:4  60:1  61:2  62:1  63:64

私の推測では、異なるスレッドが同じデータに同時にアクセスするのは問題ですが、すべてのワーク グループが入力データの異なるチャンクを処理しているため、これは当てはまりません。この問題に関するヘルプは大歓迎です!!

4

1 に答える 1

4

問題は、ワークグループ間の同期ではないバリア()に関係していると思われます。各ワークグループには独自のバリアがあり、ワークグループ自体の順序については保証されません。入力セットのサイズを 512 に変更すると、すべてのワークグループが同じマルチプロセッサで実行されるため、偶発的に同期される場合があります。

チャンク変数は get_group_id(0)/2 です。これは、2 つのワークグループ全体が同じチャンクに割り当てられていることを意味します。あなたはおそらくそれを逆に望んでいます。それらがロックステップで実行された場合、ロードとストアの依存関係が一致するため、単純に互いの作業を上書きします。そうしないと、常に値を複数回合計する方向で、干渉する場合と干渉しない場合があります。

この問題のヒントは、質問自体にあります。「各ワーク グループは、そのサイズの 2 倍のチャンクをスキャンできます。」これは、配列サイズの半分の合計作業サイズで十分であることを意味するはずです。

downsweep() のループにも奇妙な点があります。最初の反復は何もしません。lid+1>=1 で、d は 1 から始まります。これは取るに足らない余分な繰り返しかもしれませんが、計画では 1 つずれています。

于 2010-07-23T02:43:28.063 に答える