[译]在CUDA C/C++中使用共享存储器

本文翻译自NVIDIA官方博客Parallel Forall，内容仅供参考，若有疑问请访问原网站:https://devblogs.nvidia.com/p...。

在之前发布的文章中，咱们学习了被一组线程访问的全局内存如何被合并为一次事务以及对于不一样的CUDA硬件，对齐和步长如何影响合并访问。对于最近的CUDA硬件，没有对齐的数据访问并非什么大问题。然而不管是哪一代的CUDA硬件，跨越全局存储器都是个大问题，并且在不少状况下也是很难避免的，例如沿着第二和更高维度访问多维阵列中的元素时。可是，若是咱们使用共享存储器的话，也是有可能进行合并访问的。在我向你说明如何避免直接跨越全局存储器以前，我首先须要详细地介绍一下共享存储器。

共享存储器

由于它是一个片上存储器，因此共享存储器比本地存储器和全局存储器要快得多。实际上共享存储器的延迟大约比没有缓存的全局存储器低100倍(假设线程之间没有bank冲突，在以后的文章中咱们会介绍)。共享存储器被分配给每一个线程块，因此块内的线程能够访问同一个共享存储器。线程能够访问共享内存中由同一线程块中的其余线程从全局内存加载的数据。这种能力（与线程同步相结合）具备许多用途，例如用户管理的数据高速缓存，高性能并行协做算法（例如并行归约），而且在其它状况不可能的状况下促进全局存储器的合并访问 。

线程同步

当在线程之间共享数据时，咱们须要当心以免竞态条件(race conditions)，由于线程块中的线程之间虽然逻辑上是并行的，可是物理上并非同时执行的。让咱们假设线程A和线程B分别从全局存储器中加载了一个数据而且将它存到了共享存储器。而后，线程A想要从共享存储器中读取B的数据，反之亦然。咱们还要假设线程A和B位于不一样的warp。若是在A尝试读取B的数据时，B还未写入，这样就会致使未定义的行为和错误的结果。

为了保证在并行线程协做时获得正确的结果，咱们必须对线程进行同步。CUDA提供了一个简单的栅栏同步原语，__syncthreads()。每一个线程只能在块中全部的线程执行完__syncthreads()函数后，才能继续执行__syncthreads()的语句。所以咱们能够在向共享存储器存数据后以及在向共享存储器加载数据前调用__syncthreads()，这样就避免了上面所描述的竞态条件(race conditions)。咱们必需要牢记__syncthreads()被用在分支代码块中是未定义的行为，极可能会致使死锁——线程块中全部的线程必须在同一点调用__syncthreads()

共享内存的例子

在设备代码中声明共享内存要使用__shared__变量声明说明符。在核函数中有多种方式声明共享内存，这取决于你要申请的内存大小是在编译时肯定仍是在运行时肯定。下面完整的代码(能够在Github上下载)展现了使用共享内存的两种方法。

#include <stdio.h>

__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仅用于访问共享存储器，其不具备全局存储器的顺序访问限制，所以不能得到最佳性能。共享内存的惟一性能问题是bank冲突，咱们以后会作讨论。

NOTE:注意在计算能力为1.2或更高版本的设备上，内存系统仍然能够彻底地合并访问，即便是反向的保存在全局存储器中。这一技术在其余访问模式下也是颇有用的，我会在下一篇博客中介绍。

动态的共享内存

另外一个核函数使用了动态分配共享内存的方式，这主要用于共享内存的大小在编译时不能肯定的状况。在这种状况下，每一个线程块中共享内存的大小必须在核函数第三个执行配置参数中指定(以字节为单位)，以下所示：

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

该动态共享内存的核函数dynamicReverse()使用了未指定大小的extern数组语法(extern __shared__ int s[])来声明共享内存数组。

NOTE:注意中括号与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<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);

共享内存的bank冲突

