CUDA学习之第五章 共享内存和常量内存（一）

5.1 CUDA共享内存概述

GPU有两种内存：

• 板载内存
• 片上内存

全局内存就是较大的板载内存，而共享内存则是较小的片上内存。共享内存常用的用途有三：

• 块内线程通信的通道
• 用于全局内存数据的可编程管理的缓存
• 高速暂存存储器，用于转换数据以优化全局内存访问模式

5.1.1 共享内存

共享内存（shared memory SMEM）由块SM中的块的所有线程共享。
共享内存和一级缓存相较于二级缓存和全局内存来说在物理上更接近SM，所以延迟要低。
共享内存会被分配一定数量给执行的块，所以一个块中的所有线程都共享，也因此核函数使用的共享内存越多，处于并发活跃状态的线程块就越少。
每个线程束会发出共享内存访问请求，理想情况下在一个事务中完成，最坏的则是32个不同事务顺序执行，如果多个线程访问共享内存中同一个字，则一个线程读取后，多播给其他线程。

5.1.2 共享内存分配

声明可以在核函数内也可以在全局，对应作用域也不同。修饰符使用__shared__。
还分静态和动态，静态可以声明一维、二维、三维的（高维有可能不行）。动态只能声明一维的。
静态：

__shared__ float tile[size_][size_x];

其中size_x和size_y是要固定的。

动态：

extern __shared__ int tile[];

需要在核函数调用时分配共享内存，对应在尖括号内第三个参数：

kernel<<<grid,block,isize*sizeof(int)>>>(...);

表示int数组大小isize。

5.1.3 共享内存存储体和访问模式

5.1.3.1 内存存储体

共享内存被分为32个同样大小的内存模型，被称为存储体，可以被同时访问。
如果通过线程束发布共享内存加载或存储操作，且在每个存储体上只访问不多于一个的内存地址，那么该操作可由一个内存事务来完成。否则需要多个内存事务，就会降低效率。

5.1.3.2 存储体冲突

当多个访问在同一个内存体时，会产生冲突（此处不是同一个地址，而是同一个存储体）。访问方法有三种经典模式。

• 并行：多个地址访问多个存储体
• 串行：多个地址访问同一个存储体
• 广播：单一地址读取单一存储体

并行是最好的，一个线程束访问的多个地址落在多个存储体中，意味着总会有一些地址在一个事务中一起完成。最好情况就是每个地址都位于一个单独的存储体，这样就无冲突访问。
串行是最差的，如果32线程都访问同一个存储体中的不同内存地址，那么意味着是普通请求的32倍时间。
广播虽然也可以并行，但是每一个存储体中只有一小部分字节被读取，带宽利用率差。

对于多个线程访问同一存储体，如果是同一地址，就会广播，如果是不同地址，就会冲突。

5.1.3.3 访问模式

共享内存存储体宽度规定了共享内存地址与存储体的对应关系。由计算能力决定。2.x能力对应的32位宽度，3.x能力对应64位宽度。

存储体索引=（字节地址÷ 8 ）%存储体数
其中存储体数为32，如果2.x需要÷4。

如果读取64位的设备中同一64位字的两个子字（两个32位）那么算读了一个地址，不会冲突。也就是说一次读64位，但是给两个线程一人32位，这是可以的。
另一种情况是64位设备，但是使用32位模式，这时会导致存储体出现一个64位字同时存两个32位字，这俩一起读也不冲突，比如说在32位模式下存储体0会同时保留0索引和32索引的两个32位字在同一个64位中，所以他们两个可以同时读而不冲突。

5.1.3.4 内存填充

如果同时读一存储体内的32个地址，会产生32向冲突，这个时候我们可以在每32个数据后填充一个数据，这样就会错开，使得所有冲突避免。
注意：写数据此方法无效，读数据时要重新计算索引，32位设备换到64位设备时原来的填充可能无效。

5.1.3.5 访问模式设置

Kelper设备可以支持4字节（32位）或8字节（64位）共享内存访问模式。默认是4字节的，可以采用API查询访问模式：

cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);

结果返回给pCofnig，可能是两种结果：cudaSharedMemBankSizeFourByte，cudaSharedMemBankEightByte。
通过以下函数可以配置存储体大小（就是只有一列还是有两列）：

cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig *pConfig);

除了上面两种pConfig外，还可以用cudaSharedMemBankSizeDefault。

不同核函数启动之间更改存储体配置可能需要一个隐式的设备同步点。

大（两列）存储体可能会使共享内存带宽变大，但是冲突增多。

5.1.4 配置共享内存量

每个SM都有64KB的片上内存，给共享内存和一级缓存用，配置二者大小有两种方法：按设备配置或按核函数配置。
方法见全局内存（一）。
核函数用的寄存器多久给一级缓存多分点，如果共享内存多就给他多分点。

按核函数配置使用如下函数：

cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCache cheConfig);

只需要调用一次这个函数，后续此核函数就会沿用配置。

5.1.5 同步

有两种方法同步：障碍和内存栅栏。
障碍是所有线程等待其余调用线程到达障碍点，内存栅栏是等待全部内存修改对其余调用线程可见时执行。

5.1.5.1 弱排序内存模型

为了让编译器好用，内存调用顺序未必按照程序的来，所以一个线程的写入顺序对其他线程可见时，未必按照其程序里写的顺序来执行。
为此必须同步来保证其执行顺序，保证核函数的行为正确。

5.1.5.2 显式障碍

调用如下函数使用障碍点：

void __syncthreads();

可以确保所有线程到达该点后再动。也可以确保这些线程访问的所有全局和共享内存对同一块中所有线程都可见。
必须保证等待条件唯一，否则可能会导致无限等待。
这个只能解决同一块内同步，所有的都同步只能用核函数结束来解决。

5.1.5.3 内存栅栏

块内栅栏：

void __threadfence_block();

内存栅栏是不会让线程同步的，对于块内线程来说这个函数没什么必要。

网格栅栏：

void __threadfence();

挂起调用线程，直到全局内存中所有写操作对网格内线程都可见。

系统栅栏：

void __threadfence_system();

跨系统（主机和设备）设置内存栅栏，保证任何内存的写操作其他线程或主机线程可见。

5.1.5.4 Volatile修饰符

volatile修饰符声明变量可以防止编译器优化，编译器优化可能会将数据暂时缓存在寄存器或本地内存中。使用volatile修饰的变量编译器会假定随时随地会被改动，所以会直接编译到全局内存读/写指令，忽略缓存。