2

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

全球记忆

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

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

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

共享内存

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

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

4

1 回答 1

4

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

合并的内存访问

一维数组 - 一维线程网格

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

2D 数组 - 2D 线程网格

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]

共享内存库冲突

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

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

快速访问

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

访问速度慢

  • 32线程访问32同一个bank中的不同单词，使得所有的访问都被序列化；
  • 一般来说，访问共享内存的成本与同时访问单个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 访问中的第三个线程和 warp 访问myShMem[50]中的八个线程myShMem[178]，那么您将发生双向银行冲突，并且两个事务被序列化。

示例 2

考虑以下类型的访问

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

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

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];

不会导致银行冲突，因为访问了一个字节/线程，因此访问了同一个字的不同字节。

于 2014-08-14T21:35:08.457 回答