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

Posted on | In | | Visitors:

本文翻译自NVIDIA官方博客Parallel Forall，内容仅供参考，如有疑问请访问原网站:https://devblogs.nvidia.com/parallelforall/using-shared-memory-cuda-cc/。

在以前发布的文章中，我们学习了被一组线程访问的全局内存如何被合并为一次事务以及对于不同的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上下载)展示了使用共享内存的两种方法。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
#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数组：

1
2
3
4
5
6
7
8
9
__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数组，然后使用指针将它分为多个数组，如下所示：

1
2
3
4
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上实现完全地合并访问。