GPU memory 结构
来源:互联网 发布:矩阵奇异值分解 matlab 编辑:程序博客网 时间:2024/05/16 17:59
本节主要讲述GPU的memory架构。优化基于GPU device的kernel程序时,我们需要了解很多GPU的memory知识,比如内存合并,bank conflit(冲突)等等,这样才能针对具体算法做一些优化工作。
1、GPU总线寻址介绍
假定X是一个指向整数(32位整数)数组的指针,数组的首地址为0x00001232。一个线程要访问元素X[0],
int tmp = X[0];
假定memory总线宽度为256位(HD5870就是如此,即为32字节),因为基于字节地址的总线要访问memeory,必须和总线宽度对齐,也就是说按必须32字节对齐来访问memory,比如访问0x00000000,0x00000020,0x00000040,…等,所以我们要得到地址0x00001232中的数据,比如访问地址0x00001220,这时,它会同时得到0x00001220到 0x0000123F 的所有数据。因为我们只是取的一个32位整数,所以有用的数据是4个字节,其它28的字节的数据都被浪费了,白白消耗了带宽。
2、合并内存访问
为了利用总线带宽,GPU通常把多个线程的内存访问尽量合并到较少的内存请求命令中去。
假定下面的OpenCL kernel代码:int tmp = X[get_global_id(0)];
数组X的首地址和前面例子一样,也是0x00001232,则前16个线程将访问地址:0x00001232 到 0x00001272。假设每个memory访问请求都单独发送的话,则有16个request,有用的数据只有64字节,浪费掉了448字节(16*28)。
假定多个线程访问32个字节以内的地址,它们的访问可以通过一个memory request完成,这样可以大大提高带宽利用率,在专业术语描述中这样的合并访问称作coalescing。
例如上面16个线程访问地址0x00001232 到 0x00001272,我们只需要3次memory requst。
在HD5870显卡中,一个wave中16个连续线程的内存访问会被合并,称作quarter-wavefront,是重要的硬件调度单位。
下面的图是HD5870中,使用memory访问合并以及没有使用合并的bandwidth比较:
下图是GTX285中的比较:
3、Global memory的bank以及channel访问冲突
我们知道内存由bank,channel组成,bank是实际存储数据的单元,一个mc可以连接多个channel,形成单mc,多channel的连接方式。在物理上,不同bank的数据可以同时访问,相同的bank的数据则必须串行访问,channel也是同样的道理。但由于合并访问的缘故,对于global memory来说,bank conflit影响要小很多,除非是非合并问,不同线程访问同一个bank。理想情况下,我们应该做到不同的workgroup访问的不同的bank,同一个group内,最好用合并操作。
下面我简单的画一个图,不知道是否准确,仅供参考:
在HD5870中,memory地址的低8位表示一个bank中的数据,接下来的3位表示channel(共8个channel),bank位的多少依赖于显存中bank的多少。
4、local memory的bank conflit
bank访问冲突对local memory操作有更大的影响(相比于global memory),连续的local memory访问地址,应该映射到不同的bank上,
在AMD显卡中,一个产生bank访问冲突wave将会等待所有的local memory访问完成,硬件不能通过切换到另一个wave来隐藏local memory访问时延。所以对local memory访问的优化就很重要。HD5870显卡中,每个cu(simd)有32bank,每个bank 1k,按4字节对齐访问。如果没有bank conflit,每个bank能够没有延时的返回一个数据,下面的图就是这种情况。
如果多个memory访问对应到一个bank上,则conflits的数量决定时延的大小。下面的访问方式将会有3倍的时延。
但是,如果所有访问都映射到一个bank上,则系统会广播数据访问,不会产生额外时延。
原文作者:迈克老狼
- GPU memory 结构
- memory out (GPU)
- Beware of GPU memory bandwidth!!
- gpu/cuda-03-cuda memory
- Get GPU memory usage progamatically
- GPU硬件结构--CUDA
- GPU的硬件结构
- 转载:NVIDIA GPU结构
- [转]GPU硬件结构
- 一段代码搞懂 gpu memory
- OGRE - Memory , nedmalloc结构分析
- 9. CUDA shared memory使用------GPU的革命
- (CUDA 编程9).CUDA shared memory使用------GPU的革命
- CUDA6 Unified Memory,移植程序到GPU上再也不蛋碎
- How much GPU memory do iPhones and iPads have?
- Android性能优化大全(Code Memory CPU GPU UI)
- GPU 共享内存bank冲突(shared memory bank conflicts)
- GPU Memory exhausted, but no process listed by nvidia-smi
- [SCU 4500] 神舟的宝藏 (数位DP)
- yii2缓存的介绍和使用
- Spark Sql,Dataframe和数据集指南
- 你真的会二分查找吗?
- 13_JSP HTTP 状态码
- GPU memory 结构
- Spark-ML-01-小试spark分析离线商品信息
- APP开发实战7-APP的尺寸标注工作
- Spring实例化Bean的三种方式及Bean的类型(activiti实战)
- uva 361 - Cops and Robbers(凸包)
- 剑指offer之面试题11数值的整数次方
- 15_JSP 过滤器
- js检测输入的是否全部为数字
- COM开发-IDispatch