opencl - opencl の最適なグループサイズ

Question

OpenCL でマンデルブロジェネレーター (静的パラメーターからの 2D 画像) を実行しています。プログラムは簡単です:

__kernel
void mandelbrot(__global uchar * output, 
                const float xstep,
                const float xoffset,
                const float ystep,
                const float yoffset,
                const int maxiter)
{
    int gid_y = get_global_id(1);
    int gid_x = get_global_id(0);

    //calculate x and y on the fly for every pixel. 
    //This is just as fast as reading precalculated rulers from global memory.
    float x = gid_x * xstep + xoffset;
    float y = gid_y * ystep + yoffset;

    float real = 0;
    float imag = 0;

    int out = 0;

    for(int curiter = 0; curiter < maxiter; curiter++) {
        float nreal = real*real - imag*imag + x;
        imag = 2* real*imag + y;
        real = nreal;

        if (real*real + imag*imag > 4.0f) {
            out = curiter;
            break;
        }
    }

    //normalize output
    out *= 256.0 / (float)maxiter;
    output[gid_y * get_global_size(0) + gid_x] = out;

}

[編集] [完全なカーネルを投稿し、提案どおりに行と列を入れ替えました。この方法で、AMD では 18% のパフォーマンスが得られましたが、NVidia では 0% でした。元のコードは

output[get_global_id(0) * get_global_size(1) + get_global_id(1)] = out;

[/編集]

2 つの計算ユニットと 96 個の CUDA コア (計算ユニットあたり 48 コア) を備えた Nvidia Quadro 1000M で実行しています。

カーネルをキューに入れるときにローカルグループのサイズを変更して遊んでいます。これらは、400M ピクセルの画像を生成するときにさまざまなサイズで取得したパフォーマンス結果です。すべての数値は OpenCL プロファイラーからのものであり、OS への最終的なメモリコピーは除外されています。画像は 40992x10272 で、高さと幅の両方が 48 で割り切れます。

rows x columns
8x8: 397 MPixel/s
8x12: 505 MPixel/s
8x16: 523 MPixel/s
8x24: 521 MPixel/s
8x32: 520 MPixel/s
8x48: 520 MPixel/s

1x48: 321 MPixel/s
2x32: 424 MPixel/s
2x48: 523 MPixel/s
4x24: 519 MPixel/s
3x32: 525 MPixel/s
4x32: 525 MPixel/s
4x48: 525 MPixel/s

12x8: 490 MPixel/s
12x12:464 MPixel/s
12x24:505 MPixel/s
12x32:508 MPixel/s
12x48:433 MPixel/s

16x8: 499 MPixel/s
16x12:499 MPixel/s
16x16:472 MPixel/s
16x24:450 MPixel/s
16x32:440 MPixel/s
16x48:418 MPixel/s

これらの数字のいくつかは、私を困惑させます。48列で最良の結果が得られる理由は明らかですが（SIMD操作の仕組みのおかげです）、私には理解できません：

グループごとに 16 行を使用するとパフォーマンスが劇的に低下するのはなぜですか?
1x48 でパフォーマンスが低下するのはなぜですか?
3x32、4x32、8x32 で最高のパフォーマンスが得られるのはなぜですか?!? SIMD プロセッサの 33% がアイドル状態になると予想していましたが、代わりに、2 つの計算ユニットの間にワークグループが存在しているように見えますか?!?
Preferred_WORK_GROUP_SIZE_MULTIPLE が 48 ではなく 32 を返すのはなぜですか?
OpenCL情報構造から取得したものだけを考慮して、任意のGPU(ATI/Nvidia/Intel HD)で最高のパフォーマンスを得るためのジオメトリを把握するための非経験的な方法はありますか?

前もって感謝します

score 22 · Accepted Answer

ここで同様の質問に答えました。以下を読む前に興味深いと思うかもしれません。

グループごとに 16 行を使用するとパフォーマンスが劇的に低下するのはなぜですか?

実際、12行を使用すると、すでに劣化しています。メモリアクセスはトランザクション単位で機能します。トランザクションは、1 回のショットで一定数のバイトをフェッチします。いくつかの作業項目が配列内のいくつかの連続した要素にアクセスしようとすると、1 つのトランザクションですべてを処理できる可能性があります。

この方法でメモリにアクセスするため：

output[get_global_id(0) * get_global_size(1) + get_global_id(1)] = out;

これは、次元 0 のローカルサイズが大きいほど、連続していない要素 (get_global_size(1) 要素で区切られている) にアクセスする必要があるため、トランザクションの数が大きくなることを意味します。また、グローバルメモリアクセスは高価です。

したがって、12/16 行の場合、少なくとも 12/16 トランザクションが必要です。これはあなたの2番目の質問につながります：

1x48 でパフォーマンスが低下するのはなぜですか?

前に言ったことに基づくと、トランザクションの数が最小限になるため、パフォーマンスは優れているはずです。

しかし、ここでアイドルスレッドの問題が発生します。SM ごとに 48 コアに関して得た情報は、既に他の人が指摘しているように間違っています。スレッドは、NVIDIA ハードウェア上で 32 個のグループ (NVIDIA ではワープと呼ばれます) で実行されます。これらのグループはウェーブフロントと呼ばれ、AMD では最大 64 スレッドになる可能性があることに注意してください。この場合、48 個のスレッド (1 × 48) で構成されるワークグループがあるため、64 個のスレッドがスケジュールされていることを意味します。ワープの一部を実行できないため、スケジュールされるのは常に 32 の倍数のスレッド数です。

したがって、この場合、スレッドの 4 分の 1 は何も実行しません。実際に、2x32 で得られた結果と比較すると (まだ 64 スレッド - 2 ワープですが、完全に使用されています)、321 MPixel/s は 424 MPixel/s のほぼ 3/4 です。

この結果も注目に値します: 2x48: 523 MPixel/s . この場合、ワークグループのサイズは 96 (32 の倍数) です。つまり、アイドリングスレッドはありません。

3x32、4x32、8x32 で最高のパフォーマンスが得られるのはなぜですか?!?

答えは、前の 2 つの例から得られます。32 の倍数を使用し、次元 0 のスレッド数を比較的小さく保ちます。しかし、結果を詳しく見てみましょう。

2x32:  424 MPixel/s
3x32:  525 MPixel/s
4x32:  525 MPixel/s
8x32:  520 MPixel/s
16x32: 440 MPixel/s

最後の 2 行のパフォーマンスの低下は、言われたことで簡単に説明できます。ただし、最初の行と 2 番目の行の間のパフォーマンスの向上はそうではありません。

この場合、パフォーマンスの向上は別の場所からもたらされます。2 番目のケースでは、同じ SMで十分なワープが実行され、アクセスメモリのレイテンシが隠されているという事実。REFERRED_WORK_GROUP_SIZE_MULTIPLE 値は、最高のパフォーマンスを得るには、この値の MULTIPLE を使用する必要があることのみを示していることがわかります。複数のワープを同じ SM で同時にスケジュールできます。

それで、それはどのように機能しますか？3x32 の場合を考えてみましょう。3 つのワープで構成されるワークグループがあります。それらは同じワークグループに属しているため、OpenCL 標準で必要とされる同じ SM でスケジュールされます (そうでない場合、ワークグループ内のスレッド間の同期は不可能です)。

最初のワープは、メモリアクセスが必要なため、停止するまで実行を開始します。一方、warp 1 はメモリトランザクションが完了するのを待ち、warp 2 は実行を開始できます。SM には多数のレジスタがあるため、SM はコンテキストを簡単かつ迅速に切り替えて、他のワープを実行できます。ワープ 1 のすべての変数は、ワープ 1 に割り当てられたレジスタにとどまります。その後、ワープ 2 は、メモリアクセスが必要な行にヒットし、ストールします。その時点で、次の実行準備が整ったワープが実行を開始できます。ワープ 3 になる可能性がありますが、メモリアクセスが完了した場合はワープ 1 になる可能性もあります。あなたの場合、2x32 と 3x32 に違いがあるため、実行されるのはワープ 3 のようです。最初のケースでは、メモリアクセスを隠すのに十分なワープがスケジュールされていませんが、2 番目のケースではあります。

実際のところ、これは質問 2 の 1x48 サイズのパフォーマンスの低下にも影響します。

Preferred_WORK_GROUP_SIZE_MULTIPLE が 48 ではなく 32 を返すのはなぜですか?

すでに回答済み。

OpenCL情報構造から取得したものだけを考慮して、任意のGPU(ATI/Nvidia/Intel HD)で最高のパフォーマンスを得るためのジオメトリを把握するための非経験的な方法はありますか?

他の言語と同じです。内部でどのように機能するかを知っていれば、最初に適切なコードを作成するのに役立ちます。ただし、ベンチマークを行い、試行錯誤のプロセスを経て微調整する必要があります。今書いたことは、パフォーマンスにとって重要なことのほんの一部に過ぎないことを覚えておいてください。CPU / GPUの十分な理解と組み合わせてOpenCLからいくつかの情報を照会すると、明らかに役立ちます...しかしそれだけです。

パフォーマンスに影響を与えるパラメータの多くは拮抗しているため、一方で得られるものは他方で失われます。

したがって、ベンチマークを続けてください;）。

score 0 · Accepted Answer

カーネルがグローバルメモリにアクセスする方法は重要であり、ワークグループとグローバルディメンションによって決定されます。

同じワークグループ内の連続するワークアイテムによって書き込まれるアドレスは? ここで、ストライドは get_global_size(1) です。X と Y を交換したい場合があります。一般に、連続する作業項目内の連続する要素に対処する方が高速です。これが最も重要な要素です。
連続するワークグループによってどのアドレスが書き込まれるか? 連続するワークグループは、異なるコンピューティングユニットで同時にスケジュールされることがよくあります。同じチャンネル/バンクをめぐって競合することになり、パフォーマンスが低下する可能性があります。
一般に、バイトではなく 32 ビット整数を書き込むことをお勧めします。

パフォーマンスを最大化するには、より多くのボタンを導入することをお勧めします: 単一のワークアイテム内に複数のピクセル (たとえば 4x2) のブロックを計算するカーネルを記述し、(ブロックサイズ) x (ワークグループサイズ) のすべての組み合わせをベンチマークします。 x (XY スワップ) x (画像サイズ)。次に、GPU に最適なものを選択します。

opencl - opencl の最適なグループ サイズ

3 に答える 3