共享内存中的银行无冲突访问

Question

我必须使用64个元素大小的共享内存，是warp中库和线程数量的两倍。我应该如何解决这些问题，以获得无冲突的银行访问权限？

Answer 1

在32位内存访问的情况下，您可以使用默认的内存访问模式。

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

则CCD_ 1是奇数。

如果你有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是为了避免转置访问模式中的序列化。