グローバルメモリに関して私の単純なOpenCLカーネルで何が起こっているかを合理化する

StackOverflow https://stackoverflow.com/questions/3857981

  •  27-09-2019
  •  | 
  •  

質問 

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インデックスごとに、各配列の128kbチャンクを1回ロードしますか？ （4*32 = 128）グローバルなメモリ帯域幅を無駄にしていないようですか？

ところで、Compute ProfilerはGLDとGST効率が1.00003であることを示しています。 1.0以上はどうですか？

役に立ちましたか？

解決

はい、あなたのメモリアクセスパターンはかなり最適です。各半計は、16の連続した32ビット語にアクセスしています。さらに、バッファ自体が整列し、各半歩行のstartIndexは16の倍数であるため、アクセスが64バイトに並べられています。したがって、各半計は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];
}

最初は、これがCompute 1.3（GT200）ハードウェアでどのように機能するかをSEにします

これへの書き込みの場合、これはわずかに最適でないパターンを生成します（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からの読み取りの場合、スレッドは配列の要素のみにアクセスするため、半計ごとにすべてのアクセスが128byteアラインドブロックにあります（最初の要素は128b境界にあります。インデックスは32の倍数であり、4バイト要素の場合、アドレスオフセットは128bの倍数であることを意味します。 AccessPatternは128Bブロック全体に広がっているため、これにより、半分の半分ごとに128Bの転送が行われ、帯域幅の半分が再び腰を開きます。

Cからの読み取りは、各スレッドが独自の128Bブロックにインデックスする最悪のシナリオの1つを生成します。そのため、各スレッドは独自の転送が必要です。ハードウェアは転送と重複できるはずなので）。さらに悪いことに、これが各スレッドの32Bブロックを転送し、帯域幅の7/8を無駄にするという事実です（4B/スレッド、32B/4B = 8にアクセスするため、帯域幅の1/8のみが利用されます）。これはナイーブマトリックストランスポーズのアクセスパターンであるため、ローカルメモリを使用して（経験から話す）ことを行うことを非常にお勧めします。

1.0（G80）を計算する

ここで、適切なアクセスを作成する唯一のパターンは元のパターンであり、この例のすべてのパターンは完全に無視されていないアクセスを作成し、帯域幅の7/8を無駄にします（32B転送/スレッド、上記を参照）。 G80ハードウェアの場合、半計のn番目のスレッドがアクセスしないすべてのアクセスがnth要素にアクセスします。

計算2.0（フェルミ）

ここでは、メモリへのすべてのアクセスが128Bのトランザクションを作成します（すべてのデータを収集するために必要な場合、最悪の場合は16x128B）が、それらはキャッシュされているため、データが転送される場所はあまり明らかではありません。現時点では、キャッシュがすべてのデータを保持するのに十分な大きさであり、競合がないと仮定してください。したがって、128Bのキャシュラインごとに最大で1回転送されます。さらに、ハーフウォープのシリアル化された実行を想定してみましょう。そのため、決定論的なキャッシュ職業があります。

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

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

したがって、大きな配列の場合、Aへのアクセスは理論的に帯域幅をほとんど無駄にしません。この例では、現実はもちろんそれほど良くありません。Cへのアクセスはキャッシュをかなりうまく破壊するからです

プロファイラーの場合、1.0を超える効率は、浮遊点の不正確さの結果であると仮定します。

それが役立つことを願っています

ライセンス： CC-BY-SA と 帰属
所属していません StackOverflow
scroll top