4 
const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

上記のカーネルは、ループごとに 10 回実行されるベクトル加算です。プログラミング ガイドとスタック オーバーフローを使用して、グローバル メモリがどのように機能するかを理解しましたが、コードを見て、グローバル メモリに適切にアクセスしているかどうかはまだわかりません。私は連続してアクセスしており、整列した方法で推測しています。カードは、配列 a、b、および c に対して 128kb のグローバル メモリ チャンクをロードしますか? 次に、処理される 32 個の gid インデックスごとに 1 回、各配列の 128kb チャンクをロードしますか? (4*32=128) グローバル メモリ帯域幅を無駄にしていないようですよね?

ところで、コンピュート プロファイラーは gld と gst の効率が 1.00003 であることを示しています。これは奇妙に思えます。すべてのストアとロードが合体した場合、1.0 になると思いました。1.0以上はどうですか？

4

1 に答える 1

12

はい、メモリアクセスパターンはほぼ最適です。各ハーフワープは、16個の連続する32ビットワードにアクセスしています。さらに、バッファ自体が整列され、各ハーフワープのstartindexが16の倍数であるため、アクセスは64バイトに整列されます。したがって、各ハーフワープは1つの64バイトトランザクションを生成します。したがって、非合体アクセスによってメモリ帯域幅を浪費するべきではありません。

最後の質問で例を求めたので、このコードを他のコードに変更してみましょう（最適ではないアクセスパターン（ループは実際には何もしないため、無視します）：

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

最初に、これがコンピューティング1.3（GT200）ハードウェアでどのように機能するかを見てみましょう

aへの書き込みの場合、これはわずかに最適でないパターンを生成します（id範囲と対応するアクセスパターンによって識別されるハーフワープに従います）。

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

したがって、基本的に帯域幅の約半分を浪費しています（奇数のハーフワープのアクセス幅が2倍未満の場合、より多くのアクセスが生成されるため、あまり役に立ちません。いわば、より多くのバイトを浪費するよりも高速ではありません）。

bからの読み取りの場合、スレッドは配列の偶数要素にのみアクセスするため、ハーフワープごとにすべてのアクセスは128バイトに整列されたブロックにあります（最初の要素は128B境界にあります。これは、その要素のgidが16=>の倍数であるためです。インデックスは32の倍数であり、4バイト要素の場合、アドレスオフセットは128Bの倍数であることを意味します。アクセスパターンは128Bブロック全体に広がるため、これにより、ハーフワープごとに128B転送が実行され、帯域幅の半分が再び使用されます。

cからの読み取りは、各スレッドが独自の128Bブロックでインデックスを作成するという、最悪のシナリオの1つを生成します。したがって、各スレッドは独自の転送を必要とします。これは、一方ではシリアル化シナリオのビットです（ただし、通常ほど悪くはありませんが、ハードウェアは転送をオーバーラップできる必要があるため）。さらに悪いことに、これによりスレッドごとに32Bブロックが転送され、帯域幅の7/8が無駄になります（4B /スレッド、32B / 4B = 8にアクセスするため、帯域幅の1/8のみが使用されます）。これはナイーブなmatrixtransposesのアクセスパターンであるため、ローカルメモリを使用して実行することを強くお勧めします（経験から言えば）。

Compute 1.0（G80）

ここで、良好なアクセスを作成する唯一のパターンは元のパターンです。例のすべてのパターンは、完全に非合体のアクセスを作成し、帯域幅の7/8を浪費します（32B転送/スレッド、上記を参照）。G80ハードウェアの場合、ハーフワープのn番目のスレッドがn番目の要素にアクセスしないすべてのアクセスは、そのような非合体アクセスを作成します

Compute 2.0（Fermi）

ここでは、メモリにアクセスするたびに128Bのトランザクションが作成されます（すべてのデータを収集するために必要な数なので、最悪の場合は16x128B）が、それらはキャッシュされるため、データの転送先がわかりにくくなります。今のところ、キャッシュがすべてのデータを保持するのに十分な大きさであり、競合がないことを前提としています。したがって、すべての128Bキャッシュラインは最大で1回転送されます。さらにmoeにハーフワープのシリアル化された実行を想定させて、決定論的なキャッシュ占有を実現します。

bへのアクセスは、常に128Bブロックを転送します（対応するメモリ領域に他のスレッドインデックスはありません）。cにアクセスすると、スレッドごとに128Bの転送が生成されます（可能な限り最悪のアクセスパターン）。

aへのアクセスについては、次のとおりです（今のところ読み取りのように扱います）。

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

したがって、大規模なアレイの場合、へのアクセスは理論的にはほとんど帯域幅を浪費しません。この例では、cへのアクセスによってキャッシュがかなりうまく破棄されるため、現実はもちろんそれほど良くありません。

プロファイラーの場合、1.0を超える効率は、単に浮動小数点の不正確さの結果であると思います。

お役に立てば幸い

于 2010-10-04T18:46:37.187 に答える