通过使用CUDA GPU共享内存

共享内存是编写优化良好的 CUDA 代码的一个强大功能。共享内存的访问比全局内存访问快得多，因为它位于芯片上。因为共享内存由线程块中的线程共享，它为线程提供了一种协作机制。利用这种线程协作使用共享内存的一种方法是启用全局内存合并，如本文中的数组反转所示。通过使用 CUDA GPU 共享内存，我们可以在 GPU 上执行所有读操作。在下一篇文章中，我将通过使用共享内存来优化矩阵转置来继续我们的讨论。

在 上一篇文章 中，我研究了如何将一组线程访问的全局内存合并到一个事务中，以及对齐和跨步如何影响 CUDA 各代硬件的合并。对于最新版本的 CUDA 硬件，未对齐的数据访问不是一个大问题。然而，不管 CUDA 硬件是如何产生的，在全局内存中大步前进都是有问题的，而且在许多情况下似乎是不可避免的，例如在访问多维数组中沿第二个和更高维的元素时。但是，在这种情况下，如果我们使用共享内存，就可以合并内存访问。在我在下一篇文章中向您展示如何避免跨越全局内存之前，首先我需要详细描述一下共享内存。

共享内存

因为它是片上的，共享内存比本地和全局内存快得多。实际上，共享内存延迟大约比未缓存的全局内存延迟低 100 倍（前提是线程之间没有内存冲突，我们将在本文后面讨论这个问题）。共享内存是按线程块分配的，因此块中的所有线程都可以访问同一共享内存。线程可以访问由同一线程块中的其他线程从全局内存加载的共享内存中的数据。此功能（与线程同步结合）有许多用途，例如用户管理的数据缓存、高性能的协作并行算法（例如并行缩减），以及在不可能实现全局内存合并的情况下促进全局内存合并。

线程同步

在线程之间共享数据时，我们需要小心避免争用情况，因为虽然块中的线程并行运行 逻辑上 ，但并非所有线程都可以同时执行 身体上 。假设两个线程 A 和 B 分别从全局内存加载一个数据元素并将其存储到共享内存中。然后，线程 A 想从共享内存中读取 B 的元素，反之亦然。我们假设 A 和 B 是两个不同翘曲中的线。如果 B 在 A 尝试读取它之前还没有完成它的元素的编写，我们就有一个竞争条件，它可能导致未定义的行为和错误的结果。

为了保证并行线程协作时的正确结果，必须同步线程。 CUDA 提供了一个简单的屏障同步原语 __syncthreads() 。一个线程的执行只能在其块中的所有线程都执行了 __syncthreads() 之后通过 __syncthreads() 继续执行。因此，我们可以通过在存储到共享内存之后和从共享内存加载任何线程之前调用 __syncthreads() 来避免上面描述的竞争条件。需要注意的是，在发散代码中调用 __syncthreads() 是未定义的，并且可能导致死锁，线程块中的所有线程都必须在同一点调用 __syncthreads() 。

共享内存示例

使用 Clara 变量 D __shared__ 指定说明符在 CUDA C / C ++设备代码中声明共享内存。在内核中声明共享内存有多种方法，这取决于内存量是在编译时还是在运行时已知的。下面的完整代码（ 在 GitHub 上提供 ）演示了使用共享内存的各种方法。

#include __global__ void staticReverse(int *d, int n)
{ __shared__ int s[64]; int t = threadIdx.x; int tr = n-t-1; s[t] = d[t]; __syncthreads(); d[t] = s[tr];
} __global__ void dynamicReverse(int *d, int n)
{ extern __shared__ int s[]; int t = threadIdx.x; int tr = n-t-1; s[t] = d[t]; __syncthreads(); d[t] = s[tr];
} int main(void)
{ const int n = 64; int a[n], r[n], d[n]; for (int i = 0; i < n; i++) { a[i] = i; r[i] = n-i-1; d[i] = 0; } int *d_d; cudaMalloc(&d_d, n * sizeof(int)); // run version with static shared memory cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice); staticReverse<<<1,n>>>(d_d, n); cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost); for (int i = 0; i < n; i++) if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)n", i, i, d[i], r[i]); // run dynamic shared memory version cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice); dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n); cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost); for (int i = 0; i < n; i++) if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)n", i, i, d[i], r[i]); 

}此代码使用共享内存反转 64 元素数组中的数据。这两个内核非常相似，只是在共享内存数组的声明方式和内核的调用方式上有所不同。

静态共享内存

如果共享内存数组大小在编译时已知，就像在 staticReverse 内核中一样，那么我们可以显式地声明一个该大小的数组，就像我们对数组 s 所做的那样。

__global__ void staticReverse(int *d, int n)
{ __shared__ int s[64]; int t = threadIdx.x; int tr = n-t-1; s[t] = d[t]; __syncthreads(); d[t] = s[tr]; 

}在这个内核中， t 和 tr 是分别表示原始顺序和反向顺序的两个索引。线程使用语句 s[t] = d[t] 将数据从全局内存复制到共享内存，然后在两行之后使用语句 d[t] = s[tr] 完成反转。但是在执行最后一行之前，每个线程访问共享内存中由另一个线程写入的数据，请记住，我们需要通过调用 __syncthreads() 来确保所有线程都已完成对共享内存的加载。

