1

我有一个这样的数组：

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();
}

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

银行冲突

4

1 回答 1

2

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

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

G.3.3。共享内存

共享内存有 16 个存储区，这些存储区的组织方式使得连续的 32 位字映射到连续的存储区。每个存储体的带宽为每两个时钟周期 32 位。

一个warp的共享内存请求被分成两个内存请求，一个用于每个半warp，它们是独立发出的。因此，属于 warp 的前半部分的线程和属于同一 warp 的后半部分的线程之间不会发生存储库冲突。

在这些设备中，共享内存具有 16 个存储体结构，因此每个存储体具有 32 位或 4 字节的“宽度”。int例如，每个银行具有与一个或float数量相同的宽度。因此，让我们设想一下可能存储在这种共享内存中的前 32 个 4 字节数量，以及它们对应的 bank（使用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数量属于 bank 0 到 15，int共享内存中接下来的 16 个数量也属于 bank 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。这意味着indexis 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-way bank 冲突跨第一个 16 线程组。请注意，同一行代码上的其他读取和写入操作类似地存在银行冲突：

sdata[index] +=

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

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

于 2017-04-04T15:29:22.783 回答