GPU共享内存实例

Question

我有一个这样的数组：

data[16] = {10,1,8,-1,0,-2,3,5,-2,-3,2,7,0,11,0,2}

我想在 G80 GPU 上使用共享内存计算此数组的缩减量。

NVIDIA 文档中引用的内核是这样的：

__global__ void reduce1(int *g_idata, int *g_odata) {
extern __shared__ int sdata[];

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
sdata[tid] = g_idata[i];
__syncthreads();

// here the reduction :

for (unsigned int s=1; s < blockDim.x; s *= 2) {
int index = 2 * s * tid;
if (index < blockDim.x) {
sdata[index] += sdata[index + s];
}
__syncthreads();
}

论文作者说这种方法存在银行冲突的问题。我试图理解，但我不知道为什么？我知道银行冲突和广播访问的定义，但仍然无法理解。

Bank Conflicts

Answer 1

G80 处理器是一款非常古老的支持 CUDA 的 GPU，属于第一代 CUDA GPU，计算能力为 1.0。最近的 CUDA 版本（6.5 之后）不再支持这些设备，因此在线文档不再包含了解这些设备中的 bank 结构的必要信息。

因此，我将在此处从 CUDA 6.5 C 编程指南中摘录 cc 1.x 设备的必要信息：

G.3.3。共享内存

共享内存有 16 个存储体，这些存储体被组织成连续的 32 位字映射 到后来的银行。每个 bank 的带宽为每两个时钟周期 32 位。

一个warp的共享内存请求被分成两个内存请求，每个请求一个 半经线，独立发行。因此，不可能有银行 属于warp前半部分的线程和属于的线程之间的冲突 同一条经线的后半部分。

在这些设备中，共享内存具有 16 个 bank 结构，因此每个 bank 的“宽度”为 32 位或 4 字节。例如，每个 bank 的宽度与 int 或 float 数量相同。因此，让我们设想一下可能存储在这种共享内存中的前 32 个 4 字节数量，以及它们对应的存储区（使用 f 而不是 sdata 作为数组名称）：

extern __shared__ int f[];

index: f[0] f[1] f[2] f[3] ... f[15] f[16] f[17] f[18] f[19] ... f[31]
bank:    0    1    2    3  ...   15     0     1     2     3  ...   15

共享内存中的前 16 个int 数量属于银行 0 到 15，共享内存中接下来的 16 个int 数量也属于银行 0 到 15（以此类推，如果我们的数据中有更多数据） int 数组）。

现在让我们看看会触发银行冲突的代码行：

for (unsigned int s=1; s < blockDim.x; s *= 2) {
int index = 2 * s * tid;
if (index < blockDim.x) {
sdata[index] += sdata[index + s];
}

让我们考虑第一次通过上述循环，其中s 是1。这意味着index 是2*1*tid，所以对于每个线程，index 只是threadIdx.x 值的两倍：

threadIdx.x: 0 1 2 3 4  5  6  7  8  9 10 11 ...
 index:      0 2 4 6 8 10 12 14 16 18 20 22 ...
 bank:       0 2 4 6 8 10 12 14  0  2  4  6 ...

所以对于这个读取操作：

+= sdata[index + s]

我们有：

threadIdx.x: 0 1 2 3 4  5  6  7  8  9 10 11 ...
 index:      0 2 4 6 8 10 12 14 16 18 20 22 ...
 index + s:  1 3 5 7 9 11 13 15 17 19 21 23 ...
 bank:       1 3 5 7 9 11 13 15  1  3  5  7 ...

所以，在前 16 个线程中，我们有两个线程想要从 bank 1 读取，两个想要从 bank 3 读取，两个想要从 bank 5 读取，等等。因此这个读取周期遇到 2-第一个 16 线程组的方式库冲突。请注意，同一行代码上的其他读取和写入操作类似地存在银行冲突：

sdata[index] +=

因为这将读取然后写入银行 0、2、4 等。每组 16 个线程两次。

可能正在阅读此示例的其他人请注意：正如所写，它仅适用于 cc 1.x 设备。在 cc 2.x 和更新的设备上演示 bank 冲突的方法可能相似，但具体情况有所不同，这是由于 warp 执行差异以及这些新设备具有 32 路 bank 结构而不是 16 路 bank 的事实结构。