CUDAがグローバルメモリへの一貫したアクセス

Question

私はCUDAプログラミングガイドを読んだことがありますが、私は1つのことを忘れました。私はグローバルメモリに32ビットintの配列を持っていると私は共有メモリに合体アクセスでそれをコピーしたいとしましょう。 グローバル配列は0から1024までのインデックスを持ち、それぞれ256個のスレッドで4つのブロックがあるとします。

__shared__ int sData[256]; 

融合アクセスはいつ実行されますか。

1.

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

グローバルメモリ内なアドレスはは縦糸で32のスレッドでそれぞれ、0から255にコピーされたので、ここでそれは大丈夫ですされていますか？

2.

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

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

Answer 1

アクセスを統合できるルールは、やや複雑で、時間とともに変化しています。それぞれの新しいCUDAアーキテクチャは、融合可能な点でより柔軟性があります。私は最初にそれについて心配しないと言います。代わりに、どのような方法でもメモリアクセスを最も便利にしてから、CUDAプロファイラが何を言うのかを確認してください。

Answer 2

1Dグリッドとスレッドジオメトリを使用する場合は、正しい例です。あなたが使用しようとしているインデックスは[blockIdx.x*blockDim.x + threadIdx.x]だと思います。

＃1では、ワープ中の32個のスレッドが '同時に'命令を実行するため、128B（32x4）に順次整列された要求がテスラとフェルミの両方のアーキテクチャで統合されます。

＃2の場合、少しぼやけています。 someIndexが1の場合、ワープ内の32個の要求のすべてを合併することはありませんが、部分的な合体が行われる可能性があります。フェルミのデバイスは、メモリの128Bシーケンシャルセグメントの一部としてワープ内のスレッド1〜31のアクセスを統合すると考えています（最初の4Bはスレッドが不要で、無駄になります）。私はテスラアーキテクチャのデバイスは、ミスアライメントのためにそれをuncoalescedアクセスすると思いますが、私は確信していません。

someIndexとテスラのアドレスは32Bで、フェルミはそれらを32B、64B、および32Bとしてグループ化することがあります。しかし、結論は、someIndexの値とアーキテクチャに応じて、何が起こるかはぼやけており、必ずしもひどいものではありません。

Answer 3

インデックスが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]; 

だから、そのA "運" のゲームブロックが合体した場合、一般的に何

しかし、合体したメモリはルールがある読まないので、これまでのように新しいcudaバージョンでは厳しくはありません。
しかし、互換性の問題では、可能な場合は、最小のcudaバージョン用にカーネルを最適化するようにしてください。あなたが最終的にあなたの入力データが1Dまたは2D配列で、あなたのグリッドとブロックは1Dまたは2Dされているかどうかかどうかに依存して欲しい

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

Answer 4

：ここ

はいくつかの素晴らしい源です。最も簡単なケースは1D：

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

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

単純ルール（Fermi以降のGPU）では、ワープ内のすべてのスレッドのアドレスが同じ整列した128バイトの範囲に入ると、1回のメモリトランザクションが発生しますこれはデフォルトです）。 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]; 

Coa lesced、2次元グリッド、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];