CUDA学习之第四章：全局内存（三）

最新推荐文章于 2024-02-25 22:08:37 发布

网安菜13

最新推荐文章于 2024-02-25 22:08:37 发布

阅读量272

点赞数 1

分类专栏： cuda学习文章标签： gpu算力

本文链接：https://blog.csdn.net/qq_49833809/article/details/131978777

版权

cuda学习专栏收录该内容

6 篇文章 0 订阅

订阅专栏

4.3 内存访问模式

4.3.1 对齐与合并访问

全局内存通过缓存来实现加载和存储（二级缓存），核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节内存事务来实现。
所有对全局内存的访问都会通过二级缓存，也有许多会通过一级缓存，二者都用到那么内存访问是由128字节的内存事务实现的，只用到了二级缓存，就是32字节内存事务。可以在编译时选择禁用或启用一级缓存。
一行一级缓存是128字节，它映射到设备内存中一个128字节的对齐段，如果线程束的每个线程请求4个字节，那么每次就会请求128字节，正好与缓存行和设备内存段的大小相契合。

因此优化应用程序时，需要注意内存访问的两个特性：对齐内存访问和合并内存访问。
当设备内存事务的第一个地址是用于事务服务缓存粒度的偶数倍时，会出现对齐内存访问，运行非对齐的加载会浪费带宽。
当一个线程束中全部的32个线程访问一个连续的内存块时，就会出现合并内存访问。

本小节主要含义总结为如下：

内存事务就相当于内存读取原子操作，是最小粒度（就是cache每次拿过来这么多）如果不对齐合并会多使用这种操作从而浪费效率。32个线程，每个一个字，就是一个线程访问4个字节，总共128个字节，地址增长了128。
对齐指的是以原子操作最小粒度的倍数为开始地址，否则相当于跨越两个原子操作。
合并操作相当于把访问都放一起使之连续，这样会由于原子操作每次拿出来的都是固定的，连续访问就不会浪费数据。

4.3.2 全局内存读取

SM中有三种缓存路径，取决于用了哪种设备内存：

一级和二级缓存
常量内存
只读缓存

第一种是默认路径，其他两种需要显式说明。禁用缓存的编译器选项是：

-Xptxas -dlcm=cg

启用的选项是：

-Xptxas -dlcm=ca

4.3.2.1 缓存加载

经过一级缓存的加载分为对齐/非对齐和合并/非合并。
理想状态为对齐合并的一次内存加载操作只用一个事务。
对于不连续但是都在一个事务内128字节内的加载操作，利用率也是100%。
对于不连续而且分散在超过了128字节的范围，利用率会降低（需要加载多个内存事务，且有数据没用上）。
对于所有线程束访问同一地址，利用率极低，只用了一个。
对于每个线程访问一个内存事务里的一个数，是最差的，只用一个还得召唤很多内存事务。

GPU一级缓存对于时间局部性没有优化，只有空间局部性有优化。

4.3.2.2 没有缓存的加载

不用一级缓存的情况，这是更细粒度的32字节的，对非对齐或非合并的内存访问带来更好的总线利用率。
书中例子和上一小节一样，说明了如果对齐合并很差的情况下，不用一级缓存会对效率有所提升。

4.3.2.3 非对齐读取的示例。

给出了一个非对齐导致性能损失的例子。
全局加载效率=请求的全局内存加载吞吐量 / 所需的全局内存加载吞吐量。
使用nvprof可以获取gld_efficiency的指标。
最后试验了禁用一级缓存，可以发现全局加载效率提高了，所以对于非对齐非合并的程序可以禁用一级缓存。

4.3.2.4 只读缓存

计算能力3.5及以上的GPU可以用只读缓存来支持使用全局内存加载来代替一级缓存。
只读缓存的加载粒度是32字节的。
两种方法指导内存通过只读缓存读取：

使用函数__ldg
在间接引用的指针上使用修饰符。

比如在核函数中将：a[i]=b[i]换成a[i]=__ldg(&b[i])，或者对于指针在声明时使用__restrict__修饰，比如核函数定义参数时就可以加上这个修饰：__global__ void copyKernel(int * __restrict__ out, const int * __restrict__ in)。

4.3.3 全局内存写入

内存的存储操作相对简单并且与读取完全不同，一级缓存是不能存储操作的，只能通过二级缓存在32字节段粒度执行。内存事务可以同时被分为一段、两段或四段（此处一段就是32字节段）。规则是能用四段就四段，否则两段，最差的就是多个一段。

4.3.4 结构体数组和数组结构体

结构体的数组（AoS）就是普通结构体组成的数组，数组的结构体（SoA）就是成员是数组的结构体。二者可以互相替代。（本处博主认为翻译不太好，AoS应该是Array of Struct，直接翻译成数组结构体有点没法理解后文将直接使用AoS和SoA）。
并行编程其实更偏向于SoA，当结构体成员是数组的时候空间局部性更好，反之则会差。
后文又测试了二者区别，此处不赘述。