OpenCLで合計計算を並列に実行できますか？

Question

私はOpenCLの初心者です。しかし、私はC/C++の基本とOOPを理解しています。 私の質問は次のとおりです：それはどういうわけか、合計計算タスクを並列で実行することは可能ですか？理論的に可能ですか？以下は、私は私が何をしようとしたものを説明します：

タスクは、たとえば、次のとおりです。

double* values = new double[1000]; //let's pretend it has some random values inside 
double sum = 0.0; 

for(int i = 0; i < 1000; i++) { 
    sum += values[i]; 
} 

私はOpenCLのカーネルで実行しようとしましたどのような（と私はおそらくそれがアクセスするので、それは間違っていると感じ同時に異なるスレッド/タスクから同じ「合計」変数）：

__kernel void calculate2dim(__global float* vectors1dim, 
          __global float output, 
          const unsigned int count) { 
    int i = get_global_id(0); 
    output += vectors1dim[i]; 
} 

このコードが間違っています。理論的にこのようなタスクを並行して実行することができれば誰でも私に答えることができれば大変感謝しています。

Answer 1

アレイの値を並列に集計する場合は、の競合を減らしてを確認し、スレッド間でデータの依存関係がないことを確認する必要があります。

データ依存関係は、スレッドが互いに待って、競合を発生させる原因となります。これは真の並列化を避けるためです。

あなたができることの1つは、配列をN個の配列に分割することです。各配列には、元の配列のサブセクションが含まれています。そして、それぞれ異なる配列でOpenCLカーネル関数を呼び出します。

最後に、すべてのカーネルが大変な作業をしたときに、各アレイの結果を1つにまとめることができます。この操作はCPUが簡単に行うことができます。

重要なのは、各カーネルで実行される計算の間に依存関係がないことです。そのため、データを分割して処理する必要があります。

データに実際の依存関係があるかどうかはわかりませんが、それはわかります。

Answer 2

私が参照用に用意したコードは、その仕事をするべきです。

など。 N要素があり、ワークグループのサイズはWS = 64です。 Nは、2 * WSの倍数であると仮定します（これは重要ですが、1つのワークグループは2 * WS要素の合計を計算します）。その後、指定したカーネルを実行する必要があります。

globalSizeX = 2*WS*(N/(2*WS)); 

結果合計配列は2 * WS要素の部分和を持っています。 （例：sum [1]）は、インデックスが2 * WS〜4 * WS-1である要素の合計を含みます。

globalSizeXが2 * WS（つまり、ワークグループが1つだけであることを意味します）であれば、完了です。結果としての合計[0]を使用してください。 - 今度は手順を繰り返す必要がありますが、今度はの合計配列を入力配列として使用し、他の配列に出力します（2つの配列とそれらの間にピンポンを作成します）。あなたが1つのワークグループしか持てないまではそうではありません。

Hilli Steele/Blelloch並列アルゴリズムも検索してください。 This記事はここでうまく

として有用である可能性があり、実際の例です：

__kernel void par_sum(__global unsigned int* input, __global unsigned int* sum) 
{ 
    int li = get_local_id(0); 
    int groupId = get_group_id(0); 

    __local int our_h[2 * get_group_size(0)]; 
    our_h[2*li + 0] = hist[2*get_group_size(0)*blockId + 2*li + 0]; 
    our_h[2*li + 1] = hist[2*get_group_size(0)*blockId + 2*li + 1]; 

    // sweep up 
    int width = 2; 
    int num_el = 2*get_group_size(0)/width; 
    int wby2 = width>>1; 

    for(int i = 2*BLK_SIZ>>1; i>0; i>>=1) 
    { 

     barrier(CLK_LOCL_MEM_FENCE); 

     if(li < num_el) 
     { 
      int idx = width*(li+1) - 1; 
      our_h[idx] = our_h[idx] + our_h[(idx - wby2)]; 
     } 

     width<<=1; 
     wby2 = width>>1; 
     num_el>>=1; 
    } 

     barrier(CLK_LOCL_MEM_FENCE); 

    // down-sweep 
    if(0 == li) 
     sum[groupId] = our_h[2*get_group_size(0)-1]; // save sum 
}