为了得到较高的内存带宽，共享存储器被划分为多个大小相等的存储器模块，称为bank，能够被同时访问。所以任何跨越b个不一样bank的n个地址的读写操做能够被同时进行，这样就大大提升了总体带宽 ——可达到单独一个bank带宽的b倍。

然而，若是多个线程请求的地址映射到相同的内存bank，那么访问就会被顺序执行。硬件会把冲突的内存请求分为尽量多的单独的没有冲突的请求，这样就会减小必定的带宽，减小的因子与冲突的内存请求个数相等。固然，也有例外的状况：当一个warp中的全部线程访问同一个共享内存地址时，就会产生一次广播。计算能力为2.0及以上的设备还能够屡次广播共享内存访问，这意味着一个warp中任意数量的线程对于同一位置的屡次访问也能够同时进行。

译者注:这里关于warp的多播与bank冲突原文并未详细介绍，详细内容及例子能够参考CUDA programming guide。我在后续的博客中也会详细介绍这部分。

为了尽可能减小bank冲突，理解共享内存地址如何映射到bank是很是重要的。共享内存的bank是这样组织的：连续的32-bits字被分配到连续的bank中，每一个bank的带宽是每一个时钟周期32bits。

译者注:这里不一样计算能力的bank的带宽是不一样的，原文提到的带宽大小是计算能力5.0的设备，对于计算能力2.0的设备每一个bank的带宽是每两个时钟周期32bits；对于计算能力3.0的设备，每一个bank的带宽是每一个时钟周期64bits。详情请参考CUDA C programming guide。

对于计算能力1.x的设备，warp的大小是32而bank的数量是16。一个warp中线程对共享内存的请求被划分为两次请求：一个请求是前半个warp的另外一个请求时后半个warp的。注意若是每一个bank中只有一个内存地址是被半个warp中的线程访问的话，是不会有bank冲突的。

对于计算能力为2.x的设备，warp的大小是32而bank的数量也是32。一个warp中线程对共享内存的请求不会像计算能力1.x的设备那样被划分开，这就意味着同一个warp中的前半个warp中的线程与后半个warp中的线程会有可能产生bank冲突的。

计算能力为3.x的设备的bank大小是能够配置的，咱们能够经过函数cudaDeviceSetSharedMemConfig()来设置，要么设置为4字节(默认为cudaSharedMemBankSizeFourByte)，要么设置为8字节(cudaSharedMemBankSizeEightByte)。当访问双精度的数据时，将bank大小设置为8字节能够帮助避免bank冲突。

配置共享内存的数量

在计算能力为2.x和3.x的设备上，每一个多处理器有64KB的片上内存，它能够被划分为L1高速缓存和共享内存。对于计算能力为2.x的设备，总共有两种设置：48KB的共享内存/16KBL1高速缓存和16KB的共享内存/16KB的L1高速缓存。咱们能够在运行时使用cudaDeviceSetCacheConfig()在主机端为全部的核函数配置或者使用cudaFuncSetCacheConfig()为单个的核函数配置。它们有三个选项能够设置：cudaFuncCachePreferNone(在共享内存和L1中不设置首选项，即便用默认设置), cudaFuncCachePreferShared(共享内存大于L1), 和cudaFuncCachePreferL1(L1大于共享内存)。驱动程序将按照指定的首选项，除非核函数中每一个线程块须要比指定配置中更多的共享内存。在计算能力3.x的设备上容许有第三种设置选项——32KB的共享内存/32KB的L1高速缓存，能够经过cudaFuncCachePreferEqual选项设置。

总结

对于写出高性能的CUDA代码，共享内存的确是一个十分强大的特性。因为共享内存位于片上，因此访问共享内存比访问全局内存快不少。因为共享内存在线程块中能够被线程共享，因此才提供了相应的机制来保证线程的正常协做。使用共享内存来利用这种线程协做的一种方法是启用全局内存的合并访问，正如如本文中的数组逆序所演示的。在使用共享内存来使数组逆序的例子中，咱们可使用单位步长执行全部全局内存读取和写入，从而在任何CUDA GPU上实现彻底地合并访问。