Using_Shared_Memory

Using Shared Memory in CUDA C/C++

2020 Jan 29th CookieLau

Source: https://devblogs.nvidia.com/using-shared-memory-cuda-cc/

[toc]

Review

有必要再复习一下各种 memory 了：

其中 local memory 是当 thread 中的 register 用完的时候使用的，其效率和 global 一样都会最慢的。

shared memory 是给同一个 thread block 中的所有 threads 共享。

复习 blog1 复习 blog2

Thread Synchronization

介绍一个新的细粒度的同步： syncthread()，其出现是为了解决以下场景： 在使用 shared memory 前先要将所需数据从 global memory 搬运到片上，若这时候有两个 warps 在同时搬运，虽然 warps 在逻辑上是并行的，但实际上还是有时间上的先后的区别，这是如果 thread1 在thread2 还没有搬运完成的时候想去访问 thread2 正在搬运的数据，则会出现脏数据的情况。所以这个时候需要使用 syncthread() 来使得在执行同一件事的所有的 threads 进行同步等待。

Example

#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]);
}

上面的两个核函数 staticReverse() 和 dynamicReverse() 都是将数组逆序，长得几乎一模一样，唯一的区别在于调用核函数的时候的参数不同： staticReverse<<>> 和 dynamicReverse<<>>

运行得到的结果：

static shared memory 0.008128
dynamic shared memory 0.006144

static shared memory 0.008256
dynamic shared memory 0.006752

static shared memory 0.008064
dynamic shared memory 0.006464

static shared memory 0.008032
dynamic shared memory 0.006432

可以看出两者的速度几乎一样，差别在于运行的先后而已。

Static Shared Memory

当所需要的共享内存的大小在编译前就可知，我们可以采用硬编码的方式直接分配共享内存，如上面的 staticReverse() 中的 s[64]

Dynamic Shared Memory

当所需要的共享内存是我们一开始不知道的时候，就需要使用到动态的分配，在启动核函数的第三个参数上加上所需要的共享内存的大小，并在核函数中使用 extern __shared__ <dataType> var[]; 的方式获得动态分配的共享内存，如上面示例代码中的：

__global__ void dynamicReverse(int *d, int n)
{
  extern __shared__ int s[]; //**** extern 获取核函数分配的动态内存
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

Bank Conflict

当然，shared memory 也不是完全完美的，其还有可能引发一种新的冲突 ———— bank conflict: bank conflict 发生在 同一个warp中 不同的线程 访问同一个bank 不同的数据 上面四个条件缺一不可造成 bank conflict： 1. 同一个 warp，不同的 warp 一般不考虑 bank conflict，因为 warp 的分发是随机的我们是不可知的。 2. 不同的 thread，同一个 thread 一个 clock cycle 只能存取一个 data 所以这是一句废话，但是要保证在同一个 warp 里面所以不是一句废话 3. 同一个 bank，对于不同的 bank 我们管不着 4. 不同的 data，对于同一个 bank 中，如果 不同的 thread 都访存的是一个 data，cuda 提供了广播的方法，类似于 python 中的 numpy 广播，反而是对性能的提升而不是下降，只有访存不同的 data 的时候才会出现 bank conflict

Figure Explanation

上面的三种访存，左侧数字对应的同一个 warp 中 thread 的编号（在 CC 2.0 以上每个 warp 中都是 32个线程），右侧的黄色是 bank 的编号（在 CC 2.0 以上shared memory 中的 bank 的个数也是32），注意黄色的 bank 内部的橙色小矩形代表的是 data。

在 Left 和 Right 中，每个 thread 都访问的是不同的 bank，所以不满足 bank conflict 的第三个条件，没有发生 bank conflict。 在 Middle 中，相当于是 stride=16 的访问导致了偶数编号的 bank 同时被两个线程所访问，而且访问的是不同的 data，所以出现了 two-way bank conflict。

上面的左中右都没有出现 bank conflict，因为即使有两个线程访问同一个 bank，他们内部访问的都是同一个 data，不满足 bank conflict 的第四个条件，所以都没有发生 bank conflict。

Example

举矩阵乘法的例子：

$C=A*B$

global memory

__global__ 
void simpleMultiply(float *a, float* b, float *c, int N)
{
 int row = blockIdx.y * blockDim.y + threadIdx.y;
 int col = blockIdx.x * blockDim.x + threadIdx.x;
 float sum = 0.0f;
 for (int i = 0; i < TILE_DIM; i++) {
 sum += a[row*TILE_DIM+i] * b[i*N+col];
 }
 c[row*N+col] = sum;
}

partial shared memory

__global__ 
void coalescedMultiply(float *a, float* b, float *c, int N)
{
 __shared__ float aTile[TILE_DIM][TILE_DIM];
 int row = blockIdx.y * blockDim.y + threadIdx.y;
 int col = blockIdx.x * blockDim.x + threadIdx.x;
 float sum = 0.0f;
 aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
 __syncwarp(); //***
 for (int i = 0; i < TILE_DIM; i++) {
 sum += aTile[threadIdx.y][i]* b[i*N+col];
 }
 c[row*N+col] = sum;
}

注意到 * 的地方只需要 syncwarp() 不用 syncthread()，更加细粒度地实现了同步，因为我们这里从 global memory 搬运到 shared memory 的工作都是由一个 warp 上的所有线程来做，具体来讲是 threadIdx.y 相同的 32 个线程来搬运，而使用的时候这 32个线程又 仅仅使用他们32个线程所搬运的数据 —— aTile[threadIdx.y][i]**，所以只要一个 warp 中的所有线程搬运完就可以继续执行，而不用等待其他的 warp 中的其他线程。

all shared memory

__global__ 
void sharedABMultiply(float *a, float* b, float *c, int N)
{
 __shared__ float aTile[TILE_DIM][TILE_DIM],
 bTile[TILE_DIM][TILE_DIM];
 int row = blockIdx.y * blockDim.y + threadIdx.y;
 int col = blockIdx.x * blockDim.x + threadIdx.x;
 float sum = 0.0f;
 aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
 bTile[threadIdx.y][threadIdx.x] = b[threadIdx.y*N+col];
 __syncthreads(); // ***
 for (int i = 0; i < TILE_DIM; i++) {
 sum += aTile[threadIdx.y][i]* bTile[i][threadIdx.x];
 }
 c[row*N+col] = sum;
}

注意到当将所有的数据(A和B上的)所有要用的都搬到 shared memory 的时候我们又使用回了 __syncthread()，因为这个时候我们每个 warp 中的线程不只用自家搬运的数据了，我们还需要等待如 bTile[i][threadIdx.x] 这样的其他数据，所以我们一定要等所有的 thread 都搬移完成才敢继续进行。

Summary

确实使用 shared memory 会逐步带来提升，具体如下：

$C=A^TA$

global memory

__global__ 
void simpleMultiply(float *a, float *c, int M)
{
 int row = blockIdx.y * blockDim.y + threadIdx.y;
 int col = blockIdx.x * blockDim.x + threadIdx.x;
 float sum = 0.0f;
 for (int i = 0; i < TILE_DIM; i++) {
 sum += a[row*TILE_DIM+i] * a[col*TILE_DIM+i];
 }
 c[row*M+col] = sum;
}

all shared memory

__global__ 
void coalescedMultiply(float *a, float *c, int M)
{
 __shared__ float aTile[TILE_DIM][TILE_DIM];
 __shared__ float transposedTile[TILE_DIM][TILE_DIM];
 int row = blockIdx.y * blockDim.y + threadIdx.y;
 int col = blockIdx.x * blockDim.x + threadIdx.x;
 float sum = 0.0f;
 aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
 transposedTile[threadIdx.x][threadIdx.y] =
 a[(blockIdx.x*blockDim.x + threadIdx.y)*TILE_DIM +
 threadIdx.x];
 __syncthreads();
 for (int i = 0; i < TILE_DIM; i++) {
 sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x];
 }
 c[row*M+col] = sum;
}

