权威指南 第四章

概述这一章主要是从global memory角度来分析和优化算子的用transposition的例子分析maxbandwidth rate对性能的影响。重要概念内存层次全局内存位置位于DRAM上延迟和带宽都不如片上内存但是容量大成本低。固定内存是host侧专门分出的一个固定内存页用来与device侧进行数据的读写零拷贝内存位于host侧是host侧分出的固定内存页给device侧使用。device和后视图均可直接访问该内存有以下的优势·当设备内存不足时可利用主机内存·避免主机和设备间的显式数据传输提高PCIe传输率。但是由于GPU需要通过PCIe重复从主机上读取数据当有大量数据传输时使用这种方法速度较慢。统一内存寻址统一内存中创建了一个托管内存池内存池中已分配的空间可以用相同的内存地址即指针在CPU和GPU上进行访问。底层系统在统一内存空间中自动在主机和设备之间进行数据传输。相关函数如下cudaError tcudaMallocManaged(void**devPtr,size t size,unsignedintflags0);对齐与合并访问当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数倍时32字节的二级缓存或128字节的一级缓存就会出现对齐内存访问运行非对齐的加载会造成带宽浪费。当一个线程束中全部的32个线程访问一个连续的内存块时就会出现合并内存访问。AOS和SOA//AOSstructinnerStruct{floatx;floaty;}//SOAstructinnerArray{floatx[N];floaty[N];内存请求冲突对最佳性能来说被所有活跃的线程束并发访问的全局内存应该在分区中被均匀地划分。下图所示为一个简化的可视化模型它表示了使用笛卡尔坐标表示块ID时的分区冲突。在这个图中假设通过两个分区访问全局内存每个分区的宽度为256个字节并且使用一个大小为32×32的线程块启动核函数。如果每个数据块的宽度是128个字节那么需要使用两个分区为第0、1.2.3个线程块加载数据。但现实是只能使用一个分区为第0、1.2.3个线程块存储数据因此造成了分区冲突下图依旧是简化模型但这次使用对角坐标来表示块ID。在这种情况下需要使用两个分区为第0、1.2.3个线程块加载和存储数据。加载和存储请求在两个分区之间被均匀分配对于内存的分配和释放等等概念就是老生常谈的了此处不放了。transposition优化方法transposition就是对矩阵进行转置。转置的方法有多种下面分析一下不同的方法实现以及性能。naive transposition这种就直接跳过了直接进行转置分为按行转置和按列转置。unrolling 展开相邻的元素展开相邻的4个元素。这种方法下按列转置的性能较好ncu的memory性能报告这种按列转置的性能已经很不错了maxbandwidth rate已经接近80%。从书上看到其性能较好的原因是在缓存中执行了交叉读取。即使通过某一方式读入一级缓存中的数据没有都被这次访问使用到这些数据仍留在缓存中在以后的访问过程中可能发生缓存命中。而按行转置可能会发生bank conflict导致wrap中不同的线程访问同一块内存事务的地址发生排队等待。其实后续的几种方法也是尽可能避免bank conflict从而提高性能。diagonal 使用对角坐标这种方法主要是对blockIdx从笛卡尔坐标映射到对角坐标进行交叉访存尽量避免bank conflict。性能些许的提高了按行转置的性能但是相对于按列转置还是有差距。这种方法下的按列转置性能有所下降thin block unrolling 瘦块和展开很简单在unrolling基础上将blocksize从16,16改成8,32这种方法目前章节中性能超级好按行转置和按列转置的maxbandwidth rate达到了惊人的80%和91%没想到就是改了一个blocksize就能有这么大的性能提升统一内存方法这种方法优点在于把数据放在了CPU上避免了繁琐的在GPU上申请空间使得代码简洁化。但是缺点也很明显GPU需要多次从CPU上访存数据造成性能的下降。diagonalunrolling课后习题上的还没调试出来需要再熟悉一下GBD预计效果会有些许提升。此外还想到了diagonalthin blockunrolling方法真的很多就算知道理论还是需要验证一下实际的效果因为影响性能的因素真的很多。