cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑
来源:互联网 发布:基金知乎 编辑:程序博客网 时间:2024/05/16 08:18
关于共享内存(shared memory)和存储体(bank)的事实和疑惑
主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑。对于这点疑惑,网上都没有相关描述,
不管是国内还是国外的网上资料。貌似大家都是当作一个事实,一个公理,而没有对其仔细研究。还是我自己才学疏浅,不知道某些知识。
比如下面这篇讲解bank conflict的文章。
http://cuda-programming.blogspot.com/2013/02/bank-conflicts-in-shared-memory-in-cuda.html
我这里重点不在bank conflict,而是主要讨论shared memory和 memory bank的对应关系。
文中有这么一段描述:
First let’s assume your arrays are say for example of the type int (a 32-bit word). Your codesaves these ints into shared memory, acrossany half warp the Kth thread is saving to the Kth memory bank. Sofor example thread 0 of the first half warp will save to shared_a[0] which isin the first memory bank, thread 1 will save to shared_a[1], each half warp has16 threads these map to the 16 4byte banks. In the next half warp, the firstthread will now save its value into shared_a[16] which is in the first memory bankagain. So if you use a 4byte word such int, float etc, then this example willnot result in a bank conflict.
有一个数组,元素类型为整型,个数为256,开始这个数组存储在全局内存里面。现在我们一个线程块里有256个线程,我们想把这个数组拷贝到共享内存。因此每个线程负责拷贝一个元素。
shared_a[threadIdx.x] = global_a[threadIdx.x];想一下,这种访问是否会导致bank conflict呢?(看答案之前,先想想)
好的!
首先,我们假设你的数组元素是int类型的,占32位。你的代码把这些元素放进共享内存中,在任意一个half-warp,第k个线程刚好把元素放进第k个memory bank。
比如,第一个half warp中的线程0会放进shared_a[0],她刚好在第一个memory bank中,线程1把放进shared_a[1],每一个half warp有16个线程,刚好跟16个大小为4byte的bank对应。在下一个half warp中,第一个线程(线程0)会把值放进shared_a[16],她刚好也是在第一个memory bank中。所以在这个例子中,如果你使用4byte的字,比如int,float等,最后是不会产生bank confict的。
好了,回到我的讨论。
从上面描述,我们知道一些事实。
假如一个线程块有一块共享内存 int shared_a[256],该显卡设备的memory bank有16个。那么这块共享内存跟memory bank的对应关系是怎么样的?
例子说明一切,显然shared_a[0]在第1个bank中,shared_a[1]在第2个bank中,shared_a[15]在第16个bank中。
那么shared_a[16]呢?shared_a[17]呢?
根据文中的介绍,shared_a[16]在第1个bank中,shared_a[17]在第2个bank中。
规律是shared_a[index]在第(index%16+1)个bank中。
现在疑问来了,每一个bank的大小不是刚好为32位吗?(开普勒是64位)。
既然,shared_a[0]在第1个bank中,shared_a[0]已经是32位的了,那么shared_a[16]又是32位,放哪里?
shared_a[32]也是在第1个bank中,又放哪里?
一个bank怎么可以对应几个元素呢?
还是说bank只是缓存的地方,有其她地方存储,会自动切换的,类似缓存那样。
但是,貌似我没有找到任何资料有关这方面的解释。找了书,找了国内外的网上资料,都没有。
现在只好先记住这么一个事实了:shared_a[index]在第(index%16+1)个bank中。
本文作者:linger
本文链接:http://blog.csdn.net/lingerlanlan/article/details/32712749
- cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑
- 【并行计算-CUDA开发】关于共享内存(shared memory)和存储体(bank)的事实和疑惑
- GPU 共享内存bank冲突(shared memory bank conflicts)
- 共享内存(shared memory)和信号量(semaphore)简介(转)
- CUDA Programming Guide之shared memory的Bank Confict
- CUDA bank conflict in Shared Memory
- CUDA之Shared memory bank conflicts详解
- cuda shared memory 静态分配和动态分配
- 内存共享(Shared Memory)
- 共享内存: Shared Memory
- cuda中线程块共享存储(shared memory)加速较全局存储(global memory)之优势
- Android系统匿名共享内存Ashmem(Anonymous Shared Memory)简要介绍和学习计划
- Android系统匿名共享内存Ashmem(Anonymous Shared Memory)简要介绍和学习计划
- Android系统匿名共享内存Ashmem(Anonymous Shared Memory)简要介绍和学习计划
- Android系统匿名共享内存Ashmem(Anonymous Shared Memory)简要介绍和学习计划
- Android系统匿名共享内存Ashmem(Anonymous Shared Memory)简要介绍和学习计划
- Android系统匿名共享内存Ashmem(Anonymous Shared Memory)简要介绍和学习计划
- CUDA 共享内存 bank conflict
- 字符和字符串—Swift学习笔记(六)
- 1-2 Windows下启动函数(真正的入口函数) 之 寻找入口函数与_security_init_cookie
- struct linger 用法
- 啥发V大规范和那个年纪
- 稀释剂影响汽车低温烤漆
- cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑
- 取得当前日期的比较
- 行不啊
- 设计模式大总结(一)
- 语言集成查询(LINQ)
- 男子杀友获对方父亲谅解 法庭上连磕8个头
- 房间里的大象
- 开发那朵花
- Your Ride Is Here 你的飞碟在这儿