zoukankan      html  css  js  c++  java
  • gpu显存(全局内存)在使用时数据对齐的问题

    全局存储器,即普通的显存,整个网格中的随意线程都能读写全局存储器的任何位置。

    存取延时为400-600 clock cycles  很easy成为性能瓶颈。

    訪问显存时,读取和存储必须对齐,宽度为4Byte。假设没有正确的对齐,读写将被编译器拆分为多次操作,减少訪存性能。

    多个warp的读写操作假设可以满足合并訪问,则多次訪存操作会被合并成一次完毕。合并訪问的条件,1.0和1.1的设备要求较严格,1.2及更高能力的设备上放宽了合并訪问的条件。

    1.2及其更高能力的设备支持对8 bit、16 bit、32 bit、64 bit数据字的合并訪问,对应的段的大小为:32Byte 64Byte 128Byte,大于128Byte,分两次传输。

    在一次合并传输的数据中,不要求线程编号和訪问的数据字编号同样。

    当訪问128Byte数据时,假设地址没有对齐到128Byte时,在GT200会产生两次合并訪存。依据每一个区域的大小,分为两次合并訪存,如图所看到的32Byte和96Byte。

    全局存储器在使用的时候,主要注意的两个问题:

    1. 数据对齐的问题。一维数据使用cudaMalloc()开辟gpu全局内存空间,多维数据建议使用cudaMallocPitch()建立内存空间,以保证段对齐。cudaMallocPitch函数分配的内存中,数组的每一行的第一个元素的開始地址都保证是对齐的。由于每行有多少个数据是不确定的widthofx*sizeof(元素)不一定是256的倍数。故此,为保证数组的每一行的第一个元素的開始地址对齐,cudaMallocPitch在分配内存时,每行会多分配一些字节,以保证widthofx*sizeof(元素)+多分配的字节是256的倍数(对齐)。这样,y*widthofx*sizeof(元素)+x*sizeof(元素)来计算a[y][x]的地址就不对了。而应该是y*[widthofx*sizeof(元素)+多分配的字节]+x*sizeof(元素)。而函数中返回的pitch的值就是widthofx*sizeof(元素)+多分配的字节。

    2. 合并訪问。关键就是要理解,GPU是以half-warp(1.2及更高设备为warp)进行訪存时,即16个线程一起訪问存储器,到这16个线程的訪问的地址在同一块区域(指硬件上能够一起传送宽度)时,而且没有冲突产生时,则这块区域的数据能够被线程同一时候,提升了訪存的效率。

  • 相关阅读:
    jQueryMobile(二)
    (六)JavaScript之[Regular Expression]与[错误(try, catch, throw)]
    18-metaclass,是潘多拉魔盒还是阿拉丁神灯?
    17-强大的装饰器
    15-Python对象的比较、拷贝
    13-搭建积木:Python模块化
    12-面向对象(下):如何实现一个搜索引擎?
    11-面向对象(上):从生活中的类比说起
    10-简约不简单的匿名函数
    09-不可或缺的自定义函数
  • 原文地址:https://www.cnblogs.com/mfrbuaa/p/4357249.html
Copyright © 2011-2022 走看看