但是测试发现性能并没有预期的提升，是因为出现了 bank conflict： 分析 sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x] 的部分，每个 thread 自己在 transposedTile 做 TILE_DIM-stride 遍历，不同的 thread 不会碰在一起，所以是前面出了问题。 在拷贝的时候，同一个 warp 中的 threadIdx.y 是相同的，而在TransposedTile 拷贝的过程中出现了同一个 warp 中所有的 thread 都在同一列上复制，而我们的 TransposedTile 的 col=TILE_DIM=32，刚好就是每个 bank 管一个 col，那所有的 col 都在访问同一个 bank 的不同数据，造成了最严重的 32-way bank conflict !!! 所以性能必然下降，解决方法也很简单。

图示：

all shared memory without bank conflict

将 TransposedTile 的声明改为：

 __shared__ float transposedTile[TILE_DIM][TILE_DIM+1];

即可，这样一来则变成了:

[0][threadIdx.y] -> 0 -> bank 0
[1][threadIdx.y] -> 33 -> bank 1
[2][threadIdx.y] -> 66 -> bank 2
......
[31][threadIdx.y] -> 31*33 -> bank 31

完全利用到了所有的 bank，有效解决了 bank conflict。

这种方法叫做 Memory Padding

图示：

Summary

Configuring amount of Shared Memory

实际上 shared Memory 的大小是有限的，shared Memory 和 L1 Cache 共用 on-chip 的 64KB，但是也是 programable 的，具体是在 host 端使用 API函数 指定：

__host__ 
cudaError_t cudaFuncSetCacheConfig (const void *func, cudaFuncCache cacheConfig)

Effect: Sets the preferred cache configuration for **a device function**.
Params:
  1. const void *func: 必须是由 __global__ 声明的核函数
  2. cudaFuncCache cacheConfig:
    1. cudaFuncCachePreferNone: no preference for shared memory or L1 (default)
    2. cudaFuncCachePreferShared: prefer larger shared memory and smaller L1 cache
    3. cudaFuncCachePreferL1: prefer larger L1 cache and smaller shared memory
    4. cudaFuncCachePreferEqual: prefer equal size L1 cache and shared memory

__host__
cudaError_t cudaDeviceSetCacheConfig (cudaFuncCache cacheConfig)
Effect: Sets the preferred cache configuration for the **current device**.
上面的是针对某个指定的核函数 func 进行cache的分配
这个就是直接对当前的 Device 直接设置，只要在 Device 上面运行的核函数都有 cache 的 Prefer

注意： 官方文档也说了：

This is only a preference. The runtime will use the requested configuration if possible, but it is free to choose a different configuration if required to execute func.

设置的config只是prefer而已，具体的是否按照这个config去exec还要看runtime的时候具体情况具体分析。

P.S. 但是在 Maxwell 之后 L1 Cache 被舍弃了，所以 on-chip 的64KB 完全属于 shared memory，就没有 prefer 的说法了。

Reference

1. CUDA 共享内存 bank conflict

2. CUDA_C_Programming_Guide

3. CUDA_C_Best_Practices_Guide