GPU共享内存实例

Posted 2023-04-15

【中文标题】GPU共享内存实例

【英文标题】：GPU shared memory practical example

【发布时间】：2017-08-28 23:39:29

【问题描述】：

我有一个这样的数组：

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

【问题讨论】：

假设您的 blockDim.x 也是 16，在 G80 上，数据大小为 16 不会有任何银行冲突。我很确定论文的作者没有你的例子在视图中。数据大小至少为 32，blockDim.x 至少为 32，不难证明 G80 上的银行冲突是如何产生的。

***.com/q/7903566/681865

我使用的例子和本文使用的例子相同the paper（我使用的例子和本文使用的例子相同）我说的是在第 11 页（您可以在我刚刚在我的问题中添加的图片中看到它。请您演示一下 32 个元素是如何产生银行冲突的？非常感谢@Robert Crovellla

【参考方案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 的事实结构。

【讨论】：

对我花了很多时间理解的问题的解释真的很清楚。非常感谢亲爱的 Rober 先生 @Robert Crovella