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冲突的问题。我试图理解，但我无法弄清楚为什么？我知道bank conflict和broadcast access的定义，但是还是看不懂。

Bank Conflicts

Answer 1

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

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

G.3.3. Shared Memory

Shared memory has 16 banks that are organized such that successive 32-bit words map to successive banks. Each bank has a bandwidth of 32 bits per two clock cycles.

A shared memory request for a warp is split into two memory requests, one for each half-warp, that are issued independently. As a consequence, there can be no bank conflict between a thread belonging to the first half of a warp and a thread belonging to the second half of the same warp.

在这些设备中，共享内存具有 16 组结构，因此每个组都有 "width" 32 位或 4 字节。例如，每个银行的宽度与 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 个数量属于 bank 0 到 15，共享内存中接下来的 16 int 个数量也属于 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。这意味着 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 读取，等等。因此这个读取周期遇到跨第一个 16 线程组的 2-way bank 冲突。请注意，同一行代码上的其他读写操作同样存在bank-conflicted：

sdata[index] +=

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

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

GPU共享内存实例

GPU shared memory practical example

cuda

bank-conflict

gpu-shared-memory