8

CUDAプログラミングガイドを読みましたが、1つ見逃しました。グローバルメモリに32ビット整数の配列があり、それを合体アクセスで共有メモリにコピーしたいとします。グローバル配列には0から1024までのインデックスがあり、それぞれ256スレッドの4つのブロックがあるとします。

__shared__ int sData[256];

合体アクセスはいつ実行されますか？

1.1。

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y];

グローバルメモリ内のアドレスは0から255まで、それぞれワープで32スレッドずつコピーされるので、ここで問題ありませんか？

2.2。

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y + someIndex];

someIndexが32の倍数でない場合、それは合体していませんか？住所がずれていますか？あれは正しいですか？

4

4 に答える 4

16

最終的に何が必要かは、入力データが 1D 配列か 2D 配列か、およびグリッドとブロックが 1D か 2D かによって異なります。最も単純なケースは両方とも 1D です。

shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + threadIdx.x];

これは合体。私が使用する経験則は、最も急速に変化する座標 (threadIdx) がオフセットとしてブロック オフセット (blockDim * blockIdx) に追加されるというものです。最終結果は、ブロック内のスレッド間のインデックス作成ストライドが 1 になることです。ストライドが大きくなると、合体が失われます。

簡単なルール (Fermi 以降の GPU) は、ワープ内のすべてのスレッドのアドレスが同じ 128 バイトの範囲に収まる場合、単一のメモリ トランザクションが発生するというものです (ロードに対してキャッシュが有効になっていると仮定すると、これはデフォルト）。それらが 2 つのアラインされた 128 バイト範囲に入る場合、2 つのメモリ トランザクションが発生します。

GT2xx 以前の GPU では、より複雑になります。ただし、その詳細については、プログラミング ガイドを参照してください。

追加の例:

合体していない:

shmem[threadIdx.x] = gmem[blockDim.x + blockIdx.x * threadIdx.x];

合体していませんが、GT200以降ではそれほど悪くはありません:

stride = 2;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

まったく合体していない:

stride = 32;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

合体、2D グリッド、1D ブロック:

int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.x] = gmem[blockIdx.y * elementPitch + 
                          blockIdx.x * blockDim.x + threadIdx.x];

合体、2D グリッドおよびブロック:

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.y * blockDim.x + threadIdx.x] = gmem[y * elementPitch + x];
于 2012-04-26T05:29:46.480 に答える
1

1 でのインデックス作成が間違っています (または、意図的に奇妙で間違っているように見えます)。一部のブロックは各スレッドで同じ要素にアクセスするため、これらのブロックで合体アクセスする方法はありません。

証拠：

例：

Grid = dim(2,2,0)

t(blockIdx.x, blockIdx.y)

//complete block reads at 0
t(0,0) -> sData[threadIdx.x] = gData[0];
//complete block reads at 2
t(0,1) -> sData[threadIdx.x] = gData[2];
//definetly coalesced
t(1,0) -> sData[threadIdx.x] = gData[threadIdx.x];
//not coalesced since 2 is no multiple of a half of the warp size = 16
t(1,1) -> sData[threadIdx.x] = gData[threadIdx.x + 2];

ブロックが合体した場合は「運」のゲームなので、一般的にはいいえ

ただし、結合されたメモリ読み取りルールは、新しい cuda バージョンでは以前ほど厳密ではありません。
ただし、互換性の問題があるため、可能であれば、カーネルを最も低い cuda バージョンに最適化するようにしてください。

ここにいくつかの素晴らしいソースがあります：

http://mc.stanford.edu/cgi-bin/images/0/0a/M02_4.pdf

于 2012-04-26T03:17:53.340 に答える
0

アクセスを合体させることができるルールはやや複雑であり、時間の経過とともに変化しています。それぞれの新しいCUDAアーキテクチャは、合体できるものにおいてより柔軟です。最初は気にしないと思います。代わりに、最も便利な方法でメモリアクセスを実行してから、CUDAプロファイラーの内容を確認してください。

于 2012-04-25T23:53:57.973 に答える
-1

1D グリッドとスレッド ジオメトリを使用するつもりなら、あなたの例は正しいです。使用する予定のインデックスは[blockIdx.x*blockDim.x + threadIdx.x].

＃1では、ワープ内の32個のスレッドがその命令を「同時に」実行するため、シーケンシャルで128B（32 x 4）に整列されたそれらのリクエストは、TeslaとFermiの両方のアーキテクチャで合体されると思います.

＃2では、少しぼやけています。が 1 の場合someIndex、1 つのワープで 32 個のリクエストすべてを合体するわけではありませんが、部分的な合体は行う可能性があります。Fermi デバイスは、スレッド 1 から 31 へのアクセスを 128B のメモリ シーケンシャル セグメントの一部としてワープで結合すると考えています (スレッドが必要としない最初の 4B は無駄になります)。テスラ アーキテクチャ デバイスでは、位置合わせのずれにより非合体アクセスになると思いますが、よくわかりません。

たとえば 8 の場合、Tesla は 32B のアラインされたアドレスをsomeIndex持ち、Fermi はそれらを 32B、64B、および 32B としてグループ化する可能性があります。しかし、肝心なのは、 とアーキテクチャの値によってはsomeIndex、何が起こるかは曖昧であり、必ずしもひどいものではないということです。

于 2012-04-26T03:00:35.087 に答える