当前位置:网站首页>CUDA_共享内存、访存机制、访问优化

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

2020-11-09 22:41:01 一介草民李八千



共享内存简介

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

计算能力1.0、1.1、1.2、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编程权威指南》

版权声明
本文为[一介草民李八千]所创,转载请带上原文链接,感谢
https://www.cnblogs.com/TonvyLeeBlogs/p/13951344.html