CUDA_共享内存、访存机制、访问优化

---

---

共享内存简介

共享内存时受用户控制的一级缓存，共享存储器为片内高速存储器，是一块能够被同一block中的全部线程访问的可读写存储器。访问共享存储器的速度几乎和访问寄存器同样快(相对而言，不是十分严谨的说法，真实状况是，共享内存的延时极低，大约1.5T/s的带宽，远高于全局内存的190G/s,此速度是寄存器的1/10)，是实现线程间通讯的延迟最小的方法。共享存储器能够用于实现多种功能，若是用于保存共用的计数器或者block的公用结果。

计算能力1.0、1.一、1.二、1.3硬件中，每一个SM的共享存储器的大小为16KByte，被组织为16个bank，对共享存储器的动态与静态分配与初始化

int main(int argc, char** argv) 
{
    // ...
    testKernel<<<1, 10, mem_size >>>(d_idata, d_odata);
    // ...
    CUT_EXIT(argc, argv);
}

__global__ void testKernel(float* g_idata, float* g_odata)
{
    // extern声明，大小由主机端程序决定。动态声明
    extern __shared__ float sdata_dynamic[];

    // 静态声明
    __shared__ int sdata_static[16];

    // 注意shared memory不能再定义时初始化
    sdata_static[tid] = 0;
}

注意，将共享存储器中的变量声明为外部数据时，例如

extern __shared__ float shared[];

数组的大小将在kernel启动时肯定，经过其执行参数肯定。经过这种方式定义的全部变量都开始于相同的地址，所以数组中的变量的布局必须经过偏移量显示管理。例如，若是但愿在动态分配的共享存储器得到与如下代码对应的内容：

short array0[128];
float array1[64];
int array2[256];

应该按照以下的方式对应定义：

extern __shared__ char array[];
// __device__ or __global__ function
__device__ void func()
{
    short* array0 = (short*)array;
    float* array1 = (float*)&array0[128];
    int* array2 = (int*)&array1[64];
}

共享内存架构

共享内存时基于存储器切换的架构（bank-switched architecture）.为了可以在并行访问时得到高带宽，共享存储器被划分为大小相等，不能被同时访问的存储器模块，称为bank。因为不一样的存储器模块能够互不干扰的同时工做，所以对位于n个bank上的n个地址的访问可以同时进行，此时有效带宽就是只有一个bank的n倍。

若是half-warp请求访问的多个地址位于同一个bank中，就会出现bank conflict。因为存储器模块在一个时刻没法响应多个请求，所以这些请求就必须被串行的完成。硬件会将形成bank conflict的一组访存请求划分为几回不存在conflict的独立请求，此时的有效带宽会下降与拆分获得的不存在conflict的请求个数相同的倍数。例外状况：一个half-warp中的全部线程都请求访问同一个地址时，会产生一次广播，此时反而只须要一次就能够响应全部线程的请求。

bank的组织方式是：每一个bank的宽度固定为32bit，相邻的32bit字被组织在相邻的bank中，每一个bank在每一个时钟周期能够提供32bit的带宽。

在费米架构的设备上有32个存储体，而在G200与G80的硬件上只有16个存储体。每一个存储体能够存4个字节大小的数据，足以用来存储一个单精度浮点型数据，或者一个标准的32位的整型数。开普勒架构的设备还引入了64位宽的存储体，使双精度的数据无需在跨越两个存储体。不管有多少线程发起操做，每一个存储体每一个周期只执行一次操做。

若是线程束中的每一个线程访问一个存储体，那么全部线程的操做均可以在一个周期内同时执行。此时无须顺序地访问，由于每一个线程访问的存储体在共享内存中都是独立的，互不影响。实际上，在每一个存储体与线程之间有一个交叉开关将它们链接，这在字的交换中颇有用。

此外，当线程束中的全部线程同时访问相同地址的存储体时，使用共享内存会有很大帮助，同常量内存同样，当全部线程访问同一地址的存储单元时，会触发一个广播机制到线程束中的每一个线程中。一般0号线程会写一个值而后与线程束中的其余线程进行通讯。

共享存储访问优化

在访问共享存储器的时候，须要着重关注如何减小bank conflict.产生bank conflict会形成序列化访问，严重下降有效带宽。

对于计算能力1.x设备，每一个warp大小都是32个线程，而一个SM中的shared memory被划分为16个bank（0-15）。一个warp中的线程对共享存储器的访问请求会被划分为2个half-warp的访问请求，只有处于同一half-warp内的线程才可能发生bank conflict，而一个warp中位于前half-warp的线程与位于后half-warp的线程间则不会发生bank conflict。

没有bank conflic的共享存储器访问示例（线程从数组读取32bit字场景）：

产生bank conflict的共享存储器访问示例（线程从数组读取32bit字场景）：

若是每一个线程访问的数据大小不是32bit时，也会产生bank conflict。例如如下对char数组的访问会形成4way bank conflict：

__shared__ char shared[32];
char data = shared[BaseIndex + tid];

此时，shared[0]、shared[1]、shared[2]、shared[3]属于同一个bank。对一样的数组，按照下面的形式进行访问，则能够避免bank conflict问题：

char data = shared[BaseIndex + 4* tid];

---

对于一个结构体赋值会被编译为几回访存请求，例如：

__shared__ struct type shared[32];
struct type data = shared[BaseIndex + tid];

假如type的类型有以下几种：

// type1
struct type {
	float x, y, z;
};

// type2
struct type {
	float x, y;
};

// type3
struct type {
	float x;
	char c;
};

若是type定义为type1，那么type的访问会被编译为三次独立的存储器访问，每一个结构体的同一成员之间有3个32bit字的间隔，因此不存在bank conflict。（没有bank conflic的共享存储器访问示例中场景c）

若是type定义为type2，那么type的访问会被编译为两个独立的存储器访问，每一个结构体成员都有2个32bit字的间隔，线程ID相隔8的线程间就会发生bank conflict。（产生bank conflict的共享存储器访问示例中场景b）

若是type定义为type3，那么type的访问会被编译为两个独立的存储器访问，每一个结构体成员都是经过5byte的间隔来访问，因此总会存在bank conflict。

---

shared memory访存机制

shared memory采用了广播机制，在响应一个对同一个地址的读请求时，一个32bit能够被读取的同时会广播给不一样的线程。当half-warp有多个线程读取同一32bit字地址中的数据时，能够减小bank conflict的数量。而若是half-warp中的线程全都读取同一地址中的数据时，则彻底不会发生bank conflict。不过，若是half-warp内有多个线程要对同一地址进行写操做，此时则会产生不肯定的结果，发生这种状况时应该使用对shared memory 的原子操做。

对不一样地址的访存请求，会被分为若干个处理步，每两个执行单元周期完成一步，每步都只处理一个conflict-free的访存请求的子集，知道half-warp的全部线程请求均完成。在每一步中都会按照如下规则构建子集：

（1）从还没有访问的地址所指向的字中，选出一个做为广播字；

（2）继续选取访问其余bank，而且不存在bank conflict的线程，再与上一步中广播字对应的线程一块儿构建一个子集。在每一个周期中，选择哪一个字做为广播字，以及选择哪些与其余bank对应的线程，都是不肯定的。

参考：

《高性能运算之CUDA》

《CUDA并行程序设计 GPU编程指南》

《GPU高性能编程 CUDA实战》

《CUDA专家手册 GPU编程权威指南》