CUDA学习(二)矩阵转置及优化(合并访问、共享内存、bank conflict)
发布网友
发布时间：2024-10-16 20:28
共1个回答
热心网友
时间：2天前
首先介绍下如何计算cuda核函数的执行时间，为接下来的性能测试做准备。
CUDA内存组织
后续的优化过程涉及到cuda的内存使用，此处先作介绍。
cuda中的内存主要分为全局内存、共享内存、常量内存、寄存器、缓存。
其中，全局内存内存容量大，延迟高，是最常用的一种内存。全局内存的使用方式分为动态使用和静态使用。动态使用全局内存需要通过函数cudaMalloc()在主机端分配设备内存(没错，是在主机端分配而不是在设备端)，然后便可以直接在核函数中访问分配的内存。静态使用全局内存则只需要在声明变量时加上__device__修饰符。
共享内存的内存容量有限，但是其读写速度仅次于寄存器，比全局内存快很多。共享内存的使用方式也分为动态使和使用。静态使用共享内存只需要在声明变量时加上__shared__修饰符，而动态使用则需要在声明变量时加上extern __shared__修饰符，并且需要把变量声明为动态数组。
cuda实现矩阵转置
接下来使用cuda来实现矩阵转置，并探讨一些可以优化的地方。
下面是两段矩阵转置代码：
上述两个函数都能实现矩阵转置，但是性能却不同，分别对其在v100的全局内存上的性能进行测试：
可以发现，transpose2的执行时间明显比transpose1的执行时间短。
想要了解原因，首先得了解全局内存的访问模式，有合并访问和非合并访问两种方式。若一个线程束对全局内存的一次访问导致最少的数据传输，则称该访问为合并访问，否则为非合并访问。
举个例子(假设一次数据传输指的是将32字节的数据从全局内存通过32字节的缓存传输到SM，且已知从全局内存转移到缓存的首地址一定是一个最小粒度(此处为32字节)的整数倍(比如0~31、32~63、64~95这样传)，cudaMalloc分配的内存的首地址至少是256字节的整数倍)，下面这两个函数，add1是合并访问的，观察其第一次传输，第一个线程块中的线程束将访问x中的第0~31个元素，总共128字节的数据大小，这样4次传输就可以完成数据搬运，而128/32=4，说明合并度为100%。而add2则是非合并访问的，观察第一次传输，第一个线程块中的线程束将访问x中的第1~32个元素，若x的首地址为256字节，则线程束将作5次传输：256~287、288~319、320~351、352~383、384~415，其合并度为4/5=80%。
有了这个知识后，再去看矩阵转置的代码，可以发现，transpose1对矩阵A中的读取是合并的，而对矩阵B中的写入是非合并的(参见第一章多维和单维的地址转换)；而transpose2对矩阵A的读取是非合并的，而对矩阵B的写入是合并的。有人便有疑惑：为何都是一次合并访问和一次非合并访问，transpose2要更快呢？个人猜测是因为GPU架构对数据读取做了优化而未对数据写入进行优化。
共享内存使用
针对上述问题，我们尝试利用共享内存来改善全局内存的访问模式。
上面通过共享内存作为中介，使得全局内存的访问都变成了合并的，当然这个操作会增加一些共享内存的读写耗时，不过由于共享内存的访问速度比全局内存快很多，所以这样还是能提高效率的。
从上述测试结果可以看出，transpose3的执行时间比transpose1快，但还是慢于transpose2，说明优化后有一定效果，但还是有优化空间。
Bank conflicts
为了获取高带宽，共享内存在物理上被分为32个(内建变量warpsize的值)大小相同的内存bank，在每一个bank中，又可以对其中的内存地址从0开始编号。
对于bank宽度为4字节的架构(在开普勒架构中，每个bank的宽度为8字节；其他架构中bank则为4字节)，共享内存是按如下方式线性地址映射到内存bank：连续的128字节的内容分摊到32分bank的某一层中，每个bank负责4字节的内容。
当同一个线程束内的多个线程试图访问同一个bank的不同层的数据时，就会发生bank冲突。下图为bank冲突示意图。
一般情况下，可以通过改变共享内存数组的大小来消除或减轻bank冲突。
上述代码修改了共享内存的定义，就可以解决bank conflicts，这是因为改变共享内存数组大小之后，同一个线程束中的每个线程之间的地址跨度为33个字节而不是32个字节。例如，假设线程束中的第一个线程访问的是0000，则第二个线程将访问0132而不是0128，这样即可避免bank conflicts。
从上述结果可以发现，解决了共享内存的bank conflicts的transpose4的执行速度是几个矩阵转置函数中最快的。
