還元の最終的な方法

Question

this linkで説明した最初の問題の続きを取り上げます。

OpenCLで複数の合計を減らす方法を適用したいと思います（私のGPUデバイスはOpenCL 1.2のみをサポートしています）。

現在のところ、1回の合計削減（つまり、1回の繰り返し ）のバージョンを作成しました。このバージョンでは、簡単にするために、順次CPUループを使用して各部分合計の合計を計算し、合計の最終値を取得しました。

私の前例のアドバイスから、私の問題は、NDRangeKernel関数をもう一度呼び出して（つまり、2度目のカーネルコードを実行して）最終合計を実行する方法がわからないことです。

実際、2回目の呼び出しでは、部分的な合計（それ自体が最初の呼び出しNDRangeKernelから計算されたもの）の合計を得るために、常に同じ問題に直面します。再帰的な問題です。

入力配列のサイズが10240000で、WorkGroup sizeが16の場合、10000*2^10/2^4 = 10000*2^6 = 640000 WorkGroupsとなります。

だから、最初の呼び出しの後、私は640000 partial sumsを取得する：すべてのこれらの部分和の最終サメイションをどのように扱いますか？私はカーネルコード別の時間を呼び出すと、例えば、WorkGroup size = 16とグローバルsize = 640000、私はnWorkGroups = 640000/16 = 40000 partial sumsを取得しますので、私はカーネルコードをもう一度呼び出し、nWorkGroups < WorkGroup sizeまで、このプロセスを繰り返す必要があります。誰かがカーネルのものを、この上記のコードスニペットを説明できる場合

たぶん私は非常によく第二段階、「2段減速」（on this link, I think this is the case of searching for minimum into input array）

__kernel 
void reduce(__global float* buffer, 
      __local float* scratch, 
      __const int length, 
      __global float* result) { 

    int global_index = get_global_id(0); 
    float accumulator = INFINITY; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
    float element = buffer[global_index]; 
    accumulator = (accumulator < element) ? accumulator : element; 
    global_index += get_global_size(0); 
    } 

    // Perform parallel reduction 
    ... 

からカーネルコードの大部分はこの部分を理解していませんコードは行います。

還元の第2段階、すなわち最終合算との関係はありますか？

は、あなたが私の問題を理解していない場合は、私に詳細を聞いてお気軽に。

感謝のコメントで述べたように

Answer 1

：文

入力配列のサイズは10240000で、ワークグループサイズが16であれば、我々は取得10000 * 2^10月2日^ 4 = 10000 * 2^6 = 640000ワークグループ。

が正しくありません。 「任意の」ワークグループサイズと「任意の」数のワークグループを選択できます。ここで選択する番号は、ターゲットデバイスに合わせて調整することができます。例えば、デバイスは特定のローカルメモリサイズを有することができる。これはclDeviceGetInfoで問い合わせることができます。

cl_ulong localMemSize = 0; 
clDeviceGetInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, 
    sizeof(cl_ulong), &localMemSize, nullptr); 

これは、各作業グループは、ローカルメモリの

sizeof(cl_float) * workGroupSize 

バイトを必要とするという事実を考慮すると、ローカルワークグループの大きさを計算するために使用することができます。

同様に、ワークグループの数は、他のデバイス固有のパラメータから派生することがあります。

---

還元自体に関する重要な点は、ワークグループのサイズがを処理することができる配列のサイズを制限しないことです。私はまた、全体としてのアルゴリズムを理解した上で、いくつかの困難を持っていたので、私はいくつかの画像は千の言葉の価値があることを期待して、ここでそれを説明しようとした：あなたが見ることができるように、多数の

ワークグループおよびワークグループのサイズは固定されており、入力配列の長さには依存しません。例では8つのワークグループを使用しています（グローバルサイズは24です）。長さ64の配列は、処理される。これは主に入力配列をたどる最初のループと、グローバルワークサイズ（ここでは24）に等しい「ステップサイズ」が原因です。結果はとなり、24スレッドのそれぞれの累積値はとなります。 これらのは、並行して縮小されます。