CUDA最佳内存访问布局：全局内存合并和共享内存库冲突

Question

我对 CUDA 中最方便的全局和共享内存访问布局有一些疑问。

全局记忆

1) 以下内存地址(0,0)、(0,1)、(1,0)和(1,1)在CPU内存和GPU内存中是如何排列的？换句话说，它们的存储顺序是什么？

2) (m, n) 中哪个是行索引，哪个是列索引？

3) 全局内存合并是通过按列主序还是行主序访问元素实现的？

共享内存

1) 银行冲突如何产生或不产生？请使用示例/案例告诉我。

2) 在64K 中配置共享内存和 L1 的命令是什么？该命令在哪里？

Answer 1

上面的 cmets 已经回答了您的大部分问题。我只是想提供一些对您以及对下一个用户有用的规则，这些规则涉及合并内存访问、一些关于共享内存库冲突的示例以及一些关于避免共享内存库冲突的规则.

联合内存访问

一维数组 - 一维线程网格

gmem[blockDim.x * blockIdx.x + threadIdx.x]

二维数组 - 二维线程网格

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int elementPitch = blockDim.x * gridDim.x;
gmem[y][x] or gmem[y * elementPitch + x]

共享内存库冲突

为了实现高带宽，共享内存被分成独立的银行。通过这种方式，共享内存可以为线程同时访问提供服务。每个流式多处理器 (SM) 都在32 内存库中组织了共享内存。每个存储体的带宽为每两个时钟周期 32 位，并承载四个字节的字（32 位）：连续的32-位字地址分配给连续的存储体。

bank 冲突发生在两个不同的线程访问同一 bank 中的 不同字时。银行冲突会对性能产生不利影响，因为它们强制硬件对共享内存的访问进行序列化。请注意，如果不同的线程访问同一个字中的任何字节，则不会发生冲突。另请注意，之间没有银行冲突 属于不同经线的线程。

快速访问

• 如果一个 warp 的所有线程都访问不同的 bank，则不存在 bank 冲突；
• 如果一个 warp 的所有线程都访问一个相同的地址以进行 fetch 操作，则没有 银行冲突（广播）。

访问缓慢

• 32线程访问32同一个bank中的不同单词，使所有的访问序列化；
• 一般来说，访问共享内存的成本与同时访问单个存储区的最大数量成正比。

示例 1

smem[4]:   accesses bank #4  (physically, the fifth one – first row)

smem[31]:  accesses bank #31 (physically, the last one  – first row)

smem[50]:  accesses bank #18 (physically, the 19th one  – second row)

smem[128]: accesses bank #0  (physically, the first one – fifth row)

smem[178]: accesses bank #18 (physically, the 19th one  – sixth row)

如果warp 中的第三个线程访问myShMem[50]，而warp 中的第八个线程访问myShMem[178]，那么就会发生双向银行冲突，并且两个事务会被序列化。

示例 2

考虑以下类型的访问

__shared__ float smem[256];
smem[b + s * threadIdx.x]

要在同一warp的两个线程t1和t2之间发生bank冲突，必须满足以下条件

b + s * t2 = b + s * t1 + 32 * k, with k positive integer
0 <= t2 - t1 < 32

以上意思

32 * k = s * (t2 - t1)
0 <= t2 - t1 < 32

如果s 是奇数，这两个条件不成立，即没有银行冲突。

示例 3

从示例2，如下访问

smem[b + threadIdx.x]

如果smem 是32-bits 数据类型，则不会导致冲突。但也

extern __shared__ char smem[];
foo = smem[baseIndex + threadIdx.x];

和

extern __shared__ short smem[];
foo = smem[baseIndex + threadIdx.x];

不会导致存储库冲突，因为访问了一个字节/线程，因此访问了同一单词的不同个字节。