1

合体している場合はn < warpSize？

// In kernel
int x;
if (threadId < n)
  x = globalMem[threadId];

Nこのような状況は、 で割り切れないものがある場合、サイクルの最後の反復で発生しますwarpSize。これらの状況について実行し、割り切れるデバイスメモリを割り当てるwarpSizeか、そのまま合体する必要がありますか?

4

1 に答える 1

1

cuda プログラミング ガイド - スレッド階層threadIdに記載されているように が正しく計算された場合、このアクセスは結合されます - これは.threadId = threadIdx.x

コンピューティング アーキテクチャが異なると、メモリの合体​​が若干異なります。詳細については、cuda プログラミング ガイドの付録 G を参照してください。

一般に、最初のスレッドがアクセスする要素のアドレスから開始して、スレッドがメモリ内の連続する要素を取得する場合、グローバル メモリ アクセスは結合されていると言えます。

float 配列があるとします。
float array[]
あなたのメモリアクセスはそのように見えます

array[threadIdx.x == 0, threadId.x == 1, threadIdx.x == 2, ..., threadIdx.x == 31]

あなたのアクセスが結合されるよりも。

しかし、そのようにメモリにアクセスする場合（たとえばインターリーブ）

array[threadIdx.x == 0, NONE, threadId.x == 1, NONE, threadIdx.x == 2, ..., NONE,  threadIdx.x == 31]

あなたのアクセスが合体していないよりも（NONEこの配列要素がどのスレッドからもアクセスされていないことを意味します）

最初のケースでは、連続する 128 バイトのメモリを取得します。2 番目のケースでは、256 バイトを取得します。2 番目のケースでは、グローバル メモリからメモリをロードするために、最初のケースでは 1 つのワープが必要でしたが、2 つのワープが必要です。ただし、どちらの場合も、次の計算に必要な float 要素は 32 個 (つまり 128 バイト) だけです。したがって、その単純なケースでは、グローバルな負荷率は 1.0 から 0.5 に低下します。

于 2014-02-21T12:28:38.480 に答える