CUDA之静态、动态共享内存分配详解

在设备代码中声明共享内存要使用__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)>>>(...);

https://segmentfault.com/a/1190000007526655