2D共有メモリをCUDAに配置する方法

Question

私はいつもリニア共有メモリ（ロード、ストア、アクセスネイバー）を使っていましたが、バンクコンフリクトを調べるために2Dで簡単なテストを行いました。

次のコードは、1次元グローバルメモリアレイから共有メモリにデータを読み込み、共有メモリからグローバルメモリにコピーバックします。

__global__ void update(int* gIn, int* gOut, int w) { 

// shared memory space 
__shared__ int shData[16][16]; 
// map from threadIdx/BlockIdx to data position 
int x = threadIdx.x + blockIdx.x * blockDim.x; 
int y = threadIdx.y + blockIdx.y * blockDim.y; 
// calculate the global id into the one dimensional array 
int gid = x + y * w; 

// load shared memory 
shData[threadIdx.x][threadIdx.y] = gIn[gid]; 
// synchronize threads not really needed but keep it for convenience 
__syncthreads(); 
// write data back to global memory 
gOut[gid] = shData[threadIdx.x][threadIdx.y]; 
} 

ビジュアルプロファイラは、共有メモリで競合を報告しました。次のコードは競合を回避します（相違点のみを表示します）

// load shared memory 
shData[threadIdx.y][threadIdx.x] = gIn[gid]; 

// write data back to global memory 
gOut[gid] = shData[threadIdx.y][threadIdx.x]; 

この動作は、超並列プロセッサのプログラミングでは私を混乱させます。我々が読むことができるハンズオンアプローチ：CとCUDAの行列要素は、行の主要な規則に従って線形にアドレス指定された位置に配置されます：

すなわち、行列の行0の要素は、最初に連続する位置に順番に配置される。

これは共有メモリアセンブリに関連していますか？またはスレッドインデックスを使用していますか？多分私は何かを逃していますか？

カーネルの構成は以下の通りです：事前に

// kernel configuration 
dim3 dimBlock = dim3 (16, 16, 1); 
dim3 dimGrid = dim3 (64, 64); 
// Launching a grid of 64x64 blocks with 16x16 threads -> 1048576 threads 
update<<<dimGrid, dimBlock>>>(d_input, d_output, 1024); 

感謝。

Answer 1

はい、共有メモリは、期待通りに行優先順位に並べられます。 16の32ビット共有メモリバンクがプリフェルミハードウェアに存在するため、各列内のすべての整数のエントリが共有1にマップ

 bank0 .... bank15 
row 0 [ 0 .... 15 ] 
    1 [ 16 .... 31 ] 
    2 [ 32 .... 47 ] 
    3 [ 48 .... 63 ] 
    4 [ 64 .... 79 ] 
    5 [ 80 .... 95 ] 
    6 [ 96 .... 111 ] 
    7 [ 112 .... 127 ] 
    8 [ 128 .... 143 ] 
    9 [ 144 .... 159 ] 
    10 [ 160 .... 175 ] 
    11 [ 176 .... 191 ] 
    12 [ 192 .... 207 ] 
    13 [ 208 .... 223 ] 
    14 [ 224 .... 239 ] 
    15 [ 240 .... 255 ] 
     col 0 .... col 15 

：だからあなたの[16] [16]配列は、行が賢明のようなものに保存されていますメモリバンク。では、どのようにインデックス作成スキームを選択したのと相互作用しますか？

ブロック内のスレッドは、列のメジャー順と同じに番号が付けられます（技術的には、構造のxディメンションが最も速く変化し、次にy、zが続きます）。あなたは、このインデックス・スキームを使用するときに：半ワープ内

shData[threadIdx.x][threadIdx.y] 

のスレッドが同じ共有メモリバンクから読み込む暗示同じ列からの読み取りとなり、バンク競合が発生します。あなたが逆の方式を使用する場合：同じ半ワープ内

shData[threadIdx.y][threadIdx.x] 

スレッドは16個の異なる共有メモリバンクのそれぞれからの読み出し暗示同じ行からの読み出しであろうが、コンフリクトが発生しません。