理论教育 向量化访存:优化全局内存性能与Canny算法表现

向量化访存:优化全局内存性能与Canny算法表现

时间:2023-06-23 理论教育 版权反馈
【摘要】:提高global memory的访存效率,可以有效减少对global memory的访存延迟,这对于Canny算子性能的提高有着至关重要的作用。由于数据的连续访存,可以通过向量化提高global memory的访存效率。自Fermi架构以来,GPU开始引入一级缓存与二级缓存,每次访问缓存,取来的最少数据量就是一个cache line,即128B。同时,Kepler架构的GPU的cache line是由4个32B的内存片断组成的。经过向量长度为4的向量后,每个warp访存3×128B,可以将warp内的访存请求合并成3次访存请求,从而提高globalmemory的访存效率。

向量化访存:优化全局内存性能与Canny算法表现

Canny算子各个部分都是连续数据访存密集型算法。提高global memory的访存效率,可以有效减少对global memory的访存延迟,这对于Canny算子性能的提高有着至关重要的作用。

由于数据的连续访存,可以通过向量化提高global memory的访存效率。一方面,可以充分利用数据的重用,比如在blur滤波算法中,相邻点的计算可以重用6个数据,极大地减少了访问global memory的访存次数,从而提高了global memory的访存效率。另一方面,可以充分利用Kepler架构的GPU的cache line。本文全部采用向量长度为4的向量。

cache line大小为128B,并且对齐在device memory地址的128B的内存片断上。自Fermi架构以来,GPU开始引入一级缓存(L1)与二级缓存(L2),每次访问缓存,取来的最少数据量就是一个cache line,即128B。对于NVIDIA Kepler架构的GPU而言,对于global memory的访问仅仅会被缓存在L2中(L1用于缓存local memory的访问)。同时,Kepler架构的GPU的cache line是由4个32B的内存片断组成的。如果warp的每个线程访问相邻的4B的数据,并且warp对齐在128B的内存片断上,那么可以实现global memory合并访存机制如第5章图5-2所示。如果warp与内存映射没有对齐,访存机制如第5章图5-3所示。因此,实现warp与内存映射的对齐和充分利用cache line可以有效提高global memory的访存效率。(www.daowen.com)

以灰度化算法为例,原来每个线程需要处理原图像的1个像素点(3个数据),每个warp处理32×3个数据,而每个warp访存的基本单位是128B,无法合并访存。经过向量长度为4的向量后,每个warp访存3×128B,可以将warp内的访存请求合并成3次访存请求,从而提高globalmemory的访存效率。

在NVIDIA GPU上,通过CUDA Runtime API分配内存(如cudaMalloc()函数),就已经保证内存至少对齐在256B。因此,在分配线程时选择合适的block size,在warp对齐于cache line时是有利于内存的访存的。例如,block size是warp size的整数倍。如果block size不是warp size的整数倍,可能会出现一块内存同时被第二个block、第三个block甚至后面的block访问的问题。

免责声明:以上内容源自网络,版权归原作者所有,如有侵犯您的原创版权请告知,我们将尽快删除相关内容。

我要反馈