gpu显存(全局内存)在使用时数据对齐的问题
来源:互联网 发布:60级魔兽世界数据库 编辑:程序博客网 时间:2024/06/05 22:46
全局存储器,即普通的显存,整个网格中的任意线程都能读写全局存储器的任意位置。
存取延时为400-600 clock cycles 非常容易成为性能瓶颈。
访问显存时,读取和存储必须对齐,宽度为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个线程的访问的地址在同一块区域(指硬件上可以一起传送宽度)时,并且没有冲突产生时,则这块区域的数据可以被线程同时,提升了访存的效率。
- gpu显存(全局内存)在使用时数据对齐的问题
- tensorflow使用GPU训练时的显存占用问题
- tensorflow使用GPU训练时的显存占用问题
- tensorflow使用GPU训练时的显存占用问题
- 关于使用PyTorch设置多线程(threads)进行数据读取而导致GPU显存始终不释放的问题
- (转)tensorflow中使用指定的GPU及GPU显存
- tensorflow中使用指定的GPU及GPU显存
- TensorFlow使用GPU训练网络时多块显卡的显存使用问题
- Keras限制GPU显存使用
- caffe 使用LMDB数据库训练时 显存爆炸的问题
- 显卡显存和内存显存的区别
- 显卡显存和内存显存的区别
- 内存中的数据对齐问题
- gpu显存与模型文件的关系
- 内存对齐的问题
- 内存对齐的问题
- 内存对齐的问题
- 【Ubuntu-Tensorflow】程序结束掉GPU显存没有释放的问题
- VPN服务提供商为自由
- MFC_DLL(二)
- GetWindowTextA
- 计算机基础--网络攻击基础
- HDU1800 Flying to the Mars 【贪心】
- gpu显存(全局内存)在使用时数据对齐的问题
- JAVA线程池例子
- VPN服务和隐私级别
- 可怜的蒟蒻(找规律+欧几里得)
- memcached+spring集成
- VPN服务与隐私
- erlang escript使用
- 【Linux】线程同步之信号量同步
- Ubuntu12.04 安装ADB调试环境