GPU memory结构
2012-09-08 13:12
531 查看
本节主要讲述GPU的memory架构。优化基于GPU device的kernel程序时,我们需要了解很多GPU的memory知识,比如内存合并,bank conflit(冲突)等等,这样才能针对具体算法做一些优化工作。
1、GPU总线寻址介绍
![](http://articles.csdn.net/uploads/allimg/120820/145133NO-0.png)
假定X是一个指向整数(32位整数)数组的指针,数组的首地址为0x00001232。一个线程要访问元素X[0],
int tmp = X[0];
![](http://articles.csdn.net/uploads/allimg/120820/145133H39-1.png)
假定memory总线宽度为256位(HD5870就是如此,即为32字节), 因为基于字节地址的总线要访问memeory,必须和总线宽度对齐,也就是说按必须32字节对齐来访问memory,比如访问 0x00000000,0x00000020,0x00000040,…等,所以我们要得到地址0x00001232中的数据,比如访问地址 0x00001220,这时,它会同时得到0x00001220到 0x0000123F 的所有数据。因为我们只是取的一个32位整数,所以有用的数据是4个字节,其它28的字节的数据都被浪费了,白白消耗了带宽。
![](http://articles.csdn.net/uploads/allimg/120820/1451336351-2.png)
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。
![](http://articles.csdn.net/uploads/allimg/120820/1451331D7-3.png)
例如上面16个线程访问地址0x00001232 到 0x00001272,我们只需要3次memory requst。
在HD5870显卡中,一个wave中16个连续线程的内存访问会被合并,称作quarter-wavefront,是重要的硬件调度单位。
下面的图是HD5870中,使用memory访问合并以及没有使用合并的bandwidth比较:
![](http://articles.csdn.net/uploads/allimg/120820/1451336435-4.png)
下图是GTX285中的比较:
![](http://articles.csdn.net/uploads/allimg/120820/1451332129-5.png)
3、Global memory的bank以及channel访问冲突
我们知道内存由bank,channel组 成,bank是实际存储数据的单元,一个mc可以连接多个channel,形成单mc,多channel的连接方式。在物理上,不同bank的数据可以同 时访问,相同的bank的数据则必须串行访问,channel也是同样的道理。但由于合并访问的缘故,对于global memory来说,bank conflit影响要小很多,除非是非合并问,不同线程访问同一个bank。理想情况下,我们应该做到不同的workgroup访问的不同的bank,同 一个group内,最好用合并操作。
下面我简单的画一个图,不知道是否准确,仅供参考:
![](http://articles.csdn.net/uploads/allimg/120820/145133E28-6.png)
![](http://articles.csdn.net/uploads/allimg/120820/145133E08-7.png)
在HD5870中,memory地址的低8位表示一个bank中的数据,接下来的3位表示channel(共8个channel),bank位的多少依赖于显存中bank的多少。
4、local memory的bank conflit
bank访问冲突对local memory操作有更大的影响(相比于global memory),连续的local memory访问地址,应该映射到不同的bank上,
![](http://articles.csdn.net/uploads/allimg/120820/145133J92-8.png)
在AMD显卡中,一个产生bank访问冲突 wave将会等待所有的local memory访问完成,硬件不能通过切换到另一个wave来隐藏local memory访问时延。所以对local memory访问的优化就很重要。HD5870显卡中,每个cu(simd)有32bank,每个bank 1k,按4字节对齐访问。如果没有bank conflit,每个bank能够没有延时的返回一个数据,下面的图就是这种情况。
![](http://articles.csdn.net/uploads/allimg/120820/14513359C-9.png)
如果多个memory访问对应到一个bank上,则conflits的数量决定时延的大小。下面的访问方式将会有3倍的时延。
![](http://articles.csdn.net/uploads/allimg/120820/1451331B9-10.png)
但是,如果所有访问都映射到一个bank上,则系统会广播数据访问,不会产生额外时延。
![](http://articles.csdn.net/uploads/allimg/120820/145133M49-11.png)
原文作者:迈克老狼
1、GPU总线寻址介绍
![](http://articles.csdn.net/uploads/allimg/120820/145133NO-0.png)
假定X是一个指向整数(32位整数)数组的指针,数组的首地址为0x00001232。一个线程要访问元素X[0],
int tmp = X[0];
![](http://articles.csdn.net/uploads/allimg/120820/145133H39-1.png)
假定memory总线宽度为256位(HD5870就是如此,即为32字节), 因为基于字节地址的总线要访问memeory,必须和总线宽度对齐,也就是说按必须32字节对齐来访问memory,比如访问 0x00000000,0x00000020,0x00000040,…等,所以我们要得到地址0x00001232中的数据,比如访问地址 0x00001220,这时,它会同时得到0x00001220到 0x0000123F 的所有数据。因为我们只是取的一个32位整数,所以有用的数据是4个字节,其它28的字节的数据都被浪费了,白白消耗了带宽。
![](http://articles.csdn.net/uploads/allimg/120820/1451336351-2.png)
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。
![](http://articles.csdn.net/uploads/allimg/120820/1451331D7-3.png)
例如上面16个线程访问地址0x00001232 到 0x00001272,我们只需要3次memory requst。
在HD5870显卡中,一个wave中16个连续线程的内存访问会被合并,称作quarter-wavefront,是重要的硬件调度单位。
下面的图是HD5870中,使用memory访问合并以及没有使用合并的bandwidth比较:
![](http://articles.csdn.net/uploads/allimg/120820/1451336435-4.png)
下图是GTX285中的比较:
![](http://articles.csdn.net/uploads/allimg/120820/1451332129-5.png)
3、Global memory的bank以及channel访问冲突
我们知道内存由bank,channel组 成,bank是实际存储数据的单元,一个mc可以连接多个channel,形成单mc,多channel的连接方式。在物理上,不同bank的数据可以同 时访问,相同的bank的数据则必须串行访问,channel也是同样的道理。但由于合并访问的缘故,对于global memory来说,bank conflit影响要小很多,除非是非合并问,不同线程访问同一个bank。理想情况下,我们应该做到不同的workgroup访问的不同的bank,同 一个group内,最好用合并操作。
下面我简单的画一个图,不知道是否准确,仅供参考:
![](http://articles.csdn.net/uploads/allimg/120820/145133E28-6.png)
![](http://articles.csdn.net/uploads/allimg/120820/145133E08-7.png)
在HD5870中,memory地址的低8位表示一个bank中的数据,接下来的3位表示channel(共8个channel),bank位的多少依赖于显存中bank的多少。
4、local memory的bank conflit
bank访问冲突对local memory操作有更大的影响(相比于global memory),连续的local memory访问地址,应该映射到不同的bank上,
![](http://articles.csdn.net/uploads/allimg/120820/145133J92-8.png)
在AMD显卡中,一个产生bank访问冲突 wave将会等待所有的local memory访问完成,硬件不能通过切换到另一个wave来隐藏local memory访问时延。所以对local memory访问的优化就很重要。HD5870显卡中,每个cu(simd)有32bank,每个bank 1k,按4字节对齐访问。如果没有bank conflit,每个bank能够没有延时的返回一个数据,下面的图就是这种情况。
![](http://articles.csdn.net/uploads/allimg/120820/14513359C-9.png)
如果多个memory访问对应到一个bank上,则conflits的数量决定时延的大小。下面的访问方式将会有3倍的时延。
![](http://articles.csdn.net/uploads/allimg/120820/1451331B9-10.png)
但是,如果所有访问都映射到一个bank上,则系统会广播数据访问,不会产生额外时延。
![](http://articles.csdn.net/uploads/allimg/120820/145133M49-11.png)
原文作者:迈克老狼
相关文章推荐
- GPU memory 结构
- Java中的OutOfMemoryError和JVM内存结构
- GPU 共享内存bank冲突(shared memory bank conflicts)
- pyCUDA教程-系列学习(1):GPU结构、pyCUDA、numbapro安装及HelloGPU例子
- [GPU体系结构/GPU架构]深度解析AMD GPU中的Hierarchical Z以及Early-z被诸多限制的原因
- Java中的OutOfMemoryError的各种情况及解决和JVM内存结构
- Get GPU memory usage progamatically
- CUDA系列学习(三)GPU设计与结构QA & coding练习
- Java中的OutOfMemoryError和JVM内存结构
- Android平台美颜相机/Camera实时滤镜/视频编解码/影像后期/人脸技术探索——2.1 GPUImage结构简析
- GPU Memory Usage占满而GPU-Util却为0的调试
- 一段代码搞懂 gpu memory
- CUDA系列学习(三)GPU设计与结构QA & coding练习
- 深入理解java虚拟机-1 内存结构与OutOfMemory溢出异常
- chromium for android GPU进程结构分析
- Java中的OutOfMemoryError和JVM内存结构
- 内存数据表和内存索引通用结构(基于C++) --Memory Table Interface Embed In Your Application
- Android客户端性能工具3:GpuMemoryTracker分析
- OGRE - Memory , nedmalloc结构分析
- Nvidia gpu物理结构与编程模型简介