共有メモリでのバンクコンフリクトフリーアクセス

Question

32バイトの2倍のサイズの64要素の共有メモリを使用する必要があります。従って、メモリアクセスの数は、ワープ内のスレッドの数の2倍である。銀行とのコンフリクトのないアクセスを得るためにはどうすればよいですか？

Answer 1

32ビットメモリアクセスの場合、デフォルトのメモリアクセスパターンを使用できます。

__shared__ int shared[32]; 
int data = shared[base + stride * tid]; 

ここにstrideが奇数である。

あなたは64ビットアクセスを持っている場合は、このようないくつかのトリックを使用することができます。

struct type 
{ 
    int x, y, z; 
}; 
__shared__ struct type shared[32]; 
struct type data = shared[base + tid]; 

Answer 2

のは、ので、あなたの共有メモリが16のバンクがあり、各スレッドが持っている、あなたは計算能力1.xのを使用していると仮定しましょう共有メモリ内の2つの要素にアクセスします。

スレッドは、両方の要素の同じメモリバンクにアクセスする必要があります。そのため、必要な要素が互いに16離れているように整理すると、バンクの競合は避けてください。

__shared__ int shared[32]; 
int data = shared[base + stride * tid]; 
int data = shared[base + stride * tid + 16]; 

私は複雑な山車を格納するために、このパターンを使用しますが、1は、転置アクセスパターンにシリアル化を避けるためである場合は、それは

#define TILE_WIDTH 16 

__shared__ float shared[TILE_WIDTH][2*TILE_WIDTH + 1]; 
float real = shared[base + stride * tid]; 
float imag = shared[base + stride * tid + TILE_WIDTH]; 

のように見えたので、私は、複雑なfloatの配列を持っていました。