cuda - 共有メモリでの Bank-Conflict-Free アクセス

Question

サイズが 64 要素で、32 バンクの 2 倍の共有メモリを使用する必要があります。したがって、メモリアクセスの数がワープ内のスレッド数の 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

コンピューティング機能 1.x を使用していると仮定すると、共有メモリには 16 のバンクがあり、各スレッドは共有メモリ内の 2 つの要素にアクセスする必要があります。

必要なのは、スレッドが両方の要素に対して同じメモリ バンクにアクセスできるようにすることです。そのため、必要な要素が互いに 16 離れているように編成する場合は、バンクの競合を避ける必要があります。

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

複雑なフロートを格納するためにこのパターンを使用しましたが、複雑なフロートの配列があったため、次のようになりました

#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];

+1 は、転置されたアクセス パターンでシリアル化を回避するためのものです。