CUDA - 在CUDA C/C++中使用共享内存

原文链接：Using Shared Memory in CUDA C/C++

---

在上一篇文章中，我研究了如何将一组线程的全局内存访问合并为一个事务，以及对齐和步长如何影响不同代CUDA硬件的合并操作。对于最近版本的CUDA硬件来说，数据访问不一致并不是什么大问题。然而，无论CUDA硬件的版本如何，跨步访问（striding through）全局内存都是存在问题的，并且在许多情况下似乎是不可避免的，例如当沿着第二以及更高维度访问多维数组中的元素时。然而，在这种情况下如果我们使用共享内存，那么内存访问可以合并。在下一篇文章向您展示如何避免跨步访问全局内存之前，首先我需要详细介绍一下共享内存。

共享内存

共享内存因为是片上内存，所以它比本地（local）和全局（global）内存快得多。事实上，共享内存延迟大约是未缓存的全局内存延迟的1/100（前提是线程之间没有bank冲突（bank conflict），我们将在本文稍后进行研究）。共享内存是按线程块分配的，因此一个块中的所有线程可以访问相同的共享内存。线程可以访问由同一线程块内的其他线程从全局内存加载到共享内存中的数据。这种能力（与线程同步相结合）有很多用途，例如用户管理的数据缓存、高性能协作并行算法（例如parallel reductions），以及在使用其他方法无法实现全局内存合并的情况下合并全局内存。

线程同步

当在线程之间共享数据时，我们需要小心避免竞争状态（race conditions），因为虽然块中的线程在逻辑上是并行运行的，但并非所有线程都可以在物理上同时执行。假设两个线程A和B各自从全局内存加载一个数据元素，并将其存储到共享内存中。然后，线程A希望从共享内存中读取B的元素，反之亦然。让我们假设A和B是两个不同warp中的线程。如果B在A尝试读取它的元素之前还没有完成元素的写入，那么就产生了竞争状态，这可能导致未定义的行为和不正确的结果。

为了确保并行线程协作时得到正确的结果，我们必须同步线程。CUDA提供了一个简单的隔离同步原语（barrier synchronization primitive）__syncthreads()。线程只有在其所属的块中的所有线程都执行了__syncthreads()之后才能执行__syncthreads()后面的内容。因此，我们可以通过在存储到共享内存之后和任何线程从共享内存加载之前调用__syncthreads()来避免上述竞争状态。需要注意的是，在有分支结构的代码中调用__syncthreads()是未定义的，可能会导致死锁——线程块中的所有线程都必须在同一点调用__syncthreads()。

共享内存示例

可以使用__shared__变量声明说明符在CUDA C/C++设备代码中声明共享内存。有多种方法可以在内核中声明共享内存，这取决于内存大小在编译时还是在运行时已知。以下代码（可以在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冲突，我们稍后将对此进行讨论。（请注意，在计算能力1.2或更高版本的设备上，内存系统甚至可以将反向索引存储完全合并到全局内存中。但正如我将在下一篇文章中展示的那样，共享内存技术对其他访问模式仍然有用。）

动态共享内存

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

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

动态共享内存内核dynamicReverse()使用未定义大小的extern数组语法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<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);

共享内存bank冲突

为了实现并发访问的高内存带宽，共享内存被划分为可以同时访问的大小相等的内存模块（banks）。因此，跨越b个不同内存bank的n个地址的任何内存加载或存储都可以同时提供服务，从而使有效带宽是单个bank带宽的b倍。

但是，如果多个线程请求的地址映射到同一个内存bank，则访问将被序列化。硬件将冲突内存请求根据需要拆分为多个独立无冲突请求，从而将有效带宽减少为1/n，n为冲突内存请求数。一个例外是warp中的所有线程都寻址同一个共享内存地址，这会导致一次广播（broadcast）。计算能力2.0及更高版本的设备还具有多播（multicast）共享内存访问的能力，这意味着一个warp内任意数量的线程对同一位置的多路访问可以同时进行。

为了最大限度地减少bank冲突，了解内存地址如何映射到内存bank是很重要的。共享内存bank的组织方式为使得连续的32位word被分配给连续的bank，带宽是32位每时钟周期每bank。对于计算能力为1.x的设备，warp尺寸为32个线程，bank数为16。一个warp的共享内存请求被分为一个warp的前半部分的请求和一个warp的后半部分的请求。请注意，如果半个warp的线程每个bank只访问一个内存位置，则不会发生bank冲突。

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

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

配置共享内存数量

在计算能力为2.x和3.x的设备上，每个multiprocessor都有64KB的片上内存，这些片上内存可以在L1 cache和共享内存之间进行分配。对于计算能力为2.x的设备，有两种设置，48KB共享内存/16KB L1 cache（默认）和16KB共享内存/48KB L1 cache。可以在主机使用cudaDeviceSetCacheConfig()为所有内核设置分配方式，也可以使用cudaFuncSetCacheConfig()针对每个内核设置。它们接受以下三个选项之一：cudaFuncCachePreferNone、cudaFuncCachePreferShared和cudaFuncCachePreferL1。驱动程序将遵守指定的选项，除非内核每个线程块需要的共享内存比指定配置中可用的内存多。计算能力3.x的设备允许使用选项cudaFuncCachePreferEqual做第三种设置：32KB共享内存/32KB L1 cache。

总结

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