文章目录
- 前言
- 共享内存的合理使用
- 数组归约计算
- 使用全局内存的计算
- 引入线程块中的同步函数
- 使用共享内存计算
- 静态共享内存
- 使用动态共享内存
- 性能比较
- 避免共享内存的bank冲突
- 使用共享内存进行数组转置
- bank概念
- 性能比较
- 总结
前言
cuda共享内存的合理使用。
共享内存的合理使用
共享内存的主要作用：
1.减少核函数中对全局内存的访问次数，实现高效的线程块内部的通信；
2.提高全局内存访问的合并度。
数组归约计算
数组归约，即计算数组中所有元素的和。
s
u
m
=
x
[
0
]
+
x
[
1
]
+
x
[
2
]
+
.
.
.
+
x
[
N
−
1
]
sum = x[0]+x[1]+x[2]+...+x[N-1]
sum=x[0]+x[1]+x[2]+...+x[N−1]
用一个C++函数实现：
real reduce(const real *x,const int N)
{
real sum = 0.0;
for(int n=0;n<N;++n)
{
sum += x[n];
}
return sum;
}
上面程序如果计算较长的数组，比如 1 0 8 10^8 108，数组初始化每个元素为1.23，这样如果采用双精度浮点运算结果为sum=123000000.110771，前面9位有效数字都正确，第10位开始有错误。单精度输出sum=33554432.000000，结果错误（因为单精度浮点数只有6，7位精确的有效数字）。
使用全局内存的计算
使用cuda程序要比c++程序计算稳健，计算效率要高。数组归约的计算要比数组相加计算复杂，数组相加只要定义和数组元素一样多的线程，每个线程进行元素相加即可。而数组归约是在一个数组上进行相加，最终得到一个数。
拥有多线程，只需要控制每个线程进行一次相加即可。同时为了提高性能，采用折半归约（前一半与后一半相加，重复，第一个元素即为数组归约的值）的方法计算。使用折半归约的算法，线程N的数量要是2的指数次方，了解了之后，会习惯性的将函数修改成下面的核函数
// real *d_x 为全局内存
void __global__ reduce(real *d_x,int N)
{
int n = blockDIm.x * blockIdx.x + threadIdx.x;
for(int offset=N/2;offset>0;offset/=2)
{
if(n<offset)
{
d_x[n] += d_x[n+offset];
}
}
}
但这是一个错误的函数：
（1）N的数量必须是2的指数次方，内存分配不灵活。
（2）因为单指令-多线程执行的原因，线程块和线程的执行不是顺序同步的，所以读取和写入可能会产生冲突，从而造成错误。例如把这个函数的前两次迭代写出：
// offset = N/2和N/4
if(n<N/2){d_x[n] += d_x[n+N/2]};
if(n<N/4){d_x[n] +=d_x[n+N/4]};
从上面的代码可知，n=N/4时，数组d_x[N/4]是会被写入数据的。但是当n=0时，第二条语句会读取d_x[N/4]的数据。因为线程之间的执行不是顺序的，所以可能在要读取d_x[N/4]的时候，d_x[N/4]里还没写入数据。
引入线程块中的同步函数
要保证核函数中语句的执行顺序与出现顺序一致，可以使用cuda里提供的同步函数__syncthreads()，只能在核函数里使用。该函数仅保证一个线程块里的所有线程在执行语句的时候保持顺序同步的，不同线程块的线程执行顺序还是不同步。数组归约可以使用该函数让每个线程块对其中的元素进行归约来实现。
// real *d_x, real *d_y为全局内存
void __global__ reduce_global(real *d_x, real *d_y)
{
const int tid = threadIdx.x;
// 定义一个寄存器指针变量来作为临时的缓存，指向每个线程块的起始地址
real *x = d_x + blockDim.x * blockIdx.x; // blockDim.x要是2的指数次方的整数
// 折半归约，用位运算代替/2，在和函数中更高效
for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
{
if (tid < offset)
{
x[tid] += x[tid + offset];
}
__syncthreads();
}
// 寄存器内存生命周期在和函数里，将值保存到全局内存里。
// 保证一个线程块中，仅执行一次。将寄存器缓存里的每个线程块里的第一个元素值赋值给全局内存，后面主机内存进行所有线程块首元素值相加即是数组归约值。
if (tid == 0) // 保证一个线程块中，仅执行一次
{
d_y[blockIdx.x] = x[0];
}
}
上面的代码，定义的real *x，能够在不同的线程块中指向全局内存中的不同地址，使得可以在不同线程块中对数组d_x中的不同部分归约。
每个线程块内独立的对其中的数据进行归约。同步函数在每个线程块执行之后使用。每个线程块之间的计算执行不是顺序的，但这不影响结果的正确性。因为在核函数中，每个线程块是独立的处理不同的数据，相互之间没有依赖。 所以，N不用是2的指数次的倍数了，只要线程块blockDim.x是2的指数次的倍数且能被N整除就行。
使用共享内存计算
全局内存不够高效，寄存器内存仅对单个线程可见，使用对整个线程块可见的共享内存来提高性能。前面说过在核函数中，要定义一个变量为共享内存，需要在定义语句加上限定符__shared__。
静态共享内存
共享内存用来定义一个长度为线程块大小的数组。
// 在核函数里定义
__shared__ real s_y[128];
在利用共享内存进行线程块之间的通信之前，都要进行使用__syncthreads()同步，以确保共享内存变量中的数据对线程块内的所有线程来说都准备就绪。并且为了方便的定义N的大小（前面其实都有限制），在初始化共享内存定义的数组时，有要用到的内存就初始化为1.23，没用到的就为0，这样就能在一个线程块里正确的进行数组归约计算。
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int n = bid * blockDim.x + tid;
// 定义共享内存
__shared__ real s_y[128];
// 初始化
s_y[tid] = (n < N) ? d_x[n] : 0.0;
// 使用之前要使用线程块同步函数
__syncthreads();
整个使用共享内存计算数组归约的函数如下：
void __global__ reduce_shared(real *d_x, real *d_y)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int n = bid * blockDim.x + tid;
// 定义共享内存
__shared__ real s_y[128];
// 初始化
s_y[tid] = (n < N) ? d_x[n] : 0.0;
// 使用之前要使用线程块同步函数
__syncthreads();
for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
{
if (tid < offset)
{
s_y[tid] += s_y[tid + offset];
}
__syncthreads();
}
// 共享内存生命周期在和函数里，将值保存到全局内存里
if (tid == 0)
{
d_y[bid] = s_y[0];
}
}
使用动态共享内存
同样，为了方便指定共享内存定义的数组长度，使用动态共享内存。
相比较于静态的共享内存，修改2个地方：
（1）在调用核函数的时候，要在<<<grid_size, block_size>>>里加上第三个参数，表示需要的动态共享内存的字节数大小，不写默认为0；
// smem 表示需要的动态共享内存的字节数大小
const int smem = sizeof(real) * block_size;
reduce_dynamic<<<grid_size, block_size, smem>>>(d_x, d_y);
（2）在定义的语句前面还要加上限定词extern。注意这里只能是数组的格式，不用指定长度。不能使用指针。
extern __shared__ real s_y[];
// 使用指针声明是错误的
extern __shared__ real *s_y;
整个程序：
void __global__ reduce_dynamic(real *d_x, real *d_y)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int n = bid * blockDim.x + tid;
extern __shared__ real s_y[];
s_y[tid] = (n < N) ? d_x[n] : 0.0;
__syncthreads();
for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
{
if (tid < offset)
{
s_y[tid] += s_y[tid + offset];
}
__syncthreads();
}
if (tid == 0)
{
d_y[bid] = s_y[0];
}
}
性能比较
在2080Ti上的测试，查看结果和性能：
1.使用单精度：
（1）全局内存：
（2）静态共享内存：
（3）使用动态共享内存
单精度的结果都比不用核函数的（33554432.000000）要好，但是也只有三位精确。在2080Ti上的性能测试也是差不多。使用其他的架构的GPU测试，使用共享内存的性能有提升。
2.使用双精度：
（1）全局内存：
（2）静态共享内存：
（3）使用动态共享内存
双精度的结果都要精确，但是性能测试用全局内存的要好。这点在后面还有其他优化。
避免共享内存的bank冲突
使用共享内存进行数组转置
定义一个共享内存数组，大小为32x32（对应上篇博客里使用全局内存）。相当于作为一个缓存，能够避免全局内存的非合并访问。
__global__ void transpose1(const real *A, real *B, const int N)
{
__shared__ real S[TILE_DIM][TILE_DIM];
int bx = blockIdx.x * TILE_DIM;
int by = blockIdx.y * TILE_DIM;
int nx1 = bx + threadIdx.x;
int ny1 = by + threadIdx.y;
if (nx1 < N && ny1 < N)
{
S[threadIdx.y][threadIdx.x] = A[ny1 * N + nx1];
}
__syncthreads();
int nx2 = bx + threadIdx.y;
int ny2 = by + threadIdx.x;
if (nx2 < N && ny2 < N)
{
B[nx2 * N + ny2] = S[threadIdx.x][threadIdx.y];
}
}
通过共享内存数组的缓存，注意写入和读取的数组数据顺序。第一个是按照x方向写入，即正常习惯顺序的，对应全局内存数组A也是顺序的，即是合并访问的。在生命周期结束前，要将数据赋值给全局内存数组B，在B中是按照y轴写入的，不是习惯的，但是对应的S数组也是这样子读取的，这两个是对应的，即访问是合并的。
bank概念
对于共享内存，有一个内存bank的概念。为了获得较高的内存带宽，共享内存在物理上被分为32（刚好是一个线程束中的线程数目）个同样宽度，能同时访问的内存bank。将32个bank从0-31编号，将所有bank编号为0的内存称为第一内存，编号为1的称为第二内存，以此类推，在开普勒架构上，每个bank的宽度为8个字节，其他架构，每个bank的宽度为4个字节。对于bank宽度为4字节的架构，共享内存数组是按如下方式线性的映射到内存bank的。将共享内存数组中连续的128个字节的内容分摊到32个bank的某一层中。第0-31个数组元素对应32个bank中的第一层，第32-63个元素对应第二层…
什么叫共享内存里Bank冲突？
当同一个线程束内的多个线程试图访问同一个bank中不同层的数据时，如果有线程束对同一个bank中的n层数据同时访问，将导致n层内存事务，就发生了n路bank冲突。这种冲突虽然不会对结果有影响，但对性能有影响(本来访问一次就行，现在访问n次)，应要尽量避免。
解决：
前面的通过共享内存进行数组转置的方法是存在bank冲突的，怎样解决。通常可以通过改变共享内存数组大小的方式来避免或者减轻bank冲突的影响。只需如下修改数组大小：
__shared__ real S[TILE_DIM][TILE_DIM + 1];
因为这样改变之后，同一个线程束中的32个线程将对应共享内存数组中跨度为33的数据。如果第一个线程访问的是第一个bank的第0层，第二个线程访问的是第二个bank的第二层（而不是第一个bank的第二层，如上面的图）。
性能比较
相对于上一篇全局内存的数组转置，可以对比发现，使用共享内存有bank冲突的性能比全局内存写入非合并，读取合并的要好，没有使用全局内存写入合并，读取非合并（默认会使用__ldg()读取）的好。但是使用共享内存无bank冲突的性能是最好的。
总结
cuda共享内存的合理使用
参考：
如博客内容有侵权行为，可及时联系删除！
CUDA 编程：基础与实践
https://docs.nvidia.com/cuda/
https://docs.nvidia.com/cuda/cuda-runtime-api
https://github.com/brucefan1983/CUDA-Programming
