zoukankan      html  css  js  c++  java
  • AMD OpenCL大学教程(9)

        本节主要讲述GPU的memory架构。优化基于GPU device的kernel程序时,我们需要了解很多GPU的memory知识,比如内存合并,bank conflit(冲突)等等,这样才能针对具体算法做一些优化工作。

    1、GPU总线寻址介绍

    image

       假定X是一个指向整数(32位整数)数组的指针,数组的首地址为0x00001232。一个线程要访问元素X[0],

       int tmp = X[0];

    image  

        假定memory总线宽度为256位(HD5870就是如此,即为32字节),因为基于字节地址的总线要访问memory,必须和总线宽度对齐,也就是说按必须32字节对齐来访问memory,比如访问0x00000000,0x00000020,0x00000040,…等,所以我们要得到地址0x00001232中的数据,比如访问地址0x00001220,这时,它会同时得到0x00001220到 0x0000123F 的所有数据。因为我们只是取的一个32位整数,所以有用的数据是4个字节,其它28的字节的数据都被浪费了,白白消耗了带宽。

        image

    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。

    image

       例如上面16个线程访问地址0x00001232 到 0x00001272,我们只需要3次memory requst。

       在HD5870显卡中,一个wave中16个连续线程的内存访问会被合并,称作quarter-wavefront,是重要的硬件调度单位。

       下面的图是HD5870中,使用memory访问合并以及没有使用合并的bandwidth比较:

    image

       下图是GTX285中的比较:

    image

    3、Global memory的bank以及channel访问冲突

       我们知道内存由bank,channel组成,bank是实际存储数据的单元,一个mc可以连接多个channel,形成单mc,多channel的连接方式。在物理上,不同bank的数据可以同时访问,相同的bank的数据则必须串行访问,channel也是同样的道理。但由于合并访问的缘故,对于global memory来说,bank conflit影响要小很多,除非是非合访问,不同线程访问同一个bank。理想情况下,我们应该做到不同的workgroup访问的不同的bank,同一个group内,最好用合并操作。

       

    4、local memory的bank conflit

       bank访问冲突对local memory操作有更大的影响(相比于global memory),连续的local memory访问地址,应该映射到不同的bank上,

    image

         在AMD显卡中,一个产生bank访问冲突wavefront将会等待所有的local memory访问完成,硬件不能通过切换到另一个wave来隐藏local memory访问时延。所以对local memory访问的优化就很重要。HD5870显卡中,每个CU(simd)有32bank,每个bank 1k,按4字节对齐访问。如果没有bank conflit,每个bank能够没有延时的返回一个数据,下面的图就是这种情况。

    image

       如果多个memory访问对应到一个bank上,则conflits的数量决定时延的大小。下面的访问方式将会有3倍的时延。

    image

      但是,如果所有访问都映射到一个bank上,则系统会广播数据访问,不会产生额外时延。

    image

  • 相关阅读:
    用人之道(一) 如何组建软件开发队伍[转]
    用人之道(二) 何管理软件开发团队[转]
    2005年度世界500强公司名单[转]
    人类的15个欲望与游戏设计[转&收藏]
    Flash读取Cookie[转]
    高效程序员应该养成的七个习惯
    六度隔离学说,1967年,哈佛大学,Stanley Milgram
    做技术,切不可沉湎于技术[转&收藏]
    庆祝VSX团队成立,加入VSX团队申请帖
    如何把菜单添加到另外一个VSPackage的菜单里?
  • 原文地址:https://www.cnblogs.com/mikewolf2002/p/2333901.html
Copyright © 2011-2022 走看看