3

仕組みは理解し#pragma unrollていますが、次の例があるとします。

__global__ void
test_kernel( const float* B, const float* C, float* A_out)
{
  int j = threadIdx.x + blockIdx.x * blockDim.x;
  if (j < array_size) {
     #pragma unroll
     for (int i = 0; i < LIMIT; i++) {
       A_out[i] = B[i] + C[i];
     }
  }
}

LIMIT上記のカーネルでx、スレッドy数とブロック数で起動される の最適な値を決定したいと考えています。は～のLIMITいずれかになります。100 万は変数にとって非常に大きな数のように見えるため (100 万のループを展開するとレジスタ プレッシャが発生し、コンパイラがその展開を行うかどうかはわかりません)、「公正な」数とは何ですか? そして、その制限をどのように決定しますか？21<<20

4

2 に答える 2

1

CUDA は、作業を複数のスレッドに分割することで公開されるスレッド レベルの並列処理と、CUDA がコンパイル済みコード内の独立した命令を検索して検出する命令レベルの並列処理を利用します。

@talonmies の結果は、ループが 4096 から 8192 反復の間のどこかで展開される可能性があることを示しており、分岐予測や投機的手法などの手法でほとんどの反復オーバーヘッドが最適化されている最新の CPU では、ループの展開による利益が大幅に減少するため、私には驚きでした。実行。

CPU では、たとえば 10 ～ 20 回以上の反復を展開しても得られるものは多くないと思います。また、展開されたループは命令キャッシュ内でより多くのスペースを占有するため、展開にもコストがかかります。CUDA コンパイラは、実行するアンロールの量を決定する際に、コストと利点のトレードオフを考慮します。問題は、4096 回以上の反復をアンロールすることのメリットは何でしょうか? 命令レベルの並列処理を使用して、GPU が独立した命令を同時に実行できるコードをより多く提供するためだと思います。

ループの本体はA_out[i] = B[i] + C[i];. ループ内のロジックは外部変数にアクセスせず、ループの以前の反復からの結果にもアクセスしないため、各反復は他のすべての反復から独立しています。したがってi、順次増加する必要はありません。ループが の各値を完全にランダムな順序で反復した場合でも、最終結果は同じにi なります0。LIMIT - 1この特性により、ループは並列最適化の良い候補になります。

しかし、落とし穴があり、それは私がコメントで述べたことです。ループの反復は、バッファーがおよびバッファーとAは別に格納されている場合にのみ独立しています。バッファーがメモリ内のand/orバッファーと部分的または完全にオーバーラップしている場合、異なる反復間の接続が作成されます。1 つの反復は、 に書き込むことによって、別の反復のおよび入力値を変更できます。したがって、2 つの反復のどちらを最初に実行するかによって、異なる結果が得られます。BCABCBCA

メモリ内の同じ場所を指す複数のポインターは、ポインターのエイリアシングと呼ばれます。そのため、一般に、ポインター エイリアシングは、コードの 1 つのセクションによって 1 つのポインターを介して行われた書き込みが、別のポインターから読み取ったコードの別のセクションによって読み取られた値を変更する可能性があるため、コードのセクション間に「隠れた」接続を引き起こす可能性があります。デフォルトでは、CPU コンパイラは、ポインタのエイリアシングの可能性を考慮したコードを生成し、正しい結果を生成するコードを生成します。問題は CUDA が何をするかということです。なぜなら、talonmies のテスト結果に戻ると、このような大量のアンロールについて私が見ることができる唯一の理由は、命令レベルの並列処理のためにコードが開かれるからです。しかし、これは、この特定の状況では、CUDA がポインターのエイリアシングを考慮していないことを意味します。

再。複数のスレッドを実行することについての質問ですが、スレッドの数を増やしても、通常のシリアル プログラムは自動的に並列プログラムにはなりません。並行して実行できる作業の部分を特定し、それを CUDA カーネルで表現する必要があります。これはスレッドレベルの並列処理と呼ばれ、コードのパフォーマンス向上の主な原因です。さらに、CUDA は各カーネルで独立した命令を検索し、それらを同時に実行する場合があります。これが命令レベルの並列処理です。上級の CUDA プログラマーは、命令レベルの並列処理を念頭に置いて、それを容易にするコードを作成するかもしれませんが、私たち人間は、スレッド レベルの並列処理だけに集中する必要があります。つまり、コードをもう一度見て、並列に実行できる可能性があることを考慮する必要があります。

于 2015-12-30T23:35:28.003 に答える