在这个例子中使用共享内存的原因是为了在旧的 CUDA 设备（计算能力 1 . 1 或更早版本）上促进全局内存合并。由于全局内存总是通过线性对齐索引 t 访问，所以读写都可以实现最佳的全局内存合并。反向索引 tr 仅用于访问共享内存，它不具有全局内存的顺序访问限制以获得最佳性能。共享内存的唯一性能问题是银行冲突，我们将在后面讨论。（请注意，在计算能力为 1 . 2 或更高版本的设备上，内存系统甚至可以将反向索引存储完全合并到全局内存中。但是这种技术对于其他访问模式仍然有用，我将在下一篇文章中展示。）

动态共享内存

本例中的其他三个内核使用动态分配的共享内存，当编译时共享内存的数量未知时，可以使用该内存。在这种情况下，必须使用可选的第三个执行配置参数指定每个线程块的共享内存分配大小（以字节为单位），如下面的摘录所示。

dynamicReverse<<<1, n, n*sizeof(int)>>>(d_d, n);

动态共享内存内核 dynamicReverse() 使用未大小化的外部数组语法 extern shared int s[] 声明共享内存数组（注意空括号和 extern 说明符的使用）。大小在内核启动时由第三个执行配置参数隐式确定。内核代码的其余部分与 staticReverse() 内核相同。

如果在一个内核中需要多个动态大小的数组怎么办？您必须像前面一样声明一个 extern 非大小数组，并使用指向它的指针将其划分为多个数组，如下面的摘录所示。

extern __shared__ int s[];
int *integerData = s; // nI ints
float *floatData = (float*)&integerData[nI]; // nF floats
char *charData = (char*)&floatData[nF]; // nC chars

在内核中指定启动所需的总内存。

myKernel<<>>(...);

共享内存库冲突

为了实现并发访问的高内存带宽，共享内存被分成大小相等的内存模块（库），这些模块可以同时访问。因此，任何跨越 b 不同内存组的 n 地址的内存负载或存储都可以同时进行服务，从而产生的有效带宽是单个存储库带宽的 b 倍。

但是，如果多个线程的请求地址映射到同一个内存库，则访问将被序列化。硬件根据需要将冲突内存请求拆分为多个独立的无冲突请求，将有效带宽减少一个与冲突内存请求数量相等的因子。一个例外情况是，一个 warp 中的所有线程都使用同一个共享内存地址，从而导致广播。计算能力 2 . 0 及更高版本的设备具有多播共享内存访问的额外能力，这意味着在一个 warp 中通过任意数量的线程对同一个位置的多个访问同时进行。

为了最小化内存冲突，了解内存地址如何映射到内存库是很重要的。共享存储库被组织成这样，连续的 32 位字被分配给连续的存储库，带宽是每个库每个时钟周期 32 位。对于计算能力为 1 . x 的设备， warp 大小为 32 个线程，库的数量为 16 个。一个 warp 的共享内存请求被分为一个对 warp 前半部分的请求和一个对 warp 后半部分的请求。请注意，如果每个内存库只有一个内存位置被半个线程访问，则不会发生库冲突。

对于计算能力为 2 . 0 的设备， warp 大小是 32 个线程，而 bank 的数量也是 32 个。 warp 的共享内存请求不会像计算能力为 1 . x 的设备那样被拆分，这意味着 warp 前半部分的线程和同一 warp 后半部分的线程之间可能会发生库冲突。

计算能力为 3 . x 的设备具有可配置的存储大小，可以使用 CUDA Devicsetsharedmeconfig() 将其设置为四个字节（ CUDA SharedMemBankSizeFourByte ，默认值）或八个字节（ cudaSharedMemBankSizeEightByte) 。将存储大小设置为 8 字节有助于避免访问双精度数据时的共享内存库冲突。

配置共享内存量

在计算能力为 2 . x 和 3 . x 的设备上，每个多处理器都有 64KB 的片上内存，可以在一级缓存和共享内存之间进行分区。对于计算能力为 2 . x 的设备，有两个设置： 48KB 共享内存/ 16KB 一级缓存和 16KB 共享内存/ 48KB 一级缓存。默认情况下，使用 48KB 共享内存设置。这可以在运行时 API 期间使用 cudaDeviceSetCacheConfig() 为所有内核配置，也可以使用 cudaFuncSetCacheConfig() 在每个内核的基础上进行配置。它们接受以下三个选项之一： cudaFuncCachePreferNone 、 cudaFuncCachePreferShared 和 cudaFuncCachePreferL1 。驱动程序将遵循指定的首选项，除非内核每个线程块需要比指定配置中可用的共享内存更多的共享内存。计算能力为 3 . x 的设备允许使用选项 cudaFuncCachePreferEqual 获得 32KB 共享内存/ 32kbl1 缓存的第三个设置。

　　关于作者

　　Mark Harris 是 NVIDIA 杰出的工程师，致力于 RAPIDS 。 Mark 拥有超过 20 年的 GPUs 软件开发经验，从图形和游戏到基于物理的模拟，到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时，他意识到了一种新生的趋势，并为此创造了一个名字： GPGPU （图形处理单元上的通用计算）。

　　审核编辑：郭婷